Arrays of texture references?

I’d like to do something like:

texture<float, 1, cudaReadModeElementType> gpu_texture[2];

I know I could do a 2D texture instead, but I need to bind the texture to device memory rather than an array, so that won’t work.

The construct above doesn’t work. cudaBindTexture complains it’s an invalid texture reference. Even if I try something like:

texture<float, 1, cudaReadModeElementType> gpu_tex1, gpu_tex2;

texture<float, 1, cudaReadModeElementType> gpu_texture[] =

{ gpu_tex1, gpu_text2 };

This doesn’t work either.

I also apparently can’t use texture reference pointers as this causes the nvcc compiler to crash.

texture<float, 1, cudaReadModeElementType> gpu_tex1, gpu_tex2;

__global__ void func()

{

  texture<float, 1, cudaReadModeElementType> *texture =

    threadIdx.y ? &gpu_tex2 : &gpu_tex1;

...

Any suggestions? Thanks!

You can’t do arrays of texture references AFAIK External Image

Not to be a forum nazi, but do a search and you’ll find some nice threads.

In my case, I went with a big texture to hold all the little textures, bound to a cudaArray. Speed is totally not bad, check out the cudaArray stuff if your memory write cycle allows it.

It does not, hence the post.

  1. texture<…> textureRef[2] does not do what is expected, this seems post-worthy

  2. A compiler crash with the latest SDK would seem to be post worthy as well.

In any case, cudaMallocArray cannot create a 2D array large enough to fit my data (it returns “invalid argument”), but I can create two linear blocks with no problem. The fact this is a possibility would seem to indicate arrays of texture references (or texture reference pointers) would be valuable. Again, post-worthy.

If you’re going to be a post Nazi, pick on the people asking what CUDA is. Sheesh.

My tone wasn’t explicitly friendly enough - my intention was to help, not to criticise you for posting. People with more knowledge than me have covered this topic in previous threads, so I refer you to the older threads. I assumed that you knew you couldn’t go the cudaArray route, but I wanted to state that if potential speed issues were holding you back, you should look into it.

IMO, texture references are absolutely essential to have. Workarounds are available, check up on what info is out there already and I’ll be the first person to help you diagnose your problem.

Edit: BTW, the smiley in my original post - External Media - was meant for CUDA and its lack of texture references, not you. Friends? External Media

Further BTW - can you post the compiler crash output? I’ve come across one compiler bug, and that had to do with CUDA kernel variable names being too long (!) - which may need ruling out first.

CUDA doesn’t handle texture arrays, pointers, arguments, references, etc… It only works if you have the texture declared as a plain global variable and use it as such. The reason behind this limit has to do with how textures are referenced in the cubin, but I’m not a cubin expert so I can’t tell you exactly why.

It is unfortunate that the compiler doesn’t produce any sensible warnings or errors for this.

You may have already tried this, but I’ll add it anyways: Can you allocate one 1D texture large enough to hold both your textures? Then you could just index the two textures with tex1Dfetch(tex, idx + offset) and vary offset based on which texture you need to access. The 2^27 elements limit would allow up to a 512 MiB texture of floats: more if you use float4.

That problem is most likely the one I described in the “ptxas crash with testcase” thread in which case there would be no need to report it.

Being able to pass texture references as function parameters would be very nice indeed but I am not sure if the hardware can handle it (though at least some handling of arrays should be only a compiler issue in theory).

Oh wow, where’s my head. :wacko:

Thank you!