c++ - CUDA 模板内核和纹理

Question

我正在尝试在 CUDA 中实现使用纹理内存的通用内核，但我遇到了一个问题。

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
}

本质上，我需要根据模板参数 T 访问正确绑定(bind)的纹理。我知道我可以做一些复杂的事情，例如编写和调用不同的内核，或者传入一个变量来指示哪个纹理使用。我更喜欢更清洁的解决方案。有什么建议么？我宁愿避免为如此小的东西复制内核，因为这会破坏模板的目的。

编辑:

澄清一下，我有模板内核，比如数据复制内核，它在 T 类型的全局内存上运行。因此，short 数组，int 数组等。为了执行任何类型的复制。我想移动它以将纹理内存用于其他内核，但是我不确定如何正确访问正确的纹理。我已经提供了可用的全局纹理引用，适用于我希望支持的每种类型，并且我有逻辑来绑定(bind) CPU 端的正确纹理。我的问题是，告诉我的内核在 tex2D 函数调用中使用哪个纹理引用的正确方法是什么？当然，决定取决于该内核的模板参数(即我应该使用 float 纹理还是 int 纹理)。我正在寻找可遵循的模式或设计，因为我不确定解决问题的最佳方法。

Answer 1

使用纹理对象而不是纹理引用。使用纹理对象，所有纹理参数都在运行时定义，而不是在编译时定义。

如果您需要坚持使用纹理引用，另一种可能性是像这样包装纹理获取调用:

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

(代码是在没有检查的情况下在浏览器中编写的...)对于您要使用的每种类型，您都需要这些短包装器之一...

此外，如果您只需要纹理作为全局内存的缓存读取，请检查 __restrict__ 关键字是否更适合您的需求。