From 634da9f3bfbb32a6c337623d34fb74a879d147f3 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 9 Jun 2014 15:32:35 +0400 Subject: [PATCH] added norm_inf support to minmaxloc kernel --- .../core/include/opencv2/core/opencl/ocl_defs.hpp | 2 +- modules/core/src/opencl/minmaxloc.cl | 22 +++++++++-- modules/core/src/opencl/reduce.cl | 25 +----------- modules/core/src/stat.cpp | 44 +++------------------- 4 files changed, 26 insertions(+), 67 deletions(-) diff --git a/modules/core/include/opencv2/core/opencl/ocl_defs.hpp b/modules/core/include/opencv2/core/opencl/ocl_defs.hpp index 76d4f84..55f8849 100644 --- a/modules/core/include/opencv2/core/opencl/ocl_defs.hpp +++ b/modules/core/include/opencv2/core/opencl/ocl_defs.hpp @@ -5,7 +5,7 @@ // Copyright (C) 2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. -#define CV_OPENCL_RUN_ASSERT +//#define CV_OPENCL_RUN_ASSERT #ifdef HAVE_OPENCL diff --git a/modules/core/src/opencl/minmaxloc.cl b/modules/core/src/opencl/minmaxloc.cl index 56de655..eb57347 100644 --- a/modules/core/src/opencl/minmaxloc.cl +++ b/modules/core/src/opencl/minmaxloc.cl @@ -41,10 +41,15 @@ #if kercn != 3 #define loadpix(addr) *(__global const srcT *)(addr) -#define srcTSIZE (int)sizeof(srcT1) +#define srcTSIZE (int)sizeof(srcT) #else #define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) -#define srcTSIZE ((int)sizeof(srcT1)) +#define srcTSIZE ((int)sizeof(srcT1) * 3) +#endif + +#ifndef HAVE_MASK +#undef srcTSIZE +#define srcTSIZE (int)sizeof(srcT1) #endif #ifdef NEED_MINLOC @@ -106,7 +111,12 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off { int lid = get_local_id(0); int gid = get_group_id(0); - int id = get_global_id(0) * kercn; + int id = get_global_id(0) +#ifndef HAVE_MASK + * kercn; +#else + ; +#endif srcptr += src_offset; #ifdef HAVE_MASK @@ -150,7 +160,11 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off dstT temp2; #endif - for (int grain = groupnum * WGS * kercn; id < total; id += grain) + for (int grain = groupnum * WGS +#ifndef HAVE_MASK + * kercn +#endif + ; id < total; id += grain) { #ifdef HAVE_MASK #ifdef HAVE_MASK_CONT diff --git a/modules/core/src/opencl/reduce.cl b/modules/core/src/opencl/reduce.cl index 9418cec..888b5df 100644 --- a/modules/core/src/opencl/reduce.cl +++ b/modules/core/src/opencl/reduce.cl @@ -543,32 +543,9 @@ #define CALC_RESULT \ storepix(localmem[0], dstptr + dstTSIZE * gid) -// norm (NORM_INF) with cn > 1 and mask -#elif defined OP_NORM_INF_MASK - -#define DECLARE_LOCAL_MEM \ - __local srcT localmem_max[WGS2_ALIGNED] -#define DEFINE_ACCUMULATOR \ - srcT maxval = MIN_VAL, temp -#define REDUCE_GLOBAL \ - MASK_INDEX; \ - if (mask[mask_index]) \ - { \ - temp = loadpix(srcptr + src_index); \ - maxval = max(maxval, (srcT)(temp >= (srcT)(0) ? temp : -temp)); \ - } -#define SET_LOCAL_1 \ - localmem_max[lid] = maxval -#define REDUCE_LOCAL_1 \ - localmem_max[lid - WGS2_ALIGNED] = max(maxval, localmem_max[lid - WGS2_ALIGNED]) -#define REDUCE_LOCAL_2 \ - localmem_max[lid] = max(localmem_max[lid], localmem_max[lid2]) -#define CALC_RESULT \ - storepix(localmem_max[0], dstptr + dstTSIZE * gid) - #else #error "No operation" -#endif // end of norm (NORM_INF) with cn > 1 and mask +#endif #ifdef OP_DOT #undef EXTRA_PARAMS diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp index 01f50fa..79da3c6 100644 --- a/modules/core/src/stat.cpp +++ b/modules/core/src/stat.cpp @@ -1437,7 +1437,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int* kercn = haveMask ? cn : std::min(4, ocl::predictOptimalVectorWidth(_src)); CV_Assert( (cn == 1 && (!haveMask || _mask.type() == CV_8U)) || - (cn >= 1 && (!haveMask || haveSrc2) && !minLoc && !maxLoc) ); + (cn >= 1 && !minLoc && !maxLoc) ); if (ddepth < 0) ddepth = depth; @@ -1465,7 +1465,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int* if (needMinVal) needMinLoc = true; else - needMaxVal = true; + needMaxLoc = true; } char cvt[40]; @@ -1484,8 +1484,6 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int* haveSrc2 ? " -D HAVE_SRC2" : "", maxVal2 ? " -D OP_CALC2" : "", haveSrc2 && _src2.isContinuous() ? " -D HAVE_SRC2_CONT" : ""); - printf("%s\n", opts.c_str()); - ocl::Kernel k("minmaxloc", ocl::core::minmaxloc_oclsrc, opts); if (k.empty()) return false; @@ -1496,7 +1494,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int* (maxVal2 ? esz : 0)); UMat src = _src.getUMat(), src2 = _src2.getUMat(), db(1, dbsize, CV_8UC1), mask = _mask.getUMat(); - if (cn > 1) + if (cn > 1 && !haveMask) { src = src.reshape(1); src2 = src2.reshape(1); @@ -2181,39 +2179,9 @@ static bool ocl_norm( InputArray _src, int normType, InputArray _mask, double & if (normType == NORM_INF) { - if (cn == 1 || !haveMask) - ocl_minMaxIdx(_src, NULL, &result, NULL, NULL, _mask, - std::max(depth, CV_32S), depth != CV_8U && depth != CV_16U); - else - { - int dbsize = d.maxComputeUnits(); - size_t wgs = d.maxWorkGroupSize(); - - int wgs2_aligned = 1; - while (wgs2_aligned < (int)wgs) - wgs2_aligned <<= 1; - wgs2_aligned >>= 1; - - ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, - format("-D OP_NORM_INF_MASK -D HAVE_MASK -D DEPTH_%d" - " -D srcT=%s -D srcT1=%s -D WGS=%d -D cn=%d -D WGS2_ALIGNED=%d%s%s%s", - depth, ocl::typeToStr(type), ocl::typeToStr(depth), - wgs, cn, wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "", - src.isContinuous() ? " -D HAVE_CONT_SRC" : "", - _mask.isContinuous() ? " -D HAVE_MASK_CONT" : "")); - if (k.empty()) - return false; - - UMat db(1, dbsize, type), mask = _mask.getUMat(); - k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(), - dbsize, ocl::KernelArg::PtrWriteOnly(db), ocl::KernelArg::ReadOnlyNoSize(mask)); - - size_t globalsize = dbsize * wgs; - if (!k.run(1, &globalsize, &wgs, true)) - return false; - - minMaxIdx(db.getMat(ACCESS_READ), NULL, &result, NULL, NULL, noArray()); - } + if (!ocl_minMaxIdx(_src, NULL, &result, NULL, NULL, _mask, + std::max(depth, CV_32S), depth != CV_8U && depth != CV_16U)) + return false; } else if (normType == NORM_L1 || normType == NORM_L2 || normType == NORM_L2SQR) { -- 2.7.4