CUDA cosnstant使用（一）

书接上回《9.CUDA shared mem使用》讲了shared memory的使用，最近有几个朋友都在问我cosntant的使用的问题，这次首先先讲一下cosntant的使用，下一章节才讲一下cosntant的使用中性能的体现；
下面是一个简单的代码:
/********************************************************************

* cosntant_test.cu

* This is a example of the CUDA program.

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

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

#include

#include

#include

#include

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

/* Init CUDA */

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

#if __DEVICE_EMULATION__

bool InitCUDA(void){return true;}

#else

bool InitCUDA(void)

{

int count = 0;

int i = 0;

cudaGetDeviceCount(&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) {

if(prop.major >= 1) {

break;

}

}

}

if(i == count) {

fprintf(stderr, "There is no device supporting CUDA.\n");

return false;

}

cudaSetDevice(i);

printf("CUDA initialized.\n");

return true;

}

#endif

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

/* Example */

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

__constant__ char p_HelloCUDA[11];// = "Hello CUDA!";

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

__constant__ int num = 11;

__global__ static void HelloCUDA(char* result)

{

int i = 0;

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

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

}

}

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

/* HelloCUDA */

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

int main(int argc, char* argv[])

{

if(!InitCUDA()) {

return 0;

}

char helloCUDA[] = "Hdjik CUDA!";

char *device_result = 0;

char host_result[12] ={0};

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

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

unsigned int timer = 0;

CUT_SAFE_CALL( cutCreateTimer( &timer));

CUT_SAFE_CALL( cutStartTimer( timer));

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

CUT_CHECK_ERROR("Kernel execution failed\n");

CUDA_SAFE_CALL( cudaThreadSynchronize() );

CUT_SAFE_CALL( cutStopTimer( timer));

printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));

CUT_SAFE_CALL( cutDeleteTimer( timer));

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

printf("%s\n", host_result);

CUDA_SAFE_CALL( cudaFree(device_result));

CUT_EXIT(argc, argv);

return 0;

}

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

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

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

__constant__ int num = 11;

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

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

__constant__ char p_HelloCUDA[11];// = "Hello CUDA!";

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

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

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

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

文章来源：http://blog.csdn.net/OpenHero/archive/2009/07/15/4348314.aspx

来自 “ ITPUB博客 ” ，链接：http://blog.itpub.net/22785983/viewspace-619578/，如需转载，请注明出处，否则将追究法律责任。

转载于:http://blog.itpub.net/22785983/viewspace-619578/