最近初试cuda编程,作为一个新手,遇到了各种各样的问题,然后花费了大量时间解决这些匪夷所思的问题。为了避免后来人重蹈覆辙,现把自己遇到的问题总结如下。
(一) cudaMalloc
初次使用该函数,感觉没有什么困难,和c语言的malloc类似。但是在具体应用中却出了一个很难找的错误,花费了很多时间。该函数使用是需要注意的就是,它分配的内存空间单位是字节,所以需要我们在使用时用sizeof指定具体分配的变量类型,这样才能正确分配空间。例:
cudaMalloc((void**)&gpu_data,sizeof(float)*1024);
(二) 函数的执行位置
cuda程序的一大特色是程序的核心部分在GPU上执行,所以cuda函数就分为不同的类别:host、global、device三类。所以我们在编写函数时一定要分清楚当前正在编写的是哪类函数,可以调用什么库函数。
- host函数:在CPU上调用,在CPU上执行,可以调用global函数,不能调用device函数;
- global函数:只能在host函数中调用,但是执行是在GPU上执行,例如cudaMalloc之类的内存操作库函数,可以调用device函数;
- device函数:只能在GPU上调用和执行,只能被global函数引用。
__device__ int count=0; __global__ static void sum(int* data_gpu,int* block_gpu,int *sum_gpu,int length) { extern __shared__ int blocksum[]; __shared__ int islast; int offset; const int tid=threadIdx.x; const int bid=blockIdx.x; blocksum[tid]=0; //第11行,对要访问的共享内存进行初始化 for(int i=bid*THREAD_NUM+tid;i<length;i+=BLOCK_NUM*THREAD_NUM) { blocksum[tid]+=data_gpu[i]; } __syncthreads(); offset=THREAD_NUM/2; while(offset>0) { if(tid<offset) { blocksum[tid]+=blocksum[tid+offset]; } offset>>=1; __syncthreads(); } if(tid==0) { block_gpu[bid]=blocksum[0]; __threadfence(); int value=atomicAdd(&count,1); islast=(value==gridDim.x-1); } __syncthreads(); if(islast) { if(tid==0) { int s=0; for(int i=0;i<BLOCK_NUM;i++) { s+=block_gpu[i]; } *sum_gpu=s; } } }
- for循环内可以定义变量,标准C语言不支持,所以我们可以直接用(for int i=0;i<length;i++),这样的好处是可以节省一个寄存器;
- 变量定义位置无限制,可以在任意位置定义变量;
- CUDA支持多态,所以我们可以定义多个名称相同,参数不同的函数,这个没有问题;
- 有时多态可以用模版(template)来合并代码,达到简化编程的目的;
(六) block和thread号的正确使用
__global__ static
void saliencefunc(float *peaks_gpu,int *index_gpu,float *saliencebins_gpu,int framenumber) { __shared__ float peaks[HALF_PEAK_NUM]; __shared__ int index[HALF_PEAK_NUM]; int tid=threadIdx.x; int bid=blockIdx.x; for(int i=bid;i<framenumber;i+=BLOCK_NUM) { if(tid<HALF_PEAK_NUM) { peaks[tid]=peaks_gpu[HALF_PEAK_NUM*i+tid]; //第13行赋值操作 index[tid]=index_gpu[HALF_PEAK_NUM*i+tid]; } __syncthreads(); } }