CUDA编程常见问题 转

http://blog.csdn.net/yutianzuijin/article/details/8147912

分类： 编程语言2012-11-05 10:55 2521人阅读 评论(0) 收藏 举报
cudaGPU

     最近初试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函数引用。
    关于函数类别容易出现的错误就是内存分配时CPU和GPU的混淆。我们只需要记住，在host函数中可以直接使用的内存都是CPU上的内存，GPU上的内存需要通过cudaMemcpy函数调用拷贝到CPU内存空间；在global和device函数中使用的内存都是在GPU内存空间，使用之前需要分配。
 
（三）、共享内存
     共享内存是提升程序性能很重要的一部分，能不能用好共享内存是是否掌握cuda编程的一个重要依据。在此只想强调一点：共享内存没有初始化！下面是自己写的一个数组求和程序，用到了共享内存：
[cpp] view plaincopyprint?
 
  1. __device__ int count=0;  
  2.   
  3. __global__ static void sum(int* data_gpu,int* block_gpu,int *sum_gpu,int length)  
  4. {  
  5.     extern __shared__ int blocksum[];  
  6.     __shared__ int islast;  
  7.     int offset;  
  8.   
  9.     const int tid=threadIdx.x;  
  10.     const int bid=blockIdx.x;  
  11.     blocksum[tid]=0;  
  12.     for(int i=bid*THREAD_NUM+tid;i<length;i+=BLOCK_NUM*THREAD_NUM)  
  13.     {  
  14.         blocksum[tid]+=data_gpu[i];  
  15.     }   
  16.   
  17.     __syncthreads();  
  18.     offset=THREAD_NUM/2;  
  19.     while(offset>0)  
  20.     {  
  21.         if(tid<offset)  
  22.         {  
  23.             blocksum[tid]+=blocksum[tid+offset];  
  24.         }  
  25.         offset>>=1;  
  26.         __syncthreads();  
  27.     }  
  28.   
  29.     if(tid==0)  
  30.     {  
  31.         block_gpu[bid]=blocksum[0];  
  32.        __threadfence();  
  33.   
  34.         int value=atomicAdd(&count,1);  
  35.         islast=(value==gridDim.x-1);  
  36.     }  
  37.     __syncthreads();  
  38.   
  39.     if(islast)  
  40.     {  
  41.         if(tid==0)  
  42.         {  
  43.             int s=0;  
  44.   
  45.             for(int i=0;i<BLOCK_NUM;i++)  
  46.             {  
  47.                 s+=block_gpu[i];  
  48.             }  
  49.   
  50.             *sum_gpu=s;  
  51.         }  
  52.     }  
  53. }  
   特别注意第11八行代码，不对要访问的共享内存进行初始化将得不到正确的结果。
 
（四）、原子函数调用
   在调用原子函数时，需要指定当前显卡的计算能力，否则会报错“atomic*** is undefined.”。 linux下解决方案是在编译源代码时为nvcc编译器指定一个计算能力的选项。例如计算能力时1.3，则可以添加参数：-arch sm_13，这样就可以顺利编译。
 
（五）、CUDA语法
   很多参考书都介绍说CUDA采用的是C扩展语法，所以一开始我们很容易认为采用C语法就够了。但是这样也容易让我们陷入一个误区：只能是C语法，而不能是其他。其实CUDA是C和C++的混合体，有时候采用C++的语法会更便利：
  • for循环内可以定义变量，标准C语言不支持，所以我们可以直接用(for int i=0;i<length;i++)，这样的好处是可以节省一个寄存器；
  • 变量定义位置无限制，可以在任意位置定义变量；
  • CUDA支持多态，所以我们可以定义多个名称相同，参数不同的函数，这个没有问题；
  • 有时多态可以用模版（template）来合并代码，达到简化编程的目的；

（六）、block和thread号的正确使用
    为了调度不同的线程，我们通常需要利用内置变量threadIdx和blockIdx作为循环中的增量。但是切记在循环内部要正确使用内置变量，两天debug的教训！下面是一个示例代码：
[cpp] view plaincopyprint?
 
  1. __global__ static void saliencefunc(float *peaks_gpu,int *index_gpu,float *saliencebins_gpu,int framenumber)  
  2. {  
  3.     __shared__ float peaks[HALF_PEAK_NUM];  
  4.     __shared__ int index[HALF_PEAK_NUM];  
  5.   
  6.     int tid=threadIdx.x;  
  7.     int bid=blockIdx.x;  
  8.   
  9.     for(int i=bid;i<framenumber;i+=BLOCK_NUM)  
  10.     {  
  11.         if(tid<HALF_PEAK_NUM)  
  12.         {  
  13.             peaks[tid]=peaks_gpu[HALF_PEAK_NUM*i+tid];  
  14.             index[tid]=index_gpu[HALF_PEAK_NUM*i+tid];  
  15.         }  
  16.         __syncthreads();      
  17.     }  
  18. }  

注意代码第十三和十四行的赋值操作HALF_PEAK_NUM*i+tid，笔者之前的写法是HALF_PEAK_NUM*bid+tid，结果花了两天的时间找问题，所以要正确使用，在可以替换的情况下就用i或者j这样的变量，尽量少用内置变量。
（七）、空间释放
    在GPU上分配的空间，在使用完成之后要及时释放。对于运行一次的程序，不释放空间没有什么大碍，毕竟程序结束空间自动会被释放掉。但是当程序不间断运行多次的时候，不释放空间会导致非常严重的GPU内存泄露。第一个问题是随着程序的运行，GPU内存耗尽，导致后续内存分配失败；第二个问题是，程序运行会越来越慢。所以我们一定要养成用完及时释放空间的习惯。
【推广】 免费学中医，健康全家人
原文地址：https://www.cnblogs.com/huashiyiqike/p/3868358.html