在 CUDA 中不灵活地使用纹理引用是否有解决方法答案

【问题标题】：Is there a workaround to the inflexible use of texture references in CUDA在 CUDA 中不灵活地使用纹理引用是否有解决方法
【发布时间】：2012-03-31 10:25:54
【问题描述】：

我有一些我想在内核启动期间读取的纹理。文档说这些必须在全球范围内定义。问题是我想用不同的纹理源调用相同的内核函数（非常复杂）。

texture<unsigned char, 2, cudaReadModeElementType> g_tex_a;
texture<unsigned char, 2, cudaReadModeElementType> g_tex_b;

__global__ void gpu_kernel_1()
{
    // long complicated kernel
    foo = tex2D(g_tex_a, x,y);
}

__global__ void gpu_kernel_2()
{
    // long complicated kernel
    bar = tex2D(g_tex_a, x,y);
}

main()
{
    gpu_kernel_1<<<grid, block>>>();
    gpu_kernel_2<<<grid, block>>>();
}

不可能有一种方法将它应该从 tex2D 读取的纹理传递给内核，例如g_tex_a 还是 g_tex_b？似乎编译器需要在编译时知道纹理引用。我希望能够重新使用 gpu_kernel_1 和 gpu_kernel_2 的代码来作用于不同的纹理。

我不太了解主机和设备代码使用纹理参考对象的方式。目前，我有一个可怕的解决方案，即为每个内核复制所有代码，唯一的变化是 tex2D 函数使用不同的纹理引用，例如gpu_kernel_1_with_tex_a(), gpu_kernel_1_with_tex_b()。

还有其他解决办法吗？谢谢。

【问题讨论】：

你不能向内核传递一个额外的参数并使用if/switch 条件吗？如果所有线程都走相同的路径，则不会有分歧。关于您将纹理作为参数传递的原始问题，我不知道。
@pQB：模板化参数是一个更好的解决方案 - 编译器中的死代码删除将删除给定内核实例中未使用的代码路径。
谢谢。我没有尝试对其进行模板化，因为我预计编译器不会允许内核函数的模板。不过我承认我不确定。
CUDA 从一开始就支持模板。模板内核函数是编写通用内核代码的一种非常常见且高效的方式。

标签： textures cuda

【解决方案1】：

除了纹理参考 API，CUDA 5.0 现在还支持纹理对象 API。无需全局定义纹理引用，使用此新 API，您现在可以在本地定义纹理对象并将其作为函数参数传递给内核。查看 CUDA 编程指南（第 3.2.10.1.1 节）。

【讨论】：

docs.nvidia.com/cuda/cuda-c-programming-guide/…

【解决方案2】：

在使用纹理引用之前，必须将其映射到内存，例如使用 cudaBindTextureToArray()。您可以在映射中决定对哪些数据进行操作。

【讨论】：

好吧，这是有道理的。我忘记了绑定，因为它位于代码的不同部分，只是在某个 init 函数中设置了一次。谢谢！