Skip to content

O
opencv
Commit 0faac595 authored by Ilya Lavrenov's avatar Ilya Lavrenov
Browse files
Options

fixed minMaxLoc kernel (removed compilation errors)

parent 544c0240
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 40 additions and 166 deletions
modules/ocl/src/arithm.cpp
View file @ 0faac595
This diff is collapsed.
modules/ocl/src/opencl/arithm_minMaxLoc.cl
View file @ 0faac595
...@@ -142,29 +142,35 @@ ...@@ -142,29 +142,35 @@
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics:enable #pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics:enable
/**************************************Array minMax**************************************/ /**************************************Array minMax**************************************/
__kernel void arithm_op_minMaxLoc (int cols,int invalid_cols,int offset,int elemnum,int groupnum,
__kernel void arithm_op_minMaxLoc(int cols, int invalid_cols, int offset, int elemnum, int groupnum,
__global VEC_TYPE *src, __global RES_TYPE *dst) __global VEC_TYPE *src, __global RES_TYPE *dst)
{ {
unsigned int lid = get_local_id(0); unsigned int lid = get_local_id(0);
unsigned int gid = get_group_id(0); unsigned int gid = get_group_id(0);
unsigned int id = get_global_id(0); unsigned int id = get_global_id(0);
unsigned int idx = offset + id + (id / cols) * invalid_cols; unsigned int idx = offset + id + (id / cols) * invalid_cols;
__local VEC_TYPE localmem_max[128],localmem_min[128];
VEC_TYPE minval,maxval,temp; __local VEC_TYPE localmem_max[128], localmem_min[128];
__local VEC_TYPE_LOC localmem_maxloc[128],localmem_minloc[128]; VEC_TYPE minval, maxval, temp;
VEC_TYPE_LOC minloc,maxloc,temploc,negative = -1;
__local VEC_TYPE_LOC localmem_maxloc[128], localmem_minloc[128];
VEC_TYPE_LOC minloc, maxloc, temploc, negative = -1;
int idx_c; int idx_c;
if(id < elemnum)
if (id < elemnum)
{ {
temp = src[idx]; temp = src[idx];
idx_c = idx << 2; idx_c = idx << 2;
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 == 0 )
if (id % cols == 0 )
{ {
repeat_s(temp); repeat_s(temp);
repeat_s(temploc); repeat_s(temploc);
} }
if(id % cols == cols - 1) if (id % cols == cols - 1)
{ {
repeat_e(temp); repeat_e(temp);
repeat_e(temploc); repeat_e(temploc);
...@@ -181,164 +187,33 @@ __kernel void arithm_op_minMaxLoc (int cols,int invalid_cols,int offset,int elem ...@@ -181,164 +187,33 @@ __kernel void arithm_op_minMaxLoc (int cols,int invalid_cols,int offset,int elem
minloc = negative; minloc = negative;
maxloc = negative; maxloc = negative;
} }
float4 aaa;
for(id=id + (groupnum << 8); id < elemnum;id = id + (groupnum << 8)) int grainSize = (groupnum << 8);
for (id = id + grainSize; id < elemnum; id = id + grainSize)
{ {
idx = offset + id + (id / cols) * invalid_cols; idx = offset + id + (id / cols) * invalid_cols;
temp = src[idx]; temp = src[idx];
idx_c = idx << 2; idx_c = idx << 2;
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 == 0 )
if (id % cols == 0 )
{ {
repeat_s(temp); repeat_s(temp);
repeat_s(temploc); repeat_s(temploc);
} }
if(id % cols == cols - 1) if (id % cols == cols - 1)
{ {
repeat_e(temp); repeat_e(temp);
repeat_e(temploc); repeat_e(temploc);
} }
minval = min(minval,temp);
maxval = max(maxval,temp);
minloc = CONDITION_FUNC(minval == temp, temploc , minloc);
maxloc = CONDITION_FUNC(maxval == temp, temploc , maxloc);
aaa= convert_float4(maxval == temp);
maxloc = convert_int4(aaa) ? temploc : maxloc;
}
if(lid > 127)
{
localmem_min[lid - 128] = minval;
localmem_max[lid - 128] = maxval;
localmem_minloc[lid - 128] = minloc;
localmem_maxloc[lid - 128] = maxloc;
}
barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 128)
{
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]);
}
barrier(CLK_LOCAL_MEM_FENCE);
for(int lsize = 64; lsize > 0; lsize >>= 1)
{
if(lid < lsize)
{
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]);
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if( lid == 0)
{
dst[gid] = CONVERT_RES_TYPE(localmem_min[0]);
dst[gid + groupnum] = CONVERT_RES_TYPE(localmem_max[0]);
dst[gid + 2 * groupnum] = CONVERT_RES_TYPE(localmem_minloc[0]);
dst[gid + 3 * groupnum] = CONVERT_RES_TYPE(localmem_maxloc[0]);
}
}
#if defined (REPEAT_S0) minval = min(minval, temp);
#define repeat_ms(a) a = a; maxval = max(maxval, temp);
#endif minloc = CONDITION_FUNC(minval == temp, temploc, minloc);
#if defined (REPEAT_S1) maxloc = CONDITION_FUNC(maxval == temp, temploc, maxloc);
#define repeat_ms(a) a.s0 = 0;
#endif
#if defined (REPEAT_S2)
#define repeat_ms(a) a.s0 = 0;a.s1 = 0;
#endif
#if defined (REPEAT_S3)
#define repeat_ms(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;
#endif
#if defined (REPEAT_E0)
#define repeat_me(a) a = a;
#endif
#if defined (REPEAT_E1)
#define repeat_me(a) a.s3 = 0;
#endif
#if defined (REPEAT_E2)
#define repeat_me(a) a.s3 = 0;a.s2 = 0;
#endif
#if defined (REPEAT_E3)
#define repeat_me(a) a.s3 = 0;a.s2 = 0;a.s1 = 0;
#endif
/**************************************Array minMaxLoc mask**************************************/
/*
__kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int elemnum,int groupnum,__global VEC_TYPE *src,
int minvalid_cols,int moffset,__global uchar4 *mask,__global RES_TYPE *dst)
{
unsigned int lid = get_local_id(0);
unsigned int gid = get_group_id(0);
unsigned int id = get_global_id(0);
unsigned int idx = offset + id + (id / cols) * invalid_cols;
unsigned int midx = moffset + id + (id / cols) * minvalid_cols;
__local VEC_TYPE localmem_max[128],localmem_min[128];
VEC_TYPE minval,maxval,temp,max_val = MAX_VAL,min_val = MIN_VAL,zero = 0,m_temp;
__local VEC_TYPE_LOC localmem_maxloc[128],localmem_minloc[128];
VEC_TYPE_LOC minloc,maxloc,temploc,negative = -1;
if(id < elemnum)
{
temp = src[idx];
m_temp = CONVERT_TYPE(mask[midx]);
int idx_c = idx << 2;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == 0 )
{
repeat_ms(m_temp);
repeat_s(temploc);
}
if(id % cols == cols - 1)
{
repeat_me(m_temp);
repeat_e(temploc);
}
minval = m_temp > zero ? temp : max_val;
maxval = m_temp > zero ? temp : min_val;
minloc = CONDITION_FUNC(m_temp > zero, temploc , negative);
maxloc = minloc;
} }
else
{
minval = MAX_VAL;
maxval = MIN_VAL;
minloc = negative;
maxloc = negative;
}
for(id=id + (groupnum << 8); id < elemnum;id = id + (groupnum << 8))
{
idx = offset + id + (id / cols) * invalid_cols;
midx = moffset + id + (id / cols) * minvalid_cols;
temp = src[idx];
m_temp = CONVERT_TYPE(mask[midx]);
int idx_c = idx << 2;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == 0 )
{
repeat_ms(m_temp);
repeat_s(temploc);
}
if(id % cols == cols - 1)
{
repeat_me(m_temp);
repeat_e(temploc);
}
minval = min(minval,m_temp > zero ? temp : max_val);
maxval = max(maxval,m_temp > zero ? temp : min_val);
temploc = CONDITION_FUNC(m_temp > zero, temploc , negative); if (lid > 127)
minloc = CONDITION_FUNC(minval == temp, temploc , minloc);
maxloc = CONDITION_FUNC(maxval == temp, temploc , maxloc);
}
if(lid > 127)
{ {
localmem_min[lid - 128] = minval; localmem_min[lid - 128] = minval;
localmem_max[lid - 128] = maxval; localmem_max[lid - 128] = maxval;
...@@ -346,29 +221,30 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int ...@@ -346,29 +221,30 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
localmem_maxloc[lid - 128] = maxloc; localmem_maxloc[lid - 128] = maxloc;
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 128)
if (lid < 128)
{ {
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]); 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]); localmem_maxloc[lid] = CONDITION_FUNC(localmem_max[lid] == maxval, maxloc, localmem_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;
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] = localmem_minloc[lid] = CONDITION_FUNC(localmem_min[lid] == localmem_min[lid2], localmem_minloc[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]);
localmem_maxloc[lid] =
CONDITION_FUNC(localmem_max[lid] == localmem_max[lid2], localmem_maxloc[lid2] , localmem_maxloc[lid]);
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
if( lid == 0)
if ( lid == 0)
{ {
dst[gid] = CONVERT_RES_TYPE(localmem_min[0]); dst[gid] = CONVERT_RES_TYPE(localmem_min[0]);
dst[gid + groupnum] = CONVERT_RES_TYPE(localmem_max[0]); dst[gid + groupnum] = CONVERT_RES_TYPE(localmem_max[0]);
...@@ -376,5 +252,3 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int ...@@ -376,5 +252,3 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
dst[gid + 3 * groupnum] = CONVERT_RES_TYPE(localmem_maxloc[0]); dst[gid + 3 * groupnum] = CONVERT_RES_TYPE(localmem_maxloc[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