Commit 5cf41668 authored by Konstantin Matskevich's avatar Konstantin Matskevich

some fixes

parent 904dac76
...@@ -135,29 +135,27 @@ ...@@ -135,29 +135,27 @@
#define MAX_VAL DBL_MAX #define MAX_VAL DBL_MAX
#endif #endif
#define locT int
#define DECLARE_LOCAL_MEM \ #define DECLARE_LOCAL_MEM \
__local srcT localmem_min[WGS2_ALIGNED]; \ __local srcT localmem_min[WGS2_ALIGNED]; \
__local srcT localmem_max[WGS2_ALIGNED]; \ __local srcT localmem_max[WGS2_ALIGNED]; \
__local locT localmem_minloc[WGS2_ALIGNED]; \ __local int localmem_minloc[WGS2_ALIGNED]; \
__local locT localmem_maxloc[WGS2_ALIGNED] __local int localmem_maxloc[WGS2_ALIGNED]
#define DEFINE_ACCUMULATOR \ #define DEFINE_ACCUMULATOR \
srcT minval = MAX_VAL; \ srcT minval = MAX_VAL; \
srcT maxval = MIN_VAL; \ srcT maxval = MIN_VAL; \
locT negative = (locT)(-1); \ int negative = -1; \
locT minloc = negative; \ int minloc = negative; \
locT maxloc = negative; \ int maxloc = negative; \
srcT temp; \ srcT temp; \
locT temploc int temploc
#define REDUCE_GLOBAL \ #define REDUCE_GLOBAL \
temp = src[0]; \ temp = src[0]; \
temploc = (locT)id; \ temploc = id; \
srcT temp_minval = minval, temp_maxval = maxval; \ srcT temp_minval = minval, temp_maxval = maxval; \
minval = min(minval, temp); \ minval = min(minval, temp); \
maxval = max(maxval, temp); \ maxval = max(maxval, temp); \
minloc = (minval == temp_minval) ? minloc : temploc; \ minloc = (minval == temp_minval) ? (temp_minval == MAX_VAL) ? temploc : minloc : temploc; \
maxloc = (maxval == temp_maxval) ? maxloc : temploc maxloc = (maxval == temp_maxval) ? (temp_maxval == MIN_VAL) ? temploc : maxloc : temploc
#define SET_LOCAL_1 \ #define SET_LOCAL_1 \
localmem_min[lid] = minval; \ localmem_min[lid] = minval; \
localmem_max[lid] = maxval; \ localmem_max[lid] = maxval; \
...@@ -179,43 +177,43 @@ ...@@ -179,43 +177,43 @@
localmem_min[lid] = min(localmem_min[lid], localmem_min[lid2]); \ localmem_min[lid] = min(localmem_min[lid], localmem_min[lid2]); \
localmem_max[lid] = max(localmem_max[lid], localmem_max[lid2]); \ localmem_max[lid] = max(localmem_max[lid], localmem_max[lid2]); \
srcT min1 = localmem_min[lid], min2 = localmem_min[lid2]; \ srcT min1 = localmem_min[lid], min2 = localmem_min[lid2]; \
localmem_minloc[lid] = (min1 == min2) ? (min1 == oldmin) ? min(localmem_minloc[lid2],localmem_minloc[lid]) : \ 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]; \ localmem_minloc[lid2] : localmem_minloc[lid]; \
srcT max1 = localmem_max[lid], max2 = localmem_max[lid2]; \ srcT max1 = localmem_max[lid], max2 = localmem_max[lid2]; \
localmem_maxloc[lid] = (max1 == max2) ? (max1 == oldmax) ? min(localmem_maxloc[lid2],localmem_maxloc[lid]) : \ 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] localmem_maxloc[lid2] : localmem_maxloc[lid]
#define CALC_RESULT \ #define CALC_RESULT \
__global srcT * dstminval = (__global srcT *)(dstptr + (int)sizeof(srcT) * gid); \ __global srcT * dstminval = (__global srcT *)(dstptr + (int)sizeof(srcT) * gid); \
__global srcT * dstmaxval = (__global srcT *)(dstptr2 + (int)sizeof(srcT) * gid); \ __global srcT * dstmaxval = (__global srcT *)(dstptr2 + (int)sizeof(srcT) * gid); \
__global dstlocT * dstminloc = (__global dstlocT *)(dstlocptr + (int)sizeof(dstlocT) * gid); \
__global dstlocT * dstmaxloc = (__global dstlocT *)(dstlocptr2 + (int)sizeof(dstlocT) * gid); \
dstminval[0] = localmem_min[0]; \ dstminval[0] = localmem_min[0]; \
dstmaxval[0] = localmem_max[0]; \ dstmaxval[0] = localmem_max[0]; \
dstminloc[0] = localmem_minloc[0]; \ dstlocptr[gid] = localmem_minloc[0]; \
dstmaxloc[0] = localmem_maxloc[0] dstlocptr2[gid] = localmem_maxloc[0]
#if defined OP_MIN_MAX_LOC_MASK #if defined OP_MIN_MAX_LOC_MASK
#undef DEFINE_ACCUMULATOR #undef DEFINE_ACCUMULATOR
#define DEFINE_ACCUMULATOR \ #define DEFINE_ACCUMULATOR \
srcT minval = MAX_VAL; \ srcT minval = MAX_VAL; \
srcT maxval = MIN_VAL; \ srcT maxval = MIN_VAL; \
locT negative = (locT)(-1); \ int negative = -1; \
locT minloc = negative; \ int minloc = negative; \
locT maxloc = negative; \ int maxloc = negative; \
srcT temp, temp_mask, zeroVal = (srcT)(0); \ srcT temp, temp_mask, zeroVal = (srcT)(0); \
locT temploc int temploc
#undef REDUCE_GLOBAL #undef REDUCE_GLOBAL
#define REDUCE_GLOBAL \ #define REDUCE_GLOBAL \
temp = src[0]; \ temp = src[0]; \
temploc = (locT)id; \ temploc = id; \
int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols) * (int)sizeof(uchar)); \ 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); \ __global const uchar * mask = (__global const uchar *)(maskptr + mask_index); \
temp_mask = mask[0]; \ temp_mask = mask[0]; \
srcT temp_minval = minval, temp_maxval = maxval; \ srcT temp_minval = minval, temp_maxval = maxval; \
minval = (temp_mask == zeroVal) ? minval : min(minval, temp); \ minval = (temp_mask == zeroVal) ? minval : min(minval, temp); \
maxval = (temp_mask == zeroVal) ? maxval : max(maxval, temp); \ maxval = (temp_mask == zeroVal) ? maxval : max(maxval, temp); \
minloc = (temp_mask == zeroVal) ? minloc : (minval == temp_minval) ? minloc : temploc; \ minloc = (temp_mask == zeroVal) ? minloc : (minval == temp_minval) ? (temp_minval == MAX_VAL) ? temploc : minloc : temploc; \
maxloc = (temp_mask == zeroVal) ? maxloc : (maxval == temp_maxval) ? maxloc : temploc maxloc = (temp_mask == zeroVal) ? maxloc : (maxval == temp_maxval) ? (temp_maxval == MIN_VAL) ? temploc : maxloc : temploc
#endif #endif
#else #else
...@@ -224,12 +222,12 @@ ...@@ -224,12 +222,12 @@
#if defined OP_MIN_MAX_LOC #if defined OP_MIN_MAX_LOC
#undef EXTRA_PARAMS #undef EXTRA_PARAMS
#define EXTRA_PARAMS , __global uchar * dstptr2, __global uchar * dstlocptr, __global uchar * dstlocptr2 #define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2
#endif #endif
#if defined OP_MIN_MAX_LOC_MASK #if defined OP_MIN_MAX_LOC_MASK
#undef EXTRA_PARAMS #undef EXTRA_PARAMS
#define EXTRA_PARAMS , __global uchar * dstptr2, __global uchar * dstlocptr, __global uchar * dstlocptr2, \ #define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2, \
__global const uchar * maskptr, int mask_step, int mask_offset __global const uchar * maskptr, int mask_step, int mask_offset, __global int * test
#endif #endif
__kernel void reduce(__global const uchar * srcptr, int step, int offset, int cols, __kernel void reduce(__global const uchar * srcptr, int step, int offset, int cols,
......
...@@ -1180,11 +1180,12 @@ void getMinMaxRes(const Mat &minv, const Mat &maxv, const Mat &minl, const Mat & ...@@ -1180,11 +1180,12 @@ void getMinMaxRes(const Mat &minv, const Mat &maxv, const Mat &minl, const Mat &
{ {
int current_minloc = minl.at<int>(0,i); int current_minloc = minl.at<int>(0,i);
int current_maxloc = maxl.at<int>(0,i); int current_maxloc = maxl.at<int>(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; 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; maxloc = (oldmax == current_max) ? std::min(maxloc, current_maxloc) : (oldmax > current_max) ? maxloc : current_maxloc;
} }
} }
bool zero_mask = (maxloc%cols == -1); bool zero_mask = (maxloc == INT_MAX) || (minloc == INT_MAX);
if(minVal) if(minVal)
*minVal = zero_mask ? 0 : (double)min; *minVal = zero_mask ? 0 : (double)min;
if(maxVal) if(maxVal)
...@@ -1223,7 +1224,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int* ...@@ -1223,7 +1224,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int*
wgs2_aligned <<= 1; wgs2_aligned <<= 1;
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 -D dstlocT=int", 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" : ""); depth, _mask.empty() ? "" : "_MASK", (int)wgs, wgs2_aligned, doubleSupport ? "-D DOUBLE_SUPPORT" : "");
ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, opts); ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, opts);
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment