繼8個月前的 CUDA Introduction 後一直沒把其他的文章補上來，會趁這個月把文章補齊。
※ 第二章 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)：使多執行單元的進度在某點上對齊(先到的要等待還沒到的，等全部到齊後才繼續前進)，通常是為了交換共用資料，避免讀寫順序錯亂導致的資料錯誤
留言列表
OS (1)
