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;
12. texture<float> texRef1D; // 1D texture
13. texture<float, 2> texRef2D; // 2D texture
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;
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;
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));// 內存保存返回數據
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)); // 將顯存數據和紋理綁定
55. Texture1D<<<10, 1>>>(devRet1D, 10, 1); // 運行1D紋理操作函數
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] << " ";
63. cutilSafeCall( cudaUnbindTexture(texRef1D)); // 解綁定
64. cutilSafeCall( cudaFree(dev1D)); // 釋放顯存空間
65. cutilSafeCall( cudaFree(devRet1D));
66. free(host1D); // 釋放內存空間
67. free(hostRet1D);
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)); // 內存返回數據
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數組
94. dim3 threads(width, height);
95. Texture2D<<<1, threads>>>(devRet2D, width, height); // 運行2D紋理操作函數
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. }
107. cutilSafeCall( cudaUnbindTexture(texRef2D)); // 解綁定
108. cutilSafeCall( cudaFreeArray(cuArray)); // 釋放顯存空間
109. cutilSafeCall( cudaFree(devRet2D));
110. free(host2D); // 釋放內存空間
111. free(hostRet2D);
113. CUT_EXIT(argc, argv); // 退出CUDA
114. }