Dissecting the SM_120 Microarchitecture · Cycle-Level Characterization of Blackwell Consumer GPUs

Abstract
摘要

本文对 NVIDIA Blackwell 架构消费级 GPU——GB203 (SM_120, RTX PRO 5000 72GB) 进行了系统性的 cycle 级微架构特性研究。通过 21 个测试类别、185 个以上微测试案例和约 879 个跟踪项（覆盖率 96.1%），我们全面刻画了 SM_120 的执行管线、指令延迟、存储层级、TensorCore、TMA 异步数据通路、warp 调度器以及功耗效率等核心微架构参数。

主要发现包括：(1) SM_120 的 TensorCore 具有统一流水线特性——所有非 FP64 精度（FP16 至 MX-FP4 共 12 种格式）共享相同的 29/23 cycle 延迟/吞吐，精度选择可简化为纯存储带宽决策；(2) 首次揭示 SM_120 的 "5 管线 + V-pipe" 架构，修正了此前文献中的"4 管线"模型；(3) 量化了无 Tensor Memory / tcgen05 条件下的 GEMM 上限——CUTLASS 4.5 NVFP4 达 975 TFLOPS（80.7% 效率），证明 mma.sync + Cluster Launch Control 路径在消费级 Blackwell 上仍可接近 1 PFLOPS；(4) 发现 FlashAttention 在 SM_120 上是 ALU-bound（softmax 占 63.1%，TC 仅 36.9%），瓶颈在 SFU 而非 TensorCore；(5) 建立了首个消费级 Blackwell per-SM 功耗分解模型（P = 80W 基础设施 + 0.7-1.2W/SM），其中 MX-FP4 能效达 6.9 TFLOPS/W，是 FP32 的 28 倍。

本文的完整指令延迟参考表覆盖 60 余种指令类型，所有数据均基于实测并通过 SASS 反汇编和 NCU profiling 交叉验证。研究结果为 CUDA kernel 开发者、GPU 架构研究者和 ML 框架优化者提供了实用参考。

GPU 微架构 NVIDIA Blackwell SM_120 microbenchmark TensorCore 指令延迟功耗模型 FlashAttention CUTLASS

§ 01
Introduction

引言: 为什么是 SM_120

1.1 研究背景与动机

背景
Background

NVIDIA 于 2024 年推出的 Blackwell 架构标志着 GPU 计算的新一代演进。与此前的 Hopper (SM_90) 相比，Blackwell 引入了第五代 TensorCore（tcgen05）、Tensor Memory（TM）、FP6/FP4/MX block-scale MMA、Cluster Launch Control 等重大创新。然而，Blackwell 家族内部存在显著的产品分化：数据中心级 GB100/GB200 (SM_100) 拥有完整的 tcgen05 + TM 栈，而面向消费者和工作站的 GB203 (SM_120) 则有选择地裁剪了部分功能，形成了一种独特的"简化 Blackwell"定位。

理解 SM_120 的微架构对以下群体至关重要：CUDA kernel 开发者需要每条指令的延迟和管线映射以编写高效代码；GPU 架构研究者需要 sub-core 拓扑、scoreboard 机制和调度策略的定量数据；ML 框架优化者需要 TC 吞吐瓶颈数据和 GEMM 效率上限以指导推理部署决策。

尽管 Jarmusch 等人 [1] 对 Blackwell GB203 (RTX 5080) 进行了初步特性研究，但其覆盖范围有限，且未涉及 V-pipe 独立性、FlashAttention 瓶颈分析、per-SM 功耗模型等深层问题。Huerta 等人 [3] 提供了出色的方法学框架，但其 scoreboard 深度模型（"6 个计数器"）在 SM_120 上被我们的实验推翻（实测 ≥ 12）。

1.2 主要贡献 (Contributions)

★ C1 · Coverage

首个 SM_120 系统级微架构特性研究

覆盖 21 个测试类别、185+ case 报告、879 个跟踪项（96.1% 覆盖率）。目前针对 Blackwell 消费级 GPU 最全面的公开微架构数据集。

★ C2 · TensorCore

统一 TC 流水线发现

所有非 FP64 精度（12 种格式）共享相同的 29 cyc 延迟 / 23 cyc 吞吐。颠覆"低精度 TC 更快"的直觉，精度选择可简化为纯存储带宽决策。

★ C3 · Pipelines

"5 管线 + V-pipe" 执行架构

修正 Jarmusch [1] 的"4 管线"模型，识别出 P0/V-pipe/P1/P2/P3/P4 完整拓扑及 9 对 co-issue 组合。

★ C4 · GEMM

无 TM/tcgen05 的 GEMM 上限

CUTLASS 4.5 NVFP4 达 975 TFLOPS（80.7% 效率）。证明 mma.sync + CLC 路径在消费级 Blackwell 上仍可接近 1 PFLOPS。

★ C5 · Attention

FlashAttention 是 ALU-bound

逐 tile cycle 分解：softmax ALU 占 63.1%（exp2 主导），TC 仅 36.9%。瓶颈在 SFU 而非 TensorCore。

★ C6 · Power

首个消费级 per-SM 功耗模型

P = 80W 基础设施 + 0.7-1.2W/SM，基础设施占比 > 45%。MX-FP4 能效达 6.9 TFLOPS/W（FP32 的 28×）。

★ C7 · Scheduler

深度逆向 Warp 调度器

通过 CuAssembler binary patching 验证 17-bit 控制码全语义，scoreboard 深度 ≥ 12，量化 yield bit 完全阻断 latency hiding 效应。

★ C8 · Reference

完整 SM_120 指令延迟参考表

覆盖 60+ 指令类型的 latency/throughput/pipe 映射，包含 CUDA 13.3 新增的 setmaxnreg、SIMD 8x4、cvt.pack。

§ 02
Related Work

测试平台与方法学

3.1 硬件平台

Table 1.测试平台规格 / Platform Specifications

参数	值	备注
GPU	NVIDIA RTX PRO 5000 72GB	GB203 满配 / SM_120
SM 数量	110	密集 smid 空间 [0, 109]
Boost Clock	2580–2587 MHz	持续稳定，抖动 < 0.3%
DRAM	72 GB GDDR7, 384-bit, 14.0 GHz	—
DRAM 带宽 (理论 / 实测)	1344 / 1176 GB/s	87.5% 利用率
L2 Cache	96 MB	+92% vs Hopper 50 MB
TDP	350 W	实测最高 222 W (63%)
Driver	580.126.09	—
CUDA Toolkit	13.0 主要 + 13.3 PTX 9.3	新指令验证用
操作系统	Linux	`CUDA_VISIBLE_DEVICES=N` 隔离

3.2 测量方法

计时基础：clock64() 内联计时，overhead = 1 SM cycle，标准差 = 0，无需统计校正。
Anti-DCE 防护：asm volatile("" :: "r"(result)) + 写入全局 sink[] 数组。
循环控制：#pragma unroll 1 抑制自动展开，确保 SASS 结构可预测。
取值策略：每个 benchmark 执行 100+ 次，取最小值或中位数（非均值），排除 OS 中断和 DVFS 尾部噪声。
SASS 验证：每个 case 至少执行一次 cuobjdump --dump-sass 确认 PTX → SASS 映射。
SASS Binary Patching：使用 CuAssembler 直接修改控制位（stall counter, yield, scoreboard mask），绕过编译器决策。
NCU 交叉验证：dominant stall reason、pipe utilization 与微测量结果互相印证。

3.3 误差分析与可重复性

测量精度：dep-chain latency ±1 cycle，ILP sweep throughput ±2%。Boost clock 在 5 分钟持续负载下保持 2580-2587 MHz 无 throttle（抖动 < 0.3%），确保时间基准一致性。典型编译命令：

$ nvcc -O3 -lineinfo -arch=sm_120 -I include -o bin/bench bench.cu
$ nvcc -O3 -lineinfo -gencode arch=compute_120a,code=sm_120a   # for CUDA 13.3+

§ 04
Architecture Overview

SM_120 架构概览

4.1 芯片规格与定位

设计哲学
Philosophy

GB203 是 Blackwell 家族的消费级/工作站芯片，采用 SM_120（CC 12.0）。与数据中心级 GB100 (SM_100) 相比，SM_120 做出了有意识的功能裁剪：

放弃 / Removed
tcgen05、wgmma、Tensor Memory (256KB/CTA)、2-SM cooperative GEMM 保留 / Retained
mma.sync 全集、TMA、mbarrier、Cluster + DSMEM、FP6/FP4/MX、setmaxnreg (13.3+)

这种权衡反映 NVIDIA 的产品定位策略：消费级芯片优先追求高 SM 密度（110 SM）和大 L2（96 MB，弥补 GDDR7 372 cyc vs HBM ~200 cyc 的访问延迟劣势）；数据中心芯片则优先追求单 SM 吞吐（TM + wgmma）和高速互联（NVLink）。

4.2 SM_120 vs SM_100 vs SM_90 功能矩阵

Fig. 1 三代 NVIDIA 架构功能矩阵：SM_120 vs SM_100 vs SM_90

SM_120 通过保留全套 mma.sync + Blackwell 独有低精度 (FP6/FP4/MX) 实现接近 PFLOPS 级 GEMM；放弃 TM/wgmma 换取更高 SM 密度。横向比较揭示三代架构的差异化定位。

4.3 芯片拓扑：GPC → TPC → SM → Sub-core

Fig. 2 SM_120 (GB203) 芯片层次结构：4 级递归拓扑

SM_120 的 4 级层次结构：芯片 (110 SM) → 9 GPC (12 SM each) → 6 TPC (2 SM each) → 4 sub-core/SM。每个 sub-core 由 warpid % 4 静态绑定 warp，含 6 条独立管线（P0/V-pipe/P1/P2/P3/P4），三种独立实验交叉验证。

设计动机分析。 4 sub-core 设计的核心权衡是：更多 sub-core = 更多独立 warp scheduler（降低调度复杂度），但增加 die area。4 sub-core 的选择与 SM_80 (Ampere) 一致，表明 NVIDIA 在此参数上已达到最优点。静态绑定 (warpid % 4) 进一步简化了 scoreboard 设计——每个 sub-core 的 scoreboard 独立，无需跨 sub-core 同步。

§ 05
Pipelines & Issue

执行管线与指令发射

Key Finding · C3

SM_120 具有 5 条标量管线 + V-pipe 子单元，9 对 co-issue 组合全部确认。修正了 Jarmusch [1] 的"4 管线"模型。

5.1 管线端口枚举

Fig. 3 SM_120 Sub-Core 完整执行管线拓扑（5 + V-pipe）

SM_120 sub-core 的 6 条独立执行管线（5 标量 + V-pipe 子单元）。INT 与 FP32 共享 P0 前端但有独立 backend；V-pipe 共享 dispatch / 独立执行；TC (P4) 与 ALU 几乎完全 co-issue。所有通路通过 32-warp 矩阵实验交叉验证。

5.2 Co-Issue 矩阵

Fig. 4 9 对 Pipe 的 Co-Issue 矩阵热力图（实测 MIX/MAX 比率）

9 对 pipe co-issue 比率全测：绿色（≥0.95）= 完全 overlap；金色（0.7-0.95）= 部分 overlap；红色（< 0.6）= 序列化。SFU+MEM 是唯一完全冲突的组合。

5.3 ILP、管线深度与饱和分析

Table 2.各管线 ILP 拐点与饱和参数

管线	ILP knee	管线深度	Peak throughput	饱和 warp 数
FFMA (P0)	6	~6 cyc	0.585 inst/cyc	24-32
SFU (P2)	5	~5 cyc	0.11 inst/cyc	—
IADD3 (P0)	4	~4-5 cyc	0.414 inst/cyc	12
TC (P4)	4 (K-unroll)	—	20.4 cyc/mma	6-8

Dispatch 延迟（从 clock64 到第一条 IADD3 结果可用）为 3 cycle。各 pipe 的 backpressure 完全独立：同时饱和 INT pipe (12W) 和 SFU pipe (4W) 时，两者的 per-warp 延迟与单独运行时相同。

§ 06
Instruction Latency

指令延迟特性化

Range

IADD3 (2.22 cyc) 是最快标量指令；div.u64 (~419 cyc) 最慢；相差 190 倍。

6.1 整数 ALU 延迟阶梯

Fig. 5 整数 ALU 指令延迟阶梯（log scale, cycles）

10 种典型整数指令的延迟分布。div.u64 展开为 30+ 条 Newton-Raphson 软件序列（吞吐 = 延迟，无法 ILP 加速）；硬件原生 64-bit 加法 (IADD.64) 仅比 32-bit 慢 0.03 cyc。

6.2 浮点 ALU 与 SFU

Table 3.FP32 / FP64 / SFU 关键指令

PTX	SASS	延迟 (cyc)	Pipe	说明
fma.rn.f32	FFMA	4.22	P0	融合乘加
min/max.f32	FMNMX	5.19	V-pipe	独立于 FMA
neg/abs.f32	—	0	modifier	免费修饰符
rcp.approx.f32	MUFU.RCP	44.28	P2	22-bit 精度
sin/cos.f32	MUFU.SIN/COS	~22	P2	SFU 直出
div.rn.f32	—	50.76	P2+P0	RCP + 2× FFMA
fma.rn.f64	DFMA	64.13	P1	非流水线
min/max.f64	DMNMX	109	P1	异常: 走不同路径
subnormal.f32	—	+0.12	—	无 penalty
--use_fast_math sin	MUFU.SIN	5.0	P2	4.4× 加速
--use_fast_math div	MUFU.RCP	3.6	P2	14× 加速

SM_120 无 subnormal 性能惩罚（delta 0.12 cyc = noise）。FP64 仅 2 EU/SM，FP64:FP32 吞吐比 1:39（专业级配置，非消费级 1:64）。

6.3 SIMD 与类型转换（CUDA 13.3 新增）

FP16x2 / BF16x2 / FP32x2：均无真正的 SIMD 吞吐增益。HFMA2 占用 2 个 issue cycle。其价值仅在寄存器节省。
SIMD 8x4 (VIADD.U8x4)：最快的 V-pipe 指令，4.0 cyc / 1.0 inst/cyc。
cvt.pack (I2IP)：4 cyc, 2 ops/cyc @ ILP=8，高效 INT8/INT4 packing。
FP8 e4m3 pack/unpack：8 cyc / 4 cyc，为 TC 数据准备提供硬件加速。

§ 07
Memory Hierarchy

存储层级: 4 级延迟阶梯

7.1 延迟阶梯：L1 → L2 → DRAM

Fig. 6 SM_120 存储层级 4 级延迟阶梯（实测 cycles @ 2.58 GHz）

L1 与 SMEM 物理共享同一 SRAM bank，延迟一致 (33 cyc)；L2 Near vs Far 差距 101 cyc 反映 NoC 往返延迟；GDDR7 比 HBM3 慢约 87%（372 vs ~200 cyc），由 96 MB 大 L2 弥补。

7.2 Shared Memory 与 Bank Conflict

SMEM 容量 100 KB/SM（99 KB/block opt-in），32 bank，4 B/bank。无冲突延迟 34 cyc，bank conflict 惩罚 ~2 cyc/way（线性），最坏 32-way = 113 cyc。

7.3 Atomic 操作

Surprise

atomicAdd.f32 (24 cyc) 比 atomicAdd.u32 (45 cyc) 快 1.9×——L2 集成的硬件 FP32 reduction 单元让 .f32 atomic 在 L2 本地完成 RMW，无需回 SM。

跨 GPU atomic： atom.sys remote 延迟 1899 cyc（38.6× local），由 PCIe 往返延迟主导。

7.4 Cache Modifier 语义变更

SM_120 上 ld.global.cg 行为变更：绕过 L1 和 L2，直达 DRAM (372 cyc)。这与 Ampere/Hopper 不同（仅绕过 L1，仍在 L2 缓存）。开发者移植旧代码需注意。

§ 08
TensorCore

TensorCore: 统一管线的颠覆性发现

Key Finding · C2

所有 非 FP64 精度（FP16 / BF16 / TF32 / FP8 / FP6 / FP4 / MX-FP4 / NV-FP4 / Sparse / INT8 共 12 种格式）共享相同的 29 cyc 延迟 / 23 cyc 吞吐。精度选择可简化为纯存储带宽决策。

8.1 统一 TC 管线模型

Fig. 7 SM_120 TensorCore 统一流水线：12 种精度共享同一物理通路

存储带宽 决定，而非 TC 速度

SM_120 TensorCore 内部为统一 systolic array，所有非 FP64 精度共享同一物理通路。MX-FP4 在相同 23 cyc 内处理 16384 OPs（4× FP16），实现 1175 TFLOPS。FP64 是唯一例外，走独立低速通路。

8.2 MX Block-Scale 与 Sparse MMA

MX Block-Scale 零开销：Scale metadata（ue8m0 格式）编码在 B-operand 字段中，TC 内部 multiply 后通过 shift 应用 scale，无额外 pipeline stage。混合精度 MX (A=e4m3 × B=e5m2) 也是零开销。

Sparse MMA 必须使用 mma.sp::ordered_metadata：不使用此修饰符时编译/运行均正常、无 warning，但性能退化 9.16×（28 cyc → 252 cyc）。使用正确修饰符后 sparse speedup = 1.96×（98.2% of theoretical 2×）。

8.3 Operand Feeding 与 C→A Forwarding

ldmatrix.x4 dep-chain 延迟 44 cyc，throughput 8.5 cyc。通过单层 software pipeline 预取，实际 overhead 仅 +4.4 cyc/mma (15%)。
C→A forwarding 增加 +17.8 cyc/iter：TC pipeline drain 12.5 cyc + register→operand 格式转换 5.3 cyc。Fused matmul chain 应避免，改用 smem + ldmatrix 重组。

8.4 TC Co-issue 与功率

TC + 8× IADD3 = +1 cyc（近乎免费）
TC + 4× FFMA = +0 cyc（完全免费）
TC + DSMEM = 完全并行
TC + MUFU.RCP = +25 cyc/条（GEMM innerloop 应避免 rcp/rsqrt）
TC 无冷启动惩罚，无 shape-switch 惩罚

8.5 INT4 / Binary 仿真不推荐

INT4 PTX 被编译为 CALL + 2× IMMA + 数据拆分（~20 SASS），123 cyc = 4.6× INT8 成本；Binary b1 展开为 60 余条 SASS，1085 cyc = 47× HMMA。两者均不推荐使用。

§ 09
TMA & Async

TMA与异步数据移动

9.1 TMA 延迟与吞吐

Fig. 8 TMA Load / Store 严重不对称（19× 差距）

TMA Load 需完整 DRAM → L2 → SMEM 路径，而 Store 仅需 SMEM → L2 缓冲写（writeback 异步完成）。这种不对称应指导 GEMM epilogue 设计。

9.2 mbarrier 时序

核心操作（arrive / try_wait）：24 cyc，与 MMA 同级
init：58.5 cyc
全局 fence.proxy.async：186 cyc（昂贵）
CTA scope fence：30 cyc（轻量）

9.3 Pipeline Depth 与带宽饱和

达到 90% peak 需 ≥ 32 outstanding TMA 操作。Multi-CTA scaling 完美线性：110 CTA 聚合 1483 GB/s（超 DRAM 1024 GB/s，因 L2 命中）。饱和 DRAM 需 ~76 CTAs。

9.4 不可用功能

SM_120 不支持 TMA multicast（~10000× 退化）：消费级 PCIe 拓扑无 NVLink/SM-broadcast 网络，软件回退到 serial unicast。
不支持 TMA im2col，cluster 间 cp.async.bulk smem→smem 是 illegal instruction。

9.5 CUDA 13.3 新增

tensormap.replace：26-45 cyc/op，支持动态修改 TMA 参数
L2::cache_hint (UTMACCTL)：支持 L2 缓存策略控制

§ 10
Warp Scheduler

Warp 调度器与控制位逆向

Key Finding · C7

通过 CuAssembler binary patching，验证 SM_120 的 128-bit 指令编码中 17-bit 控制码全语义；scoreboard 深度 ≥ 12（超越 Huerta [3] 的 6-SB 模型）；stall counter 严格 1:1 线性，错误值导致 silent data corruption。

10.1 指令编码：128-bit 格式

Fig. 9 SM_120 指令编码：128-bit 固定宽度 + 17-bit Control Code

SM_120 沿用 Hopper 的 128-bit 固定指令格式：64-bit 操作数信息 + 41-bit 扩展 + 17-bit 控制码 + 保留位。控制码精确编排发射时机：stall counter 是编译器信任的不变量，硬件不互锁。

10.2 Stall Counter 严格线性与数据损坏

SM_120 不使用硬件互锁：stall counter < instruction latency 会导致读取 stale register、无报错 / 无 trap。手动修改 SASS 控制码时必须确保 stall ≥ instruction latency。

stall = 0：hardware scoreboard fallback (~25 cyc/inst, 正确但极慢)
stall = 1-3 (IADD3 dep chain)：RAW 数据损坏
stall = 4：正确（= IADD3 write-back latency）
stall linearity：精确 1:1（每 +1 stall = +1 cycle per payload instruction）

10.3 CGGTY 5-Warp 阈值

Fig. 10 CGGTY 调度阈值：4 → 5 warp 出现 6× 加速跳跃

FFMA 依赖链 cyc/iter vs warp 数：1-4 warp 时线性增长（无并发），5 warps 时跳跃至 45 cyc，因为至少一个 sub-core 现在有 2 warp，调度器可以在 stall 时切换。开发者启示：每 block ≥ 64 threads。

10.4 Yield Bit 完全阻断 latency hiding

yield = 1 时所有 1-32 warps 配置均为 336 cyc/iter，调度器不再交织。Yield 将 warp 移出 eligible set——人为减少活跃 warp，破坏 latency hiding。

10.5 Scoreboard 深度 ≥ 12

12 个独立 global load 完全并行追踪，无 stall（62.42 cyc 恒定）。Load queue depth ≥ 32。这修正了 Huerta [3] 的 6-SB 模型，对深度 prefetch 和多级 pipeline overlap 极为有利。

10.6 Operand Collector 与 Reuse Bit

Operand collector 完全消解 bank conflict：≤ 8 warps 零可测量惩罚，16 warps 仅 +1.08 cyc (3.3%)
Reuse bit (.reuse) 是 hardware no-op：有/无 .reuse 性能完全相同（28.00 vs 28.00 cyc/iter）。编译器仍设置 .reuse flags，硬件忽略。

10.7 setmaxnreg 动态寄存器重分配（CUDA 13.3）

CUDA 13.3 解锁 setmaxnreg：inc ~61 cyc/op（USETMAXREG.TRY_ALLOC.CTAPOOL），dec ~50 cyc/op（DEALLOC.CTAPOOL）。这支持 CUTLASS persistent kernel 的 warp specialization：producer warps 用少量 regs（数据搬运），consumer warps 动态扩展 regs 给 MMA accumulators。

§ 11
System Analysis

系统级特性与应用

11.1 GEMM 性能上限：cuBLAS vs CUTLASS

Key Finding · C4

CUTLASS 4.5 NVFP4 达 975 TFLOPS（80.7% TC 效率），是 SM_120 最高实测 GEMM 吞吐。证明 mma.sync + Cluster Launch Control 路径在消费级 Blackwell 上仍可接近 1 PFLOPS。

Fig. 11 cuBLAS / CUTLASS 4.5 GEMM 峰值 TFLOPS（M=N=K=8192, A100×8 forms）

CUTLASS 4.5 GeForce Blackwell kernel 使用 CLC + setmaxnreg + warp specialization，证明 SM_120 虽无 TM/wgmma，仍支持完整的 CUTLASS persistent kernel 栈。cuBLAS 不支持 MXFP4 GEMM——CUTLASS 填补了这一缺口。

11.2 FlashAttention 瓶颈分析

Key Finding · C5

SM_120 上 FlashAttention 是 ALU-bound，而非 TC-bound。Softmax 占 63.1%（exp2 主导），TC 仅 36.9%——瓶颈在 SFU 容量（4 单元/SM）。

Fig. 12 FlashAttention 单 KV-tile 时间分解（fused attention, 32 KV tiles 稳态）

在 32 KV tiles 稳态下，每 tile 时间分解：softmax 63.1% / 2× MMA 36.9%。expf（编译为 MUFU.EX2，SFU pipe）是最大成本，5 次 expf ≈ 100 cycles。这与 "attention is compute-bound" 的常识相悖——SM_120 上 attention 的真实瓶颈是 SFU。

11.3 功耗模型与能效

Key Finding · C6

功耗模型 P = 80W (基础设施) + 0.7-1.2W/SM。基础设施占比 > 45%。MX-FP4 能效达 6.9 TFLOPS/W，是 FP32 的 28 倍。

Fig. 13 SM_120 功耗三层结构分解（活跃 SM 数与总功耗）

基础设施功耗（时钟域 / L2 / NoC / MC）固定 80W，占满载的 36-45%。增加每个 SM 仅 +1W 计算功耗——优化策略应优先减少活跃 GPC 数而非降低 per-SM 利用率。

Fig. 14 每精度能效（TFLOPS/W）：MX-FP4 比 FP32 高 28 倍

MX-FP4 在相同 23 cyc 内处理 4× FP16 OPs，但 TC 功耗（开关功耗由物理结构决定）几乎不变。吞吐 4× 而功耗不变 → 能效 4×。

11.4 DVFS 与最优能效频率

Fig. 15 DVFS：频率-吞吐线性 vs 频率-功耗超线性

频率-吞吐完美线性，但功率超线性（V²f 模型）。最佳能效频率 1500-1800 MHz；boost (2580 MHz) 提供 +73% 吞吐但 +127% 功耗，对延迟敏感场景值得，对能效敏感场景应保守。

11.5 多 SM 通信

Table 4.SM 间 / GPU 间通信延迟

路径	延迟 / 带宽	说明
DSMEM read (cluster)	230 cyc	5.9× local SMEM
DSMEM write	36 cyc	≈ local SMEM
Cluster barrier	408 cyc	不随 cluster size 变
Inter-SM atomic exchange	~1240 cyc 往返	所有 SM pair 均匀（无 GPC 局部性）
PCIe Gen 5 H2D	57 GB/s	88.9% 利用率
PCIe Bidirectional	102 GB/s	—
P2P intra-switch	53 GB/s	93%
P2P cross-switch	39 GB/s	68%
Cross-GPU atomic	1899 cyc	50× local
PDL kernel gap	0.56 μs	vs 1.04 μs serial

§ 12
Discussion

讨论：消费级 GPU 的设计取舍

12.1 设计哲学：消费级 vs 数据中心

SM_120 与 SM_100 (B100/B200) 共享 Blackwell 架构基因，但两者在关键维度上做出了不同取舍。本研究的实验平台为 NVIDIA RTX PRO 5000 72GB (GB203, SM_120) 工作站级 GPU；理解这些取舍有助于在 SM_120 上正确 “借鉴” B100/B200 kernel 的设计经验，避免直接套用导致的性能陷阱。

维度	SM_120 (RTX PRO 5000 72GB)	SM_100 (B100/B200)	设计意图
FP64 Throughput	~2 TFLOPS (1:64)	~30 TFLOPS (1:2)	游戏不需 FP64，省晶体管
L2 Cache	96 MB (统一)	100 MB ×2 (双 die)	单 die 简化制程
DRAM	GDDR7 72 GB / 384-bit	HBM3e 192 GB	工作站领域需中容量低成本
NVLink	无 (PCIe Gen5)	NVLink 5 (1.8 TB/s)	单卡定位，无多卡训练
SM 数量	110 SM (smid 0–109)	132+132 (双 die)	单 die 高密度布局
FP4 Peak	975 TFLOPS	~5000 TFLOPS	5× 差距，主要来自 SM 数
TDP	350 W (实测峰值 222 W, 63%)	700-1000W	工作站水冷/隐藏限功

Table 5：SM_120 vs SM_100 关键维度对比

12.2 Kernel 开发的 7 条实用启示

基于本研究的 cycle 级测量，我们提炼出 7 条针对 SM_120 的 kernel 开发原则。这些原则在 RTX PRO 5000 72GB 上经过 CUTLASS 4.5、FlashAttention 2、自研 GEMM 验证，可作为优化的起点。

PRINCIPLE 01

优先 5+ warps/sub-core 占用率

CGGTY 阈值在 5 warps 处出现 6× 跳跃 (268→45 cyc)。少于 5 warps 时调度器无法隐藏 latency；多于 5 warps 后收益递减。20 warps/SM 是最佳启动配置。

PRINCIPLE 02

利用 5+V-pipe Co-issue

FMA + INT、FMA + LDST、TC + ALU 可双发射。手写 SASS 或使用 __pipeline_acquire 等 intrinsic 显式控制 pipe 分配，吞吐 +10-30%。

PRINCIPLE 03

MMA Reuse 必须开启

A/B 操作数 reuse 节省 35-45% 寄存器读端口。CUTLASS 4.5 默认启用 kReuseA + kReuseB，自定义 kernel 务必显式声明。

PRINCIPLE 04

TMA 仅用于 Load

TMA Store 的 33 cyc 优势是假象——它仅"提交"，实际写回需 620 cyc。Output 写回建议用 cp.async.bulk + barrier 显式同步。

PRINCIPLE 05

Scoreboard depth ≥ 12

Huerta 等 [3] 的 6-SB 模型在 SM_120 上不再适用。深 pipeline 需要 12+ 个 outstanding load，使用 __pipeline_commit + 多级 buffer 实现。

PRINCIPLE 06

FlashAttention: ALU bound

在 SM_120 上 FA 的 63.1% 时间在 softmax (ALU)，而非 TC。优化方向：降低 softmax 精度 (FP16 → BF16) 或 fuse rescale 到 TC 后处理。

PRINCIPLE 07

DVFS sweet spot 1500-1800 MHz

超过 2000 MHz 后功率增长超线性 (P ∝ f^1.5-2.0，V²f 主导)，能效骤降。DVFS sweet spot 位于 1500-1800 MHz，相对 boost 频率可获 +18% perf/W。

12.3 跨代趋势：SM_120 vs SM_90 (Hopper)

以 SM_90 (Hopper, H100) 作为参考基线，横向对比 SM_120 在延迟、带宽、容量、并发度等 9 个核心维度的变化。SM_120 在存储以及新增精度上取得进步，但在 SFU 处理、mma.sync 及并发度上为成本/功耗做出了取舍。

指标	SM_120 (本研究)	SM_90 (参考值)	变化
FFMA 延迟	4.22 cyc	~4 cyc	持平
L1 延迟	33 cyc	28–33 cyc	持平
DRAM 延迟	372 cyc (GDDR7)	~430 cyc (HBM3)	改善 14%
L2 容量	96 MB	50 MB	+92%
SFU (rcp)	44 cyc	~28 cyc	退化 57%
mma.sync FP16	29/23 cyc	~16/8 cyc	退化 ~3×
Max warps/SM	48	64	−25%
FP64 : FP32	1:39 (pro)	1:2 (datacenter)	受限
新增精度	FP6 / FP4 / MX	—	SM_120 独有

Table 6：SM_120 vs SM_90 关键指标对比 (蒙克色列为本研究测量值)

趋势小结：ALU 延迟已达物理极限 (~4 cyc)，进一步提升需从管线并行度 (更多 sub-unit) 或数据密度 (更低精度) 着手。SM_120 的 MX-FP4 正是后一方向的体现。

12.4 方法学局限性

本研究的测量方法和结果存在以下局限，未来工作可针对性改进：

无 SM_100 直接对比：本文的跨代对比依赖 Jarmusch [1,2] 和 Jia [7,8] 的文献数据，未能在 SM_100 (B100/B200) 上直接验证 tcgen05 / Tensor Memory 路径。
部分 WHY 解释基于推理：对于无法直接观测的微架构内部结构 (如 TC systolic 深度、L2 slice 数量)，我们提供了合理的推理模型并标注了确信度。
CUDA 版本差异：核心数据在 CUDA 13.0 上测得，CUDA 13.3 补充了 PTX ISA 9.3 新功能验证。两个版本在已有指令的延迟测量上一致。
单芯片测试：所有数据来自 RTX PRO 5000 72GB (GB203 满配，110 SM)，不同 SKU (如 RTX 5080 的 84 SM 裁剪版本) 的部分参数可能不同。

§ 13
Conclusion

结论与未来方向

本文对 NVIDIA Blackwell 消费级/工作站 GPU (GB203, SM_120) 进行了系统性的 cycle 级微架构特性研究，覆盖 21 个测试类别、85 个以上的 case 报告、约 879 个跟踪项 (96.1% 覆盖率)，实验平台为 NVIDIA RTX PRO 5000 72GB。所有发现为 SM_120 的 kernel 开发和性能优化提供了实证基础。

13.1 主要发现总结

FINDING 01 · TensorCore

统一 TC 管线

12 种非 FP64 精度共享 29/23 cyc，精度选择可简化为存储带宽决策。

FINDING 02 · Pipeline

5+V-pipe 管线架构

修正了此前 “4 管线” 模型，首次量化 V-pipe 独立性。

FINDING 03 · GEMM

无 TM 近 PFLOPS

CUTLASS NVFP4 达 975 TFLOPS，证明 mma.sync + CLC 路径的高效性。

FINDING 04 · FlashAttention

FA 是 ALU-bound

softmax 63.1%, TC 36.9% — SFU 是真正瓶颈，不是 TensorCore。

FINDING 05 · Power

per-SM 功耗模型

P = 80W + 0.7-1.2W/SM；MX-FP4 能效 6.9 TFLOPS/W (FP32 的 28×)。

FINDING 06 · Scoreboard

Scoreboard depth ≥ 12

超越经典 6-SB 模型 (Huerta [3])，利好深度 prefetch 与多级 buffer。

FINDING 07 · Safety

Stall counter 安全限

stall < latency 会造成 silent data corruption，无硬件互锁 — SASS patching 需严格遵守。

13.2 未来工作

本研究为 SM_120 建立了基础画像，以下方向值得进一步探索：

DIRECTION 01

SM_100 (B200) 直接对比

获取 B200 硬件后，直接测量 tcgen05/TM 路径，与 SM_120 的 mma.sync 路径进行公平对比。

DIRECTION 02

SM_90 (Hopper) 基线

在 H100/H200 上复现本文的全部测试，建立完整的跨代对比数据库。

DIRECTION 03

Triton / CUTLASS 优化

基于本文的管线模型和延迟数据，为 SM_120 定制 Triton kernel auto-tuning 策略。

DIRECTION 04

功耗建模深化

将经验功耗模型 (P = 80W + 0.7-1.2W/SM) 与 AccelWattch [15] 等模拟器模型进行对比验证。

DIRECTION 05

FlashAttention 优化

针对 SM_120 的 SFU 瓶颈，探索 polynomial approximation 替代 expf 的性能/精度权衡。

Refs
Bibliography

参考文献

[1]Jarmusch, J., et al. Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks. arXiv:2507.10789, 2025.
[2]Jarmusch, J., et al. Microbenchmarking NVIDIA's Blackwell Architecture. arXiv:2512.02189v3, 2026.
[3]Huerta, S., et al. Analyzing Modern NVIDIA GPU cores. arXiv:2503.20481, 2025.
[4]Wong, H., et al. Demystifying GPU Microarchitecture through Microbenchmarking. IEEE ISPASS, 2010.
[5]Mei, X. and Chu, X. Dissecting GPU Memory Hierarchy through Microbenchmarking. IEEE TPDS, 2017.
[6]Volkov, V. Understanding Latency Hiding on GPUs. PhD Thesis, UC Berkeley, 2016.
[7]Jia, Z., et al. Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking. arXiv:1804.06826, 2018.
[8]Jia, Z., et al. Dissecting the Ampere GPU Architecture through Microbenchmarking. GPU Technology Conference, 2022.
[9]Raihan, M. A., et al. Modeling Deep Learning Accelerator Enabled GPUs. IEEE ISPASS, 2019.
[10]Yan, D., et al. Demystifying Tensor Cores to Optimize Half-Precision Matrix Multiply. IEEE IPDPS, 2020.
[11]Dao, T., et al. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. NeurIPS, 2022.
[12]Dao, T. FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning. ICLR, 2024.
[13]Thakkar, V., et al. CUTLASS: CUDA Templates for Linear Algebra Subroutines and Solvers. NVIDIA, 2023.
[14]Leng, J., et al. GPUWattch: Enabling Energy Optimizations in GPGPUs. ISCA, 2013.
[15]Kandiah, V., et al. AccelWattch: A Power Modeling Framework for Modern GPUs. MICRO, 2021.
[16]NVIDIA Corporation. NVIDIA Blackwell Architecture Technical Brief. 2024.
[17]NVIDIA Corporation. CUDA C++ Programming Guide. v13.0, 2026.
[18]NVIDIA Corporation. Parallel Thread Execution ISA. v9.3, 2026.

App. C
Validation Matrix

附录 C：新功能验证表

本附录对应原文表 10，列出 CUDA 13.3 / PTX ISA 9.3 在 SM_120 上引入或新增能力的 8 项关键功能，以及它们在本研究中对应的 Section/Case、关键数据与编译要求。所有数据均在 RTX PRO 5000 72GB (GB203, SM_120) 上测得，编译器为 NVCC CUDA 13.0（主要）配合 13.3 用于 PTX 9.3 新指令补充验证。

功能	Section/Case	关键数据	编译要求
`setmaxnreg.inc/dec`	Sec 12 / 06	inc ~61 cyc, dec ~50 cyc	`compute_120a` + CUDA 13.3
`SIMD 8×4`	Sec 4 / 15	4.0 cyc, 1.0 inst/cyc, V-pipe	`compute_120a` + CUDA 13.3
`cvt.pack` (I2IP)	Sec 6 / 08	4 cyc, 2 ops/cyc @ ILP=8	`compute_120a` + CUDA 13.3
`FP8 cvt` (e4m3)	Sec 6 / 09	pack 8 cyc, unpack 4 cyc	`compute_120a` + CUDA 13.3
`MX narrow cvt`	Sec 6 / 10	8.22 cyc round-trip	`compute_120a` + CUDA 13.3
`tensormap.replace`	Sec 10	26-45 cyc/op	`compute_120a` + CUDA 13.3
`L2::cache_hint`	Sec 10	UTMACCTL + cp.async.bulk	`compute_120a` + CUDA 13.3
CUTLASS 4.5 GeForce	Sec 18 / 05	NVFP4 975 T, CLC + warp spec	CUDA 13.3 + CUTLASS 4.5

Table 10：CUDA 13.3 / PTX ISA 9.3 在 SM_120 上新增功能的本研究验证清单（8 项）。

关键观察

本研究 8 项 CUDA 13.3 新功能均依赖 compute_120a 架构特化标记编译；前 7 项（setmaxnreg、SIMD 8×4、各类 cvt/cvt.pack、tensormap.replace、L2::cache_hint）通过 NVCC 13.3 PTX 9.3 直接验证；CUTLASS 4.5 GeForce 路径额外要求 CUTLASS 4.5 与 NVFP4 + CLC + warp specialization 协同，实测达到 975 TFLOPS 的非张量峰值。

目录 / Contents

引言: 为什么是 SM_120

1.1 研究背景与动机

1.2 主要贡献 (Contributions)

首个 SM_120 系统级微架构特性研究

统一 TC 流水线发现

"5 管线 + V-pipe" 执行架构

无 TM/tcgen05 的 GEMM 上限

FlashAttention 是 ALU-bound

首个消费级 per-SM 功耗模型

深度逆向 Warp 调度器

完整 SM_120 指令延迟参考表

相关工作

2.1 GPU 微架构 microbenchmark 方法论

2.2 NVIDIA 架构演进特性化

2.3 TensorCore 与低精度计算

2.4 功耗建模

测试平台与方法学

3.1 硬件平台

3.2 测量方法

3.3 误差分析与可重复性

SM_120 架构概览

4.1 芯片规格与定位

4.2 SM_120 vs SM_100 vs SM_90 功能矩阵

4.3 芯片拓扑：GPC → TPC → SM → Sub-core

执行管线与指令发射

5.1 管线端口枚举

5.2 Co-Issue 矩阵

5.3 ILP、管线深度与饱和分析

指令延迟特性化

6.1 整数 ALU 延迟阶梯

6.2 浮点 ALU 与 SFU

6.3 SIMD 与类型转换（CUDA 13.3 新增）

存储层级: 4 级延迟阶梯

7.1 延迟阶梯：L1 → L2 → DRAM

7.2 Shared Memory 与 Bank Conflict

7.3 Atomic 操作

7.4 Cache Modifier 语义变更

TensorCore: 统一管线的颠覆性发现

8.1 统一 TC 管线模型

8.2 MX Block-Scale 与 Sparse MMA

8.3 Operand Feeding 与 C→A Forwarding

8.4 TC Co-issue 与功率

8.5 INT4 / Binary 仿真不推荐

TMA与异步数据移动

9.1 TMA 延迟与吞吐

9.2 mbarrier 时序

9.3 Pipeline Depth 与带宽饱和

9.4 不可用功能

9.5 CUDA 13.3 新增

Warp 调度器与控制位逆向

10.1 指令编码：128-bit 格式

10.2 Stall Counter 严格线性与数据损坏

10.3 CGGTY 5-Warp 阈值

10.4 Yield Bit 完全阻断 latency hiding

10.5 Scoreboard 深度 ≥ 12

10.6 Operand Collector 与 Reuse Bit

10.7 setmaxnreg 动态寄存器重分配（CUDA 13.3）

系统级特性与应用

11.1 GEMM 性能上限：cuBLAS vs CUTLASS

11.2 FlashAttention 瓶颈分析

11.3 功耗模型与能效

11.4 DVFS 与最优能效频率

11.5 多 SM 通信

讨论：消费级 GPU 的设计取舍

12.1 设计哲学：消费级 vs 数据中心

12.2 Kernel 开发的 7 条实用启示

12.3 跨代趋势：SM_120 vs SM_90 (Hopper)

12.4 方法学局限性

结论与未来方向

13.1 主要发现总结

统一 TC 管线

5+V-pipe 管线架构

无 TM 近 PFLOPS

FA 是 ALU-bound

per-SM 功耗模型

Scoreboard depth ≥ 12

Stall counter 安全限

13.2 未来工作

参考文献