CUDA template kernels and textures

Question

I am working on trying to implement generic kernels in CUDA that make use of texture memory, and I have run into a problem.

template<typename T>
__global__(void){
   tex3D( // correct texture for type T )
}

// host pseudo code
template <typename T>
__host__(void){
    if(T == 'short')
       bind(short_texture);
    else if (T == 'int')
       bind(int_texture);
    invoke_kernel<>(); // <--- How do I tell the kernel which texture was just bound
}

Essentially what I have is the need to access the correctly bound texture based on the template parameter T. I know that I can do a few convoluted things such as write and invoke different kernels, or perhaps pass in a variable indicating which texture to use. I would prefer a cleaner solution. Any suggestions? I would prefer to avoid duplicating the kernels for something so minor, as that would defeat the purpose of the templates.

EDIT:

To clarify, I have template kernels, say a data copy kernel, that operates on global memory of type T. Ergo, short array, int array, etc.. In order to perform copies of any type. I want to move this to use texture memory for other kernels however, I am not sure how I can correctly access the right texture. I have made available, global texture references, applicable to each type that I wish to support, and I have logic to bind the correct texture of the CPU side. My question is, what is the correct way to tell my kernel which texture reference to use in the tex2D function call; the decision of course depends on the template parameter of that kernel (i.e. should I use the float texture, or the int texture). I'm looking for a pattern or a design to follow as I'm not sure of the best way to approach the problem.

What problem do you face; what are the error messages, etc. etc. — Kerrek SB
I believe the problem should be clear - I don't face any error messages because I'm in the design phase, not the implementation phase. I want to have a kernel that has the possibility of using 1 of a few different textures and I'm not sure how to tell it which one to use. — SystemFun
I'm unclear how to access the correctly bound texture in the kernel. If the CPU is able to bind one of say 10 textures, I need a way to tell the kernel which one happens to be bound. — SystemFun

kunzmi kunzmi · Accepted Answer · 2015-08-02T22:35:54

Use texture objects and not texture references. Using texture objects all texture parameters are defined at runtime and not at compile time.

If you need to stick to texture references, another possibility is to wrap the texture fetch calls like this:

template <typename T>
__device__  T myTextureFetch(float x, float y, float z)
{
    return tex3D(tex_ref_to_T_type, x, y, z);
}

(Code is written in browser without checking...) For each type you want to use, you'd need one of these short wrappers...

Further, if you only need the texture as a cached read of global memory, check if the __restrict__ keyword is not better suited for your needs.

CUDA template kernels and textures

1 Answers