opencl- image object clamp to edge

Question

I have two different kernels for my sobel operator. One uses buffer object and the other one uses image object.

In my opinion, these two kernels should make the same result, but they're not.

These two codes handle edges using (clamp to edge)

Where is the problem?

Code with buffer object

__kernel void sobel_filter(__global uchar *ucGRAY, __global float *sobel, __global float *grad_max, int im_width, int im_height)
{
    float2 xt;
    int i = get_global_id(0);
    int j = get_global_id(1);


    int ii_p, jj_p, ii_n, jj_n; // ii_n,jj_n = (i,j)-1 ii_p,jj_p = (i,j)+1


    if (i == 0)
        ii_n = i;
    else if (i == im_width - 1)
        ii_p = i;
    else
    {
        ii_n = i - 1;
        ii_p = i + 1;
    }
    if (j == 0)
        jj_n = i;
    else if (j == im_height - 1)
        jj_p = j;
    else
    {
        jj_n = j - 1;
        jj_p = j + 1;
    }


    xt.x = (float)(ucGRAY[(jj_n)* im_width + (ii_p)] // 3
                + ucGRAY[j * im_width + (ii_p)] * 2 //6
                + ucGRAY[(jj_p) * im_width + (ii_p)] //9

                - ucGRAY[(jj_n)* im_width + (ii_n)] //1
                - ucGRAY[j * im_width + (ii_n)] * 2 //4
                - ucGRAY[(jj_p)* im_width + (ii_n)]) / 1020; //7

    xt.y =(float)(  ucGRAY[(jj_p)* im_width + (ii_n)] //7
                 +ucGRAY[(jj_p)* im_width + (i)] * 2 //8
                 +ucGRAY[(jj_p)* im_width + (ii_p)] //9

                - ucGRAY[(jj_n)* im_width + (ii_n)] //1
                - ucGRAY[(jj_n)* im_width + (i)] * 2 //2
                - ucGRAY[(jj_n)* im_width + (ii_p)]) / 1020; //3 


    sobel[j * im_height + i] = length(xt);

    AtomicMax(grad_max, sobel[j * im_width + i]);
}

Code with image object

const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
                      CLK_ADDRESS_CLAMP_TO_EDGE | //Clamp to edge
                      CLK_FILTER_NEAREST; //Don't interpolate

__kernel void sobel_filter_image(read_only image2d_t ucGRAY,__global float  *sobel,__global float *grad_max,int Width, int Height)

{

     int2 coord = (int2)(get_global_id(0), get_global_id(1));

    float2 xt;
    float temp;

        uchar val5=read_imageui(ucGRAY, smp, (int2)(coord.x,coord.y)).x;


        uchar val1=read_imageui(ucGRAY, smp, (int2)(coord.x-1,coord.y-1)).x;
        uchar val2=read_imageui(ucGRAY, smp, (int2)(coord.x,coord.y-1)).x;
        uchar val3=read_imageui(ucGRAY, smp, (int2)(coord.x+1,coord.y-1)).x;


        uchar val4=read_imageui(ucGRAY, smp, (int2)(coord.x-1,coord.y)).x;      
        uchar val6=read_imageui(ucGRAY, smp, (int2)(coord.x+1,coord.y)).x;

        uchar val7=read_imageui(ucGRAY, smp, (int2)(coord.x-1,coord.y+1)).x;
        uchar val8=read_imageui(ucGRAY, smp, (int2)(coord.x,coord.y+1)).x;
        uchar val9=read_imageui(ucGRAY, smp, (int2)(coord.x+1,coord.y+1)).x;


        xt.x = (float)(val3 + (val6 * 2) + val9 
                    - val1 - (val4 * 2) - val7) / 1020;
        xt.y = (float)(val7 + (val8 * 2) + val9 
                    - val1 - (val2 * 2) - val3) / 1020;
    sobel[coord.y * Width + coord.x] = length(xt);// G=sqrt(Gy^2+Gx^2)

    AtomicMax(grad_max,sobel[coord.y * Width + coord.x]);


}

Answer 1

In your buffer version, you have this:

if (j == 0)
    jj_n = i;

Presumably that should be:

if (j == 0)
    jj_n = j;