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>

24 files changed, 636 insertions, 4 deletions

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 */