MI300X · CDNA3 · GFX942 · 特性专题
异步通路 · 无 TMA 的世界
★ CDNA3 搬运主线 · LOAD→LDS 直达 · AccVGPR 直载 · 靠 32 wave/CU 并行隐藏延迟
是什么
DEFINITIONCDNA3 没有 TMA / bulk-DMA 引擎，也没有 mbarrier。它的异步武器是三类细粒度机制[2]：
① LOAD→LDS 直达位——global_load_lds_* / buffer_load_* lds=1 / scratch_load_lds_*：每线程一个 dword，数据从 global 直接落 LDS 不经过 VGPR（≈Ampere cp.async 的角色）。② AccVGPR 直载——矩阵累加器/A 操作数可从 memory 直接装载进 AccVGPR。③ 计数器同步——LGKM/VM 计数器 + s_waitcnt，等一批异步读写完成。
根本的延迟隐藏手段不在指令里，在调度模型里：每 CU 32 个 wave 的并发，内存等待时硬件切到别的 wave——「用并行度换延迟」而非「用异步流水换延迟」。
为什么（vs TMA 异步流水）
RATIONALE · VS TMA PIPELINE| 维度 | H200 TMA+mbarrier | MI300X 直达 load+waitcnt |
|---|---|---|
| 发起粒度 | 单线程 1 条指令搬整块 tile（≤smem） | 每线程 1 dword，一个 wave 256B/发 |
| 地址生成 | 硬件按 tensor-map 描述符 | 软件算地址（SGPR 间接寻址 buffer 资源） |
| 越界/填充 | 硬件 0/NaN 填充 | 软件 guard / padding |
| 完成同步 | mbarrier 记字节数 | s_waitcnt lgkmcnt(0)（指令计数） |
| 占用资源 | 几乎不占线程（1 个发射） | 占满 wave 的发射带宽 |
| 延迟隐藏主力 | 异步流水（TMA+wgmma 重叠） | wave 并发切换（32/CU）+ MFMA 与 VALU 同发射 |
迁移要点：Hopper 的「TMA 预取下一块 + wgmma 算当前块」双缓冲，在 CDNA3 上对应「前一 wave 的直达 load 与当前 wave 的 MFMA 天然重叠」——数据通路自动重叠，但搬运本身吃发射带宽，tile 不能太大（LDS 64KB 也顶不住大 tile）。
关键数字
KEY NUMBERS · CONSTRAINTS| 项 | 值 / 约束 | 注 |
|---|---|---|
| 直达 LDS 指令族 | global_load_lds_{u,s}{byte,short,dword} · buffer_load …lds=1 · scratch_load_lds_* | 每线程 1 dword；字节/短字带符号扩展[2] |
| AccVGPR 直载 | memory → AccVGPR（不经 VGPR） | MFMA A 操作数专用通路 |
| 完成计数器 | LGKM（LDS/global/消息）· VM（向量读写） | s_waitcnt 按 wave 计，非字节数 |
| LDS 写口 | 成对 SIMD 共享 · 64B/cyc/方向 | 直达 load 的落地带宽上限[1] |
| 单 wave 直达搬运 | 64 dword = 256 B/发 | 1KB tile 要 4 wave·发或多次发 |
| 重叠来源 | 32 wave/CU 切换 + MFMA/VALU 同发射 | occupancy 掉到 8 wave 以下重叠就崩 |
可编译示例
COMPILABLE EXAMPLE · VERBATIM① Global→LDS 直达（≈cp.async）+ 计数器等待
// 每线程一个 dword 直落 LDS；v[0:1] = global 地址（SGPR 基址 + lane 偏移）：
asm volatile("global_load_lds_dword v[2:3], v[0:1], off");
// …继续发若干条，然后一次等完本 wave 的 LGKM 队列：
asm volatile("s_waitcnt lgkmcnt(0)");
// LDS 内数据此后可见（workgroup 级再用 barrier 同步对齐）② HIP 层的等价写法（编译器生成直达序列）
// HIP 对齐 cp.async 的封装（ROCm 6+）：
__hip_memcpy_to_shared(lds_ptr, gmem_ptr, 4);
// 底层即 global_load_lds_dword + s_waitcnt；
// 大块搬运让编译器展开为多条直达 load③ 双缓冲模式（wave 级流水）
// wave i: 预取 tile[i+1]（直达 LDS buf1）
// wave j: MFMA 消费 tile[i] （LDS buf0 → VGPR → 累加）
// 两波天然重叠——不需要 mbarrier，
// 只需保证 buf0/buf1 各 32KB 内（LDS 64KB 上限）注意事项 / 陷阱
PITFALLS · CHECKLIST
CAUTION · 陷阱清单
- s_waitcnt 是按指令条数，不是字节数：等的是本 wave 的 LGKM 队列清空；跨 wave 的可见性要另加 barrier。
- 直达 load 占发射带宽：搬运过密会挤占 VALU/MFMA 发射——搬运量大的算子先看「直达 load 条数 vs MFMA 条数」的比例。
- LDS 64KB 顶：双缓冲两块 tile 各 ≤32KB；从 Hopper 移植的 128KB 级 tile 必须重切。
- 没有硬件越界填充——边界 tile 的 masking/padding 是软件责任（对比 TMA 的 OOB_FILL）。
- LDS 写口是成对 SIMD 共享 64B/cyc：直达 load 的跨 bank 写布局要均匀，否则落 LDS 这一跳就成瓶颈。
- 别指望「异步」掩盖低 occupancy：wave 数 < ~8 时延迟裸奔——先按 wave64 特性页的四约束把并发拉起来，再谈搬运重叠。
数据源
SOURCES- CDNA3 ISA Reference（MI300）§9 Memory Buffer Load to LDS / §13（直达指令族 + s_waitcnt 语义）
- CDNA3 White Paper（LDS 端口 / AccVGPR 直载）
- arXiv 2602.10262（MI300A 异步执行特性实测：直达 load 与计数器行为）
- Composable Kernel GEMM 指南（wave 级双缓冲实践）