#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;
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 )
{
--- /dev/null
+/*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
+}
doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
if (!k.empty())
{
+ UMat src = *this;
_dst.create( size(), _type );
UMat dst = _dst.getUMat();
float alphaf = (float)alpha, betaf = (float)beta;
- k.args(ocl::KernelArg::ReadOnlyNoSize(*this), ocl::KernelArg::WriteOnly(dst, cn), alphaf, betaf);
+ k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn), alphaf, betaf);
size_t globalsize[2] = { dst.cols * cn, dst.rows };
if (k.run(2, globalsize, NULL, false))
}
}
+//////////////////////////////////////// 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 : 5e-5;
+ 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()));
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
projects / profile / ivi / opencv.git / commitdiff