CODEWITHSUNDEEP
cudatexturescuda-arrays
user3367017
user3367017

Reputation: 21

CUDA texture as a class member?

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:

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?

Upvotes: 1

Views: 1088

Answers (2)

Roger Dahl
Roger Dahl

Reputation: 15734

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

Upvotes: 1

kunzmi
kunzmi

Reputation: 1024

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.

Upvotes: 2

Related Questions