From 5f9107280d098de719782530e19663cd655c9a71 Mon Sep 17 00:00:00 2001 From: steniu01 Date: Wed, 23 Aug 2017 10:15:22 +0100 Subject: COMPMID-513 Choose maximum local workgroup size at run time Change-Id: I9ab3cf6dc92a93b0ae5f746e078355e443b3a545 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/84906 Tested-by: Kaizen Reviewed-by: Anthony Barbier --- arm_compute/core/CL/CLKernelLibrary.h | 25 +++++++++++++++++++++---- arm_compute/core/CL/ICLKernel.h | 3 ++- src/core/CL/CLKernelLibrary.cpp | 20 +++++++++++++++++++- src/core/CL/ICLKernel.cpp | 2 +- 4 files changed, 43 insertions(+), 7 deletions(-) diff --git a/arm_compute/core/CL/CLKernelLibrary.h b/arm_compute/core/CL/CLKernelLibrary.h index 38cfad6c1d..4d4565d6fd 100644 --- a/arm_compute/core/CL/CLKernelLibrary.h +++ b/arm_compute/core/CL/CLKernelLibrary.h @@ -162,9 +162,11 @@ public: */ void init(std::string kernel_path = ".", cl::Context context = cl::Context::getDefault(), cl::Device device = cl::Device::getDefault()) { - _kernel_path = std::move(kernel_path); - _context = std::move(context); - _device = std::move(device); + _kernel_path = std::move(kernel_path); + _context = std::move(context); + _device = std::move(device); + _max_workgroup_size = 0; + max_local_workgroup_size(); } /** Sets the path that the kernels reside in. * @@ -206,6 +208,9 @@ public: { _device = cl_devices[0]; } + + _max_workgroup_size = 0; + max_local_workgroup_size(); }; /** Sets the CL device for which the programs are created. * @@ -213,7 +218,9 @@ public: */ void set_device(cl::Device device) { - _device = std::move(device); + _device = std::move(device); + _max_workgroup_size = 0; + max_local_workgroup_size(); }; /** Creates a kernel from the kernel library. * @@ -231,6 +238,15 @@ public: * */ void load_binary(); + /** Find the maximum number of local work items in a workgroup can be supported by the device + * + */ + size_t max_local_workgroup_size(); + + /** Return the default NDRange that is suitable for the device. + * + */ + cl::NDRange default_ndrange(); private: /** Load program and its dependencies. @@ -254,6 +270,7 @@ private: static const std::map _kernel_program_map; /**< Map that associates kernel names with programs. */ static const std::map _program_source_map; /**< Contains sources for all programs. Used for compile-time kernel inclusion. >*/ + size_t _max_workgroup_size; /** Maximum local workgroup size supported on the device */ }; } #endif /* __ARM_COMPUTE_CLKERNELLIBRARY_H__ */ diff --git a/arm_compute/core/CL/ICLKernel.h b/arm_compute/core/CL/ICLKernel.h index d96b9c026e..cfbf760f1e 100644 --- a/arm_compute/core/CL/ICLKernel.h +++ b/arm_compute/core/CL/ICLKernel.h @@ -24,6 +24,7 @@ #ifndef __ARM_COMPUTE_ICLKERNEL_H__ #define __ARM_COMPUTE_ICLKERNEL_H__ +#include "arm_compute/core/CL/CLKernelLibrary.h" #include "arm_compute/core/CL/CLTypes.h" #include "arm_compute/core/CL/OpenCL.h" #include "arm_compute/core/IKernel.h" @@ -175,6 +176,6 @@ protected: * * @note If any dimension of the lws is greater than the global workgroup size then no lws will be passed. */ -void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint = cl::Range_128_1); +void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint = CLKernelLibrary::get().default_ndrange()); } #endif /*__ARM_COMPUTE_ICLKERNEL_H__ */ diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index b5b331b1c1..019f3ea132 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -27,6 +27,7 @@ #include "arm_compute/core/Error.h" #include "arm_compute/core/Utils.h" +#include #include #include #include @@ -524,7 +525,7 @@ const std::map CLKernelLibrary::_program_source_map = }; CLKernelLibrary::CLKernelLibrary() - : _context(), _device(), _kernel_path("."), _programs_map(), _built_programs_map() + : _context(), _device(), _kernel_path("."), _programs_map(), _built_programs_map(), _max_workgroup_size(0) { } @@ -663,3 +664,20 @@ std::string CLKernelLibrary::get_program_source(const std::string &program_name) return program_source_it->second; } + +size_t CLKernelLibrary::max_local_workgroup_size() +{ + if(_max_workgroup_size == 0) + { + size_t err = clGetDeviceInfo(_device.get(), CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &_max_workgroup_size, nullptr); + ARM_COMPUTE_ERROR_ON_MSG(err != 0, "clGetDeviceInfo failed to return valid information"); + ARM_COMPUTE_UNUSED(err); + } + + return _max_workgroup_size; +} + +cl::NDRange CLKernelLibrary::default_ndrange() +{ + return cl::NDRange(std::min(_max_workgroup_size, 128u), 1); +} diff --git a/src/core/CL/ICLKernel.cpp b/src/core/CL/ICLKernel.cpp index bace631549..12af8c68c1 100644 --- a/src/core/CL/ICLKernel.cpp +++ b/src/core/CL/ICLKernel.cpp @@ -60,7 +60,7 @@ void arm_compute::enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Wind } ICLKernel::ICLKernel() - : _kernel(nullptr), _lws_hint(cl::Range_128_1), _target(GPUTarget::MIDGARD) + : _kernel(nullptr), _lws_hint(CLKernelLibrary::get().default_ndrange()), _target(GPUTarget::MIDGARD) { } -- cgit v1.2.1