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

看板C_and_CPP (C/C++)作者a5000ml (咖啡裡的海洋藍)時間16年前 (2008/11/12 23:10)推噓2(2推 0噓 1→)

留言3則, 3人參與討論串1/1

※ [本文轉錄自 VideoCard 看板] 作者: a5000ml (咖啡裡的海洋藍) 看板: VideoCard 標題: [分享] CUDA 程式設計(10) -- 速成篇(上) 時間: Wed Nov 12 22:53:25 2008 (1) 有學弟反應 CUDA 內容有點繁雜, 很多概念容易搞混, 而且希望多點範例, 所以這兩個禮拜把之前的文稿整理成【新手速成篇】,希望對他們有所幫助. (2) 順便幫國網打個廣告: CUDA 中文教學 DVD (免費線上版) 出現了請至國網的教育訓練網登入 https://edu.nchc.org.tw 詳情請看編號 18026 一文 ※ 第十章新手速成篇(上) ============================================================================ 前言 ============================================================================ 因為 CUDA 的一些延伸語法太繁雜，容易讓人混淆 (例如記憶體種類就有4~5種，同樣的 global memory 又有兩種寫法)，所以針對這個問題，寫成了速成篇，去除那些枝枝節節，只講最重要的，並佐以範例，務求讓初學者【七招闖天下】。第一招主機、裝置第二招使用 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_t cudaMalloc(void** ptr, size_t count); ptr 指向目的指位器之位址 count 欲配置的大小(單位 bytes) 傳回值 cudaError_t 是個 enum, 執行成功時傳回 0, 其它的錯誤代號可用 cudaGetErrorString() 來解譯. -------------------------------------------------------- (2)釋放顯示記憶體 cudaFree() [cuda.h] -------------------------------------------------------- cudaError_t cudaFree(void* ptr); ptr 指向欲釋放的位址 (device memory) -------------------------------------------------------- (3)記憶體複製 cudaMemcpy() [cuda.h] -------------------------------------------------------- cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind); dst 指向目的位址 src 指向來源位址 count 拷貝區塊大小 (單位 bytes) kind 有四種拷貝流向 cudaMemcpyHostToHost 主機 -> 主機 cudaMemcpyHostToDevice 主機 -> 裝置 cudaMemcpyDeviceToHost 裝置 -> 主機 cudaMemcpyDeviceToDevice 裝置 -> 裝置 -------------------------------------------------------- (4)錯誤字串解譯 cudaGetErrorString() [cuda.h] -------------------------------------------------------- const char* cudaGetErrorString(cudaError_t error); 傳回錯誤代號(error)所代表的字串 -------------------------------------------------------- (5)同步化 cudaThreadSynchronize() [cuda.h] -------------------------------------------------------- cudaError_t cudaThreadSynchronize(void); 使前後兩個核心時序上分離, 確保資料的前後相依性正確 //------------------------------------------------------------------------- //範例(1): 透過裝置記憶體進行複製 [081112-api.cu] // PCI-E PCI-E // 主機記憶體 a[] --------> 裝置記憶體 g[] --------> 主機記憶體 b[] //------------------------------------------------------------------------- #include<stdio.h> #include<cuda.h> int main(){ const int num=100; int* g; cudaError_t r; //主機陣列 & 初始化 int a[num], b[num]; for(int k=0; k<num; k++){ a[k]=k; b[k]=0; } //配置裝置記憶體 & 顯示錯誤訊息 r=cudaMalloc((void**) &g, sizeof(int)*num); printf("cudaMalloc : %s\n",cudaGetErrorString(r)); //複製記憶體: 主機記憶體 a[] ------> 裝置記憶體 g[] r=cudaMemcpy(g, a, sizeof(int)*num, cudaMemcpyHostToDevice); printf("cudaMemcpy a => g : %s\n",cudaGetErrorString(r)); //複製記憶體: 裝置記憶體 g[] ------> 主機記憶體 b[] r=cudaMemcpy(b, g, sizeof(int)*num, cudaMemcpyDeviceToHost); printf("cudaMemcpy g => b : %s\n",cudaGetErrorString(r)); //結果比對 bool ooo=true; for(int k=0; k<num; k++){ if(a[k]!=b[k]){ ooo=false; break; } } printf("check a==b? : %s\n",ooo?"pass":"wrong"); //釋放裝置記憶體 r=cudaFree(g); printf("cudaFree : %s\n",cudaGetErrorString(r)); return 0; } ------------------------------------------------------------- 範例(1)執行結果: ------------------------------------------------------------- cudaMalloc : no error cudaMemcpy a => g : no error cudaMemcpy g => b : no error check a==b? : pass cudaFree : no error ============================================================================ 第三招函式 & 呼叫 (主機、裝置) ============================================================================ CUDA 中，主機函式的寫法與呼叫和傳統 C/C++ 無異，而裝置核心 (kernel) 要使用延伸語法： __global__ void 函式名稱 (函式引數...){ ...函式內容... }; 多了 __global__ 這標籤來標明這道函式是核心程式碼，要編譯器特別照顧一下，注意事項如下： (1) 傳回值只能是 void (要傳東西出來請透過引數) (2) 裡面不能呼叫主機函式或 global 函式 (這兩者皆是主機用的) (3) 輸入的資料若是位址或參考時，必需指向裝置記憶體。呼叫 kernel 函式的語法比一般 C 函式多了指定網格和區塊大小的手序：函式名稱 <<<網格大小, 區塊大小>>> (函式引數...); 網格和區塊詳見第四招 //----------------------------------------------------------------------- //範例(2): hello CUDA 函式 (使用 global 函式填入字串) [081112-hello.cu] //----------------------------------------------------------------------- #include<stdio.h> #include<cuda.h> //裝置函式(核心) 在顯示卡記憶體中填入 hello CUDA 字串 __global__ void hello(char* s){ char w[50]="hello CUDA ~~~ =^.^="; int k; for(k=0; w[k]!=0; k++) s[k]=w[k]; s[k]=0; }; //主機函式 int main(){ char* d; char h[100]; //配置裝置記憶體 cudaMalloc((void**) &d, 100); //呼叫裝置核心 (只使用單一執行緒) hello<<<1,1>>>(d); //下載裝置記憶體內容到主機上 cudaMemcpy(h, d, 100, cudaMemcpyDeviceToHost); //顯示內容 printf("%s\n", h); //釋放裝置記憶體 cudaFree(d); return 0; } ------------------------------------------------------------- 範例(2)執行結果: ------------------------------------------------------------- hello CUDA ~~~ =^.^= ============================================================================ 第四招網格、區塊、執行緒 (線程群組) ============================================================================ 網格、區塊、執行緒是 CUDA 中最重要的部份, 必需熟悉 (1) GPU 是具備超多核心，能行大量平行化運算的晶片，執行緒眾多，要分群組管理：最基本的執行單位是【執行緒】(thread)，數個執行緒組成【區塊】(block)，數個區塊組成【網格】(grid)，整個網格就是所謂的【核心】(kernel)。 (2)【執行緒】是最基本的執行單位，程式設計師站在執行緒的角度，透過內建變數，定出執行緒的位置，對工作進行主動切割。 (3)【區塊】為執行緒的群組，一個區塊可包含 1~512 個執行緒，每個執行緒在區塊中擁有唯一的索引編號，記錄於內建變數 threadIdx。每個區塊中包含的執行緒數目，記錄於內建變數 blockDim。相同區塊內的執行緒可同步化，而且可透過共享記憶體交換資料 (詳見第五、六招) (4)【網格】為區塊的群組，一個網格可包含 1~65535 個區塊，每個區塊在網格中擁有唯一的索引編號，記錄於內建變數 blockIdx。每個網格中包含的區塊數目，記錄於內建變數 gridDim。網格中的區塊可能會同時或分散在不同時間執行，視硬體情況而定。 (5) 內建唯讀變數 gridDim, blockDim, blockIdx, threadIdx 皆是 3D 正整數的結構體 uint3 gridDim ：網格大小 (網格中包含的區塊數目) uint3 blockIdx ：區塊索引 (網格中區塊的索引) uint3 blockDim ：區塊大小 (區塊中包含的執行緒數目) uint3 threadIdx：執行緒索引 (區塊中執行緒的索引) 其中 uint3 為 3D 的正整數型態，定義如下 struct uint3{ unsigned int x,y,z; }; 這些唯讀變數只能在核心中使用。 (6) 核心呼叫時指定的網格和區塊大小對應的就是其中 gridDim 和 blockDim 兩變數 uint3 gridDim ：網格大小 (網格中包含的區塊數目) uint3 blockDim ：區塊大小 (區塊中包含的執行緒數目) 可以在呼叫時只指定一維，此時變數裡面的 y 和 z 成員都等於 1：核心名稱<<<int 網格大小, int 區塊大小>>>(引數...); 也可以指定三維的呼叫: 核心名稱<<<dim3 網格大小, dim3 區塊大小>>>(引數...); 或者混合使用: 核心名稱<<<dim3 網格大小, int 區塊大小>>>(引數...); 核心名稱<<<int 網格大小, dim3 區塊大小>>>(引數...); 其中 dim3 等於 uint3，只是有寫好 constructor 而己。 (7) 網格和區塊大小在設定時有一定的限制網格: max(gridDim) = 65535 區塊: max(blockDim) = 512 實際在用的時候 blockDim 還會有資源上的限制, 主要是暫存器數目, 所以有時達不到 512 這個數量, 在 3 維的情況還會有其它的限制, 建議使用 1 維的方式呼叫, 到核心中再去切, 執行緒組態比較簡單, 而且 bug 和限制也會比較少. //----------------------------------------------------------------- //範例(3): 列出在各執行緒中看到的區塊和執行緒索引 [081112-idx.cu] // 【使用一維結構】 //----------------------------------------------------------------- #include<stdio.h> #include<cuda.h> //索引用到的緒構體 struct Index{ int block, thread; }; //核心:把索引寫入裝置記憶體 __global__ void prob_idx(Index id[]){ int b=blockIdx.x; //區塊索引 int t=threadIdx.x; //執行緒索引 int n=blockDim.x; //區塊中包含的執行緒數目 int x=b*n+t; //執行緒在陣列中對應的位置 //每個執行緒寫入自己的區塊和執行緒索引. id[x].block=b; id[x].thread=t; }; //主函式 int main(){ Index* d; Index h[100]; //配置裝置記憶體 cudaMalloc((void**) &d, 100*sizeof(Index)); //呼叫裝置核心 int g=3, b=4, m=g*b; prob_idx<<<g,b>>>(d); //下載裝置記憶體內容到主機上 cudaMemcpy(h, d, 100*sizeof(Index), cudaMemcpyDeviceToHost); //顯示內容 for(int i=0; i<m; i++){ printf("h[%d]={block:%d, thread:%d}\n", i,h[i].block,h[i].thread); } //釋放裝置記憶體 cudaFree(d); return 0; } ------------------------------------------------------------- 範例(3)執行結果: ------------------------------------------------------------- h[0]={block:0, thread:0} h[1]={block:0, thread:1} h[2]={block:0, thread:2} h[3]={block:0, thread:3} h[4]={block:1, thread:0} h[5]={block:1, thread:1} h[6]={block:1, thread:2} h[7]={block:1, thread:3} h[8]={block:2, thread:0} h[9]={block:2, thread:1} h[10]={block:2, thread:2} h[11]={block:2, thread:3} //------------------------------------------------------------------- //範例(4): 列出在各執行緒中看到的區塊和執行緒索引 [081112-idx_3d.cu] // 【使用三維結構】 //------------------------------------------------------------------- #include<stdio.h> #include<cuda.h> //索引用到的緒構體 struct Index{ uint3 block, thread; }; //核心:把索引寫入裝置記憶體 __global__ void prob_idx_3d(Index* id){ //計算區塊索引 int b=(blockIdx.z*gridDim.y+blockIdx.y)*gridDim.x+blockIdx.x; //計算執行緒索引 int t=(threadIdx.z*blockDim.y+threadIdx.y)*blockDim.x+threadIdx.x; //計算區塊中包含的執行緒數目 int n=blockDim.x*blockDim.y*blockDim.z; //執行緒在陣列中對應的位置 int x=b*n+t; //每個執行緒寫入自己的區塊和執行緒索引. id[x].block=blockIdx; id[x].thread=threadIdx; } //主函式 int main(){ //網格和區塊大小設定 dim3 grid=dim3(4,1,1); dim3 block=dim3(2,3,1); printf("gridDim = dim3(%d,%d,%d)\n", grid.x,grid.y,grid.z); printf("blockDim = dim3(%d,%d,%d)\n", block.x,block.y,block.z); //計算總執行緒數 int num=grid.x*grid.y*grid.z*block.x*block.y*block.z; printf("total num of threads = %d\n", 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<<<grid,block>>>(d); //測試是否執行成功. cudaError_t r=cudaGetLastError(); printf("prob_idx_3d: %s\n", cudaGetErrorString(r)); if(r!=0) goto end; //下載裝置記憶體內容到主機上. cudaMemcpy(h, d, num*sizeof(Index), cudaMemcpyDeviceToHost); //顯示內容 for(int i=0; i<num; i++){ printf("h[%d]={block:(%d,%d,%d), thread:(%d,%d,%d)}\n", i, h[i].block.x, h[i].block.y, h[i].block.z, h[i].thread.x, h[i].thread.y, h[i].thread.z ); } end:; //釋放裝置記憶體. cudaFree(d); free(h); return 0; } ------------------------------------------------------------- 範例(4)執行結果: ------------------------------------------------------------- gridDim = dim3(4,1,1) blockDim = dim3(2,3,1) total num of threads = 24 prob_idx_3d: no error h[0]={block:(0,0,0), thread:(0,0,0)} h[1]={block:(0,0,0), thread:(1,0,0)} h[2]={block:(0,0,0), thread:(0,1,0)} h[3]={block:(0,0,0), thread:(1,1,0)} h[4]={block:(0,0,0), thread:(0,2,0)} h[5]={block:(0,0,0), thread:(1,2,0)} h[6]={block:(1,0,0), thread:(0,0,0)} h[7]={block:(1,0,0), thread:(1,0,0)} h[8]={block:(1,0,0), thread:(0,1,0)} h[9]={block:(1,0,0), thread:(1,1,0)} h[10]={block:(1,0,0), thread:(0,2,0)} h[11]={block:(1,0,0), thread:(1,2,0)} h[12]={block:(2,0,0), thread:(0,0,0)} h[13]={block:(2,0,0), thread:(1,0,0)} h[14]={block:(2,0,0), thread:(0,1,0)} h[15]={block:(2,0,0), thread:(1,1,0)} h[16]={block:(2,0,0), thread:(0,2,0)} h[17]={block:(2,0,0), thread:(1,2,0)} h[18]={block:(3,0,0), thread:(0,0,0)} h[19]={block:(3,0,0), thread:(1,0,0)} h[20]={block:(3,0,0), thread:(0,1,0)} h[21]={block:(3,0,0), thread:(1,1,0)} h[22]={block:(3,0,0), thread:(0,2,0)} h[23]={block:(3,0,0), thread:(1,2,0)} 我們可以由範例(3)和(4)看出執行緒索引的配置方式. =========================================================================== 待續... -- 。o O　○。o O　○。o O　○。o O　○。o O　○。o 國網 CUDA 中文教學 DVD 影片 (免費線上版) 請至國網的教育訓練網登入 https://edu.nchc.org.tw -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 114.45.214.93