3 files changed, 53 insertions, 14 deletions

diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index e165cf3350..6e5e802538 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -569,7 +569,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map =
 };
 
 CLKernelLibrary::CLKernelLibrary()
-    : _context(), _device(), _kernel_path("."), _programs_map(), _built_programs_map(), _max_workgroup_size(0)
+    : _context(), _device(), _kernel_path("."), _programs_map(), _built_programs_map()
 {
 }
 
@@ -709,19 +709,18 @@ std::string CLKernelLibrary::get_program_source(const std::string &program_name)
     return program_source_it->second;
 }
 
-size_t CLKernelLibrary::max_local_workgroup_size()
+size_t CLKernelLibrary::max_local_workgroup_size(const cl::Kernel &kernel) const
 {
-    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);
-    }
+    size_t result;
+
+    size_t err = kernel.getWorkGroupInfo(_device, CL_KERNEL_WORK_GROUP_SIZE, &result);
+    ARM_COMPUTE_ERROR_ON_MSG(err != 0, "clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
+    ARM_COMPUTE_UNUSED(err);
 
-    return _max_workgroup_size;
+    return result;
 }
 
-cl::NDRange CLKernelLibrary::default_ndrange()
+cl::NDRange CLKernelLibrary::default_ndrange() const
 {
-    return cl::NDRange(std::min<size_t>(_max_workgroup_size, 128u), 1);
+    return cl::NDRange(128u, 1);
 }
diff --git a/src/core/CL/ICLKernel.cpp b/src/core/CL/ICLKernel.cpp
index 1e04f00343..17b58b727f 100644
--- a/src/core/CL/ICLKernel.cpp
+++ b/src/core/CL/ICLKernel.cpp
@@ -52,18 +52,28 @@ void arm_compute::enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Wind
                     (window.y().end() - window.y().start()) / window.y().step(),
                     (window.z().end() - window.z().start()) / window.z().step());
 
+    cl::NDRange valid_lws;
+    if(lws_hint[0] * lws_hint[1] * lws_hint[2] > kernel.get_max_workgroup_size())
+    {
+        valid_lws = cl::NullRange;
+    }
+    else
+    {
+        valid_lws = lws_hint;
+    }
+
     cl::NDRange lws = cl::NullRange;
 
-    if((lws_hint[0] <= gws[0]) && (lws_hint[1] <= gws[1]) && (lws_hint[2] <= gws[2]))
+    if((valid_lws[0] <= gws[0]) && (valid_lws[1] <= gws[1]) && (valid_lws[2] <= gws[2]))
     {
-        lws = lws_hint;
+        lws = valid_lws;
     }
 
     queue.enqueueNDRangeKernel(kernel.kernel(), cl::NullRange, gws, lws);
 }
 
 ICLKernel::ICLKernel()
-    : _kernel(nullptr), _lws_hint(CLKernelLibrary::get().default_ndrange()), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id)
+    : _kernel(nullptr), _lws_hint(CLKernelLibrary::get().default_ndrange()), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id), _max_workgroup_size(0)
 {
 }
 
@@ -163,3 +173,12 @@ GPUTarget ICLKernel::get_target() const
 {
     return _target;
 }
+
+size_t ICLKernel::get_max_workgroup_size()
+{
+    if(_max_workgroup_size == 0)
+    {
+        _max_workgroup_size = CLKernelLibrary::get().max_local_workgroup_size(_kernel);
+    }
+    return _max_workgroup_size;
+}
diff --git a/src/core/CL/OpenCL.cpp b/src/core/CL/OpenCL.cpp
index cc2391977b..287c5e2f5a 100644
--- a/src/core/CL/OpenCL.cpp
+++ b/src/core/CL/OpenCL.cpp
@@ -101,6 +101,7 @@ bool CLSymbols::load(const std::string &library)
     clGetDeviceIDs            = reinterpret_cast<clGetDeviceIDs_func>(dlsym(handle, "clGetDeviceIDs"));
     clRetainEvent             = reinterpret_cast<clRetainEvent_func>(dlsym(handle, "clRetainEvent"));
     clGetPlatformIDs          = reinterpret_cast<clGetPlatformIDs_func>(dlsym(handle, "clGetPlatformIDs"));
+    clGetKernelWorkGroupInfo  = reinterpret_cast<clGetKernelWorkGroupInfo_func>(dlsym(handle, "clGetKernelWorkGroupInfo"));
 
     dlclose(handle);
 
@@ -647,3 +648,23 @@ cl_int clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms, cl_uint
         return CL_OUT_OF_RESOURCES;
     }
 }
+
+cl_int
+clGetKernelWorkGroupInfo(cl_kernel                 kernel,
+                         cl_device_id              device,
+                         cl_kernel_work_group_info param_name,
+                         size_t                    param_value_size,
+                         void                     *param_value,
+                         size_t                   *param_value_size_ret)
+{
+    arm_compute::CLSymbols::get().load_default();
+    auto func = arm_compute::CLSymbols::get().clGetKernelWorkGroupInfo;
+    if(func != nullptr)
+    {
+        return func(kernel, device, param_name, param_value_size, param_value, param_value_size_ret);
+    }
+    else
+    {
+        return CL_OUT_OF_RESOURCES;
+    }
+}