diff options
| author | steniu01 <steven.niu@arm.com> | 2017-08-23 10:15:22 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | 5f9107280d098de719782530e19663cd655c9a71 (patch) | |
| tree | c2eaff4fb8087ea055e8883b720a3f27e6d7a0ce /src/core/CL | |
| parent | 907dba8d4c26c525e362be8dee6f18461f591262 (diff) | |
| download | ComputeLibrary-5f9107280d098de719782530e19663cd655c9a71.tar.gz | |
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 <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core/CL')
| -rw-r--r-- | src/core/CL/CLKernelLibrary.cpp | 20 | ||||
| -rw-r--r-- | src/core/CL/ICLKernel.cpp | 2 |
2 files changed, 20 insertions, 2 deletions
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 <algorithm> #include <fstream> #include <iostream> #include <utility> @@ -524,7 +525,7 @@ const std::map<std::string, std::string> 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<size_t>(_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) { } |