diff options
Diffstat (limited to 'arm_compute/core/CL')
| -rw-r--r-- | arm_compute/core/CL/CLCompileContext.h | 33 | ||||
| -rw-r--r-- | arm_compute/core/CL/CLCoreRuntimeContext.h | 75 | ||||
| -rw-r--r-- | arm_compute/core/CL/CLDevice.h | 39 | ||||
| -rw-r--r-- | arm_compute/core/CL/CLHelpers.h | 80 | ||||
| -rw-r--r-- | arm_compute/core/CL/CLKernelLibrary.h | 7 | ||||
| -rw-r--r-- | arm_compute/core/CL/CLTypes.h | 51 | ||||
| -rw-r--r-- | arm_compute/core/CL/ICLArray.h | 17 | ||||
| -rw-r--r-- | arm_compute/core/CL/ICLDistribution1D.h | 102 | ||||
| -rw-r--r-- | arm_compute/core/CL/ICLHOG.h | 113 | ||||
| -rw-r--r-- | arm_compute/core/CL/ICLLut.h | 94 | ||||
| -rw-r--r-- | arm_compute/core/CL/ICLMultiHOG.h | 56 | ||||
| -rw-r--r-- | arm_compute/core/CL/ICLMultiImage.h | 59 | ||||
| -rw-r--r-- | arm_compute/core/CL/ICLTensor.h | 11 | ||||
| -rw-r--r-- | arm_compute/core/CL/OpenCL.h | 49 |
14 files changed, 168 insertions, 618 deletions
diff --git a/arm_compute/core/CL/CLCompileContext.h b/arm_compute/core/CL/CLCompileContext.h index 46a8c9b341..dcd3b45670 100644 --- a/arm_compute/core/CL/CLCompileContext.h +++ b/arm_compute/core/CL/CLCompileContext.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021 Arm Limited. + * Copyright (c) 2020-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -77,6 +77,8 @@ public: */ const StringSet &options() const; + bool operator==(const CLBuildOptions &other) const; + private: StringSet _build_opts; /**< Build options set */ }; @@ -248,8 +250,12 @@ public: * * @return The created kernel. */ - Kernel create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source, - const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const; + Kernel create_kernel(const std::string &kernel_name, + const std::string &program_name, + const std::string &program_source, + const std::string &kernel_path, + const StringSet &build_options_set, + bool is_binary) const; /** Clear the library's cache of binary programs */ @@ -302,6 +308,18 @@ public: */ bool is_wbsm_supported() const; + /** Return the DDK version. If the DDK version cannot be detected, return -1. + * + * @return The DDK version. + */ + int32_t get_ddk_version() const; + + /** Return the Gpu target of the associated device + * + * @return GPUTarget + */ + GPUTarget get_gpu_target() const; + private: /** Load program and its dependencies. * @@ -309,7 +327,8 @@ private: * @param[in] program_source Source of the program. * @param[in] is_binary Flag to indicate if the program source is binary. */ - const Program &load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const; + const Program & + load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const; /** Generates the build options given a string of user defined ones * @@ -329,11 +348,11 @@ private: */ std::string stringify_set(const StringSet &s, const std::string &kernel_path) const; - cl::Context _context; /**< Underlying CL context. */ - CLDevice _device; /**< Underlying CL device. */ + cl::Context _context; /**< Underlying CL context. */ + CLDevice _device; /**< Underlying CL device. */ mutable std::map<std::string, const Program> _programs_map; /**< Map with all already loaded program data. */ mutable std::map<std::string, cl::Program> _built_programs_map; /**< Map with all already built program data. */ - bool _is_wbsm_supported; /**< Support of worksize batch size modifier support boolean*/ + bool _is_wbsm_supported; /**< Support of worksize batch size modifier support boolean*/ }; } // namespace arm_compute #endif /* ARM_COMPUTE_CLCOMPILECONTEXT_H */ diff --git a/arm_compute/core/CL/CLCoreRuntimeContext.h b/arm_compute/core/CL/CLCoreRuntimeContext.h deleted file mode 100644 index 23f282354c..0000000000 --- a/arm_compute/core/CL/CLCoreRuntimeContext.h +++ /dev/null @@ -1,75 +0,0 @@ -/* - * Copyright (c) 2019 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_CLCORERUNTIME_CONTEXT_H -#define ARM_COMPUTE_CLCORERUNTIME_CONTEXT_H - -#include "arm_compute/core/CL/OpenCL.h" - -namespace arm_compute -{ -// Forward declarations -class CLKernelLibrary; - -/** Core runtime context for OpenCL */ -class CLCoreRuntimeContext final -{ -public: - /** Legacy constructor */ - CLCoreRuntimeContext(); - - /** Constructor */ - CLCoreRuntimeContext(CLKernelLibrary *kernel_lib, cl::Context ctx, cl::CommandQueue queue); - /** Destructor */ - ~CLCoreRuntimeContext() = default; - /** Default copy constructor */ - CLCoreRuntimeContext(const CLCoreRuntimeContext &) = default; - /** Default move constructor */ - CLCoreRuntimeContext(CLCoreRuntimeContext &&) = default; - /** Default copy assignment */ - CLCoreRuntimeContext &operator=(const CLCoreRuntimeContext &) = default; - /** Default move assignment operator */ - CLCoreRuntimeContext &operator=(CLCoreRuntimeContext &&) = default; - /** Kernel Library accessor - * - * @return The kernel library instance used by the core context - */ - CLKernelLibrary *kernel_library() const; - /** OpenCL context accessor - * - * @return The OpenCL context used by the core context - */ - cl::Context context(); - /** OpenCL command queue accessor - * - * @return The OpenCL queue used by the core context - */ - cl::CommandQueue queue(); - -private: - CLKernelLibrary *_kernel_lib{ nullptr }; - cl::Context _ctx{}; - cl::CommandQueue _queue{}; -}; -} // namespace arm_compute -#endif /*ARM_COMPUTE_CLCORERUNTIME_CONTEXT_H */ diff --git a/arm_compute/core/CL/CLDevice.h b/arm_compute/core/CL/CLDevice.h index 033bf8fa96..ded6bb8493 100644 --- a/arm_compute/core/CL/CLDevice.h +++ b/arm_compute/core/CL/CLDevice.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020 Arm Limited. + * Copyright (c) 2020-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -30,6 +30,7 @@ #include "arm_compute/core/IDevice.h" #include <set> +#include <sstream> #include <string> namespace arm_compute @@ -43,8 +44,7 @@ class CLDevice : public IDevice { public: /** Default Constructor */ - CLDevice() - : _device(cl::Device()), _options() + CLDevice() : _device(cl::Device()), _options() { } @@ -52,8 +52,7 @@ public: * * @param[in] cl_device OpenCL device */ - CLDevice(const cl::Device &cl_device) - : _device(), _options() + CLDevice(const cl::Device &cl_device) : _device(), _options() { _device = cl_device; @@ -65,13 +64,13 @@ public: std::string extensions = _device.getInfo<CL_DEVICE_EXTENSIONS>(); std::istringstream iss(extensions); - for(std::string s; iss >> s;) + for (std::string s; iss >> s;) { _options.extensions.insert(s); } // SW workaround for G76 - if(_options.gpu_target == GPUTarget::G76) + if (_options.gpu_target == GPUTarget::G76) { _options.extensions.insert("cl_arm_integer_dot_product_int8"); } @@ -142,6 +141,32 @@ public: return _options.extensions.count(extension) != 0; } + /** Returns whether non-uniform workgroup is supported and the build options. + * + * If the feature is supported, the appropriate build options will be + * appended to the specified string. + * + * @return A tuple (supported, build_options) indicating whether the feature + * is supported and the corresponding build options to enable it. + */ + std::tuple<bool, std::string> is_non_uniform_workgroup_supported() const + { + if (version() == CLVersion::CL30 && get_cl_non_uniform_work_group_supported(_device)) + { + return {true, " -cl-std=CL3.0 "}; + } + else if (version() == CLVersion::CL20) + { + return {true, " -cl-std=CL2.0 "}; + } + else if (supported("cl_arm_non_uniform_work_group_size")) + { + return {true, " -cl-arm-non-uniform-work-group-size "}; + } + + return {false, ""}; + } + private: cl::Device _device; /**< OpenCL device. */ struct CLDeviceOptions _options; /**< OpenCL device options */ diff --git a/arm_compute/core/CL/CLHelpers.h b/arm_compute/core/CL/CLHelpers.h index 0e9aa5d6e5..1a639e47f9 100644 --- a/arm_compute/core/CL/CLHelpers.h +++ b/arm_compute/core/CL/CLHelpers.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2021 Arm Limited. + * Copyright (c) 2016-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -26,18 +26,13 @@ #include "arm_compute/core/CL/CLTypes.h" #include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Types.h" #include <set> #include <string> -/* CL Device capabilities */ -#define ARM_COMPUTE_LIBRARY_OPENCL_DEVICE_CAPABILITIES_ARM 0x41E4 -/* Workgroup Batch Size Modifier */ -#define ARM_COMPUTE_LIBRARY_OPENCL_EXEC_WBSM_ARM 0x41E6 - namespace arm_compute { -class CLCoreRuntimeContext; class CLCompileContext; class CLBuildOptions; @@ -46,6 +41,9 @@ enum class DataType; /** Max vector width of an OpenCL vector */ static constexpr unsigned int max_cl_vector_width = 16; +/** Max number of manual loop unrolling */ +static constexpr int max_manual_loop_unrolling = 128; + /** Translates a tensor data type to the appropriate OpenCL type. * * @param[in] dt @ref DataType to be translated to OpenCL type. @@ -126,6 +124,14 @@ CLVersion get_cl_version(const cl::Device &device); */ size_t get_cl_image_pitch_alignment(const cl::Device &device); +/** Helper function to check whether non-uniform work group is supported + * + * @param[in] device A CL device + * + * @return True if the feature is supported + */ +bool get_cl_non_uniform_work_group_supported(const cl::Device &device); + /** Helper function to check whether a given extension is supported * * @param[in] device A CL device @@ -173,7 +179,9 @@ bool dot8_acc_supported(const cl::Device &device); * * @return True if the configuration is supported */ -bool cl_winograd_convolution_layer_supported(const Size2D &output_tile, const Size2D &kernel_size, DataLayout data_layout); +bool cl_winograd_convolution_layer_supported(const Size2D &output_tile, + const Size2D &kernel_size, + DataLayout data_layout); /** Helper function to get the preferred native vector width size for built-in scalar types that can be put into vectors * @@ -201,16 +209,6 @@ bool preferred_dummy_work_items_support(const cl::Device &device); */ bool image2d_from_buffer_supported(const cl::Device &device); -/** Creates an opencl kernel - * - * @param[in] ctx A context to be used to create the opencl kernel. - * @param[in] kernel_name The kernel name. - * @param[in] build_opts The build options to be used for the opencl kernel compilation. - * - * @return An opencl kernel - */ -cl::Kernel create_opencl_kernel(CLCoreRuntimeContext *ctx, const std::string &kernel_name, const CLBuildOptions &build_opts); - /** Creates an opencl kernel using a compile context * * @param[in] ctx A compile context to be used to create the opencl kernel. @@ -219,7 +217,9 @@ cl::Kernel create_opencl_kernel(CLCoreRuntimeContext *ctx, const std::string &ke * * @return An opencl kernel */ -cl::Kernel create_kernel(const CLCompileContext &ctx, const std::string &kernel_name, const std::set<std::string> &build_opts = std::set<std::string>()); +cl::Kernel create_kernel(const CLCompileContext &ctx, + const std::string &kernel_name, + const std::set<std::string> &build_opts = std::set<std::string>()); /** Creates a suitable LWS hint object for parallel implementations. Sets the number of WG based on the input size. * If input width is smaller than 128 we can use fewer threads than 8. @@ -246,5 +246,47 @@ bool get_wbsm_support_info(const cl::Device &device); */ void set_wbsm(cl::Kernel &kernel, cl_int wbsm_hint); +/* Helper function to check if we can export the tensor to cl_image + * + * @param[in] input tensor + * + * @return true if we can export the tensor to cl_image + */ +bool export_to_cl_image(const ITensorInfo *tensor); + +/* Helper function to force unroll with pragma when any of the input values (iterations) are greater than @ref max_manual_loop_unrolling + * + * This function passes UNROLL_WITH_PRAGMA at compile time when any of the input values are greater than @ref max_manual_loop_unrolling + * + * @param[in] built_opts OpenCL kernel build options + * @param[in] values Input values (iterations) + * + */ +void set_unroll_with_pragma(CLBuildOptions &built_opts, std::initializer_list<int> values); + +/** Helper function to check whether the cl_arm_matrix_multiply extension is supported + * + * @param[in] device A CL device + * + * @return True if the extension is supported + */ +bool arm_matrix_multiply_supported(const cl::Device &device); + +/** Check whether cl_khr_command_buffer extension is supported by the specified CL device. + * + * @param[in] device The CL device + * + * @return True if the extension is supported by the CL device. + */ +bool command_buffer_supported(const cl::Device &device); + +/** Check whether cl_khr_command_buffer_mutable_dispatch extension is supported by the specified CL device. + * + * @param[in] device The CL device + * + * @return True if the extension is supported by the CL device. + */ +bool command_buffer_mutable_dispatch_supported(const cl::Device &device); + } // namespace arm_compute #endif /* ARM_COMPUTE_CLHELPERS_H */ diff --git a/arm_compute/core/CL/CLKernelLibrary.h b/arm_compute/core/CL/CLKernelLibrary.h index 24eef639ae..527733ccf1 100644 --- a/arm_compute/core/CL/CLKernelLibrary.h +++ b/arm_compute/core/CL/CLKernelLibrary.h @@ -170,12 +170,7 @@ public: CLCompileContext &get_compile_context(); private: - CLCompileContext _compile_context; /**< Compile Context. */ - std::string _kernel_path; /**< Path to the kernels folder. */ - mutable std::map<std::string, std::string> _decompressed_source_map; /**< Map holding the decompressed files when compression is used */ - static const std::map<std::string, std::string> _kernel_program_map; /**< Map that associates kernel names with programs. */ - static const std::map<std::string, std::string> _program_source_map; /**< Contains sources for all programs. - Used for compile-time kernel inclusion. >*/ + CLCompileContext _compile_context; /**< Compile Context. */ }; } // namespace arm_compute #endif /* ARM_COMPUTE_CLKERNELLIBRARY_H */ diff --git a/arm_compute/core/CL/CLTypes.h b/arm_compute/core/CL/CLTypes.h index 0f6eb0dfa4..0f088e2b10 100644 --- a/arm_compute/core/CL/CLTypes.h +++ b/arm_compute/core/CL/CLTypes.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2020 Arm Limited. + * Copyright (c) 2017-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -41,7 +41,8 @@ enum class CLVersion CL10, /* the OpenCL 1.0 */ CL11, /* the OpenCL 1.1 */ CL12, /* the OpenCL 1.2 */ - CL20, /* the OpenCL 2.0 and above */ + CL20, /* the OpenCL 2.x */ + CL30, /* the OpenCL 3.x */ UNKNOWN /* unkown version */ }; @@ -62,53 +63,27 @@ struct CLDeviceOptions struct CLQuantization { /** Default Constructor */ - CLQuantization() - : scale(nullptr), offset(nullptr) {}; + CLQuantization() : scale(nullptr), offset(nullptr){}; /** Constructor * * @param[in] scale OpenCL scale array * @param[in] offset OpenCL offset array */ - CLQuantization(const ICLFloatArray *scale, const ICLInt32Array *offset) - : scale(scale), offset(offset) {}; + CLQuantization(const ICLFloatArray *scale, const ICLInt32Array *offset) : scale(scale), offset(offset){}; const ICLFloatArray *scale; /**< Quantization scale array */ const ICLInt32Array *offset; /**< Quantization offset array */ }; -/** Internal keypoint structure for Lucas-Kanade Optical Flow */ -struct CLLKInternalKeypoint +enum CLKernelType { - float x{ 0.f }; /**< x coordinate of the keypoint */ - float y{ 0.f }; /**< y coordinate of the keypoint */ - float tracking_status{ 0.f }; /**< the tracking status of the keypoint */ - float dummy{ 0.f }; /**< Dummy field, to make sure the data structure 128-bit align, so that GPU can use vload4 */ + UNKNOWN, /**< Unknown CL kernel type */ + DEPTHWISE, /**< Depthwise CL kernel type */ + DIRECT, /**< Direct Convolution CL kernel type */ + ELEMENTWISE, /**< Elementwise CL kernel type */ + GEMM, /**< GEMM CL kernel type */ + POOL, /**< Pool CL kernel type */ + WINOGRAD /**< Winograd CL kernel type */ }; - -/** Structure for storing Spatial Gradient Matrix and the minimum eigenvalue for each keypoint */ -struct CLCoefficientTable -{ - float A11; /**< iA11 * FLT_SCALE */ - float A12; /**< iA11 * FLT_SCALE */ - float A22; /**< iA11 * FLT_SCALE */ - float min_eig; /**< Minimum eigenvalue */ -}; - -/** Structure for storing ival, ixval and iyval for each point inside the window */ -struct CLOldValue -{ - int16_t ival; /**< ival extracts from old image */ - int16_t ixval; /**< ixval extracts from scharr Gx image */ - int16_t iyval; /**< iyval extracts from scharr Gy image */ - int16_t dummy; /**< Dummy field, to make sure the data structure 128-bit align, so that GPU can use vload4 */ -}; - -/** Interface for OpenCL Array of Internal Key Points. */ -using ICLLKInternalKeypointArray = ICLArray<CLLKInternalKeypoint>; -/** Interface for OpenCL Array of Coefficient Tables. */ -using ICLCoefficientTableArray = ICLArray<CLCoefficientTable>; -/** Interface for OpenCL Array of Old Values. */ -using ICLOldValArray = ICLArray<CLOldValue>; - } // namespace arm_compute #endif /* ARM_COMPUTE_CL_TYPES_H */ diff --git a/arm_compute/core/CL/ICLArray.h b/arm_compute/core/CL/ICLArray.h index 2fa2f34c5d..a2b2baa5b3 100644 --- a/arm_compute/core/CL/ICLArray.h +++ b/arm_compute/core/CL/ICLArray.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 Arm Limited. + * Copyright (c) 2016-2021, 2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -40,8 +40,7 @@ public: * @param[in] max_num_values Maximum size of the array. * */ - explicit ICLArray(size_t max_num_values) - : IArray<T>(max_num_values), _mapping(nullptr) + explicit ICLArray(size_t max_num_values) : IArray<T>(max_num_values), _mapping(nullptr) { } @@ -66,8 +65,6 @@ public: * @param[in] blocking If true, then the mapping will be ready to use by the time * this method returns, else it is the caller's responsibility * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - * - * @return The mapping address. */ void map(cl::CommandQueue &q, bool blocking = true) { @@ -115,14 +112,6 @@ private: uint8_t *_mapping; }; -/** Interface for OpenCL Array of Key Points. */ -using ICLKeyPointArray = ICLArray<KeyPoint>; -/** Interface for OpenCL Array of 2D Coordinates. */ -using ICLCoordinates2DArray = ICLArray<Coordinates2D>; -/** Interface for OpenCL Array of Detection Windows. */ -using ICLDetectionWindowArray = ICLArray<DetectionWindow>; -/** Interface for OpenCL Array of 2D Sizes. */ -using ICLSize2DArray = ICLArray<Size2D>; /** Interface for OpenCL Array of uint8s. */ using ICLUInt8Array = ICLArray<cl_uchar>; /** Interface for OpenCL Array of uint16s. */ @@ -135,5 +124,5 @@ using ICLInt16Array = ICLArray<cl_short>; using ICLInt32Array = ICLArray<cl_int>; /** Interface for OpenCL Array of floats. */ using ICLFloatArray = ICLArray<cl_float>; -} +} // namespace arm_compute #endif /*ARM_COMPUTE_ICLARRAY_H*/ diff --git a/arm_compute/core/CL/ICLDistribution1D.h b/arm_compute/core/CL/ICLDistribution1D.h deleted file mode 100644 index 18afabd52e..0000000000 --- a/arm_compute/core/CL/ICLDistribution1D.h +++ /dev/null @@ -1,102 +0,0 @@ -/* - * Copyright (c) 2016-2019 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_ICLDISTRIBUTION1D_H -#define ARM_COMPUTE_ICLDISTRIBUTION1D_H - -#include "arm_compute/core/IDistribution1D.h" - -#include <cstddef> -#include <cstdint> - -namespace cl -{ -class Buffer; -class CommandQueue; -} - -namespace arm_compute -{ -/** ICLDistribution1D interface class */ -class ICLDistribution1D : public IDistribution1D -{ -public: - /** Constructor: Creates a 1D CLDistribution of a consecutive interval [offset, offset + range - 1] - * defined by a start offset and valid range, divided equally into num_bins parts. - * - * @param[in] num_bins The number of bins the distribution is divided in. - * @param[in] offset The start of the values to use. - * @param[in] range The total number of the consecutive values of the distribution interval. - */ - ICLDistribution1D(size_t num_bins, int32_t offset, uint32_t range); - /** Prevent instances of this class from being copied (As this class contains pointers) */ - ICLDistribution1D(const ICLDistribution1D &) = delete; - /** Prevent instances of this class from being copied (As this class contains pointers) */ - const ICLDistribution1D &operator=(const ICLDistribution1D &) = delete; - /** Enqueue a map operation of the allocated buffer on the given queue. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - * @param[in] blocking If true, then the mapping will be ready to use by the time - * this method returns, else it is the caller's responsibility - * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - */ - void map(cl::CommandQueue &q, bool blocking = true); - /** Enqueue an unmap operation of the allocated and mapped buffer on the given queue. - * - * @note This method simply enqueues the unmap operation, it is the caller's responsibility to flush the queue and make sure the unmap is finished before - * the memory is accessed by the device. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - */ - void unmap(cl::CommandQueue &q); - /** Interface to be implemented by the child class to return a reference to the OpenCL buffer containing the distribution's data. - * - * @return A reference to an OpenCL buffer containing the distribution's data. - */ - virtual cl::Buffer &cl_buffer() = 0; - // Inherited methods overridden: - uint32_t *buffer() const override; - -protected: - /** Method to be implemented by the child class to map the OpenCL buffer - * - * @param[in,out] q The CL command queue to use for the mapping operation. - * @param[in] blocking If true, then the mapping will be ready to use by the time - * this method returns, else it is the caller's responsibility - * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - */ - virtual uint32_t *do_map(cl::CommandQueue &q, bool blocking) = 0; - /** Method to be implemented by the child class to unmap the OpenCL buffer - * - * @note This method simply enqueues the unmap operation, it is the caller's responsibility to flush the queue and make sure the unmap is finished before - * the memory is accessed by the device. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - */ - virtual void do_unmap(cl::CommandQueue &q) = 0; - -protected: - uint32_t *_mapping; /**< The distribution data. */ -}; -} -#endif /* ARM_COMPUTE_ICLDISTRIBUTION1D_H */ diff --git a/arm_compute/core/CL/ICLHOG.h b/arm_compute/core/CL/ICLHOG.h deleted file mode 100644 index da3c0c65c8..0000000000 --- a/arm_compute/core/CL/ICLHOG.h +++ /dev/null @@ -1,113 +0,0 @@ -/* - * Copyright (c) 2017-2019 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_ICLHOG_H -#define ARM_COMPUTE_ICLHOG_H - -#include "arm_compute/core/IHOG.h" - -#include <cstdint> - -namespace cl -{ -class Buffer; -class CommandQueue; -} - -namespace arm_compute -{ -/** Interface for OpenCL HOG data-object */ -class ICLHOG : public IHOG -{ -public: - /** Default constructor */ - ICLHOG(); - /** Prevent instances of this class from being copied (As this class contains pointers) */ - ICLHOG(const ICLHOG &) = delete; - /** Prevent instances of this class from being copied (As this class contains pointers) */ - ICLHOG &operator=(const ICLHOG &) = delete; - /** Allow instances of this class to be moved */ - ICLHOG(ICLHOG &&) = default; - /** Allow instances of this class to be moved */ - ICLHOG &operator=(ICLHOG &&) = default; - /** Default destructor */ - virtual ~ICLHOG() = default; - - /** Interface to be implemented by the child class to return a reference to the OpenCL buffer containing the hog's descriptor - * - * @return A reference to an OpenCL buffer containing the hog's descriptor - */ - virtual const cl::Buffer &cl_buffer() const = 0; - - /** Enqueue a map operation of the allocated buffer on the given queue. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - * @param[in] blocking If true, then the mapping will be ready to use by the time - * this method returns, else it is the caller's responsibility - * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - * - * @return The mapping address. - */ - void map(cl::CommandQueue &q, bool blocking = true); - - /** Enqueue an unmap operation of the allocated and mapped buffer on the given queue. - * - * @note This method simply enqueues the unmap operation, it is the caller's responsibility to flush the queue and make sure the unmap is finished before - * the memory is accessed by the device. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - */ - void unmap(cl::CommandQueue &q); - - /** Interface to be implemented by the child class to free the allocated cl buffer. - * - * @warning The buffer must have been allocated previously. Otherwise calling the function will fail. - */ - virtual void free() = 0; - - // Inherited methods overridden: - float *descriptor() const override; - -protected: - /** Method to be implemented by the child class to map the OpenCL buffer - * - * @param[in,out] q The CL command queue to use for the mapping operation. - * @param[in] blocking If true, then the mapping will be ready to use by the time - * this method returns, else it is the caller's responsibility - * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - */ - virtual uint8_t *do_map(cl::CommandQueue &q, bool blocking) = 0; - /** Method to be implemented by the child class to unmap the OpenCL buffer - * - * @note This method simply enqueues the unmap operation, it is the caller's responsibility to flush the queue and make sure the unmap is finished before - * the memory is accessed by the device. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - */ - virtual void do_unmap(cl::CommandQueue &q) = 0; - -private: - uint8_t *_mapping; -}; -} -#endif /*ARM_COMPUTE_ICLHOG_H */ diff --git a/arm_compute/core/CL/ICLLut.h b/arm_compute/core/CL/ICLLut.h deleted file mode 100644 index b4d7471e1f..0000000000 --- a/arm_compute/core/CL/ICLLut.h +++ /dev/null @@ -1,94 +0,0 @@ -/* - * Copyright (c) 2016-2019 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_ICLLUT_H -#define ARM_COMPUTE_ICLLUT_H - -#include "arm_compute/core/ILut.h" - -#include <cstdint> - -namespace cl -{ -class Buffer; -class CommandQueue; -} - -namespace arm_compute -{ -/** Interface for OpenCL LUT */ -class ICLLut : public ILut -{ -public: - ICLLut(); - ICLLut(const ICLLut &) = delete; - ICLLut &operator=(const ICLLut &) = delete; - - /** Interface to be implemented by the child class to return a reference to the OpenCL buffer containing the lut's data. - * - * @return A reference to an OpenCL buffer containing the lut's data. - */ - virtual const cl::Buffer &cl_buffer() const = 0; - /** Enqueue a map operation of the allocated buffer on the given queue. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - * @param[in] blocking If true, then the mapping will be ready to use by the time - * this method returns, else it is the caller's responsibility - * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - */ - void map(cl::CommandQueue &q, bool blocking = true); - /** Enqueue an unmap operation of the allocated and mapped buffer on the given queue. - * - * @note This method simply enqueues the unmap operation, it is the caller's responsibility to flush the queue and make sure the unmap is finished before - * the memory is accessed by the device. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - */ - void unmap(cl::CommandQueue &q); - - // Inherited methods overridden: - uint8_t *buffer() const override; - -protected: - /** Method to be implemented by the child class to map the OpenCL buffer - * - * @param[in,out] q The CL command queue to use for the mapping operation. - * @param[in] blocking If true, then the mapping will be ready to use by the time - * this method returns, else it is the caller's responsibility - * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - */ - virtual uint8_t *do_map(cl::CommandQueue &q, bool blocking) = 0; - /** Method to be implemented by the child class to unmap the OpenCL buffer - * - * @note This method simply enqueues the unmap operation, it is the caller's responsibility to flush the queue and make sure the unmap is finished before - * the memory is accessed by the device. - * - * @param[in,out] q The CL command queue to use for the mapping operation. - */ - virtual void do_unmap(cl::CommandQueue &q) = 0; - -private: - uint8_t *_mapping; -}; -} -#endif /*ARM_COMPUTE_ICLLUT_H */ diff --git a/arm_compute/core/CL/ICLMultiHOG.h b/arm_compute/core/CL/ICLMultiHOG.h deleted file mode 100644 index 109b4d480d..0000000000 --- a/arm_compute/core/CL/ICLMultiHOG.h +++ /dev/null @@ -1,56 +0,0 @@ -/* - * Copyright (c) 2017-2019 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_ICLMULTIHOG_H -#define ARM_COMPUTE_ICLMULTIHOG_H - -#include "arm_compute/core/CL/ICLHOG.h" -#include "arm_compute/core/IMultiHOG.h" - -namespace arm_compute -{ -/** Interface for storing multiple HOG data-objects */ -class ICLMultiHOG : public IMultiHOG -{ -public: - /** Return a pointer to the requested OpenCL HOG model - * - * @param[in] index The index of the wanted OpenCL HOG model. - * - * @return A pointer pointed to the HOG model - */ - virtual ICLHOG *cl_model(size_t index) = 0; - /** Return a constant pointer to the requested OpenCL HOG model - * - * @param[in] index The index of the wanted OpenCL HOG model. - * - * @return A constant pointer pointed to the OpenCL HOG model - */ - virtual const ICLHOG *cl_model(size_t index) const = 0; - - // Inherited methods overridden: - IHOG *model(size_t index) override; - const IHOG *model(size_t index) const override; -}; -} -#endif /*ARM_COMPUTE_ICLMULTIHOG_H */ diff --git a/arm_compute/core/CL/ICLMultiImage.h b/arm_compute/core/CL/ICLMultiImage.h deleted file mode 100644 index 23ed04a484..0000000000 --- a/arm_compute/core/CL/ICLMultiImage.h +++ /dev/null @@ -1,59 +0,0 @@ -/* - * Copyright (c) 2016-2019 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_ICLMULTIIMAGE_H -#define ARM_COMPUTE_ICLMULTIIMAGE_H - -#include "arm_compute/core/IMultiImage.h" - -namespace arm_compute -{ -class ICLTensor; -/** Interface for OpenCL images */ -using ICLImage = ICLTensor; - -/** Interface for OpenCL multi-planar images */ -class ICLMultiImage : public IMultiImage -{ -public: - /** Return a pointer to the requested OpenCL plane of the image. - * - * @param[in] index The index of the wanted planed. - * - * @return A pointer pointed to the OpenCL plane - */ - virtual ICLImage *cl_plane(unsigned int index) = 0; - /** Return a constant pointer to the requested OpenCL plane of the image. - * - * @param[in] index The index of the wanted planed. - * - * @return A constant pointer pointed to the OpenCL plane - */ - virtual const ICLImage *cl_plane(unsigned int index) const = 0; - - // Inherited methods overridden: - IImage *plane(unsigned int index) override; - const IImage *plane(unsigned int index) const override; -}; -} -#endif /*ARM_COMPUTE_ICLMULTIIMAGE_H */ diff --git a/arm_compute/core/CL/ICLTensor.h b/arm_compute/core/CL/ICLTensor.h index fd05e64732..8de5423762 100644 --- a/arm_compute/core/CL/ICLTensor.h +++ b/arm_compute/core/CL/ICLTensor.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 Arm Limited. + * Copyright (c) 2016-2019, 2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -24,9 +24,8 @@ #ifndef ARM_COMPUTE_ICLTENSOR_H #define ARM_COMPUTE_ICLTENSOR_H -#include "arm_compute/core/ITensor.h" - #include "arm_compute/core/CL/CLTypes.h" +#include "arm_compute/core/ITensor.h" #include <cstdint> @@ -34,7 +33,7 @@ namespace cl { class Buffer; class CommandQueue; -} +} // namespace cl namespace arm_compute { @@ -71,8 +70,6 @@ public: * @param[in] blocking If true, then the mapping will be ready to use by the time * this method returns, else it is the caller's responsibility * to flush the queue and wait for the mapping operation to have completed before using the returned mapping pointer. - * - * @return The mapping address. */ void map(cl::CommandQueue &q, bool blocking = true); /** Enqueue an unmap operation of the allocated and mapped buffer on the given queue. @@ -115,5 +112,5 @@ private: }; using ICLImage = ICLTensor; -} +} // namespace arm_compute #endif /*ARM_COMPUTE_ICLTENSOR_H */ diff --git a/arm_compute/core/CL/OpenCL.h b/arm_compute/core/CL/OpenCL.h index 64b24dba8f..8b5bf97099 100644 --- a/arm_compute/core/CL/OpenCL.h +++ b/arm_compute/core/CL/OpenCL.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2021 Arm Limited. + * Copyright (c) 2016-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -21,8 +21,8 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#ifndef ARM_COMPUTE_OPENCL_H -#define ARM_COMPUTE_OPENCL_H +#ifndef ACL_ARM_COMPUTE_CORE_CL_OPENCL_H +#define ACL_ARM_COMPUTE_CORE_CL_OPENCL_H #include <string> #include <utility> @@ -31,8 +31,8 @@ #ifndef ARM_COMPUTE_NO_EXCEPTIONS #define CL_HPP_ENABLE_EXCEPTIONS #endif // ARM_COMPUTE_NO_EXCEPTIONS -#define CL_TARGET_OPENCL_VERSION 200 -#define CL_HPP_TARGET_OPENCL_VERSION 110 +#define CL_TARGET_OPENCL_VERSION 300 +#define CL_HPP_TARGET_OPENCL_VERSION 110 #define CL_HPP_MINIMUM_OPENCL_VERSION 110 #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Weffc++" @@ -40,8 +40,8 @@ #pragma GCC diagnostic ignored "-Wunused-parameter" #if defined(__GNUG__) && __GNUG__ >= 8 #pragma GCC diagnostic ignored "-Wcatch-value" -#endif // defined(__GNUG__) && __GNUG__ >= 8 -#include <CL/cl2.hpp> +#endif // defined(__GNUG__) && __GNUG__ >= 8 +#include <CL/opencl.hpp> // include new hpp header instead of cl2.hpp #pragma GCC diagnostic pop namespace cl @@ -63,12 +63,6 @@ class CLSymbols final public: /** Default Constructor */ CLSymbols() noexcept(false); - /** Destructor */ - ~CLSymbols(); - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLSymbols(const CLSymbols &) = delete; - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLSymbols &operator=(const CLSymbols &) = delete; /** Load OpenCL symbols from handle * * @param[in] handle Handle to load symbols from @@ -79,25 +73,27 @@ public: * @return The static instance of CLSymbols. */ static CLSymbols &get(); - /** Load symbols from the given OpenCL library path. + /** This method attempts to load the OpenCL symbols from the first available library from the provided OpenCL libraries. * - * @param[in] library Path to the OpenCL library. + * @param[in] libraries_filenames Vector containing the filenames of the libraries to be loaded. + * @param[in] use_loader Use symbol loader function loadOpenCLPointer. * - * @return True if loading the library is successful. + * @return True if loading the library is successful. False if all the provided libraries could not be loaded. */ - bool load(const std::string &library); + bool load(const std::vector<std::string> &libraries_filenames, bool use_loader = false); /** Load symbols from any of the default OpenCL library names. + * If all the default libraries could not be loaded, this method will print a warning message and return false. * * @return True if loading any library is successful. */ bool load_default(); -#define DECLARE_FUNCTION_PTR(func_name) \ - std::function<decltype(func_name)> func_name##_ptr = nullptr +#define DECLARE_FUNCTION_PTR(func_name) std::function<decltype(func_name)> func_name##_ptr = nullptr 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); @@ -129,6 +125,7 @@ public: DECLARE_FUNCTION_PTR(clGetDeviceIDs); DECLARE_FUNCTION_PTR(clGetMemObjectInfo); DECLARE_FUNCTION_PTR(clRetainEvent); + DECLARE_FUNCTION_PTR(clGetPlatformInfo); DECLARE_FUNCTION_PTR(clGetPlatformIDs); DECLARE_FUNCTION_PTR(clGetKernelWorkGroupInfo); DECLARE_FUNCTION_PTR(clGetCommandQueueInfo); @@ -142,6 +139,17 @@ public: DECLARE_FUNCTION_PTR(clWaitForEvents); DECLARE_FUNCTION_PTR(clCreateImage); DECLARE_FUNCTION_PTR(clSetKernelExecInfo); + DECLARE_FUNCTION_PTR(clGetExtensionFunctionAddressForPlatform); + + // Command buffer and mutable dispatch command buffer extensions + DECLARE_FUNCTION_PTR(clCreateCommandBufferKHR); + DECLARE_FUNCTION_PTR(clRetainCommandBufferKHR); + DECLARE_FUNCTION_PTR(clReleaseCommandBufferKHR); + DECLARE_FUNCTION_PTR(clFinalizeCommandBufferKHR); + DECLARE_FUNCTION_PTR(clEnqueueCommandBufferKHR); + DECLARE_FUNCTION_PTR(clCommandNDRangeKernelKHR); + + DECLARE_FUNCTION_PTR(clUpdateMutableCommandsKHR); // Third-party extensions DECLARE_FUNCTION_PTR(clImportMemoryARM); @@ -150,7 +158,6 @@ public: private: std::pair<bool, bool> _loaded; - void *_handle; }; } // namespace arm_compute -#endif /* ARM_COMPUTE_OPENCL_H */ +#endif // ACL_ARM_COMPUTE_CORE_CL_OPENCL_H |