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>

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<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__ */
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 <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)
 {
 }