mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-29 22:00:25 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Remove unused kernels.

Browse Source
This commit is contained in:
peng xiao 2013-04-12 14:32:13 +08:00
parent aec7a67d0f
commit 7586145235
1 changed files with 0 additions and 182 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

182
modules/ocl/src/opencl/imgproc_canny.cl
View File

@ -360,188 +360,6 @@ __kernel
} }
} }
// non local memory version
__kernel
void calcMap_2
(
__global const int * dx,
__global const int * dy,
__global const float * mag,
__global int * map,
int rows,
int cols,
float low_thresh,
float high_thresh,
int dx_step,
int dx_offset,
int dy_step,
int dy_offset,
int mag_step,
int mag_offset,
int map_step,
int map_offset
)
{
dx_step /= sizeof(*dx);
dx_offset /= sizeof(*dx);
dy_step /= sizeof(*dy);
dy_offset /= sizeof(*dy);
mag_step /= sizeof(*mag);
mag_offset /= sizeof(*mag);
map_step /= sizeof(*map);
map_offset /= sizeof(*map);
int gidx = get_global_id(0);
int gidy = get_global_id(1);
if(gidy < rows && gidx < cols)
{
int x = dx[gidx + gidy * dx_step];
int y = dy[gidx + gidy * dy_step];
const int s = (x ^ y) < 0 ? -1 : 1;
const float m = mag[gidx + 1 + (gidy + 1) * mag_step];
x = abs(x);
y = abs(y);
// 0 - the pixel can not belong to an edge
// 1 - the pixel might belong to an edge
// 2 - the pixel does belong to an edge
int edge_type = 0;
if(m > low_thresh)
{
const int tg22x = x * TG22;
const int tg67x = tg22x + (x << (1 + CANNY_SHIFT));
y <<= CANNY_SHIFT;
if(y < tg22x)
{
if(m > mag[gidx + (gidy + 1) * mag_step] && m >= mag[gidx + 2 + (gidy + 1) * mag_step])
{
edge_type = 1 + (int)(m > high_thresh);
}
}
else if (y > tg67x)
{
if(m > mag[gidx + 1 + gidy* mag_step] && m >= mag[gidx + 1 + (gidy + 2) * mag_step])
{
edge_type = 1 + (int)(m > high_thresh);
}
}
else
{
if(m > mag[gidx + 1 - s + gidy * mag_step] && m > mag[gidx + 1 + s + (gidy + 2) * mag_step])
{
edge_type = 1 + (int)(m > high_thresh);
}
}
}
map[gidx + 1 + (gidy + 1) * map_step] = edge_type;
}
}
// [256, 1, 1] threaded, local memory version
__kernel
void calcMap_3
(
__global const int * dx,
__global const int * dy,
__global const float * mag,
__global int * map,
int rows,
int cols,
float low_thresh,
float high_thresh,
int dx_step,
int dx_offset,
int dy_step,
int dy_offset,
int mag_step,
int mag_offset,
int map_step,
int map_offset
)
{
dx_step /= sizeof(*dx);
dx_offset /= sizeof(*dx);
dy_step /= sizeof(*dy);
dy_offset /= sizeof(*dy);
mag_step /= sizeof(*mag);
mag_offset /= sizeof(*mag);
map_step /= sizeof(*map);
map_offset /= sizeof(*map);
__local float smem[18][18];
int lidx = get_local_id(0) % 16;
int lidy = get_local_id(0) / 16;
int grp_pix = get_global_id(0); // identifies which pixel is processing currently in the target block
int grp_ind = get_global_id(1); // identifies which block of pixels is currently processing
int grp_idx = (grp_ind % (cols/16)) * 16;
int grp_idy = (grp_ind / (cols/16)) * 16; //(grp_ind / (cols/16)) * 16
int gidx = grp_idx + lidx;
int gidy = grp_idy + lidy;
int tid = get_global_id(0) % 256;
int lx = tid % 18;
int ly = tid / 18;
if(ly < 14)
{
smem[ly][lx] = mag[grp_idx + lx + (grp_idy + ly) * mag_step];
}
if(ly < 4 && grp_idy + ly + 14 <= rows && grp_idx + lx <= cols)
{
smem[ly + 14][lx] = mag[grp_idx + lx + (grp_idy + ly + 14) * mag_step];
}
barrier(CLK_LOCAL_MEM_FENCE);
if(gidy < rows && gidx < cols)
{
int x = dx[gidx + gidy * dx_step];
int y = dy[gidx + gidy * dy_step];
const int s = (x ^ y) < 0 ? -1 : 1;
const float m = smem[lidy + 1][lidx + 1];
x = abs(x);
y = abs(y);
// 0 - the pixel can not belong to an edge
// 1 - the pixel might belong to an edge
// 2 - the pixel does belong to an edge
int edge_type = 0;
if(m > low_thresh)
{
const int tg22x = x * TG22;
const int tg67x = tg22x + (x << (1 + CANNY_SHIFT));
y <<= CANNY_SHIFT;
if(y < tg22x)
{
if(m > smem[lidy + 1][lidx] && m >= smem[lidy + 1][lidx + 2])
{
edge_type = 1 + (int)(m > high_thresh);
}
}
else if (y > tg67x)
{
if(m > smem[lidy][lidx + 1]&& m >= smem[lidy + 2][lidx + 1])
{
edge_type = 1 + (int)(m > high_thresh);
}
}
else
{
if(m > smem[lidy][lidx + 1 - s]&& m > smem[lidy + 2][lidx + 1 + s])
{
edge_type = 1 + (int)(m > high_thresh);
}
}
}
map[gidx + 1 + (gidy + 1) * map_step] = edge_type;
}
}
#undef CANNY_SHIFT #undef CANNY_SHIFT
#undef TG22 #undef TG22