#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
{
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
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
#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
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;
if (needMinVal)
needMinLoc = true;
else
- needMaxVal = true;
+ needMaxLoc = true;
}
char cvt[40];
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;
(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);
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)
{