Warp 排程器 (Warp Scheduler)

3 秒看懂

一句話：Warp 排程器是 GPU 流式多處理器 (SM) 內部的”指揮棒”——每個時鐘週期從數十個就緒的執行緒束 (Warp) 中挑一個下達執行，通過高頻切換隱藏訪存延遲，讓數千執行緒”看起來”同時跑滿吞吐。

核心機制：延遲隱藏 (Latency Hiding) + 指令級並行排程 (ILP Scheduling)。

3 分鐘產業解釋

為什麼 GPU 能同時跑幾千個執行緒卻不用幾千個核？

答案藏在 Warp 排程器 裡。

CPU 設計哲學是”減少每個執行緒的延遲”——深流水線、大快取、亂序執行。GPU 設計哲學截然相反：用海量執行緒掩蓋單個執行緒的慢。一個 SM 可能只有 46 個執行流水線，但同時”持有” 4864 個 Warp（每個 Warp 32 執行緒 = 1536~2048 個執行緒上下文）。當某個 Warp 在等記憶體返回資料時，Warp 排程器在納秒級切換到另一個就緒 Warp——CPU 花在等記憶體的時間，GPU 花在了算別的執行緒上。

這就是為什麼 GPU 對”頻寬敏感型”和”高並行度”負載特別擅長，也解釋了為什麼 NVIDIA 從 Fermi 時代起就把多 Warp 排程器作為架構賣點。

對投資者意味著什麼？

• 架構創新 = 效能/功耗比提升：更聰明的排程策略意味著同面積矽片能榨出更多算力，這是每代 GPU IPC 提升的核心驅動力之一。
• 軟體生態繫結：排程策略與 CUDA 程式設計模型深度耦合，形成極高遷移成本，是 NVIDIA 護城河的技術底層。
• 瓶頸轉移訊號：當排程器效率接近理論上限，瓶頸回到記憶體頻寬（→ HBM 漲價邏輯）或互聯頻寬（→ NVLink 重要性）。

15 分鐘專家深入

1. Warp 排程的完整決策鏈

┌──────────────────────────────────────────────────────┐
│                    SM 內部檢視                         │
│                                                       │
│  ┌─────────┐  ┌─────────┐  ┌─────────┐              │
│  │ Warp 0  │  │ Warp 1  │  │ Warp N  │  (48~128個)  │
│  │ 狀態機  │  │ 狀態機  │  │ 狀態機  │              │
│  └────┬────┘  └────┬────┘  └────┬────┘              │
│       │            │            │                     │
│       ▼            ▼            ▼                     │
│  ┌─────────────────────────────────────┐             │
│  │       Warp Scheduler (1~4個)        │             │
│  │  ┌─────────────────────────────┐    │             │
│  │  │ Scoreboard: 每條in-flight   │    │             │
│  │  │ 指令的暫存器依賴/就緒狀態    │    │             │
│  │  └─────────────────────────────┘    │             │
│  │  選擇策略: 就緒佇列中pick一個Warp   │             │
│  └──────────────┬──────────────────────┘             │
│                 │                                     │
│                 ▼                                     │
│  ┌─────────────────────────────────────┐             │
│  │       Issue Slot / Dispatch         │             │
│  │  ┌────┐ ┌────┐ ┌────┐ ┌────┐       │             │
│  │  │FP32│ │FP32│ │INT │ │LSU │       │             │
│  │  │pipe│ │pipe│ │pipe│ │pipe│       │             │
│  │  └────┘ └────┘ └────┘ └────┘       │             │
│  └─────────────────────────────────────┘             │
└──────────────────────────────────────────────────────┘

關鍵步驟：

步驟	動作	延遲貢獻
① 讀取 Warp 狀態	檢查 scoreboard 中每條 in-flight 指令的就緒位	0 cycle (組合邏輯)
② 選擇候選 Warp	過濾出”所有源運算元已就緒”且”無 structural hazard”的 Warp	0~1 cycle
③ Pick 策略	從候選佇列中按策略選 1 個（或多個，取決於排程器數量）	0 cycle
④ 譯碼 & 發射	向對應執行單元發射指令	1 cycle

2. 延遲隱藏的數學直覺

假設某 Warp 執行一條指令的 流水線延遲 為 L 個 cycle（例如：FP32 算術 ~6 cycle，L2 命中訪存 ~100-300+ cycle，HBM 訪存 ~200+ cycle [估算，具體取決於代際和命中率]），要完全隱藏這個延遲，需要：

就緒 Warp 數 ≥ L

實際約束：

• SM 內 Warp 總數受暫存器檔案大小限制
• 每個 Warp 的暫存器佔用 = 每執行緒暫存器數 × Warp 寬度 (32)
• 編譯器通過 maxrregcount 或 occupancy tuning 控制

Occupancy（佔用率） 是衡量”排程空間”的核心指標：

Occupancy = Active Warps / Max Supported Warps

Occupancy 不是越高越好——過度增加活躍 Warp 會擠壓每執行緒暫存器，導致 register spilling 到 local memory，反而增大延遲。存在最優點。

3. 排程策略演進

① LRR (Loose Round-Robin)：

• 最簡單，輪詢選就緒 Warp
• 優點：公平、實現簡單
• 缺點：不利用 ILP（指令級並行），可能選到有長依賴鏈的 Warp

② GTO (Greedy-Then-Oldest)：

• 優先給”最近被選中的 Warp”連續發射機會（Greedy），直到它 stall 才切換
• 切換時選最老的就緒 Warp（Oldest）
• 優點：同一 Warp 連續發射能利用 ILP，減少暫存器埠壓力
• 缺點：可能餓死其他 Warp

③ TLP-first vs ILP-first：

• 執行緒級並行優先（多 Warp 切換）vs 指令級並行優先（深挖單 Warp 的 ILP）
• 不同負載偏好不同：記憶體密集型偏好 TLP，計算密集型偏好 ILP
• 現代架構通常在硬體中做混合策略

[注]：NVIDIA 未公開揭露其生產 GPU 使用的具體排程演算法細節，上述策略來自學術研究與逆向工程推斷 [來源：公開學術論文]。

---

技術原理

4. Scoreboard 機制詳解

Scoreboard 是 Warp 排程器的”賬本”，追蹤每條 in-flight 指令對暫存器的寫依賴。

┌─────────────────────────────────────────────┐
│            Scoreboard (per Warp)             │
├──────────┬───────────┬──────────────────────┤
│ 暫存器號  │ 狀態位    │ 等待的指令/功能單元   │
├──────────┼───────────┼──────────────────────┤
│ R0       │ Ready     │ -                    │
│ R1       │ Pending   │ SFU unit, ~20 cycles │
│ R2       │ Ready     │ -                    │
│ R3       │ Pending   │ Memory, ~100 cycles  │
│ ...      │ ...       │ ...                  │
└──────────┴───────────┴──────────────────────┘

排程器決策邏輯：
  IF warp[i].src_operands ALL == Ready
     AND 所需執行單元無 structural hazard
  THEN warp[i] 進入候選佇列

當一條指令的所有源運算元都 Ready 時，該 Warp 才有資格被選中。指令發射後，目標暫存器立即被標記為 Pending，直到執行單元寫回。

5. Warp Divergence（分支發散）

SIMT（單指令多執行緒）模型要求同一 Warp 內 32 個執行緒執行相同指令。遇到條件分支時：

// 32 個執行緒，部分滿足條件，部分不滿足
if (threadIdx.x % 2 == 0) {
    path_A();  // 16 個執行緒
} else {
    path_B();  // 16 個執行緒
}

硬體執行方式：

1. 先執行 path_A，path_B 執行緒被 mask 掉（不寫回）
2. 再執行 path_B，path_A 執行緒被 mask 掉
3. 兩條路徑序列化，總時間 = T(A) + T(B)

獨立執行緒排程與SIMT Stack：Volta+ 架構引入獨立執行緒排程，每個執行緒有自己的程式計數器和呼叫棧，允許更細粒度的分支管理與重新收斂機制，但 Warp 排程仍以 Warp 粒度選擇執行。

對排程器的影響：發散的 Warp 兩次發射之間 mask 變化頻繁，增加了排程複雜度。

6. 執行埠與 Structural Hazard

Warp 排程器不僅看 Warp 是否”就緒”，還要看執行單元是否有空位：

執行單元型別	典型埠數（per SM，依架構而異）	流水線深度（估算）
FP32 ALU	2~4	~6 cycles
FP64 ALU	0~2 (消費級少, 計算卡多)	630 cycles
INT32 ALU	1~2	~6 cycles
SFU (sin/cos/rsqrt)	1~2	~20 cycles
LSU (Load/Store)	1~4	取決於命中級別
Tensor Core	依 SM 子分割槽 (Sub-partition)	~200-300 cycles [指令延遲估算]

當多個就緒 Warp 爭搶同一埠時，產生 structural hazard，排程器需要仲裁。

---

技術演進史

代際	年份	關鍵變化	Warp 排程器相關創新
Tesla (G80)	2006	首個通用計算 GPU	單 Warp 排程器，Warp 寬度 32
Fermi	2010	首代真正計算架構	雙 Warp 排程器（每 SM 2 個），每週期可發射 2 條獨立指令
Kepler	2012	SMX 擴充套件	4 個 Warp 排程器，SM 內執行緒數大幅增加
Maxwell	2014	效率優先	排程邏輯最佳化，同功耗下效能提升
Pascal	2016	HBM2 引入	排程器針對高頻寬記憶體最佳化（更積極預取）
Volta	2017	Tensor Core 誕生	獨立執行緒排程 (Independent Thread Scheduling)，引入 Thread Convergence 硬體支援；Tensor Core 指令排程新增
Turing	2018	RT Core	排程器需管理 3 種執行單元（CUDA Core / Tensor Core / RT Core）
Ampere	2020	Sparsity 支援	結構化稀疏 2:4 的排程路徑；非同步複製 (async copy) 改變訪存排程模型
Hopper	2022	非同步執行引擎	TMA (Tensor Memory Accelerator)、非同步 warp-specialized 模式，Warp 可被程式設計為 producer/consumer 角色，排程策略從硬體通用轉向軟體可程式設計
Blackwell	2024	雙 die 封裝	跨 die 排程協調需求增加

趨勢總結：

• 早期：硬體黑盒排程，簡單輪詢/貪心策略
• 中期：更多執行單元型別 → 排程器需管理更多 structural hazard
• 近期：軟體參與度提升（async copy、warp specialization），排程從”純硬體決策”向”硬體+軟體協同”演進

---

技術路線對比

NVIDIA vs AMD 排程模型

維度	NVIDIA (Warp)	AMD (Wavefront)
執行緒束寬度	32	64 (GCN/RDNA)，RDNA1起引入 Wave32 模式
排程粒度	Warp (32 執行緒)	Wavefront (64 或 32 執行緒)
SIMD 寬度	32 (per sub-core)	32 (SIMD32，一個 wavefront 分兩個 cycle 執行)
排程器數量 (per CU/SM)	依架構：1~4	GCN: 1 個非同步計算引擎 + 4 個 SIMD；RDNA: 雙 issue
分支處理	條件碼 mask + 序列化	EXEC 暫存器 mask + 序列化 (64 bit mask)
軟體排程參與度	低 (傳統)→高 (Hopper async warp)	中 (LDS/scratch memory 管理)

Wave32 vs Wave64：

• RDNA 架構引入 Wave32 模式，與 NVIDIA Warp 對齊
• Wave32：更好匹配較短 SIMD 流水線，減少 divergence 損失
• Wave64：更大粒度 → 暫存器利用率可能更高，但 divergence 成本更大

---

上下游

上游：誰給 Warp 排程器提供”原材料”

上游環節	內容	關鍵約束
編譯器 (nvcc/ptxas)	暫存器分配、指令排程、occupancy 決策	直接決定活躍 Warp 數和暫存器壓力
驅動/執行時	Grid/Block 配置、Stream 管理	Block 到 SM 的對映影響排程空間
暫存器檔案 (Register File)	儲存每個執行緒的上下文	總大小限制活躍 Warp 數
共享記憶體 (Shared Memory)	Block 內同步/通訊	大 Shared Memory 佔用 → SM 能容納的 Block 減少
L1 快取/紋理單元	資料供給速度	快取命中率直接影響 Warp stall 時長

下游：Warp 排程器的輸出去哪

下游環節	內容
執行流水線 (CUDA Core/Tensor Core/SFU/LSU)	被選中的指令發射到對應執行單元
記憶體子系統	Load/Store 指令觸發 → L1/L2/HBM 訪問 → 結果回來後喚醒等待 Warp
Scoreboard 更新	指令完成 → 更新就緒位 → Warp 重回候選佇列

---

關鍵指標

指標	定義	意義
Occupancy	活躍 Warps / 最大支援 Warps	衡量延遲隱藏空間；太低→stall 嚴重，太高→暫存器溢位
Warp Occupancy (per SM)	同時駐留的 Warp 數	受暫存器檔案大小、Shared Memory、Block 限制
IPC (Instructions Per Cycle)	SM 每週期平均發射指令數	理論上限 = 排程器數量 × 每排程器發射數
Stall 分類	暫存器依賴、記憶體依賴、執行依賴、同步等待	幫助判斷瓶頸在計算還是記憶體
Warp Stall Reasons	__cudaProfiler / Nsight Compute 可採集	工程最佳化的核心診斷手段
指令吞吐 (FLOPS)	峰值 FLOPS = 核心數 × 頻率 × 2 (FMA)	排程器效率影響實際可達比例

Occupancy 計算示例（概念性，非具體架構）：

假設:
  - SM 有 65536 個 32-bit 暫存器
  - 每執行緒使用 64 個暫存器
  - Warp 寬度 = 32

每個 Warp 消耗暫存器 = 64 × 32 = 2048
最大 Warp 數 = 65536 / 2048 = 32

如果還有 Shared Memory 和 Block 維度約束,
實際最大可能更低。

---

供需與市場資料

為什麼這個微觀架構概念影響宏觀市場？

① 排程效率 → 有效算力密度

場景	對排程器的需求
訓練大型模型 (記憶體密集)	需高 Occupancy 隱藏 HBM 延遲 → 暫存器不能太重
推論 (計算密集+低 batch)	Occupancy 可能不高，但需高效排程避免空轉
渲染 (高度發散)	Warp divergence 嚴重 → 排程器需快速 mask 切換
HPC (規整計算)	排程效率影響小，瓶頸在互聯/記憶體

② 對 GPU 市場的影響鏈：

更高效排程策略
    → 同面積更多有效算力 (FLOPS 利用率↑)
    → 單卡效能↑ 或 單效能成本↓
    → 影響採購量 × 單價 → 影響 NVIDIA/AMD 資料中心營收

③ 關鍵市場資料點（來源：公開財報/行業報告，非排程器專屬資料）：

• NVIDIA 資料中心業務 FY2024 營收超 $470 億 [來源：NVIDIA 財報]
• GPU 算力利用率（FLOPS utilization）在最佳化良好的訓練負載中通常在 40%~65% [行業估算，具體取決於模型和最佳化程度]
• 排程效率是決定”實際利用率”的核心因素之一

---

代表公司與資本對映

硬體層面

公司	與 Warp 排程的關係	投資標的/關注點
NVIDIA	CUDA Core 和 Tensor Core 的排程策略是其架構競爭力核心	NVDA (NASDAQ)
AMD	RDNA (消費) / CDNA (計算) 的 wavefront 排程	AMD (NASDAQ)
Intel	Xe/HPC 架構的 EU 排程模型 (類似 Warp 的 thread group)	INTC (NASDAQ)
華為海思	達芬奇架構的 Cube/Vector 排程	非上市 / 關注昇騰生態
寒武紀	MLU 架構的 warp 級排程設計	688256.SH

軟體/工具層面

工具/公司	關聯
Nsight Compute (NVIDIA)	可採集 Warp Stall Reasons、Occupancy 詳細資料
ROCm rocprofiler (AMD)	AMD 對應的效能分析工具
Triton (OpenAI)	編譯器層面隱式管理排程策略
CUDA Profiler API	應用級排程最佳化介面

---

投資邏輯

從 Warp 排程器看 GPU 架構投資的 3 層邏輯

第 1 層：架構創新週期

每代 GPU 架構升級的 IPC 提升中，相當比例來自排程策略最佳化（而非單純堆核心數）。關注 NVIDIA/AMD 的架構日程：

• NVIDIA：Hopper → Blackwell → 下一代（Feynman?）
• AMD：CDNA 3 → CDNA “Next”

排程策略的重大創新（如 Hopper 的 async warp specialization）往往預示著軟體生態需要重新適配，這也意味著新架構的紅利期和學習曲線。

第 2 層：利用率競爭

在 AI 訓練/推論的 TCO 競賽中：

實際每美元效能 = 峰值 FLOPS × 排程效率 × 互聯效率 / 總擁有成本

排程效率（40%~65%）是離”理論峰值”最遠的因素之一，提升空間大，且不依賴製程節點突破。誰能在排程策略上做得更好，誰就能在同製程下交付更高有效算力。

第 3 層：程式設計模型繫結

Warp 排程策略與 CUDA 程式設計範式深度耦合（Grid/Block/Warp/Thread 層級）。這意味著：

• 開發者心智模型被鎖定
• 競爭對手需要同時在硬體排程器和軟體生態上追趕
• 這種”認知稅”是比硬體本身更深的護城河

---

常見誤讀糾偏

❌ 誤讀 1：“更多 Warp 排程器 = 更好效能”

糾偏：排程器數量受限於執行埠數量。4 個排程器 × 4 個 FP32 埠 = 每週期理論 16 次發射，但如果 SM 內只有 4 個 FP32 流水線，4 個排程器中只有能有效分配的才有意義。盲目增加排程器數量會增加面積和功耗，但可能因為結構性瓶頸（structural hazard）而無法提升實際吞吐。

Kepler SMX 確實有 4 個排程器，但其後續架構（Maxwell/Pascal）又調整了排程器數量與執行埠的配比，說明這不是簡單”越多越好”的線性關係。

❌ 誤讀 2：“Occupancy 100% 就是最佳效能”

糾偏：Occupancy 越高，意味著每個執行緒分到的暫存器越少。當暫存器不夠時，編譯器會 spill 到 local memory（實際落在 L1/L2 或視訊記憶體），這會產生遠大於 Warp 切換能隱藏的額外延遲。

經驗法則：Occupancy 與效能的關係是非線性的，存在最優點。NVIDIA 的 Nsight Compute 工具提供 “Occupancy vs Performance” 分析，幫助找到該最優點。某些高效能 kernel 反而以較低 Occupancy 執行（因為每執行緒需要大量暫存器來儲存中間結果以減少訪存）。

❌ 誤讀 3：“Warp 排程器負責執行緒的並行化”

糾偏：並行化的”粒度分配”由程式設計師/編譯器決定（Grid 維度、Block 維度）。Warp 排程器只負責在已分配給 SM 的 Warp 中做分時多工排程——它不能創造更多執行緒，只能更高效地利用已有執行緒。真正的並行性來自 SIMT 模型本身（同一 Warp 內 32 執行緒天然並行執行同一指令）。

---

學習路徑

入門 → 進階 → 專家

Level 1: 概念建立
├── NVIDIA CUDA C Programming Guide (官方文件，免費)
│   └── 重點章節: Hardware Implementation, SIMT Architecture
├── "Programming Massively Parallel Processors" (Kirk & Hwu)
│   └── 第 3~4 章: 執行緒組織與執行模型
│
Level 2: 深入理解
├── NVIDIA GTC Talks (免費)
│   └── 搜尋 "CUDA Performance" / "Warp Scheduling"
├── Nsight Compute 實操
│   └── 採集 Stall Reasons，理解 Occupancy vs Performance 曲線
├── "GPU Gems" 系列 (NVIDIA, 免費線上)
│
Level 3: 架構研究
├── 論文: "Demystifying GPU Microarchitecture via Microbenchmarking"
│   └── (Ardestani & Renau, 學術論文)
├── 論文: "Anatomy of High-Performance Many-Threaded Matrix Multiplication"
│   └── (Nath et al., 關注 Tensor Core 排程)
├── 各代架構 Whitepaper (NVIDIA Developer Zone)
│   └── Volta/Turing/Ampere/Hopper/Blackwell 架構白皮書
│
Level 4: 前沿追蹤
├── NVIDIA GTC 年度 Keynote + Architecture Session
├── ISCA / MICRO / HPCA 會議論文 (體系結構頂會)
└── HOTCHIPS 會議演講 (廠商微架構揭露)

推薦實操

1. 寫一個 Occupancy-bound kernel：故意用大量暫存器，觀察效能下降
2. 用 Nsight Compute 分析：檢視 Warp Stall Reasons breakdown
3. 調優 maxrregcount：觀察 Occupancy 與效能的非線性關係
4. 對比 Wave32 vs Wave64（如果用 AMD GPU）：觀察分支發散場景的差異

---

一句話總結

Warp 排程器是 GPU “用延遲換吞吐” 哲學的硬體執行者：它不創造並行性，而是通過納秒級執行緒切換，將 SM 內有限的執行埠與海量執行緒上下文編織成一條高效的指令流水線——理解它，就理解了 GPU 效能最佳化的核心槓桿點。

---

延伸閱讀與來源

權威一手資料

來源	連結/位置	說明
NVIDIA CUDA Programming Guide	docs.nvidia.com/cuda	官方 SIMT/Warp 機制文件
NVIDIA 架構白皮書	NVIDIA Developer Zone 各代 Whitepaper	微架構級細節（Volta+ 包含排程資訊）
Nsight Compute 文件	NVIDIA Developer Zone	Stall Reasons、Occupancy 分析工具

學術論文

• 作者不詳, “Demystifying GPU Microarchitecture via Microbenchmarking” — GPU 微架構逆向研究經典
• 各代架構相關 ISCA/MICRO 論文 — 關注 scheduling policy 分析

資料來源宣告

• 本文中的流水線深度、執行埠數量等具體微架構數字：基於公開學術論文和逆向工程推斷，未充分揭露時已標註 [估算]
• 市場營收資料：來源為公司公開財報，已標註
• NVIDIA/AMD 未公開揭露其生產 GPU 中 Warp 排程器使用的具體演算法實現

---

最後更新：2025 | 概念學習頁系列