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);  // 退出CUDA114 }