Fix CL_INVALID_WORK_ITEM_SIZE for "reduce" test (#143)

* Fix CL_INVALID_WORK_ITEM_SIZE for "reduce" test

The work-group size (WGS) may exceed the maximum allowed per-dimension
limit reported by CL_DEVICE_MAX_WORK_ITEM_SIZES[0], leading to
clEnqueueNDRangeKernel failing with CL_INVALID_WORK_GROUP_SIZE.

This patch queries the device's max work-item sizes and clamps WGS to
the valid maximum for the x-dimension. This ensures correct behavior on
devices with smaller limits and improves portability.

Signed-off-by: Xin Jin <xin.jin@arm.com>

* Clamp WGS to CL_DEVICE_MAX_WORK_ITEM_SIZES[0] in reduce.cpp

The work-group size (WGS) may exceed the maximum allowed per-dimension
limit reported by CL_DEVICE_MAX_WORK_ITEM_SIZES[0], leading to
clEnqueueNDRangeKernel failing with CL_INVALID_WORK_GROUP_SIZE.

The reduce sample launches kernels with a 1-D NDRange over a flat array
of integers, so only the x-dimension limit is relevant. This patch
queries the device's max work-item sizes and clamps WGS to the valid
maximum for the x-dimension. This ensures correct behavior on devices
with smaller limits and improves portability.

Signed-off-by: Xin Jin <xin.jin@arm.com>

---------

Signed-off-by: Xin Jin <xin.jin@arm.com>

Author: xinjin01

samples/core/reduce/main.c

@@ -386,6 +386,15 @@ int main(int argc, char *argv[])
     OCLERROR_RET(clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE,
                                  sizeof(cl_ulong), &loc_mem, NULL),
                  error, red);
+
+    // Ensure WGS is valid for this device
+    size_t max_wi_dims[3];
+    OCLERROR_RET(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES,
+                                 sizeof(max_wi_dims), max_wi_dims, NULL),
+                 error, red);
+
+    wgs = wgs > max_wi_dims[0] ? max_wi_dims[0] : wgs;
+
     while (loc_mem < wgs * 2 * sizeof(cl_int)) wgs -= psm;
 
     if (wgs == 0)

samples/core/reduce/main.cpp

@@ -175,11 +175,14 @@ int main(int argc, char* argv[])
             cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::LocalSpaceArg,
                               cl_ulong, cl_int>(program, "reduce");
 
-        // Query maximum supported WGS of kernel on device based on private mem
-        // (register) constraints
-        auto wgs =
+        // Query kernel- and device-specific limits
+        size_t kernel_wgs =
             reduce.getKernel().getWorkGroupInfo<CL_KERNEL_WORK_GROUP_SIZE>(
                 device);
+        size_t device_wi_x = device.getInfo<CL_DEVICE_MAX_WORK_ITEM_SIZES>()[0];
+
+        // Clamp work-group size to both limits
+        size_t wgs = std::min(kernel_wgs, device_wi_x);
 
         // Further constrain (reduce) WGS based on shared mem size on device
         while (device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>()
@@ -191,7 +194,7 @@ int main(int argc, char* argv[])
 
         if (wgs == 0)
             throw std::runtime_error{
-                "Not enough local memory to serve a single sub-group."
+                "Not enough local memory to serve a single work-group."
             };
 
         cl_ulong factor = wgs * 2;