diff options
author | Georgios Pinitas <georgios.pinitas@arm.com> | 2021-03-18 10:59:40 +0000 |
---|---|---|
committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2021-04-28 17:05:40 +0000 |
commit | c3c352e60050f3deacad767e429a88dc24b31af0 (patch) | |
tree | ad30a0ba717a742caf5e4dcb9d89389cfdc134b0 /src | |
parent | e2535154fa34ac0290ec3daaa44545be0b2b4606 (diff) | |
download | ComputeLibrary-c3c352e60050f3deacad767e429a88dc24b31af0.tar.gz |
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>
Diffstat (limited to 'src')
-rw-r--r-- | src/c/AclQueue.cpp | 99 | ||||
-rw-r--r-- | src/c/cl/AclOpenClExt.cpp | 75 | ||||
-rw-r--r-- | src/common/IContext.h | 16 | ||||
-rw-r--r-- | src/common/IQueue.h | 100 | ||||
-rw-r--r-- | src/core/CL/OpenCL.cpp | 18 | ||||
-rw-r--r-- | src/core/cpu/kernels/CpuDirectConvolutionOutputStageKernel.cpp (renamed from src/core/cpu/kernels/CpuDirectConvolutionStageKernel.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/neon/fp16.cpp (renamed from src/core/cpu/kernels/activation/NEON/fp16.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/neon/fp32.cpp (renamed from src/core/cpu/kernels/activation/NEON/fp32.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/neon/qasymm8.cpp (renamed from src/core/cpu/kernels/activation/NEON/qasymm8.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/neon/qasymm8_signed.cpp (renamed from src/core/cpu/kernels/activation/NEON/qasymm8_signed.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/neon/qsymm16.cpp (renamed from src/core/cpu/kernels/activation/NEON/qsymm16.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/sve/fp16.cpp (renamed from src/core/cpu/kernels/activation/SVE/fp16.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/sve/fp32.cpp (renamed from src/core/cpu/kernels/activation/SVE/fp32.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/sve/qasymm8.cpp (renamed from src/core/cpu/kernels/activation/SVE/qasymm8.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/sve/qasymm8_signed.cpp (renamed from src/core/cpu/kernels/activation/SVE/qasymm8_signed.cpp) | 0 | ||||
-rw-r--r-- | src/core/cpu/kernels/activation/sve/qsymm16.cpp (renamed from src/core/cpu/kernels/activation/SVE/qsymm16.cpp) | 0 | ||||
-rw-r--r-- | src/cpu/CpuContext.cpp | 6 | ||||
-rw-r--r-- | src/cpu/CpuContext.h | 1 | ||||
-rw-r--r-- | src/cpu/CpuQueue.cpp | 48 | ||||
-rw-r--r-- | src/cpu/CpuQueue.h | 56 | ||||
-rw-r--r-- | src/gpu/cl/ClContext.cpp | 23 | ||||
-rw-r--r-- | src/gpu/cl/ClContext.h | 12 | ||||
-rw-r--r-- | src/gpu/cl/ClQueue.cpp | 102 | ||||
-rw-r--r-- | src/gpu/cl/ClQueue.h | 84 |
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 */ |