__kernel void sobel(sampler_t sampler, __read_only image2d_t image,
                    __global double *_kernel, __write_only image2d_t result,
                    int w, int h, int k)
{

    int x = get_global_id(0);
    int y = get_global_id(1);

    if (x >= w || y >= h)
        return;

    int4 sum = {0, 0, 0, 255};

    for (int kx = 0; kx < 2*k+1; kx++) {
        for (int ky = 0; ky < 2*k+1; ky++) {
            int overlap_x = x + kx - k;
            int overlap_y = y + ky - k;

            if (overlap_x < 0)
                overlap_x = 0;
            else if (overlap_x >= w)
                overlap_x = w - 1;

            if (overlap_y < 0)
                overlap_y = 0;
            else if (overlap_y >= h)
                overlap_y = h - 1;

            int idxk = kx * 2*k+1 + ky;
            int4 pixel = read_imagei(image, sampler, (int2) { overlap_x, overlap_y });

            sum.x += pixel.x * _kernel[idxk];
            sum.y += pixel.y * _kernel[idxk];
            sum.z += pixel.z * _kernel[idxk];
        }
    }

    write_imagei(result, (int2) { x, y }, sum);
}