(CUDA 编程10).CUDA cosnstant使用（一）------GPU的革命

(CUDA编程10).CUDAcosnstant使用（一）------GPU的革命

作者：赵开勇来源：http://www.hpctech.com/2009/0818/209.html

10.CUDAcosnstant使用（一）------GPU的革命

序言：最近的事情无比的多，差点就找不到回家的路了，都快忘记出发的起点的时候，冷静下来，侧夜未眠，事情再多，都要一件一件的做好，做不好的，就不接，我想别人也可
10.CUDAcosnstant使用（一）------GPU的革命

书接上回《9.CUDAsharedmem使用》讲了sharedmemory的使用，最近有几个朋友都在问我cosntant的使用的问题，这次首先先讲一下cosntant的使用，下一章节才讲一下cosntant的使用中性能的体现；

下面是一个简单的代码:

/********************************************************************

*cosntant_test.cu

*ThisisaexampleoftheCUDAprogram.

*author:zhao.kaiyong(at)gmail.com

*********************************************************************/

#include<stdio.h>

#include<stdlib.h>

#include<cuda_runtime.h>

#include<cutil.h>

/************************************************************************/

/*InitCUDA*/

/************************************************************************/

#if__DEVICE_EMULATION__

boolInitCUDA(void){returntrue;}

#else

boolInitCUDA(void)

{

intcount=0;

inti=0;

cudaGetDeviceCount(&count);

if(count==0){

fprintf(stderr,"Thereisnodevice.");

returnfalse;

}

for(i=0;i<count;i++){

cudaDevicePropprop;

if(cudaGetDeviceProperties(&prop,i)==cudaSuccess){

if(prop.major>=1){

break;

}

}

}

if(i==count){

fprintf(stderr,"ThereisnodevicesupportingCUDA.");

returnfalse;

}

cudaSetDevice(i);

printf("CUDAinitialized.");

returntrue;

}

#endif

/************************************************************************/

/*Example*/

/************************************************************************/

__constant__charp_HelloCUDA[11];//="HelloCUDA!";

__constant__intt_HelloCUDA[11]={0,1,2,3,4,5,6,7,8,9,10};

__constant__intnum=11;

__global__staticvoidHelloCUDA(char*result)

{

inti=0;

for(i=0;i<num;i++){

result[i]=p_HelloCUDA[i]+t_HelloCUDA[i];

}

}

/************************************************************************/

/*HelloCUDA*/

/************************************************************************/

intmain(intargc,char*argv[])

{

if(!InitCUDA()){

return0;

}

charhelloCUDA[]="HdjikCUDA!";

char*device_result=0;

charhost_result[12]={0};

CUDA_SAFE_CALL(cudaMalloc((void**)&device_result,sizeof(char)*11));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(p_HelloCUDA,helloCUDA,sizeof(char)*11));

unsignedinttimer=0;

CUT_SAFE_CALL(cutCreateTimer(&timer));

CUT_SAFE_CALL(cutStartTimer(timer));

HelloCUDA<<<1,1,0>>>(device_result);

CUT_CHECK_ERROR("Kernelexecutionfailed");

CUDA_SAFE_CALL(cudaThreadSynchronize());

CUT_SAFE_CALL(cutStopTimer(timer));

printf("Processingtime:%f(ms)",cutGetTimerValue(timer));

CUT_SAFE_CALL(cutDeleteTimer(timer));

CUDA_SAFE_CALL(cudaMemcpy(&host_result,device_result,sizeof(char)*11,cudaMemcpyDeviceToHost));

printf("%s",host_result);

CUDA_SAFE_CALL(cudaFree(device_result));

CUT_EXIT(argc,argv);

return0;

}

这里写了两种使用cosntant的方法：

1.一种方法是直接在定义的时候初始化constant：

__constant__intt_HelloCUDA[11]={0,1,2,3,4,5,6,7,8,9,10};

__constant__intnum=11;

在kernel里面直接使用就可以了；

2.第二种方法是定义一个cosntant数组，然后使用函数初始化它；

__constant__charp_HelloCUDA[11];//="HelloCUDA!";

CUDA_SAFE_CALL(cudaMemcpyToSymbol(p_HelloCUDA,helloCUDA,sizeof(char)*11));

前面一句是定义constant，后面一句是初始化这个常量，在kernel里面使用的时候，就按照定义的方式使用就可以了；

当然也有朋友想定义自己的结构体，当然是可以的，只是在初始化的时候，copy相应的结构体就可以了。这个只是一个初步的使用方法，希望对家有用。

下一次会对cosntant的使用中体现出来的性能优势做一个简单的分析。