CUDA texture as a class member?

Question

Attempting to define a class with a per-instance texture. Yes, the number of instances of that class will be small. To work around the restriction that CUDA texture must be a global variable, I tried the following approach:

• Define a global table of textures.
• Add an instance-ID data member to the class.
• Have a class method select a texture from the table using its ID, and pass it as an argument to a CUDA kernel.

Doesn't work. A texture cannot be passed as an argument (nor by pointer or reference), and the kernel doesn't recognize the array name, barring passing by index. I could probably do it with a switch statement, but that is ugly. Any suggestions?

Answer 1

If you have a GPU with Compute Capability >= 3.0, then you can use texture objects instead of texture references. You can then pass the texture object as a kernel/function argument or use it as a class member. See Cuda Programming Guide section B.8 or Texture objects.

In case you don’t have a device with CC 3.0 or above, I guess you’re out of luck and would need a, as you said, "ugly" switch statement in your kernel that chooses the right texture reference depending on some argument.

Answer 2

You can bind the texture that you need before calling the kernel.

So, you have a single texture reference and any number of textures stored in, for instance, cuArrays. Before calling the kernel, you bind the reference to the cuArray that you need:

texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

if (need_texture_1) {
  cudaBindTextureToArray(texRef, cuArray1, ...);
else if (need_texture_2) {
  cudaBindTextureToArray(texRef, cuArray2, ...);
}
kernel<<<>>>();

__global__ void kernel() {
  var = tex2D<float>(texRef, ...);
}