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

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

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

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

书接上回《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 <stdio.h>

#include <stdlib.h>

#include <cuda_runtime.h>

#include <cutil.h>

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

/* 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. ");

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. ");

   return false;

}

cudaSetDevice(i);

printf("CUDA initialized. ");

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 ");

CUDA_SAFE_CALL( cudaThreadSynchronize() );

CUT_SAFE_CALL( cutStopTimer( timer));

printf("Processing time: %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);

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的使用中体现出来的性能优势做一个简单的分析。