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 访问正确绑定的纹理。我知道我可以做一些复杂的事情，例如编写和调用不同的内核，或者可能传入一个指示哪个纹理的变量使用。我更喜欢更清洁的解决方案。有什么建议？我宁愿避免为很小的事情复制内核，因为这会破坏模板的目的。

编辑：

为了澄清，我有模板内核，比如数据复制内核，它在 T 类型的全局内存上运行。Ergo、短数组、int 数组等。为了执行任何类型的副本。我想移动它以将纹理内存用于其他内核，但是我不确定如何正确访问正确的纹理。我已经提供了可用的全局纹理引用，适用于我希望支持的每种类型，并且我有逻辑来绑定 CPU 端的正确纹理。我的问题是，告诉我的内核在 tex2D 函数调用中使用哪个纹理引用的正确方法是什么？决定当然取决于该内核的模板参数（即我应该使用浮点纹理还是 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__ 关键字是否更适合您的需求。