MI300X · CDNA3 · GFX942 · 特性专题
Wave64 · VGPR/SGPR 与占用度
★ CDNA3 调度主线 · 64 线程 wave 固定 · 标量/向量寄存器分家 · wave 槽位占用度
是什么
DEFINITIONCDNA 的调度单位是 wavefront（wave）= 64 线程，固定不可选（wave32 是 RDNA 游戏卡的特性，gfx942 没有）[2]。一个 workgroup 最多 16 waves = 1024 线程，每 CU 同时驻留 32 waves（4 SIMD × 8 wave 槽）——满占用恰好也是 2048 线程，与 H200 SM 相同。
寄存器体系与 CUDA 根本不同：VGPR（向量寄存器，每 lane 独立值，512KB/CU）与 SGPR（标量寄存器，整 wave 共享一个值，16–102/wave）是两个物理文件；矩阵累加再叠一层 AccVGPR（≤256/wave）[2]。
为什么（vs CUDA warp 模型）
RATIONALE · VS WARP(32)| 维度 | CUDA（H200） | CDNA3（MI300X） |
|---|---|---|
| 调度单位 | warp = 32 线程 | wave = 64 线程，固定 |
| 满占用线程 | 2,048/SM（64 warps） | 2,048/CU（32 waves）——数值巧合地相同 |
| 寄存器文件 | 统一 RF 256KB，255/thread | VGPR 512KB + SGPR 分离，512/wave（含 AccVGPR） |
| 分配粒度 | 按 warp，8 regs 步进 | VGPR 按 8 / SGPR 按 16 步进，按 wave 整体申请 |
| uniform 数据 | 无独立文件（LDU 路径弱化） | SGPR 专用：地址/循环变量/uniform 分支全走标量域 |
| 占用度公式 | threads/regs/blocks 三约束 | wave 槽位 + VGPR + LDS + SGPR 四约束 |
写 CUDA 时的「寄存器压力」迁移到 CDNA3 变成三选一：VGPR 压力（数据并行）/ SGPR 压力（地址并行）/ AccVGPR 压力（矩阵累加）。标量化（把不变的地址算术挪进 SGPR）是 AMD 侧独有的免费优化。
关键数字
KEY NUMBERS · OCCUPANCY BUDGET| 项 | 值 | 注 |
|---|---|---|
| wave 大小 | 64 线程（固定） | wave32 不存在于 gfx942[2] |
| wave 槽位 | 8/SIMD × 4 = 32/CU | 满占用 = 32 wave = 2048 线程 |
| VGPR 文件 | 128 KB/SIMD（32,768 个）· 512 KB/CU | 每 wave 用 V×64 个 |
| VGPR/wave | ≤ 256 常规 + 256 AccVGPR = 512 | 8 个粒度向上取整[2] |
| SGPR/wave | 16–102（16 粒度） | VCC 占最高 2 个 |
| workgroup 上限 | 16 waves（1024 线程） | LDS 按整 WG 分配（512B 粒度） |
| LDS | 64 KB/CU | occupancy 第三约束 |
占用度速算与标量化
FORMULA · SCALARIZATION① wave 槽位四约束（背下来）
# V = 每 wave 常规 VGPR 数（8 对齐）；LDS/G 按 workgroup：
waves_cu = min( 32, # ① wave slots
4 * floor(32768/(V*64)), # ② VGPR 文件
floor(65536/LDS_per_WG) * waves_per_WG, # ③ LDS
sgpr_waves ) # ④ SGPR（很少是瓶颈）
# 速查：V=64→100% · V=96→75%(24w) · V=128→50%(16w) · V=256→25%(8w)
# AccVGPR 另算：32×32 累加器 ≈ 128 AccVGPR/wave，与 V 合看 512 帽② 标量化：把地址算术搬进 SGPR
// HIP：指针/循环归纳变量用 const/整型，编译器自动标量化：
__global__ void saxpy(const float* x, float* y,
float a, int n) {
size_t i = (size_t)blockIdx.x * 64 + threadIdx.x;
// blockIdx 维度对整 wave 是 uniform → 地址基址走 SGPR
y[i] = a * x[i] + y[i]; // 只有数据本身占 VGPR
}
// 验证：rocprof-compute 的 SGPR/VGPR 分配报告注意事项 / 陷阱
PITFALLS · CHECKLIST
CAUTION · 陷阱清单
- 别按 32 算并行度：CUDA 的 warp-level 原语（__shfl / ballot / coalesced 32 段）在 HIP 里是 wave 级（64 段）——数组宽度、掩码位宽全要改。
- 分支粒度翻倍：wave64 里 64 线程一起分歧；32 线程粒度的 guard 带来的 divergence 行为与 CUDA 不同，对齐到 64 设计。
- workgroup 线程数要按 64 的倍数设计；尾块 63 线程也占一个整 wave 的寄存器预算。
- LDS 64KB 是硬顶（H200 228KB 的 28%）：从 NVIDIA 移植的 tile 尺寸普遍要砍半，bank/布局按 32 bank×4B（GCN 传统）重排。
- 每 CU 最大 workgroup 数无官方数字[6]——小 workgroup（如 64 线程）堆并发时先实测，别按 GCN 的 16 假设。
- occupancy 不是越高越好：MFMA 重载时 8–16 waves 高 ILP 常胜过 32 waves 满占用——与 CUDA 相同的取舍逻辑，但数字体系完全不同，别套用 H200 的调参表。
数据源
SOURCES- CDNA3 ISA Reference（MI300）§2/§3（wave64 / VGPR / SGPR / LDS 分配规则）
- ROCm Profiler CDNA pipeline（32 wave 槽 / 4×SIMD 组织）
- CDNA3 White Paper（寄存器文件容量）
- Composable Kernel GEMM 指南（occupancy 与 tile 的取舍实践）