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>

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__ */