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

[a5000ml (咖啡裡的海洋藍)]

(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
→ fanzero:push
推 levine21:新手篇  我推~   我已經嚴重落後了 XDDD
※ a5000ml:轉錄至看板 C_and_CPP                                    11/12 23:10
推 dkfum:偷問大大 BT還要發佈嗎??
→ a5000ml:有網友那邊有壓縮好的, 畫質聽說不錯, 不知道有沒有人需要
→ a5000ml:要不要表決一下 BT 分享?
推 finkel:我也想要　我的光碟有點壞
推 tingyushyu:我也想要
→ finkel:只是不知道這有沒有版權問題？
推 Lineages:請問BT檔的檔案跟訓練網的影片檔案一樣嗎?
→ a5000ml:應該比訓練網的還清楚, 版權問過了沒問題
→ a5000ml:內容是完全一樣的
推 Lineages:清楚是畫面的解析度???
→ a5000ml:是啊, 連右下角的投影片都看得清楚
推 Lineages:這樣阿~了解了 非常謝謝
推 Dissipate:推
推 VictorTom:推; 話說剛把國網的wmv抓完, 才注意到有更清楚的....Orz
推 Luciferspear:推推
推 yayax:推阿~感謝
※ uf2000uf:轉錄至某隱形看板                                       11/13 09:43
推 b24333666:推一下
推 joyfulpizza:我也想要  推一下~
推 radeon9700:推一下BT分享 XD
推 mnmnqq:推!!
推 perchik:推BT分享 ^^
推 Colaman:推BT分享~~
推 hungmao:推!!跪求BT分享Q_Q
推 vixen:求種啊
推 allenchen821:  大推，新手入門