diff --git a/projects/clr/opencl/api/opencl/amdocl/cl_execute.cpp b/projects/clr/opencl/api/opencl/amdocl/cl_execute.cpp index 15415ac3ad..0336353929 100644 --- a/projects/clr/opencl/api/opencl/amdocl/cl_execute.cpp +++ b/projects/clr/opencl/api/opencl/amdocl/cl_execute.cpp @@ -199,13 +199,6 @@ RUNTIME_ENTRY(cl_int, clEnqueueNDRangeKernel, #endif // CL_VERSION if (global_work_size == NULL) { return CL_INVALID_VALUE; - } else { - // >32bits global work size is not supported. - for (cl_uint dim = 0; dim < work_dim; ++dim) { - if (global_work_size[dim] > static_cast(0xffffffff)) { - return CL_INVALID_GLOBAL_WORK_SIZE; - } - } } if (local_work_size == NULL) { @@ -214,24 +207,28 @@ RUNTIME_ENTRY(cl_int, clEnqueueNDRangeKernel, } else { size_t numWorkItems = 1; for (cl_uint dim = 0; dim < work_dim; ++dim) { - if (local_work_size[dim] == 0 || - local_work_size[dim] > device.info().maxWorkItemSizes_[dim]) { - return CL_INVALID_WORK_ITEM_SIZE; - } - if ((local_work_size[dim] != 0) && (devKernel->workGroupInfo()->compileSize_[0] != 0) && + if ((devKernel->workGroupInfo()->compileSize_[0] != 0) && (local_work_size[dim] != devKernel->workGroupInfo()->compileSize_[dim])) { return CL_INVALID_WORK_GROUP_SIZE; } - if ((global_work_size[dim] == 0) || (((global_work_size[dim] % local_work_size[dim]) != 0) && - (!device.settings().partialDispatch_ || - devKernel->workGroupInfo()->uniformWorkGroupSize_))) { - return CL_INVALID_WORK_GROUP_SIZE; + // >32bits global work size is not supported. + if ((global_work_size[dim] == 0) || (global_work_size[dim] > static_cast(0xffffffff))) { + return CL_INVALID_GLOBAL_WORK_SIZE; } numWorkItems *= local_work_size[dim]; } - if (numWorkItems > devKernel->workGroupInfo()->size_) { + // Make sure local work size is valid + if ((numWorkItems == 0) || (numWorkItems > devKernel->workGroupInfo()->size_)) { return CL_INVALID_WORK_GROUP_SIZE; } + // Check if uniform was requested and validate dimensions + if (devKernel->workGroupInfo()->uniformWorkGroupSize_) { + for (cl_uint dim = 0; dim < work_dim; ++dim) { + if ((global_work_size[dim] % local_work_size[dim]) != 0) { + return CL_INVALID_WORK_GROUP_SIZE; + } + } + } } // Check that all parameters have been defined.