CODEWITHSUNDEEP
c++templatescuda
SystemFun
SystemFun

Reputation: 1072

CUDA template kernels and textures

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.

Upvotes: 0

Views: 1073

Answers (1)

kunzmi
kunzmi

Reputation: 1024

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.

Upvotes: 1

Related Questions