Add Queue support

Queues are responsible for scheduling operators and performing other
runtime related activities like for example tuning.

Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com>
Change-Id: I0366d9048470d277b8cbf59fa42f95c0ae57c5c9
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5487
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>

45 files changed, 1279 insertions, 141 deletions

diff --git a/Android.bp b/Android.bp
index f542b20104..78ac7f130f 100644
--- a/Android.bp
+++ b/Android.bp
@@ -52,6 +52,7 @@ cc_library_static {
     export_include_dirs: [".", "./include"],
     srcs: [
         "src/c/AclContext.cpp",
+        "src/c/AclQueue.cpp",
         "src/c/AclTensor.cpp",
         "src/c/AclTensorPack.cpp",
         "src/c/AclVersion.cpp",
@@ -300,7 +301,7 @@ cc_library_static {
         "src/core/cpu/kernels/CpuDepthwiseConvolutionNativeKernel.cpp",
         "src/core/cpu/kernels/CpuDequantizationKernel.cpp",
         "src/core/cpu/kernels/CpuDirectConvolutionKernel.cpp",
-        "src/core/cpu/kernels/CpuDirectConvolutionStageKernel.cpp",
+        "src/core/cpu/kernels/CpuDirectConvolutionOutputStageKernel.cpp",
         "src/core/cpu/kernels/CpuElementwiseKernel.cpp",
         "src/core/cpu/kernels/CpuElementwiseUnaryKernel.cpp",
         "src/core/cpu/kernels/CpuFillKernel.cpp",
@@ -315,16 +316,16 @@ cc_library_static {
         "src/core/cpu/kernels/CpuSoftmaxKernel.cpp",
         "src/core/cpu/kernels/CpuSubKernel.cpp",
         "src/core/cpu/kernels/CpuTransposeKernel.cpp",
-        "src/core/cpu/kernels/activation/NEON/fp16.cpp",
-        "src/core/cpu/kernels/activation/NEON/fp32.cpp",
-        "src/core/cpu/kernels/activation/NEON/qasymm8.cpp",
-        "src/core/cpu/kernels/activation/NEON/qasymm8_signed.cpp",
-        "src/core/cpu/kernels/activation/NEON/qsymm16.cpp",
-        "src/core/cpu/kernels/activation/SVE/fp16.cpp",
-        "src/core/cpu/kernels/activation/SVE/fp32.cpp",
-        "src/core/cpu/kernels/activation/SVE/qasymm8.cpp",
-        "src/core/cpu/kernels/activation/SVE/qasymm8_signed.cpp",
-        "src/core/cpu/kernels/activation/SVE/qsymm16.cpp",
+        "src/core/cpu/kernels/activation/neon/fp16.cpp",
+        "src/core/cpu/kernels/activation/neon/fp32.cpp",
+        "src/core/cpu/kernels/activation/neon/qasymm8.cpp",
+        "src/core/cpu/kernels/activation/neon/qasymm8_signed.cpp",
+        "src/core/cpu/kernels/activation/neon/qsymm16.cpp",
+        "src/core/cpu/kernels/activation/sve/fp16.cpp",
+        "src/core/cpu/kernels/activation/sve/fp32.cpp",
+        "src/core/cpu/kernels/activation/sve/qasymm8.cpp",
+        "src/core/cpu/kernels/activation/sve/qasymm8_signed.cpp",
+        "src/core/cpu/kernels/activation/sve/qsymm16.cpp",
         "src/core/cpu/kernels/add/neon/integer.cpp",
         "src/core/cpu/kernels/add/neon/qasymm8.cpp",
         "src/core/cpu/kernels/add/neon/qasymm8_signed.cpp",
@@ -390,8 +391,10 @@ cc_library_static {
         "src/core/utils/misc/MMappedFile.cpp",
         "src/core/utils/quantization/AsymmHelpers.cpp",
         "src/cpu/CpuContext.cpp",
+        "src/cpu/CpuQueue.cpp",
         "src/cpu/CpuTensor.cpp",
         "src/gpu/cl/ClContext.cpp",
+        "src/gpu/cl/ClQueue.cpp",
         "src/gpu/cl/ClTensor.cpp",
         "src/runtime/Allocator.cpp",
         "src/runtime/BlobLifetimeManager.cpp",
diff --git a/SConscript b/SConscript
index b09551fc82..63c2a483dc 100644
--- a/SConscript
+++ b/SConscript
@@ -188,11 +188,25 @@ runtime_files = Glob('src/runtime/*.cpp')
 runtime_files += Glob('src/runtime/CPP/ICPPSimpleFunction.cpp')
 runtime_files += Glob('src/runtime/CPP/functions/*.cpp')
 
-runtime_files += Glob('src/c/*.cpp')
-runtime_files += Glob('src/common/*.cpp')
-runtime_files += Glob('src/common/utils/*.cpp')
-runtime_files += Glob('src/cpu/*.cpp')
+# C API files
+c_api_files = ['src/c/AclContext.cpp',
+               'src/c/AclQueue.cpp',
+               'src/c/AclTensor.cpp',
+               'src/c/AclTensorPack.cpp',
+               'src/c/AclVersion.cpp',
+               ]
+if env['opencl']:
+    c_api_files += ['src/c/cl/AclOpenClExt.cpp']
+
+# Common backend files
+common_backend_files = ['src/common/utils/LegacySupport.cpp',
+                        'src/common/AllocatorWrapper.cpp',
+                        'src/common/ITensorV2.cpp',
+                        'src/common/TensorPack.cpp',
+                        ]
 
+core_files += common_backend_files
+runtime_files += c_api_files
 # CLHarrisCorners uses the Scheduler to run CPP kernels
 runtime_files += Glob('src/runtime/CPP/SingleThreadScheduler.cpp')
 
@@ -225,7 +239,6 @@ if env['opencl']:
     runtime_files += Glob('src/runtime/CL/gemm_auto_heuristics/*.cpp')
 
     runtime_files += Glob('src/gpu/cl/*.cpp')
-    runtime_files += Glob('src/c/cl/*.cpp')
 
     graph_files += Glob('src/graph/backends/CL/*.cpp')
 
@@ -278,8 +291,36 @@ if env['neon']:
     runtime_files += Glob('src/runtime/NEON/functions/*.cpp')
     runtime_files += Glob('src/runtime/NEON/functions/assembly/*.cpp')
 
-    core_files += Glob('src/core/cpu/*.cpp')
-    core_files += Glob('src/core/cpu/kernels/*.cpp')
+    cpu_kernel_hp_files = ['src/core/cpu/kernels/CpuActivationKernel.cpp',
+                           'src/core/cpu/kernels/CpuDepthwiseConvolutionNativeKernel.cpp',
+                           'src/core/cpu/kernels/CpuDirectConvolutionKernel.cpp',
+                           'src/core/cpu/kernels/CpuDirectConvolutionOutputStageKernel.cpp',
+                           'src/core/cpu/kernels/CpuPermuteKernel.cpp',
+                           'src/core/cpu/kernels/CpuPoolingAssemblyWrapperKernel.cpp',
+                           'src/core/cpu/kernels/CpuPoolingKernel.cpp',
+                           'src/core/cpu/kernels/CpuReshapeKernel.cpp',
+                          ]
+    cpu_kernel_files = ['src/core/cpu/kernels/CpuAddKernel.cpp',
+                        'src/core/cpu/kernels/CpuConcatenateBatchKernel.cpp',
+                        'src/core/cpu/kernels/CpuConcatenateDepthKernel.cpp',
+                        'src/core/cpu/kernels/CpuConcatenateHeightKernel.cpp',
+                        'src/core/cpu/kernels/CpuConcatenateWidthKernel.cpp',
+                        'src/core/cpu/kernels/CpuConvertFullyConnectedWeightsKernel.cpp',
+                        'src/core/cpu/kernels/CpuCopyKernel.cpp',
+                        'src/core/cpu/kernels/CpuDequantizationKernel.cpp',
+                        'src/core/cpu/kernels/CpuElementwiseKernel.cpp',
+                        'src/core/cpu/kernels/CpuElementwiseUnaryKernel.cpp',
+                        'src/core/cpu/kernels/CpuFillKernel.cpp',
+                        'src/core/cpu/kernels/CpuFloorKernel.cpp',
+                        'src/core/cpu/kernels/CpuPixelWiseMultiplicationKernel.cpp',
+                        'src/core/cpu/kernels/CpuQuantizationKernel.cpp',
+                        'src/core/cpu/kernels/CpuScaleKernel.cpp',
+                        'src/core/cpu/kernels/CpuSoftmaxKernel.cpp',
+                        'src/core/cpu/kernels/CpuSubKernel.cpp',
+                        'src/core/cpu/kernels/CpuTransposeKernel.cpp',
+                       ]
+    core_files += [cpu_kernel_hp_files, cpu_kernel_files]
+
     core_files += Glob('src/core/cpu/kernels/*/*.cpp')
     if any(i in env['data_type_support'] for i in ['all', 'fp16']):
         core_files += Glob('src/core/cpu/kernels/*/*/fp16.cpp')
@@ -293,12 +334,40 @@ if env['neon']:
         core_files += Glob('src/core/cpu/kernels/*/*/qsymm16.cpp')
     if any(i in env['data_type_support'] for i in ['all', 'integer']):
         core_files += Glob('src/core/cpu/kernels/*/*/integer.cpp')
-   
+
     if any(i in env['data_layout_support'] for i in ['all', 'nchw']):
         core_files += Glob('src/core/cpu/kernels/*/*/nchw/all.cpp')
 
-    runtime_files += Glob('src/runtime/cpu/*.cpp')
-    runtime_files += Glob('src/runtime/cpu/operators/*.cpp')
+    cpu_rt_files = ['src/cpu/CpuContext.cpp',
+                    'src/cpu/CpuQueue.cpp',
+                    'src/cpu/CpuTensor.cpp'
+                   ]
+    cpu_operator_hp_files = ['src/runtime/cpu/operators/CpuActivation.cpp',
+                             'src/runtime/cpu/operators/CpuDepthwiseConvolution.cpp',
+                             'src/runtime/cpu/operators/CpuDepthwiseConvolutionAssemblyDispatch.cpp',
+                             'src/runtime/cpu/operators/CpuDirectConvolution.cpp',
+                             'src/runtime/cpu/operators/CpuPermute.cpp',
+                             'src/runtime/cpu/operators/CpuPooling.cpp',
+                             'src/runtime/cpu/operators/CpuPoolingAssemblyDispatch.cpp',
+                            ]
+    cpu_operator_files = ['src/runtime/cpu/operators/CpuAdd.cpp',
+                          'src/runtime/cpu/operators/CpuConcatenate.cpp',
+                          'src/runtime/cpu/operators/CpuConvertFullyConnectedWeights.cpp',
+                          'src/runtime/cpu/operators/CpuCopy.cpp',
+                          'src/runtime/cpu/operators/CpuDequantization.cpp',
+                          'src/runtime/cpu/operators/CpuElementwise.cpp',
+                          'src/runtime/cpu/operators/CpuElementwiseUnary.cpp',
+                          'src/runtime/cpu/operators/CpuFill.cpp',
+                          'src/runtime/cpu/operators/CpuFloor.cpp',
+                          'src/runtime/cpu/operators/CpuPixelWiseMultiplication.cpp',
+                          'src/runtime/cpu/operators/CpuQuantization.cpp',
+                          'src/runtime/cpu/operators/CpuReshape.cpp',
+                          'src/runtime/cpu/operators/CpuScale.cpp',
+                          'src/runtime/cpu/operators/CpuSoftmax.cpp',
+                          'src/runtime/cpu/operators/CpuSub.cpp',
+                          'src/runtime/cpu/operators/CpuTranspose.cpp',
+                         ]
+    runtime_files += [ cpu_rt_files, cpu_operator_hp_files, cpu_operator_files ]
 
 bootcode_o = []
 if env['os'] == 'bare_metal':
diff --git a/arm_compute/Acl.hpp b/arm_compute/Acl.hpp
index 01f7179c2f..93ac2d8ed9 100644
--- a/arm_compute/Acl.hpp
+++ b/arm_compute/Acl.hpp
@@ -42,6 +42,7 @@ namespace acl
 {
 // Forward declarations
 class Context;
+class Queue;
 class Tensor;
 class TensorPack;
 
@@ -83,6 +84,7 @@ struct ObjectDeleter
     };
 
 OBJECT_DELETER(AclContext, AclDestroyContext)
+OBJECT_DELETER(AclQueue, AclDestroyQueue)
 OBJECT_DELETER(AclTensor, AclDestroyTensor)
 OBJECT_DELETER(AclTensorPack, AclDestroyTensorPack)
 
@@ -384,7 +386,7 @@ public:
         AclContext ctx;
         const auto st = detail::as_enum<StatusCode>(AclCreateContext(&ctx, detail::as_cenum<AclTarget>(target), &options.copts));
         reset(ctx);
-        report_status(st, "[Arm Compute Library] Failed to create context");
+        report_status(st, "[Compute Library] Failed to create context");
         if(status)
         {
             *status = st;
@@ -392,6 +394,92 @@ public:
     }
 };
 
+/**< Available tuning modes */
+enum class TuningMode
+{
+    Rapid      = AclRapid,
+    Normal     = AclNormal,
+    Exhaustive = AclExhaustive
+};
+
+/** Queue class
+ *
+ * Queue is responsible for the execution related aspects, with main responsibilities those of
+ * scheduling and tuning operators.
+ *
+ * Multiple queues can be created from the same context, and the same operator can be scheduled on each concurrently.
+ *
+ * @note An operator might depend on the maximum possible compute units that are provided in the context,
+ *       thus in cases where the number of the scheduling units of the queue are greater might lead to errors.
+ */
+class Queue : public detail::ObjectBase<AclQueue_>
+{
+public:
+    /**< Queue options */
+    struct Options
+    {
+        /** Default Constructor
+         *
+         * As default options, no tuning will be performed, and the number of scheduling units will
+         * depends on internal device discovery functionality
+         */
+        Options()
+            : opts{ AclTuningModeNone, 0 } {};
+        /** Constructor
+         *
+         * @param[in] mode          Tuning mode to be used
+         * @param[in] compute_units Number of scheduling units to be used
+         */
+        Options(TuningMode mode, int32_t compute_units)
+            : opts{ detail::as_cenum<AclTuningMode>(mode), compute_units }
+        {
+        }
+
+        AclQueueOptions opts;
+    };
+
+public:
+    /** Constructor
+     *
+     * @note Serves as a simpler delegate constructor
+     * @note As queue options, default conservative options will be used
+     *
+     * @param[in]  ctx    Context to create queue for
+     * @param[out] status Status information if requested
+     */
+    explicit Queue(Context &ctx, StatusCode *status = nullptr)
+        : Queue(ctx, Options(), status)
+    {
+    }
+    /** Constructor
+     *
+     * @note As queue options, default conservative options will be used
+     *
+     * @param[in]  ctx     Context from where the queue will be created from
+     * @param[in]  options Queue options to be used
+     * @param[out] status  Status information if requested
+     */
+    explicit Queue(Context &ctx, const Options &options = Options(), StatusCode *status = nullptr)
+    {
+        AclQueue   queue;
+        const auto st = detail::as_enum<StatusCode>(AclCreateQueue(&queue, ctx.get(), &options.opts));
+        reset(queue);
+        report_status(st, "[Compute Library] Failed to create queue!");
+        if(status)
+        {
+            *status = st;
+        }
+    }
+    /** Block until all the tasks of the queue have been marked as finished
+     *
+     * @return Status code
+     */
+    StatusCode finish()
+    {
+        return detail::as_enum<StatusCode>(AclQueueFinish(_object.get()));
+    }
+};
+
 /**< Data type enumeration */
 enum class DataType
 {
@@ -519,7 +607,7 @@ public:
         AclTensor  tensor;
         const auto st = detail::as_enum<StatusCode>(AclCreateTensor(&tensor, ctx.get(), desc.get(), allocate));
         reset(tensor);
-        report_status(st, "[Arm Compute Library] Failed to create tensor!");
+        report_status(st, "[Compute Library] Failed to create tensor!");
         if(status)
         {
             *status = st;
@@ -533,7 +621,7 @@ public:
     {
         void      *handle = nullptr;
         const auto st     = detail::as_enum<StatusCode>(AclMapTensor(_object.get(), &handle));
-        report_status(st, "[Arm Compute Library] Failed to map the tensor and extract the tensor's backing memory!");
+        report_status(st, "[Compute Library] Failed to map the tensor and extract the tensor's backing memory!");
         return handle;
     }
     /** Unmaps tensor's memory
@@ -545,7 +633,7 @@ public:
     StatusCode unmap(void *handle)
     {
         const auto st = detail::as_enum<StatusCode>(AclUnmapTensor(_object.get(), handle));
-        report_status(st, "[Arm Compute Library] Failed to unmap the tensor!");
+        report_status(st, "[Compute Library] Failed to unmap the tensor!");
         return st;
     }
     /** Import external memory to a given tensor object
@@ -558,7 +646,7 @@ public:
     StatusCode import(void *handle, ImportType type)
     {
         const auto st = detail::as_enum<StatusCode>(AclTensorImport(_object.get(), handle, detail::as_cenum<AclImportMemoryType>(type)));
-        report_status(st, "[Arm Compute Library] Failed to import external memory to tensor!");
+        report_status(st, "[Compute Library] Failed to import external memory to tensor!");
         return st;
     }
     /** Get the size of the tensor in byte
@@ -571,7 +659,7 @@ public:
     {
         uint64_t   size{ 0 };
         const auto st = detail::as_enum<StatusCode>(AclGetTensorSize(_object.get(), &size));
-        report_status(st, "[Arm Compute Library] Failed to get the size of the tensor");
+        report_status(st, "[Compute Library] Failed to get the size of the tensor");
         return size;
     }
     /** Get the descriptor of this tensor
@@ -582,7 +670,7 @@ public:
     {
         AclTensorDescriptor desc;
         const auto          st = detail::as_enum<StatusCode>(AclGetTensorDescriptor(_object.get(), &desc));
-        report_status(st, "[Arm Compute Library] Failed to get the descriptor of the tensor");
+        report_status(st, "[Compute Library] Failed to get the descriptor of the tensor");
         return TensorDescriptor(desc);
     }
 };
@@ -623,7 +711,7 @@ public:
         AclTensorPack pack;
         const auto    st = detail::as_enum<StatusCode>(AclCreateTensorPack(&pack, ctx.get()));
         reset(pack);
-        report_status(st, "[Arm Compute Library] Failure during tensor pack creation");
+        report_status(st, "[Compute Library] Failure during tensor pack creation");
         if(status)
         {
             *status = st;
diff --git a/arm_compute/AclEntrypoints.h b/arm_compute/AclEntrypoints.h
index cd974341c2..cf4a237a44 100644
--- a/arm_compute/AclEntrypoints.h
+++ b/arm_compute/AclEntrypoints.h
@@ -62,6 +62,49 @@ AclStatus AclCreateContext(AclContext              *ctx,
  */
 AclStatus AclDestroyContext(AclContext ctx);
 
+/** Create an operator queue
+ *
+ * Queue is responsible for any scheduling related activities
+ *
+ * @param[in, out] queue   A valid non-zero queue object is not failures occur
+ * @param[in]      ctx     Context to be used
+ * @param[in]      options Queue options to be used for the operators using the queue
+ *
+ * @return Status code
+ *
+ * Returns:
+ *  - @ref AclSuccess if function was completed successfully
+ *  - @ref AclOutOfMemory if there was a failure allocating memory resources
+ *  - @ref AclUnsupportedTarget if the requested target is unsupported
+ *  - @ref AclInvalidArgument if a given argument is invalid
+ */
+AclStatus AclCreateQueue(AclQueue *queue, AclContext ctx, const AclQueueOptions *options);
+
+/** Wait until all elements on the queue have been completed
+ *
+ * @param[in] queue Queue to wait on completion
+ *
+ * @return Status code
+ *
+ * Returns:
+ *  - @ref AclSuccess if functions was completed successfully
+ *  - @ref AclInvalidArgument if the provided queue is invalid
+ *  - @ref AclRuntimeError on any other runtime related error
+ */
+AclStatus AclQueueFinish(AclQueue queue);
+
+/** Destroy a given queue object
+ *
+ * @param[in] queue A valid context object to destroy
+ *
+ * @return Status code
+ *
+ * Returns:
+ *  - @ref AclSuccess if functions was completed successfully
+ *  - @ref AclInvalidArgument if the provided context is invalid
+ */
+AclStatus AclDestroyQueue(AclQueue queue);
+
 /** Create a Tensor object
  *
  * Tensor is a generalized matrix construct that can represent up to ND dimensionality (where N = 6 for Compute Library)
diff --git a/arm_compute/AclOpenClExt.h b/arm_compute/AclOpenClExt.h
index 15b233ca12..b9080dabf2 100644
--- a/arm_compute/AclOpenClExt.h
+++ b/arm_compute/AclOpenClExt.h
@@ -43,7 +43,6 @@ extern "C" {
 /** Extract the underlying OpenCL context used by a given Compute Library context object
  *
  * @note @ref AclContext should be of an OpenCL backend target
- * @note @ref AclContext refcount should be 0, meaning not used by other objects
  *
  * @param[in]  ctx            A valid non-zero context
  * @param[out] opencl_context Underlying OpenCL context used
@@ -52,7 +51,18 @@ extern "C" {
  */
 AclStatus AclGetClContext(AclContext ctx, cl_context *opencl_context);
 
-/** Set the underlying OpenCL context used by a given Compute Library context object
+/** Extract the underlying OpenCL device id used by a given Compute Library context object
+ *
+ * @note @ref AclContext should be of an OpenCL backend target
+ *
+ * @param[in]  ctx           A valid non-zero context
+ * @param[out] opencl_device Underlying OpenCL device used
+ *
+ * @return Status code
+ */
+AclStatus AclGetClDevice(AclContext ctx, cl_device_id *opencl_device);
+
+/** Set the underlying OpenCL context to be used by a given Compute Library context object
  *
  * @note @ref AclContext should be of an OpenCL backend target
  *
@@ -63,6 +73,30 @@ AclStatus AclGetClContext(AclContext ctx, cl_context *opencl_context);
  */
 AclStatus AclSetClContext(AclContext ctx, cl_context opencl_context);
 
+/** Extract the underlying OpenCL queue used by a given Compute Library queue object
+ *
+ * @note @ref AclQueue should be of an OpenCL backend target
+ * @note @ref AclQueue refcount should be 0, meaning not used by other objects
+ *
+ * @param[in]  queue        A valid non-zero queue
+ * @param[out] opencl_queue Underlying OpenCL queue used
+ *
+ * @return Status code
+ */
+AclStatus AclGetClQueue(AclQueue queue, cl_command_queue *opencl_queue);
+
+/** Set the underlying OpenCL queue to be used by a given Compute Library queue object
+ *
+ * @note @ref AclQueue should be of an OpenCL backend target
+ * @note opecl_queue needs to be created from the same context that the AclContext that the queue will use
+ *
+ * @param[in]  queue        A valid non-zero queue object
+ * @param[out] opencl_queue Underlying OpenCL queue to be used
+ *
+ * @return Status code
+ */
+AclStatus AclSetClQueue(AclQueue queue, cl_command_queue opencl_queue);
+
 /** Extract the underlying OpenCL memory object by a given Compute Library tensor object
  *
  * @param[in]  tensor     A valid non-zero tensor
diff --git a/arm_compute/AclTypes.h b/arm_compute/AclTypes.h
index 69717ec8a8..902a508b91 100644
--- a/arm_compute/AclTypes.h
+++ b/arm_compute/AclTypes.h
@@ -33,6 +33,8 @@ extern "C" {
 
 /**< Opaque Context object */
 typedef struct AclContext_ *AclContext;
+/**< Opaque Queue object */
+typedef struct AclQueue_ *AclQueue;
 /**< Opaque Tensor object */
 typedef struct AclTensor_ *AclTensor;
 /**< Opaque Tensor pack object */
@@ -138,6 +140,22 @@ typedef struct AclContextOptions
     AclAllocator         *allocator;          /**< Allocator to be used by all the memory internally */
 } AclContextOptions;
 
+/**< Supported tuning modes */
+typedef enum
+{
+    AclTuningModeNone = 0, /**< No tuning */
+    AclRapid          = 1, /**< Fast tuning mode, testing a small portion of the tuning space */
+    AclNormal         = 2, /**< Normal tuning mode, gives a good balance between tuning mode and performance */
+    AclExhaustive     = 3, /**< Exhaustive tuning mode, increased tuning time but with best results */
+} AclTuningMode;
+
+/**< Queue options */
+typedef struct
+{
+    AclTuningMode mode;          /**< Tuning mode */
+    int32_t       compute_units; /**< Compute Units that the queue will deploy */
+} AclQueueOptions;
+
 /**< Supported data types */
 typedef enum AclDataType
 {
diff --git a/arm_compute/core/CL/OpenCL.h b/arm_compute/core/CL/OpenCL.h
index 1e6b04c042..bbe469f1a8 100644
--- a/arm_compute/core/CL/OpenCL.h
+++ b/arm_compute/core/CL/OpenCL.h
@@ -92,6 +92,7 @@ public:
     DECLARE_FUNCTION_PTR(clCreateContext);
     DECLARE_FUNCTION_PTR(clCreateContextFromType);
     DECLARE_FUNCTION_PTR(clCreateCommandQueue);
+    DECLARE_FUNCTION_PTR(clCreateCommandQueueWithProperties);
     DECLARE_FUNCTION_PTR(clGetContextInfo);
     DECLARE_FUNCTION_PTR(clBuildProgram);
     DECLARE_FUNCTION_PTR(clEnqueueNDRangeKernel);
diff --git a/docs/01_library.dox b/docs/01_library.dox
index 722a07f958..25535d111a 100644
--- a/docs/01_library.dox
+++ b/docs/01_library.dox
@@ -508,7 +508,27 @@ However this process takes quite a lot of time, which is why it cannot be enable
 
 But, when the @ref CLTuner is disabled ( Target = 1 for the graph examples), the @ref graph::Graph will try to reload the file containing the tuning parameters, then for each executed kernel the Compute Library will use the fine tuned LWS if it was present in the file or use a default LWS value if it's not.
 
-@section S4_10_weights_manager Weights Manager
+@section S4_10_cl_queue_prioritites OpenCL Queue Priorities
+
+OpenCL 2.1 exposes the `cl_khr_priority_hints` extensions that if supported by an underlying implementation allows the user to specify priority hints to the created command queues.
+Is important to note that this does not specify guarantees or the explicit scheduling behavior, this is something that each implementation needs to expose.
+
+In some cases, priority queues can be used when there is an implicit internal priority between graphics and compute queues and thus allow some level of priority control between them.
+At the moment three priority level can be specified:
+- CL_QUEUE_PRIORITY_HIGH_KHR
+- CL_QUEUE_PRIORITY_MED_KHR
+- CL_QUEUE_PRIORITY_LOW_KHR
+
+Compute Library allows extraction of the internal OpenCL queue or the ability to inject directly a user-defined queue to the @ref CLScheduler.
+This way the user can utilize this extension to define priorities between the queues and setup the OpenCL scheduler mechanism to utilize them.
+
+@code{.cpp}
+cl_queue_properties queue_properties[] = {CL_QUEUE_PRIORITY_KHR, CL_QUEUE_PRIORITY_HIGH_KHR, 0};
+cl_command_queue priority_queue = clCreateCommandQueueWithProperties(ctx, dev, queue_properties, &error);
+CLScheduler::get().set_queue(::cl::CommandQueue(priority_queue));
+@endcode
+
+@section S4_11_weights_manager Weights Manager
 
 @ref IWeightsManager is a weights managing interface that can be used to reduce the memory requirements of a given pipeline by reusing transformed weights across multiple function executions.
 @ref IWeightsManager is responsible for managing weight tensors alongside with their transformations.
diff --git a/src/c/AclQueue.cpp b/src/c/AclQueue.cpp
new file mode 100644
index 0000000000..020c6ed531
--- /dev/null
+++ b/src/c/AclQueue.cpp
@@ -0,0 +1,99 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/AclEntrypoints.h"
+
+#include "src/common/IQueue.h"
+#include "src/common/utils/Macros.h"
+#include "src/common/utils/Validate.h"
+
+namespace
+{
+/** Check if queue options are valid
+ *
+ * @param[in] options Queue options
+ *
+ * @return true in case of success else false
+ */
+bool is_mode_valid(const AclQueueOptions *options)
+{
+    ARM_COMPUTE_ASSERT_NOT_NULLPTR(options);
+    return arm_compute::utils::is_in(options->mode, { AclTuningModeNone, AclRapid, AclNormal, AclExhaustive });
+}
+} // namespace
+
+extern "C" AclStatus AclCreateQueue(AclQueue *external_queue, AclContext external_ctx, const AclQueueOptions *options)
+{
+    using namespace arm_compute;
+
+    auto ctx = get_internal(external_ctx);
+
+    StatusCode status = detail::validate_internal_context(ctx);
+    ARM_COMPUTE_RETURN_CENUM_ON_FAILURE(status);
+
+    if(options != nullptr && !is_mode_valid(options))
+    {
+        ARM_COMPUTE_LOG_ERROR_ACL("Queue options are invalid");
+        return AclInvalidArgument;
+    }
+
+    auto queue = ctx->create_queue(options);
+    if(queue == nullptr)
+    {
+        ARM_COMPUTE_LOG_ERROR_ACL("Couldn't allocate internal resources");
+        return AclOutOfMemory;
+    }
+
+    *external_queue = queue;
+
+    return AclSuccess;
+}
+
+extern "C" AclStatus AclQueueFinish(AclQueue external_queue)
+{
+    using namespace arm_compute;
+
+    auto queue = get_internal(external_queue);
+
+    StatusCode status = detail::validate_internal_queue(queue);
+    ARM_COMPUTE_RETURN_CENUM_ON_FAILURE(status);
+
+    status = queue->finish();
+    ARM_COMPUTE_RETURN_CENUM_ON_FAILURE(status);
+
+    return AclSuccess;
+}
+
+extern "C" AclStatus AclDestroyQueue(AclQueue external_queue)
+{
+    using namespace arm_compute;
+
+    auto queue = get_internal(external_queue);
+
+    StatusCode status = detail::validate_internal_queue(queue);
+    ARM_COMPUTE_RETURN_CENUM_ON_FAILURE(status);
+
+    delete queue;
+
+    return AclSuccess;
+}
diff --git a/src/c/cl/AclOpenClExt.cpp b/src/c/cl/AclOpenClExt.cpp
index ce6d2969de..e72babcae8 100644
--- a/src/c/cl/AclOpenClExt.cpp
+++ b/src/c/cl/AclOpenClExt.cpp
@@ -26,6 +26,7 @@
 #include "src/common/ITensorV2.h"
 #include "src/common/Types.h"
 #include "src/gpu/cl/ClContext.h"
+#include "src/gpu/cl/ClQueue.h"
 
 #include "arm_compute/core/CL/ICLTensor.h"
 
@@ -85,6 +86,80 @@ extern "C" AclStatus AclSetClContext(AclContext external_ctx, cl_context opencl_
     return AclStatus::AclSuccess;
 }
 
+extern "C" AclStatus AclGetClDevice(AclContext external_ctx, cl_device_id *opencl_device)
+{
+    using namespace arm_compute;
+    IContext *ctx = get_internal(external_ctx);
+
+    if(detail::validate_internal_context(ctx) != StatusCode::Success)
+    {
+        return AclStatus::AclInvalidArgument;
+    }
+
+    if(ctx->type() != Target::GpuOcl)
+    {
+        return AclStatus::AclInvalidTarget;
+    }
+
+    if(opencl_device == nullptr)
+    {
+        return AclStatus::AclInvalidArgument;
+    }
+
+    *opencl_device = utils::cast::polymorphic_downcast<arm_compute::gpu::opencl::ClContext *>(ctx)->cl_dev().get();
+
+    return AclStatus::AclSuccess;
+}
+
+extern "C" AclStatus AclGetClQueue(AclQueue external_queue, cl_command_queue *opencl_queue)
+{
+    using namespace arm_compute;
+    IQueue *queue = get_internal(external_queue);
+
+    if(detail::validate_internal_queue(queue) != StatusCode::Success)
+    {
+        return AclStatus::AclInvalidArgument;
+    }
+
+    if(queue->header.ctx->type() != Target::GpuOcl)
+    {
+        return AclStatus::AclInvalidTarget;
+    }
+
+    if(opencl_queue == nullptr)
+    {
+        return AclStatus::AclInvalidArgument;
+    }
+
+    *opencl_queue = utils::cast::polymorphic_downcast<arm_compute::gpu::opencl::ClQueue *>(queue)->cl_queue().get();
+
+    return AclStatus::AclSuccess;
+}
+
+extern "C" AclStatus AclSetClQueue(AclQueue external_queue, cl_command_queue opencl_queue)
+{
+    using namespace arm_compute;
+    IQueue *queue = get_internal(external_queue);
+
+    if(detail::validate_internal_queue(queue) != StatusCode::Success)
+    {
+        return AclStatus::AclInvalidArgument;
+    }
+
+    if(queue->header.ctx->type() != Target::GpuOcl)
+    {
+        return AclStatus::AclInvalidTarget;
+    }
+
+    auto cl_queue = utils::cast::polymorphic_downcast<arm_compute::gpu::opencl::ClQueue *>(queue);
+    if(!cl_queue->set_cl_queue(::cl::CommandQueue(opencl_queue)))
+    {
+        return AclStatus::AclRuntimeError;
+    }
+
+    return AclStatus::AclSuccess;
+}
+
 extern "C" AclStatus AclGetClMem(AclTensor external_tensor, cl_mem *opencl_mem)
 {
     using namespace arm_compute;
diff --git a/src/common/IContext.h b/src/common/IContext.h
index ee234795cf..31f39da06d 100644
--- a/src/common/IContext.h
+++ b/src/common/IContext.h
@@ -43,6 +43,7 @@ namespace arm_compute
 {
 // Forward declarations
 class ITensorV2;
+class IQueue;
 
 /**< Context interface */
 class IContext : public AclContext_
@@ -52,11 +53,13 @@ public:
         : AclContext_(), _target(target), _refcount(0)
     {
     }
+
     /** Virtual Destructor */
     virtual ~IContext()
     {
         header.type = detail::ObjectType::Invalid;
     };
+
     /** Target type accessor
      *
      * @return Target that the context is associated with
@@ -65,16 +68,19 @@ public:
     {
         return _target;
     }
+
     /** Increment context refcount */
     void inc_ref() const
     {
         ++_refcount;
     }
+
     /** Decrement context refcount */
     void dec_ref() const
     {
         --_refcount;
     }
+
     /** Reference counter accessor
      *
      * @return The number of references pointing to this object
@@ -83,6 +89,7 @@ public:
     {
         return _refcount;
     }
+
     /** Checks if an object is valid
      *
      * @return True if sucessful otherwise false
@@ -91,6 +98,7 @@ public:
     {
         return header.type == detail::ObjectType::Context;
     }
+
     /** Create a tensor object
      *
      * @param[in] desc     Descriptor to use
@@ -100,6 +108,14 @@ public:
      */
     virtual ITensorV2 *create_tensor(const AclTensorDescriptor &desc, bool allocate) = 0;
 
+    /** Create a queue object
+     *
+     * @param[in] options Queue options to be used
+     *
+     * @return A pointer to the created queue object
+     */
+    virtual IQueue *create_queue(const AclQueueOptions *options) = 0;
+
 private:
     Target                   _target;   /**< Target type of context */
     mutable std::atomic<int> _refcount; /**< Reference counter */
diff --git a/src/common/IQueue.h b/src/common/IQueue.h
new file mode 100644
index 0000000000..6a0cbc75da
--- /dev/null
+++ b/src/common/IQueue.h
@@ -0,0 +1,100 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef SRC_COMMON_IQUEUE_H_
+#define SRC_COMMON_IQUEUE_H_
+
+#include "src/common/IContext.h"
+
+struct AclQueue_
+{
+    arm_compute::detail::Header header{ arm_compute::detail::ObjectType::Queue, nullptr };
+
+protected:
+    AclQueue_()  = default;
+    ~AclQueue_() = default;
+};
+
+namespace arm_compute
+{
+/** Base class specifying the queue interface */
+class IQueue : public AclQueue_
+{
+public:
+    /** Explict Operator Constructor
+     *
+     * @param[in] ctx Context to be used by the operator
+     */
+    explicit IQueue(IContext *ctx)
+    {
+        this->header.ctx = ctx;
+        this->header.ctx->inc_ref();
+    }
+    /** Destructor */
+    virtual ~IQueue()
+    {
+        this->header.ctx->dec_ref();
+        this->header.type = detail::ObjectType::Invalid;
+    };
+    /** Checks if a queue is valid
+     *
+     * @return True if successful otherwise false
+     */
+    bool is_valid() const
+    {
+        return this->header.type == detail::ObjectType::Queue;
+    };
+    virtual StatusCode finish() = 0;
+};
+
+/** Extract internal representation of a Queue
+ *
+ * @param[in] queue Opaque queue pointer
+ *
+ * @return The internal representation as an IQueue
+ */
+inline IQueue *get_internal(AclQueue queue)
+{
+    return static_cast<IQueue *>(queue);
+}
+
+namespace detail
+{
+/** Check if an internal queue is valid
+ *
+ * @param[in] queue Internal queue to check
+ *
+ * @return A status code
+ */
+inline StatusCode validate_internal_queue(const IQueue *queue)
+{
+    if(queue == nullptr || !queue->is_valid())
+    {
+        ARM_COMPUTE_LOG_ERROR_ACL("[IQueue]: Invalid queue object");
+        return StatusCode::InvalidArgument;
+    }
+    return StatusCode::Success;
+}
+} // namespace detail
+} // namespace arm_compute
+#endif /* SRC_COMMON_IQUEUE_H_ */
diff --git a/src/core/CL/OpenCL.cpp b/src/core/CL/OpenCL.cpp
index a7be534397..d8c2736ef7 100644
--- a/src/core/CL/OpenCL.cpp
+++ b/src/core/CL/OpenCL.cpp
@@ -91,6 +91,7 @@ bool CLSymbols::load(const std::string &library)
     LOAD_FUNCTION_PTR(clCreateContext, handle);
     LOAD_FUNCTION_PTR(clCreateContextFromType, handle);
     LOAD_FUNCTION_PTR(clCreateCommandQueue, handle);
+    LOAD_FUNCTION_PTR(clCreateCommandQueueWithProperties, handle);
     LOAD_FUNCTION_PTR(clGetContextInfo, handle);
     LOAD_FUNCTION_PTR(clBuildProgram, handle);
     LOAD_FUNCTION_PTR(clEnqueueNDRangeKernel, handle);
@@ -293,6 +294,23 @@ cl_command_queue clCreateCommandQueue(cl_context                  context,
     }
 }
 
+cl_command_queue clCreateCommandQueueWithProperties(cl_context                 context,
+                                                    cl_device_id               device,
+                                                    const cl_queue_properties *properties,
+                                                    cl_int                    *errcode_ret)
+{
+    arm_compute::CLSymbols::get().load_default();
+    auto func = arm_compute::CLSymbols::get().clCreateCommandQueueWithProperties_ptr;
+    if(func != nullptr)
+    {
+        return func(context, device, properties, errcode_ret);
+    }
+    else
+    {
+        return nullptr;
+    }
+}
+
 cl_context clCreateContext(
     const cl_context_properties *properties,
     cl_uint                      num_devices,
diff --git a/src/core/cpu/kernels/CpuDirectConvolutionStageKernel.cpp b/src/core/cpu/kernels/CpuDirectConvolutionOutputStageKernel.cpp
index 5f7a574e5a..5f7a574e5a 100644
--- a/src/core/cpu/kernels/CpuDirectConvolutionStageKernel.cpp
+++ b/src/core/cpu/kernels/CpuDirectConvolutionOutputStageKernel.cpp
diff --git a/src/core/cpu/kernels/activation/NEON/fp16.cpp b/src/core/cpu/kernels/activation/neon/fp16.cpp
index 6f2d5d8533..6f2d5d8533 100644
--- a/src/core/cpu/kernels/activation/NEON/fp16.cpp
+++ b/src/core/cpu/kernels/activation/neon/fp16.cpp
diff --git a/src/core/cpu/kernels/activation/NEON/fp32.cpp b/src/core/cpu/kernels/activation/neon/fp32.cpp
index 54301d45ad..54301d45ad 100644
--- a/src/core/cpu/kernels/activation/NEON/fp32.cpp
+++ b/src/core/cpu/kernels/activation/neon/fp32.cpp
diff --git a/src/core/cpu/kernels/activation/NEON/qasymm8.cpp b/src/core/cpu/kernels/activation/neon/qasymm8.cpp
index a1217435b6..a1217435b6 100644
--- a/src/core/cpu/kernels/activation/NEON/qasymm8.cpp
+++ b/src/core/cpu/kernels/activation/neon/qasymm8.cpp
diff --git a/src/core/cpu/kernels/activation/NEON/qasymm8_signed.cpp b/src/core/cpu/kernels/activation/neon/qasymm8_signed.cpp
index 8b40bf8e72..8b40bf8e72 100644
--- a/src/core/cpu/kernels/activation/NEON/qasymm8_signed.cpp
+++ b/src/core/cpu/kernels/activation/neon/qasymm8_signed.cpp
diff --git a/src/core/cpu/kernels/activation/NEON/qsymm16.cpp b/src/core/cpu/kernels/activation/neon/qsymm16.cpp
index 54b41820f2..54b41820f2 100644
--- a/src/core/cpu/kernels/activation/NEON/qsymm16.cpp
+++ b/src/core/cpu/kernels/activation/neon/qsymm16.cpp
diff --git a/src/core/cpu/kernels/activation/SVE/fp16.cpp b/src/core/cpu/kernels/activation/sve/fp16.cpp
index bf31fd7d93..bf31fd7d93 100644
--- a/src/core/cpu/kernels/activation/SVE/fp16.cpp
+++ b/src/core/cpu/kernels/activation/sve/fp16.cpp
diff --git a/src/core/cpu/kernels/activation/SVE/fp32.cpp b/src/core/cpu/kernels/activation/sve/fp32.cpp
index 75f9f8a4c3..75f9f8a4c3 100644
--- a/src/core/cpu/kernels/activation/SVE/fp32.cpp
+++ b/src/core/cpu/kernels/activation/sve/fp32.cpp
diff --git a/src/core/cpu/kernels/activation/SVE/qasymm8.cpp b/src/core/cpu/kernels/activation/sve/qasymm8.cpp
index 228b4ae530..228b4ae530 100644
--- a/src/core/cpu/kernels/activation/SVE/qasymm8.cpp
+++ b/src/core/cpu/kernels/activation/sve/qasymm8.cpp
diff --git a/src/core/cpu/kernels/activation/SVE/qasymm8_signed.cpp b/src/core/cpu/kernels/activation/sve/qasymm8_signed.cpp
index 989f825eb9..989f825eb9 100644
--- a/src/core/cpu/kernels/activation/SVE/qasymm8_signed.cpp
+++ b/src/core/cpu/kernels/activation/sve/qasymm8_signed.cpp
diff --git a/src/core/cpu/kernels/activation/SVE/qsymm16.cpp b/src/core/cpu/kernels/activation/sve/qsymm16.cpp
index 66974875da..66974875da 100644
--- a/src/core/cpu/kernels/activation/SVE/qsymm16.cpp
+++ b/src/core/cpu/kernels/activation/sve/qsymm16.cpp
diff --git a/src/cpu/CpuContext.cpp b/src/cpu/CpuContext.cpp
index d62c1b6310..b9a6999f84 100644
--- a/src/cpu/CpuContext.cpp
+++ b/src/cpu/CpuContext.cpp
@@ -24,6 +24,7 @@
 #include "src/cpu/CpuContext.h"
 
 #include "arm_compute/core/CPP/CPPTypes.h"
+#include "src/cpu/CpuQueue.h"
 #include "src/cpu/CpuTensor.h"
 #include "src/runtime/CPUUtils.h"
 
@@ -196,5 +197,10 @@ ITensorV2 *CpuContext::create_tensor(const AclTensorDescriptor &desc, bool alloc
     }
     return tensor;
 }
+
+IQueue *CpuContext::create_queue(const AclQueueOptions *options)
+{
+    return new CpuQueue(this, options);
+}
 } // namespace cpu
 } // namespace arm_compute
diff --git a/src/cpu/CpuContext.h b/src/cpu/CpuContext.h
index d2062e4bdd..e909767a7b 100644
--- a/src/cpu/CpuContext.h
+++ b/src/cpu/CpuContext.h
@@ -69,6 +69,7 @@ public:
 
     // Inherrited methods overridden
     ITensorV2 *create_tensor(const AclTensorDescriptor &desc, bool allocate) override;
+    IQueue *create_queue(const AclQueueOptions *options) override;
 
 private:
     AllocatorWrapper _allocator;
diff --git a/src/cpu/CpuQueue.cpp b/src/cpu/CpuQueue.cpp
new file mode 100644
index 0000000000..0f0097b3f4
--- /dev/null
+++ b/src/cpu/CpuQueue.cpp
@@ -0,0 +1,48 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "src/cpu/CpuQueue.h"
+
+#include "arm_compute/runtime/Scheduler.h"
+
+namespace arm_compute
+{
+namespace cpu
+{
+CpuQueue::CpuQueue(IContext *ctx, const AclQueueOptions *options)
+    : IQueue(ctx)
+{
+    ARM_COMPUTE_UNUSED(options);
+}
+
+arm_compute::IScheduler &CpuQueue::scheduler()
+{
+    return arm_compute::Scheduler::get();
+}
+
+StatusCode CpuQueue::finish()
+{
+    return StatusCode::Success;
+}
+} // namespace cpu
+} // namespace arm_compute
diff --git a/src/cpu/CpuQueue.h b/src/cpu/CpuQueue.h
new file mode 100644
index 0000000000..871a36c85b
--- /dev/null
+++ b/src/cpu/CpuQueue.h
@@ -0,0 +1,56 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef SRC_CPU_CPUQUEUE_H
+#define SRC_CPU_CPUQUEUE_H
+
+#include "src/common/IQueue.h"
+
+#include "arm_compute/runtime/IScheduler.h"
+
+namespace arm_compute
+{
+namespace cpu
+{
+/** CPU queue implementation class */
+class CpuQueue final : public IQueue
+{
+public:
+    /** Construct a new CpuQueue object
+     *
+     * @param[in] ctx     Context to be used
+     * @param[in] options Command queue options
+     */
+    CpuQueue(IContext *ctx, const AclQueueOptions *options);
+    /** Return legacy scheduler
+     *
+     * @return arm_compute::IScheduler&
+     */
+    arm_compute::IScheduler &scheduler();
+
+    // Inherited functions overridden
+    StatusCode finish() override;
+};
+} // namespace cpu
+} // namespace arm_compute
+#endif /* SRC_CPU_CPUQUEUE_H */
diff --git a/src/gpu/cl/ClContext.cpp b/src/gpu/cl/ClContext.cpp
index 2e04e1d593..d8ef18e62e 100644
--- a/src/gpu/cl/ClContext.cpp
+++ b/src/gpu/cl/ClContext.cpp
@@ -23,8 +23,11 @@
  */
 #include "src/gpu/cl/ClContext.h"
 
+#include "src/gpu/cl/ClQueue.h"
 #include "src/gpu/cl/ClTensor.h"
 
+#include "arm_compute/core/CL/CLKernelLibrary.h"
+
 namespace arm_compute
 {
 namespace gpu
@@ -49,12 +52,15 @@ mlgo::MLGOHeuristics populate_mlgo(const char *filename)
 ClContext::ClContext(const AclContextOptions *options)
     : IContext(Target::GpuOcl),
       _mlgo_heuristics(),
-      _cl_context()
+      _cl_ctx(),
+      _cl_dev()
 {
     if(options != nullptr)
     {
         _mlgo_heuristics = populate_mlgo(options->kernel_config_file);
     }
+    _cl_ctx = CLKernelLibrary::get().context();
+    _cl_dev = CLKernelLibrary::get().get_device();
 }
 
 const mlgo::MLGOHeuristics &ClContext::mlgo() const
@@ -64,14 +70,20 @@ const mlgo::MLGOHeuristics &ClContext::mlgo() const
 
 ::cl::Context ClContext::cl_ctx()
 {
-    return _cl_context;
+    return _cl_ctx;
+}
+
+::cl::Device ClContext::cl_dev()
+{
+    return _cl_dev;
 }
 
 bool ClContext::set_cl_ctx(::cl::Context ctx)
 {
     if(this->refcount() == 0)
     {
-        _cl_context = ctx;
+        _cl_ctx = ctx;
+        CLScheduler::get().set_context(ctx);
         return true;
     }
     return false;
@@ -86,6 +98,11 @@ ITensorV2 *ClContext::create_tensor(const AclTensorDescriptor &desc, bool alloca
     }
     return tensor;
 }
+
+IQueue *ClContext::create_queue(const AclQueueOptions *options)
+{
+    return new ClQueue(this, options);
+}
 } // namespace opencl
 } // namespace gpu
 } // namespace arm_compute
diff --git a/src/gpu/cl/ClContext.h b/src/gpu/cl/ClContext.h
index dd6699a0c9..2a0d4ee1c8 100644
--- a/src/gpu/cl/ClContext.h
+++ b/src/gpu/cl/ClContext.h
@@ -44,6 +44,7 @@ public:
      * @param[in] options Creational options
      */
     explicit ClContext(const AclContextOptions *options);
+
     /** Extract MLGO heuristics
      *
      * @return Heuristics tree
@@ -55,6 +56,13 @@ public:
      * @return the cl context used
      */
     ::cl::Context cl_ctx();
+
+    /** Underlying cl device accessor
+     *
+     * @return the cl device used
+     */
+    ::cl::Device cl_dev();
+
     /** Update/inject an underlying cl context object
      *
      * @warning Context will be able to set if the object doesn't have any pending reference to other objects
@@ -67,10 +75,12 @@ public:
 
     // Inherrited methods overridden
     ITensorV2 *create_tensor(const AclTensorDescriptor &desc, bool allocate) override;
+    IQueue *create_queue(const AclQueueOptions *options) override;
 
 private:
     mlgo::MLGOHeuristics _mlgo_heuristics;
-    ::cl::Context        _cl_context;
+    ::cl::Context        _cl_ctx;
+    ::cl::Device         _cl_dev;
 };
 } // namespace opencl
 } // namespace gpu
diff --git a/src/gpu/cl/ClQueue.cpp b/src/gpu/cl/ClQueue.cpp
new file mode 100644
index 0000000000..2123adcf39
--- /dev/null
+++ b/src/gpu/cl/ClQueue.cpp
@@ -0,0 +1,102 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "src/gpu/cl/ClQueue.h"
+
+#include "arm_compute/runtime/CL/CLScheduler.h"
+#include "arm_compute/runtime/CL/CLTuner.h"
+
+namespace arm_compute
+{
+namespace gpu
+{
+namespace opencl
+{
+namespace
+{
+CLTunerMode map_tuner_mode(AclTuningMode mode)
+{
+    switch(mode)
+    {
+        case AclRapid:
+            return CLTunerMode::RAPID;
+            break;
+        case AclNormal:
+            return CLTunerMode::NORMAL;
+            break;
+        case AclExhaustive:
+            return CLTunerMode::EXHAUSTIVE;
+            break;
+        default:
+            ARM_COMPUTE_ERROR("Invalid tuner mode");
+            break;
+    }
+}
+
+std::unique_ptr<CLTuner> populate_tuner(const AclQueueOptions *options)
+{
+    if(options == nullptr || options->mode == AclTuningModeNone)
+    {
+        return nullptr;
+    }
+
+    CLTuningInfo tune_info;
+    tune_info.tuner_mode = map_tuner_mode(options->mode);
+    tune_info.tune_wbsm  = false;
+
+    return std::make_unique<CLTuner>(true /* tune_new_kernels */, tune_info);
+}
+} // namespace
+
+ClQueue::ClQueue(IContext *ctx, const AclQueueOptions *options)
+    : IQueue(ctx), _tuner(nullptr)
+{
+    _tuner = populate_tuner(options);
+}
+
+arm_compute::CLScheduler &ClQueue::scheduler()
+{
+    return arm_compute::CLScheduler::get();
+}
+
+::cl::CommandQueue ClQueue::cl_queue()
+{
+    return arm_compute::CLScheduler::get().queue();
+}
+
+bool ClQueue::set_cl_queue(::cl::CommandQueue queue)
+{
+    // TODO: Check queue is from the same context
+    arm_compute::CLScheduler::get().set_queue(queue);
+    return true;
+}
+
+StatusCode ClQueue::finish()
+{
+    arm_compute::CLScheduler::get().queue().finish();
+    return StatusCode::Success;
+}
+
+} // namespace opencl
+} // namespace gpu
+} // namespace arm_compute
diff --git a/src/gpu/cl/ClQueue.h b/src/gpu/cl/ClQueue.h
new file mode 100644
index 0000000000..b16a0f4e83
--- /dev/null
+++ b/src/gpu/cl/ClQueue.h
@@ -0,0 +1,84 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef SRC_GPU_CLQUEUE_H
+#define SRC_GPU_CLQUEUE_H
+
+#include "src/common/IQueue.h"
+
+#include "arm_compute/runtime/CL/CLScheduler.h"
+
+#include <memory>
+
+namespace arm_compute
+{
+// Forward declarations
+class CLTuner;
+
+namespace gpu
+{
+namespace opencl
+{
+/** OpenCL queue implementation class */
+class ClQueue final : public IQueue
+{
+public:
+    /** Construct a new CpuQueue object
+     *
+     * @param[in] ctx     Context to be used
+     * @param[in] options Command queue options
+     */
+    ClQueue(IContext *ctx, const AclQueueOptions *options);
+
+    /** Return legacy scheduler
+     *
+     * @return arm_compute::IScheduler&
+     */
+    arm_compute::CLScheduler &scheduler();
+
+    /** Underlying cl command queue accessor
+     *
+     * @return the cl command queue used
+     */
+    ::cl::CommandQueue cl_queue();
+
+    /** Update/inject an underlying cl command queue object
+     *
+     * @warning Command queue needs to come from the same context as the AclQueue
+     *
+     * @param[in] queue Underlying cl command queue to be used
+     *
+     * @return true if the queue was set successfully else falseS
+     */
+    bool set_cl_queue(::cl::CommandQueue queue);
+
+    // Inherited functions overridden
+    StatusCode finish() override;
+
+private:
+    std::unique_ptr<CLTuner> _tuner;
+};
+} // namespace opencl
+} // namespace gpu
+} // namespace arm_compute
+#endif /* SRC_GPU_CLQUEUE_H */
diff --git a/tests/framework/Macros.h b/tests/framework/Macros.h
index 23c826657d..a6ba137b59 100644
--- a/tests/framework/Macros.h
+++ b/tests/framework/Macros.h
@@ -224,6 +224,11 @@
 #define DISABLED_FIXTURE_TEST_CASE(TEST_NAME, FIXTURE, MODE) \
     FIXTURE_TEST_CASE_IMPL(TEST_NAME, FIXTURE, MODE, arm_compute::test::framework::TestCaseFactory::Status::DISABLED)
 
+#define EMPTY_BODY_FIXTURE_TEST_CASE(TEST_NAME, FIXTURE, MODE) \
+    FIXTURE_TEST_CASE(TEST_NAME, FIXTURE, MODE)                \
+    {                                                          \
+    }
+
 #define FIXTURE_DATA_TEST_CASE_IMPL(TEST_NAME, FIXTURE, MODE, STATUS, DATASET)                                                      \
     template <typename T>                                                                                                           \
     class TEST_NAME;                                                                                                                \
diff --git a/tests/validation/cpu/unit/Context.cpp b/tests/validation/cpu/unit/Context.cpp
index 519a7bee5f..57ca866032 100644
--- a/tests/validation/cpu/unit/Context.cpp
+++ b/tests/validation/cpu/unit/Context.cpp
@@ -21,7 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "tests/validation/fixtures/UNIT/Context.h"
+#include "tests/validation/fixtures/UNIT/ContextFixture.h"
 
 #include "src/cpu/CpuContext.h"
 
@@ -74,18 +74,10 @@ TEST_CASE(CreateContextWithInvalidOptions, framework::DatasetMode::ALL)
     ARM_COMPUTE_ASSERT(ctx == nullptr);
 }
 
-FIXTURE_TEST_CASE(DestroyInvalidContext, DestroyInvalidContextFixture<AclTarget::AclCpu>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(SimpleContextCApi, SimpleContextCApiFixture<AclTarget::AclCpu>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(SimpleContextCppApi, SimpleContextCppApiFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(MultipleContexts, MultipleContextsFixture<AclTarget::AclCpu>, framework::DatasetMode::ALL)
-{
-}
+EMPTY_BODY_FIXTURE_TEST_CASE(DestroyInvalidContext, DestroyInvalidContextFixture<AclTarget::AclCpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleContextCApi, SimpleContextCApiFixture<AclTarget::AclCpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleContextCppApi, SimpleContextCppApiFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MultipleContexts, MultipleContextsFixture<AclTarget::AclCpu>, framework::DatasetMode::ALL)
 
 /** Test-case for CpuCapabilities
  *
diff --git a/tests/validation/cpu/unit/Queue.cpp b/tests/validation/cpu/unit/Queue.cpp
new file mode 100644
index 0000000000..7d977cc48e
--- /dev/null
+++ b/tests/validation/cpu/unit/Queue.cpp
@@ -0,0 +1,48 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "tests/validation/fixtures/UNIT/QueueFixture.h"
+
+#include "src/cpu/CpuQueue.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CPU)
+TEST_SUITE(UNIT)
+TEST_SUITE(Queue)
+
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateQueueWithInvalidContext, CreateQueueWithInvalidContextFixture, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateQueuerWithInvalidOptions, CreateQueuerWithInvalidOptionsFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(DestroyInvalidQueue, DestroyInvalidQueueFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleQueue, SimpleQueueFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+
+TEST_SUITE_END() // Queue
+TEST_SUITE_END() // UNIT
+TEST_SUITE_END() // CPU
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/cpu/unit/Tensor.cpp b/tests/validation/cpu/unit/Tensor.cpp
index aa2e3abdf1..cc0c55758f 100644
--- a/tests/validation/cpu/unit/Tensor.cpp
+++ b/tests/validation/cpu/unit/Tensor.cpp
@@ -21,7 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "tests/validation/fixtures/UNIT/Tensor.h"
+#include "tests/validation/fixtures/UNIT/TensorFixture.h"
 
 namespace arm_compute
 {
@@ -33,26 +33,19 @@ TEST_SUITE(CPU)
 TEST_SUITE(UNIT)
 TEST_SUITE(Tensor)
 
-#define TENSOR_TESE_CASE(name, fixture)                           \
-    FIXTURE_TEST_CASE(name, fixture, framework::DatasetMode::ALL) \
-    {                                                             \
-    }
-
-TENSOR_TESE_CASE(CreateTensorWithInvalidContext, CreateTensorWithInvalidContextFixture)
-TENSOR_TESE_CASE(CreateTensorWithInvalidDescriptor, CreateTensorWithInvalidDescriptorFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(DestroyInvalidTensor, DestroyInvalidTensorFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(SimpleTensor, SimpleTensorFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(TensorStress, TensorStressFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(MapInvalidTensor, MapInvalidTensorFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(MapNotAllocatedTensor, MapNotAllocatedTensorFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(MapAllocatedTensor, MapAllocatedTensorFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(ImportMemory, ImportMemoryFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(GetSize, TensorSizeFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(GetInvalidSize, InvalidTensorSizeFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(GetDescriptor, DescriptorConversionFixture<acl::Target::Cpu>)
-TENSOR_TESE_CASE(GetInvalidDescriptor, InvalidDescriptorConversionFixture<acl::Target::Cpu>)
-
-#undef TENSOR_TEST_CASE
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateTensorWithInvalidContext, CreateTensorWithInvalidContextFixture, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateTensorWithInvalidDescriptor, CreateTensorWithInvalidDescriptorFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(DestroyInvalidTensor, DestroyInvalidTensorFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleTensor, SimpleTensorFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(TensorStress, TensorStressFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MapInvalidTensor, MapInvalidTensorFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MapNotAllocatedTensor, MapNotAllocatedTensorFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MapAllocatedTensor, MapAllocatedTensorFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(ImportMemory, ImportMemoryFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetSize, TensorSizeFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetInvalidSize, InvalidTensorSizeFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetDescriptor, DescriptorConversionFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetInvalidDescriptor, InvalidDescriptorConversionFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
 
 TEST_SUITE_END() // Tensor
 TEST_SUITE_END() // UNIT
diff --git a/tests/validation/cpu/unit/TensorPack.cpp b/tests/validation/cpu/unit/TensorPack.cpp
index 5436ceb0c1..f019e8e3c4 100644
--- a/tests/validation/cpu/unit/TensorPack.cpp
+++ b/tests/validation/cpu/unit/TensorPack.cpp
@@ -21,7 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "tests/validation/fixtures/UNIT/TensorPack.h"
+#include "tests/validation/fixtures/UNIT/TensorPackFixture.h"
 
 namespace arm_compute
 {
@@ -33,21 +33,11 @@ TEST_SUITE(CPU)
 TEST_SUITE(UNIT)
 TEST_SUITE(TensorPack)
 
-FIXTURE_TEST_CASE(CreateTensorPackWithInvalidContext, CreateTensorPackWithInvalidContextFixture, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(DestroyInvalidTensorPack, DestroyInvalidTensorPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(AddInvalidObjectToTensorPack, AddInvalidObjectToTensorPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(SimpleTensorPack, SimpleTensorPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(MultipleTensorsInPack, MultipleTensorsInPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
-{
-}
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateTensorPackWithInvalidContext, CreateTensorPackWithInvalidContextFixture, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(DestroyInvalidTensorPack, DestroyInvalidTensorPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(AddInvalidObjectToTensorPack, AddInvalidObjectToTensorPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleTensorPack, SimpleTensorPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MultipleTensorsInPack, MultipleTensorsInPackFixture<acl::Target::Cpu>, framework::DatasetMode::ALL)
 
 TEST_SUITE_END() // Tensor
 TEST_SUITE_END() // UNIT
diff --git a/tests/validation/fixtures/UNIT/Context.h b/tests/validation/fixtures/UNIT/ContextFixture.h
index afa49e00e0..77cbc12320 100644
--- a/tests/validation/fixtures/UNIT/Context.h
+++ b/tests/validation/fixtures/UNIT/ContextFixture.h
@@ -21,8 +21,8 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#ifndef ARM_COMPUTE_TEST_UNIT_CONTEXT
-#define ARM_COMPUTE_TEST_UNIT_CONTEXT
+#ifndef ARM_COMPUTE_TEST_UNIT_CONTEXT_FIXTURE
+#define ARM_COMPUTE_TEST_UNIT_CONTEXT_FIXTURE
 
 #include "arm_compute/Acl.hpp"
 #include "tests/framework/Asserts.h"
@@ -145,4 +145,4 @@ public:
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
-#endif /* ARM_COMPUTE_TEST_UNIT_CONTEXT */
+#endif /* ARM_COMPUTE_TEST_UNIT_CONTEXT_FIXTURE */
diff --git a/tests/validation/fixtures/UNIT/QueueFixture.h b/tests/validation/fixtures/UNIT/QueueFixture.h
new file mode 100644
index 0000000000..bc93f5f120
--- /dev/null
+++ b/tests/validation/fixtures/UNIT/QueueFixture.h
@@ -0,0 +1,144 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_UNIT_QUEUE_FIXTURE
+#define ARM_COMPUTE_TEST_UNIT_QUEUE_FIXTURE
+
+#include "arm_compute/Acl.hpp"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/framework/Macros.h"
+#include "tests/validation/Validation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+/** Test case for AclCreateQueue
+ *
+ * Validate that AclCreateQueue behaves as expected with invalid context
+ *
+ * Test Steps:
+ *  - Call AclCreateQueue with an invalid context
+ *  - Confirm that AclInvalidArgument is reported
+ *  - Confirm that the queue is still nullptr
+ */
+class CreateQueueWithInvalidContextFixture : public framework::Fixture
+{
+public:
+    void setup()
+    {
+        AclQueue queue = nullptr;
+        ARM_COMPUTE_ASSERT(AclCreateQueue(&queue, nullptr, nullptr) == AclStatus::AclInvalidArgument);
+        ARM_COMPUTE_ASSERT(queue == nullptr);
+    };
+};
+
+/** Test-case for AclCreateQueue
+ *
+ * Validate that AclCreateQueue behaves as expected with invalid options
+ *
+ * Test Steps:
+ *  - Call AclCreateQueue with valid context but invalid options
+ *  - Confirm that AclInvalidArgument is reported
+ *  - Confirm that queue is still nullptr
+ */
+template <acl::Target Target>
+class CreateQueuerWithInvalidOptionsFixture : public framework::Fixture
+{
+public:
+    void setup()
+    {
+        acl::Context ctx(Target);
+
+        // Check invalid tuning mode
+        AclQueueOptions invalid_queue_opts;
+        invalid_queue_opts.mode = static_cast<AclTuningMode>(-1);
+
+        AclQueue queue = nullptr;
+        ARM_COMPUTE_ASSERT(AclCreateQueue(&queue, ctx.get(), &invalid_queue_opts) == AclStatus::AclInvalidArgument);
+        ARM_COMPUTE_ASSERT(queue == nullptr);
+    };
+};
+
+/** Test case for AclDestroyQueue
+*
+* Validate that AclDestroyQueue behaves as expected when an invalid queue is given
+*
+* Test Steps:
+*  - Call AclDestroyQueue with null queue
+*  - Confirm that AclInvalidArgument is reported
+*  - Call AclDestroyQueue on empty array
+*  - Confirm that AclInvalidArgument is reported
+*  - Call AclDestroyQueue on an ACL object other than AclQueue
+*  - Confirm that AclInvalidArgument is reported
+*  - Confirm that queue is still nullptr
+*/
+template <acl::Target Target>
+class DestroyInvalidQueueFixture : public framework::Fixture
+{
+public:
+    void setup()
+    {
+        acl::Context ctx(Target);
+
+        std::array<char, 256> empty_array{};
+        AclQueue queue = nullptr;
+
+        ARM_COMPUTE_ASSERT(AclDestroyQueue(queue) == AclStatus::AclInvalidArgument);
+        ARM_COMPUTE_ASSERT(AclDestroyQueue(reinterpret_cast<AclQueue>(ctx.get())) == AclStatus::AclInvalidArgument);
+        ARM_COMPUTE_ASSERT(AclDestroyQueue(reinterpret_cast<AclQueue>(empty_array.data())) == AclStatus::AclInvalidArgument);
+        ARM_COMPUTE_ASSERT(queue == nullptr);
+    };
+};
+
+/** Test case for AclCreateQueue
+ *
+ * Validate that a queue can be created successfully
+ *
+ * Test Steps:
+ *  - Create a valid context
+ *  - Create a valid queue
+ *  - Confirm that AclSuccess is returned
+ */
+template <acl::Target Target>
+class SimpleQueueFixture : public framework::Fixture
+{
+public:
+    void setup()
+    {
+        acl::StatusCode err = acl::StatusCode::Success;
+
+        acl::Context ctx(Target, &err);
+        ARM_COMPUTE_ASSERT(err == acl::StatusCode::Success);
+
+        acl::Queue queue(ctx, &err);
+        ARM_COMPUTE_ASSERT(err == acl::StatusCode::Success);
+    };
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_UNIT_QUEUE_FIXTURE */
diff --git a/tests/validation/fixtures/UNIT/Tensor.h b/tests/validation/fixtures/UNIT/TensorFixture.h
index 32260cb431..bfe115b3ed 100644
--- a/tests/validation/fixtures/UNIT/Tensor.h
+++ b/tests/validation/fixtures/UNIT/TensorFixture.h
@@ -21,8 +21,8 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#ifndef ARM_COMPUTE_TEST_UNIT_TENSOR
-#define ARM_COMPUTE_TEST_UNIT_TENSOR
+#ifndef ARM_COMPUTE_TEST_UNIT_TENSOR_FIXTURE
+#define ARM_COMPUTE_TEST_UNIT_TENSOR_FIXTURE
 
 #include "arm_compute/Acl.hpp"
 #include "tests/framework/Asserts.h"
@@ -149,7 +149,7 @@ public:
 /** Test case for AclTensor
  *
  * Validate that multiple tensors can be created successfully
- * Possibly stress the possibility of memory leaks
+ * Stress the possibility of memory leaks
  *
  * Test Steps:
  *  - Create a valid context
@@ -421,4 +421,4 @@ public:
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
-#endif /* ARM_COMPUTE_TEST_UNIT_TENSOR */
+#endif /* ARM_COMPUTE_TEST_UNIT_TENSOR_FIXTURE */
diff --git a/tests/validation/fixtures/UNIT/TensorPack.h b/tests/validation/fixtures/UNIT/TensorPackFixture.h
index 98bffb1665..bc14631936 100644
--- a/tests/validation/fixtures/UNIT/TensorPack.h
+++ b/tests/validation/fixtures/UNIT/TensorPackFixture.h
@@ -21,8 +21,8 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#ifndef ARM_COMPUTE_TEST_UNIT_TENSORPACK
-#define ARM_COMPUTE_TEST_UNIT_TENSORPACK
+#ifndef ARM_COMPUTE_TEST_UNIT_TENSORPACK_FIXTURE
+#define ARM_COMPUTE_TEST_UNIT_TENSORPACK_FIXTURE
 
 #include "arm_compute/Acl.hpp"
 #include "tests/framework/Asserts.h"
@@ -181,4 +181,4 @@ public:
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
-#endif /* ARM_COMPUTE_TEST_UNIT_TENSORPACK */
+#endif /* ARM_COMPUTE_TEST_UNIT_TENSORPACK_FIXTURE */
diff --git a/tests/validation/gpu/unit/Context.cpp b/tests/validation/gpu/unit/Context.cpp
index 523a0283a7..598e219fd9 100644
--- a/tests/validation/gpu/unit/Context.cpp
+++ b/tests/validation/gpu/unit/Context.cpp
@@ -21,7 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "tests/validation/fixtures/UNIT/Context.h"
+#include "tests/validation/fixtures/UNIT/ContextFixture.h"
 
 #include "src/gpu/cl/ClContext.h"
 
@@ -37,15 +37,9 @@ TEST_SUITE(CL)
 TEST_SUITE(UNIT)
 TEST_SUITE(Context)
 
-FIXTURE_TEST_CASE(SimpleContextCApi, SimpleContextCApiFixture<AclTarget::AclGpuOcl>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(SimpleContextCppApi, SimpleContextCppApiFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(MultipleContexts, MultipleContextsFixture<AclTarget::AclGpuOcl>, framework::DatasetMode::ALL)
-{
-}
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleContextCApi, SimpleContextCApiFixture<AclTarget::AclGpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleContextCppApi, SimpleContextCppApiFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MultipleContexts, MultipleContextsFixture<AclTarget::AclGpuOcl>, framework::DatasetMode::ALL)
 
 /** Test-case for MLGO kernel configuration file
  *
diff --git a/tests/validation/gpu/unit/Queue.cpp b/tests/validation/gpu/unit/Queue.cpp
new file mode 100644
index 0000000000..8154a7954f
--- /dev/null
+++ b/tests/validation/gpu/unit/Queue.cpp
@@ -0,0 +1,81 @@
+/*
+ * Copyright (c) 2021 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "tests/validation/fixtures/UNIT/QueueFixture.h"
+
+#include "arm_compute/AclOpenClExt.h"
+#include "src/gpu/cl/ClQueue.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(UNIT)
+TEST_SUITE(Queue)
+
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateQueueWithInvalidContext, CreateQueueWithInvalidContextFixture, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateQueuerWithInvalidOptions, CreateQueuerWithInvalidOptionsFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(DestroyInvalidQueue, DestroyInvalidQueueFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleQueue, SimpleQueueFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+
+TEST_CASE(KhrQueuePriorities, framework::DatasetMode::ALL)
+{
+    acl::StatusCode err = acl::StatusCode::Success;
+
+    acl::Context ctx(acl::Target::GpuOcl, &err);
+    ARM_COMPUTE_ASSERT(err == acl::StatusCode::Success);
+
+    acl::Queue queue(ctx, &err);
+    ARM_COMPUTE_ASSERT(err == acl::StatusCode::Success);
+
+    cl_device_id cl_dev;
+    auto         status = AclGetClDevice(ctx.get(), &cl_dev);
+    ARM_COMPUTE_ASSERT(status == AclSuccess);
+
+    std::string extensions = cl::Device(cl_dev).getInfo<CL_DEVICE_EXTENSIONS>();
+    if(extensions.find("cl_khr_priority_hints") != std::string::npos)
+    {
+        cl_int error = CL_SUCCESS;
+
+        cl_context cl_ctx;
+        auto       status = AclGetClContext(ctx.get(), &cl_ctx);
+        ARM_COMPUTE_ASSERT(status == AclSuccess);
+
+        /* Check a queue with high priority */
+        cl_queue_properties queue_properties[] = { CL_QUEUE_PRIORITY_KHR, CL_QUEUE_PRIORITY_HIGH_KHR, 0 };
+        cl_command_queue    priority_queue     = clCreateCommandQueueWithProperties(cl_ctx, cl_dev, queue_properties, &error);
+        ARM_COMPUTE_ASSERT(error == CL_SUCCESS);
+
+        clReleaseCommandQueue(priority_queue);
+    }
+}
+
+TEST_SUITE_END() // Queue
+TEST_SUITE_END() // UNIT
+TEST_SUITE_END() // CL
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/gpu/unit/Tensor.cpp b/tests/validation/gpu/unit/Tensor.cpp
index b40d6264f5..18102733ce 100644
--- a/tests/validation/gpu/unit/Tensor.cpp
+++ b/tests/validation/gpu/unit/Tensor.cpp
@@ -21,7 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "tests/validation/fixtures/UNIT/Tensor.h"
+#include "tests/validation/fixtures/UNIT/TensorFixture.h"
 
 namespace arm_compute
 {
@@ -33,24 +33,17 @@ TEST_SUITE(CL)
 TEST_SUITE(UNIT)
 TEST_SUITE(Tensor)
 
-#define TENSOR_TESE_CASE(name, fixture)                           \
-    FIXTURE_TEST_CASE(name, fixture, framework::DatasetMode::ALL) \
-    {                                                             \
-    }
-
-TENSOR_TESE_CASE(CreateTensorWithInvalidContext, CreateTensorWithInvalidContextFixture)
-TENSOR_TESE_CASE(CreateTensorWithInvalidDescriptor, CreateTensorWithInvalidDescriptorFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(DestroyInvalidTensor, DestroyInvalidTensorFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(SimpleTensor, SimpleTensorFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(TensorStress, TensorStressFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(MapInvalidTensor, MapInvalidTensorFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(MapAllocatedTensor, MapAllocatedTensorFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(GetSize, TensorSizeFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(GetInvalidSize, InvalidTensorSizeFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(GetDescriptor, DescriptorConversionFixture<acl::Target::GpuOcl>)
-TENSOR_TESE_CASE(GetInvalidDescriptor, InvalidDescriptorConversionFixture<acl::Target::GpuOcl>)
-
-#undef TENSOR_TEST_CASE
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateTensorWithInvalidContext, CreateTensorWithInvalidContextFixture, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateTensorWithInvalidDescriptor, CreateTensorWithInvalidDescriptorFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(DestroyInvalidTensor, DestroyInvalidTensorFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleTensor, SimpleTensorFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(TensorStress, TensorStressFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MapInvalidTensor, MapInvalidTensorFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MapAllocatedTensor, MapAllocatedTensorFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetSize, TensorSizeFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetInvalidSize, InvalidTensorSizeFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetDescriptor, DescriptorConversionFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(GetInvalidDescriptor, InvalidDescriptorConversionFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
 
 TEST_SUITE_END() // Tensor
 TEST_SUITE_END() // UNIT
diff --git a/tests/validation/gpu/unit/TensorPack.cpp b/tests/validation/gpu/unit/TensorPack.cpp
index b057db44ae..b62426d056 100644
--- a/tests/validation/gpu/unit/TensorPack.cpp
+++ b/tests/validation/gpu/unit/TensorPack.cpp
@@ -21,7 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "tests/validation/fixtures/UNIT/TensorPack.h"
+#include "tests/validation/fixtures/UNIT/TensorPackFixture.h"
 
 namespace arm_compute
 {
@@ -33,21 +33,11 @@ TEST_SUITE(CL)
 TEST_SUITE(UNIT)
 TEST_SUITE(TensorPack)
 
-FIXTURE_TEST_CASE(CreateTensorPackWithInvalidContext, CreateTensorPackWithInvalidContextFixture, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(DestroyInvalidTensorPack, DestroyInvalidTensorPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(AddInvalidObjectToTensorPack, AddInvalidObjectToTensorPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(SimpleTensorPack, SimpleTensorPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
-{
-}
-FIXTURE_TEST_CASE(MultipleTensorsInPack, MultipleTensorsInPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
-{
-}
+EMPTY_BODY_FIXTURE_TEST_CASE(CreateTensorPackWithInvalidContext, CreateTensorPackWithInvalidContextFixture, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(DestroyInvalidTensorPack, DestroyInvalidTensorPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(AddInvalidObjectToTensorPack, AddInvalidObjectToTensorPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(SimpleTensorPack, SimpleTensorPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
+EMPTY_BODY_FIXTURE_TEST_CASE(MultipleTensorsInPack, MultipleTensorsInPackFixture<acl::Target::GpuOcl>, framework::DatasetMode::ALL)
 
 TEST_SUITE_END() // Tensor
 TEST_SUITE_END() // UNIT