CUDA学习（四）

一．纹理存储器介绍

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

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

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

二．纹理存储器的使用步骤：

1.       声明CUDA数组，分配空间

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

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

   cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

    cudaArray* cu_array;

    cutilSafeCall( cudaMallocArray( &cu_array, &channelDesc, width, height ));                        //为cu_array开辟空间

cutilSafeCall( cudaMemcpyToArray( cu_array, 0, 0, h_data, size, cudaMemcpyHostToDevice));   //复制到设备端的数组

2.       声明纹理参照系

// declare texture reference for 2D float texture        声明纹理参照系

texture<float, 2, cudaReadModeElementType> tex;

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

tex.addressMode[0] = cudaAddressModeWrap;        //采用循环模式

tex.addressMode[1] = cudaAddressModeWrap;        //说明寻址模式

tex.filterMode = cudaFilterModeLinear;           //设置纹理的滤泼模式

    tex.normalized = true;    // access with normalized texture coordinates设置是否对纹理坐标进行归一化

4.       纹理绑定

// Bind the array to the texture             //纹理绑定

    cutilSafeCall( cudaBindTextureToArray( tex, cu_array, channelDesc));

5.       纹理拾取

g_odata[y*width + x] = tex2D(tex, tu, tv);       //纹理拾取