纹理内存

定义

纹理缓存是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计。只读。

在kernel中访问纹理存储器的操作称为纹理拾取(texture fetching)。

将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding).

　　显存中可以绑定到纹理的数据有两种，分别是普通的线性存储器和cuda数组。

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

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

纹理缓存

Texture mem的访问速度快于Global mem，是因为纹理内存有缓存，只有当缓存没有命中的时候才会去访问设备内存，否则访问纹理缓存具有很小​​的延迟。有如下特点：

（1）纹理缓存中的数据可以被重复利用

（2）纹理缓存一次预取拾取坐标对应位置附近的几个象元，可以实现滤波模式。

应用

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

注意：CUDA数组为纹理访问进行了优化，并且在Device端只能通过纹理拾取访问。

实例如下：

texture<float, 2, cudaReadModeElementType> texRef;// 2D float texture

__global__ void transformKernel(float* output, int width, int height, float theta)

{

     

     unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; // 根据tid bid计算归一化的拾取坐标

     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数组

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 parameters

     texRef.addressMode[0] = cudaAddressModeWrap; //循环寻址方式

     texRef.addressMode[1] = cudaAddressModeWrap;

     texRef.filterMode = cudaFilterModeLinear; //因为这里是一个图像。若要保持原来的值则千万不要用线性滤波

     texRef.normalized = true; //归一化坐标

     

     cudaBindTextureToArray(texRef, cuArray, channelDesc);   // Bind the array to the texture

    

     float* output;    // Allocate result of transformation in device memory

     cudaMalloc((void**)&output, width * height * sizeof(float));

   

     dim3 dimBlock(16, 16);    // Invoke kernel

     dim3 dimGrid((width + dimBlock.x –1) / dimBlock.x,(height + dimBlock.y –1) / dimBlock.y);

     transformKernel<<<dimGrid, dimBlock>>>(output, width, height,angle);

   

     cudaFreeArray(cuArray);    // Free device memory

     cudaFree(output);

}