CUDA中纹理存储器的使用.pdf

CUDA 中纹理存储器的使用 纹理存储器（texture memory ）是一种只读存储器，由GPU 用于纹理渲染的的 图形专用单元发展而来，因此也提供了一些特殊功能。纹理存储器中的数据位 于显存，但可以通过纹理缓存加速读取。在纹理存储器中可以绑定的数据比在 常量存储器可以声明的 64K 大很多，并且支持一维、二维或者三维纹理。在通 用计算中，纹理存储器十分适合用于实现图像处理或查找表，并且对数据量较 大时的随机数据访问或者非对齐访问也有良好的加速效果。 纹理存储器在硬件中并不对应一块专门的存储器，而实际上是牵涉到显存、两 级纹理缓存、纹理抓取单元的纹理流水线。纹理存储器提供了地址映射、数据 滤波、缓存等功能，这些功能都是围绕着纹理渲染的需求设计的。在 CUDA 编 程模型中，纹理缓存是透明的，编程人员不用去了解它的实现机制。 从 CUDA 的内核函数访问纹理存储器的操作被称为纹理拾取(texture fetching) 。 纹理拾取使用的坐标与数据在显存中的地址可以不同，两者通过纹理参照系 texture reference （ ）约定从数据的地址到纹理坐标的映射方式。将显存中的数 据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding) 。显存中 可以绑定到纹理的数据有两种，分别是普通的线性内存（Linear Memroy）和 CUDA 数组（CUDA Array ）。CUDA 数组则为纹理访问进行了优化，并且在 Device 端中只能通过纹理拾取访问。 纹理缓存有两个作用。首先，纹理缓存中的数据可以被重复利用，当一次访问 需要的数据已经存在于纹理缓存中时，就可以避免对显存进行读取。数据重用 过滤了一部分对显存的访问，节约了带宽，也不必按照显存对齐的要求读取。 第二，纹理缓存可以缓存拾取坐标附近几个像元的数据，可以实现滤波模式， 也能提高具有一定局部性的访问的效率。 纹理存储器是只读的，不需要关心缓存数据一致性问题。这意味着如果更改了 绑定到纹理存储器的数据，纹理缓存中的数据可能并没有被更新，此时通过纹 理拾取得到的数据可能是错误的。因此，在每次修改了绑定到纹理的数据以后， 都需要对纹理进行重新绑定。由于不能从设备端修改 CUDA 数组，因此只有在 对绑定到纹理的线性内存进行修改时才需要注意这一点。 使用过程说明： （一）host 端 1）声明实际纹理数据 float2 *hvfield = NULL; float2 *dvfield = NULL; 2) 对 host 数据和 device 数据进行初始化 hvfield = (cData*)malloc(sizeof(cData) * DS); memset(hvfield, 0, sizeof(cData) * DS); // Allocate and initialize device data cudaMallocPitch((void**)dvfield, tPitch, sizeof(cData)*DIM, DIM); cudaMemcpy(dvfield, hvfield, sizeof(cData) * DS, cudaMemcpyHostToDevice); 3 ）调用 device 端函数，建立并绑定纹理数据 setupTexture(DIM, DIM); bindTexture(); // , array texref 注意 其中将 和纹理标识 绑定 （二）device 端 1)声明纹理对象标识，以及纹理实际数据数组 texturefloat2, 2 texref; static cudaArray *array = NULL; 2 ）实现纹理建立、绑定函数 void setupTexture(int x, int y) { // Wrap mode appears to be the new default texref.filterMode = cudaFilterModeLinear; cudaChannelFormatDesc desc = cudaCreateChannelDescfloat2(); cudaMal