Microbenchmarking NVIDIA's Blackwell Architecture — L2 #

Jarmusch & Chandrasekaran, 2025-12 | cs.AR |

§1 TL;DR #

首篇 Blackwell B200 PTX 微基准实测：TMEM 降低 58% 延迟 (420 vs 1000 cycle)；tcgen05 空间阵列设计实现 ~11 cycle 恒定 MMA 延迟 (Hopper wgmma 线性增长 32→128)；FP4 达 7700 TFLOPS (96% 峰值)；端到端推理 1.6–2×、训练 1.55×、能效 +32%。

§2 Q1 痛点 / Q2 方法 / Q3 结果 #

Q1 痛点 #

GPU 架构迭代远超社区深度理解。Blackwell 引入三大架构创新——Tensor Memory (TMEM)、第五代 Tensor Core (tcgen05 指令)、硬件解压引擎 (DE)——但 NVIDIA 仅发布高层 whitepaper，指令延迟、TMEM 带宽、DE 吞吐等关键微架构参数完全未知。现有方法均无法填补空白：

• Application profiler (nvprof/nsight)：粒度不足，无法区分 TMEM 访问与 L2 cache 延迟
• Roofline model：无法捕捉 TMEM、DE 等新硬件单元的行为
• 模拟器 (Accel-Sim, GCoM)：需要微架构参数作为输入——正是本文要测量的东西

核心技术壁垒：对 PTX→SASS 翻译的逆向映射理解。tcgen05.mma 根据精度编译成不同 SASS 指令——FP16/FP32→HMMA，FP8→QMMA，FP4/FP6→OMMA (全新)，INT→IMMA，FP64→DMMA (独立路径)。其中 OMMA 是 Blackwell 首创的 octal-byte MMA 指令，专为 sub-byte 格式设计。不理解这个映射就无法设计正确的 dependency chain 来隔离单指令延迟。

Q2 方法 #

自研 PTX 微基准测试套件，三个核心测量技术：

1. Dependency chain 隔离延迟——每条 MMA 的 accumulator $\mathbf{D}$ 作为下一条 MMA 的输入，阻止 pipeline overlap，测真单指令延迟
2. Back-to-back issuing 测吞吐——独立 MMA 指令饱和 tensor core pipeline，测峰值吞吐
3. Pointer-chase 测 TMEM/cache 延迟——依赖型 load 阻止预取和 pipeline 重叠

所有测量在 B200 和 H200 上对照执行，1000 次迭代取平均 (100 次 warmup)。通过 CUTLASS disassembly 验证 PTX→SASS 编译无意外优化。

Q3 结果 #

维度	关键数据
TMEM 延迟	420 cycle cache-miss (−58% vs Hopper 1000 cycle)
TMEM 带宽	16 TB/s read, 8 TB/s write per SM
tcgen05 延迟	11.0–11.4 cycle 恒定 (空间阵列)；Hopper wgmma 32–128 cycle (时域流水线)
FP4 吞吐	7700.2 TFLOPS (96.2% peak)
FP6 吞吐	5134.4 TFLOPS (96.0% peak)
DE 吞吐	Bitcomp 462 GB/s, LZ4 173 GB/s, output-bandwidth-bounded
LLM 推理	Mistral-7B FP16: 1.97×, FP8: 1.16×; Mixtral-8x7B FP16: 1.71×
训练	ResNet-50 1.85×, GPT-1.3B 1.55×, 能效 +32%

---

§3 架构 / 方法图 #

B200 双 Die 拓扑 #



Paper's Figure 1, verbatim (caption: "NVIDIA Blackwell GPU dual-die design interconnected via NV-HBI").

B200 是 NVIDIA 首款数据中心 dual-die GPU：两个 GPU die 共 208B 晶体管、148 SM、8 GPC，通过 NV-HBI 高带宽接口互联为统一设备。每个 die 拥有独立的 L2 cache 分区（共 4 个，Hopper 的 2 倍）和 HBM3e 内存堆栈（共 8 个，192 GB）。NV-HBI 使双 die 对软件完全透明——统一内存空间、无显式 NUMA 管理。图中可见两个 die 呈对称布局，中间由 NV-HBI 桥接，每个 die 内部的 4 个 GPC 块均匀排列。

Tensor Core 指令 Pipeline 演进 #



Paper's Figure 2, verbatim (caption: "Tensor Core instruction pipeline for tcgen05, wgmma, and Volta/Ampere architectures").

三代演进的核心数据通路变化清晰可见：

• Volta/Ampere (左列)：mma.sync warp-synchronous，8 线程 (Volta) 或 1 warp (Ampere) 锁步执行，操作数全部经由 SMEM→RF→Tensor Core→RF
• Hopper (中列)：wgmma warp-group 级 (×4 warps = 128 线程)，operand B 从 SMEM 直读，operand A 从 RMEM，accumulator D 写回 RF
• Blackwell (右列)：tcgen05.mma 单线程发射，TMEM 替代 RF 作为 accumulator 存储——operand A 从 TMEM，operand B 从 SMEM，结果 D 写回 TMEM

这一变化不仅降低执行粒度 (128→32 线程)，更关键的是 TMEM 使 accumulator 驻留在专用片上内存中，消除了 chained MMA 场景下中间结果经 RF 到全局内存的回写。

微基准测量方法论 #

flowchart TB subgraph latency["延迟测量 — Dependency Chain"] L1["tcgen05.mma #1
D₁ = A×B + C"] --> L2["tcgen05.mma #2
D₂ = A'×B' + D₁"] L2 --> LN["tcgen05.mma #N
Dₙ = ...+ Dₙ₋₁"] LN --> LR["Δclock / N = SI-LAT"] end subgraph throughput["吞吐测量 — Independent Issue"] T1["tcgen05.mma"] ~~~ T2["tcgen05.mma"] T2 ~~~ TN["tcgen05.mma ×N"] TN --> TR["total ops / time = TFLOPS"] end subgraph tmem["TMEM 延迟 — Pointer Chase"] M1["tcgen05.ld addr₀ → addr₁"] --> M2["tcgen05.ld addr₁ → addr₂"] M2 --> MN["...addr_N"] MN --> MR["Δclock / N = access latency"] end

PTX→SASS 指令映射 (Table IV) #

精度	PTX (tcgen05.mma)	SASS	Hopper wgmma SASS
FP16, BF16	kind::f16	HMMA	HGMMA
FP32, TF32	kind::tf32	HMMA	HGMMA
FP8	kind::mxf8 / f8f6f4	QMMA	QGMMA
FP6	kind::mxf6	QMMA	N/A
FP4	kind::mxf4 / mxf4nvf4	OMMA	N/A
INT4, INT8	kind::i8	IMMA	IGMMA
FP64	不支持 tcgen05	—	DMMA

OMMA 是 Blackwell 全新指令，专为 FP4/FP6 设计。FP64 不走 tcgen05 路径，使用独立的翻倍 FP64 单元 (DMMA)。

---

§4 作者证明 #

无形式化数学证明 — 仅实证。 以下为 6 项实证一致性检查：

#	检查项	论文声称	验证方式	判定
1	TMEM 延迟	420 cycle cache-miss，−58% vs Hopper	Pointer-chase 微基准 (§IV-A1)，dependent load 阻止 pipeline overlap	✓ 降低来自 TMEM 绕过 L2 分区竞争的专用仲裁逻辑
2	tcgen05 恒定延迟	11.0–11.4 cycle (tile 64→256)	Accumulator dependency chain (§IV-A3)，三种 tile size 对照	✓ 0.4 cycle 方差支持空间阵列假说；Hopper 对照组线性增长 (32→128) 排除测量伪影
3	FP4 峰值利用率	7700.2 TFLOPS = 96.2% peak	Back-to-back issuing vs NVIDIA 官方 spec 8000 TFLOPS	✓ 96% 利用率与其他精度 (96.0–99.6%) 一致
4	DE output-bounded	170–220 GB/s output 不随压缩率变化	4 种数据模式 (1×→245×) 对比 (Table II)：input 随 $1/C$ 下降，output 稳定	✓ $1/C$ 趋势清晰
5	1.27× 一致加速比	B200/H200 共享精度 (FP8–INT8)	Table VII 六种精度全部 1.27×	✓ 一致性极强，推测为 SM 数量 × 频率乘积
6	训练加速分解	$1.09 \times 1.27 \times 1.26 \approx 1.75$	但实测仅 1.55× (Table XI)	⚠ 因子非独立——乘积超出实测 13%，存在未量化的 diminishing returns

额外可疑点：

• Mistral-7B FP8 加速仅 1.16× (Table VIII)，远低于 FP16 的 1.97×——H200 的 FP8 路径已被深度优化，Blackwell 在 FP8 上的增量优势被压缩
• STREAM Triad 小工作集 (4–16 GB) B200 仅达 52% 峰值 BW，低于 H200 的 60%——论文归因于"H200 更适配小工作集"但未给出具体机理
• TMEM 15% 功耗节省在小矩阵上反转为 3–5% 功耗增加，crossover 矩阵尺寸未精确量化

---

§5 实验与数据 #

5.1 Tensor Core 单指令延迟 — 空间阵列 vs 时域流水线 (Table V) #

Instruction	Tile Shape	Scope	SI-LAT (cycles)
wgmma	m64n64k16	Warp-group	32.0
wgmma	m64n128k16	Warp-group	64.0
wgmma	m64n256k16	Warp-group	128.0
tcgen05.mma	m64n64k16	Warp	11.0
tcgen05.mma	m128n128k16	Warp	11.3
tcgen05.mma	m256n256k16	Warp	11.4

本文最重要的微架构发现：Hopper wgmma 延迟随 tile 宽度线性增长 (32→128 cycle)，Blackwell tcgen05 恒定 ~11 cycle。Tile size 影响吞吐而非延迟——证明 Blackwell tensor core 采用空间展开的计算阵列，不同于 Hopper 的时序流水线复用。

5.2 多精度 Tensor Core 吞吐 (Table VII) #

Precision	B200 (TFLOPS)	% Peak	H200 (TFLOPS)	Speedup
FP64	44.8	99.6%	34.0	1.32×
FP32	482.0	96.4%	378.4	1.27×
TF32	964.5	96.5%	756.9	1.27×
BF16	1926.4	96.3%	1513.5	1.27×
FP16	1929.6	96.5%	1515.2	1.27×
FP8	3850.6	96.3%	3026.9	1.27×
FP6	5134.4	96.0%	N/A	New
FP4	7700.2	96.2%	N/A	New
INT8	3928.5	98.2%	3088.4	1.27×

三个关键观察：(1) 所有精度均达 96–99% 峰值利用率——tensor core 不是 workload 瓶颈；(2) 共享精度一致 1.27× 加速比，FP64 例外为 1.32× (FP64 单元翻倍)；(3) FP4→FP64 吞吐差 177× 但延迟仅差 1.27× (11.2→14.2 cycle)，证明吞吐提升来自更宽数据通路而非更深 pipeline。

5.3 逐精度 Tensor Core 特性 (Table VI) #

Input (A/B)	Accum (C/D)	Latency (cycles)	Throughput (TFLOPS)	SASS
FP16	FP16	11.2	964.8	HMMA
FP16	FP32	11.5	482.4	HMMA
BF16	FP32	11.4	481.6	HMMA
FP8	FP16	11.8	1925.3	QMMA
FP8	FP32	12.1	1912.8	QMMA
FP6	FP16	12.3	2567.2	QMMA
FP4	FP16	12.6	3850.1	OMMA
INT8	INT32	11.9	3928.5	IMMA

FP32 accumulator 瓶颈：FP16→FP16 吞吐 964.8 TFLOPS，切换到 FP16→FP32 直接减半至 482.4 TFLOPS。瓶颈在 accumulator datapath 而非乘法单元——推理应优先 FP16 accumulator，训练无法回避 FP32 的吞吐代价。

5.4 LLM 推理性能 (Table VIII) #

Model	Precision	B200 tok/s	H200 tok/s	Speedup	Perplexity	ΔPPL
Mistral-7B	FP16	56,028	28,500	1.97×	6.82	—
Mistral-7B	FP8	57,125	49,200	1.16×	6.95	+1.9%
Mistral-7B	FP4	112,800	N/A	N/A	7.38	+8.2%
Mixtral-8x7B	FP16	31,033	18,100	1.71×	5.94	—
Mixtral-8x7B	FP8	51,200	32,400	1.58×	6.08	+2.4%
Mixtral-8x7B	FP4	76,900	N/A	N/A	6.48	+9.1%

FP4 在 Mistral-7B 上实现 2.01× 吞吐提升 (vs FP16 baseline)，代价是 perplexity +8.2%。BW utilization 从 FP16 的 67–76% 降至 FP4 的 47–49%——量化使 workload 从 memory-bound 向 compute-bound 迁移。注意 FP8 推理的 B200/H200 加速仅 1.16×——H200 的 FP8 路径已被充分优化，Blackwell 在 FP8 上的增量优势极小。

5.5 Latency vs Batch Size (Table IX) #

Batch Size	B200 (ms)	H200 (ms)	Ratio	B200 tok/s
1	12.3	18.7	1.52×	166,504
2	14.8	22.1	1.49×	276,757
4	19.2	28.4	1.48×	426,667
8	28.6	41.3	1.44×	572,727
16	47.1	67.8	1.44×	696,178
32	89.3	128.4	1.44×	734,264

小 batch 延迟优势 (1.48–1.52×) 超出纯计算比——论文推测 Blackwell 自动 pipeline 重配置将处理阶段从 ~20 缩减至 8–10。尾延迟 p99/median 从 H200 的 1.23–1.38 改善至 B200 的 1.12–1.14。

5.6 端到端训练 (Table XI) #

Model	B200	H200	Ratio	B200 Energy Eff
ResNet-50	2,928 img/s	1,580 img/s	1.85×	5.09 img/s/W
GPT-1.3B (BS=128)	14,363 tok/s	9,240 tok/s	1.55×	20.63 tok/s/W
GPT-1.3B (BS=64)	14,121 tok/s	9,070 tok/s	1.55×	20.27 tok/s/W

训练加速分解为三个因子：SM 增量 (1.09×) × CTA pairing (1.27×) × TMEM (1.26×)。能效提升 32% 尽管功耗高 ~14%，主要得益于 TMEM 减少 L2 thrashing 降低动态功耗。

5.7 DE 格式吞吐 (Table I) #

Format	Comp. Ratio	Output (GB/s)	Latency (ms)	用途
LZ4	1.00×	172.55	0.608	通用基线
Snappy	1.91×	117.24	0.894	实时场景
Zstd	2.00×	154.94	0.677	通用最优
GZIP	2.00×	83.83	1.251	旧系统兼容
Bitcomp	3.00×	462.37	0.227	科学计算
ANS	N/A	539.21	0.194	熵编码

所有格式 sub-millisecond 延迟。Bitcomp/ANS 输出吞吐远超其他格式 (462–539 vs 84–173 GB/s)，推测有专用整数/熵编码硬件路径。

5.8 DE 压缩率敏感性 (Table II) #

Data Pattern	Comp. Ratio	Input (GB/s)	Output (GB/s)	Latency (ms)
Random	1.00×	173.23	172.55	0.608
Mixed	1.98×	80.11	158.94	0.660
Repetitive	15.02×	14.63	219.80	0.477
Zeros	245.45×	0.85	209.83	0.500

输出吞吐稳定在 170–220 GB/s 不随压缩率变化，输入吞吐从 173→0.85 GB/s 随 $1/C$ 下降。DE 是 output-bandwidth-bounded 设计——对不可压缩数据充当 pass-through，对高压缩率数据内部解压不是瓶颈。

5.9 跨 Workload 汇总 (Table X) #

Workload	Metric	B200	H200	Speedup	Key Enabler
Attention Block	Latency (μs)	284	468	1.65×	TMEM
HPC DGEMM (FP64)	TFLOPS	36.3	18.9	1.92×	Doubled FP64 units
STREAM Triad (4–16 GB)	BW (TB/s)	4.14	—	—	HBM3e
SpMV (compressed)	GFLOPS	5.08	3.2	1.58×	Decomp engine
GPT Training (1.3B)	tok/s	14,363	9,240	1.55×	CTA pairs + TMEM + TC
ResNet Training	img/s	2,928	1,580	1.85×	5th Gen TC + mem BW
Energy Eff. (Training)	tok/s/W	20.63	15.6	1.32×	Process + TMEM

---

§6 论证链 #

Step	论据	来源	结论
1	Blackwell 引入 TMEM (256KB/SM)、tcgen05 (warp-level MMA)、DE (硬件解压) 三大创新，微架构参数全部未知；profiler/Roofline/模拟器均无法覆盖	§I, §II, §III	需要 PTX 级微基准来填补参数空白
2	PTX dependency chain (accumulator 传递) 隔离单指令延迟，pointer-chase 隔离 TMEM 访存延迟，back-to-back issuing 饱和 pipeline——CUTLASS disassembly 验证 PTX→SASS 正确性	§IV	测量方法可信地隔离单硬件单元行为
3	TMEM cache-miss 420 cycle (−58% vs Hopper)，16 TB/s 读带宽；chained GEMM $\mathbf{D} = (\mathbf{A} \times \mathbf{B}) \times \mathbf{C}$ 中间结果驻留消除每 SM ~12 TB/s 全局内存回写	§V-A	TMEM 对矩阵密集 workload 提供显著延迟+带宽优势
4	tcgen05 延迟恒定 11.0–11.4 cycle 不随 tile 从 m64→m256 变化；Hopper wgmma 线性增长 32→128 cycle	§VI-A, Table V	空间阵列设计确认（tile size 影响吞吐不影响延迟），区别于 Hopper 时域流水线
5	所有精度均达 96–99% 峰值利用率；FP4 (7700 TFLOPS) vs FP64 (44.8 TFLOPS) 吞吐差 177× 但延迟仅差 1.27×	§VI-A, Table VI–VII	吞吐提升来自更宽数据通路（更高并行度），tensor core 本身不是瓶颈——内存带宽和 kernel launch overhead 才是
6	端到端验证：推理 1.58–1.97×, 训练 1.55–1.85×, 能效 +32%；加速可分解为 SM (1.09×) × CTA (1.27×) × TMEM (1.26×)，但乘积 1.75 > 实测 1.55——因子非独立	§VII, Table X–XI	微基准结论在真实 workload 中得到验证；三大创新的贡献可近似独立量化，但存在 diminishing returns

---

§7 实现 cross-reference #

[实现未公开] — 论文声称开发了开源 PTX 微基准测试套件，但因审稿限制未发布代码：

"unable to share the code at this time due to double-blind" (§I)

关键实现细节 #

1. PTX→SASS 验证是可信性基础：所有微基准通过 CUTLASS disassembly 验证 PTX 编译到预期 SASS 指令 (Table IV)。tcgen05.mma 的 kind::mxf4nvf4 映射到 OMMA 而非 QMMA——精度参数错误会导致编译到完全不同的硬件单元，整个测量无效。复现时必须在 SASS 层面验证指令映射。
1. TMEM 最优 tile 64×64：TMEM 的 2D 结构 (512 列 × 128 lanes × 32-bit cells) 在 $64 \times 64$ tile (对 FP8 为 4KB) 时完全利用 1024-bit 接口宽度。Tiles $< 32 \times 32$ 接口利用率不足；tiles $> 128 \times 128$ 触发多阶段传输。所有迁移到 Blackwell 的 GEMM/Attention kernel 都应将 tiling 策略从 Hopper 的 $32 \times 32$ 调整为 $64 \times 64$。

TMEM 内存层次定位 #

层级	容量/SM	带宽	延迟	用途
Register File	~256KB	—	0–1 cycle	通用
L1 / SMEM	256KB	—	~30 cycle	通用 + TC operand B
TMEM	256KB	16 TB/s R, 8 TB/s W	420 cycle (miss)	TC 专用 (accumulator + operand A)
L2	~100MB (4 分区)	—	—	共享
HBM3e	192 GB	8.0 TB/s	~1000 cycle	全局

算力-带宽比 #

$$\text{FP16: } 1930 \text{ TFLOPS} / 8.0 \text{ TB/s} = 241 \text{ Ops/Byte}$$

$$\text{FP8: } 3851 \text{ TFLOPS} / 8.0 \text{ TB/s} = 481 \text{ Ops/Byte}$$

$$\text{FP4: } 7700 \text{ TFLOPS} / 8.0 \text{ TB/s} = 963 \text{ Ops/Byte}$$

更低精度使更多 workload 从 memory-bound 转向 compute-bound。实测验证：FP16 推理的 BW utilization 为 67.3% (memory-bound)，FP4 降至 47.6% (转向 compute-bound)。

架构设计约束 #

• TMEM 需要 kernel 重写：传统 ldmatrix/wmma.load 无法访问 TMEM，必须使用 tcgen05.cp/tcgen05.ld/tcgen05.st 全新指令族
• FP64 不走 TMEM/tcgen05 路径：FP64 DGEMM 使用独立的翻倍 FP64 单元 (DMMA)，TMEM 不改善科学计算
• DE 是 output-bounded 设计：高压缩率数据的 compressed input 吞吐随 $1/C$ 下降——DE 最适合中等压缩率场景 (2–4×)
• CTA Pair 执行：两个相邻 CTA 共享操作数通过 TPC 内通信网络，减少冗余数据搬运——这是 1.27× 训练加速的关键因子之一
• FP32 accumulator 使吞吐减半：硬件设计决定，无法通过软件绕过——推理用 FP16 accumulator，训练被迫承受 2× 吞吐损失