Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

added norm_inf support to minmaxloc kernel

Browse Source
pull/2837/head
Ilya Lavrenov 11 years ago
parent 7f2662b310
commit 634da9f3bf
4 changed files with 26 additions and 67 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 2
      modules/core/include/opencv2/core/opencl/ocl_defs.hpp
  2. 22
      modules/core/src/opencl/minmaxloc.cl
  3. 25
      modules/core/src/opencl/reduce.cl
  4. 44
      modules/core/src/stat.cpp

2
modules/core/include/opencv2/core/opencl/ocl_defs.hpp
Unescape View File

@ -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

22
modules/core/src/opencl/minmaxloc.cl
Unescape View File

@ -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

25
modules/core/src/opencl/reduce.cl
Unescape View File

@ -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

44
modules/core/src/stat.cpp
Unescape View File

@ -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)
{

Write Preview
Loading…
Cancel
Save