CUDA Texture纹理存储器 示例程序


纹理存储器非常适合实现图像处理和查找表，对大量数据的随机访问或非随机访问也有良好的加速效果。第一次接触纹理存储器，写了以下一个小程序。

相关阅读：

Ubuntu 11.10 上安装CUDA开发环境 

Ubuntu 11.04 安装 NVIDIA CUDA 4.0 RC2 见

  1. /* 
  2. * Copyright 徐洪志(西北农林科技大学.信息工程学院).  All rights reserved. 
  3. * Data: 2012-4-20 
  4. */  
  5. //   
  6. // 此程序是演示了1D和2D纹理存储器的使用   
  7. #include <stdio.h>   
  8. #include <cutil_inline.h>   
  9. #include <iostream>   
  10. using namespace std;  
  11.   
  12. texture<float> texRef1D;     // 1D texture   
  13. texture<float, 2> texRef2D;  // 2D texture   
  14.   
  15. // 1D 纹理操作函数   
  16. __global__ void Texture1D(float *dst, int w, int h)  
  17. {  
  18.     int x = threadIdx.x + blockIdx.x * blockDim.x;  
  19.     int y = threadIdx.y + blockIdx.y * blockDim.y;  
  20.     int offset = x + y * blockDim.x * gridDim.x;  
  21.   
  22.     if(x<w && y<h)  
  23.         dst[offset] = tex1Dfetch(texRef1D, offset);  
  24. }  
  25. // 2D 纹理操作函数   
  26. __global__ void Texture2D(float *dst, int w, int h)  
  27. {  
  28.     int x = threadIdx.x + blockIdx.x * blockDim.x;  
  29.     int y = threadIdx.y + blockIdx.y * blockDim.y;  
  30.     int offset = x + y * blockDim.x * gridDim.x;  
  31.   
  32.     dst[offset] = tex2D(texRef2D, x, y);  
  33. }  
  34. int main(int argc, char **argv)  
  35. {  
  36.     CUT_DEVICE_INIT(argc, argv);  // 启动 CUDA   
  37.     /// 1D 纹理内存   
  38.     cout << "1D texture" << endl;  
  39.     float *host1D = (float*)calloc(10, sizeof(float)); // 内存原数据   
  40.     float *hostRet1D = (float*)calloc(10, sizeof(float));// 内存保存返回数据   
  41.   
  42.     float *dev1D, *devRet1D; // 显存数据   
  43.     int i;  
  44.     cout << " host1D:" << endl;  
  45.     for(i = 0; i < 10; ++i)  // 初始化内存原数据   
  46.     {  
  47.         host1D[i] = i * 2;  
  48.         cout << "  " << host1D[i] << " ";  
  49.     }  
  50.     cutilSafeCall( cudaMalloc((void**)&dev1D, sizeof(float)*10));  // 申请显存空间   
  51.     cutilSafeCall( cudaMalloc((void**)&devRet1D, sizeof(float)*10));  
  52.     cutilSafeCall( cudaMemcpy(dev1D, host1D, sizeof(float)*10, cudaMemcpyHostToDevice)); // 将内存数据拷贝入显存   
  53.     cutilSafeCall( cudaBindTexture(NULL, texRef1D, dev1D, sizeof(float)*10));  // 将显存数据和纹理绑定   
  54.       
  55.     Texture1D<<<10, 1>>>(devRet1D, 10, 1);  // 运行1D纹理操作函数   
  56.   
  57.     cutilSafeCall( cudaMemcpy(hostRet1D, devRet1D, sizeof(float)*10, cudaMemcpyDeviceToHost)); // 将显存数据拷贝入内存   
  58.     // 打印内存数据   
  59.     cout << endl << " hostRet1D:" << endl;  
  60.     for(i = 0; i < 10; ++i)  
  61.         cout << "  " << hostRet1D[i] << " ";  
  62.   
  63.     cutilSafeCall( cudaUnbindTexture(texRef1D));  // 解绑定   
  64.     cutilSafeCall( cudaFree(dev1D));  // 释放显存空间   
  65.     cutilSafeCall( cudaFree(devRet1D));   
  66.     free(host1D);  // 释放内存空间   
  67.     free(hostRet1D);  
  68.       
  69.     /// 2D 纹理内存    
  70.     cout << endl << "2D texture" << endl;  
  71.     int width = 5, height = 3;  
  72.     float *host2D = (float*)calloc(width*height, sizeof(float));    // 内存原数据   
  73.     float *hostRet2D = (float*)calloc(width*height, sizeof(float)); // 内存返回数据    
  74.   
  75.     cudaArray *cuArray;  // CUDA数组   
  76.     float *devRet2D;     // 显存数据   
  77.     int row, col;  
  78.     cout << " host2D:" << endl;  
  79.     for(row = 0; row < height; ++row)  // 初始化内存原数据   
  80.     {  
  81.         for(col = 0; col < width; ++col)  
  82.         {  
  83.             host2D[row*width + col] = row + col;  
  84.             cout << "  " << host2D[row*width + col] << " ";  
  85.         }  
  86.         cout << endl;  
  87.     }  
  88.     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();  
  89.     cutilSafeCall( cudaMallocArray(&cuArray, &channelDesc, width, height));  // 申请显存空间   
  90.     cutilSafeCall( cudaMalloc((void**) &devRet2D, sizeof(float)*width*height));  
  91.     cutilSafeCall( cudaBindTextureToArray(texRef2D, cuArray)); // 将显存数据和纹理绑定   
  92.     cutilSafeCall( cudaMemcpyToArray(cuArray, 0, 0, host2D, sizeof(float)*width*height, cudaMemcpyHostToDevice)); // 将内存数据拷贝入CUDA数组   
  93.   
  94.     dim3 threads(width, height);  
  95.     Texture2D<<<1, threads>>>(devRet2D, width, height);  // 运行2D纹理操作函数   
  96.   
  97.     cutilSafeCall( cudaMemcpy(hostRet2D, devRet2D, sizeof(float)*width*height, cudaMemcpyDeviceToHost)); // 将显存数据拷贝入内存   
  98.     // 打印内存数据   
  99.     cout << " hostRet2D:" << endl;  
  100.     for(row = 0; row < height; ++row)  
  101.     {  
  102.         for(col = 0; col < width; ++col)  
  103.             cout << "  " << hostRet2D[row*width + col] << " ";  
  104.         cout << endl;  
  105.     }  
  106.   
  107.     cutilSafeCall( cudaUnbindTexture(texRef2D)); // 解绑定   
  108.     cutilSafeCall( cudaFreeArray(cuArray));  // 释放显存空间   
  109.     cutilSafeCall( cudaFree(devRet2D));  
  110.     free(host2D);  // 释放内存空间   
  111.     free(hostRet2D);  
  112.   
  113.     CUT_EXIT(argc, argv);  // 退出CUDA   
  114. }  

相关内容