Understanding texture fetching in CUDA

Question

I executed following cuda kernel, to understand the texture fetching mechanism but looks like A[i][j] th entry in array should be fetched by tex2D(tex_ref,j,i) Isn't this counter-intuitive ? Am I missing anything here ?

 texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
 surface<void,cudaSurfaceType2D> surfRef;

__global__ void transformKernel(float device_array[3][3],
                                int width, int height,
                                float theta)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    printf("\n  Array : %d %d = %f",x,y,device_array[x][y]);
    printf("\n  Texture : %d %d = %f",x,y,tex2D(texRef,x,y));
    float temp;
    surf2Dread(&temp, surfRef, x*sizeof(float),y);
    printf("\n  Surface : %d %d = %f",x,y,temp);

}

Array is populated by setting host_array[i][j] to (i-j), and copied to device_array and same array was bound to texRef,surfRef.

The output looks like :

Array : 0 0 = 0.000000

Array : 1 0 = 1.000000

Array : 0 1 = -1.000000

Array : 1 1 = 0.000000

Texture : 0 0 = 0.000000

Texture : 1 0 = -1.000000

Texture : 0 1 = 1.000000

Texture : 1 1 = 0.000000

Reads from surface returns similar values as that returned after fetching texture.

Robert Crovella · Accepted Answer · 2016-07-18 21:54:24Z

You're being confused by this:

printf("\n  Array : %d %d = %f",x,y,device_array[x][y]);
                                                 ^^^^

The x texture dimension is the horizontal (i.e. width) dimension in a 2D texturing operation. That means it is effectively selecting a column in the data array. In your printf statement, however, you are using x to select a row.

You may wish to read the documentation which actually gives an example of 2D texturing (note that "texture fetching" is a distinct term from "texturing" but we can ignore that distinction here). In the 2D texture example there, you will see that the "width" dimension index is passed first (i.e. as x) and the height dimension index is passed second (i.e. as y).

The documentation also indicates that the order of texture dimensions in a multidimensional texturing operation is width, height, depth.

You should be able to create comparable results with a reversal of indices on your printout:

printf("\n  Array : %d %d = %f",x,y,device_array[y][x]);
                                                 ^^^^

Collectives™ on Stack Overflow

Understanding texture fetching in CUDA

1 Answer 1

Comments

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

1 Answer 1

Comments

Your Answer

Sign up or log in

Post as a guest

Related