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 /arm_compute/core/CL/CLKernelLibrary.h | |
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 'arm_compute/core/CL/CLKernelLibrary.h')
-rw-r--r-- | arm_compute/core/CL/CLKernelLibrary.h | 25 |
1 files changed, 21 insertions, 4 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<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. >*/ + size_t _max_workgroup_size; /** Maximum local workgroup size supported on the device */ }; } #endif /* __ARM_COMPUTE_CLKERNELLIBRARY_H__ */ |