Opencl Sobel filter implementation washed out in Float compared to Int

Question

I'm looking at two implementations of a 3x3 vertical sobel filter.One does the computation in float, the other does it as long. For some reason the float version's edges show up in the output visibly washed out. I get that there'll be some precision difference, but seems way to much a of a delta.

This is the long version:

__kernel void sobel(read_only image2d_t inputImage, write_only image2d_t outputImage)
{
    const sampler_t sampler=CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_CLAMP_TO_EDGE|CLK_FILTER_LINEAR;
    int i = get_global_id(0);
    int j = get_global_id(1);
    int accumWeights = 0;
    uint4 pixel_orig=read_imageui(inputImage,sampler,(int2)(i,j));
    long4 pixel_orig_long=convert_long4(pixel_orig);
    long4 pixel_new = (long4)(0, 0, 0, 255);
    long4 pixel0 = convert_long4(read_imageui(inputImage,sampler,(int2)(i+1,j+0)));
    long4 pixel1 = convert_long4(read_imageui(inputImage,sampler,(int2)(i+1,j+1)));
    long4 pixel2 = convert_long4(read_imageui(inputImage,sampler,(int2)(i+1,j-1)));
    long4 pixel3 = convert_long4(read_imageui(inputImage,sampler,(int2)(i-1,j+0)));
    long4 pixel4 = convert_long4(read_imageui(inputImage,sampler,(int2)(i-1,j+1)));
    long4 pixel5 = convert_long4(read_imageui(inputImage,sampler,(int2)(i-1,j-1)));
    pixel_new += (pixel0)*2; 
    pixel_new += (pixel1)*1; 
    pixel_new += (pixel2)*1; 
    pixel_new -= (pixel3)*2; 
    pixel_new -= (pixel4)*1; 
    pixel_new -= (pixel5)*1; 
    if (pixel_new.x<0) pixel_new.x = 0-pixel_new.x; ;
    if (pixel_new.y<0) pixel_new.y = 0-pixel_new.y; ;
    if (pixel_new.z<0) pixel_new.z = 0-pixel_new.z; ;
    pixel_new.w = 255;
    write_imageui(outputImage,(int2)(i,j), convert_uint4(pixel_new));
};

And this is the float version:

__kernel void sobel(read_only image2d_t inputImage, write_only image2d_t outputImage)
{
    const sampler_t sampler=CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_CLAMP_TO_EDGE|CLK_FILTER_LINEAR;
    int i = get_global_id(0);
    int j = get_global_id(1);
    float4 pixel_orig=read_imagef(inputImage,sampler,(int2)(i,j));
    float4 pixel_new = (float4)(0, 0, 0, 255);
    float4 pixel0 = read_imagef(inputImage,sampler,(int2)(i+1,j+0));
    float4 pixel1 = read_imagef(inputImage,sampler,(int2)(i+1,j+1));
    float4 pixel2 = read_imagef(inputImage,sampler,(int2)(i+1,j-1));
    float4 pixel3 = read_imagef(inputImage,sampler,(int2)(i-1,j+0));
    float4 pixel4 = read_imagef(inputImage,sampler,(int2)(i-1,j+1));
    float4 pixel5 = read_imagef(inputImage,sampler,(int2)(i-1,j-1));
    pixel_new += (pixel0)*2;
    pixel_new += (pixel1)*1;
    pixel_new += (pixel2)*1;
    pixel_new -= (pixel3)*2;
    pixel_new -= (pixel4)*1;
    pixel_new -= (pixel5)*1;
    pixel_new = fabs(pixel_new);
    pixel_new.w = 255;
    write_imagef(outputImage,(int2)(i,j), pixel_new);
};

Don't have enough reputation to post the images, but using Lenna as a test the whiteness on the edges are clearly washed out in the float version. Can't seem to figure out why.

Answer 1

A couple things: 1) You're using CLK_FILTER_LINEAR but the OpenCL specification says "The read_imagef calls that take integer coordinates must use a sampler with filter mode set to CLK_FILTER_NEAREST". 2) You're using read_imagef expecting 0-255 but it returns normalize 0.0-1.0 values.