2011-07-11 41 views
7

当我结合的阵列,以在CUDA纹理,CUDA纹理存储器空间

  1. 是阵列拷贝到纹理空间?或者,
  2. 是数组引用作为纹理?

如果答案是1,那么我可以在纹理存储器空间中绑定纹理和安全提取数据,同时将结果写入在全局存储器中分配的数组。

如果答案是2,那么纹理内存是一个全局内存空间,其中的数据被缓存和空间提取?

我想知道这个话题,因为我已经看到一些与这个话题有关的问题,我现在还没有明确的答案。

在此先感谢。

回答

13

答案是第二种选择,但从那里事情会变得更复杂一点。没有“纹理存储器”这样的东西,只有全局存储器可以通过专用硬件访问,其中包括一个GPU读取缓存(根据卡的规格,每MP 6-8kb,参见Cuda编程指南附录F中的表F-2)和一些硬件加速滤波/插值操作。有两种方法可以在CUDA中使用纹理硬件:

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

您可能会发现由David Kanter编写的GT200体系结构概述,以便更好地了解实际体系结构如何实现API公开的内存层次结构。

+0

令人惊叹的答案:)未记录和传闻是我想知道的;)所以,如果我写入原始数组并从绑定纹理中获取数据,行为是不确定的,对吗?谢谢。 – pQB

+3

在CUDA中,纹理和数组是不透明的只读对象。不允许写入(甚至不保证内部纹理布局在不同硬件之间相同)。如果您有费米卡并且正在使用CUDA 3.2或更高版本,则可以使用曲面。虽然没有筛选支持,但[表面API](http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/online/group__CUDART__SURFACE.html)同时提供对“纹理”的读取和写入。目前,AFIAK。 – talonmies

+3

在内核调用中,纹理高速缓存不保持与底层全局内存存储的一致性,请参见编程指南的第3.2.10.4节。使用滤波的2D纹理不必绑定到不透明的cudaArray,它们也可以绑定到使用cudaMallocPitch()分配的音调线性内存。 – njuffa