Merge pull request #2045 from SpecLad:merge-2.4
authorRoman Donchenko <roman.donchenko@itseez.com>
Wed, 25 Dec 2013 12:18:00 +0000 (16:18 +0400)
committerOpenCV Buildbot <buildbot@opencv.org>
Wed, 25 Dec 2013 12:18:00 +0000 (16:18 +0400)
modules/core/src/arithm.cpp
modules/core/src/convert.cpp
modules/core/src/opencl/arithm.cl
modules/core/src/opencl/copyset.cl
modules/core/src/umatrix.cpp
modules/core/test/ocl/test_arithm.cpp
modules/core/test/ocl/test_matrix_operation.cpp
modules/video/src/camshift.cpp

index 359d272..449303c 100644 (file)
@@ -915,11 +915,12 @@ void convertAndUnrollScalar( const Mat& sc, int buftype, uchar* scbuf, size_t bl
 
 enum { OCL_OP_ADD=0, OCL_OP_SUB=1, OCL_OP_RSUB=2, OCL_OP_ABSDIFF=3, OCL_OP_MUL=4,
        OCL_OP_MUL_SCALE=5, OCL_OP_DIV_SCALE=6, OCL_OP_RECIP_SCALE=7, OCL_OP_ADDW=8,
-       OCL_OP_AND=9, OCL_OP_OR=10, OCL_OP_XOR=11, OCL_OP_NOT=12, OCL_OP_MIN=13, OCL_OP_MAX=14 };
+       OCL_OP_AND=9, OCL_OP_OR=10, OCL_OP_XOR=11, OCL_OP_NOT=12, OCL_OP_MIN=13, OCL_OP_MAX=14,
+       OCL_OP_RDIV_SCALE=15 };
 
 static const char* oclop2str[] = { "OP_ADD", "OP_SUB", "OP_RSUB", "OP_ABSDIFF",
     "OP_MUL", "OP_MUL_SCALE", "OP_DIV_SCALE", "OP_RECIP_SCALE",
-    "OP_ADDW", "OP_AND", "OP_OR", "OP_XOR", "OP_NOT", "OP_MIN", "OP_MAX", 0 };
+    "OP_ADDW", "OP_AND", "OP_OR", "OP_XOR", "OP_NOT", "OP_MIN", "OP_MAX", "OP_RDIV_SCALE", 0 };
 
 static bool ocl_binary_op(InputArray _src1, InputArray _src2, OutputArray _dst,
                           InputArray _mask, bool bitwise, int oclop, bool haveScalar )
@@ -1301,25 +1302,27 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
 
     int kercn = haveMask || haveScalar ? cn : 1;
 
-    char cvtstr[3][32], opts[1024];
+    char cvtstr[4][32], opts[1024];
     sprintf(opts, "-D %s%s -D %s -D srcT1=%s -D srcT2=%s "
-            "-D dstT=%s -D workT=%s -D convertToWT1=%s "
+            "-D dstT=%s -D workT=%s -D scaleT=%s -D convertToWT1=%s "
             "-D convertToWT2=%s -D convertToDT=%s%s",
             (haveMask ? "MASK_" : ""), (haveScalar ? "UNARY_OP" : "BINARY_OP"),
             oclop2str[oclop], ocl::typeToStr(CV_MAKETYPE(depth1, kercn)),
             ocl::typeToStr(CV_MAKETYPE(depth2, kercn)),
             ocl::typeToStr(CV_MAKETYPE(ddepth, kercn)),
             ocl::typeToStr(CV_MAKETYPE(wdepth, kercn)),
+            ocl::typeToStr(CV_MAKETYPE(wdepth, 1)),
             ocl::convertTypeStr(depth1, wdepth, kercn, cvtstr[0]),
             ocl::convertTypeStr(depth2, wdepth, kercn, cvtstr[1]),
             ocl::convertTypeStr(wdepth, ddepth, kercn, cvtstr[2]),
             doubleSupport ? " -D DOUBLE_SUPPORT" : "");
 
+    size_t usrdata_esz = CV_ELEM_SIZE(wdepth);
     const uchar* usrdata_p = (const uchar*)usrdata;
     const double* usrdata_d = (const double*)usrdata;
     float usrdata_f[3];
     int i, n = oclop == OCL_OP_MUL_SCALE || oclop == OCL_OP_DIV_SCALE ||
-        oclop == OCL_OP_RECIP_SCALE ? 1 : oclop == OCL_OP_ADDW ? 3 : 0;
+        oclop == OCL_OP_RDIV_SCALE || oclop == OCL_OP_RECIP_SCALE ? 1 : oclop == OCL_OP_ADDW ? 3 : 0;
     if( n > 0 && wdepth == CV_32F )
     {
         for( i = 0; i < n; i++ )
@@ -1352,13 +1355,20 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
         ocl::KernelArg scalararg = ocl::KernelArg(0, 0, 0, buf, esz);
 
         if( !haveMask )
-            k.args(src1arg, dstarg, scalararg);
+        {
+            if(n == 0)
+                k.args(src1arg, dstarg, scalararg);
+            else if(n == 1)
+                k.args(src1arg, dstarg, scalararg,
+                       ocl::KernelArg(0, 0, 0, usrdata_p, usrdata_esz));
+            else
+                CV_Error(Error::StsNotImplemented, "unsupported number of extra parameters");
+        }
         else
             k.args(src1arg, maskarg, dstarg, scalararg);
     }
     else
     {
-        size_t usrdata_esz = CV_ELEM_SIZE(wdepth);
         src2 = _src2.getUMat();
         ocl::KernelArg src2arg = ocl::KernelArg::ReadOnlyNoSize(src2, cscale);
 
@@ -1439,6 +1449,8 @@ static void arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
             swapped12 = true;
             if( oclop == OCL_OP_SUB )
                 oclop = OCL_OP_RSUB;
+            if ( oclop == OCL_OP_DIV_SCALE )
+                oclop = OCL_OP_RDIV_SCALE;
         }
         else if( !checkScalar(*psrc2, type1, kind2, kind1) )
             CV_Error( CV_StsUnmatchedSizes,
index 05c1a6e..6259a7a 100644 (file)
@@ -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 );
     }
 }
 
index 9c86057..1647e8d 100644 (file)
 
 #elif defined OP_MUL_SCALE
 #undef EXTRA_PARAMS
-#define EXTRA_PARAMS , workT scale
-#define PROCESS_ELEM dstelem = convertToDT(srcelem1 * srcelem2 * scale)
+#ifdef UNARY_OP
+#define EXTRA_PARAMS , workT srcelem2, scaleT scale
+#else
+#define EXTRA_PARAMS , scaleT scale
+#endif
+#define PROCESS_ELEM dstelem = convertToDT(srcelem1 * scale * srcelem2)
 
 #elif defined OP_DIV
 #define PROCESS_ELEM \
 
 #elif defined OP_DIV_SCALE
 #undef EXTRA_PARAMS
-#define EXTRA_PARAMS , workT scale
+#ifdef UNARY_OP
+#define EXTRA_PARAMS , workT srcelem2, scaleT scale
+#else
+#define EXTRA_PARAMS , scaleT scale
+#endif
 #define PROCESS_ELEM \
         workT e2 = srcelem2, zero = (workT)(0); \
-        dstelem = convertToDT(e2 != zero ? srcelem1 * scale / e2 : zero)
+        dstelem = convertToDT(e2 == zero ? zero : (srcelem1 * (workT)(scale) / e2))
+
+#elif defined OP_RDIV_SCALE
+#undef EXTRA_PARAMS
+#ifdef UNARY_OP
+#define EXTRA_PARAMS , workT srcelem2, scaleT scale
+#else
+#define EXTRA_PARAMS , scaleT scale
+#endif
+#define PROCESS_ELEM \
+        workT e1 = srcelem1, zero = (workT)(0); \
+        dstelem = convertToDT(e1 == zero ? zero : (srcelem2 * (workT)(scale) / e1))
 
 #elif defined OP_RECIP_SCALE
 #undef EXTRA_PARAMS
-#define EXTRA_PARAMS , workT scale
+#define EXTRA_PARAMS , scaleT scale
 #define PROCESS_ELEM \
         workT e1 = srcelem1, zero = (workT)(0); \
         dstelem = convertToDT(e1 != zero ? scale / e1 : zero)
 
 #elif defined OP_ADDW
 #undef EXTRA_PARAMS
-#define EXTRA_PARAMS , workT alpha, workT beta, workT gamma
+#define EXTRA_PARAMS , scaleT alpha, scaleT beta, scaleT gamma
 #define PROCESS_ELEM dstelem = convertToDT(srcelem1*alpha + srcelem2*beta + gamma)
 
 #elif defined OP_MAG
@@ -260,7 +279,8 @@ dstelem = v > (dstT)(0) ? log(v) : log(-v)
 #undef srcelem2
 #if defined OP_AND || defined OP_OR || defined OP_XOR || defined OP_ADD || defined OP_SAT_ADD || \
     defined OP_SUB || defined OP_SAT_SUB || defined OP_RSUB || defined OP_SAT_RSUB || \
-    defined OP_ABSDIFF || defined OP_CMP || defined OP_MIN || defined OP_MAX || defined OP_POW
+    defined OP_ABSDIFF || defined OP_CMP || defined OP_MIN || defined OP_MAX || defined OP_POW || \
+    defined OP_MUL || defined OP_DIV
     #undef EXTRA_PARAMS
     #define EXTRA_PARAMS , workT srcelem2
 #endif
index 8fb5a00..05cde8e 100644 (file)
 //
 //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
index 0b61374..95e203b 100644 (file)
@@ -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;
index 7a24f31..58edcec 100644 (file)
@@ -293,7 +293,7 @@ OCL_TEST_P(Mul, Mat)
     }
 }
 
-OCL_TEST_P(Mul, DISABLED_Scalar)
+OCL_TEST_P(Mul, Scalar)
 {
     for (int j = 0; j < test_loop_times; j++)
     {
@@ -306,7 +306,7 @@ OCL_TEST_P(Mul, DISABLED_Scalar)
     }
 }
 
-OCL_TEST_P(Mul, DISABLED_Mat_Scale)
+OCL_TEST_P(Mul, Mat_Scale)
 {
     for (int j = 0; j < test_loop_times; j++)
     {
@@ -319,6 +319,20 @@ OCL_TEST_P(Mul, DISABLED_Mat_Scale)
     }
 }
 
+OCL_TEST_P(Mul, Mat_Scalar_Scale)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::multiply(src1_roi, val, dst1_roi, val[0]));
+        OCL_ON(cv::multiply(usrc1_roi, val, udst1_roi, val[0]));
+
+        Near(udst1_roi.depth() >= CV_32F ? 1e-2 : 1);
+    }
+}
+
+
 //////////////////////////////// Div /////////////////////////////////////////////////
 
 typedef ArithmTestBase Div;
@@ -335,7 +349,7 @@ OCL_TEST_P(Div, Mat)
     }
 }
 
-OCL_TEST_P(Div, DISABLED_Scalar)
+OCL_TEST_P(Div, Scalar)
 {
     for (int j = 0; j < test_loop_times; j++)
     {
@@ -348,6 +362,19 @@ OCL_TEST_P(Div, DISABLED_Scalar)
     }
 }
 
+OCL_TEST_P(Div, Scalar2)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::divide(src1_roi, val, dst1_roi));
+        OCL_ON(cv::divide(usrc1_roi, val, udst1_roi));
+
+        Near(udst1_roi.depth() >= CV_32F ? 1e-3 : 1);
+    }
+}
+
 OCL_TEST_P(Div, Mat_Scale)
 {
     for (int j = 0; j < test_loop_times; j++)
@@ -361,8 +388,7 @@ OCL_TEST_P(Div, Mat_Scale)
     }
 }
 
-
-OCL_TEST_P(Div, DISABLED_Mat_Scalar_Scale)
+OCL_TEST_P(Div, Mat_Scalar_Scale)
 {
     for (int j = 0; j < test_loop_times; j++)
     {
@@ -375,6 +401,19 @@ OCL_TEST_P(Div, DISABLED_Mat_Scalar_Scale)
     }
 }
 
+OCL_TEST_P(Div, Recip)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::divide(val[0], src1_roi, dst1_roi));
+        OCL_ON(cv::divide(val[0], usrc1_roi, udst1_roi));
+
+        Near(udst1_roi.depth() >= CV_32F ? 1e-3 : 1);
+    }
+}
+
 //////////////////////////////// Min/Max /////////////////////////////////////////////////
 
 typedef ArithmTestBase Min;
@@ -1180,6 +1219,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 +1275,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
 
index aabbb3f..77c5dad 100644 (file)
@@ -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
 
index 9ba0238..5449a1b 100644 (file)
 
 int cv::meanShift( InputArray _probImage, Rect& window, TermCriteria criteria )
 {
-    Mat mat = _probImage.getMat();
+    Size size;
+    int cn;
+    Mat mat;
+    UMat umat;
+    bool isUMat = _probImage.isUMat();
+
+    if (isUMat)
+        umat = _probImage.getUMat(), cn = umat.channels(), size = umat.size();
+    else
+        mat = _probImage.getMat(), cn = mat.channels(), size = mat.size();
+
     Rect cur_rect = window;
 
-    CV_Assert( mat.channels() == 1 );
+    CV_Assert( cn == 1 );
 
     if( window.height <= 0 || window.width <= 0 )
         CV_Error( Error::StsBadArg, "Input window has non-positive sizes" );
 
-    window = window & Rect(0, 0, mat.cols, mat.rows);
+    window = window & Rect(0, 0, size.width, size.height);
 
     double eps = (criteria.type & TermCriteria::EPS) ? std::max(criteria.epsilon, 0.) : 1.;
     eps = cvRound(eps*eps);
@@ -59,16 +69,16 @@ int cv::meanShift( InputArray _probImage, Rect& window, TermCriteria criteria )
 
     for( i = 0; i < niters; i++ )
     {
-        cur_rect = cur_rect & Rect(0, 0, mat.cols, mat.rows);
+        cur_rect = cur_rect & Rect(0, 0, size.width, size.height);
         if( cur_rect == Rect() )
         {
-            cur_rect.x = mat.cols/2;
-            cur_rect.y = mat.rows/2;
+            cur_rect.x = size.width/2;
+            cur_rect.y = size.height/2;
         }
         cur_rect.width = std::max(cur_rect.width, 1);
         cur_rect.height = std::max(cur_rect.height, 1);
 
-        Moments m = moments(mat(cur_rect));
+        Moments m = isUMat ? moments(umat(cur_rect)) : moments(mat(cur_rect));
 
         // Calculating center of mass
         if( fabs(m.m00) < DBL_EPSILON )
@@ -77,8 +87,8 @@ int cv::meanShift( InputArray _probImage, Rect& window, TermCriteria criteria )
         int dx = cvRound( m.m10/m.m00 - window.width*0.5 );
         int dy = cvRound( m.m01/m.m00 - window.height*0.5 );
 
-        int nx = std::min(std::max(cur_rect.x + dx, 0), mat.cols - cur_rect.width);
-        int ny = std::min(std::max(cur_rect.y + dy, 0), mat.rows - cur_rect.height);
+        int nx = std::min(std::max(cur_rect.x + dx, 0), size.width - cur_rect.width);
+        int ny = std::min(std::max(cur_rect.y + dy, 0), size.height - cur_rect.height);
 
         dx = nx - cur_rect.x;
         dy = ny - cur_rect.y;
@@ -99,9 +109,17 @@ cv::RotatedRect cv::CamShift( InputArray _probImage, Rect& window,
                               TermCriteria criteria )
 {
     const int TOLERANCE = 10;
-    Mat mat = _probImage.getMat();
+    Size size;
+    Mat mat;
+    UMat umat;
+    bool isUMat = _probImage.isUMat();
+
+    if (isUMat)
+        umat = _probImage.getUMat(), size = umat.size();
+    else
+        mat = _probImage.getMat(), size = mat.size();
 
-    meanShift( mat, window, criteria );
+    meanShift( _probImage, window, criteria );
 
     window.x -= TOLERANCE;
     if( window.x < 0 )
@@ -112,15 +130,15 @@ cv::RotatedRect cv::CamShift( InputArray _probImage, Rect& window,
         window.y = 0;
 
     window.width += 2 * TOLERANCE;
-    if( window.x + window.width > mat.cols )
-        window.width = mat.cols - window.x;
+    if( window.x + window.width > size.width )
+        window.width = size.width - window.x;
 
     window.height += 2 * TOLERANCE;
-    if( window.y + window.height > mat.rows )
-        window.height = mat.rows - window.y;
+    if( window.y + window.height > size.height )
+        window.height = size.height - window.y;
 
     // Calculating moments in new center mass
-    Moments m = moments( mat(window) );
+    Moments m = isUMat ? moments(umat(window)) : moments(mat(window));
 
     double m00 = m.m00, m10 = m.m10, m01 = m.m01;
     double mu11 = m.mu11, mu20 = m.mu20, mu02 = m.mu02;
@@ -164,19 +182,19 @@ cv::RotatedRect cv::CamShift( InputArray _probImage, Rect& window,
     int t1 = cvRound( fabs( width * sn ));
 
     t0 = MAX( t0, t1 ) + 2;
-    window.width = MIN( t0, (mat.cols - _xc) * 2 );
+    window.width = MIN( t0, (size.width - _xc) * 2 );
 
     t0 = cvRound( fabs( length * sn ));
     t1 = cvRound( fabs( width * cs ));
 
     t0 = MAX( t0, t1 ) + 2;
-    window.height = MIN( t0, (mat.rows - _yc) * 2 );
+    window.height = MIN( t0, (size.height - _yc) * 2 );
 
     window.x = MAX( 0, _xc - window.width / 2 );
     window.y = MAX( 0, _yc - window.height / 2 );
 
-    window.width = MIN( mat.cols - window.x, window.width );
-    window.height = MIN( mat.rows - window.y, window.height );
+    window.width = MIN( size.width - window.x, window.width );
+    window.height = MIN( size.height - window.y, window.height );
 
     RotatedRect box;
     box.size.height = (float)length;