Merge pull request #1969 from KonstantinMatskevich:ocl_tapi_minmaxloc

modules/core/src/opencl/reduce.cl

@@ -51,6 +51,7 @@
 #endif
 
 #define noconvert
+#define EXTRA_PARAMS
 
 #if defined OP_SUM || defined OP_SUM_ABS || defined OP_SUM_SQR
 #if OP_SUM
@@ -60,41 +61,183 @@
 #elif OP_SUM_SQR
 #define FUNC(a, b) a += b * b
 #endif
+#define DECLARE_LOCAL_MEM \
+    __local dstT localmem[WGS2_ALIGNED]
 #define DEFINE_ACCUMULATOR \
     dstT accumulator = (dstT)(0)
 #define REDUCE_GLOBAL \
     dstT temp = convertToDT(src[0]); \
     FUNC(accumulator, temp)
+#define SET_LOCAL_1 \
+        localmem[lid] = accumulator
 #define REDUCE_LOCAL_1 \
     localmem[lid - WGS2_ALIGNED] += accumulator
 #define REDUCE_LOCAL_2 \
     localmem[lid] += localmem[lid2]
+#define CALC_RESULT \
+    __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid); \
+    dst[0] = localmem[0]
 
 #elif defined OP_COUNT_NON_ZERO
 #define dstT int
+#define DECLARE_LOCAL_MEM \
+    __local dstT localmem[WGS2_ALIGNED]
 #define DEFINE_ACCUMULATOR \
     dstT accumulator = (dstT)(0); \
     srcT zero = (srcT)(0), one = (srcT)(1)
 #define REDUCE_GLOBAL \
     accumulator += src[0] == zero ? zero : one
+#define SET_LOCAL_1 \
+        localmem[lid] = accumulator
 #define REDUCE_LOCAL_1 \
     localmem[lid - WGS2_ALIGNED] += accumulator
 #define REDUCE_LOCAL_2 \
     localmem[lid] += localmem[lid2]
+#define CALC_RESULT \
+    __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid); \
+    dst[0] = localmem[0]
+
+#elif defined OP_MIN_MAX_LOC || defined OP_MIN_MAX_LOC_MASK
+
+#if defined (DEPTH_0)
+#define srcT uchar
+#define MIN_VAL 0
+#define MAX_VAL 255
+#endif
+#if defined (DEPTH_1)
+#define srcT char
+#define MIN_VAL -128
+#define MAX_VAL 127
+#endif
+#if defined (DEPTH_2)
+#define srcT ushort
+#define MIN_VAL 0
+#define MAX_VAL 65535
+#endif
+#if defined (DEPTH_3)
+#define srcT short
+#define MIN_VAL -32768
+#define MAX_VAL 32767
+#endif
+#if defined (DEPTH_4)
+#define srcT int
+#define MIN_VAL INT_MIN
+#define MAX_VAL INT_MAX
+#endif
+#if defined (DEPTH_5)
+#define srcT float
+#define MIN_VAL (-FLT_MAX)
+#define MAX_VAL FLT_MAX
+#endif
+#if defined (DEPTH_6)
+#define srcT double
+#define MIN_VAL (-DBL_MAX)
+#define MAX_VAL DBL_MAX
+#endif
+
+#define DECLARE_LOCAL_MEM \
+    __local srcT localmem_min[WGS2_ALIGNED]; \
+    __local srcT localmem_max[WGS2_ALIGNED]; \
+    __local int localmem_minloc[WGS2_ALIGNED]; \
+    __local int localmem_maxloc[WGS2_ALIGNED]
+#define DEFINE_ACCUMULATOR \
+    srcT minval = MAX_VAL; \
+    srcT maxval = MIN_VAL; \
+    int negative = -1; \
+    int minloc = negative; \
+    int maxloc = negative; \
+    srcT temp; \
+    int temploc
+#define REDUCE_GLOBAL \
+    temp = src[0]; \
+    temploc = id; \
+    srcT temp_minval = minval, temp_maxval = maxval; \
+    minval = min(minval, temp); \
+    maxval = max(maxval, temp); \
+    minloc = (minval == temp_minval) ? (temp_minval == MAX_VAL) ? temploc : minloc : temploc; \
+    maxloc = (maxval == temp_maxval) ? (temp_maxval == MIN_VAL) ? temploc : maxloc : temploc
+#define SET_LOCAL_1 \
+    localmem_min[lid] = minval; \
+    localmem_max[lid] = maxval; \
+    localmem_minloc[lid] = minloc; \
+    localmem_maxloc[lid] = maxloc
+#define REDUCE_LOCAL_1 \
+    srcT oldmin = localmem_min[lid-WGS2_ALIGNED]; \
+    srcT oldmax = localmem_max[lid-WGS2_ALIGNED]; \
+    localmem_min[lid - WGS2_ALIGNED] = min(minval,localmem_min[lid-WGS2_ALIGNED]); \
+    localmem_max[lid - WGS2_ALIGNED] = max(maxval,localmem_max[lid-WGS2_ALIGNED]); \
+    srcT minv = localmem_min[lid - WGS2_ALIGNED], maxv = localmem_max[lid - WGS2_ALIGNED]; \
+    localmem_minloc[lid - WGS2_ALIGNED] = (minv == minval) ? (minv == oldmin) ? \
+        min(minloc, localmem_minloc[lid-WGS2_ALIGNED]) : minloc : localmem_minloc[lid-WGS2_ALIGNED]; \
+    localmem_maxloc[lid - WGS2_ALIGNED] = (maxv == maxval) ? (maxv == oldmax) ? \
+        min(maxloc, localmem_maxloc[lid-WGS2_ALIGNED]) : maxloc : localmem_maxloc[lid-WGS2_ALIGNED]
+#define REDUCE_LOCAL_2 \
+    srcT oldmin = localmem_min[lid]; \
+    srcT oldmax = localmem_max[lid]; \
+    localmem_min[lid] = min(localmem_min[lid], localmem_min[lid2]); \
+    localmem_max[lid] = max(localmem_max[lid], localmem_max[lid2]); \
+    srcT min1 = localmem_min[lid], min2 = localmem_min[lid2]; \
+    localmem_minloc[lid] = (localmem_minloc[lid] == negative) ? localmem_minloc[lid2] : (localmem_minloc[lid2] == negative) ? \
+        localmem_minloc[lid] : (min1 == min2) ? (min1 == oldmin) ? min(localmem_minloc[lid2],localmem_minloc[lid]) : \
+        localmem_minloc[lid2] : localmem_minloc[lid]; \
+    srcT max1 = localmem_max[lid], max2 = localmem_max[lid2]; \
+    localmem_maxloc[lid] = (localmem_maxloc[lid] == negative) ? localmem_maxloc[lid2] : (localmem_maxloc[lid2] == negative) ? \
+        localmem_maxloc[lid] : (max1 == max2) ? (max1 == oldmax) ? min(localmem_maxloc[lid2],localmem_maxloc[lid]) : \
+        localmem_maxloc[lid2] : localmem_maxloc[lid]
+#define CALC_RESULT \
+    __global srcT * dstminval = (__global srcT *)(dstptr + (int)sizeof(srcT) * gid); \
+    __global srcT * dstmaxval = (__global srcT *)(dstptr2 + (int)sizeof(srcT) * gid); \
+    dstminval[0] = localmem_min[0]; \
+    dstmaxval[0] = localmem_max[0]; \
+    dstlocptr[gid] = localmem_minloc[0]; \
+    dstlocptr2[gid] = localmem_maxloc[0]
+
+#if defined OP_MIN_MAX_LOC_MASK
+#undef DEFINE_ACCUMULATOR
+#define DEFINE_ACCUMULATOR \
+    srcT minval = MAX_VAL; \
+    srcT maxval = MIN_VAL; \
+    int negative = -1; \
+    int minloc = negative; \
+    int maxloc = negative; \
+    srcT temp, temp_mask, zeroVal = (srcT)(0); \
+    int temploc
+#undef REDUCE_GLOBAL
+#define REDUCE_GLOBAL \
+    temp = src[0]; \
+    temploc = id; \
+    int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols) * (int)sizeof(uchar)); \
+    __global const uchar * mask = (__global const uchar *)(maskptr + mask_index); \
+    temp_mask = mask[0]; \
+    srcT temp_minval = minval, temp_maxval = maxval; \
+    minval = (temp_mask == zeroVal) ? minval : min(minval, temp); \
+    maxval = (temp_mask == zeroVal) ? maxval : max(maxval, temp); \
+    minloc = (temp_mask == zeroVal) ? minloc : (minval == temp_minval) ? (temp_minval == MAX_VAL) ? temploc : minloc : temploc; \
+    maxloc = (temp_mask == zeroVal) ? maxloc : (maxval == temp_maxval) ? (temp_maxval == MIN_VAL) ? temploc : maxloc : temploc
+#endif
 
 #else
 #error "No operation"
+#endif
 
+#if defined OP_MIN_MAX_LOC
+#undef EXTRA_PARAMS
+#define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2
+#endif
+#if defined OP_MIN_MAX_LOC_MASK
+#undef EXTRA_PARAMS
+#define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2, \
+    __global const uchar * maskptr, int mask_step, int mask_offset, __global int * test
 #endif
 
 __kernel void reduce(__global const uchar * srcptr, int step, int offset, int cols,
-                     int total, int groupnum, __global uchar * dstptr)
+                     int total, int groupnum, __global uchar * dstptr EXTRA_PARAMS)
 {
     int lid = get_local_id(0);
     int gid = get_group_id(0);
     int  id = get_global_id(0);
 
-    __local dstT localmem[WGS2_ALIGNED];
+    DECLARE_LOCAL_MEM;
     DEFINE_ACCUMULATOR;
 
     for (int grain = groupnum * WGS; id < total; id += grain)
@@ -105,11 +248,15 @@ __kernel void reduce(__global const uchar * srcptr, int step, int offset, int co
     }
 
     if (lid < WGS2_ALIGNED)
-        localmem[lid] = accumulator;
+    {
+        SET_LOCAL_1;
+    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
-    if (lid >= WGS2_ALIGNED)
+    if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
+    {
         REDUCE_LOCAL_1;
+    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
@@ -124,7 +271,6 @@ __kernel void reduce(__global const uchar * srcptr, int step, int offset, int co
 
     if (lid == 0)
     {
-        __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid);
+        CALC_RESULT;
-        dst[0] = localmem[0];
     }
 }

modules/core/src/stat.cpp

@@ -1159,16 +1159,137 @@ static void ofs2idx(const Mat& a, size_t ofs, int* idx)
 
 }
 
+namespace cv
+{
+
+template <typename T>
+void getMinMaxRes(const Mat &minv, const Mat &maxv, const Mat &minl, const Mat &maxl, double* minVal,
+                  double* maxVal, int* minLoc, int* maxLoc, const int groupnum, const int cn, const int cols)
+{
+    T min = std::numeric_limits<T>::max();
+    T max = std::numeric_limits<T>::min() > 0 ? -std::numeric_limits<T>::max() : std::numeric_limits<T>::min();
+    int minloc = INT_MAX, maxloc = INT_MAX;
+    for( int i = 0; i < groupnum; i++)
+    {
+        T current_min = minv.at<T>(0,i);
+        T current_max = maxv.at<T>(0,i);
+        T oldmin = min, oldmax = max;
+        min = std::min(min, current_min);
+        max = std::max(max, current_max);
+        if (cn == 1)
+        {
+            int current_minloc = minl.at<int>(0,i);
+            int current_maxloc = maxl.at<int>(0,i);
+            if(current_minloc < 0 || current_maxloc < 0) continue;
+            minloc = (oldmin == current_min) ? std::min(minloc, current_minloc) : (oldmin < current_min) ? minloc : current_minloc;
+            maxloc = (oldmax == current_max) ? std::min(maxloc, current_maxloc) : (oldmax > current_max) ? maxloc : current_maxloc;
+        }
+    }
+    bool zero_mask = (maxloc == INT_MAX) || (minloc == INT_MAX);
+    if(minVal)
+        *minVal = zero_mask ? 0 : (double)min;
+    if(maxVal)
+        *maxVal = zero_mask ? 0 : (double)max;
+    if(minLoc)
+    {
+        minLoc[0] = zero_mask ? -1 : minloc/cols;
+        minLoc[1] = zero_mask ? -1 : minloc%cols;
+    }
+    if(maxLoc)
+    {
+        maxLoc[0] = zero_mask ? -1 : maxloc/cols;
+        maxLoc[1] = zero_mask ? -1 : maxloc%cols;
+    }
+}
+
+typedef void (*getMinMaxResFunc)(const Mat &minv, const Mat &maxv, const Mat &minl, const Mat &maxl, double *minVal,
+                                 double *maxVal, int *minLoc, int *maxLoc, const int gropunum, const int cn, const int cols);
+
+static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int* minLoc, int* maxLoc, InputArray _mask)
+{
+    CV_Assert( (_src.channels() == 1 && (_mask.empty() || _mask.type() == CV_8U)) ||
+        (_src.channels() >= 1 && _mask.empty() && !minLoc && !maxLoc) );
+
+    int type = _src.type(), depth = CV_MAT_DEPTH(type);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if (depth == CV_64F && !doubleSupport)
+        return false;
+
+    int groupnum = ocl::Device::getDefault().maxComputeUnits();
+    size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
+
+    int wgs2_aligned = 1;
+    while (wgs2_aligned < (int)wgs)
+        wgs2_aligned <<= 1;
+    wgs2_aligned >>= 1;
+
+    String opts = format("-D DEPTH_%d -D OP_MIN_MAX_LOC%s -D WGS=%d -D WGS2_ALIGNED=%d %s",
+        depth, _mask.empty() ? "" : "_MASK", (int)wgs, wgs2_aligned, doubleSupport ? "-D DOUBLE_SUPPORT" : "");
+
+    ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, opts);
+    if (k.empty())
+        return false;
+
+    UMat src = _src.getUMat(), minval(1, groupnum, src.type()),
+        maxval(1, groupnum, src.type()), minloc( 1, groupnum, CV_32SC1),
+        maxloc( 1, groupnum, CV_32SC1), mask;
+    if(!_mask.empty())
+        mask = _mask.getUMat();
+
+    if(src.channels()>1)
+        src = src.reshape(1);
+
+    if(mask.empty())
+        k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(),
+            groupnum, ocl::KernelArg::PtrWriteOnly(minval), ocl::KernelArg::PtrWriteOnly(maxval),
+            ocl::KernelArg::PtrWriteOnly(minloc), ocl::KernelArg::PtrWriteOnly(maxloc));
+    else
+        k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(), groupnum,
+            ocl::KernelArg::PtrWriteOnly(minval), ocl::KernelArg::PtrWriteOnly(maxval),
+            ocl::KernelArg::PtrWriteOnly(minloc), ocl::KernelArg::PtrWriteOnly(maxloc), ocl::KernelArg::ReadOnlyNoSize(mask));
+
+    size_t globalsize = groupnum * wgs;
+    if (!k.run(1, &globalsize, &wgs, true))
+        return false;
+
+    Mat minv = minval.getMat(ACCESS_READ), maxv = maxval.getMat(ACCESS_READ),
+        minl = minloc.getMat(ACCESS_READ), maxl = maxloc.getMat(ACCESS_READ);
+
+    static getMinMaxResFunc functab[7] =
+    {
+        getMinMaxRes<uchar>,
+        getMinMaxRes<char>,
+        getMinMaxRes<ushort>,
+        getMinMaxRes<short>,
+        getMinMaxRes<int>,
+        getMinMaxRes<float>,
+        getMinMaxRes<double>
+    };
+
+    getMinMaxResFunc func;
+
+    func = functab[depth];
+    func(minv, maxv, minl, maxl, minVal, maxVal, minLoc, maxLoc, groupnum, src.channels(), src.cols);
+
+    return true;
+}
+}
+
 void cv::minMaxIdx(InputArray _src, double* minVal,
                    double* maxVal, int* minIdx, int* maxIdx,
                    InputArray _mask)
 {
+    CV_Assert( (_src.channels() == 1 && (_mask.empty() || _mask.type() == CV_8U)) ||
+        (_src.channels() >= 1 && _mask.empty() && !minIdx && !maxIdx) );
+
+     if( ocl::useOpenCL() && _src.isUMat() && _src.dims() <= 2  && ( _mask.empty() || _src.size() == _mask.size() )
+         && ocl_minMaxIdx(_src, minVal, maxVal, minIdx, maxIdx, _mask) )
+        return;
+
     Mat src = _src.getMat(), mask = _mask.getMat();
     int depth = src.depth(), cn = src.channels();
 
-    CV_Assert( (cn == 1 && (mask.empty() || mask.type() == CV_8U)) ||
-               (cn >= 1 && mask.empty() && !minIdx && !maxIdx) );
-
 #if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
     size_t total_size = src.total();
     int rows = src.size[0], cols = (int)(total_size/rows);
@@ -1289,8 +1410,7 @@ void cv::minMaxIdx(InputArray _src, double* minVal,
 void cv::minMaxLoc( InputArray _img, double* minVal, double* maxVal,
                     Point* minLoc, Point* maxLoc, InputArray mask )
 {
-    Mat img = _img.getMat();
+    CV_Assert(_img.dims() <= 2);
-    CV_Assert(img.dims <= 2);
 
     minMaxIdx(_img, minVal, maxVal, (int*)minLoc, (int*)maxLoc, mask);
     if( minLoc )

modules/core/test/ocl/test_arithm.cpp

@@ -999,6 +999,67 @@ OCL_TEST_P(Flip, BOTH)
         Near(0);
     }
 }
+//////////////////////////////////////// minMaxIdx /////////////////////////////////////////
+
+typedef ArithmTestBase MinMaxIdx;
+
+OCL_TEST_P(MinMaxIdx, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        int p1[2], p2[2], up1[2], up2[2];
+        double minv, maxv, uminv, umaxv;
+
+        if(src1_roi.channels() > 1)
+        {
+            OCL_OFF(cv::minMaxIdx(src2_roi, &minv, &maxv) );
+            OCL_ON(cv::minMaxIdx(usrc2_roi, &uminv, &umaxv));
+
+            EXPECT_DOUBLE_EQ(minv, uminv);
+            EXPECT_DOUBLE_EQ(maxv, umaxv);
+        }
+        else
+        {
+            OCL_OFF(cv::minMaxIdx(src2_roi, &minv, &maxv, p1, p2, noArray()));
+            OCL_ON(cv::minMaxIdx(usrc2_roi, &uminv, &umaxv, up1, up2, noArray()));
+
+            EXPECT_DOUBLE_EQ(minv, uminv);
+            EXPECT_DOUBLE_EQ(maxv, umaxv);
+            for( int i = 0; i < 2; i++)
+            {
+                EXPECT_EQ(p1[i], up1[i]);
+                EXPECT_EQ(p2[i], up2[i]);
+            }
+        }
+    }
+}
+
+typedef ArithmTestBase MinMaxIdx_Mask;
+
+OCL_TEST_P(MinMaxIdx_Mask, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        int p1[2], p2[2], up1[2], up2[2];
+        double minv, maxv, uminv, umaxv;
+
+        OCL_OFF(cv::minMaxIdx(src2_roi, &minv, &maxv, p1, p2, mask_roi));
+        OCL_ON(cv::minMaxIdx(usrc2_roi, &uminv, &umaxv, up1, up2, umask_roi));
+
+        EXPECT_DOUBLE_EQ(minv, uminv);
+        EXPECT_DOUBLE_EQ(maxv, umaxv);
+        for( int i = 0; i < 2; i++)
+        {
+            EXPECT_EQ(p1[i], up1[i]);
+            EXPECT_EQ(p2[i], up2[i]);
+        }
+
+    }
+}
 
 //////////////////////////////// Norm /////////////////////////////////////////////////
 
@@ -1149,6 +1210,8 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64
 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, Flip, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx_Mask, Combine(OCL_ALL_DEPTHS, ::testing::Values(Channels(1)), Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sqrt, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));