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