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

分類： 編程語言2012-11-05 10:55 2521人閱讀 評論(0) 收藏 舉報cudaGPU    最近初試cuda編程，作為一個新手，遇到了各種各樣的問題，然后花費(fèi)了大量時間解決這些匪夷所思的問題。為了避免后來人重蹈覆轍，現(xiàn)把自己遇到的問題總結(jié)如下。

（一）、cudaMalloc

創(chuàng)新互聯(lián)公司-成都網(wǎng)站建設(shè)公司，專注成都做網(wǎng)站、網(wǎng)站建設(shè)、外貿(mào)營銷網(wǎng)站建設(shè)、網(wǎng)站營銷推廣，空間域名，網(wǎng)頁空間，網(wǎng)站運(yùn)營有關(guān)企業(yè)網(wǎng)站制作方案、改版、費(fèi)用等問題，請聯(lián)系創(chuàng)新互聯(lián)公司。

   初次使用該函數(shù)，感覺沒有什么困難，和c語言的malloc類似。但是在具體應(yīng)用中卻出了一個很難找的錯誤，花費(fèi)了很多時間。該函數(shù)使用是需要注意的就是，它分配的內(nèi)存空間單位是字節(jié)，所以需要我們在使用時用sizeof指定具體分配的變量類型，這樣才能正確分配空間。例：

              cudaMalloc((void**)&gpu_data,sizeof(float)*1024);

（二）、函數(shù)的執(zhí)行位置

  cuda程序的一大特色是程序的核心部分在GPU上執(zhí)行，所以cuda函數(shù)就分為不同的類別：host、global、device三類。所以我們在編寫函數(shù)時一定要分清楚當(dāng)前正在編寫的是哪類函數(shù)，可以調(diào)用什么庫函數(shù)。

• host函數(shù)：在CPU上調(diào)用，在CPU上執(zhí)行，可以調(diào)用global函數(shù)，不能調(diào)用device函數(shù)；
• global函數(shù)：只能在host函數(shù)中調(diào)用，但是執(zhí)行是在GPU上執(zhí)行，例如cudaMalloc之類的內(nèi)存操作庫函數(shù)，可以調(diào)用device函數(shù)；
• device函數(shù)：只能在GPU上調(diào)用和執(zhí)行，只能被global函數(shù)引用。

  關(guān)于函數(shù)類別容易出現(xiàn)的錯誤就是內(nèi)存分配時CPU和GPU的混淆。我們只需要記住，在host函數(shù)中可以直接使用的內(nèi)存都是CPU上的內(nèi)存，GPU上的內(nèi)存需要通過cudaMemcpy函數(shù)調(diào)用拷貝到CPU內(nèi)存空間；在global和device函數(shù)中使用的內(nèi)存都是在GPU內(nèi)存空間，使用之前需要分配。   （三）、共享內(nèi)存    共享內(nèi)存是提升程序性能很重要的一部分，能不能用好共享內(nèi)存是是否掌握cuda編程的一個重要依據(jù)。在此只想強(qiáng)調(diào)一點(diǎn)：共享內(nèi)存沒有初始化！下面是自己寫的一個數(shù)組求和程序，用到了共享內(nèi)存： [cpp] view plaincopyprint?  

1. __device__ int count=0;
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;
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.     }
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.     }
29.     if(tid==0)
30.     {
31.         block_gpu[bid]=blocksum[0];
32.        __threadfence();
34.         int value=atomicAdd(&count,1);
35.         islast=(value==gridDim.x-1);
36.     }
37.     __syncthreads();
39.     if(islast)
40.     {
41.         if(tid==0)
42.         {
43.             int s=0;
45.             for(int i=0;i<BLOCK_NUM;i++)
46.             {
47.                 s+=block_gpu[i];
48.             }
50.             *sum_gpu=s;
51.         }
52.     }
53. }

  特別注意第11八行代碼，不對要訪問的共享內(nèi)存進(jìn)行初始化將得不到正確的結(jié)果。   （四）、原子函數(shù)調(diào)用   在調(diào)用原子函數(shù)時，需要指定當(dāng)前顯卡的計(jì)算能力，否則會報錯“atomic*** is undefined.”。 linux下解決方案是在編譯源代碼時為nvcc編譯器指定一個計(jì)算能力的選項(xiàng)。例如計(jì)算能力時1.3，則可以添加參數(shù)：-arch sm_13，這樣就可以順利編譯。   （五）、CUDA語法   很多參考書都介紹說CUDA采用的是C擴(kuò)展語法，所以一開始我們很容易認(rèn)為采用C語法就夠了。但是這樣也容易讓我們陷入一個誤區(qū)：只能是C語法，而不能是其他。其實(shí)CUDA是C和C++的混合體，有時候采用C++的語法會更便利：

• for循環(huán)內(nèi)可以定義變量，標(biāo)準(zhǔn)C語言不支持，所以我們可以直接用(for int i=0;i<length;i++)，這樣的好處是可以節(jié)省一個寄存器；
• 變量定義位置無限制，可以在任意位置定義變量；
• CUDA支持多態(tài)，所以我們可以定義多個名稱相同，參數(shù)不同的函數(shù)，這個沒有問題；
• 有時多態(tài)可以用模版（template）來合并代碼，達(dá)到簡化編程的目的；

（六）、block和thread號的正確使用   為了調(diào)度不同的線程，我們通常需要利用內(nèi)置變量threadIdx和blockIdx作為循環(huán)中的增量。但是切記在循環(huán)內(nèi)部要正確使用內(nèi)置變量，兩天debug的教訓(xùn)！下面是一個示例代碼： [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];
6.     int tid=threadIdx.x;
7.     int bid=blockIdx.x;
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，結(jié)果花了兩天的時間找問題，所以要正確使用，在可以替換的情況下就用i或者j這樣的變量，盡量少用內(nèi)置變量。 （七）、空間釋放   在GPU上分配的空間，在使用完成之后要及時釋放。對于運(yùn)行一次的程序，不釋放空間沒有什么大礙，畢竟程序結(jié)束空間自動會被釋放掉。但是當(dāng)程序不間斷運(yùn)行多次的時候，不釋放空間會導(dǎo)致非常嚴(yán)重的GPU內(nèi)存泄露。第一個問題是隨著程序的運(yùn)行，GPU內(nèi)存耗盡，導(dǎo)致后續(xù)內(nèi)存分配失敗；第二個問題是，程序運(yùn)行會越來越慢。所以我們一定要養(yǎng)成用完及時釋放空間的習(xí)慣。

分享文章：CUDA編程常見問題轉(zhuǎn)-創(chuàng)新互聯(lián)
文章分享：http://www.rwnh.cn/article26/ehsjg.html

成都網(wǎng)站建設(shè)公司_創(chuàng)新互聯(lián)，為您提供外貿(mào)建站、網(wǎng)頁設(shè)計(jì)公司、虛擬主機(jī)、軟件開發(fā)、搜索引擎優(yōu)化、移動網(wǎng)站建設(shè)

廣告

聲明：本網(wǎng)站發(fā)布的內(nèi)容（圖片、視頻和文字）以用戶投稿、用戶轉(zhuǎn)載內(nèi)容為主，如果涉及侵權(quán)請盡快告知，我們將會在第一時間刪除。文章觀點(diǎn)不代表本網(wǎng)站立場，如需處理請聯(lián)系客服。電話：028-86922220；郵箱：631063699@qq.com。內(nèi)容未經(jīng)允許不得轉(zhuǎn)載，或轉(zhuǎn)載時需注明來源： 創(chuàng)新互聯(lián)