Binding textures to global variables?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi Everyone,

I’m trying to convert a program currently using global memory to one using texture fetches. Is it possible to bind textures to global device variables? As an example, I have the following code:

__device__ float cell_vars[num_cells*num_hydro_vars];

texture<float> tex;

...

cudaBindTexture(tex, cell_vars, (num_cells*num_hydro_vars*sizeof(float)), 0);

But the cudaBindTexture() call fails. I can get the call to work if I do instead

void *devPtr;

cudaGetSymbolAddress(&devPtr, cell_vars[0]);

cudaBindTexture(tex,devPtr,(num_cells*num_hydro_vars*sizeof(float)), 0);

But then, texfetch does not return the same values as direct access to cell_vars does… Does anyone have a suggestion for what might be wrong?

Thanks,

Peter

2

The texunits expect a special memory layout for the texture. The CUDA array memory is such a layout. You need to use it for the texfetch. So you either store your global data in this layout (see driver API description for strides etc) or you do a GPU->GPU copy from the linear global mem to an additional allocated array mem.

Peter