CUDA中texture的使用

[size=medium]任务一：第一个texture程序
我有一点始终搞不明白，编程大牛们为什么总喜欢把简单的事情都说的那么含蓄，让读者看了总是心虚。我也一样，看了很多大牛们对CUDA中texture的讲解，我稀里糊涂的，心里没底，今天我终于发现了一个能让我马上能接受的讲解，所以很有必要做个笔录。
什么是texture，其实就是GPU内存，它只可读，但是在性能上具有很多优势（我现在也没搞明白在哪，一会再去查查）,怎么用呢？首先你的在cpu上定义一个texture的参考，也就是绑定到texture的代号（就像我们给陈水扁取名为“阿扁”一样）；然后就是绑定texture参考到纹理（gpu内存），利用一些纹理获取函数取得纹理中的数据；最后记得调用unbind函数删除texture。。就这么简单，看个例子：
main.cu
#include <stdio.h>
extern "C" void Blend_GPU( unsigned char* aImg1, unsigned char* aImg2, unsigned char* aImg3, int width, int height, int channel );//声明要调用的Blend_GPU函数

//这个Blend_CPU函数是在CPU上执行，完成和GPU上运行一样的功能，以作对比，功能很简单吧，就不介绍了
void Blend_CPU( unsigned char* aImg1, unsigned char* aImg2,
unsigned char* aRS,int width, int height, int channel )
{
for( int i = 0; i < width * height * channel; ++ i )
aRS[i] = (unsigned char)( 0.5 * aImg1[i] + 0.5 * aImg2[i] );
}
//main函数，入口，初始化数据
void main( int argc, char** argv )
{
int width = 1920,height = 1200, channel = 3;

// 分配4个空间，类型都是unsinged char型，都写在一行了
unsigned char *aImg1 = new unsigned char[ width*height*channel ],
*aImg2 = new unsigned char[ width*height*channel ],
*aRS1 = new unsigned char[ width*height*channel ],
*aRS2 = new unsigned char[ width*height*channel ];
//初始化数据，aImg1数组里都放0，aImg2里都放200
for( int i = 0; i < width * height * channel; ++ i )
{
aImg1[i] = 0;
aImg2[i] = 200;
}

// 调用CPU端程序
Blend_CPU( aImg1, aImg2, aRS1, width, height, channel );
// 调用Blend_GPU函数，Blend_GPU中会调用gpu端的kernel函数
Blend_GPU( aImg1, aImg2, aRS2, width, height, channel );
// 测试CPU端和GPU端执行的结果是不是一样，不是一样给出错误提示
for( int i = 0; i < width * height * channel; ++ i )
if( aRS1[i] != aRS2[i] )
{
printf( "Error!!!!\n" );
break;
}
printf("success!!");
}

blend_gpu.cu

#define BLOCK_DIM 512
//声明纹理参考，用来绑定纹理，其实也就是个纹理标识
texture<unsigned char, 1, cudaReadModeElementType> rT1;
texture<unsigned char, 1, cudaReadModeElementType> rT2;
//声明函数
extern "C" void Blend_GPU( unsigned char* aImg1, unsigned char* aImg2,unsigned char* aRS,int width, int height, int channel );

//核心代码，在gpu端执行的kernel，
__global__ void Blending_Texture( unsigned char* aRS, int size )
{
//通过线程ID得到数组下标
int index = blockIdx.x * blockDim.x + threadIdx.x;
//通过纹理获取函数得到数据再运算，把结果放到aRS数组中去
if( index < size )
aRS[index] = 0.5 * tex1Dfetch( rT1, index )+ 0.5 * tex1Dfetch( rT2, index );
}

void Blend_GPU( unsigned char* aImg1, unsigned char* aImg2,unsigned char* aRS,int width, int height, int channel )
{
int size = height * width * channel;
int data_size = size * sizeof( unsigned char );

//开辟3个空间
cudaMalloc( (void**)&dev_A, data_size );
cudaMalloc( (void**)&dev_B, data_size );
cudaMalloc( (void**)&dev_C, data_size );

//将host端的数据拷贝到device端
cudaMemcpy( dev_A, aImg1, data_size, cudaMemcpyHostToDevice );
cudaMemcpy( dev_B, aImg2, data_size, cudaMemcpyHostToDevice );

//将纹理参考绑定到device端的两数组 ☆device就是pgu
cudaBindTexture(0, rT1, dev_A );
cudaBindTexture(0, rT2, dev_B );

//调用kernel
Blending_Texture<<< ceil( (float)size / BLOCK_DIM ), BLOCK_DIM >>>( dev_C, size );

//将结果拷贝到host端 ☆host就是CPU
cudaMemcpy( aRS, dev_C, data_size, cudaMemcpyDeviceToHost );

//取消绑定
cudaUnbindTexture(rT1);
cudaUnbindTexture(rT2);

//释放内存空间
cudaFree(dev_A);
cudaFree(dev_B);
cudaFree(dev_C);
}

任务二：关于运行程序时显示无法打开“包括文件cutil_inline.h”的问题解答
广州一CUDA爱好者朋友问我这个问题，我刚好前两天也遇到过类似问题，还是记录下吧，其实这个问题就是因为环境变量没配置好的原因，cutil_inline.h头文件是在sdk中common下的inc目录，有两种方法解决：
方法一：首先在环境变量中配置CUDA_COMMON_INC_PATH变量，值为sdk下的common下的inc目录，然后右键项目，选择CUDAGeneralAdditional Include Directories中添加;$(CUDA_COMMON_INC_PATH)，再执行程序就ok了，
测试程序如下：
方法二：在环境变量CUDA_INC_PATH中添加sdk中common下的inc目录，其原理都一样，都是让程序能够找到那个头文件
测试程序：
main.cu
/********************************************************************
* main.cu
* This is a example of the CUDA test program.
*********************************************************************/
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
//为么不能包含这个头文件？
#include <cutil_inline.h>
bool InitCUDA(void)
{
int count = 0;
int i = 0;
cudaGetDeviceCount(&count);
printf("devivce count : %d\n",count);
if ( count == 0 ) {
fprintf( stderr , "There is no device.\n");
return false;
}
for ( i = 0 ; i < count ; i++) {
cudaDeviceProp prop;
if ( cudaGetDeviceProperties(&prop,i) == cudaSuccess ) {
printf("Name : %s\n", prop.name);
printf("totalGlobalMem : %u MB \n" , prop.totalGlobalMem / (1024 * 1024));
printf("sharedMemPerBlock : %u KB \n" , prop.sharedMemPerBlock / 1024 );
printf("regsPerBlock:%d \n", prop.regsPerBlock);
printf("warpSize : %d \n" , prop.warpSize);
printf("memPitch : %u \n", prop.memPitch);
printf("maxThreadPerBlock %d \n" , prop.maxThreadsPerBlock ) ;
printf("maxThreadsDim:x %d, y %d, z %d\n",prop.maxThreadsDim[0],prop.maxThreadsDim[1] , prop.maxThreadsDim[2]);
printf("maxGridSize:x %d, y %d, z%d\n", prop.maxGridSize[0],prop.maxGridSize[0] , prop.maxGridSize[1], prop.maxGridSize[2]);
printf("totalConstMem:%u\n" , prop.totalConstMem);
printf("major:%d\n",prop.major);
printf("minor:%d\n",prop.minor);
printf("clockRate:%d\n",prop.clockRate);
printf("textureAlignment:%u\n",prop.textureAlignment);
if ( prop.major >= 1 ) {
break;
}
}
}
if ( i == count) {
fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
return false;
}
cudaSetDevice(i);
printf("CUDA initialized. \n");
return true;
}
int main(int argc , char ** argv)
{
InitCUDA();
//cutilDeviceInit(argc, argv);
return 0;
}[/size]