Bilinear interpolation problem 2D texture fetch linear interpolation won't work

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,
I wrote code to use the CUDA bilinear interpolation to compute a tranlated version of an image (grayscale) with no integer shifts.
The problem is that the produced interpolated image is of poor quality and even worse when there is no shift
the resulting image is not the same with the initial! For example a pixel with the value 16 becomes 16.25 where it should remain the same.
Any help would be valuable.

My code (most significant part):

texture<float, 2, cudaReadModeElementType> my_tex;

void global shift(…,float dx,float dy)
{
…
output[i*Nx+j]=tex2D(texInput,j+dx,i+dy);
}
main()
{
…
Im = image_read(name);//image load

cudaArray *input; //2D cuda array to save the image on the device
cudaChannelFormatDesc tex_descriptor = cudaCreateChannelDesc<float>();

cudaMallocArray(&input, &, Nx, Ny);//2D array allocation
cudaMemcpyToArray(input, 0, 0, Im, Nx*Ny*sizeof(float), cudaMemcpyHostToDevice);//copy to image's array

texInput.filterMode = cudaFilterModeLinear;//use bilinear interpolation
texInput.normalized = false;

cudaBindTextureToArray(my_tex,input);

…
dim3 dimBlock2(64,64);
dim3 dimGrid2( (Nx + dimBlock2.x - 1) / dimBlock2.x,(Nx + dimBlock2.y - 1) / dimBlock2.y);
shift<<<dimGrid2,dimBlock2>>>(out_dev,Nx,Ny,dx,dy);//kernel execute

…

}

2

You probably want to use cudaFilterModeNearest (I think that’s what it’s called…) rather than cudaFilterModeLinear. Linear filtering can do funky things when your samples are not exactly in the texel center, and worse, the actual coords of the center of a texel seem to vary depending on the API, since some of them put 0,0 at the corner of the corner texel, and some of them at the center of the corner texel.

3

Thank you for your imediate reply.

So, you propose that there is no hope to make it work as it should be. Also, nearest neighbor (cudaFilterModePoint) is not efficient for my application (I need linear operation).

4

The problem is solved. From your directions I realised that CUDA places a sample in the center of a texel, so I add 0.5 to each coordinate *. Then interpolation is performed more accurately. And when there is no shift the image remains the same.

(By the way I just found that 2D sinc interpolation outperforms bilinear).

Thank you :D

*this means I use: tex2D(j+dx+0.5,i+dy+05)