CUDA模板内核和纹理

CUDA template kernels and textures

本文关键字：纹理 内核 CUDA      更新时间：2023-10-16

我正在努力在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类型的全局内存上操作，因此，短数组，int数组等。来执行任何类型的拷贝。我想将此移动到使用其他内核的纹理内存，但是，我不确定如何才能正确访问正确的纹理。我已经提供了可用的全局纹理引用，适用于我希望支持的每种类型，并且我有逻辑来绑定CPU端的正确纹理。我的问题是，什么是正确的方法来告诉我的内核纹理参考使用在tex2D函数调用;当然，这个决定取决于内核的模板参数(例如，我应该使用float纹理，还是int纹理)。我正在寻找一个模式或设计遵循，因为我不确定最好的方法来解决这个问题。

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

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

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

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

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

相关文章：