[分享] CUDA 程式設計(2) -- SIMT概觀

看板C_and_CPP (C/C++)作者 (咖啡裡的海洋藍)時間16年前 (2008/10/15 20:59), 編輯推噓2(200)
留言2則, 2人參與, 最新討論串1/1
※ [本文轉錄自 VideoCard 看板] 作者: 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

10/02 01:27,
先搶頭推再來看!!
10/02 01:27

10/02 01:30,
雖然我看不懂還是推認真分享文
10/02 01:30

10/02 01:30,
10/02 01:30
※ 編輯: a5000ml 來自: 114.45.208.87 (10/02 02:06)

10/02 02:21,
快推...不然別人以為我看不懂...
10/02 02:21

10/02 03:05,
真的是做研究才有機會用到 XDD 大學的話 幾乎沒機會用到
10/02 03:05

10/02 09:31,
推一下...
10/02 09:31

10/02 09:53,
推推~看了頗久,真用心
10/02 09:53

10/02 10:33,
專業推
10/02 10:33

10/02 10:59,
大推心得分享文
10/02 10:59

10/02 11:08,
如果有大學做出來就有一台超級電腦了XDD
10/02 11:08

10/02 11:10,
sdk:另外要cuda在windows上有5秒的限制 建議用linux會
10/02 11:10

10/02 11:11,
比較好 有5秒限制是啥意思?
10/02 11:11

10/02 12:14,
windows/xwinow每5秒watchdog會check顯卡有沒有死掉..如果這5
10/02 12:14

10/02 12:15,
秒內你的程式一直在跑..他會判定GPU掛了而reset GPU...
10/02 12:15

10/02 12:15,
windows上不能把window manager關掉..但linux上可以XD
10/02 12:15

10/02 12:16,
其實不管是N社還是A社就是想搞HPC on desktop啊...
10/02 12:16

10/02 12:17,
另外補充一下不管grid/block...基本單元就是thread..在GPU上
10/02 12:17

10/02 12:18,
他們使用massive threaded architecture..也就是說一次可以執
10/02 12:18

10/02 12:18,
行上萬個thread...才是最有效率的...(既使只有240個core..但是
10/02 12:18

10/02 12:19,
GPU上的context switch overhead幾乎是0..)..這樣的設計是為了
10/02 12:19

10/02 12:19,
hide global memory access latency..
10/02 12:19

10/02 12:20,
(早知道就回文了...= =)
10/02 12:20

10/02 13:14,
可以用這個做電路的Place&Route的EDA TOOL嗎?
10/02 13:14

10/02 13:25,
回頁上..應該ok..只是國外已經有start-up在做這個了XD
10/02 13:25

10/02 13:26,
另外我也做過floorplan的部份..雖然只optimize area...
10/02 13:26

10/02 15:58,
推文一起推:)
10/02 15:58

10/02 16:18,
推!!!!!
10/02 16:18

10/03 11:23,
推 專業文 受教
10/03 11:23

10/03 12:42,
原PO真強者阿
10/03 12:42

10/03 22:05,
從網路上看到cuda執行以warp為單位,那warp指的是??
10/03 22:05

10/04 11:09,
要是有範例會更好 : )
10/04 11:09

10/04 11:50,
warp 之後會介紹哦, 它和硬體組成有關, 是 32 threads
10/04 11:50

10/04 11:52,
並起來執行的單位, 用 8 個執行單元管線以 4 個週期執行
10/04 11:52

10/04 11:53,
所以達成平均 1 個週期 1 個指令的效果
10/04 11:53

10/04 11:55,
至於範例後面就會有, 因為現在只是剛開始而己, 介紹的
10/04 11:55

10/04 11:57,
比較傾向於概念部份, 等寫完安裝部份後, 開始進入
10/04 11:57

10/04 11:58,
寫程式的正題, 就會有很多範例 =^.^=
10/04 11:58
AE2001:轉錄至看板 NTUT_MMRE86 10/05 01:05 -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 114.45.209.219

10/16 12:56, , 1F
10/16 12:56, 1F

07/01 15:10, , 2F
push
07/01 15:10, 2F
文章代碼(AID): #18zUb8zU (C_and_CPP)
文章代碼(AID): #18zUb8zU (C_and_CPP)