Skip to content

O
opencv
Commit df6bce37 authored by Roman Donchenko's avatar Roman Donchenko Committed by OpenCV Buildbot
Browse files
Options

Merge pull request #1789 from ilya-lavrenov:ocl_minMaxLoc

parents a81efdbb 0bf73506
Hide whitespace changes
Inline Side-by-side
Showing with 38 additions and 26 deletions
modules/ocl/src/opencl/arithm_minMaxLoc.cl
View file @ df6bce37
...@@ -222,8 +222,9 @@ __kernel void arithm_op_minMaxLoc(int cols, int invalid_cols, int offset, int el ...@@ -222,8 +222,9 @@ __kernel void arithm_op_minMaxLoc(int cols, int invalid_cols, int offset, int el
{ {
localmem_min[lid] = min(minval,localmem_min[lid]); localmem_min[lid] = min(minval,localmem_min[lid]);
localmem_max[lid] = max(maxval,localmem_max[lid]); localmem_max[lid] = max(maxval,localmem_max[lid]);
localmem_minloc[lid] = CONDITION_FUNC(localmem_min[lid] == minval, minloc, localmem_minloc[lid]); VEC_TYPE minVal = localmem_min[lid], maxVal = localmem_max[lid];
localmem_maxloc[lid] = CONDITION_FUNC(localmem_max[lid] == maxval, maxloc, localmem_maxloc[lid]); localmem_minloc[lid] = CONDITION_FUNC(minVal == minval, minloc, localmem_minloc[lid]);
localmem_maxloc[lid] = CONDITION_FUNC(maxVal == maxval, maxloc, localmem_maxloc[lid]);
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
...@@ -234,8 +235,10 @@ __kernel void arithm_op_minMaxLoc(int cols, int invalid_cols, int offset, int el ...@@ -234,8 +235,10 @@ __kernel void arithm_op_minMaxLoc(int cols, int invalid_cols, int offset, int el
int lid2 = lsize + lid; int lid2 = lsize + lid;
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]);
localmem_minloc[lid] = CONDITION_FUNC(localmem_min[lid] == localmem_min[lid2], localmem_minloc[lid2], localmem_minloc[lid]); VEC_TYPE min1 = localmem_min[lid], min2 = localmem_min[lid2];
localmem_maxloc[lid] = CONDITION_FUNC(localmem_max[lid] == localmem_max[lid2], localmem_maxloc[lid2], localmem_maxloc[lid]); localmem_minloc[lid] = CONDITION_FUNC(min1 == min2, localmem_minloc[lid2], localmem_minloc[lid]);
VEC_TYPE max1 = localmem_max[lid], max2 = localmem_max[lid2];
localmem_maxloc[lid] = CONDITION_FUNC(max1 == max2, localmem_maxloc[lid2], localmem_maxloc[lid]);
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
......
modules/ocl/src/opencl/arithm_minMaxLoc_mask.cl
View file @ df6bce37
...@@ -152,24 +152,26 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int ...@@ -152,24 +152,26 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
int id = get_global_id(0); int id = get_global_id(0);
int idx = id + (id / cols) * invalid_cols; int idx = id + (id / cols) * invalid_cols;
int midx = id + (id / cols) * minvalid_cols; int midx = id + (id / cols) * minvalid_cols;
__local VEC_TYPE lm_max[128],lm_min[128]; __local VEC_TYPE lm_max[128],lm_min[128];
VEC_TYPE minval,maxval,temp,m_temp; VEC_TYPE minval, maxval, temp, m_temp, zeroVal = (VEC_TYPE)(0);
__local VEC_TYPE_LOC lm_maxloc[128],lm_minloc[128]; __local VEC_TYPE_LOC lm_maxloc[128], lm_minloc[128];
VEC_TYPE_LOC minloc,maxloc,temploc,negative = -1,one = 1,zero = 0; VEC_TYPE_LOC minloc, maxloc, temploc, negative = -1, one = 1, zero = 0;
if(id < elemnum) if(id < elemnum)
{ {
temp = vload4(idx, &src[offset]); temp = vload4(idx, &src[offset]);
m_temp = CONVERT_TYPE(vload4(midx,&mask[moffset])); m_temp = CONVERT_TYPE(vload4(midx,&mask[moffset]));
int idx_c = (idx << 2) + offset; int idx_c = (idx << 2) + offset;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3); temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == cols - 1) if (id % cols == cols - 1)
{ {
repeat_me(m_temp); repeat_me(m_temp);
repeat_e(temploc); repeat_e(temploc);
} }
minval = m_temp != (VEC_TYPE)0 ? temp : (VEC_TYPE)MAX_VAL; minval = m_temp != zeroVal ? temp : (VEC_TYPE)MAX_VAL;
maxval = m_temp != (VEC_TYPE)0 ? temp : (VEC_TYPE)MIN_VAL; maxval = m_temp != zeroVal ? temp : (VEC_TYPE)MIN_VAL;
minloc = CONDITION_FUNC(m_temp != (VEC_TYPE)0, temploc , negative); minloc = CONDITION_FUNC(m_temp != zeroVal, temploc , negative);
maxloc = minloc; maxloc = minloc;
} }
else else
...@@ -179,6 +181,7 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int ...@@ -179,6 +181,7 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
minloc = negative; minloc = negative;
maxloc = negative; maxloc = negative;
} }
for(id=id + (groupnum << 8); id < elemnum;id = id + (groupnum << 8)) for(id=id + (groupnum << 8); id < elemnum;id = id + (groupnum << 8))
{ {
idx = id + (id / cols) * invalid_cols; idx = id + (id / cols) * invalid_cols;
...@@ -187,17 +190,18 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int ...@@ -187,17 +190,18 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
m_temp = CONVERT_TYPE(vload4(midx,&mask[moffset])); m_temp = CONVERT_TYPE(vload4(midx,&mask[moffset]));
int idx_c = (idx << 2) + offset; int idx_c = (idx << 2) + offset;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3); temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == cols - 1) if (id % cols == cols - 1)
{ {
repeat_me(m_temp); repeat_me(m_temp);
repeat_e(temploc); repeat_e(temploc);
} }
minval = min(minval,m_temp != (VEC_TYPE)0 ? temp : minval); minval = min(minval, m_temp != zeroVal ? temp : minval);
maxval = max(maxval,m_temp != (VEC_TYPE)0 ? temp : maxval); maxval = max(maxval, m_temp != zeroVal ? temp : maxval);
minloc = CONDITION_FUNC((minval == temp) && (m_temp != (VEC_TYPE)0), temploc , minloc); minloc = CONDITION_FUNC(minval == temp && m_temp != zeroVal, temploc , minloc);
maxloc = CONDITION_FUNC((maxval == temp) && (m_temp != (VEC_TYPE)0), temploc , maxloc); maxloc = CONDITION_FUNC(maxval == temp && m_temp != zeroVal, temploc , maxloc);
} }
if(lid > 127) if(lid > 127)
{ {
lm_min[lid - 128] = minval; lm_min[lid - 128] = minval;
...@@ -206,32 +210,37 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int ...@@ -206,32 +210,37 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
lm_maxloc[lid - 128] = maxloc; lm_maxloc[lid - 128] = maxloc;
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 128) if(lid < 128)
{ {
lm_min[lid] = min(minval,lm_min[lid]); lm_min[lid] = min(minval, lm_min[lid]);
lm_max[lid] = max(maxval,lm_max[lid]); lm_max[lid] = max(maxval, lm_max[lid]);
VEC_TYPE con_min = CONVERT_TYPE(minloc != negative ? one : zero); VEC_TYPE con_min = CONVERT_TYPE(minloc != negative ? one : zero);
VEC_TYPE con_max = CONVERT_TYPE(maxloc != negative ? one : zero); VEC_TYPE con_max = CONVERT_TYPE(maxloc != negative ? one : zero);
lm_minloc[lid] = CONDITION_FUNC((lm_min[lid] == minval) && (con_min != (VEC_TYPE)0), minloc , lm_minloc[lid]); VEC_TYPE lmMinVal = lm_min[lid], lmMaxVal = lm_max[lid];
lm_maxloc[lid] = CONDITION_FUNC((lm_max[lid] == maxval) && (con_max != (VEC_TYPE)0), maxloc , lm_maxloc[lid]); lm_minloc[lid] = CONDITION_FUNC(lmMinVal == minval && con_min != zeroVal, minloc , lm_minloc[lid]);
lm_maxloc[lid] = CONDITION_FUNC(lmMaxVal == maxval && con_max != zeroVal, maxloc , lm_maxloc[lid]);
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
for(int lsize = 64; lsize > 0; lsize >>= 1) for(int lsize = 64; lsize > 0; lsize >>= 1)
{ {
if(lid < lsize) if(lid < lsize)
{ {
int lid2 = lsize + lid; int lid2 = lsize + lid;
lm_min[lid] = min(lm_min[lid] , lm_min[lid2]); lm_min[lid] = min(lm_min[lid], lm_min[lid2]);
lm_max[lid] = max(lm_max[lid] , lm_max[lid2]); lm_max[lid] = max(lm_max[lid], lm_max[lid2]);
VEC_TYPE con_min = CONVERT_TYPE(lm_minloc[lid2] != negative ? one : zero); VEC_TYPE con_min = CONVERT_TYPE(lm_minloc[lid2] != negative ? one : zero);
VEC_TYPE con_max = CONVERT_TYPE(lm_maxloc[lid2] != negative ? one : zero); VEC_TYPE con_max = CONVERT_TYPE(lm_maxloc[lid2] != negative ? one : zero);
lm_minloc[lid] =
CONDITION_FUNC((lm_min[lid] == lm_min[lid2]) && (con_min != (VEC_TYPE)0), lm_minloc[lid2] , lm_minloc[lid]); VEC_TYPE lmMinVal1 = lm_min[lid], lmMinVal2 = lm_min[lid2];
lm_maxloc[lid] = VEC_TYPE lmMaxVal1 = lm_max[lid], lmMaxVal2 = lm_max[lid2];
CONDITION_FUNC((lm_max[lid] == lm_max[lid2]) && (con_max != (VEC_TYPE)0), lm_maxloc[lid2] , lm_maxloc[lid]); lm_minloc[lid] = CONDITION_FUNC(lmMinVal1 == lmMinVal2 && con_min != zeroVal, lm_minloc[lid2] , lm_minloc[lid]);
lm_maxloc[lid] = CONDITION_FUNC(lmMaxVal1 == lmMaxVal2 && con_max != zeroVal, lm_maxloc[lid2] , lm_maxloc[lid]);
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
if( lid == 0) if( lid == 0)
{ {
dst[gid] = CONVERT_RES_TYPE(lm_min[0]); dst[gid] = CONVERT_RES_TYPE(lm_min[0]);
......
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