#define noconvert
+#ifdef HAVE_MASK_CONT
+#define MASK_INDEX int mask_index = id + mask_offset;
+#else
+#define MASK_INDEX int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols))
+#endif
+
#if cn != 3
#define loadpix(addr) *(__global const srcT *)(addr)
#define storepix(val, addr) *(__global dstT *)(addr) = val
#ifdef HAVE_MASK
#define REDUCE_GLOBAL \
- int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols)); \
+ MASK_INDEX; \
if (mask[mask_index]) \
{ \
dstT temp = convertToDT(loadpix(srcptr + src_index)); \
FUNC(accumulator, temp); \
}
#elif defined OP_DOT
+
+#ifdef HAVE_SRC2_CONT
+#define SRC2_INDEX int src2_index = mad24(id, srcTSIZE, src2_offset);
+#else
+#define SRC2_INDEX int src2_index = mad24(id / cols, src2_step, mad24(id % cols, srcTSIZE, src2_offset))
+#endif
+
#define REDUCE_GLOBAL \
- int src2_index = mad24(id / cols, src2_step, mad24(id % cols, srcTSIZE, src2_offset)); \
+ SRC2_INDEX; \
dstT temp = convertToDT(loadpix(srcptr + src_index)), temp2 = convertToDT(loadpix(src2ptr + src2_index)); \
FUNC(accumulator, temp, temp2)
#else
#define DEFINE_ACCUMULATOR \
srcT maxval = MIN_VAL, temp
#define REDUCE_GLOBAL \
- int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols)); \
+ MASK_INDEX; \
if (mask[mask_index]) \
{ \
temp = loadpix(srcptr + src_index); \
#define REDUCE_GLOBAL \
temp = loadpix(srcptr + src_index); \
temploc = id; \
- int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols) * (int)sizeof(uchar)); \
+ MASK_INDEX; \
__global const uchar * mask = (__global const uchar *)(maskptr + mask_index); \
temp_mask = mask[0]; \
srcT temp_minval = minval, temp_maxval = maxval; \
int gid = get_group_id(0);
int id = get_global_id(0);
+ srcptr += src_offset;
+
DECLARE_LOCAL_MEM;
DEFINE_ACCUMULATOR;
for (int grain = groupnum * WGS; id < total; id += grain)
{
- int src_index = mad24(id / cols, src_step, mad24(id % cols, srcTSIZE, src_offset));
+#ifdef HAVE_SRC_CONT
+ int src_index = mul24(id, srcTSIZE);
+#else
+ int src_index = mad24(id / cols, src_step, mul24(id % cols, srcTSIZE));
+#endif
REDUCE_GLOBAL;
}
char cvt[40];
ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
format("-D srcT=%s -D srcT1=%s -D dstT=%s -D dstT1=%s -D ddepth=%d -D cn=%d"
- " -D convertToDT=%s -D %s -D WGS=%d -D WGS2_ALIGNED=%d%s%s",
+ " -D convertToDT=%s -D %s -D WGS=%d -D WGS2_ALIGNED=%d%s%s%s%s",
ocl::typeToStr(type), ocl::typeToStr(depth),
ocl::typeToStr(dtype), ocl::typeToStr(ddepth), ddepth, cn,
ocl::convertTypeStr(depth, ddepth, cn, cvt),
opMap[sum_op], (int)wgs, wgs2_aligned,
doubleSupport ? " -D DOUBLE_SUPPORT" : "",
- haveMask ? " -D HAVE_MASK" : ""));
+ haveMask ? " -D HAVE_MASK" : "",
+ _src.isContinuous() ? " -D HAVE_SRC_CONT" : "",
+ _mask.isContinuous() ? " -D HAVE_MASK_CONT" : ""));
if (k.empty())
return false;
wgs2_aligned >>= 1;
ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
- format("-D srcT=%s -D OP_COUNT_NON_ZERO -D WGS=%d -D WGS2_ALIGNED=%d%s",
+ format("-D srcT=%s -D OP_COUNT_NON_ZERO -D WGS=%d "
+ "-D WGS2_ALIGNED=%d%s%s",
ocl::typeToStr(type), (int)wgs,
- wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+ wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "",
+ _src.isContinuous() ? " -D HAVE_SRC_CONT" : ""));
if (k.empty())
return false;
wgs2_aligned <<= 1;
wgs2_aligned >>= 1;
- String opts = format("-D DEPTH_%d -D srcT=%s -D OP_MIN_MAX_LOC%s -D WGS=%d -D WGS2_ALIGNED=%d%s",
+ String opts = format("-D DEPTH_%d -D srcT=%s -D OP_MIN_MAX_LOC%s -D WGS=%d -D WGS2_ALIGNED=%d%s%s%s",
depth, ocl::typeToStr(depth), _mask.empty() ? "" : "_MASK", (int)wgs,
- wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+ wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "",
+ _src.isContinuous() ? " -D HAVE_SRC_CONT" : "",
+ _mask.isContinuous() ? " -D HAVE_MASK_CONT" : "");
ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, opts);
if (k.empty())
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",
+ " -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" : ""));
+ 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;
char cvt[40];
ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
- format("-D srcT=%s -D dstT=%s -D ddepth=%d -D convertToDT=%s -D OP_DOT -D WGS=%d -D WGS2_ALIGNED=%d%s",
+ format("-D srcT=%s -D dstT=%s -D ddepth=%d -D convertToDT=%s -D OP_DOT -D WGS=%d -D WGS2_ALIGNED=%d%s%s%s",
ocl::typeToStr(depth), ocl::typeToStr(ddepth), ddepth, ocl::convertTypeStr(depth, ddepth, 1, cvt),
- (int)wgs, wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+ (int)wgs, wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "",
+ _src1.isContinuous() ? " -D HAVE_SRC_CONT" : "",
+ _src2.isContinuous() ? " -D HAVE_SRC2_CONT" : ""));
if (k.empty())
return false;