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?

score 2 · Answer 1 · edited Jun 20 '20 at 09:12

2

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.

edited Jun 20 '20 at 09:12

Community

1
1

answered Mar 01 '14 at 10:58

kunzmi

1,024
1
6
8

A switch statement is needed, but not within the kernel. – Roger Dahl Mar 02 '14 at 19:26
You’re right, when answering the question I thought of using multiple textures in the kernel and switch upon data on a per thread basis. But actually there’s no evidence the questioner wants to do this, why switching to the right texture reference before kernel launch is definitely the better choice. – kunzmi Mar 03 '14 at 08:17
To switch on a per thread basis, the best option would be to combine the textures into a large texture and then manipulate sampling positions within each thread. – Roger Dahl Mar 03 '14 at 09:27

score 1 · Answer 2 · answered Mar 02 '14 at 19:24

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, ...);
}

CUDA texture as a class member?

2 Answers2