CUDA纹理存储器的特性及其使用

纹理存储器（texture memory）是一种只读存储器，由GPU用于纹理渲染的的图形专用单元发展而来，因此也提供了一些特殊功能。纹理存储器中的数据位于显存，但可以通过纹理缓存加速读取。在纹理存储器中可以绑定的数据比在常量存储器可以声明的64K大很多，并且支持一维、二维或者三维纹理。在通用计算中，纹理存储器十分适合用于实现图像处理或查找表，并且对数据量较大时的随机数据访问或者非对齐访问也有良好的加速效果。

纹理存储器在硬件中并不对应一块专门的存储器，而实际上是牵涉到显存、两级纹理缓存、纹理抓取单元的纹理流水线。纹理存储器提供了地址映射、数据滤波、缓存等功能，这些功能都是围绕着纹理渲染的需求设计的。关于GPU纹理流水线的介绍可以参考本书3.3.3节。在CUDA编程模型中，纹理缓存是透明的，编程人员不用去了解它的实现机制。

从CUDA的内核函数访问纹理存储器的操作被称为纹理拾取(texture fetching)。纹理拾取使用的坐标与数据在显存中的地址可以不同，两者通过纹理参照系（texture reference）约定从数据的地址到纹理坐标的映射方式。将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding)。显存中可以绑定到纹理的数据有两种，分别是普通的线性内存（Linear Memroy）和CUDA数组（CUDA Array）。CUDA数组则为纹理访问进行了优化，并且在Device端中只能通过纹理拾取访问。

绑定到纹理的线性内存和数组中的元素被称为像元（texels），是texture elements的缩写。像元的数据类型可以是其中的元素可以是CUDA中规定的1，2或者4元组（不能是3元组）的有符号或者无符号8-,16-，32-bit整型或者16-bit(目前只能通过driver API支持)整型，以及32-bit浮点型数据。与CUDA数组绑定的纹理参照系中的元素使用的N-元组数据中的组件数量必须与CUDA数组相同。

纹理缓存有两个作用。首先，纹理缓存中的数据可以被重复利用，当一次访问需要的数据已经存在于纹理缓存中时，就可以避免对显存进行读取。数据重用过滤了一部分对显存的访问，节约了带宽，也不必按照显存对齐的要求读取。第二，纹理缓存可以缓存拾取坐标附近几个像元的数据，可以实现滤波模式，也能提高具有一定局部性的访问的效率。

纹理存储器是只读的，不需要关心缓存数据一致性问题。这意味着如果更改了绑定到纹理存储器的数据，纹理缓存中的数据可能并没有被更新，此时通过纹理拾取得到的数据可能是错误的。因此，在每次修改了绑定到纹理的数据以后，都需要对纹理进行重新绑定。由于不能从设备端修改CUDA数组，因此只有在对绑定到纹理的线性内存进行修改时才需要注意这一点。

线性内存中的数据只能与一维纹理绑定，并且纹理拾取坐标是定点型，坐标的值也与数据在线性内存中的偏移量相同；而CUDA数组可以与一维、二维或者三维纹理绑定，纹理拾取坐标是浮点型，并且支持许多特殊功能。纹理存储器的特殊功能有：

浮点型纹理拾取坐标：使用浮点型的纹理拾取坐标对纹理进行寻址，只对与CUDA数组绑定的存储器有效。地址映射的方式可以是归一化或者非归一化的：使用归一化纹理时，纹理在每个维度上的坐标被映射到浮点数[0.0, 1.0)范围内；使用非归一化纹理坐标时，各个维度上的坐标则被映射到浮点数[0.0, N)的范围内，其中N是纹理在该维度上像元的数量。由于在GPU中通常用浮点计算点的坐标，因此使用浮点数作为纹理拾取坐标更加自然；使用归一化的纹理拾取坐标可以不用关心纹理的实际尺寸，简化了渲染程序的编写。

       寻址模式：寻址模式规定了纹理拾取的输入坐标超出纹理寻址范围时的行为，有钳位模式和循环模式两种。使用钳位模式时，当输入的坐标超出了寻址范围，输入的值将被“钳位”到寻址范围的最大值或者最小值；循环模式只对归一化坐标有效，此时要对超出寻址范围的纹理坐标作求模等处理。例如，对映射到[0.0, 1.0)的归一化纹理坐标，输入拾取坐标1.25，钳位模式会将输入按照0.99999…处理，而循环模式会将输入0.25处理。

       类型转换：如果像元中的数据是8-bit或者16-bit定点型，类型转换功能对拾取的返回值进行类型转换，将其映射到归一化的浮点范围[0.0f, 1.0f]（对无符号整型）或者[ -1.0f, 1.0f]（对有符号整型）。

       滤波：如果将返回类型是浮点型的CUDA数组与纹理绑定，那么就可以对返回的值进行滤波。滤波模式可以是最近点取样模式或者线性滤波模式两种。最近点模式返回与浮点型的纹理抓取坐标最近像元的值，而线性滤波模式则会先取出附近几个像元，然后按照抓取坐标与这几个像元的距离进行线性插值，返回线性插值得到的值。线性滤波可以使纹理渲染得到的画面更加平滑自然。线性滤波需要的插值计算不需要可编程单元参与，提供了额外的浮点处理能力，但精度较低。使用线性滤波模式返回的值经过了插值处理，适合用于图像处理；使用最近点取样模式的返回值不会改变纹理中像元的值，适合用于实现查找表。

       关于纹理拾取模式的详细描述，可以参考附录F。

       使用纹理存储器时，首先要在主机端声明需要绑定到纹理的线性存储器或CUDA数组，并设置好纹理参照系，然后将纹理参照系与线性内存或者CUDA数组绑定。在主机端完成配置工作后，就可以在内核函数中通过纹理抓取函数访问纹理存储器了。

3.2.4.1 CUDA数组

在显存中可以分配的空间有两种：CUDA 数组和线性内存。此外，常数存储器中通过缓存加速读取的数据实际也存在于显存中。CUDA数组和线性内存都可以与纹理参照系绑定，但CUDA数组对纹理拾取访问进行了优化，在设备端也只能通过纹理拾取访问。

声明CUDA数组之前，必须先以结构体channelDesc描述CUDA数组中的数据类型。

struct cudaChannelFormatDesc {      int x, y, z, w;      enum cudaChannelFormatKind f; };

其中，x, y, z和w分别是每个返回值成员的位数，而f是一个枚举变量，可以取一下几个值：

n         cudaChannelFormatKindSigned，如果这些成员是有符号整型；

n         cudaChannelFormatKindUnsigned，如果这些成员是无符号整型；

n          cudaChannelFormatKindFloat，如果这些成员是浮点型；

然后，我们要确定CUDA数组的维度和尺寸。CUDA数组可以通过cudaMalloc3DArray()或cudaMallocArray()函数分配。用cudaMalloc3DArray可以分配一维、二维或者三维的CUDA数组，而cudaMallocArray()一般用于分配二维CUDA数组。在使用完CUDA数组后，要使用cudaFreeArray函数释放显存。

由cudaMalloc3DArray分配的CUDA数组使用cudaMemcpy3D()完成与其他CUDA数组或者线性内存的数据传输。CUDA API中使用结构体cudaExtent描述3D Array和3D线性内存在三个维度上的尺寸，在描述一维、二维和三维数组分别用以下的形式：

cudaextent extent = make_cudaextent([1,8192],0,0);

cudaextent extent = make_cudaextent([1,65535],[1,32768],0);

cudaextent extent = make_cudaextent([1,2048],[1,2048],[1,2048]);

其中方括号[]内为允许的寻址范围。注意到二维CUDA数组的第一个维度的寻址范围大于一维CUDA数组的寻址范围，因此在一维CUDA数组的尺寸不够用时，将二维CUDA数组的第二个维度设为1代替一维CUDA数组，获得更大的寻址范围。

下面是声明一个数据类型为char2型，宽×高×深为64×32×16的CUDA 3D数组，对其初始化，最后释放数组的示例代码：

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 0, 0,cudaChannelFormatKindunsigned); //每个像元由两个char构成

cudaExtent extent = make_cudaextent(64,32,16);//建立cudaExtent结构体，描述CUDA数组的维度和尺寸

cudaArray* cuArray;

cudaMalloc3DArray(&cuArray, &channelDesc, extent); //为cuArray开辟空间

缺

cudaFreeArray(cuArray);

下面则是使用cudaMallocArray声明一个由float型构成，尺寸为64×32的CUDA数组，对其赋值，并最后释放的示例代码：

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindunsigned); //每个像元由一个float构成

cudaArray* cuArray;

cudaMallocArray(&cuArray, &channelDesc, 64, 32); //为cuArray开辟空间

cudaMemcpyToArray(cuArray, 0, 0, h_data, &channelDesc);//第二和第三个参数分别表示在宽度和高度上的偏移量，假设h_data中的数据已经初始化

cudaFreeArray(cuArray);

用于在CUDA数组和主机端或者设备端线性内存，以及在CUDA数组间传输数据的函数还有很多，这些还是还有一些异步调用版本，关于这些函数的具体使用方法请参考CUDA Reference mannual。

3.2.4.1 声明纹理参照系

纹理参照系中的一些属性必须在编译时之前被显示声明。纹理参照系通过一个作用范围为全文件的texture型变量声明：

texture<Type, Dim, ReadMode> texRef;

 

其中，

?         Type 确定了由纹理拾取返回的数据类型；Type可以是B3.1节中描述的任意一种由基本整型或者单精度浮点型组成能的1-，2-或者4-元组向量类型。

?         Dim 确定了纹理参照系的维度，默认为1。

?         ReadMode 可以是cudaReadModeNormalizedFloat或者cudaReadModeElementType。如果ReadMode是cudaReadModeNomalizedFloat，并且Type是16-或者8-bit整型，那么返回的值将是一个浮点数。此时，原来整形的值域会被映射到[0.0，1.0]（对无符号整型），或者[-1.0，1.0]（对有符号整型）。例如，一个值为0xff的8-bit无符号整型会被映射为1.0f。如果使用cudaReadModeElementType,那么就不会对输出进行转换。ReadMode是一个可选参数，如果不写，那么默认就是cudaReadModeElementType。

例如，下面的代码声明了一个二维，像元数据为unsigned char型，但将返回值转换为float型的纹理参照系：

texture<unsigned char, 2, cudaReadModeNormalizedFloat> texRef;

3.2.4.2 设置运行时纹理参照系属性

纹理参照系中的其它属性可以不必声明，并在运行时进行修改。这些参数规定了纹理的寻址模式，是否进行归一化，以及纹理滤波。runtimeAPI拥有底层的C风格和高层的C++风格两种接口。高层API中的texture类型是从底层的textureReference中派生而来的。TextureReference是一个下面的代码描述的结构体。

struct textureReference {      int normalized;      enum cudaTextureFilterMode filterMode;      enum cudaTextureAddressMode addressMode[3];      struct cudaChannelFormatDesc channelDesc; }

 

?        normalized 设置是否对纹理坐标是否进行归一化。如果normalized是一个非零值，那么就会使用归一化到[0，1)的坐标进行寻址，否则对尺寸为width, height, depth的纹理使用坐标[0,width-1], [0,height-1], [0,depth-1]寻址。例如，一个尺寸为64×32的纹理可以通过x维度范围为[0，63]，y维度范围[0,31]的坐标寻址。如果采用归一化方式对尺寸为64×32的纹理进行寻址，在x和y维度上的坐标就都是[0.0,1.0)。这样就可以保证纹理的坐标与纹理的尺寸无关。

?        filterMode用于设置纹理的滤波模式，即如何根据坐标计算返回的纹理值。滤波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。滤波模式为CudaFilterModePoint时，返回值是与坐标最接近的像元的值。CudaFilterModeLinear模式只能对返回值为浮点型的纹理使用，启用这一种模式时将拾取纹理坐标周围的像元，然后根据坐标与这些像元之间的距离进行插值计算。对一维纹理可以使用线性滤波，对二维纹理可以使用双线性滤波。返回值会是对最接近纹理坐标的两个像元（对一维纹理），四个像元（对二维纹理）或者八个像元（对三维纹理）进行插值后得到的值。

?        addressmode说明了寻址模式，即如何处理超出寻址范围的纹理坐标；addressmode是一个大小为3的数组，三个元素分别说明对第一、二、三个纹理坐标的取址模式；取址模式可以是cudaAddressModeClamp或cudaAddressModeWrap中的一种，前者将超出寻址范围的纹理坐标”钳位”到寻址范围内的最大或最小值，后者将超出寻址范围的纹理坐标“折叠”进合理范围。cudaAddressModeWrap只支持归一化的纹理坐标。

对非归一化的坐标，如果寻址的坐标超过了范围[0，N]，大于N的坐标将被钳位，设为N-1。

对归一化的坐标，有钳位和循环两种处理方式，在钳位方式下，超过[0.0，1.0)范围的坐标将被钳位到[0.0，1.0)；循环方式一般用于周期循环纹理，它只使用了纹理坐标中有用的小数部分，例如1.25会被当作0.25处理，而-1.25则会被当成0.75处理。

?        channelDesc描述纹理获取返回值类型，我们已经在3.2.4.1小节讲解CUDA array时介绍过这个结构体。纹理参照系的返回值类型描述必须和与之绑定的CUDA array的数据类型描述相同，或者和与之绑定的线性内存中的元素类型相同。

normalized, addressMode和filterMode可以直接在主机端代码中修改。它们只适用于与CUDA数组绑定的纹理参照系。

 

 

附录D中列出了关于纹理拾取的更多信息。

3.2.4.3 纹理绑定

       在kernel能用纹理参照系从纹理内存中读数据前，纹理参照系必须通过cudaBindTexture()或cudaBindTextureToArray()绑定到纹理上。cudaUnbindTexture()用于解除纹理参照系的绑定。

       以下代码示例绑定一个纹理参照系到devPtr指向的线性内存：

?        使用低级API：

texture<float, 2, cudaReadModeElementType> texRef; textureReference* texRefPtr; cudaGetTextureReference(&texRefPtr, “texRef”); cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); cudaBindTexture2D(0, texRefPtr, devPtr, &channelDesc, width, height, pitch);

?        使用高级API

texture<float, 2, cudaReadModeElementType> texRef; cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); cudaBindTexture2D(0, texRef, devPtr, &channelDesc, width, height, pitch);

 

       以下代码示例绑定纹理参照系到一个CUDA数组cuArray：

?        使用低级API：

texture<float, 2, cudaReadModeElementType> texRef; textureReference* texRefPtr; cudaGetTextureReference(&texRefPtr, “texRef”); cudaChannelFormatDesc channelDesc; cudaGetChannelDesc(&channelDesc, cuArray); cudaBindTextureToArray(texRef, cuArray, &channelDesc);

?        使用高级API

texture<float, 2, cudaReadModeElementType> texRef; cudaBindTextureToArray(texRef, cuArray);

      

       当绑定一个纹理到纹理参照系时，格式必须与声明纹理参照系时的参数匹配；否则，纹理获取的结果是undefined的。

 

3.2.4.4 纹理拾取

纹理拾取函数采用纹理坐标对纹理存储器进行访问。

对与线性内存绑定的纹理，使用texfetch1D函数访问，采用的纹理坐标是整型。由cudaMallocPitch或者cudaMalloc3D分配的线性空间实际上仍然是经过填充、对齐的一维线性空间，因此也用texfetch1D()函数访问。

对与一维、二维和三维CUDA数组绑定的问哪里，分别使用tex1D()、tex2D()和tex3D()函数访问，并且使用浮点型纹理坐标。

关于纹理拾取函数的更多讨论，请见本书附录D.8

 

3.2.4.5 例子分析：Simple texture

// 2D float texture texture<float, 2, cudaReadModeElementType> texRef; // Simple transformation kernel __global__ void transformKernel(float* output, int width, int height, float theta) {      // 根据tid bid计算归一化的拾取坐标      unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;      unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;      float u = x / (float)width;      float v = y / (float)height;      // 旋转拾取坐标      u -= 0.5f;      v -= 0.5f;      float tu = u * cosf(theta) –v * sinf(theta) + 0.5f;      float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;      //从纹理存储器中拾取数据，并写入显存      output[y * width + x] = tex2D(tex, tu, tv); }   // Host code int main() {      // 分配CUDA数组      cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat);      cudaArray* cuArray;      cudaMallocArray(&cuArray, &channelDesc, width, height);      // Copy to device memory some data located at address h_data      // in host memory      cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);      // Set texture parameters      texRef.addressMode[0] = cudaAddressModeWrap; //循环寻址方式      texRef.addressMode[1] = cudaAddressModeWrap;      texRef.filterMode = cudaFilterModeLinear;   //线性滤波，因为这里是一个图像。如果要保持原来的值则千万不要用线性滤波      texRef.normalized = true; //归一化坐标      // Bind the array to the texture      cudaBindTextureToArray(texRef, cuArray, channelDesc);      // Allocate result of transformation in device memory      float* output;      cudaMalloc((void**)&output, width * height * sizeof(float));      // Invoke kernel      dim3 dimBlock(16, 16);      dim3 dimGrid((width + dimBlock.x –1) / dimBlock.x,(height + dimBlock.y –1) / dimBlock.y);      transformKernel<<<dimGrid, dimBlock>>>(output, width, height,angle);      // Free device memory      cudaFreeArray(cuArray);      cudaFree(output); }