added cv::normalize to T-API; implemented missed UMat::copyTo with mask

modules/core/src/convert.cpp

@@ -1357,43 +1357,65 @@ void cv::LUT( InputArray _src, InputArray _lut, OutputArray _dst )
         func(ptrs[0], lut.data, ptrs[1], len, cn, lutcn);
 }
 
+namespace cv {
+
+static bool ocl_normalize( InputArray _src, OutputArray _dst, InputArray _mask, int rtype,
+                           double scale, double shift )
+{
+    UMat src = _src.getUMat(), dst = _dst.getUMat();
+
+    if( _mask.empty() )
+        src.convertTo( dst, rtype, scale, shift );
+    else
+    {
+        UMat temp;
+        src.convertTo( temp, rtype, scale, shift );
+        temp.copyTo( dst, _mask );
+    }
+
+    return true;
+}
+
+}
 
 void cv::normalize( InputArray _src, OutputArray _dst, double a, double b,
                     int norm_type, int rtype, InputArray _mask )
 {
-    Mat src = _src.getMat(), mask = _mask.getMat();
-
     double scale = 1, shift = 0;
     if( norm_type == CV_MINMAX )
     {
         double smin = 0, smax = 0;
         double dmin = MIN( a, b ), dmax = MAX( a, b );
-        minMaxLoc( _src, &smin, &smax, 0, 0, mask );
+        minMaxLoc( _src, &smin, &smax, 0, 0, _mask );
         scale = (dmax - dmin)*(smax - smin > DBL_EPSILON ? 1./(smax - smin) : 0);
         shift = dmin - smin*scale;
     }
     else if( norm_type == CV_L2 || norm_type == CV_L1 || norm_type == CV_C )
     {
-        scale = norm( src, norm_type, mask );
+        scale = norm( _src, norm_type, _mask );
         scale = scale > DBL_EPSILON ? a/scale : 0.;
         shift = 0;
     }
     else
         CV_Error( CV_StsBadArg, "Unknown/unsupported norm type" );
 
+    int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
     if( rtype < 0 )
-        rtype = _dst.fixedType() ? _dst.depth() : src.depth();
+        rtype = _dst.fixedType() ? _dst.depth() : depth;
+    _dst.createSameSize(_src, CV_MAKETYPE(rtype, cn));
 
-    _dst.create(src.dims, src.size, CV_MAKETYPE(rtype, src.channels()));
-    Mat dst = _dst.getMat();
+    if (ocl::useOpenCL() && _dst.isUMat() &&
+            ocl_normalize(_src, _dst, _mask, rtype, scale, shift))
+        return;
 
-    if( !mask.data )
+    Mat src = _src.getMat(), dst = _dst.getMat();
+    if( _mask.empty() )
         src.convertTo( dst, rtype, scale, shift );
     else
     {
         Mat temp;
         src.convertTo( temp, rtype, scale, shift );
-        temp.copyTo( dst, mask );
+        temp.copyTo( dst, _mask );
     }
 }
 

modules/core/src/opencl/copyset.cl

@@ -41,6 +41,52 @@
 //
 //M*/
 
+#ifdef COPY_TO_MASK
+
+#define DEFINE_DATA \
+    int src_index = mad24(y, src_step, x*(int)sizeof(T)*scn + src_offset); \
+    int dst_index = mad24(y, dst_step, x*(int)sizeof(T)*scn + dst_offset); \
+     \
+    __global const T * src = (__global const T *)(srcptr + src_index); \
+    __global T * dst = (__global T *)(dstptr + dst_index)
+
+__kernel void copyToMask(__global const uchar * srcptr, int src_step, int src_offset,
+                         __global const uchar * maskptr, int mask_step, int mask_offset,
+                         __global uchar * dstptr, int dst_step, int dst_offset,
+                         int dst_rows, int dst_cols)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < dst_cols && y < dst_rows)
+    {
+        int mask_index = mad24(y, mask_step, x * mcn + mask_offset);
+        __global const uchar * mask = (__global const uchar *)(maskptr + mask_index);
+
+#if mcn == 1
+        if (mask[0])
+        {
+            DEFINE_DATA;
+
+            #pragma unroll
+            for (int c = 0; c < scn; ++c)
+                dst[c] = src[c];
+        }
+#elif scn == mcn
+        DEFINE_DATA;
+
+        #pragma unroll
+        for (int c = 0; c < scn; ++c)
+            if (mask[c])
+                dst[c] = src[c];
+#else
+#error "(mcn == 1 || mcn == scn) should be true"
+#endif
+    }
+}
+
+#else
+
 __kernel void setMask(__global const uchar* mask, int maskstep, int maskoffset,
                       __global uchar* dstptr, int dststep, int dstoffset,
                       int rows, int cols, dstT value )
@@ -71,3 +117,5 @@ __kernel void set(__global uchar* dstptr, int dststep, int dstoffset,
         *(__global dstT*)(dstptr + dst_index) = value;
     }
 }
+
+#endif

modules/core/src/umatrix.cpp

@@ -661,6 +661,45 @@ void UMat::copyTo(OutputArray _dst) const
     }
 }
 
+void UMat::copyTo(OutputArray _dst, InputArray _mask) const
+{
+    if( _mask.empty() )
+    {
+        copyTo(_dst);
+        return;
+    }
+
+    int cn = channels(), mtype = _mask.type(), mdepth = CV_MAT_DEPTH(mtype), mcn = CV_MAT_CN(mtype);
+    CV_Assert( mdepth == CV_8U && (mcn == 1 || mcn == cn) );
+
+    if (ocl::useOpenCL() && _dst.isUMat() && dims <= 2)
+    {
+        UMatData * prevu = _dst.getUMat().u;
+        _dst.create( dims, size, type() );
+
+        UMat dst = _dst.getUMat();
+
+        if( prevu != dst.u ) // do not leave dst uninitialized
+            dst = Scalar(0);
+
+        ocl::Kernel k("copyToMask", ocl::core::copyset_oclsrc,
+                      format("-D COPY_TO_MASK -D T=%s -D scn=%d -D mcn=%d",
+                             ocl::memopTypeToStr(depth()), cn, mcn));
+        if (!k.empty())
+        {
+            k.args(ocl::KernelArg::ReadOnlyNoSize(*this), ocl::KernelArg::ReadOnlyNoSize(_mask.getUMat()),
+                   ocl::KernelArg::WriteOnly(dst));
+
+            size_t globalsize[2] = { cols, rows };
+            if (k.run(2, globalsize, NULL, false))
+                return;
+        }
+    }
+
+    Mat src = getMat(ACCESS_READ);
+    src.copyTo(_dst, _mask);
+}
+
 void UMat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const
 {
     bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON;

modules/core/test/ocl/test_arithm.cpp

@@ -1180,6 +1180,28 @@ OCL_TEST_P(Sqrt, Mat)
     }
 }
 
+//////////////////////////////// Normalize ////////////////////////////////////////////////
+
+typedef ArithmTestBase Normalize;
+
+OCL_TEST_P(Normalize, Mat)
+{
+    static int modes[] = { CV_MINMAX, CV_L2, CV_L1, CV_C };
+
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        for (int i = 0, size = sizeof(modes) / sizeof(modes[0]); i < size; ++i)
+        {
+            OCL_OFF(cv::normalize(src1_roi, dst1_roi, 10, 110, modes[i], src1_roi.type(), mask_roi));
+            OCL_ON(cv::normalize(usrc1_roi, udst1_roi,  10, 110, modes[i], src1_roi.type(), umask_roi));
+
+            Near(1);
+        }
+    }
+}
+
 //////////////////////////////////////// Instantiation /////////////////////////////////////////
 
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
@@ -1214,6 +1236,8 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx, Combine(OCL_ALL_DEPTHS, OCL_ALL_C
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx_Mask, Combine(OCL_ALL_DEPTHS, ::testing::Values(Channels(1)), Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sqrt, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Normalize, Combine(OCL_ALL_DEPTHS, Values(Channels(1)), Bool()));
+
 
 } } // namespace cvtest::ocl
 

modules/core/test/ocl/test_matrix_operation.cpp

@@ -54,7 +54,7 @@ namespace ocl {
 
 ////////////////////////////////converto/////////////////////////////////////////////////
 
-PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool)
+PARAM_TEST_CASE(ConvertTo, MatDepth, MatDepth, Channels, bool)
 {
     int src_depth, cn, dstType;
     bool use_roi;
@@ -85,8 +85,6 @@ PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool)
     }
 };
 
-typedef MatrixTestBase ConvertTo;
-
 OCL_TEST_P(ConvertTo, Accuracy)
 {
     for (int j = 0; j < test_loop_times; j++)
@@ -103,7 +101,51 @@ OCL_TEST_P(ConvertTo, Accuracy)
     }
 }
 
-typedef MatrixTestBase CopyTo;
+//////////////////////////////// CopyTo /////////////////////////////////////////////////
+
+PARAM_TEST_CASE(CopyTo, MatDepth, Channels, bool, bool)
+{
+    int depth, cn;
+    bool use_roi, use_mask;
+
+    TEST_DECLARE_INPUT_PARAMETER(src)
+    TEST_DECLARE_INPUT_PARAMETER(mask)
+    TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+    virtual void SetUp()
+    {
+        depth = GET_PARAM(0);
+        cn = GET_PARAM(1);
+        use_roi = GET_PARAM(2);
+        use_mask = GET_PARAM(3);
+    }
+
+    void generateTestData()
+    {
+        const int type = CV_MAKE_TYPE(depth, cn);
+
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        if (use_mask)
+        {
+            Border maskBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+            int mask_cn = randomDouble(0.0, 2.0) > 1.0 ? cn : 1;
+            randomSubMat(mask, mask_roi, roiSize, maskBorder, CV_8UC(mask_cn), 0, 2);
+            cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
+        }
+
+        Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 16);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src)
+        if (use_mask)
+            UMAT_UPLOAD_INPUT_PARAMETER(mask)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+    }
+};
+
 
 OCL_TEST_P(CopyTo, Accuracy)
 {
@@ -111,8 +153,16 @@ OCL_TEST_P(CopyTo, Accuracy)
     {
         generateTestData();
 
-        OCL_OFF(src_roi.copyTo(dst_roi));
-        OCL_ON(usrc_roi.copyTo(udst_roi));
+        if (use_mask)
+        {
+            OCL_OFF(src_roi.copyTo(dst_roi, mask_roi));
+            OCL_ON(usrc_roi.copyTo(udst_roi, umask_roi));
+        }
+        else
+        {
+            OCL_OFF(src_roi.copyTo(dst_roi));
+            OCL_ON(usrc_roi.copyTo(udst_roi));
+        }
 
         OCL_EXPECT_MATS_NEAR(dst, 0);
     }
@@ -122,7 +172,7 @@ OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine(
                             OCL_ALL_DEPTHS, OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, CopyTo, Combine(
-                                OCL_ALL_DEPTHS, Values((MatDepth)0), OCL_ALL_CHANNELS, Bool()));
+                                OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
 
 } } // namespace cvtest::ocl