CUDA 纹理内存空间答案

【问题标题】：CUDA texture memory spaceCUDA 纹理内存空间
【发布时间】：2011-10-02 15:19:34
【问题描述】：

当我将数组绑定到 CUDA 中的纹理时，

该数组是否复制到纹理空间？或者，
该数组引用是纹理吗？

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

如果答案是 2，那么纹理内存是否是一个全局内存空间，用于缓存和空间获取数据？

我想知道这个话题，因为我看到了一些与这个话题相关的问题，但我现在还不清楚答案。

提前致谢。

【问题讨论】：

标签： memory cuda hierarchy texture-mapping

【解决方案1】：

答案是第二种选择，但从那里开始，事情变得更加复杂。没有“纹理内存”之类的东西，只有通过专用硬件访问的全局内存，其中包括 GPU 读取缓存（每个 MP 6-8kb，具体取决于卡，请参阅 Cuda 编程指南附录 F 中的表 F-2）以及一些硬件加速的过滤/插值操作。在 CUDA 中可以通过两种方式使用纹理硬件：

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

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

【讨论】：

惊人的答案 :) 无证和传闻是我想知道的 ;) 所以，如果我写入原始数组并从绑定的纹理中获取数据，行为是未定义的，对吧？谢谢你。
在 CUDA 中，纹理和数组是不透明的只读对象。不允许写入（甚至不能保证内部纹理布局在不同硬件之间是相同的）。如果您有 Fermi 卡并使用 CUDA 3.2 或更高版本，则可以使用表面。 surfaces API 提供对“纹理”的读取和写入，尽管 AFIAK 目前不支持过滤。
在内核调用中，纹理缓存不保持与底层全局内存存储的一致性，请参阅编程指南的第 3.2.10.4 节。使用过滤的 2D 纹理不必绑定到不透明的 cudaArray，它们也可以绑定到使用 cudaMallocPitch() 分配的音高线性内存。