繼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):使多執行單元的進度在某點上對齊(先到的要等待還沒到的,等全部到齊後才繼續前進),通常是為了交換共用資料,避免讀寫順序錯亂導致的資料錯誤


 

創作者介紹

Philippe's Expérience Note

philip 發表在 痞客邦 PIXNET 留言(0) 人氣()