From 39cb77e158b806461507758d1fdbd7895b75ca0d Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 20 Oct 2026 10:20:06 +0200 Subject: [PATCH 1/3] Added negative test for non-uniform work-group size --- ci/pocl/golden.json | 1 + .../api/negative_enqueue_ndrange.cpp | 60 +++++++++++++++++++ 2 files changed, 61 insertions(+) diff --git a/ci/pocl/golden.json b/ci/pocl/golden.json index 144dd102ae..eeac7f2543 100644 --- a/ci/pocl/golden.json +++ b/ci/pocl/golden.json @@ -20,6 +20,7 @@ "negative_enqueue_zero_work_group": "pass", "negative_set_kernel_arg_invalid_image_msaa": "skip", "negative_set_kernel_arg_invalid_image": "pass", + "negative_enqueue_non_uniform_work_group_size": "pass", "set_kernel_arg_struct_array": "pass" } }, diff --git a/test_conformance/api/negative_enqueue_ndrange.cpp b/test_conformance/api/negative_enqueue_ndrange.cpp index 1082fa711d..e771c4d397 100644 --- a/test_conformance/api/negative_enqueue_ndrange.cpp +++ b/test_conformance/api/negative_enqueue_ndrange.cpp @@ -14,6 +14,7 @@ // limitations under the License. // #include "testBase.h" +#include "harness/parseParameters.h" // Verify that clEnqueueNDRangeKernel returns CL_INVALID_WORK_GROUP_SIZE // when local_work_size is not NULL and the total number of work-items in @@ -86,3 +87,62 @@ REGISTER_TEST(negative_enqueue_zero_work_group) return TEST_PASS; } + +// Verify that clEnqueueNDRangeKernel returns CL_INVALID_WORK_GROUP_SIZE when +// the global size is not evenly divisible by the local size and the kernel +// comes from a program built from source without -cl-std. Such a program is +// compiled for OpenCL C 1.x, which requires uniform work-groups +REGISTER_TEST(negative_enqueue_non_uniform_work_group_size) +{ + cl_int error = CL_SUCCESS; + + // Offline compilation rules do not apply + if (gCompilationMode != kOnline) + { + log_info("Test requires online compilation, skipping test\n"); + return TEST_SKIPPED_ITSELF; + } + + // Build without options (including -cl-std), so the program is compiled for + // the highest OpenCL C 1.x version supported by the device. Thus, + // create_single_kernel_helper is not used - it adds -cl-std automatically + const char* kernel_test = "__kernel void test(){}\n"; + + clProgramWrapper program = + clCreateProgramWithSource(context, 1, &kernel_test, NULL, &error); + test_error(error, "clCreateProgramWithSource failed"); + + error = clBuildProgram(program, 1, &device, NULL, NULL, NULL); + test_error(error, "clBuildProgram failed"); + + clKernelWrapper kernel = clCreateKernel(program, "test", &error); + test_error(error, "clCreateKernel failed"); + + size_t maxWorkGroupSize = 0; + error = get_max_allowed_1d_work_group_size_on_device(device, kernel, + &maxWorkGroupSize); + test_error(error, "get_max_allowed_1d_work_group_size_on_device failed"); + + // A local size of 2 with a global size of 3 gives a non-uniform + // work-group, while the local size stays within all limits. + const size_t localWorkSize = 2; + const size_t globalWorkSize = 3; + + if (maxWorkGroupSize < localWorkSize) + { + log_info("Device does not allow a work-group size of %zu, skipping " + "test\n", + localWorkSize); + return TEST_SKIPPED_ITSELF; + } + + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalWorkSize, + &localWorkSize, 0, NULL, NULL); + test_failure_error_ret(error, CL_INVALID_WORK_GROUP_SIZE, + "Program built without -cl-std requires uniform " + "work-groups but clEnqueueNDRangeKernel did not " + "return CL_INVALID_WORK_GROUP_SIZE as expected", + TEST_FAIL); + + return TEST_PASS; +} From 1170d5089443ee6787f11a2d27f422e3ea0ed3e6 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 20 Oct 2026 11:03:44 +0200 Subject: [PATCH 2/3] corrected configuration due to CI/CD printout --- ci/pocl/golden.json | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ci/pocl/golden.json b/ci/pocl/golden.json index eeac7f2543..d924a0fd2a 100644 --- a/ci/pocl/golden.json +++ b/ci/pocl/golden.json @@ -20,7 +20,7 @@ "negative_enqueue_zero_work_group": "pass", "negative_set_kernel_arg_invalid_image_msaa": "skip", "negative_set_kernel_arg_invalid_image": "pass", - "negative_enqueue_non_uniform_work_group_size": "pass", + "negative_enqueue_non_uniform_work_group_size": "skip", "set_kernel_arg_struct_array": "pass" } }, From a852383a027aa3c3ee252f27a38bdfece65c46b4 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 20 Oct 2026 11:49:19 +0200 Subject: [PATCH 3/3] corrected configuration related to code review --- ci/pocl/golden.json | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ci/pocl/golden.json b/ci/pocl/golden.json index d924a0fd2a..eeac7f2543 100644 --- a/ci/pocl/golden.json +++ b/ci/pocl/golden.json @@ -20,7 +20,7 @@ "negative_enqueue_zero_work_group": "pass", "negative_set_kernel_arg_invalid_image_msaa": "skip", "negative_set_kernel_arg_invalid_image": "pass", - "negative_enqueue_non_uniform_work_group_size": "skip", + "negative_enqueue_non_uniform_work_group_size": "pass", "set_kernel_arg_struct_array": "pass" } },