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