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 }