mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-06-17 07:10:51 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

ocl: update kernel global size adjustment

Browse Source 
Prevents 10000x1 => 10000x8 transformation after getContinuousSize() call
This commit is contained in:
Alexander Alekhin 2017-09-18 19:04:46 +03:00
parent 0624411875
commit 2e68f89225
1 changed files with 6 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
modules/core/src/ocl.cpp
View File

@ -2269,7 +2269,7 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
return false; return false;
cl_command_queue qq = getQueue(q); cl_command_queue qq = getQueue(q);
size_t offset[CV_MAX_DIM] = {0}, globalsize[CV_MAX_DIM] = {1,1,1}; size_t globalsize[CV_MAX_DIM] = {1,1,1};
size_t total = 1; size_t total = 1;
CV_Assert(_globalsize != 0); CV_Assert(_globalsize != 0);
for (int i = 0; i < dims; i++) for (int i = 0; i < dims; i++)
@ -2278,15 +2278,16 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
dims == 1 ? 64 : dims == 2 ? (i == 0 ? 256 : 8) : dims == 3 ? (8>>(int)(i>0)) : 1; dims == 1 ? 64 : dims == 2 ? (i == 0 ? 256 : 8) : dims == 3 ? (8>>(int)(i>0)) : 1;
CV_Assert( val > 0 ); CV_Assert( val > 0 );
total *= _globalsize[i]; total *= _globalsize[i];
globalsize[i] = ((_globalsize[i] + val - 1)/val)*val; if (_globalsize[i] == 1)
val = 1;
globalsize[i] = divUp(_globalsize[i], (unsigned int)val) * val;
} }
if( total == 0 ) CV_Assert(total > 0);
return true;
if( p->haveTempDstUMats ) if( p->haveTempDstUMats )
sync = true; sync = true;
cl_event asyncEvent = 0; cl_event asyncEvent = 0;
cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims, cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims,
offset, globalsize, _localsize, 0, 0, NULL, globalsize, _localsize, 0, 0,
sync ? 0 : &asyncEvent); sync ? 0 : &asyncEvent);
#if CV_OPENCL_SHOW_RUN_ERRORS #if CV_OPENCL_SHOW_RUN_ERRORS
if (retval != CL_SUCCESS) if (retval != CL_SUCCESS)