CUDA Texture纹理存储器 示例程序

  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 }