SIMT 對 SIMD(單指令多執行緒 對 單指令多資料)
/ SIM-tee versus SIM-dee /
現代 CPU 與 GPU 都能一次把一個運算套用到多筆資料上——這正是平行硬體的全部重點。但它們呈現給程式設計師的,是兩種不同的心智模型,把它們搞混會導致對各自效能的錯誤直覺。SIMD 與 SIMT 就是這兩種模型:兩種述說「單一指令、多筆資料」的方式。
SIMD(單指令多資料)是 CPU 向量的觀點。你撰寫的程式碼操作的是一個寬暫存器,裡頭比方說一次裝了 8 個浮點數;一條 add 指令在單一步驟內把全部 8 條通道相加。這些通道不是執行緒——只有一個程式計數器、一條指令流,平行性明明白白寫在資料寬度裡。你想的是「我有一個含 8 個元素的向量,一個運算打中它全部」。遮罩(masking)處理那種你只想讓部分通道作用的尷尬情況。SIMT(單指令多執行緒)是 GPU 的觀點(NVIDIA 的用語)。在這裡你把程式碼寫得像是給單一純量執行緒用的——「為我的 i 計算 c[i] = a[i] + b[i]」——而硬體會讓許多這樣的執行緒、一個 32 個的 warp,在共享的指令發派邏輯上步調一致地執行。它感覺像普通執行緒,但它們是齊步前進的;當一個 warp 內的執行緒走向不同的分支方向時,硬體必須把這些分支序列化(這就是 warp 發散),先執行走 if 的那群、讓走 else 的那群閒置,再對調——所以一個劈開 warp 的分支可能把吞吐量砍半。
弄清楚這個區別之所以重要,是因為它告訴你效能從哪裡來、又從哪裡漏掉。用 SIMD 時你以向量寬度來思考,而你這位程式設計師(或自動向量化器)必須塑形資料;用 SIMT 時你以獨立執行緒來思考,但分支與分散記憶體存取的代價被隱藏著,直到它咬你一口。它們確實是近親——SIMT 本質上就是配上較友善的逐執行緒程式設計模型與硬體管理遮罩的 SIMD——而誠實的總結是:相同的底層想法(單一指令施加於多筆資料)、不同的抽象、不同的失敗模式。
// SIMD(CPU):一條指令施加於一個由多通道組成的向量 // __m256 v = _mm256_add_ps(a, b); // 一次相加 8 個浮點數;一個 PC、8 條通道 // SIMT(GPU):寫一個純量執行緒;warp 以 32 個步調一致地執行 // c[i] = a[i] + b[i]; // 你的 'i';32 個執行緒一起發派這同一條 add
SIMD 在一條指令流中把向量明白攤開;SIMT 看起來像 32 個普通執行緒,卻偷偷以一個 warp 的形式齊步前進。
兩者是近親而非對立面:SIMT 本質上就是配上逐執行緒程式設計模型與硬體遮罩的 SIMD。關鍵的實務差異是 warp 發散——在 SIMT 中,一個把同一 warp 的執行緒送往不同路徑的分支,會把那些路徑序列化並浪費通道,這種代價被看似純量的程式碼隱藏了起來。