AMD · INSTINCT · CDNA3 · 8 XCD + 4 IOD · OAM
MI300X· 数字速查
CUDA 算子 / 框架工程师的 AMD 侧对照视角 · 括号内为相对 ×H200 倍数
CUDA ↔ ROCm 术语对照
TERMINOLOGY MAP · FOR CUDA ENGINEERS本页以下全部使用 AMD 原生术语；这张表是两套词汇的 Rosetta 石。量级直觉：1 个 CU ≈ ½ 个 SM 的通用算力，但 CU 数量是 SM 的 2.3 倍；1 wave = 2 warp。
| CUDA (NVIDIA) | ROCm (AMD · CDNA3) | 注解 | |
|---|---|---|---|
| SM | ↔ | CU（Compute Unit） | MI300X 304 CU vs H200 132 SM；单 CU 64 lane |
| Warp（32 线程） | ↔ | Wavefront / wave（64 线程，固定） | wave32 是 RDNA 特性，gfx942 没有 |
| Thread block | ↔ | Workgroup（≤16 waves = 1024 线程） | __syncthreads 在 HIP 里同名 |
| Thread block cluster / DSMEM | ↔ | — 无对应 | CDNA3 无集群层，跨 CU 共享靠 L2/global |
| Shared memory | ↔ | LDS（Local Data Share，64 KB/CU） | 与 L1 分离，不是合池 |
| L1 cache | ↔ | vL1（32 KB/CU，向量数据缓存） | + 16KB 标量 sL1（只读/uniform 路径）[3] |
| Register file | ↔ | VGPR（512/wave 上限）+ SGPR（分离） | 向量/标量分家 + AccVGPR 专属矩阵核 |
| Tensor Core / wgmma·tcgen05 | ↔ | Matrix Core / MFMA（同步指令） | 无异步 MMA；靠 wave 并行 + ILP 隐藏 |
| cp.async | ↔ | GLOBAL/BUFFER_LOAD … LDS=1（直达 LDS） | 每线程小粒度，无 bulk 引擎 |
| TMA | ↔ | — 无等价 | 没有 bulk-DMA / 描述符引擎 |
| mbarrier | ↔ | —（workgroup barrier / 计数器） | 异步完成靠 LGKM/VM 计数器 |
| Occupancy（warps/SM） | ↔ | Occupancy（waves/CU，≤32） | 限制维度不同：见 §2 折叠舱 |
| cc 9.0 / sm_90 | ↔ | gfx942（LLVM target） | ROCm 6+ / hipcc |
| PTX | ↔ | GCN ISA（AMDGPU 汇编） | v_mfma / v_global_load … |
| Nsight Compute | ↔ | rocprofiler-compute（Omniperf） | 计数器体系对应 SIMT 章节命名不同 |
整机级
ACCELERATOR-LEVEL · FULL-DEVICE SPEC| 指标 | MI300X OAM | 相对 H200 · 刻度 0–4× | 直觉注解 |
|---|---|---|---|
| 核心规模 | |||
| 发布年份 | 2023.12（Advancing AI） | H200 · 2023.11（SC23） | 相隔 1 个月同期对位；两者均 2024 起规模交付 |
| 架构 / 封装 | CDNA3 · 8 XCD + 4 IOD | Hopper · 单 GH100 | chiplet：XCD 5nm + IOD 6nm，3.5D 封装[1] |
| 工艺 / 晶体管 | N5+N6 · 153B | 4N · 80B | 良率优先的拆 die 路线 |
| CU 数量 | 304（8 XCD × 38） | ×2.30 | 每 XCD 造 40 开 38；单 CU ≈ ½ SM[1] |
| 流处理器（FP32 lane） | 19,456 (64/CU) | ×1.15 | vs H200 16,896 CUDA core |
| Matrix Cores | 1,216 (4/CU) | ×2.30 | vs 528 Tensor Core；MFMA 见 §5 |
| FP64（vector） | 81.7 TFLOPS | ×2.40 | AMD 传统强项；HPC 口径看这行[1] |
| 存储与带宽 | |||
| HBM 容量 | 192 GB HBM3 | ×1.36 | 8×24GB 堆叠；单卡塞大模型的意义[1] |
| HBM 带宽 | 5.3 TB/s (5325 GB/s) | ×1.10 | 8192-bit @5.2Gbps/pin；HBM3 非 3e |
| L2（每 XCD 内） | 32 MB（4 MB × 8 XCD） | ×0.64 | XCD 内一致；跨 XCD 走 fabric |
| ★ MALL（Infinity Cache） | 256 MB（4 IOD × 64MB） | ≈ L2 ×5.1 | memory-side、全片共享、不存脏数据[1] |
| 频率 / 功耗 / 互联 | |||
| 峰值引擎频率 | 2.1 GHz | ×1.15 | 官方峰值均按 2.1GHz 推[1]；1 cyc ≈ 0.48 ns |
| TDP | 750 W | ×1.07 | OAM 被动散热 |
| Infinity Fabric（scale-up） | 896 GB/s（7×128 GB/s 链路） | ×1.00 | vs NVLink4 900；8 卡全互联[1] |
| PCIe | Gen5 · 128 GB/s | 同 | 单 x16；其余 7 链路给 fabric |
| 目标 / 软件栈 | gfx942 · ROCm 6+ | cc 9.0 · CUDA | hipcc 近源码兼容迁移 |
单 CU 级
PER-CU RESOURCES · GFX942 VS CC 9.0 SM| 资源 | MI300X CU (gfx942) | H200 SM (cc 9.0) | 直觉注解 |
|---|---|---|---|
| SIMD 组织 | 4 SIMD × 16-lane | 4 分区 × 32-lane | CU 每 cycle 吐 64 lane；SM 128 |
| Matrix Cores | 4（MFMA，挂 VALU 发射） | 4 Tensor Core | MFMA 可与 VALU 指令同发射[1] |
| VGPR 文件 | 512 KB/CU（128KB×4 SIMD） | 256 KB RF | ×2；矩阵累加另有 AccVGPR |
| VGPR / wave 上限 | 512 = 256 常规 + 256 AccVGPR | 255 regs/thread | AccVGPR 专属矩阵核，可从 memory 直载[2] |
| VGPR 分配粒度 | 8 个/步（8 dword 组） | 8（warp 级） | 向上取整到 8 的倍数 |
| SGPR | 16–102/wave · 16 粒度 | — 无分离 | 地址/循环变量走标量域，省 VGPR |
| LDS（≈smem） | 64 KB | 228 KB | ×0.28——tiling 预算紧得多 |
| LDS 分配粒度 | 512 B 块 | 128 B | 按 workgroup 分配 |
| vL1（≈L1） | 32 KB（128B line） | 合池内 28–228KB | CDNA3 L1 与 LDS 分离 |
| 每 CU 并发 wave | 32（8/SIMD × 4）= 2,048 线程 | 64 warps = 2,048 线程 | 满占用线程数恰好相同[2] |
| Workgroup 上限 | 16 waves = 1,024 线程 | 1,024 线程 | 同；但 64 线程粒度（1 wave） |
| 每 CU 最大 workgroup 数 | 无官方数字第三方 | 32 blocks | GCN 传统 16，未见 CDNA3 官方值[6] |
OCCUPANCY 速算 · wave 槽位四约束
AMD occupancy 的单位是 wave（不是线程）。每 CU 并发 wave 数 = 四个约束的最小值[2]：
# V = 每 wave 的 VGPR 数（8 对齐）；每 SIMD 文件 32768 VGPR
waves_per_simd = floor(32768 / (ceil64(V) * 64))
waves_cu = min( 32, # ① wave 槽位（8/SIMD×4）
4 * waves_per_simd, # ② VGPR 文件
lds_waves, # ③ LDS：64KB / 每 WG 占用
sgpr_waves ) # ④ SGPR 文件
# 例：V=64 → 32768/(64*64)=8/SIMD → 32 waves → 100% 占用
# V=128 → 4/SIMD → 16 waves → 1024 线程 → 50%
# V=256 → 2/SIMD → 8 waves → 25%（AccVGPR 另算）与 CUDA 最大的心智差异：寄存器按 wave 整体申请（64 线程 × V），且标量/向量分家——把地址和循环归纳变量放进 SGPR、uniform 分支用标量域，等效「免费的」VGPR 优化。GCN 时代传下来的 wave occupancy 纪律在 CDNA3 全部适用。
CDNA3 架构图
MICROARCHITECTURE SCHEMATIC · 8 XCD + 4 IODFIG. 03 — CDNA3 BLOCK DIAGRAM · XCD ×8 → INFINITY FABRIC → IOD/MALL ×4 → HBM
存储层级
MEMORY HIERARCHY · CAPACITY × BANDWIDTH| 层级 | 容量（每 CU / 全片） | 带宽 / 延迟（第三方实测） | 谁管理 / 注解 |
|---|---|---|---|
| VGPR（+ AccVGPR） | 512 KB/CU · 156 MB/片 | 无公开延迟数据[6] | 编译器分配；occupancy 主约束 |
| LDS（≈smem） | 64 KB/CU · 19 MB/片 | ~28.7 ns（~60 cyc）[4] | 程序员显式；512B 分配粒度 |
| vL1（≈L1） | 32 KB/CU | ~62 ns（~130 cyc）[4] | 128B line；relaxed 一致性 |
| L2（XCD 内） | 4 MB/XCD · 32 MB/片 | ~104 ns（~219 cyc）[4] | XCD 内硬件一致；聚合读 34.4TB/s[1] |
| ★ MALL（全片共享） | 256 MB（4×64MB） | ~218 ns（~458 cyc）[3][4] | memory-side；≈ H200 L2 的角色但大 5× |
| HBM3（global） | 192 GB | 5.3 TB/s · ~378–412 ns（~800 cyc）[3] | TLB miss 惩罚 ~47ns：64MB 处延迟跳变[4] |
精度算力对照
PRECISION × THROUGHPUT · TFLOPS单位 TFLOPS（INT8 为 TOPS）。MI300X 官方峰值即 dense；2:4 sparse（SMFMAC 指令）在 datasheet 另列一列[1]。roofline 用 dense。
DENSE · 主读数（官方口径）SPARSE · 副读数（2:4 SMFMAC，非基准）
| 精度 | MI300X DENSE | MI300X SPARSE | H200 DENSE | H200 SPARSE | ×H200 (dense) |
|---|---|---|---|---|---|
| FP64（vector） | 81.7 | — | 34 | — | ×2.40 |
| FP64 Matrix（MFMA） | 163.4 | — | 67 | — | ×2.44 |
| FP32（vector） | 163.4 | — | 67 | — | ×2.44 |
| TF32 Matrix | 653.7 | 1,307.4 | 494 | 989 | ×1.32 |
| BF16 Matrix | 1,307.4 | 2,614.9 | 990 | 1,979 | ×1.32 |
| FP16 Matrix | 1,307.4 | 2,614.9 | 990 | 1,979 | ×1.32 |
| ★ FP8 Matrix（含 BF8） | 2,614.9 | 5,229.8 | 1,979 | 3,958 | ×1.32 |
| INT8 Matrix | 2,614.9 | 5,229.8 | 1,979 | 3,958 | ×1.32 |
| FP4 / FP6 | — 无 | — | — | — | B200 独有（9 PF FP4） |
① dense 是官方主数字：与 NVIDIA 页的「sparse 主数字」相反，跨厂对比时先统一到 dense。② 2:4 稀疏 MI300X 是有的（SMFMAC 指令族，INT8/FP8/FP16/BF16/TF32）[1]——「AMD 无结构化稀疏」是 MI200 时代的过时结论。③ 无 FP4/FP6：对标 B200 的 9 PF NVFP4 是代际差（CDNA3 vs Blackwell），不是同代对比对象。
关键延迟
LATENCY CROSS-SECTION · 第三方实测 @2.1 GHZAMD 从未发布官方延迟数字，本节全部来自公开微基准[3][4]（两文方法不同：cycles 为标量 pointer-chase、ns 为向量路径，比值量级互相印证）。1 cyc ≈ 0.476 ns。
| 访问类型 | ~cycles | ~ns | 来源 / 注解 |
|---|---|---|---|
| VGPR | ~1 | ~0.5 | 无公开实测估算；依赖链 ~4cyc 量级同 GPU |
| LDS（无 bank conflict） | ~60 | ~28.7 | pointer-chase 实测（H100 对照 16.7ns）[4] |
| vL1（global 命中） | ~130 | ~62 | 单线程 load 实测[4] |
| L2（XCD 内命中） | ~219 | ~104 | XCD 内一致；跨 XCD 另算[4] |
| MALL（Infinity Cache 命中） | ~458 | ~218 | memory-side；NPS4 最多再降 1.42×[3] |
| HBM（global miss） | ~800 | ~378–412 | 大数组实测（TLB 干净后）[3] |
| TLB miss 惩罚 | ~99 | ~47.1 | 64MB 工作集处延迟跳变（~16K 项 TLB）[4] |
| 全局 atomic（跨 CU） | ~245–426 | ~116–202 | 高值 ≈ 跨 XCD 往返[4] |
实测方法与可信度说明 — 见页脚脚注 [3] [4] [6]
MI300X 的层级是干净的翻倍链：L2(~220) → MALL(~460) → HBM(~800)，每级约 2×。一次 HBM ≈ 13 次 LDS。两个 AMD 特有陷阱：LDS 延迟（~60cyc）是 H200 smem（~30cyc）的 2 倍——tiling 收益打折，tiling 深度要重算；64MB 处的 TLB 跳变——工作集跨过 64MB 界限时延迟阶跃，分块策略要避开。
Roofline 临界点
ROOFLINE RIDGE POINT · FLOP/BYTEridge 是硬件常数、AI 是算子变量；AI > ridge → compute-bound，AI < ridge → memory-bound。MI300X 算力 ×1.32 但带宽只 ×1.10，ridge 全面高于 H200 ~20%——同算子更容易落在 memory-bound 侧。
| 精度 (dense) | 峰值 | 临界 ridge (FLOP/byte) | vs H200 · 典型 kernel 归属 |
|---|---|---|---|
| FP8 Matrix | 2,614.9 T | ~493 | H200 ~412 → 更难 compute-bound |
| FP16 / BF16 Matrix | 1,307.4 T | ~247 | H200 ~206；GEMM / attention 达标 |
| TF32 Matrix | 653.7 T | ~123 | H200 ~103；大 GEMM 达标 |
| FP32（vector） | 163.4 T | ~31 | H200 ~14——FP32 kernel 更易 compute-bound |
| FP64 Matrix | 163.4 T | ~31 | H200 ~14；HPC 双精度显著友好 |
| FP64（vector） | 81.7 T | ~15 | H200 ~7；不稀疏、不依赖 TC 的 FP64 强 |
实操：rocprofiler-compute（Omniperf）取 achieved AI 对照 ridge。迁移注意：NVIDIA 上 memory-bound 的 FP32 elementwise 到 MI300X 可能翻成 compute-bound（ridge 14→31）；反过来 LDS 延迟 ×2 又会拖 tile 复用重的算子——别照搬 H200 的调参结论，两台仪器的甜点区不同。
CDNA3 关键特性
FEATURE MODULES × 5 · OPERATOR OPTIMIZATION五大算子优化特性，每张卡 = 定义 + 关键数字 + 极简 snippet。点 [详情→] 看机制 / 完整示例 / 陷阱。
Chiplet · 8 XCD + 4 IOD
153B 晶体管 · N5 计算片 + N6 IO 片 · 3.5DGPU 拆成 8 颗计算 chiplet + 4 颗 IO chiplet。L2 只在 XCD 内一致（4MB/颗），全片共享靠 256MB memory-side MALL——局部性行为与单片 GPU 完全不同。
// XCD 归属软件不可控；观测靠 rocprof：
// L2 hit vs MALL hit 比例 = XCD 局部性
// workgroup 尽量自包含（数据在本地
// L2 命中），跨 XCD 共享走 MALLMFMA · Matrix Core
4/CU · 1307 TFLOPS BF16 dense · 2:4 稀疏 ×2AMD 的矩阵指令族（≈wgmma 的角色但同步执行）。累加器在 AccVGPR（可从 memory 直载），FP64 也有 MFMA 路径（163T——NVIDIA 的 2.4×）。形状：16×16×16 / 32×32×8（BF16）。
// HIP 内联汇编，一个 wave 发一条：
asm("v_mfma_f32_16x16x16_f16 "
"v[0:15], v[0:7], v[8:15], v[0:15]");
// D/A/B 都在 VGPR/AccVGPR，同步语义Wave64 · VGPR/SGPR 分家
wave=64 线程固定 · 32 wave/CU · VGPR 512/waveCDNA 的 wave 是 64 线程且不可选。寄存器分向量（VGPR）与标量（SGPR）两个文件——地址/循环变量放 SGPR 等于省 VGPR；occupancy 按 wave 槽位算，公式与 CUDA 不同。
// occupancy：V 个 VGPR/wave →
waves = min(32, 4*floor(32768/(V*64)), lds, sgpr)
// V=64 → 32 waves（100%）; V=128 → 16（50%）Infinity Fabric · 8 卡全互联
7×128 GB/s = 896 GB/s · 无交换机直连每卡 7 条链路与其余 7 卡全互联（无 NVSwitch 式交换机）。896 GB/s 对标 NVLink4 的 900；192GB×8=1.5TB 单机显存是 MI300X 集群的卖点。ROCm 通信走 RCCL。
// RCCL 与 NCCL API 同构（hip 替换）：
ncclAllReduce(buf, buf, n, ncclBfloat16,
ncclSum, comm, stream);
// 8 卡 xgmi 全互联：P2P 直达不走 host异步通路 · 无 TMA 的世界
LOAD→LDS 直达 · AccVGPR 直载 · 无 bulk 引擎CDNA3 没有 TMA/bulk-DMA：异步靠 GLOBAL/BUFFER_LOAD 的 LDS=1 直达位、AccVGPR 从 memory 直载，以及最根本的——32 wave/CU 的并行度本身。写惯 Hopper 的要换思路。
// 每线程 1 dword 直达 LDS（≈cp.async）：
asm("global_load_lds_dword v[2:3], v[0:1], off");
// 等待靠 LGKM 计数器，非 mbarrier
asm("s_waitcnt lgkmcnt(0)");