From 997ec60839df42213384975decf2ab1bfd5dda1a Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sat, 30 Nov 2013 16:36:31 +0400
Subject: [PATCH 01/13] added OCL-TAPI arithm tests

---
 modules/core/src/arithm.cpp           |  36 +-
 modules/core/src/opencl/arithm.cl     |   8 +-
 modules/core/test/ocl/test_arithm.cpp | 643 +++++++++++++++++++++++++-
 3 files changed, 660 insertions(+), 27 deletions(-)

diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp
index c3ba8c6d6..0d510b8ad 100644
--- a/modules/core/src/arithm.cpp
+++ b/modules/core/src/arithm.cpp
@@ -929,12 +929,9 @@ static bool ocl_binary_op(InputArray _src1, InputArray _src2, OutputArray _dst,
     int srcdepth = CV_MAT_DEPTH(srctype);
     int cn = CV_MAT_CN(srctype);
 
-    if( oclop < 0 || ((haveMask || haveScalar) && cn > 4) )
+    if( oclop < 0 || ((haveMask || haveScalar) && cn > 4) || cn == 3)
         return false;
 
-    UMat src1 = _src1.getUMat(), src2;
-    UMat dst = _dst.getUMat(), mask = _mask.getUMat();
-
     char opts[1024];
     int kercn = haveMask || haveScalar ? cn : 1;
     sprintf(opts, "-D %s%s -D %s -D dstT=%s",
@@ -946,6 +943,9 @@ static bool ocl_binary_op(InputArray _src1, InputArray _src2, OutputArray _dst,
     if( k.empty() )
         return false;
 
+    UMat src1 = _src1.getUMat(), src2;
+    UMat dst = _dst.getUMat(), mask = _mask.getUMat();
+
     int cscale = cn/kercn;
     ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1, cscale);
     ocl::KernelArg dstarg = haveMask ? ocl::KernelArg::ReadWrite(dst, cscale) :
@@ -1280,24 +1280,28 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
                           void* usrdata, int oclop,
                           bool haveScalar )
 {
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
     int type1 = _src1.type(), depth1 = CV_MAT_DEPTH(type1), cn = CV_MAT_CN(type1);
     bool haveMask = !_mask.empty();
 
-    if( ((haveMask || haveScalar) && cn > 4) || cn == 3) // TODO need fix for 3 channels
+    if( ((haveMask || haveScalar) && cn > 4) || cn == 3)
         return false;
 
     int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), wdepth = std::max(CV_32S, CV_MAT_DEPTH(wtype));
+    if (!doubleSupport)
+        wdepth = std::min(wdepth, CV_32F);
+
     wtype = CV_MAKETYPE(wdepth, cn);
     int type2 = haveScalar ? wtype : _src2.type(), depth2 = CV_MAT_DEPTH(type2);
-    int kercn = haveMask || haveScalar ? cn : 1;
+    if (!doubleSupport && (depth2 == CV_64F || depth1 == CV_64F))
+        return false;
 
-    UMat src1 = _src1.getUMat(), src2;
-    UMat dst = _dst.getUMat(), mask = _mask.getUMat();
+    int kercn = haveMask || haveScalar ? cn : 1;
 
     char cvtstr[3][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 convertToWT2=%s -D convertToDT=%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)),
@@ -1305,7 +1309,8 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
             ocl::typeToStr(CV_MAKETYPE(wdepth, kercn)),
             ocl::convertTypeStr(depth1, wdepth, kercn, cvtstr[0]),
             ocl::convertTypeStr(depth2, wdepth, kercn, cvtstr[1]),
-            ocl::convertTypeStr(wdepth, ddepth, kercn, cvtstr[2]));
+            ocl::convertTypeStr(wdepth, ddepth, kercn, cvtstr[2]),
+            doubleSupport ? " -D DOUBLE_SUPPORT" : "");
 
     const uchar* usrdata_p = (const uchar*)usrdata;
     const double* usrdata_d = (const double*)usrdata;
@@ -1323,6 +1328,9 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
     if( k.empty() )
         return false;
 
+    UMat src1 = _src1.getUMat(), src2;
+    UMat dst = _dst.getUMat(), mask = _mask.getUMat();
+
     int cscale = cn/kercn;
 
     ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1, cscale);
@@ -1337,9 +1345,7 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
         Mat src2sc = _src2.getMat();
 
         if( !src2sc.empty() )
-        {
             convertAndUnrollScalar(src2sc, wtype, (uchar*)buf, 1);
-        }
         ocl::KernelArg scalararg = ocl::KernelArg(0, 0, 0, buf, esz);
 
         if( !haveMask )
@@ -1369,12 +1375,10 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
                 CV_Error(Error::StsNotImplemented, "unsupported number of extra parameters");
         }
         else
-        {
             k.args(src1arg, src2arg, maskarg, dstarg);
-        }
     }
 
-    size_t globalsize[] = { src1.cols*cscale, src1.rows };
+    size_t globalsize[] = { src1.cols * cscale, src1.rows };
     return k.run(2, globalsize, NULL, false);
 }
 
@@ -2075,7 +2079,7 @@ void cv::multiply(InputArray src1, InputArray src2,
                   OutputArray dst, double scale, int dtype)
 {
     arithm_op(src1, src2, dst, noArray(), dtype, getMulTab(),
-              true, &scale, scale == 1. ? OCL_OP_MUL : OCL_OP_MUL_SCALE);
+              true, &scale, std::abs(scale - 1.0) < DBL_EPSILON ? OCL_OP_MUL : OCL_OP_MUL_SCALE);
 }
 
 void cv::divide(InputArray src1, InputArray src2,
diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl
index a9c23645a..87a9dece6 100644
--- a/modules/core/src/opencl/arithm.cl
+++ b/modules/core/src/opencl/arithm.cl
@@ -57,7 +57,7 @@
           -D workDepth=<work depth> [-D cn=<num channels>]" - for mixed-type operations
 */
 
-#if defined (DOUBLE_SUPPORT)
+#ifdef DOUBLE_SUPPORT
 #ifdef cl_khr_fp64
 #pragma OPENCL EXTENSION cl_khr_fp64:enable
 #elif defined (cl_amd_fp64)
@@ -69,7 +69,7 @@
 #define CV_32F 5
 
 #define dstelem *(__global dstT*)(dstptr + dst_index)
-#define noconvert(x) x
+#define noconvert
 
 #ifndef workT
 
@@ -99,7 +99,9 @@
 #define PROCESS_ELEM dstelem = convertToDT(srcelem2 - srcelem1)
 
 #elif defined OP_ABSDIFF
-#define PROCESS_ELEM dstelem = abs_diff(srcelem1, srcelem2)
+#define PROCESS_ELEM \
+    workT v = srcelem1 - srcelem2; \
+    dstelem = convertToDT(v >= (workT)(0) ? v : -v);
 
 #elif defined OP_AND
 #define PROCESS_ELEM dstelem = srcelem1 & srcelem2
diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp
index c574d004f..ed6414bf4 100644
--- a/modules/core/test/ocl/test_arithm.cpp
+++ b/modules/core/test/ocl/test_arithm.cpp
@@ -119,7 +119,6 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
     bool use_roi;
     cv::Scalar val;
 
-    // declare Mat + UMat mirrors
     TEST_DECLARE_INPUT_PARAMETER(src1)
     TEST_DECLARE_INPUT_PARAMETER(src2)
     TEST_DECLARE_INPUT_PARAMETER(mask)
@@ -281,6 +280,614 @@ OCL_TEST_P(Subtract, Scalar_Mask)
     }
 }
 
+//////////////////////////////// Mul /////////////////////////////////////////////////
+
+typedef ArithmTestBase Mul;
+
+OCL_TEST_P(Mul, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::multiply(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::multiply(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Mul, DISABLED_Scalar)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::multiply(src1_roi, val, dst1_roi));
+        OCL_ON(cv::multiply(val, usrc1_roi, udst1_roi));
+
+        Near(udst1_roi.depth() >= CV_32F ? 1e-3 : 1);
+    }
+}
+
+OCL_TEST_P(Mul, DISABLED_Mat_Scale)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::multiply(src1_roi, src2_roi, dst1_roi, val[0]));
+        OCL_ON(cv::multiply(usrc1_roi, usrc2_roi, udst1_roi, val[0]));
+
+        Near(udst1_roi.depth() >= CV_32F ? 1e-3 : 1);
+    }
+}
+
+//////////////////////////////// Div /////////////////////////////////////////////////
+
+typedef ArithmTestBase Div;
+
+OCL_TEST_P(Div, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::divide(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::divide(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(1);
+    }
+}
+
+OCL_TEST_P(Div, DISABLED_Scalar)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::divide(val, src1_roi, dst1_roi));
+        OCL_ON(cv::divide(val, usrc1_roi, 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++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::divide(src1_roi, src2_roi, dst1_roi, val[0]));
+        OCL_ON(cv::divide(usrc1_roi, usrc2_roi, udst1_roi, val[0]));
+
+        Near(udst1_roi.depth() >= CV_32F ? 4e-3 : 1);
+    }
+}
+
+
+OCL_TEST_P(Div, DISABLED_Mat_Scalar_Scale)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::divide(src1_roi, val, dst1_roi, val[0]));
+        OCL_ON(cv::divide(usrc1_roi, val, udst1_roi, val[0]));
+
+        Near(udst1_roi.depth() >= CV_32F ? 4e-3 : 1);
+    }
+}
+
+//////////////////////////////// Min/Max /////////////////////////////////////////////////
+
+typedef ArithmTestBase Min;
+
+OCL_TEST_P(Min, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::max(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::max(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+typedef ArithmTestBase Max;
+
+OCL_TEST_P(Max, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::min(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::min(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+//////////////////////////////// Absdiff /////////////////////////////////////////////////
+
+typedef ArithmTestBase Absdiff;
+
+OCL_TEST_P(Absdiff, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::absdiff(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::absdiff(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Absdiff, Scalar)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::absdiff(src1_roi, val, dst1_roi));
+        OCL_ON(cv::absdiff(usrc1_roi, val, udst1_roi));
+        Near(1e-5);
+    }
+}
+
+//////////////////////////////// CartToPolar /////////////////////////////////////////////////
+
+typedef ArithmTestBase CartToPolar;
+
+OCL_TEST_P(CartToPolar, angleInDegree)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::cartToPolar(src1_roi, src2_roi, dst1_roi, dst2_roi, true));
+        OCL_ON(cv::cartToPolar(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi, true));
+        Near(0.5);
+        Near1(0.5);
+    }
+}
+
+OCL_TEST_P(CartToPolar, angleInRadians)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::cartToPolar(src1_roi, src2_roi, dst1_roi, dst2_roi));
+        OCL_ON(cv::cartToPolar(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi));
+        Near(0.5);
+        Near1(0.5);
+    }
+}
+
+//////////////////////////////// PolarToCart /////////////////////////////////////////////////
+
+typedef ArithmTestBase PolarToCart;
+
+OCL_TEST_P(PolarToCart, angleInDegree)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::polarToCart(src1_roi, src2_roi, dst1_roi, dst2_roi, true));
+        OCL_ON(cv::polarToCart(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi, true));
+
+        Near(0.5);
+        Near1(0.5);
+    }
+}
+
+OCL_TEST_P(PolarToCart, angleInRadians)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::polarToCart(src1_roi, src2_roi, dst1_roi, dst2_roi));
+        OCL_ON(cv::polarToCart(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi));
+
+        Near(0.5);
+        Near1(0.5);
+    }
+}
+
+//////////////////////////////// Transpose /////////////////////////////////////////////////
+
+typedef ArithmTestBase Transpose;
+
+OCL_TEST_P(Transpose, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::transpose(src1_roi, dst1_roi));
+        OCL_ON(cv::transpose(usrc1_roi, udst1_roi));
+
+        Near(1e-5);
+    }
+}
+
+OCL_TEST_P(Transpose, SquareInplace)
+{
+    const int type = CV_MAKE_TYPE(depth, cn);
+
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        Size roiSize = randomSize(1, MAX_VALUE);
+        roiSize.height = roiSize.width; // make it square
+
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src1, src1_roi, roiSize, srcBorder, type, 5, 16);
+
+        UMAT_UPLOAD_OUTPUT_PARAMETER(src1);
+
+        OCL_OFF(cv::transpose(src1_roi, src1_roi));
+        OCL_ON(cv::transpose(usrc1_roi, usrc1_roi));
+
+        EXPECT_MAT_NEAR(src1, usrc1, 0.0);
+        EXPECT_MAT_NEAR(src1_roi, usrc1_roi, 0.0);
+    }
+}
+
+//////////////////////////////// Bitwise_and /////////////////////////////////////////////////
+
+typedef ArithmTestBase Bitwise_and;
+
+OCL_TEST_P(Bitwise_and, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_and(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::bitwise_and(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Bitwise_and, Mat_Mask)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_and(src1_roi, src2_roi, dst1_roi, mask_roi));
+        OCL_ON(cv::bitwise_and(usrc1_roi, usrc2_roi, udst1_roi, umask_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Bitwise_and, Scalar)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_and(src1_roi, val, dst1_roi));
+        OCL_ON(cv::bitwise_and(usrc1_roi, val, udst1_roi));
+        Near(1e-5);
+    }
+}
+
+OCL_TEST_P(Bitwise_and, Scalar_Mask)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_and(src1_roi, val, dst1_roi, mask_roi));
+        OCL_ON(cv::bitwise_and(usrc1_roi, val, udst1_roi, umask_roi));
+        Near(1e-5);
+    }
+}
+
+//////////////////////////////// Bitwise_or /////////////////////////////////////////////////
+
+typedef ArithmTestBase Bitwise_or;
+
+OCL_TEST_P(Bitwise_or, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_or(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::bitwise_or(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Bitwise_or, Mat_Mask)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_or(src1_roi, src2_roi, dst1_roi, mask_roi));
+        OCL_ON(cv::bitwise_or(usrc1_roi, usrc2_roi, udst1_roi, umask_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Bitwise_or, Scalar)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_or(src1_roi, val, dst1_roi));
+        OCL_ON(cv::bitwise_or(usrc1_roi, val, udst1_roi));
+        Near(1e-5);
+    }
+}
+
+OCL_TEST_P(Bitwise_or, Scalar_Mask)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_or(src1_roi, val, dst1_roi, mask_roi));
+        OCL_ON(cv::bitwise_or(val, usrc1_roi, udst1_roi, umask_roi));
+        Near(1e-5);
+    }
+}
+
+//////////////////////////////// Bitwise_xor /////////////////////////////////////////////////
+
+typedef ArithmTestBase Bitwise_xor;
+
+OCL_TEST_P(Bitwise_xor, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_xor(src1_roi, src2_roi, dst1_roi));
+        OCL_ON(cv::bitwise_xor(usrc1_roi, usrc2_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Bitwise_xor, Mat_Mask)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_xor(src1_roi, src2_roi, dst1_roi, mask_roi));
+        OCL_ON(cv::bitwise_xor(usrc1_roi, usrc2_roi, udst1_roi, umask_roi));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Bitwise_xor, Scalar)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_xor(src1_roi, val, dst1_roi));
+        OCL_ON(cv::bitwise_xor(usrc1_roi, val, udst1_roi));
+        Near(1e-5);
+    }
+}
+
+OCL_TEST_P(Bitwise_xor, Scalar_Mask)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_xor(src1_roi, val, dst1_roi, mask_roi));
+        OCL_ON(cv::bitwise_xor(usrc1_roi, val, udst1_roi, umask_roi));
+        Near(1e-5);
+    }
+}
+
+//////////////////////////////// Bitwise_not /////////////////////////////////////////////////
+
+typedef ArithmTestBase Bitwise_not;
+
+OCL_TEST_P(Bitwise_not, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::bitwise_not(src1_roi, dst1_roi));
+        OCL_ON(cv::bitwise_not(usrc1_roi, udst1_roi));
+        Near(0);
+    }
+}
+
+//////////////////////////////// Compare /////////////////////////////////////////////////
+
+typedef ArithmTestBase Compare;
+
+OCL_TEST_P(Compare, Mat)
+{
+    int cmp_codes[] = { CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE };
+    int cmp_num = sizeof(cmp_codes) / sizeof(int);
+
+    for (int i = 0; i < cmp_num; ++i)
+        for (int j = 0; j < test_loop_times; j++)
+        {
+            generateTestData();
+
+            OCL_OFF(cv::compare(src1_roi, src2_roi, dst1_roi, cmp_codes[i]));
+            OCL_ON(cv::compare(usrc1_roi, usrc2_roi, udst1_roi, cmp_codes[i]));
+
+            Near(0);
+        }
+}
+
+//////////////////////////////// Pow /////////////////////////////////////////////////
+
+typedef ArithmTestBase Pow;
+
+OCL_TEST_P(Pow, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+        double p = 4.5;
+
+        OCL_OFF(cv::pow(src1_roi, p, dst1_roi));
+        OCL_ON(cv::pow(usrc1_roi, p, udst1_roi));
+
+        Near(1);
+    }
+}
+
+//////////////////////////////// AddWeighted /////////////////////////////////////////////////
+
+typedef ArithmTestBase AddWeighted;
+
+OCL_TEST_P(AddWeighted, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        const double alpha = 2.0, beta = 1.0, gama = 3.0;
+
+        OCL_OFF(cv::addWeighted(src1_roi, alpha, src2_roi, beta, gama, dst1_roi));
+        OCL_ON(cv::addWeighted(usrc1_roi, alpha, usrc2_roi, beta, gama, udst1_roi));
+
+        Near(3e-4);
+    }
+}
+
+//////////////////////////////// setIdentity /////////////////////////////////////////////////
+
+typedef ArithmTestBase SetIdentity;
+
+OCL_TEST_P(SetIdentity, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::setIdentity(dst1_roi, val));
+        OCL_ON(cv::setIdentity(udst1_roi, val));
+
+        Near(0);
+    }
+}
+
+//// Repeat
+
+struct RepeatTestCase :
+        public ArithmTestBase
+{
+    int nx, ny;
+
+    virtual void generateTestData()
+    {
+        const int type = CV_MAKE_TYPE(depth, cn);
+
+        nx = 2;//randomInt(1, 4);
+        ny = 2;//randomInt(1, 4);
+
+        Size srcRoiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src1, src1_roi, srcRoiSize, srcBorder, type, 2, 11);
+
+        Size dstRoiSize(srcRoiSize.width * nx, srcRoiSize.height * ny);
+        Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst1, dst1_roi, dstRoiSize, dst1Border, type, 5, 16);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src1)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst1)
+    }
+};
+
+typedef RepeatTestCase Repeat;
+
+OCL_TEST_P(Repeat, DISABLED_Mat)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::repeat(src1_roi, ny, nx, dst1_roi));
+        OCL_ON(cv::repeat(usrc1_roi, ny, nx, udst1_roi));
+
+        Near();
+    }
+}
+
+//////////////////////////////// CountNonZero /////////////////////////////////////////////////
+
+typedef ArithmTestBase CountNonZero;
+
+OCL_TEST_P(CountNonZero, MAT)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        int cpures, gpures;
+        OCL_OFF(cpures = cv::countNonZero(src1_roi));
+        OCL_ON(gpures = cv::countNonZero(usrc1_roi));
+
+        EXPECT_EQ(cpures, gpures);
+    }
+}
+
+//////////////////////////////// Sum /////////////////////////////////////////////////
+
+typedef ArithmTestBase Sum;
+
+OCL_TEST_P(Sum, MAT)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        Scalar cpures, gpures;
+        OCL_OFF(cpures = cv::sum(src1_roi));
+        OCL_ON(gpures = cv::sum(usrc1_roi));
+
+        for (int i = 0; i < cn; ++i)
+            EXPECT_NEAR(cpures[i], gpures[i], 0.1);
+    }
+}
+
+//////////////////////////////// meanStdDev /////////////////////////////////////////////////
+
+typedef ArithmTestBase MeanStdDev;
+
+OCL_TEST_P(MeanStdDev, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        Scalar cpu_mean, cpu_stddev;
+        Scalar gpu_mean, gpu_stddev;
+
+        OCL_OFF(cv::meanStdDev(src1_roi, cpu_mean, cpu_stddev));
+        OCL_ON(cv::meanStdDev(usrc1_roi, gpu_mean, gpu_stddev));
+
+        for (int i = 0; i < cn; ++i)
+        {
+            EXPECT_NEAR(cpu_mean[i], gpu_mean[i], 0.1);
+            EXPECT_NEAR(cpu_stddev[i], gpu_stddev[i], 0.1);
+        }
+    }
+}
+
+
 //////////////////////////////////////// Log /////////////////////////////////////////
 
 typedef ArithmTestBase Log;
@@ -359,13 +966,33 @@ OCL_TEST_P(Magnitude, Mat)
 
 //////////////////////////////////////// Instantiation /////////////////////////////////////////
 
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, ::testing::Values(1, 2, 3, 4), Bool(), Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(OCL_ALL_DEPTHS, ::testing::Values(1, 2, 4), Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Subtract, Combine(OCL_ALL_DEPTHS, ::testing::Values(1, 2, 4), Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Subtract, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Mul, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Div, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Min, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Max, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Absdiff, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, CartToPolar, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, PolarToCart, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Transpose, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_and, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_not, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_xor, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_or, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Pow, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Compare, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, SetIdentity, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Repeat, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, CountNonZero, Combine(OCL_ALL_DEPTHS, testing::Values(Channels(1)), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sum, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 
 } } // namespace cvtest::ocl
 

From 3e586f42a0dea12f34378884d68a090cec42b59e Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sat, 30 Nov 2013 17:20:45 +0400
Subject: [PATCH 02/13] added cv::pow to T-API

---
 modules/core/src/mathfuncs.cpp    | 31 +++++++++++++++++++++++++++++++
 modules/core/src/opencl/arithm.cl |  5 ++++-
 2 files changed, 35 insertions(+), 1 deletion(-)

diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp
index 29601bec2..2f146c855 100644
--- a/modules/core/src/mathfuncs.cpp
+++ b/modules/core/src/mathfuncs.cpp
@@ -1955,9 +1955,40 @@ static IPowFunc ipowTab[] =
     (IPowFunc)iPow32s, (IPowFunc)iPow32f, (IPowFunc)iPow64f, 0
 };
 
+static bool ocl_pow(InputArray _src, double power, OutputArray _dst)
+{
+    int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if ( !(_src.dims() <= 2 && (depth == CV_32F || depth == CV_64F)) ||
+         (depth == CV_64F && !doubleSupport) )
+        return false;
+
+    UMat src = _src.getUMat();
+    _dst.create(src.size(), type);
+    UMat dst = _dst.getUMat();
+
+    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
+                  format("-D dstT=%s -D OP_POW -D UNARY_OP%s", ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
+                         doubleSupport ? "-D DOUBLE_SUPPORT" : ""));
+
+    ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
+            dstarg = ocl::KernelArg::WriteOnly(dst, cn);
+
+    if (depth == CV_32F)
+        k.args(srcarg, dstarg, (float)power);
+    else
+        k.args(srcarg, dstarg, power);
+
+    size_t globalsize[2] = { dst.cols *  cn, dst.rows };
+    return k.run(2, globalsize, NULL, false);
+}
 
 void pow( InputArray _src, double power, OutputArray _dst )
 {
+    if (ocl::useOpenCL() && _dst.isUMat() && ocl_pow(_src, power, _dst))
+        return;
+
     Mat src = _src.getMat();
     int type = src.type(), depth = src.depth(), cn = src.channels();
 
diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl
index 87a9dece6..3b0fb99cb 100644
--- a/modules/core/src/opencl/arithm.cl
+++ b/modules/core/src/opencl/arithm.cl
@@ -171,6 +171,9 @@
 #elif defined OP_EXP
 #define PROCESS_ELEM dstelem = exp(srcelem1)
 
+#elif defined OP_POW
+#define PROCESS_ELEM dstelem = pow(srcelem1, srcelem2)
+
 #elif defined OP_SQRT
 #define PROCESS_ELEM dstelem = sqrt(srcelem1)
 
@@ -198,7 +201,7 @@ 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_ABSDIFF || defined OP_CMP || defined OP_MIN || defined OP_MAX || defined OP_POW
     #undef EXTRA_PARAMS
     #define EXTRA_PARAMS , workT srcelem2
 #endif

From 6ab297718d705502855ea530afe2eda08a6e0bd2 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sat, 30 Nov 2013 17:40:03 +0400
Subject: [PATCH 03/13] added setIdentity to T-API

---
 modules/core/src/matrix.cpp             | 28 +++++++++++-
 modules/core/src/opencl/set_identity.cl | 59 +++++++++++++++++++++++++
 2 files changed, 86 insertions(+), 1 deletion(-)
 create mode 100644 modules/core/src/opencl/set_identity.cl

diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp
index 919452dc9..689b57d88 100644
--- a/modules/core/src/matrix.cpp
+++ b/modules/core/src/matrix.cpp
@@ -41,6 +41,7 @@
 //M*/
 
 #include "precomp.hpp"
+#include "opencl_kernels.hpp"
 
 /****************************************************************************************\
 *                           [scaled] Identity matrix initialization                      *
@@ -2368,10 +2369,35 @@ void cv::vconcat(InputArray _src, OutputArray dst)
 }
 
 //////////////////////////////////////// set identity ////////////////////////////////////////////
+
+namespace cv {
+
+static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s )
+{
+    int type = _m.type(), cn = CV_MAT_CN(type);
+    if (cn == 3)
+        return false;
+
+    UMat m = _m.getUMat();
+
+    ocl::Kernel k("setIdentity", ocl::core::set_identity_oclsrc,
+                  format("-D T=%s", ocl::memopTypeToStr(type)));
+    k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, type, s)));
+
+    size_t globalsize[2] = { m.cols, m.rows };
+    return k.run(2, globalsize, NULL, false);
+}
+
+}
+
 void cv::setIdentity( InputOutputArray _m, const Scalar& s )
 {
+    CV_Assert( _m.dims() <= 2 );
+
+    if (ocl::useOpenCL() && _m.isUMat() && ocl_setIdentity(_m, s))
+        return;
+
     Mat m = _m.getMat();
-    CV_Assert( m.dims <= 2 );
     int i, j, rows = m.rows, cols = m.cols, type = m.type();
 
     if( type == CV_32FC1 )
diff --git a/modules/core/src/opencl/set_identity.cl b/modules/core/src/opencl/set_identity.cl
new file mode 100644
index 000000000..de8caaf85
--- /dev/null
+++ b/modules/core/src/opencl/set_identity.cl
@@ -0,0 +1,59 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Jin Ma jin@multicorewareinc.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+__kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols,
+                          T scalar)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols && y < rows)
+    {
+        int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(T));
+        __global T * src = (__global T *)(srcptr + src_index);
+
+        src[0] = x == y ? scalar : (T)(0);
+    }
+}

From 435028ccee31b3792d166ab1a608e0bd92e29008 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sat, 30 Nov 2013 18:42:09 +0400
Subject: [PATCH 04/13] added cv::transpose to T-API

---
 modules/core/src/matrix.cpp          |  50 ++++++++++-
 modules/core/src/opencl/transpose.cl | 124 +++++++++++++++++++++++++++
 2 files changed, 172 insertions(+), 2 deletions(-)
 create mode 100644 modules/core/src/opencl/transpose.cl

diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp
index 689b57d88..94ce7aff9 100644
--- a/modules/core/src/matrix.cpp
+++ b/modules/core/src/matrix.cpp
@@ -2574,18 +2574,63 @@ static TransposeInplaceFunc transposeInplaceTab[] =
     0, 0, 0, 0, 0, 0, 0, transposeI_32sC6, 0, 0, 0, 0, 0, 0, 0, transposeI_32sC8
 };
 
+static inline int divUp(int a, int b)
+{
+    return (a + b - 1) / b;
+}
+
+static bool ocl_transpose( InputArray _src, OutputArray _dst )
+{
+    const int TILE_DIM = 32, BLOCK_ROWS = 8;
+    int type = _src.type(), cn = CV_MAT_CN(type);
+
+    if (cn == 3)
+        return false;
+
+    UMat src = _src.getUMat();
+    _dst.create(src.cols, src.rows, type);
+    UMat dst = _dst.getUMat();
+
+    String kernelName("transpose");
+    bool inplace = dst.u == src.u;
+
+    if (inplace)
+    {
+        CV_Assert(dst.cols == dst.rows);
+        kernelName += "_inplace";
+    }
+
+    ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc,
+                  format("-D T=%s -D TILE_DIM=%d -D BLOCK_ROWS=%d",
+                         ocl::memopTypeToStr(type), TILE_DIM, BLOCK_ROWS));
+    if (inplace)
+        k.args(ocl::KernelArg::ReadWriteNoSize(dst), dst.rows);
+    else
+        k.args(ocl::KernelArg::ReadOnly(src),
+               ocl::KernelArg::WriteOnlyNoSize(dst));
+
+    size_t localsize[3]  = { TILE_DIM, BLOCK_ROWS, 1 };
+    size_t globalsize[3] = { src.cols, inplace ? src.rows : divUp(src.rows, TILE_DIM) * BLOCK_ROWS, 1 };
+
+    return k.run(2, globalsize, localsize, false);
+}
+
 }
 
 void cv::transpose( InputArray _src, OutputArray _dst )
 {
+    int type = _src.type(), esz = CV_ELEM_SIZE(type);
+    CV_Assert( _src.dims() <= 2 && esz <= 32 );
+
+    if (ocl::useOpenCL() && _dst.isUMat() && ocl_transpose(_src, _dst))
+        return;
+
     Mat src = _src.getMat();
     if( src.empty() )
     {
         _dst.release();
         return;
     }
-    size_t esz = src.elemSize();
-    CV_Assert( src.dims <= 2 && esz <= (size_t)32 );
 
     _dst.create(src.cols, src.rows, src.type());
     Mat dst = _dst.getMat();
@@ -2602,6 +2647,7 @@ void cv::transpose( InputArray _src, OutputArray _dst )
     {
         TransposeInplaceFunc func = transposeInplaceTab[esz];
         CV_Assert( func != 0 );
+//        CV_Assert( dst.cols == dst.rows );
         func( dst.data, dst.step, dst.rows );
     }
     else
diff --git a/modules/core/src/opencl/transpose.cl b/modules/core/src/opencl/transpose.cl
new file mode 100644
index 000000000..da9608c0d
--- /dev/null
+++ b/modules/core/src/opencl/transpose.cl
@@ -0,0 +1,124 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Jia Haipeng, jiahaipeng95@gmail.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#define LDS_STEP      TILE_DIM
+
+__kernel void transpose(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
+                        __global uchar * dstptr, int dst_step, int dst_offset)
+{
+    int gp_x = get_group_id(0),   gp_y = get_group_id(1);
+    int gs_x = get_num_groups(0), gs_y = get_num_groups(1);
+
+    int groupId_x, groupId_y;
+
+    if (src_rows == src_cols)
+    {
+        groupId_y = gp_x;
+        groupId_x = (gp_x + gp_y) % gs_x;
+    }
+    else
+    {
+        int bid = gp_x + gs_x * gp_y;
+        groupId_y =  bid % gs_y;
+        groupId_x = ((bid / gs_y) + groupId_y) % gs_x;
+    }
+
+    int lx = get_local_id(0);
+    int ly = get_local_id(1);
+
+    int x = groupId_x * TILE_DIM + lx;
+    int y = groupId_y * TILE_DIM + ly;
+
+    int x_index = groupId_y * TILE_DIM + lx;
+    int y_index = groupId_x * TILE_DIM + ly;
+
+    __local T title[TILE_DIM * LDS_STEP];
+
+    if (x < src_cols && y < src_rows)
+    {
+        int index_src = mad24(y, src_step, x * (int)sizeof(T) + src_offset);
+
+        for (int i = 0; i < TILE_DIM; i += BLOCK_ROWS)
+            if (y + i < src_rows)
+            {
+                __global const T * src = (__global const T *)(srcptr + index_src);
+                title[(ly + i) * LDS_STEP + lx] = src[0];
+                index_src = mad24(BLOCK_ROWS, src_step, index_src);
+            }
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (x_index < src_rows && y_index < src_cols)
+    {
+        int index_dst = mad24(y_index, dst_step, x_index * (int)sizeof(T) + dst_offset);
+
+        for (int i = 0; i < TILE_DIM; i += BLOCK_ROWS)
+            if ((y_index + i) < src_cols)
+            {
+                __global T * dst = (__global T *)(dstptr + index_dst);
+                dst[0] = title[lx * LDS_STEP + ly + i];
+                index_dst = mad24(BLOCK_ROWS, dst_step, index_dst);
+            }
+    }
+}
+
+__kernel void transpose_inplace(__global uchar * srcptr, int src_step, int src_offset, int src_rows)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < src_rows && x < y)
+    {
+        int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(T));
+        int dst_index = mad24(x, src_step, src_offset + y * (int)sizeof(T));
+
+        __global T * src = (__global T *)(srcptr + src_index);
+        __global T * dst = (__global T *)(srcptr + dst_index);
+
+        T tmp = dst[0];
+        dst[0] = src[0];
+        src[0] = tmp;
+    }
+}

From ba850f0b64305acb5b9e34c0e725304e2ca49efa Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sat, 30 Nov 2013 20:05:54 +0400
Subject: [PATCH 05/13] added cv::cartToPolar to T-API

---
 modules/core/src/mathfuncs.cpp    | 47 +++++++++++++++++++++++++++++--
 modules/core/src/matrix.cpp       |  2 +-
 modules/core/src/ocl.cpp          |  2 +-
 modules/core/src/opencl/arithm.cl | 36 +++++++++++++++++++++--
 4 files changed, 80 insertions(+), 7 deletions(-)

diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp
index 2f146c855..d216589e7 100644
--- a/modules/core/src/mathfuncs.cpp
+++ b/modules/core/src/mathfuncs.cpp
@@ -497,10 +497,50 @@ void phase( InputArray src1, InputArray src2, OutputArray dst, bool angleInDegre
     }
 }
 
+static bool ocl_cartToPolar( InputArray _src1, InputArray _src2,
+                             OutputArray _dst1, OutputArray _dst2, bool angleInDegrees )
+{
+    int type = _src1.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if ( !(_src1.dims() <= 2 && _src2.dims() <= 2 &&
+           (depth == CV_32F || depth == CV_64F) && type == _src2.type()) ||
+         (depth == CV_64F && !doubleSupport) )
+        return false;
+
+    UMat src1 = _src1.getUMat(), src2 = _src2.getUMat();
+    Size size = src1.size();
+    CV_Assert( size == src2.size() );
+
+    _dst1.create(size, type);
+    _dst2.create(size, type);
+    UMat dst1 = _dst1.getUMat(), dst2 = _dst2.getUMat();
+
+    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
+                  format("-D BINARY_OP -D dstT=%s -D OP_CTP_%s%s",
+                         ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
+                         angleInDegrees ? "AD" : "AR",
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+
+    k.args(ocl::KernelArg::ReadOnlyNoSize(src1),
+           ocl::KernelArg::ReadOnlyNoSize(src2),
+           ocl::KernelArg::WriteOnly(dst1, cn),
+           ocl::KernelArg::WriteOnlyNoSize(dst2));
+
+    size_t globalsize[2] = { dst1.cols * cn, dst1.rows };
+    return k.run(2, globalsize, NULL, false);
+}
 
 void cartToPolar( InputArray src1, InputArray src2,
                   OutputArray dst1, OutputArray dst2, bool angleInDegrees )
 {
+    if (ocl::useOpenCL() && dst1.isUMat() && dst2.isUMat() /*&&
+            ocl_cartToPolar(src1, src2, dst1, dst2, angleInDegrees)*/)
+    {
+        CV_Assert(ocl_cartToPolar(src1, src2, dst1, dst2, angleInDegrees));
+        return;
+    }
+
     Mat X = src1.getMat(), Y = src2.getMat();
     int type = X.type(), depth = X.depth(), cn = X.channels();
     CV_Assert( X.size == Y.size && type == Y.type() && (depth == CV_32F || depth == CV_64F));
@@ -1970,7 +2010,7 @@ static bool ocl_pow(InputArray _src, double power, OutputArray _dst)
 
     ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
                   format("-D dstT=%s -D OP_POW -D UNARY_OP%s", ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
-                         doubleSupport ? "-D DOUBLE_SUPPORT" : ""));
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
 
     ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
             dstarg = ocl::KernelArg::WriteOnly(dst, cn);
@@ -1986,8 +2026,11 @@ static bool ocl_pow(InputArray _src, double power, OutputArray _dst)
 
 void pow( InputArray _src, double power, OutputArray _dst )
 {
-    if (ocl::useOpenCL() && _dst.isUMat() && ocl_pow(_src, power, _dst))
+    if (ocl::useOpenCL() && _dst.isUMat() /*&& ocl_pow(_src, power, _dst)*/)
+    {
+        CV_Assert(ocl_pow(_src, power, _dst));
         return;
+    }
 
     Mat src = _src.getMat();
     int type = src.type(), depth = src.depth(), cn = src.channels();
diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp
index 94ce7aff9..b2b164e45 100644
--- a/modules/core/src/matrix.cpp
+++ b/modules/core/src/matrix.cpp
@@ -2647,7 +2647,7 @@ void cv::transpose( InputArray _src, OutputArray _dst )
     {
         TransposeInplaceFunc func = transposeInplaceTab[esz];
         CV_Assert( func != 0 );
-//        CV_Assert( dst.cols == dst.rows );
+        CV_Assert( dst.cols == dst.rows );
         func( dst.data, dst.step, dst.rows );
     }
     else
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index 64460efb0..4ef3fdad2 100644
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -1893,7 +1893,7 @@ Context2& Context2::getDefault()
         // First, try to retrieve existing context of the same type.
         // In its turn, Platform::getContext() may call Context2::create()
         // if there is no such context.
-        ctx.create(Device::TYPE_ACCELERATOR);
+        ctx.create(Device::TYPE_CPU);
         if(!ctx.p)
             ctx.create(Device::TYPE_DGPU);
         if(!ctx.p)
diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl
index 3b0fb99cb..48d5f6fb4 100644
--- a/modules/core/src/opencl/arithm.cl
+++ b/modules/core/src/opencl/arithm.cl
@@ -63,12 +63,15 @@
 #elif defined (cl_amd_fp64)
 #pragma OPENCL EXTENSION cl_amd_fp64:enable
 #endif
+#define CV_EPSILON DBL_EPSILON
+#define CV_PI M_PI
+#else
+#define CV_EPSILON FLT_EPSILON
+#define CV_PI M_PI_F
 #endif
 
-#define CV_32S 4
-#define CV_32F 5
-
 #define dstelem *(__global dstT*)(dstptr + dst_index)
+#define dstelem2 *(__global dstT*)(dstptr2 + dst_index2)
 #define noconvert
 
 #ifndef workT
@@ -88,6 +91,7 @@
 #endif
 
 #define EXTRA_PARAMS
+#define EXTRA_INDEX
 
 #if defined OP_ADD
 #define PROCESS_ELEM dstelem = convertToDT(srcelem1 + srcelem2)
@@ -193,10 +197,35 @@ dstelem = v > (dstT)(0) ? log(v) : log(-v)
 #define EXTRA_PARAMS , workT alpha, workT beta
 #define PROCESS_ELEM dstelem = convertToDT(srcelem1*alpha + beta)
 
+#elif defined OP_CTP_AD || defined OP_CTP_AR
+#ifdef OP_CTP_AD
+#define TO_DEGREE cartToPolar *= (180 / CV_PI);
+#elif defined OP_CTP_AR
+#define TO_DEGREE
+#endif
+#define PROCESS_ELEM \
+    dstT x = srcelem1, y = srcelem2; \
+    dstT x2 = x * x, y2 = y * y; \
+    dstT magnitude = sqrt(x2 + y2); \
+    dstT tmp = y >= 0 ? 0 : CV_PI * 2; \
+    tmp = x < 0 ? CV_PI : tmp; \
+    dstT tmp1 = y >= 0 ? CV_PI * 0.5f : CV_PI * 1.5f; \
+    dstT cartToPolar = y2 <= x2 ? x * y / (x2 + 0.28f * y2 + CV_EPSILON) + tmp : (tmp1 - x * y / (y2 + 0.28f * x2 + CV_EPSILON)); \
+    TO_DEGREE \
+    dstelem = magnitude; \
+    dstelem2 = cartToPolar
+
 #else
 #error "unknown op type"
 #endif
 
+#if defined OP_CTP_AD || defined OP_CTP_AR
+    #undef EXTRA_PARAMS
+    #define EXTRA_PARAMS , __global uchar* dstptr2, int dststep2, int dstoffset2
+    #undef EXTRA_INDEX
+    #define EXTRA_INDEX int dst_index2 = mad24(y, dststep2, x*(int)sizeof(dstT) + dstoffset2)
+#endif
+
 #if defined UNARY_OP || defined MASK_UNARY_OP
 #undef srcelem2
 #if defined OP_AND || defined OP_OR || defined OP_XOR || defined OP_ADD || defined OP_SAT_ADD || \
@@ -222,6 +251,7 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1,
         int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1);
         int src2_index = mad24(y, srcstep2, x*(int)sizeof(srcT2) + srcoffset2);
         int dst_index  = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset);
+        EXTRA_INDEX;
 
         PROCESS_ELEM;
     }

From a180cebb2dbf920205ecf9881994bda0f01d8768 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sun, 1 Dec 2013 00:17:23 +0400
Subject: [PATCH 06/13] added cv::polarToCart to T-API

---
 modules/core/src/mathfuncs.cpp    | 40 +++++++++++++++++++++++++++++--
 modules/core/src/opencl/arithm.cl | 18 +++++++++++++-
 2 files changed, 55 insertions(+), 3 deletions(-)

diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp
index d216589e7..540bb87f5 100644
--- a/modules/core/src/mathfuncs.cpp
+++ b/modules/core/src/mathfuncs.cpp
@@ -684,12 +684,48 @@ static void SinCos_32f( const float *angle, float *sinval, float* cosval,
 }
 
 
+static bool ocl_polarToCart( InputArray _mag, InputArray _angle,
+                             OutputArray _dst1, OutputArray _dst2, bool angleInDegrees )
+{
+    int type = _angle.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if ( _mag.empty() || (!doubleSupport && depth == CV_64F) )
+        return false;
+
+    UMat mag = _mag.getUMat(), angle = _angle.getUMat();
+    Size size = angle.size();
+    CV_Assert(mag.size() == size);
+
+    _dst1.create(size, type);
+    _dst2.create(size, type);
+    UMat dst1 = _dst1.getUMat(), dst2 = _dst2.getUMat();
+
+    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
+                  format("-D dstT=%s -D BINARY_OP -D OP_PTC_%s%s",
+                         ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
+                         angleInDegrees ? "AD" : "AR",
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+
+    k.args(ocl::KernelArg::ReadOnlyNoSize(mag), ocl::KernelArg::ReadOnlyNoSize(angle),
+           ocl::KernelArg::WriteOnly(dst1, cn), ocl::KernelArg::WriteOnlyNoSize(dst2));
+
+    size_t globalsize[2] = { dst1.cols * cn, dst1.rows };
+    return k.run(2, globalsize, NULL, false);
+}
+
 void polarToCart( InputArray src1, InputArray src2,
                   OutputArray dst1, OutputArray dst2, bool angleInDegrees )
 {
+    int type = src2.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+    CV_Assert((depth == CV_32F || depth == CV_64F) && (src1.empty() || src1.type() == type));
+
+    if (ocl::useOpenCL() && dst1.isUMat() && dst2.isUMat() &&
+            ocl_polarToCart(src1, src2, dst1, dst2, angleInDegrees))
+        return;
+
     Mat Mag = src1.getMat(), Angle = src2.getMat();
-    int type = Angle.type(), depth = Angle.depth(), cn = Angle.channels();
-    CV_Assert( Mag.empty() || (Angle.size == Mag.size && type == Mag.type() && (depth == CV_32F || depth == CV_64F)));
+    CV_Assert( Mag.empty() || Angle.size == Mag.size);
     dst1.create( Angle.dims, Angle.size, type );
     dst2.create( Angle.dims, Angle.size, type );
     Mat X = dst1.getMat(), Y = dst2.getMat();
diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl
index 48d5f6fb4..b9f47a795 100644
--- a/modules/core/src/opencl/arithm.cl
+++ b/modules/core/src/opencl/arithm.cl
@@ -215,11 +215,26 @@ dstelem = v > (dstT)(0) ? log(v) : log(-v)
     dstelem = magnitude; \
     dstelem2 = cartToPolar
 
+#elif defined OP_PTC_AD || defined OP_PTC_AR
+#ifdef OP_PTC_AD
+#define FROM_DEGREE \
+    dstT ascale = CV_PI/180.0f; \
+    dstT alpha = y * ascale
+#else
+#define FROM_DEGREE \
+    dstT alpha = y
+#endif
+#define PROCESS_ELEM \
+    dstT x = srcelem1, y = srcelem2; \
+    FROM_DEGREE; \
+    dstelem = cos(alpha) * x; \
+    dstelem2 = sin(alpha) * x
+
 #else
 #error "unknown op type"
 #endif
 
-#if defined OP_CTP_AD || defined OP_CTP_AR
+#if defined OP_CTP_AD || defined OP_CTP_AR || defined OP_PTC_AD || defined OP_PTC_AR
     #undef EXTRA_PARAMS
     #define EXTRA_PARAMS , __global uchar* dstptr2, int dststep2, int dstoffset2
     #undef EXTRA_INDEX
@@ -295,6 +310,7 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1,
     {
         int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1);
         int dst_index  = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset);
+        EXTRA_INDEX;
 
         PROCESS_ELEM;
     }

From e17107e864663ed6f11a5915daef91694827ab4a Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sun, 1 Dec 2013 00:59:55 +0400
Subject: [PATCH 07/13] added cv::compare to T-API

---
 modules/core/src/arithm.cpp       | 45 +++++++++++++++++++++++++++++--
 modules/core/src/opencl/arithm.cl |  4 +++
 2 files changed, 47 insertions(+), 2 deletions(-)

diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp
index 0d510b8ad..f18f74b2a 100644
--- a/modules/core/src/arithm.cpp
+++ b/modules/core/src/arithm.cpp
@@ -929,7 +929,7 @@ static bool ocl_binary_op(InputArray _src1, InputArray _src2, OutputArray _dst,
     int srcdepth = CV_MAT_DEPTH(srctype);
     int cn = CV_MAT_CN(srctype);
 
-    if( oclop < 0 || ((haveMask || haveScalar) && cn > 4) || cn == 3)
+    if( oclop < 0 || ((haveMask || haveScalar) && (cn > 4 || cn == 3)) )
         return false;
 
     char opts[1024];
@@ -1284,7 +1284,7 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
     int type1 = _src1.type(), depth1 = CV_MAT_DEPTH(type1), cn = CV_MAT_CN(type1);
     bool haveMask = !_mask.empty();
 
-    if( ((haveMask || haveScalar) && cn > 4) || cn == 3)
+    if( ((haveMask || haveScalar) && (cn > 4 || cn == 3)) )
         return false;
 
     int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), wdepth = std::max(CV_32S, CV_MAT_DEPTH(wtype));
@@ -2585,6 +2585,40 @@ static double getMaxVal(int depth)
     return tab[depth];
 }
 
+static bool ocl_compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op)
+{
+    if ( !((_src1.isMat() || _src1.isUMat()) && (_src2.isMat() || _src2.isUMat())) )
+        return false;
+
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+    int type = _src1.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), type2 = _src2.type();
+    if (!doubleSupport && (depth == CV_64F || _src2.depth() == CV_64F))
+        return false;
+
+    CV_Assert(type == type2);
+    UMat src1 = _src1.getUMat(), src2 = _src2.getUMat();
+    Size size = src1.size();
+    CV_Assert(size == src2.size());
+
+    _dst.create(size, CV_8UC(cn));
+    UMat dst = _dst.getUMat();
+
+    const char * const operationMap[] = { "==", ">", ">=", "<", "<=", "!=" };
+    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
+                  format("-D BINARY_OP -D srcT1=%s -D workT=srcT1"
+                         " -D OP_CMP -D CMP_OPERATOR=%s%s",
+                         ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
+                         operationMap[op],
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+
+    k.args(ocl::KernelArg::ReadOnlyNoSize(src1),
+           ocl::KernelArg::ReadOnlyNoSize(src2),
+           ocl::KernelArg::WriteOnly(dst, cn));
+
+    size_t globalsize[2] = { dst.cols * cn, dst.rows };
+    return k.run(2, globalsize, NULL, false);
+}
+
 }
 
 void cv::compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op)
@@ -2592,6 +2626,13 @@ void cv::compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op)
     CV_Assert( op == CMP_LT || op == CMP_LE || op == CMP_EQ ||
                op == CMP_NE || op == CMP_GE || op == CMP_GT );
 
+    if (ocl::useOpenCL() && _dst.isUMat() /*&&
+            ocl_compare(_src1, _src2, _dst, op)*/)
+    {
+        CV_Assert(ocl_compare(_src1, _src2, _dst, op));
+        return;
+    }
+
     int kind1 = _src1.kind(), kind2 = _src2.kind();
     Mat src1 = _src1.getMat(), src2 = _src2.getMat();
 
diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl
index b9f47a795..b4cdb53f2 100644
--- a/modules/core/src/opencl/arithm.cl
+++ b/modules/core/src/opencl/arithm.cl
@@ -187,6 +187,10 @@ dstT v = (dstT)(srcelem1);\
 dstelem = v > (dstT)(0) ? log(v) : log(-v)
 
 #elif defined OP_CMP
+#define dstT uchar
+#define srcT2 srcT1
+#define convertToWT1
+#define convertToWT2
 #define PROCESS_ELEM dstelem = convert_uchar(srcelem1 CMP_OPERATOR srcelem2 ? 255 : 0)
 
 #elif defined OP_CONVERT

From fed7e618798935d54d3dcb0839ca9b1bc43740fd Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sun, 1 Dec 2013 01:57:13 +0400
Subject: [PATCH 08/13] added cv::repeat to T-API (fix for copyTo -> UMat is
 needed)

---
 modules/core/src/arithm.cpp    |  7 ++-----
 modules/core/src/mathfuncs.cpp | 12 +++---------
 modules/core/src/ocl.cpp       |  2 +-
 3 files changed, 6 insertions(+), 15 deletions(-)

diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp
index f18f74b2a..c1e27174b 100644
--- a/modules/core/src/arithm.cpp
+++ b/modules/core/src/arithm.cpp
@@ -2626,12 +2626,9 @@ void cv::compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op)
     CV_Assert( op == CMP_LT || op == CMP_LE || op == CMP_EQ ||
                op == CMP_NE || op == CMP_GE || op == CMP_GT );
 
-    if (ocl::useOpenCL() && _dst.isUMat() /*&&
-            ocl_compare(_src1, _src2, _dst, op)*/)
-    {
-        CV_Assert(ocl_compare(_src1, _src2, _dst, op));
+    if (ocl::useOpenCL() && _dst.isUMat() &&
+            ocl_compare(_src1, _src2, _dst, op))
         return;
-    }
 
     int kind1 = _src1.kind(), kind2 = _src2.kind();
     Mat src1 = _src1.getMat(), src2 = _src2.getMat();
diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp
index 540bb87f5..d75fad122 100644
--- a/modules/core/src/mathfuncs.cpp
+++ b/modules/core/src/mathfuncs.cpp
@@ -534,12 +534,9 @@ static bool ocl_cartToPolar( InputArray _src1, InputArray _src2,
 void cartToPolar( InputArray src1, InputArray src2,
                   OutputArray dst1, OutputArray dst2, bool angleInDegrees )
 {
-    if (ocl::useOpenCL() && dst1.isUMat() && dst2.isUMat() /*&&
-            ocl_cartToPolar(src1, src2, dst1, dst2, angleInDegrees)*/)
-    {
-        CV_Assert(ocl_cartToPolar(src1, src2, dst1, dst2, angleInDegrees));
+    if (ocl::useOpenCL() && dst1.isUMat() && dst2.isUMat() &&
+            ocl_cartToPolar(src1, src2, dst1, dst2, angleInDegrees))
         return;
-    }
 
     Mat X = src1.getMat(), Y = src2.getMat();
     int type = X.type(), depth = X.depth(), cn = X.channels();
@@ -2062,11 +2059,8 @@ static bool ocl_pow(InputArray _src, double power, OutputArray _dst)
 
 void pow( InputArray _src, double power, OutputArray _dst )
 {
-    if (ocl::useOpenCL() && _dst.isUMat() /*&& ocl_pow(_src, power, _dst)*/)
-    {
-        CV_Assert(ocl_pow(_src, power, _dst));
+    if (ocl::useOpenCL() && _dst.isUMat() && ocl_pow(_src, power, _dst))
         return;
-    }
 
     Mat src = _src.getMat();
     int type = src.type(), depth = src.depth(), cn = src.channels();
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index 4ef3fdad2..64460efb0 100644
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -1893,7 +1893,7 @@ Context2& Context2::getDefault()
         // First, try to retrieve existing context of the same type.
         // In its turn, Platform::getContext() may call Context2::create()
         // if there is no such context.
-        ctx.create(Device::TYPE_CPU);
+        ctx.create(Device::TYPE_ACCELERATOR);
         if(!ctx.p)
             ctx.create(Device::TYPE_DGPU);
         if(!ctx.p)

From 28575c19699c03814e8bab5634fec51308f979f2 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sun, 1 Dec 2013 03:02:13 +0400
Subject: [PATCH 09/13] added cv::countNonZero to T-API

---
 modules/core/src/opencl/count_non_zero.cl | 100 ++++++++++++++++++++++
 modules/core/src/stat.cpp                 |  42 ++++++++-
 modules/core/test/ocl/test_arithm.cpp     |   8 +-
 3 files changed, 145 insertions(+), 5 deletions(-)
 create mode 100644 modules/core/src/opencl/count_non_zero.cl

diff --git a/modules/core/src/opencl/count_non_zero.cl b/modules/core/src/opencl/count_non_zero.cl
new file mode 100644
index 000000000..cad89eb81
--- /dev/null
+++ b/modules/core/src/opencl/count_non_zero.cl
@@ -0,0 +1,100 @@
+////////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Shengen Yan,yanshengen@gmail.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+
+#ifdef DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
+/**************************************Count NonZero**************************************/
+
+__kernel void count_non_zero(__global const uchar * srcptr, int step, int offset, int cols,
+                            int total, int groupnum, __global uchar * dstptr)
+{
+    int lid = get_local_id(0);
+    int gid = get_group_id(0);
+    int  id = get_global_id(0);
+
+    __local int localmem[WGS2_ALIGNED];
+    if (lid < WGS2_ALIGNED)
+        localmem[lid] = 0;
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    int nonzero = (int)(0), src_index;
+    srcT zero = (srcT)(0), one = (srcT)(1);
+
+    for (int grain = groupnum * WGS; id < total; id += grain)
+    {
+        src_index = mad24(id / cols, step, offset + (id % cols) * (int)sizeof(srcT));
+        __global const srcT * src = (__global const srcT *)(srcptr + src_index);
+        nonzero += src[0] == zero ? zero : one;
+    }
+
+    if (lid >= WGS2_ALIGNED)
+        localmem[lid - WGS2_ALIGNED] = nonzero;
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (lid < WGS2_ALIGNED)
+        localmem[lid] = nonzero + localmem[lid];
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
+    {
+        if (lid < lsize)
+        {
+           int lid2 = lsize + lid;
+           localmem[lid] = localmem[lid] + localmem[lid2];
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+
+    if (lid == 0)
+    {
+        __global int * dst = (__global int *)(dstptr + (int)sizeof(int) * gid);
+        dst[0] = localmem[0];
+    }
+}
diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp
index bb2e1f493..46ec20a97 100644
--- a/modules/core/src/stat.cpp
+++ b/modules/core/src/stat.cpp
@@ -41,6 +41,7 @@
 //M*/
 
 #include "precomp.hpp"
+#include "opencl_kernels.hpp"
 #include <climits>
 #include <limits>
 
@@ -542,12 +543,51 @@ cv::Scalar cv::sum( InputArray _src )
     return s;
 }
 
+namespace cv {
+
+static bool ocl_countNonZero( InputArray _src, int & res )
+{
+    int depth = _src.depth();
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if (depth == CV_64F && !doubleSupport)
+        return false;
+
+    int dbsize = ocl::Device::getDefault().maxComputeUnits();
+    size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
+    UMat src = _src.getUMat(), db(1, dbsize, CV_32SC1);
+
+    int wgs2_aligned = 1;
+    while (wgs2_aligned < (int)wgs)
+        wgs2_aligned <<= 1;
+    wgs2_aligned >>= 1;
+
+    ocl::Kernel k("count_non_zero", ocl::core::count_non_zero_oclsrc,
+                  format("-D srcT=%s -D WGS=%d -D WGS2_ALIGNED=%d%s", ocl::typeToStr(src.type()), (int)wgs,
+                         wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(),
+           dbsize, ocl::KernelArg::PtrWriteOnly(db));
+
+    size_t globalsize = dbsize * wgs;
+    if (k.run(1, &globalsize, &wgs, true))
+        return res = cv::sum(db.getMat(ACCESS_READ))[0], true;
+    return false;
+}
+
+}
+
 int cv::countNonZero( InputArray _src )
 {
+    CV_Assert( _src.channels() == 1 );
+
+    int res = -1;
+    if (ocl::useOpenCL() && _src.isUMat() && ocl_countNonZero(_src, res))
+        return res;
+
     Mat src = _src.getMat();
     CountNonZeroFunc func = getCountNonZeroTab(src.depth());
 
-    CV_Assert( src.channels() == 1 && func != 0 );
+    CV_Assert( func != 0 );
 
     const Mat* arrays[] = {&src, 0};
     uchar* ptrs[1];
diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp
index ed6414bf4..d1f2c0170 100644
--- a/modules/core/test/ocl/test_arithm.cpp
+++ b/modules/core/test/ocl/test_arithm.cpp
@@ -969,10 +969,6 @@ OCL_TEST_P(Magnitude, Mat)
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Subtract, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Mul, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Div, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Min, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
@@ -993,6 +989,10 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, Repeat, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHAN
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, CountNonZero, Combine(OCL_ALL_DEPTHS, testing::Values(Channels(1)), Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sum, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
 
 } } // namespace cvtest::ocl
 

From b8edc2cc4d4c37784736022b954d0dcdab3bac3f Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sun, 1 Dec 2013 14:23:07 +0400
Subject: [PATCH 10/13] added cv::sum to T-API

---
 .../opencl/{count_non_zero.cl => reduce.cl}   | 60 ++++++++++++----
 modules/core/src/stat.cpp                     | 69 ++++++++++++++++++-
 2 files changed, 114 insertions(+), 15 deletions(-)
 rename modules/core/src/opencl/{count_non_zero.cl => reduce.cl} (68%)

diff --git a/modules/core/src/opencl/count_non_zero.cl b/modules/core/src/opencl/reduce.cl
similarity index 68%
rename from modules/core/src/opencl/count_non_zero.cl
rename to modules/core/src/opencl/reduce.cl
index cad89eb81..2ba36e859 100644
--- a/modules/core/src/opencl/count_non_zero.cl
+++ b/modules/core/src/opencl/reduce.cl
@@ -50,36 +50,70 @@
 #endif
 #endif
 
-/**************************************Count NonZero**************************************/
+#define noconvert
 
-__kernel void count_non_zero(__global const uchar * srcptr, int step, int offset, int cols,
-                            int total, int groupnum, __global uchar * dstptr)
+#if defined OP_SUM || defined OP_SUM_ABS || defined OP_SUM_SQR
+#if OP_SUM
+#define FUNC(a, b) a += b
+#elif OP_SUM_ABS
+#define FUNC(a, b) a += b >= (dstT)(0) ? b : -b
+#elif OP_SUM_SQR
+#define FUNC(a, b) a += b * b
+#endif
+#define DEFINE_ACCUMULATOR \
+    dstT accumulator = (dstT)(0)
+#define REDUCE_GLOBAL \
+    dstT temp = convertToDT(src[0]); \
+    FUNC(accumulator, temp)
+#define REDUCE_LOCAL_1 \
+    localmem[lid] += accumulator
+#define REDUCE_LOCAL_2 \
+    localmem[lid] += localmem[lid2]
+
+#elif defined OP_COUNT_NON_ZERO
+#define dstT int
+#define DEFINE_ACCUMULATOR \
+    dstT accumulator = (dstT)(0); \
+    srcT zero = (srcT)(0), one = (srcT)(1)
+#define REDUCE_GLOBAL \
+    accumulator += src[0] == zero ? zero : one
+#define REDUCE_LOCAL_1 \
+    localmem[lid] += accumulator
+#define REDUCE_LOCAL_2 \
+    localmem[lid] += localmem[lid2]
+
+#else
+#error "No operation"
+
+#endif
+
+__kernel void reduce(__global const uchar * srcptr, int step, int offset, int cols,
+                     int total, int groupnum, __global uchar * dstptr)
 {
     int lid = get_local_id(0);
     int gid = get_group_id(0);
     int  id = get_global_id(0);
 
-    __local int localmem[WGS2_ALIGNED];
+    __local dstT localmem[WGS2_ALIGNED];
     if (lid < WGS2_ALIGNED)
-        localmem[lid] = 0;
+        localmem[lid] = (dstT)(0);
     barrier(CLK_LOCAL_MEM_FENCE);
 
-    int nonzero = (int)(0), src_index;
-    srcT zero = (srcT)(0), one = (srcT)(1);
+    DEFINE_ACCUMULATOR;
 
     for (int grain = groupnum * WGS; id < total; id += grain)
     {
-        src_index = mad24(id / cols, step, offset + (id % cols) * (int)sizeof(srcT));
+        int src_index = mad24(id / cols, step, offset + (id % cols) * (int)sizeof(srcT));
         __global const srcT * src = (__global const srcT *)(srcptr + src_index);
-        nonzero += src[0] == zero ? zero : one;
+        REDUCE_GLOBAL;
     }
 
     if (lid >= WGS2_ALIGNED)
-        localmem[lid - WGS2_ALIGNED] = nonzero;
+        localmem[lid - WGS2_ALIGNED] = accumulator;
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (lid < WGS2_ALIGNED)
-        localmem[lid] = nonzero + localmem[lid];
+        REDUCE_LOCAL_1;
     barrier(CLK_LOCAL_MEM_FENCE);
 
     for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
@@ -87,14 +121,14 @@ __kernel void count_non_zero(__global const uchar * srcptr, int step, int offset
         if (lid < lsize)
         {
            int lid2 = lsize + lid;
-           localmem[lid] = localmem[lid] + localmem[lid2];
+           REDUCE_LOCAL_2;
         }
         barrier(CLK_LOCAL_MEM_FENCE);
     }
 
     if (lid == 0)
     {
-        __global int * dst = (__global int *)(dstptr + (int)sizeof(int) * gid);
+        __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid);
         dst[0] = localmem[0];
     }
 }
diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp
index 46ec20a97..f8ff0c29d 100644
--- a/modules/core/src/stat.cpp
+++ b/modules/core/src/stat.cpp
@@ -449,10 +449,74 @@ static SumSqrFunc getSumSqrTab(int depth)
     return sumSqrTab[depth];
 }
 
+template <typename T> Scalar ocl_part_sum(Mat m)
+{
+    CV_Assert(m.rows == 1);
+
+    Scalar s = Scalar::all(0);
+    int cn = m.channels();
+    const T * const ptr = m.ptr<T>(0);
+
+    for (int x = 0, w = m.cols * cn; x < w; )
+        for (int c = 0; c < cn; ++c, ++x)
+            s[c] += ptr[x];
+
+    return s;
+}
+
+enum { OP_SUM = 0, OP_SUM_ABS =  1, OP_SUM_SQR = 2 };
+
+static bool ocl_sum( InputArray _src, Scalar & res, int sum_op )
+{
+    CV_Assert(sum_op == OP_SUM || sum_op == OP_SUM_ABS || sum_op == OP_SUM_SQR);
+
+    int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if ( (!doubleSupport && depth == CV_64F) || cn > 4 || cn == 3 )
+        return false;
+
+    int dbsize = ocl::Device::getDefault().maxComputeUnits();
+    size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
+
+    int ddepth = std::max(CV_32S, depth), dtype = CV_MAKE_TYPE(ddepth, cn);
+    UMat src = _src.getUMat(), db(1, dbsize, dtype);
+
+    int wgs2_aligned = 1;
+    while (wgs2_aligned < (int)wgs)
+        wgs2_aligned <<= 1;
+    wgs2_aligned >>= 1;
+
+    static const char * const opMap[3] = { "OP_SUM", "OP_SUM_ABS", "OP_SUM_SQR" };
+    char cvt[40];
+    ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
+                  format("-D srcT=%s -D dstT=%s -D convertToDT=%s -D %s -D WGS=%d -D WGS2_ALIGNED=%d%s",
+                         ocl::typeToStr(type), ocl::typeToStr(dtype), ocl::convertTypeStr(depth, ddepth, cn, cvt),
+                         opMap[sum_op], (int)wgs, wgs2_aligned,
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(),
+           dbsize, ocl::KernelArg::PtrWriteOnly(db));
+
+    size_t globalsize = dbsize * wgs;
+    if (k.run(1, &globalsize, &wgs, true))
+    {
+        typedef Scalar (*part_sum)(Mat m);
+        part_sum funcs[3] = { ocl_part_sum<int>, ocl_part_sum<float>, ocl_part_sum<double> },
+                func = funcs[ddepth - CV_32S];
+        res = func(db.getMat(ACCESS_READ));
+        return true;
+    }
+    return false;
+}
+
 }
 
 cv::Scalar cv::sum( InputArray _src )
 {
+    Scalar _res;
+    if (ocl::useOpenCL() && _src.isUMat() && ocl_sum(_src, _res, OP_SUM))
+        return _res;
+
     Mat src = _src.getMat();
     int k, cn = src.channels(), depth = src.depth();
 
@@ -562,8 +626,9 @@ static bool ocl_countNonZero( InputArray _src, int & res )
         wgs2_aligned <<= 1;
     wgs2_aligned >>= 1;
 
-    ocl::Kernel k("count_non_zero", ocl::core::count_non_zero_oclsrc,
-                  format("-D srcT=%s -D WGS=%d -D WGS2_ALIGNED=%d%s", ocl::typeToStr(src.type()), (int)wgs,
+    ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
+                  format("-D srcT=%s -D OP_COUNT_NON_ZERO -D WGS=%d -D WGS2_ALIGNED=%d%s",
+                         ocl::typeToStr(src.type()), (int)wgs,
                          wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
     k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(),
            dbsize, ocl::KernelArg::PtrWriteOnly(db));

From 88bceb5060f0e71fc5c44a85884db2ead14c6603 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sun, 1 Dec 2013 14:45:48 +0400
Subject: [PATCH 11/13] added cv::meanStdDev to T-API

---
 modules/core/src/stat.cpp | 53 ++++++++++++++++++++++++++++++++++++---
 1 file changed, 49 insertions(+), 4 deletions(-)

diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp
index f8ff0c29d..b3ae4d77e 100644
--- a/modules/core/src/stat.cpp
+++ b/modules/core/src/stat.cpp
@@ -464,16 +464,16 @@ template <typename T> Scalar ocl_part_sum(Mat m)
     return s;
 }
 
-enum { OP_SUM = 0, OP_SUM_ABS =  1, OP_SUM_SQR = 2 };
+enum { OCL_OP_SUM = 0, OCL_OP_SUM_ABS =  1, OCL_OP_SUM_SQR = 2 };
 
 static bool ocl_sum( InputArray _src, Scalar & res, int sum_op )
 {
-    CV_Assert(sum_op == OP_SUM || sum_op == OP_SUM_ABS || sum_op == OP_SUM_SQR);
+    CV_Assert(sum_op == OCL_OP_SUM || sum_op == OCL_OP_SUM_ABS || sum_op == OCL_OP_SUM_SQR);
 
     int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
     bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
 
-    if ( (!doubleSupport && depth == CV_64F) || cn > 4 || cn == 3 )
+    if ( (!doubleSupport && depth == CV_64F) || cn > 4 || cn == 3 || _src.dims() > 2 )
         return false;
 
     int dbsize = ocl::Device::getDefault().maxComputeUnits();
@@ -514,7 +514,7 @@ static bool ocl_sum( InputArray _src, Scalar & res, int sum_op )
 cv::Scalar cv::sum( InputArray _src )
 {
     Scalar _res;
-    if (ocl::useOpenCL() && _src.isUMat() && ocl_sum(_src, _res, OP_SUM))
+    if (ocl::useOpenCL() && _src.isUMat() && ocl_sum(_src, _res, OCL_OP_SUM))
         return _res;
 
     Mat src = _src.getMat();
@@ -798,9 +798,54 @@ cv::Scalar cv::mean( InputArray _src, InputArray _mask )
     return s*(nz0 ? 1./nz0 : 0);
 }
 
+namespace cv {
+
+static bool ocl_meanStdDev( InputArray _src, OutputArray _mean, OutputArray _sdv )
+{
+    Scalar mean, stddev;
+    if (!ocl_sum(_src, mean, OCL_OP_SUM))
+        return false;
+    if (!ocl_sum(_src, stddev, OCL_OP_SUM_SQR))
+        return false;
+
+    double total = 1.0 / _src.total();
+    int k, j, cn = _src.channels();
+    for (int i = 0; i < cn; ++i)
+    {
+        mean[i] *= total;
+        stddev[i] = std::sqrt(std::max(stddev[i] * total - mean[i] * mean[i] , 0.));
+    }
+
+    for( j = 0; j < 2; j++ )
+    {
+        const double * const sptr = j == 0 ? &mean[0] : &stddev[0];
+        _OutputArray _dst = j == 0 ? _mean : _sdv;
+        if( !_dst.needed() )
+            continue;
+
+        if( !_dst.fixedSize() )
+            _dst.create(cn, 1, CV_64F, -1, true);
+        Mat dst = _dst.getMat();
+        int dcn = (int)dst.total();
+        CV_Assert( dst.type() == CV_64F && dst.isContinuous() &&
+                   (dst.cols == 1 || dst.rows == 1) && dcn >= cn );
+        double* dptr = dst.ptr<double>();
+        for( k = 0; k < cn; k++ )
+            dptr[k] = sptr[k];
+        for( ; k < dcn; k++ )
+            dptr[k] = 0;
+    }
+
+    return true;
+}
+
+}
 
 void cv::meanStdDev( InputArray _src, OutputArray _mean, OutputArray _sdv, InputArray _mask )
 {
+    if (ocl::useOpenCL() && _src.isUMat() && _mask.empty() && ocl_meanStdDev(_src, _mean, _sdv))
+        return;
+
     Mat src = _src.getMat(), mask = _mask.getMat();
     CV_Assert( mask.empty() || mask.type() == CV_8U );
 

From 49e729f89db1dabccc6d8128e7e866db86e44a4b Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Sun, 1 Dec 2013 15:37:56 +0400
Subject: [PATCH 12/13] ability to use cv::PSNR with transparent cv::norm

---
 modules/core/src/arithm.cpp           | 7 +++++--
 modules/core/src/mathfuncs.cpp        | 4 ++--
 modules/core/src/ocl.cpp              | 2 +-
 modules/core/src/stat.cpp             | 7 +++----
 modules/core/test/ocl/test_arithm.cpp | 8 ++++----
 5 files changed, 15 insertions(+), 13 deletions(-)

diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp
index c1e27174b..72c27c539 100644
--- a/modules/core/src/arithm.cpp
+++ b/modules/core/src/arithm.cpp
@@ -929,7 +929,10 @@ static bool ocl_binary_op(InputArray _src1, InputArray _src2, OutputArray _dst,
     int srcdepth = CV_MAT_DEPTH(srctype);
     int cn = CV_MAT_CN(srctype);
 
-    if( oclop < 0 || ((haveMask || haveScalar) && (cn > 4 || cn == 3)) )
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if( oclop < 0 || ((haveMask || haveScalar) && (cn > 4 || cn == 3)) ||
+            (!doubleSupport && srcdepth == CV_64F))
         return false;
 
     char opts[1024];
@@ -2626,7 +2629,7 @@ void cv::compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op)
     CV_Assert( op == CMP_LT || op == CMP_LE || op == CMP_EQ ||
                op == CMP_NE || op == CMP_GE || op == CMP_GT );
 
-    if (ocl::useOpenCL() && _dst.isUMat() &&
+    if (ocl::useOpenCL() && _src1.dims() <= 2 && _src2.dims() <= 2 && _dst.isUMat() &&
             ocl_compare(_src1, _src2, _dst, op))
         return;
 
diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp
index d75fad122..aa94e03d0 100644
--- a/modules/core/src/mathfuncs.cpp
+++ b/modules/core/src/mathfuncs.cpp
@@ -687,7 +687,7 @@ static bool ocl_polarToCart( InputArray _mag, InputArray _angle,
     int type = _angle.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
     bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
 
-    if ( _mag.empty() || (!doubleSupport && depth == CV_64F) )
+    if ( !doubleSupport && depth == CV_64F )
         return false;
 
     UMat mag = _mag.getUMat(), angle = _angle.getUMat();
@@ -717,7 +717,7 @@ void polarToCart( InputArray src1, InputArray src2,
     int type = src2.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
     CV_Assert((depth == CV_32F || depth == CV_64F) && (src1.empty() || src1.type() == type));
 
-    if (ocl::useOpenCL() && dst1.isUMat() && dst2.isUMat() &&
+    if (ocl::useOpenCL() && !src1.empty() && src2.dims() <= 2 && dst1.isUMat() && dst2.isUMat() &&
             ocl_polarToCart(src1, src2, dst1, dst2, angleInDegrees))
         return;
 
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index 64460efb0..f733dd11f 100644
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -3145,7 +3145,7 @@ const char* memopTypeToStr(int t)
         "ushort", "ushort2", "ushort3", "ushort4",
         "int", "int2", "int3", "int4",
         "int", "int2", "int3", "int4",
-        "long", "long2", "long3", "long4",
+        "int2", "int4", "?", "int8",
         "?", "?", "?", "?"
     };
     int cn = CV_MAT_CN(t);
diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp
index b3ae4d77e..3feb2db6e 100644
--- a/modules/core/src/stat.cpp
+++ b/modules/core/src/stat.cpp
@@ -635,7 +635,7 @@ static bool ocl_countNonZero( InputArray _src, int & res )
 
     size_t globalsize = dbsize * wgs;
     if (k.run(1, &globalsize, &wgs, true))
-        return res = cv::sum(db.getMat(ACCESS_READ))[0], true;
+        return res = saturate_cast<int>(cv::sum(db.getMat(ACCESS_READ))[0]), true;
     return false;
 }
 
@@ -2752,9 +2752,8 @@ void cv::findNonZero( InputArray _src, OutputArray _idx )
 
 double cv::PSNR(InputArray _src1, InputArray _src2)
 {
-    Mat src1 = _src1.getMat(), src2 = _src2.getMat();
-    CV_Assert( src1.depth() == CV_8U );
-    double diff = std::sqrt(norm(src1, src2, NORM_L2SQR)/(src1.total()*src1.channels()));
+    CV_Assert( _src1.depth() == CV_8U );
+    double diff = std::sqrt(norm(_src1, _src2, NORM_L2SQR)/(_src1.total()*_src1.channels()));
     return 20*log10(255./(diff+DBL_EPSILON));
 }
 
diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp
index d1f2c0170..10cec7bc0 100644
--- a/modules/core/test/ocl/test_arithm.cpp
+++ b/modules/core/test/ocl/test_arithm.cpp
@@ -977,10 +977,10 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, Absdiff, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHA
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, CartToPolar, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, PolarToCart, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Transpose, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_and, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_not, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_xor, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
-OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_or, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_and, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_not, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_xor, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_or, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Pow, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Compare, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));

From 49474903cb6eebf37a68f7f6f43e3c67db33bafd Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Tue, 3 Dec 2013 00:41:07 +0400
Subject: [PATCH 13/13] refactoring

---
 modules/core/src/arithm.cpp       | 18 ++++++++-------
 modules/core/src/mathfuncs.cpp    | 38 ++++++++++++++++++-------------
 modules/core/src/matrix.cpp       |  6 +++--
 modules/core/src/opencl/reduce.cl | 14 ++++--------
 modules/core/src/stat.cpp         | 14 ++++++++----
 5 files changed, 51 insertions(+), 39 deletions(-)

diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp
index 72c27c539..bcd11d256 100644
--- a/modules/core/src/arithm.cpp
+++ b/modules/core/src/arithm.cpp
@@ -2598,14 +2598,6 @@ static bool ocl_compare(InputArray _src1, InputArray _src2, OutputArray _dst, in
     if (!doubleSupport && (depth == CV_64F || _src2.depth() == CV_64F))
         return false;
 
-    CV_Assert(type == type2);
-    UMat src1 = _src1.getUMat(), src2 = _src2.getUMat();
-    Size size = src1.size();
-    CV_Assert(size == src2.size());
-
-    _dst.create(size, CV_8UC(cn));
-    UMat dst = _dst.getUMat();
-
     const char * const operationMap[] = { "==", ">", ">=", "<", "<=", "!=" };
     ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
                   format("-D BINARY_OP -D srcT1=%s -D workT=srcT1"
@@ -2613,6 +2605,16 @@ static bool ocl_compare(InputArray _src1, InputArray _src2, OutputArray _dst, in
                          ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
                          operationMap[op],
                          doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
+    CV_Assert(type == type2);
+    UMat src1 = _src1.getUMat(), src2 = _src2.getUMat();
+    Size size = src1.size();
+    CV_Assert(size == src2.size());
+
+    _dst.create(size, CV_8UC(cn));
+    UMat dst = _dst.getUMat();
 
     k.args(ocl::KernelArg::ReadOnlyNoSize(src1),
            ocl::KernelArg::ReadOnlyNoSize(src2),
diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp
index aa94e03d0..79959435d 100644
--- a/modules/core/src/mathfuncs.cpp
+++ b/modules/core/src/mathfuncs.cpp
@@ -508,6 +508,14 @@ static bool ocl_cartToPolar( InputArray _src1, InputArray _src2,
          (depth == CV_64F && !doubleSupport) )
         return false;
 
+    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
+                  format("-D BINARY_OP -D dstT=%s -D OP_CTP_%s%s",
+                         ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
+                         angleInDegrees ? "AD" : "AR",
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
     UMat src1 = _src1.getUMat(), src2 = _src2.getUMat();
     Size size = src1.size();
     CV_Assert( size == src2.size() );
@@ -516,12 +524,6 @@ static bool ocl_cartToPolar( InputArray _src1, InputArray _src2,
     _dst2.create(size, type);
     UMat dst1 = _dst1.getUMat(), dst2 = _dst2.getUMat();
 
-    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
-                  format("-D BINARY_OP -D dstT=%s -D OP_CTP_%s%s",
-                         ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
-                         angleInDegrees ? "AD" : "AR",
-                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
-
     k.args(ocl::KernelArg::ReadOnlyNoSize(src1),
            ocl::KernelArg::ReadOnlyNoSize(src2),
            ocl::KernelArg::WriteOnly(dst1, cn),
@@ -690,6 +692,14 @@ static bool ocl_polarToCart( InputArray _mag, InputArray _angle,
     if ( !doubleSupport && depth == CV_64F )
         return false;
 
+    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
+                  format("-D dstT=%s -D BINARY_OP -D OP_PTC_%s%s",
+                         ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
+                         angleInDegrees ? "AD" : "AR",
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
     UMat mag = _mag.getUMat(), angle = _angle.getUMat();
     Size size = angle.size();
     CV_Assert(mag.size() == size);
@@ -698,12 +708,6 @@ static bool ocl_polarToCart( InputArray _mag, InputArray _angle,
     _dst2.create(size, type);
     UMat dst1 = _dst1.getUMat(), dst2 = _dst2.getUMat();
 
-    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
-                  format("-D dstT=%s -D BINARY_OP -D OP_PTC_%s%s",
-                         ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
-                         angleInDegrees ? "AD" : "AR",
-                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
-
     k.args(ocl::KernelArg::ReadOnlyNoSize(mag), ocl::KernelArg::ReadOnlyNoSize(angle),
            ocl::KernelArg::WriteOnly(dst1, cn), ocl::KernelArg::WriteOnlyNoSize(dst2));
 
@@ -2037,13 +2041,15 @@ static bool ocl_pow(InputArray _src, double power, OutputArray _dst)
          (depth == CV_64F && !doubleSupport) )
         return false;
 
-    UMat src = _src.getUMat();
-    _dst.create(src.size(), type);
-    UMat dst = _dst.getUMat();
-
     ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
                   format("-D dstT=%s -D OP_POW -D UNARY_OP%s", ocl::typeToStr(CV_MAKE_TYPE(depth, 1)),
                          doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
+    UMat src = _src.getUMat();
+    _dst.create(src.size(), type);
+    UMat dst = _dst.getUMat();
 
     ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
             dstarg = ocl::KernelArg::WriteOnly(dst, cn);
diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp
index b2b164e45..4e9be9807 100644
--- a/modules/core/src/matrix.cpp
+++ b/modules/core/src/matrix.cpp
@@ -2378,10 +2378,12 @@ static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s )
     if (cn == 3)
         return false;
 
-    UMat m = _m.getUMat();
-
     ocl::Kernel k("setIdentity", ocl::core::set_identity_oclsrc,
                   format("-D T=%s", ocl::memopTypeToStr(type)));
+    if (k.empty())
+        return false;
+
+    UMat m = _m.getUMat();
     k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, type, s)));
 
     size_t globalsize[2] = { m.cols, m.rows };
diff --git a/modules/core/src/opencl/reduce.cl b/modules/core/src/opencl/reduce.cl
index 2ba36e859..526cc5101 100644
--- a/modules/core/src/opencl/reduce.cl
+++ b/modules/core/src/opencl/reduce.cl
@@ -66,7 +66,7 @@
     dstT temp = convertToDT(src[0]); \
     FUNC(accumulator, temp)
 #define REDUCE_LOCAL_1 \
-    localmem[lid] += accumulator
+    localmem[lid - WGS2_ALIGNED] += accumulator
 #define REDUCE_LOCAL_2 \
     localmem[lid] += localmem[lid2]
 
@@ -78,7 +78,7 @@
 #define REDUCE_GLOBAL \
     accumulator += src[0] == zero ? zero : one
 #define REDUCE_LOCAL_1 \
-    localmem[lid] += accumulator
+    localmem[lid - WGS2_ALIGNED] += accumulator
 #define REDUCE_LOCAL_2 \
     localmem[lid] += localmem[lid2]
 
@@ -95,10 +95,6 @@ __kernel void reduce(__global const uchar * srcptr, int step, int offset, int co
     int  id = get_global_id(0);
 
     __local dstT localmem[WGS2_ALIGNED];
-    if (lid < WGS2_ALIGNED)
-        localmem[lid] = (dstT)(0);
-    barrier(CLK_LOCAL_MEM_FENCE);
-
     DEFINE_ACCUMULATOR;
 
     for (int grain = groupnum * WGS; id < total; id += grain)
@@ -108,11 +104,11 @@ __kernel void reduce(__global const uchar * srcptr, int step, int offset, int co
         REDUCE_GLOBAL;
     }
 
-    if (lid >= WGS2_ALIGNED)
-        localmem[lid - WGS2_ALIGNED] = accumulator;
+    if (lid < WGS2_ALIGNED)
+        localmem[lid] = accumulator;
     barrier(CLK_LOCAL_MEM_FENCE);
 
-    if (lid < WGS2_ALIGNED)
+    if (lid >= WGS2_ALIGNED)
         REDUCE_LOCAL_1;
     barrier(CLK_LOCAL_MEM_FENCE);
 
diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp
index 3feb2db6e..b19be3b47 100644
--- a/modules/core/src/stat.cpp
+++ b/modules/core/src/stat.cpp
@@ -480,7 +480,6 @@ static bool ocl_sum( InputArray _src, Scalar & res, int sum_op )
     size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
 
     int ddepth = std::max(CV_32S, depth), dtype = CV_MAKE_TYPE(ddepth, cn);
-    UMat src = _src.getUMat(), db(1, dbsize, dtype);
 
     int wgs2_aligned = 1;
     while (wgs2_aligned < (int)wgs)
@@ -494,6 +493,10 @@ static bool ocl_sum( InputArray _src, Scalar & res, int sum_op )
                          ocl::typeToStr(type), ocl::typeToStr(dtype), ocl::convertTypeStr(depth, ddepth, cn, cvt),
                          opMap[sum_op], (int)wgs, wgs2_aligned,
                          doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
+    UMat src = _src.getUMat(), db(1, dbsize, dtype);
     k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(),
            dbsize, ocl::KernelArg::PtrWriteOnly(db));
 
@@ -611,7 +614,7 @@ namespace cv {
 
 static bool ocl_countNonZero( InputArray _src, int & res )
 {
-    int depth = _src.depth();
+    int type = _src.type(), depth = CV_MAT_DEPTH(type);
     bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
 
     if (depth == CV_64F && !doubleSupport)
@@ -619,7 +622,6 @@ static bool ocl_countNonZero( InputArray _src, int & res )
 
     int dbsize = ocl::Device::getDefault().maxComputeUnits();
     size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
-    UMat src = _src.getUMat(), db(1, dbsize, CV_32SC1);
 
     int wgs2_aligned = 1;
     while (wgs2_aligned < (int)wgs)
@@ -628,8 +630,12 @@ static bool ocl_countNonZero( InputArray _src, int & res )
 
     ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
                   format("-D srcT=%s -D OP_COUNT_NON_ZERO -D WGS=%d -D WGS2_ALIGNED=%d%s",
-                         ocl::typeToStr(src.type()), (int)wgs,
+                         ocl::typeToStr(type), (int)wgs,
                          wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
+    UMat src = _src.getUMat(), db(1, dbsize, CV_32SC1);
     k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(),
            dbsize, ocl::KernelArg::PtrWriteOnly(db));