Cuda編程系列-Cuda編程基本概念&編程模型

原文連接

系列文章：

• Cuda編程101：Cuda編程的基本概念及編程模型
• Cuda編程102：Cuda程序性能相關話題
• Cuda編程103：Cuda多卡編程
• Cuda tips: nvcc的-code、-arch、-gencode選項

基本想法

在介紹編碼相關內容以前，一個更重要的話題是什麼類型的問題適合用GPU進行解決。

GPU於CPU相比，有着驚人的核數、運算單元及內存帶寬。對於給定問題，若是有辦法把它分解爲多個獨立的子問題並行解決，那麼GPU頗有可能提供比CPU更好的性能。所謂「獨立」，指的是所分解的子問題知足：

• 子問題之間儘量避免同步
• 子問題之間儘量依賴使用全局內存同步狀態
• 子問題之間儘量避免同步關係

矩陣相乘就是一個很好的例子，對矩陣相乘結果中各個元素的計算之間沒有任何依賴關係，可以很好地經過GPU進行並行。固然對於一些問題，可能沒辦法馬上想出並行的辦法，可是卻存在可高效並行的問題分解辦法，比方說：

(思考題)

• 歸併兩個有序數組
• 對一個數組求前綴和

對於手頭的問題，若是可以順利對問題進行分解，那麼就有可能利用GPU提供的硬件特性及編程模型對其進行高效解決。

編程模型

硬件視角

• 一塊GPU上由多個Streaming Multiprocessors組成，簡稱SM。
• 每一個SM中包含多個core，即實際完成計算的單元。
• 以下圖所示，在一塊1080ti上有28個SM，每一個SM上有128個core，合計3584個cuda core。

編程視角

• 程序員編寫一個在GPU由多個thread並行執行的函數，並從CPU代碼對其調用。這樣的函數咱們將其稱爲一個kernel。
• 多個GPU thread組成一個thread block
• 對於一個kernel函數，程序員來指定啓動多少個thread block，每一個thread block裏有多少thread
• 每一個thread可以獲取本身在哪一個block中，以及本身是本block的第幾個thread。對於一個並行處理任務，thread可根據這些信息肯定本身應處理哪部分子問題。

執行視角

• 每一個thread block會被調度到其中一個SM上執行
• 對於一個thread block中的各個thread，每32個thread組成一個warp，SM以warp爲單位進行調度。在一個warp中，全部thread執行同一個指令流，即Single Instruction Multiple Thread(SIMT)。若是執行過程當中有分支語句，那麼執行不一樣分支的thread須要互相等待。比方說對於下列語句，任意時刻同一個warp中只能有一半的thread進行操做，而不是各自獨立執行本身所在的分支。在寫kernel時，不當的分支語句可能會致使性能降低。

if (threadIdx.x % 2 == 0) {
    // Some work
} else {
    // Other work
}
複製代碼

說點別的

32個thread組成的調度單元爲何叫warp？緣由是thread有線的意思，而warp是織布機相關的一個把多個thread固定注的裝置，因而就取了這個比喻：

資源限制

就像寫CPU代碼時會受到CPU核數、內存空間、訪存速度的限制同樣，GPU編程模型裏也須要留意相關的資源限制：

• 每一個thread block中的thread數量，1080ti的上限是1024
• 啓動kernel時thread block數量（這個涉及到所起的thread block多是多維的狀況，先暫時認爲是2147483647吧）
• 每一個SM能同時處理的thread block數量，1080ti的上限是32
• Shared memory的大小，1080ti的上限是96kB
• GPU的訪存速度，1080ti上限是484GB/s，若是真的受到了這個限制說明代碼寫得很是好了
• ...（還有好多）

來點代碼

CUDA編程中的常見流程是：

• 把CPU數據搬運到GPU中
• 寫一個kernel定義咱們想完成的計算
• 啓動kernel
• 把運算結果從GPU搬運回CPU中

Cuda樣例代碼中的vectorAdd完成的任務是對長爲numElements的兩個數組h_A、h_B進行對應元素加合，並將結果存入h_C中。接下來咱們以vectorAdd爲例，說明這一流程：

• 首先是把CPU數據搬運到GPU中

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
複製代碼

• 而後咱們定義咱們的加合計算

__global__ void
vectorAdd(const float *A, const float *B, float *C, int numElements)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < numElements)
    {
        C[i] = A[i] + B[i];
    }
}
複製代碼

作的事情就是每一個thread負責根據本身所在的thread block及threadIdx計算出本身所應處理的數組下標，並對這一下標對應的元素完成一次加合計算。

• 接下來咱們啓動kernel，其中<<<blocksPerGrid, threadsPerBlock>>>指定了thread block數量及每一個block中的thread數量。

vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
複製代碼

• 最後咱們把運算結果搬運回CPU中。

cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
複製代碼

一個cuda程序最重要的部分就完成了。完整代碼中還包含了內存的分配、cuda調用的錯誤檢查等內容，完整代碼可見cuda安裝目錄下的samples/0_Simple/vectorAdd。

一些須要留意的地方

• kernel的執行是異步的，啓動後會當即返回CPU代碼中。若是計時的話會發現時間極短，其實這個時間僅僅是kernel啓動的時間。
• 儘管kernel執行是異步的，然而cudaMemcpy又是阻塞的。
• 樣例代碼中的一次從CPU到GPU的cudaMemcpy調用其實完成了兩次內存拷貝，一次從CPU原內存拷貝到了CPU中一段page-lock內存中，再從這段內存拷貝到GPU內存。

更多話題

• GPU的內存層級
• GPU訪存pattern對性能的影響
• GPU的分支語句對性能的影響
• GPU中的同步操做、原子操做
• CPU、GPU間數據傳輸，PCIe，page-lock內存
• CUDA的debugger及profiler

思考題答案

• 歸併兩個有序數組:對於數組長度爲n、m的有序數組及t個thread，能夠對長爲n的數組進行t等分，並對每一子數組的起終點，二分找到數組m中對應的上界及下界，並基於此進行並行歸併。
• 對一個數組求前綴和: Parallel Prefix Sum (Scan) with CUDA

Reference

• CUDA C Programming Guide
• Udacity Intro to Parallel Programming