CUDA Texture紋理存儲器 示例程序

紋理存儲器非常適合實現圖像處理和查找表，對大量數據的隨機訪問或非隨機訪問也有良好的加速效果。第一次接觸紋理存儲器，寫了以下一個小程序。

相關閱讀：

Ubuntu 11.10 上安裝CUDA開發環境  http://www.linuxidc.com/Linux/2012-04/58913.htm

Ubuntu 11.04 安裝 NVIDIA CUDA 4.0 RC2 見 http://www.linuxidc.com/Linux/2011-10/46304.htm

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. }