diff --git a/modules/core/src/opencl/reduce.cl b/modules/core/src/opencl/reduce.cl index 526cc5101..4f0d80670 100644 --- a/modules/core/src/opencl/reduce.cl +++ b/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); - dst[0] = localmem[0]; + CALC_RESULT; } } diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp index d04857ce8..2806efeb3 100644 --- a/modules/core/src/stat.cpp +++ b/modules/core/src/stat.cpp @@ -1159,16 +1159,137 @@ static void ofs2idx(const Mat& a, size_t ofs, int* idx) } +namespace cv +{ + +template +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::max(); + T max = std::numeric_limits::min() > 0 ? -std::numeric_limits::max() : std::numeric_limits::min(); + int minloc = INT_MAX, maxloc = INT_MAX; + for( int i = 0; i < groupnum; i++) + { + T current_min = minv.at(0,i); + T current_max = maxv.at(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(0,i); + int current_maxloc = maxl.at(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, + getMinMaxRes, + getMinMaxRes, + getMinMaxRes, + getMinMaxRes, + getMinMaxRes, + getMinMaxRes + }; + + 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 ) diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp index 42fc34749..7a24f317a 100644 --- a/modules/core/test/ocl/test_arithm.cpp +++ b/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()));