作者: a5000ml (咖啡裡的海洋藍) 看板: VideoCard
標題: [分享] CUDA 程式設計(2) -- SIMT概觀
時間: Thu Oct  2 01:24:58 2008

感謝大家的支持，這禮拜讓我們來談談 CUDA 多執行緒的程式模型(SIMT)，好讓大家
對這個平行化的 C++ 有更清楚的輪廓。


※ 第二章 SIMT 概觀 ※

所謂 SIMT (single instruction multi threads) 指的是單一指令對應多執行緒的
計算機架構，利用硬體的 thread 來隱藏 I/O 的延遲 (效果有點類似之前 Intel 的
hyper-threading，不過那不是 single instruction)，nVidia 進一步地讓這些執行緒
可由程式控制，用群組的方式讓一堆執行緒執行相同的指令，並利用超多核心來強化它
(例如 8800 GTX 有 128 顆、GTX280 有 240 顆)。

簡而言之，它是把超級電腦的平行架構，濃縮到單晶片中，所以產生這樣的效能
(例如我實驗室裡的 kernel，在 GTX280上跑的效能是 Intel Q9300 的 30 多倍，
這測量的時間是實際跑完的時間，用 CPU 的高精度 timer 測量出來的，對照的是用
intel 自家的 compiler 進行 SSE3 最佳化過的)。

不過剛開始進入這多執行緒的模型，還真的有點不太習慣哩。

◆ CUDA 的平行化程式設計模型

名詞定義
        網格(Grid)    ：包含數個區塊的執行單元
        區塊(Block)   ：包含數個執行緒的執行單元
        執行緒(Thread)：最小的處理單元 (實際寫程式的環境)

CUDA 的平行化模型是將核心交由一組網格執行，再將網格切成數個區塊，然後每個區塊
再分成數個執行緒，依次分發進行平行運算，如果用軍隊來比喻，將核心視為連任務，
那網格就是連隊，區塊就是排或班，執行緒就是小兵。














   任務(kernel)
        |
        |               +--> 區塊(排or班) +--> 執行緒(小兵)
        |               |                 +--> 執行緒(小兵)
        |               |                 +--> 執行緒(小兵)
        |               |                 +--> 執行緒(小兵)
        |               |
        +--> 網格(連隊) +--> 區塊(排or班) +--> 執行緒(小兵)
                        |                 +--> 執行緒(小兵)
                        |                 +--> 執行緒(小兵)
                        |                 +--> 執行緒(小兵)
                        |
                        +--> 區塊(排or班) +--> 執行緒(小兵)
                                          +--> 執行緒(小兵)
                                          +--> 執行緒(小兵)
                                          +--> 執行緒(小兵)

        (圖一) kernel、網格、區塊、執行緒和軍隊的類比




◆ 內建變數

我們可以透過內建變數來辨識每個執行緒，讓每個小兵弄清楚要執行那一部份的任務，
基本的內建變數如下，它們只可以使用在 kernel 的程式碼中：

        uint3 gridDim  ：網格大小   (網格包含的區塊數目)
        uint3 blockIdx ：區塊索引   (區塊的ID)
        uint3 blockDim ：區塊大小   (每個區塊包含的執行緒數目)
        uint3 threadIdx：執行緒索引 (執行緒的ID)

其中 uint3 為 3D 的正整數型態

        struct uint3{
                unsigned int x,y,z;
        };

可以運用它來實做更高層次的平行運算結構，不過現階段，先不要管這種複雜的結構，
把它當成單一正整數即可，也就是 y 和 z 都當成是 1，只用 uint3 的 x。

ps. 其實我平常在寫的時候，也很少用到3D結構，因為我們的研究是4D或5D的 ~>_<~
    只好用1D載入kernel再自已去切。


◆ 網格 & 區塊大小 (gridDim, blockDim)

CUDA 透過指定網格和區塊的大小形成平行化的程式陣列，總執行緒數目為網格大小和
區塊大小的乘積，而 gridDim, blockDim 這兩個變數在 kernel 函式中為內建的唯讀
變數，可直接讀取

        總執行緒數目 = 網格大小(gridDim) x 區塊大小(blockDim)

例如下圖為 (網格大小=3, 區塊大小=4) 所形成的核心，它具有 12 個獨立的執行緒

        +-----------+-----------+--------------------+
        |           |           | thread 0   (id  0) |
        |           |           +--------------------+
        |           |           | thread 1   (id  1) |
        |           |  block 0  +--------------------+
        |           |           | thread 2   (id  2) |
        |           |           +--------------------+
        |           |           | thread 3   (id  3) |
        |           +-----------+--------------------+
        |           |           | thread 0   (id  4) |
        |           |           +--------------------+
        |           |           | thread 1   (id  5) |
        |  grid     |  block 1  +--------------------+
        | (kernel)  |           | thread 2   (id  6) |
        |           |           +--------------------+
        |           |           | thread 3   (id  7) |
        |           +-----------+--------------------+
        |           |           | thread 0   (id  8) |
        |           |           +--------------------+
        |           |           | thread 1   (id  9) |
        |           |  block 2  +--------------------+
        |           |           | thread 2   (id 10) |
        |           |           +--------------------+
        |           |           | thread 3   (id 11) |
        +-----------+-----------+--------------------+

        (圖二) 網格、區塊、執行緒 ID 的劃分









◆ 呼叫 kernel 的語法

在 CUDA 中呼叫 kernel 函式的語法和呼叫一般 C 函式並沒什麼太大的差異，
只是多了延伸的語法來指定網格和區塊大小而已：

        kernel_name <<<gridDim,blockDim>>> (arg1, arg2, ...);
        ^^^^^^^^^^^     ^^^^^   ^^^^^^      ^^^^^^^^^^^^^^^
        核心函式名   網格大小   區塊大小    函式要傳的引數(和C相同)

        所以只是多了 <<<gridDim,blockDim>>> 指定大小而已 ^^y

其中 gridDim 和 blockDim 可以是固定數字或動態變數，例如

    (1) 固定數字
        ooxx_kernel<<<123,32>>>(result, in1, in2);

    (2) 動態變數
        int grid  = some_function_g(); //計算網格大小
        int block = some_function_b(); //計算區塊大小
        ooxx_kernel<<<grid,block>>>(result, in1, in2);



◆ 區塊 & 執行緒索引 (blockIdx, threadIdx)

我們可以用區塊和執行緒索引來定出正在執行的程式位置，以決定該載入什麼樣的資料，
而 blockIdx, threadIdx 這兩個變數和 gridDim, blockDim 一樣，在 kernel 中也是
內建的唯讀變數，可直接讀取

例如在(圖二)中，我們要定出每一個小兵的唯一的 ID，可用下面這段程式碼

        int id = blockIdx.x*blockDim.x + threadIdx.x;

要產生(圖二)配置的 kernel 呼叫為

        kernel<<<3,4>>>(arguments);

其行為如下
(1) 傳入的網格和區塊大小為 1D 正整數，所以 uint3 中只有 x 有用到，其它 y=z=1
(2) 網格大小   gridDim.x   = 3       (每個網格包含 3 個區塊)
(3) 區塊大小   blockDim.x  = 4       (每個區塊包含 4 個執行緒)
(4) 區塊索引   blockIdx.x  = 0,1,2   (每個 thread 看到的不一樣)
(5) 執行緒索引 threadIdx.x = 0,1,2,3 (每個 thread 看到的不一樣)
(6) 區塊基底   blockIdx.x*blockDim.x = 0,4,8
(7) 區塊基底加上執行緒索引  id = blockIdx.x*blockDim.x + threadIdx.x
    = 0,1,2,3,  4,5,6,7,  8,9,10,11

所以我們便可得到一個全域的索引，即每一個小兵的唯一的 ID (圖二中的 id 欄)。





















◆ kernel 函式的語法

用 CUDA 寫 kernel 函式寫一般 C 函式也是沒什麼太大的差異，只是多了延伸語法來
加入一些特殊功能，並且標明這個函式是 kernel 而已：

        __global__ void kernel_name(type1 arg1, type2 arg2, ...){
                ...函式內容...
        };

其中
(1) __global__  : 標明這是 kernel 的延伸語法
(2) void        : kernel 傳回值只能是 void (要傳東西出來請透過引數)
(3) kernel_name : 函式名稱
(4) type1 arg1, type2 arg2, ... : 函式引數 (和 C 完全相同)
(5) 函式內容    : 跟寫 C 或 C++ 一樣 (但不能夠呼叫主機函式)
(6) global 函式只能在 host 函式中呼叫，不能在其它 global 中呼叫。







◆ 小結

以上是 CUDA 平行化程式設計的概觀，和傳統 C/C++ 的差異便是它這種的 SIMT 結構，
也許你會覺得奇怪，為什麼要分成兩層的 grid/block 結構，直接一層就配多個 thread
不是更簡單，這牽涉到它硬體上的細節以及成本問題(後面章節會解釋)，再者單層結構
不見得有效率，會增加同步化時執行緒等待的問題，使用兩層結構，可以使 block 單元
彈性的選擇同時或者循序執行，增加往後硬體發展和軟體重用的彈性。




※ 後續章節 ※
CUDA 安裝
簡易 kernel 範例
CUDA 的記憶體分類
CUDA 的函式種類
CUDA API介紹
GPGPU 的硬體介紹

(順序還在研究中... >_<)



※ 名詞解釋 ※
(1)SIMT(single instruction multi threads)：單一指令對應多執行緒的架構。
(2)網格  (Grid)  ：包含數個區塊的執行單元。
(3)區塊  (Block) ：包含數個執行緒的執行單元。
(4)執行緒(Thread)：最小的處理單元 (實際寫程式的環境)。
(5)核心  (Kernal)：並非執行單元，比較像是要執行某種任務的抽象歸類。
(6)網格大小(gridDim,  grid  dimension)：網格包含的區塊數目。
(7)區塊大小(blockDim, block dimension)：區塊包含的執行緒數目。
(8)區塊索引(blockIdx, block index)    ：區塊在網格中的位置。
(9)執行緒索引(threadIdx, thread index)：執行緒在區塊中的位置。
(10)唯讀變數(read-only variable)：只可讀取，不可寫入的變數。
(11)延伸語法(extension)：在標準C/C++語法之外，外加的功能性語法。
(12)函式引數(arguments)：函式呼叫時傳遞的變數。
(13)基底(base) ：計算位址時的基準點，就像座標的原點一樣。
(14)索引(index)：位址相對於基準點的偏移。
(15)同步化(synchronize)：使多執行單元的進度在某點上對齊(先到的要等待還沒到的，
    等全部到齊後才繼續前進)，通常是為了交換共用資料，避免讀寫順序錯亂導致的
    資料錯誤。



--
※ 發信站: 批踢踢實業坊(ptt.cc)
◆ From: 114.45.208.87
推 levine21:先搶頭推再來看!!                                       10/02 01:27
推 Luciferspear:雖然我看不懂還是推認真分享文                       10/02 01:30
推 CDavid:推                                                       10/02 01:30
※ 編輯: a5000ml         來自: 114.45.208.87        (10/02 02:06)
推 l23456:快推...不然別人以為我看不懂...                           10/02 02:21
推 nanie:真的是做研究才有機會用到  XDD 大學的話 幾乎沒機會用到     10/02 03:05
推 lavatar:推一下...                                               10/02 09:31
推 yayax:推推~看了頗久，真用心                                     10/02 09:53
推 b24333666:專業推                                                10/02 10:33
推 brighton16:大推心得分享文                                       10/02 10:59
→ b24333666:如果有大學做出來就有一台超級電腦了XDD                 10/02 11:08
推 b24333666:sdk:另外要cuda在windows上有5秒的限制 建議用linux會    10/02 11:10
→ b24333666:比較好　有５秒限制是啥意思?                           10/02 11:11
推 sdk:windows/xwinow每5秒watchdog會check顯卡有沒有死掉..如果這5   10/02 12:14
→ sdk:秒內你的程式一直在跑..他會判定GPU掛了而reset GPU...         10/02 12:15
→ sdk:windows上不能把window manager關掉..但linux上可以XD          10/02 12:15
→ sdk:其實不管是N社還是A社就是想搞HPC on desktop啊...             10/02 12:16
→ sdk:另外補充一下不管grid/block...基本單元就是thread..在GPU上    10/02 12:17
→ sdk:他們使用massive threaded architecture..也就是說一次可以執   10/02 12:18
→ sdk:行上萬個thread...才是最有效率的...(既使只有240個core..但是  10/02 12:18
→ sdk:GPU上的context switch overhead幾乎是0..)..這樣的設計是為了  10/02 12:19
→ sdk:hide global memory access latency..                         10/02 12:19
→ sdk:(早知道就回文了...= =)                                      10/02 12:20
推 finalhaven:可以用這個做電路的Place&Route的EDA TOOL嗎?           10/02 13:14
推 sdk:回頁上..應該ok..只是國外已經有start-up在做這個了XD          10/02 13:25
→ sdk:另外我也做過floorplan的部份..雖然只optimize area...         10/02 13:26
推 VictorTom:推文一起推:)                                          10/02 15:58
推 powertodream:推!!!!!                                            10/02 16:18
推 markfang:推 專業文  受教                                        10/03 11:23
推 airwin:原PO真強者阿                                             10/03 12:42
推 amd3dnow:從網路上看到cuda執行以warp為單位,那warp指的是??        10/03 22:05
推 netsphere:要是有範例會更好  : )                                 10/04 11:09
→ a5000ml:warp 之後會介紹哦, 它和硬體組成有關, 是 32 threads      10/04 11:50
→ a5000ml:並起來執行的單位, 用 8 個執行單元管線以 4 個週期執行    10/04 11:52
→ a5000ml:所以達成平均 1 個週期 1 個指令的效果                    10/04 11:53
→ a5000ml:至於範例後面就會有, 因為現在只是剛開始而己, 介紹的      10/04 11:55
→ a5000ml:比較傾向於概念部份, 等寫完安裝部份後, 開始進入          10/04 11:57
→ a5000ml:寫程式的正題, 就會有很多範例  =^.^=                     10/04 11:58
※ AE2001:轉錄至看板 NTUT_MMRE86                                   10/05 01:05
※ a5000ml:轉錄至看板 C_and_CPP                                    10/15 20:59
※ uf2000uf:轉錄至某隱形看板                                       10/16 21:48