B200CUDA QUICK REFQR-100

NVIDIA · BLACKWELL · GB100 双 DIE · SXM

B200· 数字速查

面向 CUDA 算子 / 框架工程师 · 写 kernel 时一眼查到的关键数字 · 括号内为相对 ×H200 倍数

MAIN READOUT · 主显示屏 DISP · 01
HBM3E 容量 · CAPACITY 180GB H200 141GB · ×1.28 · 8 堆叠
HBM 带宽 · BANDWIDTH 8TB/s ×1.67 H200 · STREAM 实测 7.48
SM 数量 · SM COUNT 148 ×1.12 · 双 die 74+74[3]
BF16 算力 · DENSE 2,250TFLOPS SPARSE 4,500 · 非基准 · ×2.27
TMEM / SM · 新存储层 256KB tcgen05 累加器 · 不占寄存器
COMPUTE CAPABILITY 10.0 BLACKWELL · tcgen05 / NVFP4 / NV-HBI
RDY · 口径 B200 SXM 1000W · DENSE=SPARSE÷2(官方规则) · ×H200 为整量比值 · 延迟含实测/估算分级标注
SCALE 1:1 · GB100 ×2 · 180 GB
01

整机级

CHIP-LEVEL · FULL-DEVICE SPEC
指标B200 SXM相对 H200 · 刻度 0–4×直觉注解
核心规模
发布年份2024.3(GTC24)H200 · 2023.11宣布仅隔 4 个月:H200 是 Blackwell 前的显存过渡档
架构 / 封装Blackwell · 双 GB100 dieHopper · 单 GH100两颗 reticle die 合一,软件见单 GPU[1]
工艺 / 晶体管TSMC 4NP · 208B4N · 80B晶体管 ×2.6;跨 die 10TB/s NV-HBI
SM 数量148(74/die)第三方×1.12nvidia-smi 可见,官方未文档化[3]
FP32 CUDA Cores18,944 (128/SM)第三方×1.12每 SM 128 与 Hopper 同[3]
FP64 Cores~9,472 (64/SM)推算×1.12FP64=½FP32 rate(实测印证)[5]
Tensor Cores(第五代)592 (4/SM)第三方×1.12数量同,架构换代:tcgen05 + TMEM[4]
存储与带宽
HBM 容量180 GB HBM3e×1.288 堆叠 · 8192-bit;GB200 口径 186GB;B300 已至 288GB(见 §9)[1]
HBM 带宽8 TB/s (8000 GB/s)×1.674800→8000;STREAM 实测 7.48 (93.5%)[6]
L2 Cache126 MB(每 die 63MB)×2.5250→126MB,缓存更大模型切片[2]
NV-HBI(die 间)10 TB/sBLACKWELL 新跨 die 一致性互联;L2 跨 die 惩罚 ~300cyc[4]
频率 / 功耗 / 互联
GPU Boost Clock~1.83 GHz第三方×1.00未官宣;1 cyc ≈ 0.55ns[3]
TDP1,000 W×1.43GB200 内 B200 为 1200W 口径[1]
NVLink(第五代)1.8 TB/s 双向 · 18 links×2.00NVL72 单域 72 GPU · 130TB/s[1]
PCIeGen6 · 256 GB/s×2.00GPU 能力 Gen6;HGX 系统到 host 实际 Gen5[1]
Compute Capability10.0(sm_100)9.0tcgen05 / TMEM / NVFP4 / gather4 门槛[2]
02

单 SM 级

PER-SM RESOURCES · CC 10.0 VS CC 9.0
资源B200 (cc 10.0)H200 (cc 9.0)直觉注解
FP32 Cores128128同;datapath 不变
FP64 Cores6464同(但 FP64 TC 吞吐弱化,见 §5)
Tensor Cores4(第五代 · tcgen05)4(第四代 · wgmma)wgmma 在 sm_100 已移除[2]
Register File256 KB = 65,536 × 32-bit256 KB[2]
Tensor Memory(TMEM)256 KB = 128 lanes × 512 col × 32btcgen05 累加器专属;按列 2 的幂分配 32–512[2]
Shared Memory(max)228 KB(每 block 上限 227 KB)228 KB同;carveout 档位同 Hopper[2]
L1 + SMEM 合池256 KB256 KB
Max threads / SM2,0482,048同;64 warps
Max thread blocks / SM3232
Max registers / thread255255同(硬件保留 1)
每 thread 平均寄存器 @满占用3232同;加 reg 必降占用度
Max thread blocks / cluster16(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 双约束」。

03

Blackwell 架构图

MICROARCHITECTURE SCHEMATIC · GB100 DUAL-DIE
FIG. 03 · 设备线框原理图
HBM3e · 180 GB · 8 TB/s · 8 堆叠(每 die 专属 4 栈) global memory · ~420 cyc/访(Hopper ~650,实测降 58%) ×1.67 H200 带宽 DIE 0(GB100) L2 · 63 MB(半片 126MB)· 分区 crossbar SM ×74 4×TC/SM GPC… cluster ≤8 TMEM 256KB/SM TMA +replace/g4 DIE 1(GB100 · 镜像) L2 · 63 MB · 独立分区 SM ×74 4×TC/SM GPC… cluster ≤8 TMEM 256KB/SM TMA +replace/g4 NV-HBI ★ 10 TB/s die 间一致性 跨 die L2 +~300cyc 单 SM 内部结构(两 die 内同构) 4× Sub-partition · 每 partition: 32 FP32 + 1 TC 4× Tensor Core(第五代 · tcgen05 单线程发射) TMEM 256KB · 累加器/矩阵操作数(128 lane×512 col) Register File 256KB · max 2048 threads Shared Mem / L1 合池 256KB(smem ≤228KB) TMA unit · +tensormap.replace / gather4 编程层级(自底向上) Thread(32/warp) → Warp → Thread Block(单 SM) → Cluster(≤8 portable / 16 nonportable) → Grid(软件单 GPU = 2 die) Blackwell 三大异步 TMA(搬运,可 device 端改描述符) tcgen05(计算,累加器驻 TMEM) mbarrier(同步,同 Hopper) wgmma 已移除:sm_90a 专属 NV-HBI 对 CUDA 编程模型完全透明(无 device 分裂、无显式跨 die 指令)
HBM(实心) 片上存储(L2 / SMEM / RF) SM / 计算单元 Blackwell 新特性(NV-HBI / TMEM / tcgen05 / TMA+) 数据通路 / 标注引线

FIG. 03 — GB100 ×2 BLOCK DIAGRAM · HBM → L2(每 die) → NV-HBI → SM(TMEM)

04

存储层级

MEMORY HIERARCHY · CAPACITY × BANDWIDTH
层级容量(每 SM / 全片)峰值带宽谁管理 / 注解
Register File256 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 / L1228 KB/SM(合池 256KB)128 B/cyc/SM · 全片 ~35 TB/s(量级)[4]程序员显式;bank conflict 防护
L2 Cache126 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 GB8 TB/sSTREAM 实测 7.48;TMA 隐藏[6]

带宽与延迟的实测来源 / 估算分级 — 见页脚脚注 [4] [6]

05

精度算力对照

PRECISION × THROUGHPUT · TFLOPS

单位 TFLOPS(INT8 为 TOPS)。官方只发 sparse(2:4)值,dense = sparse ÷ 2(官方脚注规则[1])。roofline / 性能分析用 dense。

DENSE · 主读数(基准口径)SPARSE · 副读数(2:4 参考,非基准)

精度B200 DENSEB200 SPARSE H200 DENSEH200 SPARSE×H200 (dense)
FP64(vector,非 TC)3734×1.09
FP64 Tensor Core3767×0.55
FP32(非 TC)7567×1.12
TF32 Tensor Core1,1252,250494989×2.28
BF16 Tensor Core2,2504,5009901,979×2.27
FP16 Tensor Core2,2504,5009901,979×2.27
FP8 Tensor Core4,5009,0001,9793,958×2.27
FP6 Tensor Core4,5009,000BLACKWELL 新
INT8 Tensor Core4,5009,0001,9793,958×2.27
FP4 / NVFP4 Tensor Core9,00018,000BLACKWELL 新
NOTE · 三条口径红线

① 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 口径的前提(见特性页)。

06

关键延迟

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–230pointer-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]

CAUTION · 延迟直觉锚点(量级)
HBM ~420 cyc
×1 · 14 等分
SMEM ~30 cyc
×14 ≈ 1 次 HBM
REGISTER ~1 cyc
×420 ≈ 1 次 HBM

HBM 延迟比 Hopper(~650 cyc)降约 1/3,但层级比例不变:一次 HBM ≈ 14 次 smem ≈ 420 次 寄存器。tiling + TMA 异步预取的纪律仍然成立;新增注意点是跨 die L2(~575 cyc)比 HBM 更慢——数据落在远端 die 的 L2 分区时,局部性调优(L2 residency / 数据布局)在 Blackwell 上更重要。

07

Roofline 临界点

ROOFLINE RIDGE POINT · FLOP/BYTE
RIDGE POINT · 脊点 = 峰值算力 ÷ HBM 带宽
ridge [FLOP/byte] = Peak_TFLOPS(dense) ÷ 8 TB/s

ridge 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 TC9,000 T~1,125只有大 batch 推理 GEMM 能达
FP8 TC4,500 T~562H200 ~412 → 更难达 compute-bound
FP16 / BF16 TC2,250 T~281H200 ~206;GEMM / FA 可达
TF32 TC1,125 T~141H200 ~103;大 GEMM 达标
FP32(非 TC)75 T~9.4elementwise / 归一化多在此
FP6437 T~4.6H200 ~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 完全吃不到。

08

Blackwell 关键特性

FEATURE MODULES × 5 · OPERATOR OPTIMIZATION

五大算子优化特性,每张卡 = 定义 + 关键数字 + 极简 snippet。点 [详情→] 看机制 / 完整可编译示例 / 陷阱

★ 特性 1 / 计算

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 模型 / 完整示例 / 陷阱 →
★ 特性 2 / 精度

NVFP4 · 原生 4-bit

e2m1 + 每 16 元素 E4M3 block scale · 2× FP8

Blackwell 主打精度: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
格式对比 / 精度策略 / 陷阱 →
★ 特性 3 / 封装

双 die · NV-HBI 10TB/s

两颗 GB100 合一 · 软件单 GPU · 跨 die +~300cyc

reticle 极限下的双 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 分区行为 / 调优 →
★ 特性 4 / 互联

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 拓扑 / 带宽账 →
★ 特性 5 / 搬运

TMA 增强 + Cluster 16

device 端改 tensormap · gather4/scatter4

TMA 基础能力继承 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 / 陷阱 →
AUXRESERVED · 预留扩展位
09

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 · BlackwellB300 · Blackwell Ultra直觉注解
架构Blackwell · 双 die 4NP · 208BBlackwell 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/s8 TB/s(不变)容量涨带宽平:每 GB 带宽 −33%
Dense NVFP4 推理10 PF(基准)15 PF约 +50%——Ultra 的主打提升[8]
Sparse NVFP420 PF20 PF(不变)稀疏上限持平;增量全在 dense 档
FP8(dense/sparse)5 | 10 PF5 | 10 PF(不变)FP8 训练负载在 Ultra 上无算力红利
NVLink / PCIe1.8 TB/s / Gen61.8 TB/s / Gen6(不变)互联代际不动
TGP最高 1,200 W最高 1,400 W+200W 换 dense FP4 与容量
NOTE · 换代没换的三件事

① 带宽不涨、容量 +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 只赚容量,不赚算力——预算该花在哪档精度,先想清楚。