Read from cudaBindTexture2D

Question

I have been trying store a 2D array in texture memory and read from it via cudaBindTexture2D but the value returned is 0, but I'm not sure if this is the right use of cudaBindTexture2D, and tex2D();

I made a pretty simple code to try it out :

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
texture<uint, cudaTextureType2D, cudaReadModeElementType> tex;
__global__ 
void texture2DTest(int *x){
*x = tex2D(tex,0,0);

}

void initTable(int textureTable[][9]){
int i=0;
int j=0;

for(i=0; i<10; i++){
    for(j=0; j<9; j++){
        textureTable[i][j]=0;
    }
}

textureTable[0][0] = 12;

}

int main (int argc, char ** argv){

int textureTable[10][9];

int *d_x;
int x=2;

size_t pitch;

initTable(textureTable);

cudaMalloc(&d_x, sizeof(int)); 
cudaMemcpy(d_x, &x, sizeof(int), cudaMemcpyHostToDevice);

cudaMallocPitch( (void**)textureTable,&pitch, 9, 10);
cudaChannelFormatDesc desc = cudaCreateChannelDesc<uint>(); 
cudaBindTexture2D(NULL, tex, textureTable, desc, 9, 10, pitch) ;

texture2DTest<<<1,1>>>(d_x);

cudaThreadSynchronize();

cudaMemcpy(&x,d_x, sizeof(int), cudaMemcpyDeviceToHost);

printf(" \n %d \n",x);

cudaUnbindTexture(tex);

return 0;
}

Thank you.

there are a variety of problems with this code (you don't do `cudaMallocPitch` on your host array that presumably contains your data, you are not initializing the data - you never call `initTable`, you are not doing proper cuda error checking...). Perhaps you should study the cuda [pitch linear texture sample code](http://docs.nvidia.com/cuda/cuda-samples/index.html#pitch-linear-texture) which demonstrates usage of `cudaBindTexture2D`. Also, here's a [fully worked simple 3D texturing example](http://stackoverflow.com/questions/25591045/cuda-3d-texture-interpolation). — Robert Crovella, Jan 15 '15 at 15:27
Thank you for the reply, there is no error checking as I am trying to reproduce the problem, and I did forgot to call initTable in the sample code displayed here (i added it now). I do one cudaMallocPitch on line 45, is that part wrong ? thank you. I will however read the links you posted. — Anoracx, Jan 15 '15 at 15:33

sgarizvi · Accepted Answer · 2015-01-15T18:38:20.627

There are quite a few issues in the provided code.

The device memory allocation using cudaMallocPitch is totally broken. You are trying to allocate device memory to a 2D array which is already allocated on the host. Trying to do so will result in memory corruption and undefined behavior. A separate pointer variable is required for device memory allocation and memory should be copied from host to device after allocation.
The third argument of cudaMallocPitch expects width of memory in bytes; not elements.
Textures can only be bound to device memory, so cudaBindTexture expects device memory pointer as input.

Fixing all of the above issues, your final main will look something like this:

int main (int argc, char ** argv)
{

int textureTable[10][9];

int *d_x;
int x = 2;

size_t pitch;

initTable(textureTable);

cudaMalloc(&d_x, sizeof(int)); 
cudaMemcpy(d_x, &x, sizeof(int), cudaMemcpyHostToDevice);

int* d_textureTable; //Device texture table

//Allocate pitch linear memory to device texture table
cudaMallocPitch((void**)&d_textureTable,&pitch, 9 * sizeof(int), 10);

//Use Memcpy2D as the pitch of host and device memory may be different
cudaMemcpy2D(d_textureTable, pitch, textureTable, 9 * sizeof(int), 9 *sizeof(int), 10, cudaMemcpyHostToDevice); 

cudaChannelFormatDesc desc = cudaCreateChannelDesc<uint>(); 
cudaBindTexture2D(NULL, tex, d_textureTable, desc, 9, 10, pitch) ;

texture2DTest<<<1,1>>>(d_x);

cudaThreadSynchronize();

cudaMemcpy(&x,d_x, sizeof(int), cudaMemcpyDeviceToHost);

printf(" \n %d \n",x);

cudaUnbindTexture(tex);
//Don't forget to free the allocated memory
cudaFree(d_textureTable);
cudaFree(d_x);
return 0;
}

Read from cudaBindTexture2D

1 Answers1