NVIDIA · BLACKWELL · GB100 双 DIE · SXM
B200· 数字速查
面向 CUDA 算子 / 框架工程师 · 写 kernel 时一眼查到的关键数字 · 括号内为相对 ×H200 倍数
整机级
CHIP-LEVEL · FULL-DEVICE SPEC| 指标 | B200 SXM | 相对 H200 · 刻度 0–4× | 直觉注解 |
|---|---|---|---|
| 核心规模 | |||
| 发布年份 | 2024.3(GTC24) | H200 · 2023.11 | 宣布仅隔 4 个月:H200 是 Blackwell 前的显存过渡档 |
| 架构 / 封装 | Blackwell · 双 GB100 die | Hopper · 单 GH100 | 两颗 reticle die 合一,软件见单 GPU[1] |
| 工艺 / 晶体管 | TSMC 4NP · 208B | 4N · 80B | 晶体管 ×2.6;跨 die 10TB/s NV-HBI |
| SM 数量 | 148(74/die)第三方 | ×1.12 | nvidia-smi 可见,官方未文档化[3] |
| FP32 CUDA Cores | 18,944 (128/SM)第三方 | ×1.12 | 每 SM 128 与 Hopper 同[3] |
| FP64 Cores | ~9,472 (64/SM)推算 | ×1.12 | FP64=½FP32 rate(实测印证)[5] |
| Tensor Cores(第五代) | 592 (4/SM)第三方 | ×1.12 | 数量同,架构换代:tcgen05 + TMEM[4] |
| 存储与带宽 | |||
| HBM 容量 | 180 GB HBM3e | ×1.28 | 8 堆叠 · 8192-bit;GB200 口径 186GB;B300 已至 288GB(见 §9)[1] |
| HBM 带宽 | 8 TB/s (8000 GB/s) | ×1.67 | 4800→8000;STREAM 实测 7.48 (93.5%)[6] |
| L2 Cache | 126 MB(每 die 63MB) | ×2.52 | 50→126MB,缓存更大模型切片[2] |
| ★ NV-HBI(die 间) | 10 TB/s | BLACKWELL 新 | 跨 die 一致性互联;L2 跨 die 惩罚 ~300cyc[4] |
| 频率 / 功耗 / 互联 | |||
| GPU Boost Clock | ~1.83 GHz第三方 | ×1.00 | 未官宣;1 cyc ≈ 0.55ns[3] |
| TDP | 1,000 W | ×1.43 | GB200 内 B200 为 1200W 口径[1] |
| NVLink(第五代) | 1.8 TB/s 双向 · 18 links | ×2.00 | NVL72 单域 72 GPU · 130TB/s[1] |
| PCIe | Gen6 · 256 GB/s | ×2.00 | GPU 能力 Gen6;HGX 系统到 host 实际 Gen5[1] |
| Compute Capability | 10.0(sm_100) | 9.0 | tcgen05 / TMEM / NVFP4 / gather4 门槛[2] |
单 SM 级
PER-SM RESOURCES · CC 10.0 VS CC 9.0| 资源 | B200 (cc 10.0) | H200 (cc 9.0) | 直觉注解 |
|---|---|---|---|
| FP32 Cores | 128 | 128 | 同;datapath 不变 |
| FP64 Cores | 64 | 64 | 同(但 FP64 TC 吞吐弱化,见 §5) |
| Tensor Cores | 4(第五代 · tcgen05) | 4(第四代 · wgmma) | wgmma 在 sm_100 已移除[2] |
| Register File | 256 KB = 65,536 × 32-bit | 256 KB | 同[2] |
| ★ Tensor Memory(TMEM) | 256 KB = 128 lanes × 512 col × 32b | — | tcgen05 累加器专属;按列 2 的幂分配 32–512[2] |
| Shared Memory(max) | 228 KB(每 block 上限 227 KB) | 228 KB | 同;carveout 档位同 Hopper[2] |
| L1 + SMEM 合池 | 256 KB | 256 KB | 同 |
| Max threads / SM | 2,048 | 2,048 | 同;64 warps |
| Max thread blocks / SM | 32 | 32 | 同 |
| Max registers / thread | 255 | 255 | 同(硬件保留 1) |
| 每 thread 平均寄存器 @满占用 | 32 | 32 | 同;加 reg 必降占用度 |
| Max thread blocks / cluster | 16(portable 8) | 16(portable 8) | 同;nonportable 16 需显式 opt-in[2] |
OCCUPANCY 速算 · 寄存器 × TMEM 双约束
每 SM 65,536 寄存器、最多 2,048 线程——寄存器/占用度取舍公式与 Hopper 完全相同。Blackwell 新增第四维:TMEM 列分配。
# 寄存器维度(同 Hopper):
threads = min(2048, floor(65536 / R) * blocksize)
# TMEM 维度(新):tcgen05.alloc 按列 2 的幂申请 32–512 列
# 每 SM 共 512 列(=256KB)。例:m64n256 累加器占 128 列
# → 最多 4 个 CTA 并发持累加器,TMEM 成为占用度瓶颈
Hopper 上 wgmma 的累加器占寄存器(warpgroup 每人一大块),Blackwell 把它挪进 TMEM——寄存器压力骤减,但换成了 TMEM 列压力。调参从「寄存器 vs 占用度」变成「寄存器 + TMEM 双约束」。
Blackwell 架构图
MICROARCHITECTURE SCHEMATIC · GB100 DUAL-DIEFIG. 03 — GB100 ×2 BLOCK DIAGRAM · HBM → L2(每 die) → NV-HBI → SM(TMEM)
存储层级
MEMORY HIERARCHY · CAPACITY × BANDWIDTH| 层级 | 容量(每 SM / 全片) | 峰值带宽 | 谁管理 / 注解 |
|---|---|---|---|
| Register File | 256 KB/SM · 38 MB/片 | ~13 TB/s/SM(量级) | 编译器分配;最快,决定占用度 |
| ★ TMEM(Tensor Memory) | 256 KB/SM · 38 MB/片 | 读 ~16 TB/s · 写 ~8 TB/s/SM[6] | tcgen05 专属;累加器不再吃 RF[2] |
| Shared Memory / L1 | 228 KB/SM(合池 256KB) | 128 B/cyc/SM · 全片 ~35 TB/s(量级)[4] | 程序员显式;bank conflict 防护 |
| L2 Cache | 126 MB(每 die 63MB) | 本地 ~21 TB/s(实测)[4] | 硬件管理;分区跨两 die |
| ★ NV-HBI(跨 die L2 访问) | — | 10 TB/s(die 间)[1] | 惩罚 ~300 cyc;远端 L2 ≈ 近端 2×[4] |
| HBM(global) | 180 GB | 8 TB/s | STREAM 实测 7.48;TMA 隐藏[6] |
精度算力对照
PRECISION × THROUGHPUT · TFLOPS单位 TFLOPS(INT8 为 TOPS)。官方只发 sparse(2:4)值,dense = sparse ÷ 2(官方脚注规则[1])。roofline / 性能分析用 dense。
DENSE · 主读数(基准口径)SPARSE · 副读数(2:4 参考,非基准)
| 精度 | B200 DENSE | B200 SPARSE | H200 DENSE | H200 SPARSE | ×H200 (dense) |
|---|---|---|---|---|---|
| FP64(vector,非 TC) | 37 | — | 34 | — | ×1.09 |
| FP64 Tensor Core | 37 | — | 67 | — | ×0.55 |
| FP32(非 TC) | 75 | — | 67 | — | ×1.12 |
| TF32 Tensor Core | 1,125 | 2,250 | 494 | 989 | ×2.28 |
| BF16 Tensor Core | 2,250 | 4,500 | 990 | 1,979 | ×2.27 |
| FP16 Tensor Core | 2,250 | 4,500 | 990 | 1,979 | ×2.27 |
| FP8 Tensor Core | 4,500 | 9,000 | 1,979 | 3,958 | ×2.27 |
| FP6 Tensor Core | 4,500 | 9,000 | — | — | BLACKWELL 新 |
| INT8 Tensor Core | 4,500 | 9,000 | 1,979 | 3,958 | ×2.27 |
| ★ FP4 / NVFP4 Tensor Core | 9,000 | 18,000 | — | — | BLACKWELL 新 |
① dense = sparse ÷ 2:官方 datasheet 的 TC 峰值默认带 sparsity[1],别拿 sparse 值算 roofline。② FP64 tensor 弱化:67→37 TFLOPS(×0.55),FP64 密集 HPC 负载别指望 Blackwell 升级。③ FP4 的 9 PF 需要 block-scale(NVFP4)格式——e2m1 裸 4-bit 数值范围不够,每 16 元素带 E4M3 scale factor 才是 18 PF 口径的前提(见特性页)。
关键延迟
LATENCY CROSS-SECTION · 实测/估算分级B200 公开微基准远少于 Hopper,本节逐行标注来源与可信度:估算 = 无公开实测、按 Hopper 量级推断;其余为公开实测(cycles 按 ~1.83 GHz第三方换算,1 cyc ≈ 0.55 ns)。
| 访问类型 | ~cycles | ~ns | 来源 / 注解 |
|---|---|---|---|
| Register 读 | ~1 | ~0.5 | 量级同 Hopper估算 |
| Shared Memory | ~30 | ~16 | 无公开实测;带宽 128 B/cyc/SM[4]估算 |
| L1(global 命中) | ~39 | ~19.6 | 实测 19.6 ns[5] |
| L2(近端 die 分区) | ~275 | ~150 | 实测 ~150 ns[5] |
| L2(跨 die 远端分区) | ~575 | ~310 | 近端 + ~300 cyc 跨 die 惩罚[4] |
| HBM(global,L2 miss) | ~420 | ~210–230 | pointer-chase 实测;Hopper ~1000 cyc[6] |
| ★ tcgen05.mma 单指令 | ~11 | ~6 | 实测 11.0–11.4 cyc,几乎不随 tile 增大(wgmma 32–128)[6] |
| SMEM bank conflict | ×2 / ×4+ | 倍增 | 量级同 Hopper估算 |
实测来源与可信度说明 — 见页脚脚注 [4] [5] [6] [7]
HBM 延迟比 Hopper(~650 cyc)降约 1/3,但层级比例不变:一次 HBM ≈ 14 次 smem ≈ 420 次 寄存器。tiling + TMA 异步预取的纪律仍然成立;新增注意点是跨 die L2(~575 cyc)比 HBM 更慢——数据落在远端 die 的 L2 分区时,局部性调优(L2 residency / 数据布局)在 Blackwell 上更重要。
Roofline 临界点
ROOFLINE RIDGE POINT · FLOP/BYTEridge point 是硬件常数;AI(arithmetic intensity)是算子变量。AI > ridge → compute-bound;AI < ridge → memory-bound。用 dense 列(§5)。带宽 ×1.67、算力 ×2.27 同时涨,脊点整体抬升 ~35%。
| 精度 (dense) | 峰值 | 临界 ridge (FLOP/byte) | vs H200 · 典型 kernel 归属 |
|---|---|---|---|
| FP4 TC | 9,000 T | ~1,125 | 只有大 batch 推理 GEMM 能达 |
| FP8 TC | 4,500 T | ~562 | H200 ~412 → 更难达 compute-bound |
| FP16 / BF16 TC | 2,250 T | ~281 | H200 ~206;GEMM / FA 可达 |
| TF32 TC | 1,125 T | ~141 | H200 ~103;大 GEMM 达标 |
| FP32(非 TC) | 75 T | ~9.4 | elementwise / 归一化多在此 |
| FP64 | 37 T | ~4.6 | H200 ~7 → HPC kernel 更易 compute-bound |
实操同 Hopper:Nsight Compute achieved AI 对照 ridge。注意两个迁移陷阱:① H200 上 compute-bound 的 FP8 kernel 到 B200 更可能变 memory-bound(ridge 412→562);② 用 FP4 时 ridge 高达 1125——必须把权重和激活都压到 4-bit 并保持大 tile,否则 FP4 的 9 PF 完全吃不到。
Blackwell 关键特性
FEATURE MODULES × 5 · OPERATOR OPTIMIZATION五大算子优化特性,每张卡 = 定义 + 关键数字 + 极简 snippet。点 [详情→] 看机制 / 完整可编译示例 / 陷阱。
tcgen05 · 五代 Tensor Core + TMEM
单线程发射 · 累加器驻 256KB TMEM · ~11 cyc/指令取代 wgmma:MMA 由一个线程发射,累加器放进独立的 Tensor Memory 而不是寄存器。cta_group::2 让两个 CTA 跨 SM 协同一个 128×256 大 tile。
// 分配 64 列 TMEM(2 的幂,32–512)
asm("tcgen05.alloc.cta_group::1.sync.aligned"
".shared::cta.b32 [addr], 64;");
// 单线程发射异步 MMA,累加器在 TMEM
asm("tcgen05.mma.cta_group::1.kind::f16"
" [d_tmem], a_desc, b_desc, idesc, p;");
TMEM 模型 / 完整示例 / 陷阱 →
NVFP4 · 原生 4-bit
e2m1 + 每 16 元素 E4M3 block scale · 2× FP8Blackwell 主打精度:e2m1 裸格式范围不够,NVFP4 每 16 个元素配一个 E4M3 缩放因子,精度损失可控。9 PF dense 的前提就是 block-scale 格式。
// NVFP4 = e2m1 (4b) × scale16 (E4M3, /16 元素)
__nv_fp4_e2m1 x[16];
__nv_fp8_e4m3 sf; // block scale
// tcgen05 idesc 选 .kind::mxf8f6f4 + nvfp4
// 路径;scale 经 SMEM 旁路进 TC
格式对比 / 精度策略 / 陷阱 →
双 die · NV-HBI 10TB/s
两颗 GB100 合一 · 软件单 GPU · 跨 die +~300cycreticle 极限下的双 die 方案:10TB/s 一致性互联把两颗 die 缝成单 GPU。对 CUDA 完全透明,但 L2 分区跨 die——远端分区延迟翻倍是主要性能 Cliff。
// 无任何跨 die API:cudaMalloc/launch 如常
// 但页放置影响 L2 归属:
// cudaMemAdvise(ptr, ReadMostly, ...)
// Nsight Compute 看 lts__t_sectors_op_read
// 的 near/far 分区占比定位跨 die 流量
L2 分区行为 / 调优 →
NVLink 5 · 1.8TB/s
900 GB/s 单向 · NVL72 单域 72 GPU第五代 NVLink 双向 1.8TB/s(H200 ×2),NVL72 机柜把 72 颗 GPU 缝成单一 NVLink 域(130TB/s 聚合),二级交换可扩到 576 GPU。all-reduce / TP 通信的底盘。
// NCCL 自动用 NVL72 all-to-all 路径:
ncclCommInitRank(&comm, 72, ...);
// TP=8 训练:all-reduce 走 NVLink
// 而非 IB——latency 低一个量级
ncclAllReduce(buf, buf, n, ncclFloat, ...);
NVL72 拓扑 / 带宽账 →
TMA 增强 + Cluster 16
device 端改 tensormap · gather4/scatter4TMA 基础能力继承 Hopper,新增三件:tensor map 可在 kernel 内原地改写(免 host 往返);tile::gather4/scatter4 按索引搬非连续行(MoE/embedding);cluster nonportable 上限 8→16。
// device 端原地换 tensor map 基址:
asm("tensormap.replace.tile.global_address"
.global.b1024.b64 [map], newPtr;");
// 按索引数组搬 4 行(MoE 行聚合):
asm("cp.async.bulk.tensor.2d.tile::gather4"
".shared::cluster.global.mbarrier...");
gather4 语义 / cluster 16 / 陷阱 →
B300 · Blackwell Ultra 对比
NEXT-GEN DELTA · 官方 ULTRA 口径升级路线速览,数字来自 NVIDIA Blackwell Ultra 官方材料[8]。注意口径:本表按官方 Ultra 对比口径(B200 取 1200W 超芯片档、HBM 取物理 192GB);本页 §1–§7 仍是 1000W HGX 单卡口径(180GB / FP4 dense 9 PF)。
| 指标 | B200 · Blackwell | B300 · Blackwell Ultra | 直觉注解 |
|---|---|---|---|
| 架构 | Blackwell · 双 die 4NP · 208B | Blackwell Ultra(同代微调) | 同 die 血统:4NP / 双 die / 208B 不变[8] |
| HBM3e 容量 | 192 GB(物理) | 288 GB | +50%;12-Hi 堆叠(vs B200 8-Hi) |
| HBM 容量提升 | 基准 | +50% | vs HGX B200 可用 180GB 则为 +60% |
| HBM 带宽 | 8 TB/s | 8 TB/s(不变) | 容量涨带宽平:每 GB 带宽 −33% |
| Dense NVFP4 推理 | 10 PF(基准) | 15 PF | 约 +50%——Ultra 的主打提升[8] |
| Sparse NVFP4 | 20 PF | 20 PF(不变) | 稀疏上限持平;增量全在 dense 档 |
| FP8(dense/sparse) | 5 | 10 PF | 5 | 10 PF(不变) | FP8 训练负载在 Ultra 上无算力红利 |
| NVLink / PCIe | 1.8 TB/s / Gen6 | 1.8 TB/s / Gen6(不变) | 互联代际不动 |
| TGP | 最高 1,200 W | 最高 1,400 W | +200W 换 dense FP4 与容量 |
① 带宽不涨、容量 +50%:每 GB 显存的带宽 −33%——大 KV cache / 长 context / 大 batch 推理最受益(装得下 + 流得动),但纯带宽受限的算子(embedding lookup、小 GEMM)不会变快。② dense FP4 +50% 而 sparse 上限不变:Ultra 抬的是真实负载能吃到的 dense 档位(10→15 PF),选型时别按 sparse 20 PF 外推。③ FP8 零变化:FP8 训练集群换 Ultra 只赚容量,不赚算力——预算该花在哪档精度,先想清楚。