diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 995c10b5be..07da4ad724 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -126,7 +126,7 @@ void MatAllocator::upload(UMatData* u, const void* srcptr, int dims, const size_ void MatAllocator::copy(UMatData* usrc, UMatData* udst, int dims, const size_t sz[], const size_t srcofs[], const size_t srcstep[], - const size_t dstofs[], const size_t dststep[], bool sync) const + const size_t dstofs[], const size_t dststep[], bool /*sync*/) const { if(!usrc || !udst) return; diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 22d8022439..9032546e40 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -2258,25 +2258,21 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[], return false; cl_command_queue qq = getQueue(q); - size_t offset[CV_MAX_DIM] = {0}, globalsize[CV_MAX_DIM] = {1,1,1}, localsize[CV_MAX_DIM] = {1,1,1}; + size_t offset[CV_MAX_DIM] = {0}, globalsize[CV_MAX_DIM] = {1,1,1}; size_t total = 1; + CV_Assert(_globalsize != 0); for (int i = 0; i < dims; i++) { - size_t val0 = _localsize ? _localsize[i] : - dims == 1 ? 64 : dims == 2 ? 16>>i : dims == 3 ? 8>>(i>0) : 1; - size_t val = 1; - while( val*2 < val0 ) - val *= 2; - if( _localsize ) - localsize[i] = val; - CV_Assert(_globalsize && _globalsize[i] >= 0); + size_t val = _localsize ? _localsize[i] : + dims == 1 ? 64 : dims == 2 ? (16>>i) : dims == 3 ? (8>>(int)(i>0)) : 1; + CV_Assert( val > 0 ); total *= _globalsize[i]; globalsize[i] = ((_globalsize[i] + val - 1)/val)*val; } if( total == 0 ) return true; cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims, - offset, globalsize, _localsize ? localsize : 0, 0, 0, + offset, globalsize, _localsize, 0, 0, sync ? 0 : &p->e); if( sync || retval < 0 ) { @@ -2976,7 +2972,7 @@ public: return; // there should be no user-visible CPU copies of the UMat which we are going to copy to - CV_Assert(u->refcount == 0); + CV_Assert(u->refcount == 0 || u->tempUMat()); size_t total = 0, new_sz[] = {0, 0, 0}; size_t srcrawofs = 0, new_srcofs[] = {0, 0, 0}, new_srcstep[] = {0, 0, 0}; @@ -3028,7 +3024,7 @@ public: void copy(UMatData* src, UMatData* dst, int dims, const size_t sz[], const size_t srcofs[], const size_t srcstep[], - const size_t dstofs[], const size_t dststep[], bool sync) const + const size_t dstofs[], const size_t dststep[], bool _sync) const { if(!src || !dst) return; @@ -3072,14 +3068,15 @@ public: cl_int retval; CV_Assert( (retval = clEnqueueCopyBufferRect(q, (cl_mem)src->handle, (cl_mem)dst->handle, new_srcofs, new_dstofs, new_sz, - new_srcstep[0], new_srcstep[1], new_dststep[0], new_dststep[1], + new_srcstep[0], new_srcstep[1], + new_dststep[0], new_dststep[1], 0, 0, 0)) >= 0 ); } dst->markHostCopyObsolete(true); dst->markDeviceCopyObsolete(false); - if( sync ) + if( _sync ) clFinish(q); } diff --git a/modules/core/test/test_umat.cpp b/modules/core/test/test_umat.cpp index 54df893bfe..7cd523d85c 100644 --- a/modules/core/test/test_umat.cpp +++ b/modules/core/test/test_umat.cpp @@ -201,7 +201,7 @@ void CV_UMatTest::run( int /* start_from */) TEST(Core_UMat, base) { CV_UMatTest test; test.safe_run(); } -TEST(Core_UMat, simple) +TEST(Core_UMat, getUMat) { { int a[3] = { 1, 2, 3 }; @@ -216,7 +216,7 @@ TEST(Core_UMat, simple) { uchar * const ptr = m.ptr(y); for (int x = 0; x < m.cols; ++x) - ptr[x] = x + y * 2; + ptr[x] = (uchar)(x + y * 2); } ref = m.clone(); diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 85c0ca6e34..883e2eb4d4 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -124,16 +124,17 @@ __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset, #endif } #else - const int x0 = get_global_id(0)*STRIPE_SIZE; - const int x1 = min(x0 + STRIPE_SIZE, cols); + const int x_min = get_global_id(0)*STRIPE_SIZE; + const int x_max = min(x_min + STRIPE_SIZE, cols); const int y = get_global_id(1); if( y < rows ) { - __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset)) + x0*scn; + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + + mad24(y, srcstep, srcoffset)) + x_min*scn; __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset)); int x; - for( x = x0; x < x1; x++, src += scn ) + for( x = x_min; x < x_max; x++, src += scn ) #ifdef DEPTH_5 dst[x] = src[bidx] * 0.114f + src[1] * 0.587f + src[(bidx^2)] * 0.299f; #else diff --git a/modules/ocl/test/test_api.cpp b/modules/ocl/test/test_api.cpp index a95e60e65c..68888a7c08 100644 --- a/modules/ocl/test/test_api.cpp +++ b/modules/ocl/test/test_api.cpp @@ -136,4 +136,3 @@ TEST(OCL_TestTAPI, performance) cv::destroyWindow("result0"); cv::destroyWindow("result1");*/ } -