From ac5db2d42eaa2adff7f4beea44c38afdacbe2558 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Tue, 5 May 2026 18:04:51 +0200 Subject: [PATCH] rusticl/kernel: return CL_INVALID_WORK_GROUP_SIZE in clEnqueueNDRangeKernel for an explicit 0 workgroup See https://github.com/KhronosGroup/OpenCL-Docs/pull/1542 Part-of: --- src/gallium/frontends/rusticl/api/kernel.rs | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/src/gallium/frontends/rusticl/api/kernel.rs b/src/gallium/frontends/rusticl/api/kernel.rs index 588b8ae214f..ce4e2ed1961 100644 --- a/src/gallium/frontends/rusticl/api/kernel.rs +++ b/src/gallium/frontends/rusticl/api/kernel.rs @@ -699,7 +699,7 @@ fn enqueue_ndrange_kernel( work_dim: cl_uint, global_work_offset: *const usize, global_work_size: *const usize, - local_work_size: *const usize, + local_work_size_ptr: *const usize, num_events_in_wait_list: cl_uint, event_wait_list: *const cl_event, event: *mut cl_event, @@ -732,9 +732,15 @@ fn enqueue_ndrange_kernel( // we assume the application gets it right and doesn't pass shorter arrays then actually needed. let global_work_size = unsafe { kernel_work_arr_or_default(global_work_size, work_dim) }; - let local_work_size = unsafe { kernel_work_arr_or_default(local_work_size, work_dim) }; + let local_work_size = unsafe { kernel_work_arr_or_default(local_work_size_ptr, work_dim) }; let global_work_offset = unsafe { kernel_work_arr_or_default(global_work_offset, work_dim) }; + // CL_INVALID_WORK_GROUP_SIZE if local_work_size is not NULL and if the total number of + // work-items in the work-group is zero + if !local_work_size_ptr.is_null() && local_work_size.iter().any(|&l| l == 0) { + return Err(CL_INVALID_WORK_GROUP_SIZE); + } + let device_bits = q.device.address_bits(); let device_max = u64::MAX >> (u64::BITS - device_bits);