波前（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 性能调优的第一指标。

产业含义：

1. 芯片设计：Wavefront 宽度和调度策略直接决定 GPU 架构的吞吐效率，是 NVIDIA vs AMD 架构差异的关键维度之一。
2. AI 框架/编译器：Triton、XLA、TVM 等编译器后端需要按 Wavefront/Warp 粒度进行指令调度和内存合并（coalescing）优化。
3. 算子优化：CUDA kernel 中的 __shfl_sync、Warp-level reduction 等编程范式直接操作 Warp/Wavefront 内的线程通信，是 Flash Attention、LayerNorm 等核心算子的性能基石。
4. 国产 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）。

关键设计参数对比（已知公开信息，未充分披露处已标注）：

为什么 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 宽度稳定：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 做同一件事”的传统假设。

技术路线对比（量化表）

注意：上表部分数据来源于厂商公开白皮书和 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）的效果                            │
└─────────────────────────────────────────────────────────────┘

关键指标

供需与市场数据

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 协作模式）各有差异，生态兼容性是核心挑战。

代表公司与资本映射

投资逻辑

直接逻辑

1. 架构差异化壁垒：Wavefront/Warp 调度器是 GPU 微架构中最难被逆向工程的部分之一。它不是简单的宽度数字，而是包括调度策略、依赖追踪、寄存器分配、分支处理的完整状态机。这是 NVIDIA 和 AMD 的核心 IP 护城河。
2. 编译器生态绑定：CUDA 的 __shfl_sync、mma.sync 等 Warp 级 API 被数十万个 kernel 使用，迁移成本极高。这构成了 NVIDIA 的软件 lock-in。
3. 国产替代维度：评估国产 GPU 不仅要看 Tensor Core 的 peak TFLOPS，更要看其 Warp 调度器的实际效率（occupancy 上限、stall pattern、shuffle 延迟）。这些是纸面参数看不到、但实测性能差异巨大的因素。

间接逻辑

4. **HBM 带宽与