Warps、SIMD 硬體、Control Divergence 與 Warp Scheduling

#gpu-architecture #warps-simd #control-divergence #latency-tolerance

重點總覽 (Overview)

概念	核心定義	關鍵數字 / 規則	影響
Warp	SM 中執行緒排程的基本單位	32 個連續 `threadIdx` 的執行緒	一個 block 拆成 `⌈blockDim/32⌉` 個 warp
SIMD / SIMT	同一 warp 在同一時刻執行同一條指令	多核共用一個 fetch/dispatch unit	控制硬體成本攤提，增加算術吞吐
Control Divergence	同 warp 執行緒走不同控制流路徑	每條路徑各跑一個 pass	多 pass → 效能損失，inactive thread 浪費資源
Warp Scheduling	以 warp 為單位切換以隱藏長延遲	等待中的 warp 被換下，就緒的換上	latency tolerance / latency hiding
Zero-overhead Scheduling	全部 warp 狀態常駐硬體暫存器	無需 save/restore，無閒置 cycle	GPU 不需大量 cache 與分支預測
Oversubscription	指派遠多於核心數的執行緒	A100：64 cores 但可駐 2048 threads（32×）	提高找到就緒 warp 的機率

Important

Warp、SIMD、divergence、scheduling 都是硬體實作概念。程式正確性絕不可依賴 warp 內執行緒同步執行的假設；需要同步時必須用 __syncthreads()（block 層級）或 __syncwarp()（warp 層級）。

Warp 與 SIMD 硬體 (Warps and SIMD Hardware)

Block 如何切成 Warp

Block 指派到 SM 後，進一步切分成 32-thread 的 warp（warp size 是 implementation-specific，目前所有 CUDA device 都是 32）。
依 threadIdx 線性化後連續切割：

warp n 包含 thread 32·n 到 32·(n+1) − 1

Block size 不是 32 的倍數時，最後一個 warp 以 inactive thread 補滿 32 格（例：48 threads → 2 warps，第二個 warp 有 16 個 inactive thread）。

Block (256 threads) → 256/32 = 8 warps
SM 同時駐 3 個 block → 8 × 3 = 24 warps in the SM

Block 1: [W0:t0..31][W1:t32..63] ... [W7:t224..255]
              └ 32 ┘   └  32  ┘          └   32   ┘

多維 Block 的線性化（row-major）

多維 block 先投影成 row-major 線性順序，再切 warp：z 大的排在後、同 z 內 y 大的排後、同 y 內 x 遞增。
索引記法 T(z,y,x)：

2D block 8×8 (x=0..7, y=0..7)  →  64 threads = 2 warps
  Warp 0 = T(0,0) .. T(3,7)   ← y = 0,1,2,3 各 8 個 = 32
  Warp 1 = T(4,0) .. T(7,7)   ← y = 4,5,6,7

3D block z=2,y=8,x=4  →  64 threads = 2 warps
  Warp 0 = T(0,0,0) .. T(0,7,3)   (z=0 整層)
  Warp 1 = T(1,0,0) .. T(1,7,3)   (z=1 整層)

SIMD：多核共用一個控制單元

SM 的 cores 分組成 processing block（每組共用一個 instruction fetch/dispatch unit）。例：Ampere A100 SM 有 64 cores → 4 個 processing block × 16 cores。
同一 warp 的執行緒指派到同一 processing block，硬體取一條指令、同時對 warp 內所有執行緒執行，只是資料（register 內容）不同 → add r1,r2,r3 對每個 core 的 r2/r3 值不同。

        ┌──────────── Control Unit ────────────┐
        │  PC → IR (one instruction: add r1,r2,r3) │
        └───┬──────┬──────┬──────┬──────┬───────┘
   同一控制訊號廣播給所有 core ↓ (SIMD)
        Core0   Core1   Core2  ...  Core31
         t0      t1      t2          t31   ← 同指令、不同資料

Tip

SIMD 的好處：控制硬體（fetch/dispatch、instruction cache port）成本被眾多執行單元攤提，使更大比例的晶片面積用於算術運算而非控制 → 高吞吐。NVIDIA 把這種行為稱為 SIMT（Single-Instruction, Multiple-Thread）。

控制分歧 (Control Divergence)

定義與多 pass 執行

當同一 warp 內的執行緒走不同 control flow 路徑，即發生 control divergence。
SIMD 硬體會對每條路徑各跑一個 pass；某 pass 中走其他路徑的執行緒被設為 inactive（不生效）。

Warp t0..t31 遇到 if-else:  t0..t23 走 then(A)，t24..t31 走 else(B)

Pass 1:  [t0..t23 執行 A]  [t24..t31 INACTIVE]
Pass 2:  [t0..t23 INACTIVE] [t24..t31 執行 B]
Reconverge:  t0..t31 一起執行 C

架構	多 pass 執行方式
Pascal 及之前	passes 循序執行（一條跑完再跑另一條）
Volta 及之後	passes 可並行/交錯執行 → Independent Thread Scheduling

for-loop 也會 divergence：若各執行緒迭代次數不同（如 4~8 次），前 4 次全 active，之後部分執行緒 inactive。

何時會發生？如何判斷？

判斷準則：control 條件基於 threadIdx 時就可能 divergence。
- if (threadIdx.x > 2) {...} → 第一個 warp 內 thread 0,1,2 與其餘走不同路徑 → divergent。
最常見來源：處理 boundary condition（執行緒數須為 block size 倍數，但資料大小任意），例如 vector add 的 if (i < n)。

Divergence 的成本與資料規模

SIMD efficiency = (active threads) / 32

成本 = 額外 pass + inactive thread 佔用的執行資源。影響隨資料規模增大而下降：

範例 (block size 64)	總 warp 數	divergent warp	大致影響
vector length 100	4	1 (1/4)	顯著
vector length 1000	32	1 (1/32)	≈ 3%
vector length 10,000	313	1	< 1%

2D 例（62×76 影像、16×16 block）：160 warps 中 31 個 divergent；放大到 200×150 → 1040 warps 中僅 80 個 divergent（< 8%）；真實 >1000 寬度影像 < 2%。

Warning

Control divergence 代表不能假設 warp 內所有執行緒有相同的執行時序。若 warp 內需「全部完成某階段才繼續」，必須用 __syncwarp() 確保正確性（注意：這與 block 層級的 __syncthreads() 不同層級）。

Warp 排程與延遲容忍 (Warp Scheduling and Latency Tolerance)

為什麼要 oversubscription？

指派到 SM 的執行緒數遠多於 cores 數；任一時刻 SM 只執行其中一小部分 warp 的指令。
目的：容忍長延遲操作（如 global memory access、pipelined FP、branch）。

為何要塞這麼多 warp？→ Latency Tolerance（延遲容忍 / 隱藏）

時間軸 →
 W0 |■■■─────(等 global memory)────■■|
 W1 |   ■■■────(等記憶體)───────■■   |
 W2 |      ■■■─────────────■■        |
 W3 |         ■■■───────■■           |
     └ 任一時刻硬體都能找到一個「就緒」的 warp 來填滿執行單元

排程機制

某 warp 的下一條指令需等待先前發出的長延遲操作結果 → 該 warp 不被選中。
改選另一個已就緒（不再等結果）的 resident warp 執行。
多個就緒時，由 priority 機制挑選。

Zero-overhead Scheduling（零開銷排程）

GPU 把所有 resident warp 的執行狀態（PC、register）常駐於硬體暫存器，切換 warp 無需 save/restore，也不引入任何 idle cycle。
對比傳統 CPU context-switching：需把 outgoing thread 的 register 存到 memory、載入 incoming thread 狀態 → 有顯著開銷與閒置 cycle。

Tip

正因能用 warp 排程隱藏長延遲，GPU 不需像 CPU 投入大量晶片面積在 cache 與 branch prediction，可改用於 FP 執行單元與記憶體存取通道資源。

量化：oversubscription 倍率

oversubscription = (max threads/SM) / (cores/SM)

Ampere A100：64 cores/SM，但可駐 2048 threads/SM → 32× 超額訂閱。
駐留執行緒越多 → 越可能在某 warp 卡延遲時找到另一就緒 warp → latency tolerance 越有效（這也連結到 occupancy）。

考試/面試重點 (Exam / Test Patterns)

情境 / 關鍵字	答案 / 技巧
「一個 block 有幾個 warp？」	`⌈blockDim / 32⌉`；不足 32 的最後 warp 補 inactive thread
「warp n 含哪些 thread？」	thread `32·n` ~ `32·(n+1)−1`（依 row-major 線性 index）
「2D/3D block 怎麼切 warp？」	先 row-major 線性化（z→y→x），再每 32 個切一 warp
「SIMD efficiency 計算」	active threads / 32 × 100%（divergent 時某 pass 的有效比例）
「if (threadIdx.x > 2) 是否 divergence？」	會；條件基於 `threadIdx` → warp 內分裂成兩路徑，各跑一 pass
「`if(i<n)` boundary check 哪些 warp divergent？」	只有跨越資料邊界的那個 warp；vector 1003 + block 64 → 32 warps 中僅最後一個 divergent
「divergence 對效能影響大小？」	隨資料規模增大而下降（1/總warp 數）；大資料時可忽略
「Pascal vs Volta divergence」	Pascal 循序跑 passes；Volta+ 可交錯（Independent Thread Scheduling）
「為何 SM 要塞超量執行緒？」	latency tolerance：卡長延遲時換上就緒 warp，隱藏延遲
「Zero-overhead scheduling 為何零開銷？」	warp 狀態常駐硬體暫存器，切換不需 save/restore、無 idle cycle
「32 threads/block 可省略 `__syncthreads()`？」	不建議：warp 同步是硬體行為，正確性不可依賴；需要時應用 `__syncwarp()`（且 Volta 後 ITS 使此假設更不安全）
「為何 GPU cache/分支預測少？」	因用 warp 排程隱藏延遲，面積改投 FP 與記憶體通道