Remove unused kernels.
authorpeng xiao <hisenxpress@gmail.com>
Fri, 12 Apr 2013 06:32:13 +0000 (14:32 +0800)
committerpeng xiao <hisenxpress@gmail.com>
Fri, 12 Apr 2013 06:32:13 +0000 (14:32 +0800)
modules/ocl/src/opencl/imgproc_canny.cl

index 5ec4465..839c16d 100644 (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 TG22