# Dissecting the SM_120 Microarchitecture

> Source: <https://zartbot.github.io/micro_arch/nvidia/sm_120/paper.html>
> Published: 2026-05-29 12:53:23+00:00

A Cycle-Level Characterization of NVIDIA Blackwell Consumer GPUs

—— 21 个测试类别 · 879 项 trace · 96.1% 覆盖率 · RTX PRO 5000
72GB

本文对 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 框架优化者提供了实用参考。

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）。

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

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

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

CUTLASS 4.5 NVFP4 达 **975 TFLOPS**（80.7%
效率）。证明 `mma.sync + CLC`

路径在消费级 Blackwell
上仍可接近 1 PFLOPS。

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

`P = 80W 基础设施 + 0.7-1.2W/SM`

，基础设施占比
**> 45%**。MX-FP4 能效达 6.9 TFLOPS/W（FP32 的
28×）。

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

覆盖 **60+ 指令类型**的 latency/throughput/pipe
映射，包含 CUDA 13.3 新增的 `setmaxnreg`

、SIMD
8x4、`cvt.pack`

。

系统性 GPU 微架构特性化源于 Wong 等人 [4] 对 GT200
的开创性工作，建立了 pointer-chase 延迟、ILP sweep、bank conflict
探测等标准技术。Mei 和 Chu [5] 深化了存储层级探测；Volkov [6]
从理论层面分析 ILP/TLP 与 latency hiding
的权衡。本文在此基础上引入
**CuAssembler binary patching**、** multi-warp co-issue 矩阵测量**、** NCU stall 交叉验证**。

Jia 等人 [7,8] 系统特性化了 Volta 和 Ampere，建立了跨代对比基准（Volta FFMA ~4 cyc，Ampere L1 ~28 cyc）。针对 Blackwell：

Raihan [9]、Yan [10] 研究了 Volta/Turing TC
管线模型；FlashAttention [11,12] 使 TC 利用率成为推理瓶颈；CUTLASS
[13] 持续推动 GEMM 效率向理论峰值靠近。本文首次在**消费级 Blackwell**
上量化了 TC 统一管线特性和 FlashAttention 的 ALU-bound 瓶颈。

GPUWattch [14]、AccelWattch [15]
基于模拟器建立功耗模型，但缺乏真实硬件验证。本文通过**直接测量**（nvidia-smi）建立
SM_120 经验功耗模型，补充模拟器方法的盲区。

| 参数 | 值 | 备注 |
|---|---|---|
| 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` 隔离 |

`clock64()`

内联计时，overhead
= 1 SM cycle，标准差 = 0，`asm volatile("" :: "r"(result))`

+ 写入全局 `sink[]`

数组。
`#pragma unroll 1`

抑制自动展开，确保 SASS 结构可预测。
`cuobjdump --dump-sass`

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

``` bash
$ 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+
```

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 sub-core 设计的核心权衡是：更多
sub-core = 更多独立 warp scheduler（降低调度复杂度），但增加 die
area。4 sub-core 的选择与 SM_80 (Ampere) 一致，表明 NVIDIA
在此参数上已达到最优点。静态绑定 (`warpid % 4`

)
进一步简化了 scoreboard 设计——每个 sub-core 的 scoreboard 独立，无需跨
sub-core 同步。

| 管线 | 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 延迟与单独运行时相同。

`IADD3`

(2.22 cyc) 是最快标量指令；`div.u64`

(~419 cyc) 最慢；`div.u64`

展开为 30+ 条
Newton-Raphson 软件序列（吞吐 = 延迟，无法 ILP 加速）；硬件原生
64-bit 加法 (`IADD.64`

) 仅比 32-bit 慢 0.03 cyc。
| 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）。

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。

`atomicAdd.f32`

(24 cyc) 比
`atomicAdd.u32`

(45 cyc) `atom.sys`

remote 延迟
`ld.global.cg`

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

`mma.sp::ordered_metadata`

`ldmatrix.x4`

dep-chain 延迟 44 cyc，throughput 8.5
cyc。通过单层 software pipeline 预取，实际 overhead 仅
INT4 PTX 被编译为 CALL + 2× IMMA + 数据拆分（~20 SASS），**123 cyc = 4.6× INT8 成本**；Binary b1 展开为 60 余条 SASS，** 1085 cyc = 47× HMMA**。两者均不推荐使用。

`fence.proxy.async`

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

`tensormap.replace`

：26-45 cyc/op，支持动态修改 TMA 参数
`L2::cache_hint`

(UTMACCTL)：支持 L2 缓存策略控制
yield = 1 时所有 1-32 warps 配置均为
**336 cyc/iter**，调度器不再交织。** Yield 将 warp 移出 eligible set**——人为减少活跃 warp，破坏 latency hiding。

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

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**。

`mma.sync + Cluster Launch Control`

路径在`expf`

（编译为 MUFU.EX2，SFU pipe）是最大成本，5
次 expf ≈ 100 cycles。这与 "attention is compute-bound"
的常识相悖——SM_120 上 attention 的真实瓶颈是 SFU。
`P = 80W (基础设施) + 0.7-1.2W/SM`

。基础设施占比
| 路径 | 延迟 / 带宽 | 说明 |
|---|---|---|
| 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 |

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 | 工作站水冷/隐藏限功 |

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

`__pipeline_acquire`

等 intrinsic 显式控制 pipe 分配，吞吐 +10-30%。`kReuseA + kReuseB`

，自定义 kernel 务必显式声明。`cp.async.bulk`

+ barrier 显式同步。`__pipeline_commit`

+ 多级 buffer 实现。以 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 独有 |

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

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

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

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

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

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

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

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

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

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

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

本附录对应原文表 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 |

本研究 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 的非张量峰值。