CUDA内存(五) 纹理内存

纹理内存的特性

纹理内存缓存在芯片上，因此在某些情况中，它能够减少对内存的请求并提供更高效的内存带宽。纹理缓存是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。在某个计算应用程序中，这意味着一个线程读取的位置可能与邻近线程的读取位置“非常接近”，如下图所示。


纹理存储器是一种只读存储器，如果一个kernel中更改了绑定到纹理存储器的数据，纹理缓存中的数据可能并没有被更新，后来读取的数据就可能是错误的，只能重新启动一个kernel，才会更新纹理缓存。对纹理引用的声明不能写成数组!

API

最简单的用法：

//函数外声明： 
texture<float> texConstSrc;
//----------函数内-----
cudaBindTexture(NULL, texConstSrc,d_data, imageSize);
tex1Dfetch(texConstSrc, i);
cudaUnbindTexture(texConstSrc);

Simpletexture例程

// 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数组
     cudaChannelFormatDescchannelDesc = cudaCreateChannelDesc(32, 0, 0,0,cudaChannelFormatKindFloat);
     cudaArray*cuArray;
     cudaMallocArray(&cuArray,&channelDesc, width, height);
     //Copy to device memory some data located at addressh_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 devicememory
     float*output;
     cudaMalloc((void**)&output,width * height * sizeof(float));
     //Invoke kernel
     dim3dimBlock(16, 16);
     dim3dimGrid((width + dimBlock.x –1) / dimBlock.x,(height + dimBlock.y–1) / dimBlock.y);
     transformKernel<<>>(output,width, height,angle);
     //Free device memory
     cudaFreeArray(cuArray);
     cudaFree(output);
}

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

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

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

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

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

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

举个栗子：一个数据类型为char2型，宽×高×深为64×32×16的CUDA3D数组

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

2. 声明纹理参照系

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

exture<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。

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

texture texRef;

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

纹理参照系中的其它属性可以不必声明，并在运行时进行修改。这些参数规定了纹理的寻址模式，是否进行归一化，以及纹理滤波。有C和C++两种风格的接口：

struct textureReference{
     int normalized;
     enum cudaTextureFilterModefilterMode;
     enum cudaTextureAddressModeaddressMode[3];
     struct cudaChannelFormatDescchannelDesc;
}

• 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只支持归一化的纹理坐标。
  a. 对非归一化的坐标，如果寻址的坐标超过了范围[0，N]，大于N的坐标将被钳位，设为N-1。
  b. 对归一化的坐标，有钳位和循环两种处理方式，在钳位方式下，超过[0.0，1.0)范围的坐标将被钳位到[0.0，1.0)；循环方式一般用于周期循环纹理，它只使用了纹理坐标中有用的小数部分，例如1.25会被当作0.25处理，而-1.25则会被当成0.75处理。

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

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

4. 纹理绑定

如果有CUDA数组想要使用纹理内存，需要用cudaBindTexture()或cudaBindTextureToArray()绑定到纹理上。cudaUnbindTexture()用于解除纹理参照系的绑定。

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

cudaMalloc((void**)&devPTr,w1*h1*sizeof(float)); 

• 使用低级API：

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

• 使用高级API

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

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

checkCudaErrors(cudaMallocArray(&cuArray,&channelDesc,width,height));

• 使用低级API：

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

• 使用高级API

texture<<span style="color: blue;">float, 2,cudaReadModeElementType> texRef;
cudaBindTextureToArray(texRef, cuArray);

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

5 纹理拾取

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

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

例如：

output[y* width + x] = tex2D(tex, tu, tv);

---

参考：https://blog.csdn.net/moonboat0331/article/details/10589825