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 }
