[分享] CUDA 程式設計(10) -- 速成篇(上) - PPT 短網址

(2) 兩者皆有【中央處理器】，主機上為CPU，裝置上為GPU，指令集不同： 主機上的程式碼使用傳統C/C++ 語法撰寫成，實作與呼叫和一般函式無異， 裝置上的程式碼稱 ... ------文章開始------ 作者:a5000ml(咖啡裡的海洋藍)看板:VideoCard 標題:[分享]CUDA程式設計(10)--速成篇(上) 時間:WedNov1222:53:252008 (1)有學弟反應CUDA內容有點繁雜,很多概念容易搞混,而且希望多點範例, 所以這兩個禮拜把之前的文稿整理成【新手速成篇】,希望對他們有所幫助. (2)順便幫國網打個廣告:CUDA中文教學DVD(免費線上版)出現了 請至國網的教育訓練網登入https://edu.nchc.org.tw 詳情請看編號18026一文 ※第十章新手速成篇(上) ============================================================================ 前言 ============================================================================ 因為CUDA的一些延伸語法太繁雜，容易讓人混淆(例如記憶體種類就有4~5種， 同樣的globalmemory又有兩種寫法)，所以針對這個問題，寫成了速成篇， 去除那些枝枝節節，只講最重要的，並佐以範例，務求讓初學者【七招闖天下】。

第一招主機、裝置 第二招使用API(配置裝置記憶體&主機和裝置間資料搬移) 第三招函式&呼叫(主機、裝置) 第四招網格、區塊、執行緒(線程群組) 第五招記憶體(主機、裝置、共享) 第六招執行緒同步(網格、區塊) 第七招合併讀取(最佳化) 函式部份，只介紹__global__標籤，記憶體部份，只介紹__shared__標籤， 配置顯示記憶體以及資料搬移的方式，也只使用一種，簡單來說，這份速成篇 並不是完整的CUDA，只是刪減後的正交子集合，用來突顯主要概念，以及避免 初學者常犯的錯誤，熟悉之後，務必再深入了解其它延伸語法。

============================================================================ 第一招主機、裝置 ============================================================================ (1)區分主機和裝置的不同： 【主機】就是PC。

【裝置】就是顯示卡。

(2)兩者皆有【中央處理器】，主機上為CPU，裝置上為GPU，指令集不同： 主機上的程式碼使用傳統C/C++語法撰寫成，實作與呼叫和一般函式無異， 裝置上的程式碼稱為【核心】(kernel)，需使用CUDA的延伸語法(函式前加 __global__等標籤)來撰寫，並於呼叫時指定執行緒群組大小(詳見第三招) (3)兩者皆有【各自的記憶體】(DRAM)，擁有獨立的定址空間： 主機上的透過malloc()、free()、new、delete等函式配置與釋放， 裝置上的透過cudaMalloc()、cudaFree()等API配置與釋放， 主機和裝置之間的資料搬移，使用cudaMemcpy()這個API(詳見第二招) (4)因為主機和裝置的不同，C/C++的標準函式庫不能在kernel中直接使用， 例如要秀出計算結果，必需使用cudaMemcpy()先將資料搬移至主機， 再呼叫printf或cout等標準輸出函式。

(5)使用時先在主機記憶體設好資料的初始值，然後傳入裝置記憶體，接著執行核心， 如果可以的話就儘量讓資料保留在裝置中，進行一連串的kernel操作， 避免透過PCI-E搬移造成效能下降，最後再將結果傳回主機中顯示。

============================================================================ 第二招使用API(配置裝置記憶體&主機和裝置間資料搬移) ============================================================================ 最基本的API有5個 (1)配置裝置記憶體cudaMalloc()[cuda.h] (2)釋放裝置記憶體cudaFree()[cuda.h] (3)記憶體複製cudaMemcpy()[cuda.h] (4)錯誤字串解譯cudaGetErrorString()[cuda.h] (5)同步化cudaThreadSynchronize()[cuda.h] 用法如下 -------------------------------------------------------- (1)配置顯示記憶體cudaMalloc()[cuda.h] -------------------------------------------------------- cudaError_tcudaMalloc(void**ptr,size_tcount); ptr指向目的指位器之位址 count欲配置的大小(單位bytes) 傳回值cudaError_t是個enum,執行成功時傳回0,其它的錯誤代號可用 cudaGetErrorString()來解譯. -------------------------------------------------------- (2)釋放顯示記憶體cudaFree()[cuda.h] -------------------------------------------------------- cudaError_tcudaFree(void*ptr); ptr指向欲釋放的位址(devicememory) -------------------------------------------------------- (3)記憶體複製cudaMemcpy()[cuda.h] -------------------------------------------------------- cudaError_tcudaMemcpy(void*dst,constvoid*src,size_tcount, enumcudaMemcpyKindkind); dst指向目的位址 src指向來源位址 count拷貝區塊大小(單位bytes) kind有四種拷貝流向 cudaMemcpyHostToHost主機->主機 cudaMemcpyHostToDevice主機->裝置 cudaMemcpyDeviceToHost裝置->主機 cudaMemcpyDeviceToDevice裝置->裝置 -------------------------------------------------------- (4)錯誤字串解譯cudaGetErrorString()[cuda.h] -------------------------------------------------------- constchar*cudaGetErrorString(cudaError_terror); 傳回錯誤代號(error)所代表的字串 -------------------------------------------------------- (5)同步化cudaThreadSynchronize()[cuda.h] -------------------------------------------------------- cudaError_tcudaThreadSynchronize(void); 使前後兩個核心時序上分離,確保資料的前後相依性正確 //------------------------------------------------------------------------- //範例(1):透過裝置記憶體進行複製[081112-api.cu] //PCI-EPCI-E //主機記憶體a[]-------->裝置記憶體g[]-------->主機記憶體b[] //------------------------------------------------------------------------- #include #include intmain(){ constintnum=100; int*g; cudaError_tr; //主機陣列&初始化 inta[num],b[num]; for(intk=0;k裝置記憶體g[] r=cudaMemcpy(g,a,sizeof(int)*num,cudaMemcpyHostToDevice); printf("cudaMemcpya=>g:%s ",cudaGetErrorString(r)); //複製記憶體:裝置記憶體g[]------>主機記憶體b[] r=cudaMemcpy(b,g,sizeof(int)*num,cudaMemcpyDeviceToHost); printf("cudaMemcpyg=>b:%s ",cudaGetErrorString(r)); //結果比對 boolooo=true; for(intk=0;kg:noerror cudaMemcpyg=>b:noerror checka==b?:pass cudaFree:noerror ============================================================================ 第三招函式&呼叫(主機、裝置) ============================================================================ CUDA中，主機函式的寫法與呼叫和傳統C/C++無異，而裝置核心(kernel)要使用 延伸語法： __global__void函式名稱(函式引數...){ ...函式內容... }; 多了__global__這標籤來標明這道函式是核心程式碼，要編譯器特別照顧一下， 注意事項如下： (1)傳回值只能是void(要傳東西出來請透過引數) (2)裡面不能呼叫主機函式或global函式(這兩者皆是主機用的) (3)輸入的資料若是位址或參考時，必需指向裝置記憶體。

呼叫kernel函式的語法比一般C函式多了指定網格和區塊大小的手序： 函式名稱<<>>(函式引數...); 網格和區塊詳見第四招 //----------------------------------------------------------------------- //範例(2):helloCUDA函式(使用global函式填入字串)[081112-hello.cu] //----------------------------------------------------------------------- #include #include //裝置函式(核心)在顯示卡記憶體中填入helloCUDA字串 __global__voidhello(char*s){ charw[50]="helloCUDA~~~=^.^="; intk; for(k=0;w[k]!=0;k++)s[k]=w[k]; s[k]=0; }; //主機函式 intmain(){ char*d; charh[100]; //配置裝置記憶體 cudaMalloc((void**)&d,100); //呼叫裝置核心(只使用單一執行緒) hello<<<1,1>>>(d); //下載裝置記憶體內容到主機上 cudaMemcpy(h,d,100,cudaMemcpyDeviceToHost); //顯示內容 printf("%s ",h); //釋放裝置記憶體 cudaFree(d); return0; } ------------------------------------------------------------- 範例(2)執行結果: ------------------------------------------------------------- helloCUDA~~~=^.^= ============================================================================ 第四招網格、區塊、執行緒(線程群組) ============================================================================ 網格、區塊、執行緒是CUDA中最重要的部份,必需熟悉 (1)GPU是具備超多核心，能行大量平行化運算的晶片，執行緒眾多，要分群組管理： 最基本的執行單位是【執行緒】(thread)， 數個執行緒組成【區塊】(block)， 數個區塊組成【網格】(grid)， 整個網格就是所謂的【核心】(kernel)。

(2)【執行緒】是最基本的執行單位，程式設計師站在執行緒的角度，透過內建變數， 定出執行緒的位置，對工作進行主動切割。

(3)【區塊】為執行緒的群組，一個區塊可包含1~512個執行緒， 每個執行緒在區塊中擁有唯一的索引編號，記錄於內建變數threadIdx。

每個區塊中包含的執行緒數目，記錄於內建變數blockDim。

相同區塊內的執行緒可同步化，而且可透過共享記憶體交換資料(詳見第五、六招) (4)【網格】為區塊的群組，一個網格可包含1~65535個區塊， 每個區塊在網格中擁有唯一的索引編號，記錄於內建變數blockIdx。

每個網格中包含的區塊數目，記錄於內建變數gridDim。

網格中的區塊可能會同時或分散在不同時間執行，視硬體情況而定。

(5)內建唯讀變數gridDim,blockDim,blockIdx,threadIdx皆是3D正整數的結構體 uint3gridDim：網格大小(網格中包含的區塊數目) uint3blockIdx：區塊索引(網格中區塊的索引) uint3blockDim：區塊大小(區塊中包含的執行緒數目) uint3threadIdx：執行緒索引(區塊中執行緒的索引) 其中uint3為3D的正整數型態，定義如下 structuint3{ unsignedintx,y,z; }; 這些唯讀變數只能在核心中使用。

(6)核心呼叫時指定的網格和區塊大小對應的就是其中gridDim和blockDim兩變數 uint3gridDim：網格大小(網格中包含的區塊數目) uint3blockDim：區塊大小(區塊中包含的執行緒數目) 可以在呼叫時只指定一維，此時變數裡面的y和z成員都等於1： 核心名稱<<>>(引數...); 也可以指定三維的呼叫: 核心名稱<<>>(引數...); 或者混合使用: 核心名稱<<>>(引數...); 核心名稱<<>>(引數...); 其中dim3等於uint3，只是有寫好constructor而己。

(7)網格和區塊大小在設定時有一定的限制 網格:max(gridDim)=65535 區塊:max(blockDim)=512 實際在用的時候blockDim還會有資源上的限制,主要是暫存器數目, 所以有時達不到512這個數量,在3維的情況還會有其它的限制, 建議使用1維的方式呼叫,到核心中再去切,執行緒組態比較簡單, 而且bug和限制也會比較少. //----------------------------------------------------------------- //範例(3):列出在各執行緒中看到的區塊和執行緒索引[081112-idx.cu] //【使用一維結構】 //----------------------------------------------------------------- #include #include //索引用到的緒構體 structIndex{ intblock,thread; }; //核心:把索引寫入裝置記憶體 __global__voidprob_idx(Indexid[]){ intb=blockIdx.x;//區塊索引 intt=threadIdx.x;//執行緒索引 intn=blockDim.x;//區塊中包含的執行緒數目 intx=b*n+t;//執行緒在陣列中對應的位置 //每個執行緒寫入自己的區塊和執行緒索引. id[x].block=b; id[x].thread=t; }; //主函式 intmain(){ Index*d; Indexh[100]; //配置裝置記憶體 cudaMalloc((void**)&d,100*sizeof(Index)); //呼叫裝置核心 intg=3,b=4,m=g*b; prob_idx<<>>(d); //下載裝置記憶體內容到主機上 cudaMemcpy(h,d,100*sizeof(Index),cudaMemcpyDeviceToHost); //顯示內容 for(inti=0;i #include //索引用到的緒構體 structIndex{ uint3block,thread; }; //核心:把索引寫入裝置記憶體 __global__voidprob_idx_3d(Index*id){ //計算區塊索引 intb=(blockIdx.z*gridDim.y+blockIdx.y)*gridDim.x+blockIdx.x; //計算執行緒索引 intt=(threadIdx.z*blockDim.y+threadIdx.y)*blockDim.x+threadIdx.x; //計算區塊中包含的執行緒數目 intn=blockDim.x*blockDim.y*blockDim.z; //執行緒在陣列中對應的位置 intx=b*n+t; //每個執行緒寫入自己的區塊和執行緒索引. id[x].block=blockIdx; id[x].thread=threadIdx; } //主函式 intmain(){ //網格和區塊大小設定 dim3grid=dim3(4,1,1); dim3block=dim3(2,3,1); printf("gridDim=dim3(%d,%d,%d) ",grid.x,grid.y,grid.z); printf("blockDim=dim3(%d,%d,%d) ",block.x,block.y,block.z); //計算總執行緒數 intnum=grid.x*grid.y*grid.z*block.x*block.y*block.z; printf("totalnumofthreads=%d ",num); //配置主機記憶體&清空 Index*h=(Index*)malloc(num*sizeof(Index)); memset(h,0,num*sizeof(Index)); //配置裝置記憶體&清空 Index*d; cudaMalloc((void**)&d,num*sizeof(Index)); cudaMemcpy(d,h,num*sizeof(Index),cudaMemcpyHostToDevice); //呼叫裝置核心. prob_idx_3d<<>>(d); //測試是否執行成功. cudaError_tr=cudaGetLastError(); printf("prob_idx_3d:%s ",cudaGetErrorString(r)); if(r!=0)gotoend; //下載裝置記憶體內容到主機上. cudaMemcpy(h,d,num*sizeof(Index),cudaMemcpyDeviceToHost); //顯示內容 for(inti=0;i