Merge pull request #3088 from vbystricky:ocl_enableNormEtc

modules/core/src/convert.cpp

@@ -1765,13 +1765,12 @@ static bool ocl_convertScaleAbs( InputArray _src, OutputArray _dst, double alpha
         kercn = ocl::predictOptimalVectorWidth(_src, _dst), rowsPerWI = d.isIntel() ? 4 : 1;
     bool doubleSupport = d.doubleFPConfig() > 0;
 
-    if (depth == CV_32F || depth == CV_64F)
+    if (!doubleSupport && depth == CV_64F)
         return false;
 
     char cvt[2][50];
     int wdepth = std::max(depth, CV_32F);
-    ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
+    String build_opt = format("-D OP_CONVERT_SCALE_ABS -D UNARY_OP -D dstT=%s -D srcT1=%s"
-                  format("-D OP_CONVERT_SCALE_ABS -D UNARY_OP -D dstT=%s -D srcT1=%s"
                          " -D workT=%s -D wdepth=%d -D convertToWT1=%s -D convertToDT=%s"
                          " -D workT1=%s -D rowsPerWI=%d%s",
                          ocl::typeToStr(CV_8UC(kercn)),
@@ -1780,7 +1779,8 @@ static bool ocl_convertScaleAbs( InputArray _src, OutputArray _dst, double alpha
                          ocl::convertTypeStr(depth, wdepth, kercn, cvt[0]),
                          ocl::convertTypeStr(wdepth, CV_8U, kercn, cvt[1]),
                          ocl::typeToStr(wdepth), rowsPerWI,
-                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+    ocl::Kernel k("KF", ocl::core::arithm_oclsrc, build_opt);
     if (k.empty())
         return false;
 

modules/core/src/opencl/minmaxloc.cl

@@ -98,7 +98,7 @@
 
 #ifdef OP_CALC2
 #define CALC_MAX2(p) \
-    maxval2 = MAX(maxval2, temp.p);
+    maxval2 = MAX(maxval2, temp2.p);
 #else
 #define CALC_MAX2(p)
 #endif
@@ -196,7 +196,7 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
 
 #ifdef HAVE_SRC2
 #ifdef HAVE_SRC2_CONT
-            src2_index = mul24(id, srcTSIZE);
+            src2_index = id * srcTSIZE; //mul24(id, srcTSIZE);
 #else
             src2_index = mad24(id / cols, src2_step, mul24(id % cols, srcTSIZE));
 #endif

modules/core/src/stat.cpp

@@ -1444,7 +1444,10 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int*
     bool doubleSupport = dev.doubleFPConfig() > 0, haveMask = !_mask.empty(),
         haveSrc2 = _src2.kind() != _InputArray::NONE;
     int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
-            kercn = haveMask ? cn : std::min(4, ocl::predictOptimalVectorWidth(_src));
+            kercn = haveMask ? cn : std::min(4, ocl::predictOptimalVectorWidth(_src, _src2));
+
+    if (haveMask && dev.isAMD())
+        return false;
 
     CV_Assert( (cn == 1 && (!haveMask || _mask.type() == CV_8U)) ||
               (cn >= 1 && !minLoc && !maxLoc) );
@@ -1536,7 +1539,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int*
     }
 
     size_t globalsize = groupnum * wgs;
-    if (!k.run(1, &globalsize, &wgs, false))
+    if (!k.run(1, &globalsize, &wgs, true))
         return false;
 
     static const getMinMaxResFunc functab[7] =
@@ -2190,9 +2193,6 @@ static bool ocl_norm( InputArray _src, int normType, InputArray _mask, double &
          (!doubleSupport && depth == CV_64F))
         return false;
 
-    if( depth == CV_32F && (!_mask.empty() || normType == NORM_INF) )
-        return false;
-
     UMat src = _src.getUMat();
 
     if (normType == NORM_INF)
@@ -2548,9 +2548,6 @@ static bool ocl_norm( InputArray _src1, InputArray _src2, int normType, InputArr
     normType &= ~NORM_RELATIVE;
     bool normsum = normType == NORM_L1 || normType == NORM_L2 || normType == NORM_L2SQR;
 
-    if ( !normsum || !_mask.empty() )
-        return false;
-
     if (normsum)
     {
         if (!ocl_sum(_src1, sc1, normType == NORM_L2 || normType == NORM_L2SQR ?

modules/core/test/ocl/test_arithm.cpp

@@ -341,7 +341,7 @@ OCL_TEST_P(Mul, Mat_Scalar_Scale)
         OCL_OFF(cv::multiply(src1_roi, val, dst1_roi, val[0]));
         OCL_ON(cv::multiply(usrc1_roi, val, udst1_roi, val[0]));
 
-        Near(udst1_roi.depth() >= CV_32F ? 2e-2 : 1);
+        Near(udst1_roi.depth() >= CV_32F ? 1e-2 : 1);
     }
 }
 
@@ -397,7 +397,7 @@ OCL_TEST_P(Div, Mat_Scale)
         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 ? 2e-2 : 1);
+        Near(udst1_roi.depth() >= CV_32F ? 4e-3 : 1);
     }
 }
 
@@ -1303,7 +1303,7 @@ OCL_TEST_P(Norm, NORM_INF_2args)
             OCL_OFF(const double cpuRes = cv::norm(src1_roi, src2_roi, type));
             OCL_ON(const double gpuRes = cv::norm(usrc1_roi, usrc2_roi, type));
 
-            EXPECT_NEAR(cpuRes, gpuRes, 0.2);
+            EXPECT_PRED3(relativeError, cpuRes, gpuRes, 1e-6);
         }
 }
 
@@ -1323,7 +1323,7 @@ OCL_TEST_P(Norm, NORM_INF_2args_mask)
             OCL_OFF(const double cpuRes = cv::norm(src1_roi, src2_roi, type, mask_roi));
             OCL_ON(const double gpuRes = cv::norm(usrc1_roi, usrc2_roi, type, umask_roi));
 
-            EXPECT_NEAR(cpuRes, gpuRes, 0.1);
+            EXPECT_PRED3(relativeError, cpuRes, gpuRes, 1e-6);
         }
 }
 
@@ -1547,7 +1547,49 @@ OCL_TEST_P(InRange, Scalar)
 
 //////////////////////////////// ConvertScaleAbs ////////////////////////////////////////////////
 
-typedef ArithmTestBase ConvertScaleAbs;
+PARAM_TEST_CASE(ConvertScaleAbs, MatDepth, Channels, bool)
+{
+    int depth;
+    int cn;
+    bool use_roi;
+    cv::Scalar val;
+
+    TEST_DECLARE_INPUT_PARAMETER(src);
+    TEST_DECLARE_OUTPUT_PARAMETER(dst);
+
+    virtual void SetUp()
+    {
+        depth = GET_PARAM(0);
+        cn = GET_PARAM(1);
+        use_roi = GET_PARAM(2);
+    }
+
+    virtual void generateTestData()
+    {
+        const int stype = CV_MAKE_TYPE(depth, cn);
+        const int dtype = CV_MAKE_TYPE(CV_8U, cn);
+
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, stype, 2, 11); // FIXIT: Test with minV, maxV
+
+        Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, dstBorder, dtype, 5, 16);
+
+        val = cv::Scalar(rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0),
+                         rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0));
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src);
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst);
+    }
+
+    void Near(double threshold = 0.)
+    {
+        OCL_EXPECT_MATS_NEAR(dst, threshold);
+    }
+
+};
+
 
 OCL_TEST_P(ConvertScaleAbs, Mat)
 {
@@ -1555,10 +1597,10 @@ OCL_TEST_P(ConvertScaleAbs, Mat)
     {
         generateTestData();
 
-        OCL_OFF(cv::convertScaleAbs(src1_roi, dst1_roi, val[0], val[1]));
+        OCL_OFF(cv::convertScaleAbs(src_roi, dst_roi, val[0], val[1]));
-        OCL_ON(cv::convertScaleAbs(usrc1_roi, udst1_roi, val[0], val[1]));
+        OCL_ON(cv::convertScaleAbs(usrc_roi, udst_roi, val[0], val[1]));
 
-        Near(depth <= CV_32S ? 1 : 1e-6);
+        Near(1);
     }
 }