diff --git a/modules/core/src/opencl/copyset.cl b/modules/core/src/opencl/copyset.cl index 42796eac37..a2538b2ab0 100644 --- a/modules/core/src/opencl/copyset.cl +++ b/modules/core/src/opencl/copyset.cl @@ -101,32 +101,39 @@ __kernel void copyToMask(__global const uchar * srcptr, int src_step, int src_of __kernel void setMask(__global const uchar* mask, int maskstep, int maskoffset, __global uchar* dstptr, int dststep, int dstoffset, - int rows, int cols, dstST value_ ) + int rows, int cols, dstST value_) { int x = get_global_id(0); - int y = get_global_id(1); + int y0 = get_global_id(1) * rowsPerWI; - if (x < cols && y < rows) + if (x < cols) { - int mask_index = mad24(y, maskstep, x + maskoffset); - if( mask[mask_index] ) + int mask_index = mad24(y0, maskstep, x + maskoffset); + int dst_index = mad24(y0, dststep, mad24(x, (int)sizeof(dstT1) * cn, dstoffset)); + + for (int y = y0, y1 = min(rows, y0 + rowsPerWI); y < y1; ++y) { - int dst_index = mad24(y, dststep, mad24(x, (int)sizeof(dstT1) * cn, dstoffset)); - storedst(value); + if( mask[mask_index] ) + storedst(value); + + mask_index += maskstep; + dst_index += dststep; } } } __kernel void set(__global uchar* dstptr, int dststep, int dstoffset, - int rows, int cols, dstST value_ ) + int rows, int cols, dstST value_) { int x = get_global_id(0); - int y = get_global_id(1); + int y0 = get_global_id(1) * rowsPerWI; - if (x < cols && y < rows) + if (x < cols) { - int dst_index = mad24(y, dststep, mad24(x, (int)sizeof(dstT1) * cn, dstoffset)); - storedst(value); + int dst_index = mad24(y0, dststep, mad24(x, (int)sizeof(dstT1) * cn, dstoffset)); + + for (int y = y0, y1 = min(rows, y0 + rowsPerWI); y < y1; ++y, dst_index += dststep) + storedst(value); } } diff --git a/modules/core/src/umatrix.cpp b/modules/core/src/umatrix.cpp index 0060492541..aa794ef838 100644 --- a/modules/core/src/umatrix.cpp +++ b/modules/core/src/umatrix.cpp @@ -765,27 +765,27 @@ UMat& UMat::setTo(InputArray _value, InputArray _mask) { Mat value = _value.getMat(); CV_Assert( checkScalar(value, type(), _value.kind(), _InputArray::UMAT) ); - double buf[4]={0,0,0,0}; - convertAndUnrollScalar(value, tp, (uchar*)buf, 1); + double buf[4] = { 0, 0, 0, 0 }; + convertAndUnrollScalar(value, tp, (uchar *)buf, 1); - int scalarcn = cn == 3 ? 4 : cn; - char opts[1024]; - sprintf(opts, "-D dstT=%s -D dstST=%s -D dstT1=%s -D cn=%d", ocl::memopTypeToStr(tp), - ocl::memopTypeToStr(CV_MAKETYPE(tp,scalarcn)), - ocl::memopTypeToStr(CV_MAT_DEPTH(tp)), cn); + int scalarcn = cn == 3 ? 4 : cn, rowsPerWI = ocl::Device::getDefault().isIntel() ? 4 : 1; + String opts = format("-D dstT=%s -D rowsPerWI=%d -D dstST=%s -D dstT1=%s -D cn=%d", + ocl::memopTypeToStr(tp), rowsPerWI, + ocl::memopTypeToStr(CV_MAKETYPE(tp, scalarcn)), + ocl::memopTypeToStr(CV_MAT_DEPTH(tp)), cn); ocl::Kernel setK(haveMask ? "setMask" : "set", ocl::core::copyset_oclsrc, opts); if( !setK.empty() ) { - ocl::KernelArg scalararg(0, 0, 0, 0, buf, CV_ELEM_SIZE1(tp)*scalarcn); + ocl::KernelArg scalararg(0, 0, 0, 0, buf, CV_ELEM_SIZE1(tp) * scalarcn); UMat mask; if( haveMask ) { mask = _mask.getUMat(); - CV_Assert( mask.size() == size() && mask.type() == CV_8U ); - ocl::KernelArg maskarg = ocl::KernelArg::ReadOnlyNoSize(mask); - ocl::KernelArg dstarg = ocl::KernelArg::ReadWrite(*this); + CV_Assert( mask.size() == size() && mask.type() == CV_8UC1 ); + ocl::KernelArg maskarg = ocl::KernelArg::ReadOnlyNoSize(mask), + dstarg = ocl::KernelArg::ReadWrite(*this); setK.args(maskarg, dstarg, scalararg); } else @@ -794,8 +794,8 @@ UMat& UMat::setTo(InputArray _value, InputArray _mask) setK.args(dstarg, scalararg); } - size_t globalsize[] = { cols, rows }; - if( setK.run(2, globalsize, 0, false) ) + size_t globalsize[] = { cols, (rows + rowsPerWI - 1) / rowsPerWI }; + if( setK.run(2, globalsize, NULL, false) ) return *this; } }