MI300X · CDNA3 · GFX942 · 特性专题
Infinity Fabric · 8 卡全互联
★ CDNA3 互联主线 · 7×128 GB/s 无交换机直连 · 8×192GB = 1.5TB 单机显存
是什么
DEFINITION每张 MI300X 出 7 条 16-lane Infinity Fabric 链路（每条双向 128 GB/s），在 UBB 主板上与其余 7 张卡全互联（无交换机）：合计 896 GB/s[1]。PCIe Gen5 x16 是留给 host 的第 8 条通路。卡间 P2P 与 XGMI（AMD 的 GPU 直连协议）让 8 卡表现为一个 1.5TB 显存的近节点域。
软件栈：通信走 RCCL（NCCL 的 ROCm 移植，API 同构），P2P 走 hipMemcpyPeer / hipHostRegister 直达。
为什么（vs NVLink / IB）
RATIONALE · COMPARISON| 维度 | H200 NVLink4 | B200 NVLink5/NVL72 | MI300X IF/XGMI |
|---|---|---|---|
| 单卡双向带宽 | 900 GB/s | 1.8 TB/s | 896 GB/s（7×128）[1] |
| 拓扑 | NVSwitch（基板交换） | NVSwitch4（72 卡单域） | 全互联直连，无交换机 |
| 单域规模 | 8 | 72 → 576 | 8（再往上走 IB/以太） |
| 单机显存 | 8×141GB = 1.13TB | 72×186GB = 13.4TB（NVL72） | 8×192GB = 1.5TB |
| 通信库 | NCCL | NCCL | RCCL（API 同构） |
MI300X 集群的故事是「单节点显存最大」：8 卡 1.5TB HBM3，一笔 70B 权重 + 长 context + KV cache 都塞得下。代价是 8 卡之外没有 scale-up 域——超 8 卡的 TP/EP 通信预算要按网络（IB/以太）重算。
关键数字
KEY NUMBERS · BANDWIDTH BUDGET| 项 | 值 | 注 |
|---|---|---|
| 链路数 × 带宽 | 7 × 128 GB/s（16-lane @32Gbps） | 双向口径[1] |
| 合计 | 896 GB/s | 第 8 条通路 = PCIe Gen5 x16 |
| 单机显存 | 1.5 TB（8×192GB HBM3） | vs HGX H200 1.13TB |
| 通信:显存带宽比 | 896 : 5300 ≈ 1 : 5.9 | H200 为 900:4800 ≈ 1:5.3；B200 1:4.4 |
| TP=8 all-reduce 预算 | ≈ bytes×2×7/8 ÷ 128GB/s/环 | 全互联环上每跳 128GB/s |
代码与账本
RCCL PATH · BANDWIDTH MATH① RCCL（NCCL API 同构，hipcc 编译）
// 8 卡 XGMI 全互联：RCCL 检测 P2P 直达，不走 host：
ncclCommInitRank(&comm, 8, devs, rank);
ncclAllReduce(buf, buf, n, ncclBfloat16, ncclSum, comm, stream);
// 验证拓扑：rocm-smi --showtopo / rccl 调试日志看 XGMI 链路② TP all-reduce 时延速算
# 70B 模型 FP8 梯度 70GB，ring all-reduce 通信量 2·N·(P-1)/P：
bytes = 2 * 70e9 * 7/8 # ≈ 122.5 GB / step
t_fabric = bytes / 122.5e9 # 全互联平摊 ≈ 1 s 量级（分层后更低）
# 结论：8 卡域内通信量级与 NVLink4 HGX 持平（900 vs 896），
# 超出 8 卡的部分必须重新按 IB/以太预算注意事项 / 陷阱
PITFALLS · CHECKLIST
CAUTION · 陷阱清单
- 896 GB/s 是 7 条链路的合计双向值：点到点单链路 128 GB/s——8 卡 all-to-all 满载时每对卡只有 128，不是 896。
- 无交换机全互联不能扩展：每卡链路数固定 7，8 卡是天然上限；9 卡以上回落网络，TP>8 的通信账完全不同。
- host 通路只有 PCIe Gen5 x16（~128GB/s 理论、实测更低）：CPU 侧喂数据（数据加载、非 GPUDirect 场景）容易先在这里饱和。
- RCCL 与 NCCL 参数语义兼容但默认算法不同：调优时用 rccl 的调试输出确认 ring/tree 与通道绑定，别直接抄 NCCL 环境变量结论。
- 跨卡 P2P 访问延迟（XGMI 读邻居显存）显著高于本地 HBM——kernel 里直接跨卡指针解引用只适合低频元数据，大块搬运交给 RCCL。
- 封装内 fabric 带宽的官方口径互相冲突（4.8/6TB/s）[6]：本页 896GB/s 指卡间 XGMI，与封装内 XCD↔IOD fabric 是两回事，引用别混。
数据源
SOURCES- MI300X Datasheet（7×128GB/s XGMI / PCIe Gen5 / 8 卡全互联）
- CDNA3 White Paper（Infinity Fabric 拓扑与总 I/O）
- Hot Chips 2024 MI300X（链路构成与口径冲突说明）
- ROCm RCCL 文档（API 与拓扑检测）