OpenCL image3d linear sampling

Question

I'm doing some linear interpolation with OpenCL, but the result is not as expected. So I did a simple test, kernel code shown below:

const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP_TO_EDGE;
// Kernel block.
kernel void interpolate(
                   global float4*input,
                   image3d_t image,
                   global float4*output)
{
    size_t i = get_global_id(0);
    float4 coord = input[i];
    float4 tap = read_imagef(image, sampler, coord);
    output[i] = tap;
}

Pixels (in RGBA) of the 2x2x2 image is as below:

cl_float4 image_data[8] = {
    {0, 0, 0, 0},
    {100, 0, 0, 0},
    {0, 100, 0, 0},
    {100, 100, 0, 0},
    {0, 0, 100, 0},
    {100, 0, 100, 0},
    {0, 100, 100, 0},
    {100, 100, 100, 0},
};

I used 11 coordinates ((0, 0, 0), (0.1, 0.1, 0.1)...(1, 1, 1), from 0 to 1 with step 0.1) to read the image, and I expect the results to be (0, 0, 0), (10, 10, 10)...(100, 100, 100), but I got:

coordinate:0.000000, result: 0.000000
coordinate:0.100000, result: 0.000000
coordinate:0.200000, result: 0.000000
coordinate:0.300000, result: 10.156250
coordinate:0.400000, result: 30.078125
coordinate:0.500000, result: 50.000000
coordinate:0.600000, result: 69.921875
coordinate:0.700000, result: 89.843750
coordinate:0.800000, result: 100.000000
coordinate:0.900000, result: 100.000000
coordinate:1.000000, result: 100.000000

It just returns edge value when the coordinates are smaller than 0.25 or greater than 0.75.

Anyone could explain this? Thanks.

Answer 1

From: http://www.khronos.org/registry/cl/specs/opencl-1.x-latest.pdf#page=213

"If any of the selected Tijk or Tij in the above equations refers to a location outside the image, the border color is used as the color value for Tijk or Tij."

Below 0.25 and above 0.75 in your case it is performing the interpolation from pixel values outside the image, thus it clamps them to the edge. Therefore all the values are just the edge values.

Why is this so? Because the pixel center of an edge pixel does not lie at 0, instead in your case it lies at 0.25 (0 is the left edge and 0.5 is the right edge). Therefore the values sampled between 0 and 0.25 cause the pixel to be interpolated with itself. You just need to access the array from 0.25 to 0.75 and you are good to go. In case of more pixels you just need to access it differently yet again. As an example with 4 pixels you start it with 0.125 (1/4/2)