Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

fixed compilation of ocl::minMaxLoc for Intel device

Browse Source
pull/1789/head
Ilya Lavrenov 11 years ago
parent a81efdbb25
commit 0bf7350615
2 changed files with 38 additions and 26 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. 11
      modules/ocl/src/opencl/arithm_minMaxLoc.cl
  2. 53
      modules/ocl/src/opencl/arithm_minMaxLoc_mask.cl

11
modules/ocl/src/opencl/arithm_minMaxLoc.cl
Unescape View File

@ -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_max[lid] = max(maxval,localmem_max[lid]);
localmem_minloc[lid] = CONDITION_FUNC(localmem_min[lid] == minval, minloc, localmem_minloc[lid]);
localmem_maxloc[lid] = CONDITION_FUNC(localmem_max[lid] == maxval, maxloc, localmem_maxloc[lid]);
VEC_TYPE minVal = localmem_min[lid], maxVal = localmem_max[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);
@ -234,8 +235,10 @@ __kernel void arithm_op_minMaxLoc(int cols, int invalid_cols, int offset, int el
int lid2 = lsize + lid;
localmem_min[lid] = min(localmem_min[lid], localmem_min[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]);
localmem_maxloc[lid] = CONDITION_FUNC(localmem_max[lid] == localmem_max[lid2], localmem_maxloc[lid2], localmem_maxloc[lid]);
VEC_TYPE min1 = localmem_min[lid], min2 = localmem_min[lid2];
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);
}

53
modules/ocl/src/opencl/arithm_minMaxLoc_mask.cl
Unescape View File

@ -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 idx = id + (id / cols) * invalid_cols;
int midx = id + (id / cols) * minvalid_cols;
__local VEC_TYPE lm_max[128],lm_min[128];
VEC_TYPE minval,maxval,temp,m_temp;
__local VEC_TYPE_LOC lm_maxloc[128],lm_minloc[128];
VEC_TYPE_LOC minloc,maxloc,temploc,negative = -1,one = 1,zero = 0;
VEC_TYPE minval, maxval, temp, m_temp, zeroVal = (VEC_TYPE)(0);
__local VEC_TYPE_LOC lm_maxloc[128], lm_minloc[128];
VEC_TYPE_LOC minloc, maxloc, temploc, negative = -1, one = 1, zero = 0;
if(id < elemnum)
{
temp = vload4(idx, &src[offset]);
m_temp = CONVERT_TYPE(vload4(midx,&mask[moffset]));
int idx_c = (idx << 2) + offset;
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_e(temploc);
}
minval = m_temp != (VEC_TYPE)0 ? temp : (VEC_TYPE)MAX_VAL;
maxval = m_temp != (VEC_TYPE)0 ? temp : (VEC_TYPE)MIN_VAL;
minloc = CONDITION_FUNC(m_temp != (VEC_TYPE)0, temploc , negative);
minval = m_temp != zeroVal ? temp : (VEC_TYPE)MAX_VAL;
maxval = m_temp != zeroVal ? temp : (VEC_TYPE)MIN_VAL;
minloc = CONDITION_FUNC(m_temp != zeroVal, temploc , negative);
maxloc = minloc;
}
else
@ -179,6 +181,7 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
minloc = negative;
maxloc = negative;
}
for(id=id + (groupnum << 8); id < elemnum;id = id + (groupnum << 8))
{
idx = id + (id / cols) * invalid_cols;
@ -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]));
int idx_c = (idx << 2) + offset;
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_e(temploc);
}
minval = min(minval,m_temp != (VEC_TYPE)0 ? temp : minval);
maxval = max(maxval,m_temp != (VEC_TYPE)0 ? temp : maxval);
minval = min(minval, m_temp != zeroVal ? temp : minval);
maxval = max(maxval, m_temp != zeroVal ? temp : maxval);
minloc = CONDITION_FUNC((minval == temp) && (m_temp != (VEC_TYPE)0), temploc , minloc);
maxloc = CONDITION_FUNC((maxval == temp) && (m_temp != (VEC_TYPE)0), temploc , maxloc);
minloc = CONDITION_FUNC(minval == temp && m_temp != zeroVal, temploc , minloc);
maxloc = CONDITION_FUNC(maxval == temp && m_temp != zeroVal, temploc , maxloc);
}
if(lid > 127)
{
lm_min[lid - 128] = minval;
@ -206,32 +210,37 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
lm_maxloc[lid - 128] = maxloc;
}
barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 128)
{
lm_min[lid] = min(minval,lm_min[lid]);
lm_max[lid] = max(maxval,lm_max[lid]);
lm_min[lid] = min(minval, lm_min[lid]);
lm_max[lid] = max(maxval, lm_max[lid]);
VEC_TYPE con_min = CONVERT_TYPE(minloc != 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]);
lm_maxloc[lid] = CONDITION_FUNC((lm_max[lid] == maxval) && (con_max != (VEC_TYPE)0), maxloc , lm_maxloc[lid]);
VEC_TYPE lmMinVal = lm_min[lid], lmMaxVal = lm_max[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);
for(int lsize = 64; lsize > 0; lsize >>= 1)
{
if(lid < lsize)
{
int lid2 = lsize + lid;
lm_min[lid] = min(lm_min[lid] , lm_min[lid2]);
lm_max[lid] = max(lm_max[lid] , lm_max[lid2]);
lm_min[lid] = min(lm_min[lid], lm_min[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_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]);
lm_maxloc[lid] =
CONDITION_FUNC((lm_max[lid] == lm_max[lid2]) && (con_max != (VEC_TYPE)0), lm_maxloc[lid2] , lm_maxloc[lid]);
VEC_TYPE lmMinVal1 = lm_min[lid], lmMinVal2 = lm_min[lid2];
VEC_TYPE lmMaxVal1 = lm_max[lid], lmMaxVal2 = lm_max[lid2];
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);
}
if( lid == 0)
{
dst[gid] = CONVERT_RES_TYPE(lm_min[0]);

Write Preview
Loading…
Cancel
Save