KhronosGroup / KhronosGroup/OpenCL-Docs
Clarify whether implementations may reject grids with local dimensions not dividing overall dimensions
- Dominant language
- Python
- Stars
- 420
- Forks
- 131
- Avg merge
- 5d 13h
- Merged PRs (30d)
- 11
Description
I was recently surprised to notice both my NVIDIA and Intel OpenCL drivers rejected launch grids due to the workgroup/local dimensions not dividing the overall dimensions. After all, the very fact that we specify overall dimensions rather than "grid-of-workgroups" dimensions (a-la-CUDA) implies that divisibility by the workgroup dimensions is not necessary. And yet here I was.
So, I checked the docs, namely https://registry.khronos.org/OpenCL/specs/unified/refpages/man/html/clEnqueueNDRangeKernel.html (relevant part quoted below) ... and - it says nothing. Neither clarifying that local work dims do not need to divide the global ones, nor saying that implementations may choose whether or not that's allowed. There should be more clarity about this.
> *global_work_size* points to an array of work_dim unsigned values that describe the number of global work-items in work_dim dimensions that will execute the kernel function. The total number of global work-items is computed as global_work_size[0] × … × global_work_size[work_dim - 1]. If the device associated with command_queue is an OpenCL 2.1 or newer device, and global_work_size is NULL or the value in any passed dimension is zero, then the kernel command will trivially succeed after its event dependencies are satisfied and subsequently update its completion event. The behavior in this situation is similar to that of an enqueued marker, except that unlike a marker, an enqueued kernel with no events passed to event_wait_list may run at any time.
>
>*local_work_size* points to an array of work_dim unsigned values that describe the number of work-items that make up a work-group (also referred to as the size of the work-group) that will execute the kernel specified by kernel. The total number of work-items in a work-group is computed as local_work_size[0] × … × local_work_size[work_dim - 1]. The total number of work-items in the work-group must be less than or equal to the CL_KERNEL_WORK_GROUP_SIZE value specified in the [Kernel Object Device Queries](https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_API.html#kernel-workgroup-info-table) table, and the number of work-items specified in local_work_size[0], …, local_work_size[work_dim - 1] must be less than or equal to the corresponding values specified by CL_DEVICE_MAX_WORK_GROUP_SIZES[0], …, CL_DEVICE_MAX_WORK_GROUP_SIZES[work_dim - 1]. The explicitly specified local_work_size will be used to determine how to break the global work-items specified by global_work_size into appropriate work-group instances.
Contributor guide
Research direction
Start with the clEnqueueNDRangeKernel reference page and the quoted descriptions of global_work_size and local_work_size. Check the surrounding OpenCL specification text for the rule governing non-dividing dimensions, then clarify the documentation so it explicitly states whether such launches are required or may be rejected.
Written by the indexing model from the issue text.
Assessment
- Domain
- documentation
- Issue type
- Documentation
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 48/100