CUDA纹理内存空间

Question

当我在CUDA中将数组绑定到纹理时，这个数组是被复制到纹理空间中呢，还是作为纹理的引用使用？

如果答案是1.，那么我可以绑定一个纹理，并且在我写入全局内存中的结果时，安全地从纹理内存空间获取数据。

如果答案是2.，那么纹理内存是一个全局内存空间，其中数据被缓存并进行空间提取吗？

我想了解这个主题，因为我看到了一些与此相关的问题，但我现在还没有清楚的答案。

提前致谢。

- pQB

Answer 1

答案是第二个选项，但从那里开始事情变得更加复杂。 "纹理内存"这样的东西不存在，只有通过专用硬件访问的全局内存，其中包括 GPU 上的读取缓存（每个 MP 为 6-8kb，具体取决于卡片，请参见 CUDA 编程指南的附录 F 的表 F-2）和多个硬件加速的过滤/插值操作。CUDA 中可以使用纹理硬件的两种方法：

1. 将线性内存绑定到纹理，并在内核中使用 1D 获取 API 从中读取。在这种情况下，纹理硬件实际上只是充当了一个读取缓存，（如果我没记错的话）没有可用的过滤操作。
2. 创建 CUDA 数组，将线性内存的内容复制到该数组中，并将其绑定到纹理。生成的 CUDA 数组包含线性源的空间有序版本，以某种（未记录的）空间填充曲线存储在全局内存中。纹理硬件提供对该数组的缓存访问，包括硬件加速的同时内存读取。

您可能会发现 David Kanter 写的 GT200 架构概述值得阅读，以更好地了解实际架构如何实现 API 公开的内存层次结构。

- talonmies

---

很棒的回答 :) 我想知道的是未记录和传闻的内容 ;) 因此，如果我写入原始数组并从绑定的纹理中获取数据，则行为未定义，对吗？谢谢。 - pQB

在CUDA中，纹理和数组是不透明的只读对象，不允许写入（甚至不能保证不同硬件之间的内部纹理布局相同）。如果您有一张Fermi显卡并且正在使用CUDA 3.2或更高版本，则可以使用surfaces。surfaces API提供了对“纹理”的读写支持，尽管目前还没有过滤器支持，据我所知。 - talonmies

在内核调用中，纹理缓存与底层全局内存存储不保持一致性，请参见《编程指南》第3.2.10.4节。使用过滤器的二维纹理不必绑定到不透明的cudaArray，也可以绑定到使用cudaMallocPitch()分配的pitch-linear内存。 - njuffa