CUDA内存--纹理内存的说明与使用

纹理介绍

     纹理存储器（texture memory）是一种只读存储器，由GPU用于纹理渲染的的图形专用单元发展而来，因此也提供了一些特殊功能。
     纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性插值。纹理缓存针对二维空间的局部性访问进行了优化，所以通过纹理缓存访问二维矩阵的邻域会获得加速。纹理缓存不需要满足全局内存的合并访问条件。
     纹理可以是一段连续的设备内存，也可以是一个CUDA数组。但是CUDA数组对局部寻址有优化，称为“块线性”，原理是将邻域元素缓存在同一条cache线上，这将加快邻域内的寻址，但是对于设备内存，并没有“块线性”。所以，选择采用CUDA数组，还是设备内存，需要根据实际情况决定，将数据copy至CUDA数组是很耗时的。
     纹理的创建通过texture object或者texture reference，注意只有计算能力3.0以上的设备支持texture object，本文仅介绍texture reference。

纹理参考系的声明

     纹理参考的一些属性不可变并且在编译时必须知道；它们在声明纹理参考时指定。纹理参考必须在文件域内声明，变量类型为texture；

texture<DataType, Type, ReadMode> texRef;

其中：

 -DataType指定纹理获取时的返回的数据类型，DataType限制为基本的整形和单精度浮点型，或向量类型（具体可查手册）。

-Type指定纹理参考的类型，且等于cudaTextureType1D（一维纹理）,cudaTextureType2D（二维纹理）或cudaTextureType3D（三维纹理），或者cudaTextureType1Dlayered（一维层次纹理）或cudaTextureType2Dlayered（二维层次纹理），Type是可选的，默认为cudaTextureType1D；
-ReadMode等于cudaReadModeNormalizedFloat或cudaReadModeElementType；如果它是cudaReadModeNormalizedFloat且DataType是16位或者8位整形，实际返回值是浮点类型，对于无符号整型，整形全范围被映射到[0.0，1.0]，对于有符号整型，映射成[-1.0，1.0]；例如，无符号八位值为0xff的纹理元素映射为1；如果ReadMode是cudaReadModeElementType，不会进行转换；ReadMode是个可选参数，默认为cudaReadModeElementType。

 

 纹理参考只能被声明为全局静态变量，且不能作为函数的参数传递。

运行时纹理参考属性

     纹理参考的其它属性是可变的，并且能够在运行时通过主机运行时改变。这些属性指定纹理坐标是否归一化、寻址模式和纹理滤波，细节如下。

readmode：设置是否对纹理坐标是否进行归一化。
-归一化浮点模式，normalized=true，fetch将返回【0.0,1.0】的归一化之后的数值，输入数据类型为8位、16位
整型或浮点数；归一化坐标在某些情况下非常方便，如放大缩小操作。
-元素类型模式，normalized=false，fetch返回原始数值。【默认模式】
 addressingmode：寻址模式定义了当纹理坐标越界时发生了什么了。一共四种模式，后两个只能在归一化坐标时使用。
-cudaAddressModeClamp：超出范围就用边界值代替，示意： AA | ABCDE | EE【默认模式】
-cudaAddressModeBorder：超出范围就用零代替，示意： 00 | ABCDE | 00
-cudaAddressModeWrap：重叠模式（循环），示意： DE | ABCDE || AB
-cudaAddressModeMirror：镜像模式，示意： BA | ABCDE | ED
filteringmode：滤波模式，定义了fetch返回结果的计算方式。有两种模式：

-cudaFilterModePoint：点模式，返回最接近的一个点，即最近邻插值。
-cudaFilterModeLinear：线性模式，即线性插值，对于一维纹理，两点插值；对于二维纹理，四点插值；对于三维纹理，八点插值。线性纹理滤波只能对返回值为浮点型的纹理配置起作用。

纹理绑定

     channelDesc描述获取纹理时返回值的格式；channelDesc类型定义如下：

struct cudaChannelFormatDesc{
int x;     //通道0的数据位深度（比特数）
int y;     //通道1的数据位深度
int z;     //通道2的数据位深度
int w;
enum cudaChannelFormatKind f;     //数据类型
};

  其中x、y、z和w 是返回值各组件的位数，而f为：

-cudaChannelFormatKindSigned，如果这些组件是有符号整型；
-cudaChannelFormatKindUnsigned，如果这些组件是无符号整型；
-cudaChannelFormatKindFloat，如果这些组件是浮点类型。
在内核中使用纹理参考从纹理存储器中读取数据之前，对于线性存储器必须使用 cudaBindTexture() 或cudaBindTexture2D()，对于CUDA数组，必须使用cudaBindTextureToArray()，将纹理参考绑定到纹理。cudaUnbindTexture()用于解绑定纹理参考。

示例：

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

 声明纹理参考时指定的参数必须与将纹理绑定到纹理参考时指定的格式匹配；否则纹理获取的结果没有定义。

纹理拾取

     对于一维、二维和三维的CUDA数组绑定的纹理，分别使用tex1D( )、tex2D()和tex3D()函数访问，并且使用浮点型纹理坐标。在算法实现中，这相当于是线性插值。

tex2D()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex2D(texture<DataType, 2, readMode> texRef, float x, float y);

使用纹理坐标x和y获取绑定到纹理参考texRef的CUDA数组或线性存储器区域。

官方示例代码

（来源：CUDA C Programming GuidePG-02829-001_v8.0 | June 2017）

// 2D float texture
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
// Simple transformation kernel
__global__ void transformKernel(float* output,int width, int height,float theta)
{
// Calculate normalized texture coordinates
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;
// Transform coordinates
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;
// Read from texture and write to global memory
output[y * width + x] = tex2D(texRef, tu, tv);
}

// Host code
int main()
{
// Allocate CUDA array in device memory
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 reference parameters
texRef.addressMode[0] = cudaAddressModeWrap;
texRef.addressMode[1] = cudaAddressModeWrap;
texRef.filterMode = cudaFilterModeLinear;
texRef.normalized = true;
// Bind the array to the texture reference
cudaBindTextureToArray(texRef, cuArray, channelDesc);
// Allocate result of transformation in device memory
float* output;
cudaMalloc(&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);
return 0;
}