diff options
| author | Pan Xiuli <xiuli.pan@intel.com> | 2016-03-16 07:16:50 +0800 |
|---|---|---|
| committer | Yang Rong <rong.r.yang@intel.com> | 2016-11-08 20:38:22 +0800 |
| commit | b8e07f6f6ff63e4d34e9d49f224ae123fbd043cb (patch) | |
| tree | c8db4fcc8fa8a24eb2ccfef534f545c1d7e59654 /src/cl_api_kernel.c | |
| parent | d182f461268b415b500278bb5d788ae906d8df93 (diff) | |
Runtime: Add support for non uniform group size
Enqueue multiple times if the the size is not uniform, at most 2
times for 1D, 4times for 2D and 8 times for 3D. Using the workdim
offset of walker in batch buffer to keep work groups in series.
TODO: handle events for the flush between multiple enqueues
Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Diffstat (limited to 'src/cl_api_kernel.c')
| -rw-r--r-- | src/cl_api_kernel.c | 13 |
1 files changed, 0 insertions, 13 deletions
diff --git a/src/cl_api_kernel.c b/src/cl_api_kernel.c index 1fd687bb..70140b23 100644 --- a/src/cl_api_kernel.c +++ b/src/cl_api_kernel.c @@ -89,19 +89,6 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, } } - /* Local sizes must be non-null and divide global sizes */ - if (local_work_size != NULL) { - for (i = 0; i < work_dim; ++i) { - if (UNLIKELY(local_work_size[i] == 0 || global_work_size[i] % local_work_size[i])) { - err = CL_INVALID_WORK_GROUP_SIZE; - break; - } - } - if (err != CL_SUCCESS) { - break; - } - } - /* Queue and kernel must share the same context */ assert(kernel->program); if (command_queue->ctx != kernel->program->ctx) { |