波前（Wavefront / Wave）
3 秒看懂
波前（Wavefront）是 GPU/加速器中一組被硬體捆綁、在同一時鐘週期內執行相同指令的執行緒集合。 NVIDIA 稱之為 Warp（32 執行緒），AMD 稱之為 Wavefront（GCN 時代 64 執行緒，RDNA 可 32 或 64，CDNA 僅 64）。它是一切 GPU 平行計算——包括 AI 訓練與推論——的最小排程粒度。理解 Wavefront，就是理解 GPU 吞吐的”齒輪”如何咬合。
3 分鐘產業解釋
為什麼 Wavefront 對 AI 產業鏈重要？
AI 大型模型訓練的本質是海量矩陣乘法 + 高頻寬訪存，而 GPU 的核心設計哲學是：用成千上萬個小執行緒的並行執行來掩蓋訪存延遲。這些執行緒並不是一個一個獨立排程的——那硬體開銷太大。GPU 的做法是把執行緒打包，以 Wavefront/Warp 為單位統一派發、統一執行、統一切換。
- 一個 Wavefront 就是一個”執行槽位”：SM（NVIDIA）/ CU（AMD）內的排程器每個週期決定”本週期執行哪個 Wavefront”。
- Wavefront 數量決定了”隱藏延遲的能力”：當一個 Wavefront 因記憶體訪問而停頓時，排程器可以瞬間切換到另一個就緒的 Wavefront，實現零開銷切換。這就是 latency hiding 的核心機制。
- Occupancy（佔用率）= 實際駐留 Wavefront 數 / 最大可駐留 Wavefront 數，是 GPU 效能調優的第一指標。
產業含義：
- 晶片設計：Wavefront 寬度和排程策略直接決定 GPU 架構的吞吐效率，是 NVIDIA vs AMD 架構差異的關鍵維度之一。
- AI 架構/編譯器：Triton、XLA、TVM 等編譯器後端需要按 Wavefront/Warp 粒度進行指令排程和記憶體合併（coalescing）最佳化。
- 運算元最佳化：CUDA kernel 中的
__shfl_sync、Warp-level reduction 等程式設計範式直接操作 Warp/Wavefront 內的執行緒通訊，是 Flash Attention、LayerNorm 等核心運算元的效能基石。 - 國產 GPU 生態：摩爾線程、壁仞、天數智芯等公司的架構設計中，如何定義和實現自己的 Wavefront/Warp 粒度，是衡量其 ISA 成熟度的關鍵標誌。
15 分鐘專家深入
Wavefront 的架構級地位
在現代 GPU 的執行層次中，Wavefront/Warp 處於執行緒 → Wavefront/Warp → Block → Grid 這一對映鏈的核心位置：
Grid (整個 Kernel)
└─ Block / Thread Block (軟體概念,分配到一個 SM/CU)
└─ Warp / Wavefront (硬體排程最小單位)
└─ Thread (單個執行緒,有自己的暫存器狀態)SIMT 執行模型：同一指令，不同資料
Wavefront 內所有執行緒在每個週期執行同一條指令（Single Instruction, Multiple Thread）。這與純 SIMD 不同——SIMT 允許執行緒有不同的控制流，但 divergent 分支需要序列化執行兩個路徑（通過 active mask 控制），效能代價是分支發散懲罰（branch divergence penalty）。
關鍵設計引數對比（已知公開資訊，未充分揭露處已標註）：
| 維度 | NVIDIA（Warp） | AMD GCN/RDNA/CDNA | Intel Xe（Sub-Slice） |
|---|---|---|---|
| 執行緒粒度 | 32 執行緒（自 Fermi 以來穩定） | GCN: 64; RDNA: 支援 wave32/wave64; CDNA: wave64 [AMD 架構白皮書] | 硬體 SIMD 寬度因代際而異 [未充分揭露統一命名] |
| 排程器 | Warp Scheduler（每 SM 多個） | Wavefront Scheduler（每 CU 一個或多個） | Thread Director 相關 [定性] |
| Warp 內通訊 | __shfl_sync() 系列 | ds_bpermute / ds_swizzle | 類似機制 [定性] |
| 分支處理 | Active mask 序列化 | Exec mask 序列化（64-bit） | 類似機制 [定性] |
為什麼 Warp=32 / Wavefront=64？
這不是隨意選擇——它是硬體資源分攤和延遲隱藏之間的平衡：
- 更寬的 Wavefront（如 64）：單個排程決策覆蓋更多執行緒，排程器簡單，但分支發散時浪費的執行緒槽位更多。
- 更窄的 Warp（如 32）：分支發散懲罰更小，wave32 模式下暫存器壓力可降低（因為不需要為 64 執行緒分配同一組暫存器 bank），但排程器負擔相對更高。
- AMD RDNA 引入 wave32 模式，一個核心動機就是圖形渲染中大量存在小三角形（small primitive），wave32 可減少執行緒浪費；AI 計算中矩陣運算通常高度一致，wave64 仍有優勢。
- NVIDIA 從 Fermi 到 Blackwell 一直維持 Warp=32，說明 32 執行緒在通用計算和 AI 負載中是足夠好的平衡點。
Wavefront 的生命週期
1. Block 被分配到 SM/CU
2. Block 內的執行緒被分組為若干 Wavefront/Warp
3. Wavefront 註冊到排程器的 wavefront pool / warp slot
4. 每週期，排程器從就緒（ready）佇列中選擇一個 Wavefront 發射
5. 若 Wavefront 遇到長延遲操作（如 global memory load），
標記為 stall，排程器立即切換到其他就緒 Wavefront
6. 資料到達後，該 Wavefront 重新變為 ready
7. 所有指令執行完畢，Wavefront 退出關鍵瓶頸——暫存器壓力與佔用率：
每個 Wavefront 中的每個執行緒都需要獨立的暫存器組。假設一個 SM 有 65536 個 32-bit 暫存器（NVIDIA Ampere 的公開資料），Warp=32 執行緒：
- 每執行緒使用 32 個暫存器 → 每 Warp 需 32×32 = 1024 個暫存器
- 最多駐留 65536 / 1024 = 64 個 Warp
- 每執行緒使用 128 個暫存器 → 每 Warp 需 32×128 = 4096 個暫存器
- 最多駐留 65536 / 4096 = 16 個 Warp
暫存器用得越多，能同時駐留的 Wavefront 越少，延遲隱藏能力越差。 這就是為什麼 kernel 最佳化中”暫存器 spill”和”佔用率”是永恆的權衡主題。
技術原理
1. SIMT 執行引擎的微架構
以 NVIDIA SM（Streaming Multiprocessor）為例，展示 Warp 在硬體中的執行路徑：
┌─────────────────────────────────────────────────────┐
│ SM (Streaming Multiprocessor) │
│ │
│ ┌──────────────┐ ┌──────────────┐ │
│ │ Warp Scheduler│ │ Warp Scheduler│ (Ampere: 4個) │
│ │ + Dispatch │ │ + Dispatch │ │
│ └──────┬───────┘ └──────┬───────┘ │
│ │ │ │
│ ▼ ▼ │
│ ┌──────────────────────────────────────┐ │
│ │ 執行單元 (Execution Units) │ │
│ │ ┌─────┐ ┌─────┐ ┌─────┐ ┌────────┐ │ │
│ │ │ FP32│ │ FP32│ │ INT │ │ SFU/LDX│ │ │
│ │ │ (16)│ │ (16)│ │ (16)│ │ │ │ │
│ │ └─────┘ └─────┘ └─────┘ └────────┘ │ │
│ │ 每週期: 4個 Warp 可併發發射 │ │
│ └──────────────────────────────────────┘ │
│ │
│ ┌──────────────────────────────────────┐ │
│ │ Register File (64K × 32-bit 典型值) │ │
│ │ → 所有活躍 Warp 共享物理暫存器池 │ │
│ └──────────────────────────────────────┘ │
│ │
│ ┌──────────┐ ┌──────────┐ │
│ │ Shared Mem│ │ L1 Cache │ │
│ │ (可配置) │ │ │ │
│ └──────────┘ └──────────┘ │
└─────────────────────────────────────────────────────┘關鍵機制說明：
- 多發射（Multi-Issue）：Ampere/Hopper 的 SM 中，一個 Warp 的 32 執行緒被拆成兩個 16 執行緒的 sub-warp（或稱 half-warp），分別在兩組執行單元上併發。這是”32”這個數字與硬體執行寬度匹配的結果。
- 運算元收集器（Operand Collector）：從暫存器檔案中為即將發射的 Warp 收集源運算元的硬體模組，需處理 bank conflict。
- Scoreboard：追蹤每個 Warp 中每條指令的依賴關係，決定哪些 Warp 處於 ready 狀態。
2. Wavefront 內的執行緒通訊
Wavefront/Warp 內的執行緒可以不經共享記憶體，直接通過 shuffle 指令交換資料：
// NVIDIA CUDA: Warp 內 shuffle
__shfl_sync(0xFFFFFFFF, val, src_lane); // 從指定 lane 讀取值
__shfl_down_sync(0xFFFFFFFF, val, delta); // 從 lane_id+delta 讀取
__shfl_xor_sync(0xFFFFFFFF, val, mask); // Butterfly 模式; AMD GCN: ds_bpermute
v_bpermute_b32 vdst, vsrc, vaddr ; 按 vaddr 指定的 lane 索引讀取應用場景：
- Warp-level reduction：LayerNorm、Softmax 中的求和/求最大值，一次 shuffle 即可完成 Warp 內 32/64 執行緒的歸約，無需 Shared Memory。
- Flash Attention：核心的 online softmax 更新中，Warp 間通過 Shared Memory +
__syncthreads()通訊，Warp 內通過 shuffle 通訊。 - Warp Matrix Multiply Accumulate (WMAA / WMMA)：Ampere 的
mma.sync指令讓整個 Warp 協作完成一個小矩陣乘（如 16×16×16），執行緒間通過暫存器檔案的隱式版面配置共享資料。
3. Branch Divergence 機制詳解
程式碼:
if (threadIdx.x % 2 == 0)
A(); // 偶數執行緒
else
B(); // 奇數執行緒
執行時間線 (Warp=8 為例):
Cycle 1: [exec_mask = 10101010] 執行 A() — 奇數執行緒空轉
Cycle 2: [exec_mask = 01010101] 執行 B() — 偶數執行緒空轉
→ 總耗時 = 2×，即"分支發散懲罰"為什麼對 AI 計算影響大？
- 大規模矩陣乘（GEMM）中幾乎沒有分支發散，Wavefront 效率極高（接近 100%）。
- 但 MoE 路由、動態稀疏計算、TopK 選擇 等操作會引入顯著的分支發散，降低 Wavefront 利用率。這也是硬體稀疏加速（如 NVIDIA 的結構化稀疏 2:4）試圖繞開的問題。
4. 記憶體合併（Memory Coalescing）與 Wavefront
Wavefront/Warp 內的 32/64 個執行緒如果訪問的全域性記憶體地址連續且對齊，硬體可以將它們合併為極少的記憶體事務：
理想情況（合併度=100%）:
執行緒 0 → addr[0], 執行緒 1 → addr[1], ..., 執行緒 31 → addr[31]
→ 合併為 1 次 128-byte 事務
最差情況（合併度=1/32）:
執行緒 0 → addr[0], 執行緒 1 → addr[1024], 執行緒 2 → addr[2048]...
→ 32 次獨立事務，頻寬利用率暴跌對 AI 的意義：Tensor Core 的輸入版面配置（如 HMMA 的 fragment layout）就是為了讓 Warp 內執行緒的暫存器訪問模式與矩陣的記憶體版面配置匹配，最大化合並效率。
技術演進史
| 時間 | 事件 | Wavefront/Warp 維度的變化 |
|---|---|---|
| 2006 | NVIDIA G80（Tesla 初代） | 引入 Warp=32 概念（前身是 NV4x 的 8-wide SIMD），確立 SIMT 模型 |
| 2007 | AMD/ATI R600（TeraScale） | 5-wide VLIW，每執行緒組 16 個執行緒並行，非現代 Wavefront 概念 |
| 2011 | AMD GCN（Southern Islands） | Wavefront=64 自 R600 時代延續，GCN 改進了排程和執行方式：4 個週期×16-wide SIMD 執行一個 Wavefront；同時釋出 HSA 異構計算規範 |
| 2012 | NVIDIA Kepler（GK110） | SMX 擁有 4 個 Warp Scheduler；支援 Warp Shuffle（__shfl），大幅減少 Shared Memory 使用 |
| 2014 | NVIDIA Maxwell | 進一步最佳化排程效率，L1/Shared Memory 重新劃分 |
| 2016 | NVIDIA Pascal (GP100) | 引入 FP16 半精度支援，Warp 級半精度吞吐翻倍；NVLink 提升多 GPU 通訊 |
| 2017 | NVIDIA Volta (GV100) | 獨立執行緒排程（Independent Thread Scheduling）：每個執行緒有獨立的 PC 和棧，但 Warp 仍然是執行粒度；引入 Tensor Core v1（Warp 協作完成 4×4×4 矩陣乘） |
| 2020 | AMD RDNA（RX 5000 系列） | 引入 wave32 模式，與 wave64 並存；CU 結構重新設計為 Work Group Processor (WGP) |
| 2020 | NVIDIA Ampere (A100) | Tensor Core v3，mma.sync 指令支援更大矩陣 tile（如 16×16×16）；結構化稀疏 2:4 對 Warp 級資料流提出新要求 |
| 2022 | AMD CDNA 2 (MI250X) | Wavefront=64 為預設，針對 HPC/AI 最佳化；引入 Matrix Core（對標 Tensor Core） |
| 2022 | NVIDIA Hopper (H100) | Warp Specialization：同一個 Block 內不同 Warp 可以執行不同角色（如 producer warp 負責資料搬運，consumer warp 負責計算），這是 TMA（Tensor Memory Accelerator）的軟體模型基礎；Tensor Core v4（FP8 支援） |
| 2023 | AMD CDNA 3 (MI300X) | Wavefront=64 繼續；Matrix Core 擴充套件支援 FP8/INT8；MI300X 為純 GPU 加速卡，不支援 APU 模式（統一記憶體架構和 APU 模式為 MI300A 特性；MI300X 可通過 CXL 等支援統一記憶體訪問機制） |
| 2024 | NVIDIA Blackwell (B100/B200) | Tensor Core 進一步擴充套件到 FP4/FP6；Warp 級協作機制未見公開架構級變化 [基於已公開資訊] |
| 2024 | Intel Gaudi 3 | 非傳統 GPU，使用 VPU（Vector Processing Unit）+ TPC（Tensor Processing Core）架構，排程粒度概念不同於 Warp/Wavefront [定性] |
趨勢總結：
- Wavefront/Warp 寬度穩定：NVIDIA 堅守 32 超過 15 年，AMD 在 AI/HPC 場景維持 64、圖形場景增加 wave32 靈活性。
- 執行語義越來越”粗”：從逐條標量指令，到 Warp-level 矩陣指令（
mma.sync），再到 Block-level 非同步操作（TMA +wgmma），排程粒度逐級上移。 - Warp Specialization 成趨勢：Hopper 的 producer/consumer warp 模型讓一個 Block 內不同 Warp 承擔不同角色，打破了”所有 Warp 做同一件事”的傳統假設。
技術路線對比（量化表）
| 指標 | NVIDIA Warp (Ampere A100) | NVIDIA Warp (Hopper H100) | AMD Wavefront (CDNA2 MI250X) | AMD Wavefront (CDNA3 MI300X) |
|---|---|---|---|---|
| 執行緒粒度 | 32 | 32 | 64 | 64 |
| 每 SM/CU 最大駐留 Warp/Wave 數 | 64 warps [NVIDIA 白皮書] | 64+（支援更多 concurrent warp 與 TMA 非同步）[NVIDIA 白皮書] | 每 CU 最多 40 wavefronts [AMD CDNA2 ISA] | 40 wavefronts per CU [AMD CDNA3 公開資料，待確認] |
| 每 SM/CU 暫存器檔案 | 65536 × 32-bit | 65536 × 32-bit [NVIDIA 白皮書] | 每 CU: 65536 × 32-bit [AMD 架構手冊] | 65536 × 32-bit [供應鏈估算] |
| Tensor/Matrix Core Warp 協作 | mma.sync m16n8k16（FP16） | mma.sync + wgmma（Block 級） | MFMA 32×32×8（FP16） | MFMA 擴充套件至 FP8 [AMD 公開資料] |
| Warp Shuffle 頻寬 | 32 執行緒 × 4B = 128B/週期 | 同上 | 64 執行緒 × 4B = 256B/週期 [AMD ISA] | 同 CDNA2 [待確認] |
| Warp Specialization | 不支援（Block 內 Warp 同質） | 支援（producer/consumer warp） | 不支援（截至目前公開資訊） | 不支援 [截至目前公開資訊] |
| 結構化稀疏 | 2:4（Warp 級後設資料） | 2:4 | 不適用 | 不適用 |
注意：上表部分資料來源於廠商公開白皮書和 ISA 手冊，部分為行業通用認知，具體數字可能因型號配置（如 A100 40GB vs 80GB 的 SM 數差異）而不同。所有數值請以廠商官方文件為準。
上下游
上游（影響 Wavefront 設計的因素）
┌─────────────────────────────────────────────────────────────┐
│ 製程工藝 (5nm/4nm/3nm) │
│ → 決定 SM/CU 內可整合多少執行單元和暫存器 │
├─────────────────────────────────────────────────────────────┤
│ 記憶體子系統 (HBM3/HBM3E, GDDR7) │
│ → 決定 Wavefront 訪存延遲，影響需多少 Wave 做 latency hiding │
├─────────────────────────────────────────────────────────────┤
│ 編譯器 (NVCC, ROCm/hipcc, Triton, XLA) │
│ → 決定 Warp/Wave 的指令排程和暫存器分配 │
├─────────────────────────────────────────────────────────────┤
│ AI 架構 (PyTorch, JAX, TensorFlow) │
│ → 運算元實現的 Warp 級最佳化程度 │
└─────────────────────────────────────────────────────────────┘下游（Wavefront 效率影響的層面）
┌─────────────────────────────────────────────────────────────┐
│ Kernel 效能 (單運算元吞吐) │
│ → 佔用率、分支發散、記憶體合併直接影響 │
├─────────────────────────────────────────────────────────────┤
│ 大型模型訓練吞吐 (Tokens/sec) │
│ → 所有 GEMM/Attention kernel 的 Wavefront 利用率疊加 │
├─────────────────────────────────────────────────────────────┤
│ 推論延遲 (Time-to-First-Token, Inter-Token Latency) │
│ → batch size 小時佔用率低，Wavefront 利用率是關鍵瓶頸 │
├─────────────────────────────────────────────────────────────┤
│ 多 GPU 並行效率 │
│ → Tensor/Pipeline Parallelism 中每個 GPU 內部的 Wavefront 效率│
│ 決定了通訊/計算重疊（overlap）的效果 │
└─────────────────────────────────────────────────────────────┘關鍵指標
| 指標 | 定義 | 典型優秀值 | 對 AI 的意義 |
|---|---|---|---|
| Occupancy（佔用率） | 活躍 Warp 數 / 最大 Warp 數 | ≥50%（視 kernel 而定） | 太低→延遲隱藏不足→SM 利用率低 |
| Warp Execution Efficiency | 非 divergent 週期 / 總週期 | ≥90%（GEMM kernel） | 分支發散直接降低 |
| Memory Coalescing Rate | 有效位元組 / 事務位元組 | ≥80% | 決定視訊記憶體頻寬實際利用率 |
| Warp Stall Reasons | 各類停頓原因佔比 | 應以 MIO/Long Scoreboard 為主（說明在等計算或視訊記憶體） | ncu/rocprof 核心診斷指標 |
| Register Per Thread | 每執行緒暫存器使用數 | 32–128（視 kernel） | 決定最大駐留 Warp 數 |
| Shared Memory Per Block | 每 Block 共享記憶體 | 決定 Block 級 Wavefront 可用的快速通訊資源 | 影響 Wave 間同步開銷 |
| Tensor Core Utilization | Tensor Core 活躍週期佔比 | ≥60%（成熟 kernel） | Wavefront 級矩陣指令的有效發射率 |
供需與市場資料
Wavefront/Warp 不是一個獨立的產品市場，而是 GPU 架構的內在機制。其商業影響通過以下鏈條傳導：
需求側
- AI 訓練晶片：全球 AI 訓練 GPU 市場 2024 年預估超 $80B [綜合多家行業報告，口徑差異較大]。每顆 GPU 的核心競爭力之一是其 Wavefront/Warp 排程效率對 kernel 吞吐的貢獻。
- 推論晶片：推論場景 batch size 小，佔用率普遍較低，Wavefront 利用率是推論晶片需要特別最佳化的維度。
- 編譯器/工具鏈生態：Triton、XLA、TVM 等中間表示層需要深度理解目標硬體的 Wavefront/Warp 語義，這直接決定了這些工具的實際效能上限。
供給側
- NVIDIA：CUDA 生態對 Warp 語義的抽象和最佳化已積累 15+ 年，
ncu（Nsight Compute）提供完整的 Warp 級分析指標。 - AMD：ROCm 生態正在快速補齊，
rocprof+ OmniPerf 提供 Wavefront 級分析，但工具鏈成熟度仍有差距（截至 2024 年行業共識）。 - 國產 GPU：多數國產 AI 晶片廠商借鑑 NVIDIA 的 Warp=32 模型，但 ISA 細節（如 shuffle 指令、Tensor Core/Wave 協作模式）各有差異，生態相容性是核心挑戰。
代表公司與資本對映
| 公司 | 與 Wavefront 的關聯 | 資本標的 |
|---|---|---|
| NVIDIA | Warp=32 SIMT 模型的定義者；從 Fermi 到 Blackwell 持續迭代 | NVDA |
| AMD | Wavefront=64（GCN/CDNA）/ wave32+64（RDNA）雙模式；MI300X 為 HBM3 APUs 旗艦 | AMD |
| Intel | Xe 架構的執行緒排程粒度設計；Gaudi 3 的 TPC/VPU 架構對 Warp 概念的替代方案 | INTC |
| 摩爾線程 | MUSA 架構借鑑 Warp 模型，具體粒度未完全公開 [未充分揭露] | 一級市場 |
| 壁仞科技 | BR100 採用自研 SIMT 架構，Wave/Warp 粒度設計是 ISA 核心 | 一級市場 |
| 天數智芯 | 天垓系列 GPU，SIMT 執行模型 [定性] | 一級市場 |
| 寒武紀 | MLU 架構的並行排程粒度不同於傳統 GPU Wavefront，但核心思想類似 | 688256.SH |
| Triton（OpenAI） | 編譯器後端需要為每個目標硬體生成 Wavefront/Warp 級最佳化程式碼 | 未上市 |
| TVM / Apache | 自動調優（AutoTVM）的搜尋空間包含 Wavefront/Warp 相關引數 | 開源 |
投資邏輯
直接邏輯
- 架構差異化壁壘：Wavefront/Warp 排程器是 GPU 微架構中最難被逆向工程的部分之一。它不是簡單的寬度數字，而是包括排程策略、依賴追蹤、暫存器分配、分支處理的完整狀態機。這是 NVIDIA 和 AMD 的核心 IP 護城河。
- 編譯器生態繫結：CUDA 的
__shfl_sync、mma.sync等 Warp 級 API 被數十萬個 kernel 使用，遷移成本極高。這構成了 NVIDIA 的軟體 lock-in。 - 國產替代維度：評估國產 GPU 不僅要看 Tensor Core 的 peak TFLOPS，更要看其 Warp 排程器的實際效率（occupancy 上限、stall pattern、shuffle 延遲）。這些是紙面引數看不到、但實測效能差異巨大的因素。
間接邏輯
- **HBM 頻寬與