CUDA编程常见问题转-创新互联

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

分类: 编程语言2012-11-05 10:55 2521人阅读 评论(0) 收藏 举报cudaGPU    最近初试cuda编程,作为一个新手,遇到了各种各样的问题,然后花费了大量时间解决这些匪夷所思的问题。为了避免后来人重蹈覆辙,现把自己遇到的问题总结如下。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. __global__ static void sum(int* data_gpu,int* block_gpu,int *sum_gpu,int length)
  3. {
  4.     extern __shared__ int blocksum[];
  5.     __shared__ int islast;
  6.     int offset;
  7.     const int tid=threadIdx.x;
  8.     const int bid=blockIdx.x;
  9.     blocksum[tid]=0;
  10.     for(int i=bid*THREAD_NUM+tid;i<length;i+=BLOCK_NUM*THREAD_NUM)
  11.     {
  12.         blocksum[tid]+=data_gpu[i];
  13.     }
  14.     __syncthreads();
  15.     offset=THREAD_NUM/2;
  16.     while(offset>0)
  17.     {
  18.         if(tid<offset)
  19.         {
  20.             blocksum[tid]+=blocksum[tid+offset];
  21.         }
  22.         offset>>=1;
  23.         __syncthreads();
  24.     }
  25.     if(tid==0)
  26.     {
  27.         block_gpu[bid]=blocksum[0];
  28.        __threadfence();
  29.         int value=atomicAdd(&count,1);
  30.         islast=(value==gridDim.x-1);
  31.     }
  32.     __syncthreads();
  33.     if(islast)
  34.     {
  35.         if(tid==0)
  36.         {
  37.             int s=0;
  38.             for(int i=0;i<BLOCK_NUM;i++)
  39.             {
  40.                 s+=block_gpu[i];
  41.             }
  42.             *sum_gpu=s;
  43.         }
  44.     }
  45. }
  特别注意第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.     int tid=threadIdx.x;
  6.     int bid=blockIdx.x;
  7.     for(int i=bid;i<framenumber;i+=BLOCK_NUM)
  8.     {
  9.         if(tid<HALF_PEAK_NUM)
  10.         {
  11.             peaks[tid]=peaks_gpu[HALF_PEAK_NUM*i+tid];
  12.             index[tid]=index_gpu[HALF_PEAK_NUM*i+tid];
  13.         }
  14.         __syncthreads();
  15.     }
  16. }

注意代码第十三和十四行的赋值操作HALF_PEAK_NUM*i+tid,笔者之前的写法是HALF_PEAK_NUM*bid+tid,结果花了两天的时间找问题,所以要正确使用,在可以替换的情况下就用i或者j这样的变量,尽量少用内置变量。 (七)、空间释放   在GPU上分配的空间,在使用完成之后要及时释放。对于运行一次的程序,不释放空间没有什么大碍,毕竟程序结束空间自动会被释放掉。但是当程序不间断运行多次的时候,不释放空间会导致非常严重的GPU内存泄露。第一个问题是随着程序的运行,GPU内存耗尽,导致后续内存分配失败;第二个问题是,程序运行会越来越慢。所以我们一定要养成用完及时释放空间的习惯。

分享文章:CUDA编程常见问题转-创新互联
文章分享:https://www.cdcxhl.com/article26/ehsjg.html

成都网站建设公司_创新互联,为您提供外贸建站网页设计公司虚拟主机软件开发搜索引擎优化移动网站建设

广告

声明:本网站发布的内容(图片、视频和文字)以用户投稿、用户转载内容为主,如果涉及侵权请尽快告知,我们将会在第一时间删除。文章观点不代表本网站立场,如需处理请联系客服。电话:028-86922220;邮箱:631063699@qq.com。内容未经允许不得转载,或转载时需注明来源: 创新互联

成都网站建设公司