Merge pull request #2156 from ilya-lavrenov:tapi_reduce

modules/core/perf/opencl/perf_arithm.cpp

@@ -912,6 +912,70 @@ OCL_PERF_TEST_P(PSNRFixture, PSNR,
     SANITY_CHECK(psnr, 1e-4, ERROR_RELATIVE);
 }
 
+///////////// Reduce ////////////////////////
+
+CV_ENUM(ReduceMinMaxOp, CV_REDUCE_MIN, CV_REDUCE_MAX)
+
+typedef tuple<Size, std::pair<MatType, MatType>, int, ReduceMinMaxOp> ReduceMinMaxParams;
+typedef TestBaseWithParam<ReduceMinMaxParams> ReduceMinMaxFixture;
+
+OCL_PERF_TEST_P(ReduceMinMaxFixture, Reduce,
+                ::testing::Combine(OCL_TEST_SIZES,
+                                   OCL_PERF_ENUM(std::make_pair<MatType, MatType>(CV_8UC1, CV_8UC1),
+                                                 std::make_pair<MatType, MatType>(CV_32FC4, CV_32FC4)),
+                                   OCL_PERF_ENUM(0, 1),
+                                   ReduceMinMaxOp::all()))
+{
+    const ReduceMinMaxParams params = GetParam();
+    const std::pair<MatType, MatType> types = get<1>(params);
+    const int stype = types.first, dtype = types.second,
+            dim = get<2>(params), op = get<3>(params);
+    const Size srcSize = get<0>(params),
+            dstSize(dim == 0 ? srcSize.width : 1, dim == 0 ? 1 : srcSize.height);
+    const double eps = CV_MAT_DEPTH(dtype) <= CV_32S ? 1 : 1e-5;
+
+    checkDeviceMaxMemoryAllocSize(srcSize, stype);
+    checkDeviceMaxMemoryAllocSize(srcSize, dtype);
+
+    UMat src(srcSize, stype), dst(dstSize, dtype);
+    declare.in(src, WARMUP_RNG).out(dst);
+
+    OCL_TEST_CYCLE() cv::reduce(src, dst, dim, op, dtype);
+
+    SANITY_CHECK(dst, eps);
+}
+
+CV_ENUM(ReduceAccOp, CV_REDUCE_SUM, CV_REDUCE_AVG)
+
+typedef tuple<Size, std::pair<MatType, MatType>, int, ReduceAccOp> ReduceAccParams;
+typedef TestBaseWithParam<ReduceAccParams> ReduceAccFixture;
+
+OCL_PERF_TEST_P(ReduceAccFixture, Reduce,
+                ::testing::Combine(OCL_TEST_SIZES,
+                                   OCL_PERF_ENUM(std::make_pair<MatType, MatType>(CV_8UC4, CV_32SC4),
+                                                 std::make_pair<MatType, MatType>(CV_32FC1, CV_32FC1)),
+                                   OCL_PERF_ENUM(0, 1),
+                                   ReduceAccOp::all()))
+{
+    const ReduceAccParams params = GetParam();
+    const std::pair<MatType, MatType> types = get<1>(params);
+    const int stype = types.first, dtype = types.second,
+            dim = get<2>(params), op = get<3>(params);
+    const Size srcSize = get<0>(params),
+            dstSize(dim == 0 ? srcSize.width : 1, dim == 0 ? 1 : srcSize.height);
+    const double eps = CV_MAT_DEPTH(dtype) <= CV_32S ? 1 : 3e-4;
+
+    checkDeviceMaxMemoryAllocSize(srcSize, stype);
+    checkDeviceMaxMemoryAllocSize(srcSize, dtype);
+
+    UMat src(srcSize, stype), dst(dstSize, dtype);
+    declare.in(src, WARMUP_RNG).out(dst);
+
+    OCL_TEST_CYCLE() cv::reduce(src, dst, dim, op, dtype);
+
+    SANITY_CHECK(dst, eps);
+}
+
 } } // namespace cvtest::ocl
 
 #endif // HAVE_OPENCL

modules/core/src/matrix.cpp

@@ -3007,23 +3007,80 @@ typedef void (*ReduceFunc)( const Mat& src, Mat& dst );
 #define reduceMinC32f reduceC_<float, float, OpMin<float> >
 #define reduceMinC64f reduceC_<double,double,OpMin<double> >
 
+namespace cv {
+
+static bool ocl_reduce(InputArray _src, OutputArray _dst,
+                       int dim, int op, int op0, int stype, int dtype)
+{
+    int sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype),
+            ddepth = CV_MAT_DEPTH(dtype), ddepth0 = ddepth;
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F))
+        return false;
+
+    if (op == CV_REDUCE_AVG)
+    {
+        op = CV_REDUCE_SUM;
+        if (sdepth < CV_32S && ddepth < CV_32S)
+            ddepth = CV_32S;
+    }
+
+    const char * const ops[4] = { "OCL_CV_REDUCE_SUM", "OCL_CV_REDUCE_AVG",
+                                  "OCL_CV_REDUCE_MAX", "OCL_CV_REDUCE_MIN" };
+    char cvt[40];
+    ocl::Kernel k("reduce", ocl::core::reduce2_oclsrc,
+                  format("-D %s -D dim=%d -D cn=%d -D ddepth=%d -D srcT=%s -D dstT=%s -D convertToDT=%s%s",
+                         ops[op], dim, cn, ddepth, ocl::typeToStr(sdepth), ocl::typeToStr(ddepth),
+                         ocl::convertTypeStr(sdepth, ddepth, 1, cvt),
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
+    UMat src = _src.getUMat();
+    Size dsize(dim == 0 ? src.cols : 1, dim == 0 ? 1 : src.rows);
+    _dst.create(dsize, dtype);
+    UMat dst = _dst.getUMat(), temp = dst;
+
+    if (op0 == CV_REDUCE_AVG && sdepth < CV_32S && ddepth0 < CV_32S)
+        temp.create(dsize, CV_32SC(cn));
+
+    size_t globalsize = std::max(dsize.width, dsize.height);
+
+    k.args(ocl::KernelArg::ReadOnly(src),
+           ocl::KernelArg::WriteOnlyNoSize(temp));
+    if (!k.run(1, &globalsize, NULL, false))
+        return false;
+
+    if (op0 == CV_REDUCE_AVG)
+        temp.convertTo(dst, ddepth0, 1. / (dim == 0 ? src.rows : src.cols));
+
+    return true;
+}
+
+}
+
 void cv::reduce(InputArray _src, OutputArray _dst, int dim, int op, int dtype)
 {
-    Mat src = _src.getMat();
-    CV_Assert( src.dims <= 2 );
+    CV_Assert( _src.dims() <= 2 );
     int op0 = op;
-    int stype = src.type(), sdepth = src.depth(), cn = src.channels();
+    int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype);
     if( dtype < 0 )
         dtype = _dst.fixedType() ? _dst.type() : stype;
+    dtype = CV_MAKETYPE(dtype >= 0 ? dtype : stype, cn);
     int ddepth = CV_MAT_DEPTH(dtype);
 
-    _dst.create(dim == 0 ? 1 : src.rows, dim == 0 ? src.cols : 1,
-                CV_MAKETYPE(dtype >= 0 ? dtype : stype, cn));
-    Mat dst = _dst.getMat(), temp = dst;
-
+    CV_Assert( cn == CV_MAT_CN(dtype) );
     CV_Assert( op == CV_REDUCE_SUM || op == CV_REDUCE_MAX ||
                op == CV_REDUCE_MIN || op == CV_REDUCE_AVG );
-    CV_Assert( src.channels() == dst.channels() );
+
+    if (ocl::useOpenCL() && _dst.isUMat() &&
+            ocl_reduce(_src, _dst, dim, op, op0, stype, dtype))
+        return;
+
+    Mat src = _src.getMat();
+    _dst.create(dim == 0 ? 1 : src.rows, dim == 0 ? src.cols : 1, dtype);
+    Mat dst = _dst.getMat(), temp = dst;
 
     if( op == CV_REDUCE_AVG )
     {

modules/core/src/opencl/reduce2.cl

+/*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.
+// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// 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 copyright holders 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*/
+
+#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
+
+#if ddepth == 0
+#define MIN_VAL 0
+#define MAX_VAL 255
+#elif ddepth == 1
+#define MIN_VAL -128
+#define MAX_VAL 127
+#elif ddepth == 2
+#define MIN_VAL 0
+#define MAX_VAL 65535
+#elif ddepth == 3
+#define MIN_VAL -32768
+#define MAX_VAL 32767
+#elif ddepth == 4
+#define MIN_VAL INT_MIN
+#define MAX_VAL INT_MAX
+#elif ddepth == 5
+#define MIN_VAL (-FLT_MAX)
+#define MAX_VAL FLT_MAX
+#elif ddepth == 6
+#define MIN_VAL (-DBL_MAX)
+#define MAX_VAL DBL_MAX
+#else
+#error "Unsupported depth"
+#endif
+
+#define noconvert
+
+#ifdef OCL_CV_REDUCE_SUM
+#define INIT_VALUE 0
+#define PROCESS_ELEM(acc, value) acc += value
+#elif defined(OCL_CV_REDUCE_MAX)
+#define INIT_VALUE MIN_VAL
+#define PROCESS_ELEM(acc, value) acc = value > acc ? value : acc
+#elif defined(OCL_CV_REDUCE_MIN)
+#define INIT_VALUE MAX_VAL
+#define PROCESS_ELEM(acc, value) acc = value < acc ? value : acc
+#elif defined(OCL_CV_REDUCE_AVG)
+#error "This operation should be implemented through OCL_CV_REDUCE_SUM"
+#else
+#error "No operation is specified"
+#endif
+
+__kernel void reduce(__global const uchar * srcptr, int src_step, int src_offset, int rows, int cols,
+                     __global uchar * dstptr, int dst_step, int dst_offset)
+{
+#if dim == 0 // reduce to a single row
+    int x = get_global_id(0);
+    if (x < cols)
+    {
+        int src_index = x * (int)sizeof(srcT) * cn + src_offset;
+        __global dstT * dst = (__global dstT *)(dstptr + dst_offset) + x * cn;
+        dstT tmp[cn] = { INIT_VALUE };
+
+        for (int y = 0; y < rows; ++y, src_index += src_step)
+        {
+            __global const srcT * src = (__global const srcT *)(srcptr + src_index);
+            #pragma unroll
+            for (int c = 0; c < cn; ++c)
+            {
+                dstT value = convertToDT(src[c]);
+                PROCESS_ELEM(tmp[c], value);
+            }
+        }
+
+        #pragma unroll
+        for (int c = 0; c < cn; ++c)
+            dst[c] = tmp[c];
+    }
+#elif dim == 1 // reduce to a single column
+    int y = get_global_id(0);
+    if (y < rows)
+    {
+        int src_index = mad24(y, src_step, src_offset);
+        int dst_index = mad24(y, dst_step, dst_offset);
+
+        __global const srcT * src = (__global const srcT *)(srcptr + src_index);
+        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
+        dstT tmp[cn] = { INIT_VALUE };
+
+        for (int x = 0; x < cols; ++x, src += cn)
+        {
+            #pragma unroll
+            for (int c = 0; c < cn; ++c)
+            {
+                dstT value = convertToDT(src[c]);
+                PROCESS_ELEM(tmp[c], value);
+            }
+        }
+
+        #pragma unroll
+        for (int c = 0; c < cn; ++c)
+            dst[c] = tmp[c];
+    }
+#else
+#error "Dims must be either 0 or 1"
+#endif
+}

modules/core/test/ocl/test_arithm.cpp

@@ -1566,6 +1566,106 @@ OCL_TEST_P(Psnr, Mat)
     }
 }
 
+//////////////////////////////////////// Reduce /////////////////////////////////////////////
+
+PARAM_TEST_CASE(Reduce, std::pair<MatDepth, MatDepth>, Channels, int, bool)
+{
+    int sdepth, ddepth, cn, dim, dtype;
+    bool use_roi;
+
+    TEST_DECLARE_INPUT_PARAMETER(src)
+    TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+    virtual void SetUp()
+    {
+        const std::pair<MatDepth, MatDepth> p = GET_PARAM(0);
+        sdepth = p.first;
+        ddepth = p.second;
+        cn = GET_PARAM(1);
+        dim = GET_PARAM(2);
+        use_roi = GET_PARAM(3);
+    }
+
+    virtual void generateTestData()
+    {
+        const int stype = CV_MAKE_TYPE(sdepth, cn);
+        dtype = CV_MAKE_TYPE(ddepth, cn);
+
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, stype, -40, 40);
+
+        Size dstRoiSize = Size(dim == 0 ? roiSize.width : 1, dim == 0 ? 1 : roiSize.height);
+        Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, dstRoiSize, dstBorder, dtype, 5, 16);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+    }
+};
+
+typedef Reduce ReduceSum;
+
+OCL_TEST_P(ReduceSum, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::reduce(src_roi, dst_roi, dim, CV_REDUCE_SUM, dtype));
+        OCL_ON(cv::reduce(usrc_roi, udst_roi, dim, CV_REDUCE_SUM, dtype));
+
+        double eps = ddepth <= CV_32S ? 1 : 1e-4;
+        OCL_EXPECT_MATS_NEAR(dst, eps)
+    }
+}
+
+typedef Reduce ReduceMax;
+
+OCL_TEST_P(ReduceMax, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::reduce(src_roi, dst_roi, dim, CV_REDUCE_MAX, dtype));
+        OCL_ON(cv::reduce(usrc_roi, udst_roi, dim, CV_REDUCE_MAX, dtype));
+
+        OCL_EXPECT_MATS_NEAR(dst, 0)
+    }
+}
+
+typedef Reduce ReduceMin;
+
+OCL_TEST_P(ReduceMin, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::reduce(src_roi, dst_roi, dim, CV_REDUCE_MIN, dtype));
+        OCL_ON(cv::reduce(usrc_roi, udst_roi, dim, CV_REDUCE_MIN, dtype));
+
+        OCL_EXPECT_MATS_NEAR(dst, 0)
+    }
+}
+
+typedef Reduce ReduceAvg;
+
+OCL_TEST_P(ReduceAvg, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::reduce(src_roi, dst_roi, dim, CV_REDUCE_AVG, dtype));
+        OCL_ON(cv::reduce(usrc_roi, udst_roi, dim, CV_REDUCE_AVG, dtype));
+
+        double eps = ddepth <= CV_32S ? 1 : 5e-6;
+        OCL_EXPECT_MATS_NEAR(dst, eps)
+    }
+}
+
 //////////////////////////////////////// Instantiation /////////////////////////////////////////
 
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
@@ -1607,6 +1707,42 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, ScaleAdd, Combine(OCL_ALL_DEPTHS, OCL_ALL_CH
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, PatchNaNs, Combine(OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Psnr, Combine(::testing::Values((MatDepth)CV_8U), OCL_ALL_CHANNELS, Bool()));
 
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, ReduceSum, Combine(testing::Values(std::make_pair<MatDepth, MatDepth>(CV_8U, CV_32S),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_8U, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_8U, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16U, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16U, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16S, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16S, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_32F, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_32F, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_64F, CV_64F)),
+                                                       OCL_ALL_CHANNELS, testing::Values(0, 1), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, ReduceAvg, Combine(testing::Values(std::make_pair<MatDepth, MatDepth>(CV_8U, CV_32S),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_8U, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_8U, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16U, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16U, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16S, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16S, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_32F, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_32F, CV_64F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_64F, CV_64F)),
+                                                       OCL_ALL_CHANNELS, testing::Values(0, 1), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, ReduceMax, Combine(testing::Values(std::make_pair<MatDepth, MatDepth>(CV_8U, CV_8U),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16U, CV_16U),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16S, CV_16S),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_32F, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_64F, CV_64F)),
+                                                       OCL_ALL_CHANNELS, testing::Values(0, 1), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, ReduceMin, Combine(testing::Values(std::make_pair<MatDepth, MatDepth>(CV_8U, CV_8U),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16U, CV_16U),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_16S, CV_16S),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_32F, CV_32F),
+                                                                       std::make_pair<MatDepth, MatDepth>(CV_64F, CV_64F)),
+                                                       OCL_ALL_CHANNELS, testing::Values(0, 1), Bool()));
+
+
 } } // namespace cvtest::ocl
 
 #endif // HAVE_OPENCL