MI300X · CDNA3 · GFX942 · 特性专题
Chiplet · 8 XCD + 4 IOD
★ CDNA3 封装主线 · L2 只在 XCD 内一致 · 全片共享靠 256MB memory-side MALL
是什么
DEFINITIONMI300X 把 GPU 拆成 8 颗 XCD（计算 chiplet，N5，每颗 38 CU + 4MB L2）和 4 颗 IOD（IO chiplet，N6，HBM 控制器 + 64MB MALL），SoIC 3D 堆叠 + 2.5D 转接板封装（AMD 称 3.5D）[1]。对软件是单 GPU：一个 gfx942 device、统一地址空间。
与单片 GPU 的本质区别在缓存层级：L2 是每 XCD 私有的 4MB，XCD 之间不做 L2 一致；全片共享缓存是 MALL（Infinity Cache，256MB，挂在 IOD 内存侧）。数据落在哪个 XCD 的 L2、由哪个 XCD 的 CU 访问，直接决定延迟——但软件无法指定。
为什么（vs 单片 / vs B200 双 die）
RATIONALE · VS MONOLITHIC / DUAL-DIE| 维度 | H200（单片） | B200（双 die NV-HBI） | MI300X（8+4 chiplet） |
|---|---|---|---|
| 拆分粒度 | 不拆 | 2 颗 die | 8 XCD + 4 IOD |
| 一致 L2 范围 | 全片 50MB | 全片 126MB（跨 die 一致） | 仅 XCD 内 4MB[1] |
| 全片共享缓存 | L2 本身 | L2 本身 | MALL 256MB（memory-side） |
| 远端访问代价 | L2 far 分区 ~500cyc | 跨 die L2 ~575cyc | 跨 XCD → MALL ~458cyc / HBM ~800cyc[4] |
| 软件控制 | L2 residency hint | 无（观测 near/far） | 无（观测 L2/MALL 命中） |
chiplet 是良率/成本路线（N5 只用在计算片上），代价是局部性管理变难：workgroup 与它的数据是否落在同一 XCD，决定走 4MB 私有 L2（~220cyc）还是 256MB MALL（~460cyc）——差一倍。
关键数字
KEY NUMBERS · TOPOLOGY| 项 | 值 | 注 |
|---|---|---|
| XCD 构成 | 8 ×（38 CU 启用 / 40 造）+ 4MB L2 | 每 XCD 4 个 ACE[1] |
| IOD 职责 | HBM 控制器 + MALL 64MB ×4 + PCIe/IF 链路 | N6 工艺 |
| MALL 总量 | 256 MB | memory-side：不存脏数据、带 snoop filter |
| MALL 带宽 | 17.2 TB/s（WP）或 14.7 TB/s（HC24）[6] | 官方材料冲突，引用需注明 |
| XCD↔IOD 接口 | 2.1 TB/s[1] | 口径（单 XCD/聚合）官方未注明 |
| L2 聚合读带宽 | 34.4 TB/s（2KB/clk × 16 通道） | XCD 内路径[1] |
| 延迟阶梯 | L2 ~220 → MALL ~460 → HBM ~800 cyc | 每级 ~2×[3][4] |
调优（观测 XCD 局部性）
OBSERVABILITY · ROCPROFILER没有「绑 XCD」的 API。能做的是让 workgroup 数据局部性自包含，并观测 L2/MALL 命中分布。
① rocprofiler-compute 看 L2 vs MALL 命中
# L2（XCD 内）命中 vs MALL 命中占比：
rocprof-compute -m lts__t_sectors_srcunit_tex_op_read.sum,\
SQC_L2_HIT_RATE,L2_cache_hit_rate ./kernel
# 经验：L2 命中低 + MALL 命中高
# → workgroup 与数据跨 XCD，考虑加大
# workgroup 内数据自包含度（分块按 XCD 尺寸对齐）② 布局策略：按 XCD 尺度做数据亲和
// 每 workgroup 自带 tile（HIP，同 CUDA 思路）：
__global__ void gemm_tile(const float* A, ...) {
__shared__ float a_tile[64*64]; // LDS 64KB 上限内
// 关键：一个 WG 读的 A/B tile 越聚焦，
// 越可能整块落进本 XCD 的 4MB L2
}注意事项 / 陷阱
PITFALLS · CHECKLIST
CAUTION · 陷阱清单
- 不要假设「L2 32MB」是全片一致的——它是 8 个 4MB 私有片的算术和；跨 XCD 的一致最低层其实是 MALL（memory-side）。
- MALL 不缓存脏数据（write-through 定位），写密集 workload 的 MALL 收益远低于读密集[1]。
- 64MB TLB 跳变：工作集 >64MB 后 HBM 延迟阶跃（~47ns 惩罚）[4]——大 batch 推理的分块别骑在界上。
- workgroup 调度到哪个 XCD 由驱动决定且不透明；同 kernel 跨 XCD 的 atomic 热点实测 116–202ns，高竞争时改为 per-workgroup 归约再合并。
- 官方 fabric/IC 带宽数字自相冲突（4.8/6TB/s、14.7/17.2TB/s）[6]——写报告引用时必须注明出处与口径。
- MI300A（APU）的 XCD 数、CU 数都不同（6 XCD/228 CU + Zen4），别把 MI300X 结论平移。
数据源
SOURCES- AMD CDNA3 White Paper（chiplet 构成 / L2 / MALL / 带宽）
- MI300X Datasheet（规模 / TDP / 峰值）
- SC'25 "The MALL is Open"（MALL/L2/HBM 延迟与行为实测）
- Chips and Cheese MI300X 测试（延迟阶梯 / TLB / atomics）
- ROCm MI300 微架构页（XCD/ACE/缓存组织）