diff --git a/modules/ocl/src/opencl/imgproc_canny.cl b/modules/ocl/src/opencl/imgproc_canny.cl index 5ec4465238509b3dd04e076a09511ca4f4d37b27..839c16d92c4dec4d8c6b97ea1cd0a181f39bfb83 100644 --- a/modules/ocl/src/opencl/imgproc_canny.cl +++ b/modules/ocl/src/opencl/imgproc_canny.cl @@ -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 TG22