I think when we did #1501 we missed an error condition, or perhaps the error condition was missing all along and it's just easier to see after the error code consistency changes.
The specific error conditions I'm looking at are:
CL_INVALID_WORK_GROUP_SIZE
- if local_work_size is not NULL, if the work-group size must be uniform, and if the global_work_size is not evenly divisible by the local_work_size
- if local_work_size is not NULL and if the total number of work-items in the work-group is greater than the maximum work-group size supported for kernel on the device associated with command_queue (is greater than the value returned for CL_KERNEL_WORK_GROUP_SIZE
- if local_work_size is not NULL and if the local_work_size does not match the required work-group size for kernel
- if local_work_size is not NULL and if the local_work_size is not consistent with the required number of sub-groups for kernel
The specific error condition that I believe is missing is the case where the local_work_size is not NULL, the work-group size is zero, and the work-group size may be non-uniform. I think this case should be an error, but it is not covered by any of these CL_INVALID_WORK_GROUP_SIZE error conditions, or by any other clEnqueueNDRangeKernel error conditions.
Note also, it's kind of only disallowed for the "work-group size must be uniform" case indirectly, since it's impossible for the global work size to be evenly evenly divisible by zero (the local work size).
One way to fix this is to add an explicit error condition, something like:
- if local_work_size is not NULL and the total number of work-items in the work-group is zero
Do we agree that this should be an error? If so, is this the best way to document this error condition?
I think when we did #1501 we missed an error condition, or perhaps the error condition was missing all along and it's just easier to see after the error code consistency changes.
The specific error conditions I'm looking at are:
The specific error condition that I believe is missing is the case where the local_work_size is not NULL, the work-group size is zero, and the work-group size may be non-uniform. I think this case should be an error, but it is not covered by any of these CL_INVALID_WORK_GROUP_SIZE error conditions, or by any other clEnqueueNDRangeKernel error conditions.
Note also, it's kind of only disallowed for the "work-group size must be uniform" case indirectly, since it's impossible for the global work size to be evenly evenly divisible by zero (the local work size).
One way to fix this is to add an explicit error condition, something like:
Do we agree that this should be an error? If so, is this the best way to document this error condition?