optimized UMat::setTo

diff --git a/modules/core/src/opencl/copyset.cl b/modules/core/src/opencl/copyset.cl
index 42796eac37a134291f58bff3dc5b5a135051ade9..a2538b2ab08db84c165fd707979c132e5816ca25 100644 (file)
--- 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 006049254109a5819a4d6961c4401cc301dc2c32..aa794ef83821af6f7ece7aa75be41de528618385 100644 (file)
--- 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;
         }
     }