B200 · BLACKWELL GB100 · 特性专题

双 die · NV-HBI 10 TB/s

★ Blackwell 封装主线 · 两颗 GB100 缝成单 GPU · 对 CUDA 全透明，对 L2 延迟不透明

是什么

DEFINITION

B200 用两颗 reticle 极限的 GB100 die（各 ~104B 晶体管）拼成一颗 GPU，die 间用 NV-HBI（NV High-Bandwidth Interface）连接：10 TB/s 的片间一致性互联，维持统一 L2、统一地址空间、软件单 GPU的编程模型^[1]。148 个 SM、126MB L2、180GB HBM 对 CUDA 完全无感——没有 device 分裂、没有显式跨 die 指令、没有 MIG 式切分。

但硬件上 L2 是每 die 独立 63MB：一块数据落在哪个 die 的 L2 分区，决定了访问它的 SM 是「近端」还是「跨 die 远端」。

为什么（vs 其他双芯片方案）

RATIONALE · VS ALTERNATIVES

方案	die 间带宽	编程模型	代价
B200 NV-HBI	10 TB/s，一致性	单 GPU（cudaSetDevice(0)）	跨 die L2 惩罚 ~300 cyc^[4]
两卡 NVLink 5	1.8 TB/s（双向）	2 device + NCCL/P2P	带宽仅 1/5，需显式通信
MI300X XCD 间	5.3 TB/s（4-XCD 组内）	单 GPU，CU 有 XCD 亲和性	波调度跨 XCD 延迟敏感

reticle 光刻极限卡住单 die 面积后，「10TB/s 缝合 + 透明单 GPU」是唯一不改变 CUDA 编程模型的做法。代价藏在微架构里：远端 L2 访问 ≈ 近端 2×（~575 vs ~275 cyc），而 SM 落在哪个 die、数据首次被哪个 die 的 L2 缓存，软件只能间接影响。

关键数字

KEY NUMBERS · CONSTRAINTS

项	值	注
NV-HBI 带宽	10 TB/s（die↔die）^[1]	≈ 5.6× 本卡 HBM 带宽
SM 分布	74 / die × 2 = 148第三方	CTA 落 die 由调度器决定
L2 分布	63 MB / die × 2 = 126 MB^[2]	数据按访问来源落分区
近端 L2 延迟	~150 ns（~275 cyc）^[5]	与 Hopper 量级相当
跨 die 惩罚	+~300 cyc（远端 ≈ ~575 cyc）^[4]	SemiAnalysis L2 sweep 两簇分布
跨 die 带宽差	本地 21 TB/s vs 跨 die 16.8 TB/s^[4]	atomics：同分区 90–100ns / 跨 190–220ns
显式控制 API	无	页/数据放置不可直接指定

调优（观测而非控制）

OBSERVABILITY · NO EXPLICIT API

NV-HBI 没有任何用户 API——能做的是观测跨 die 流量占比，再用 L2 residency 控制间接影响放置。

① Nsight Compute 定位 near/far L2 流量

# 远端（跨 die）L2 sector 占比：
ncu --metrics lts__t_sectors_srcunit_tex_op_read.sum,\
lts__t_sectors_gblpc_read.sum ./gemm
# 经验：far 占比 > ~30% 且 kernel 是 L2 敏感型
# （GEMM 主循环 / attention K·V 复读）时值得调放置

② L2 residency hint 间接影响放置

// 让热数据尽量钉在访问它的 SM 的本地分区：
cudaStreamAttrValue attr{};
attr.accessPolicyWindow.base_ptr  = hot_tensor;
attr.accessPolicyWindow.num_bytes = min(len, l2CacheSize); // ≤63MB 保守值
attr.accessPolicyWindow.hitRatio  = 0.6f;
attr.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

注意事项 / 陷阱

PITFALLS · CHECKLIST

CAUTION · 陷阱清单

远端 L2（~575 cyc）比 HBM（~420 cyc）还慢——「数据在 L2 里」不等于快，落错 die 比不缓存更糟。
同一 kernel 的 CTA 天然分布在两 die；工作分解跨 die 共享越少越好（避免用全局标记做跨 CTA 细粒度握手）。
原子热点：跨 die 分区的 atomic 延迟 ~2×（190–220ns vs 90–100ns）^[4]——高竞争 atomic 尽量按 die 聚簇或改为 per-warp 归约。
页放置策略 NVIDIA 未文档化，不要写死假设；用 Nsight 实测你的 kernel 而不是抄结论。
别把 B200 当「2×74 SM 的 NUMA 机器」手工平衡——调度器不给控制权，过度设计只增加复杂度；先测 far 占比再动手。
MIG / 多进程场景下 die 归属边界同样不透明，MIG 实例性能波动排查优先看 L2 分区冲突。

数据源

SOURCES

NVIDIA Blackwell 架构页 + 技术白皮书（双 die / NV-HBI 10TB/s / 统一 GPU 语义）
Blackwell Tuning Guide（L2 126MB 分区、缓存行为）
SemiAnalysis · Dissecting Blackwell Tensor（跨 die 惩罚 ~300cyc、分区带宽/atomics 实测）
Chips and Cheese B200 分析（近端 L2 延迟、SM 分布）