CUDA学习之纹理内存

1、纹理存储器的特性

    纹理存储器中的数据以一维、二维或者三维数组的形式存储在显存中，可以通过缓存加速访问，并且可以声明大小比常数存储器要大的多。在kernel中访问纹理存储器的操作称为纹理拾取(texture fetching)。将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding).　显存中可以绑定到纹理的数据有两种，分别是普通的线性存储器和cuda数组。线性存储器是用cudaMalloc()声明的一段连续的线性内存空间，而CUDA数组则是通过cudaMallocArray()进行声明，并且是线性对齐的。CUDA数组对于纹理拾取访问进行了优化，但是设备端只能通过纹理拾取访问， 即我们不能直接操作CUDA数组。可以说，CUDA数组就是为纹理拾取而生的。

在通用计算中，纹理存储器十分适合用于实现图像处理或查找表，并且对数据量较大时的随机数据访问或者非对齐访问也有良好的加速效果。

注：线性存储器只能与一维或二维纹理绑定,采用整型纹理拾取坐标，坐标值与数据在存储器中的位置相同；

        CUDA数组可以与一维、二维、三维纹理绑定，纹理拾取坐标为归一化或者非归一化的浮点型，并且支持许多特殊功能。

属性设置：

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

•TextureReference是如下代码所描述的结构体

struct   textureReference{

int normalized;//是否归一化 true or false

enum cudaTextureFilterMode   filterMode;   //滤波模式：cudaFilterModepoint和cudaFilterModeLinear

enum cudaTextureAddressMode   addressMode[3];   //寻址模式：cudaAddressModeClump和cudaAddressModeWarp

struct cudaChannelFormatDesc   channelDesc;

}

例如：

2、纹理缓存的作用

1）纹理内存中的数据可以被重复利用，当需要的数据存在于纹理缓存中，就不用再去显存读取了

2）纹理拾取可以读取纹理坐标附近的几个象元，提高局部性的访问效率，实现滤波模式。换言之，对于图像滤波具有很好的性能。

3、一维纹理存储器的使用方法

1）在host端声明线性内存空间

分配线性内存空间的方法常见有cudaMalloc(), 由cudaMallocPitch()或者cudaMalloc3D()分配的线性空间是经过填充对齐的线性内存。

2）声明纹理参考系

texture<Type, Dim, ReadMode> texRef;//Type指定数据类型，特别注意：不支持3元组//Dim指定纹理参考系的维度，默认为1//ReadMode可以是cudaReadModelNormalizedFloat或cudaReadModelElementType(默认)

注意：type仅限于基本整型、单精度浮点类型和CUDA运行时提供的1元组、2元组和四元祖。

texture<int, 1, cudaReaModeElementType> texRef;这样就声明了一个一维整数的纹理。

3）绑定数据到纹理

通过cudaBindTexture()函数将纹理参考连接到内存。将1维线性内存绑定到1维纹理。

cudaError_t cudaBindTexture(    size_t *offset, const struct textureReference *texref, const void *devPtr,const struct cudaChannelFormatDesc * desc,size_t  size = UINT_MAX )  

例如：cudaBindTexture(0, texRef, d_data);

4）设备端的纹理拾取

在kernel中对纹理存储器进行访问，要通过tex1Dfetch()函数,通过纹理参考和给出的坐标位置就可以取得数据了。

type tex1Dfetch(texture<type,1,ReadMode> texRef, int x);

通过以上四步，我们就可以在程序中利用纹理内存进行数据处理了。例子如下：

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#define N 256//声明纹理参考系texture<int, 1, cudaReadModeElementType> texRef;//CPU串行向量复制void copyCPU(int* source, int* target, int size) {for (int i = 0; i < size; i++) {target[i] = source[i];}}//GPU使用全局内存实现向量复制__global__ void kernel(int* source, int* target, int size) {int index = threadIdx.x + blockDim.x * blockIdx.x;if (index < size) {target[index] = source[index];}}void copyGPU(int* source, int* target, int size) {int Size = size * sizeof(int);int* d_data = NULL;//分配显存，并传递数据到显存上cudaMalloc((void**)&d_data, Size);cudaMemcpy(d_data, source, Size, cudaMemcpyHostToDevice);int* d_Result = NULL;cudaMalloc((void**)&d_Result, Size);//创建Block Grid size dim3 blockNumber(16);dim3 gridNumber((size + blockNumber.x - 1) / blockNumber.x);kernel<< <gridNumber, blockNumber >> > (d_data, d_Result, size);//拷贝数据回主机端cudaMemcpy(target, d_Result, Size, cudaMemcpyDeviceToHost);cudaFree(d_data);cudaFree(d_Result);}//GPU使用纹理内存实现向量复制__global__ void kernelTexture(int* target, int size) {int index = threadIdx.x + blockDim.x * blockIdx.x;if (index < size) {//纹理拾取 tex1Dfetch（纹理参考对象， 位置）target[index] = tex1Dfetch(texRef, index);}}void copyGPUTexture(int* source, int* target, int size) {int Size = N * sizeof(int);//locate Device memory and copy host data to device dataint* d_data = NULL;cudaMalloc((void**)&d_data, Size);cudaMemcpy(d_data, source, Size, cudaMemcpyHostToDevice);int* d_Result = NULL;cudaMalloc((void**)&d_Result, Size);//绑定纹理内存的数据，从全局内存到纹理内存的关联cudaBindTexture(0, texRef, d_data);dim3 blockNumber(16);dim3 gridNumber((size + blockNumber.x - 1) / blockNumber.x);kernelTexture << <gridNumber, blockNumber >> > (d_Result, N);cudaMemcpy(target, d_Result, Size, cudaMemcpyDeviceToHost);//解除纹理绑定cudaUnbindTexture(texRef);cudaFree(d_data);cudaFree(d_Result);}int main(){int *source, *target;source = (int *)malloc(sizeof(int) * N);target = (int *)malloc(sizeof(int) * N);for (int i = 0; i < N; i++) {source[i] = i + 1;target[i] = 0;}cudaEvent_t t1, t2;cudaEventCreate(&t1);cudaEventCreate(&t2);cudaEventRecord(t1, 0);int number = 2;switch (number) {case 0:copyCPU(source, target, N);break;case 1:copyGPU(source, target, N);break;case 2:copyGPUTexture(source, target, N);break;}cudaEventRecord(t2, 0);cudaEventSynchronize(t2);printf("\nSource data:\n");for (int i = 0; i < N; i++) {printf("%5d", source[i]);}printf("\nTarget data:\n");for (int i = 0; i < N; i++) {printf("%5d", target[i]);}printf("\n");float gpuTime = 0;cudaEventElapsedTime(&gpuTime, t1, t2);printf("GPU Time is：%f\n", gpuTime);    return 0;}

4、二维纹理存储器的使用方法

1）声明CUDA数组

使用CUDA数组主要通过三个函数使用：cudaMallocArray(), cudaMemcpyToArray(),cudaFreeArray(). 在声明CUDA数组之前，必须先描述结构体cudaChannelFormatDes()的组件数量和数据类型。

结构体定义：

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

x, y, z和w分别是每个返回值成员的位数，而f是一个枚举变量，可以取一下几个值：
–cudaChannelFormatKindSigned，如果这些成员是有符号整型
· –cudaChannelFormatKindUnsigned，如果这些成员是无符号整型
–cudaChannelFormatKindFloat，如果这些成员是浮点型

CUDA数组的创建方法：

cudaChannelFormatDesc  channelDesc = cudaCreateChannelDesc<float>();//声明数据类型
cudaArray *cuArray;
//分配大小为W*H的CUDA数组
cudaMallocArray(&cuArray, &channelDesc, Width, Height);

CUDA 数组的复制：

        cudaMemcpyToArray(struct cuArray* dstArray, size_t dstX, size_t dstY,  const void *src, size_t  count, enum cudaMemcpyKind kind);  

        函数功能是：把数据src复制到CUDA 数组dstArray中，复制的数据字节大小为count，从src的左上角(dstX,dstY)开始复制。cudaMemcpyKind 复制方向有：hostTohost, hostTodevice, deviceTohost,device todevice(简写方式).

2）声明纹理参考系

//声明一个float 类型的2维的纹理，读取模式为cudaReadModeElementType，也可以声明其他的读取模式
       texture<float, 2, cudaReadModeElementType> texRef;

3）绑定CUDA数组到纹理

调用cudaBindTextureToArray()函数把CUDA数组和纹理连接起来。

4）设备端的纹理拾取

和一维纹理内存的tex1Dfetch()不同，需要使用tex1D()、tex2D()、tex3D()这三个函数，分别是用在1D/2D/3D的纹理。

二维纹理内存使步骤和一维区别不大，主要是利用CUDA数组来进行绑定。具体实例如下：

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#define W 16#define H 16//声明一个float 类型的2维的纹理，读取模式为cudaReadModeElementTypetexture<float, 2, cudaReadModeElementType> texRef;__global__ void addKernel(float *array, const int w, const int h){int idx = threadIdx.x + blockDim.x * blockIdx.x;int idy = threadIdx.y + blockDim.y * blockIdx.y;if (idx < w && idy < h) {//二维纹理索取，根据坐标idx idy进行索引array[idx + idy*w] = tex2D(texRef, idx, idy);}}int main(){float *h_data = (float*)malloc(W*H * sizeof(int));float *h_result = (float *)malloc(W*H * sizeof(int));for (int i = 0; i < H; i++) {for (int j = 0; j < W; j++) {h_data[i*W + j] = i*W + j;}}   //分配CUDA数组cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();//声明数据类型cudaArray *cuArray;//分配大小为W*H的CUDA数组cudaMallocArray(&cuArray, &channelDesc, W, H);//从数据（0，0）开始复制数据，大小为W*H*sizeof(float)cudaMemcpyToArray(cuArray, 0, 0, h_data, W*H*sizeof(float), cudaMemcpyHostToDevice);//设置纹理属性texRef.addressMode[0] = cudaAddressModeWrap;//寻址方式texRef.addressMode[1] = cudaAddressModeWrap;//寻址方式 如果是三维数组则设置texRef.addressMode[2]texRef.normalized = false;//是否对纹理坐标归一化texRef.filterMode = cudaFilterModePoint;//纹理的滤波模式：最近点取样和线性滤波  cudaFilterModeLinear//纹理绑定，CUDA数组和纹理参考的连接cudaBindTextureToArray(&texRef, cuArray, &channelDesc);//设备内存结果float *d_data = NULL;cudaMalloc((void**)&d_data, W*H * sizeof(float));dim3 dimBlock(16, 16);dim3 dimGrid((W + dimBlock.x - 1) / dimBlock.x, (H + dimBlock.y - 1) / dimBlock.y);addKernel << <dimGrid, dimBlock >> > (d_data, W ,H);cudaMemcpy(h_result, d_data, W*H * sizeof(float), cudaMemcpyDeviceToHost);//解除绑定cudaUnbindTexture(texRef);//释放CDUA数组cudaFreeArray(cuArray);//释放显存cudaFree(d_data);printf("Origin data:\n");for (int i = 0; i < W*H; i++) {printf("%8.1f", h_data[i]);}printf("\nResult:\n");for (int i = 0; i < W*H; i++) {printf("%8.1f", h_result[i]);}    return 0;}

总结：网上看了很多博客，但是都讲的不是很有条理，而且实例较少。个人觉得首先搞懂一维纹理内存，对于纹理内存使用的基本步骤熟悉了，然后再去实现二维纹理内存就比较简单了。本文对于归一化坐标没有作说明，由于没有做实验，所以就没写，以后补上，，，

参考：http://blog.csdn.net/augusdi/article/details/12187159

         《基于CUDA的并行程序设计》刘金硕