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%），确保时间基准一致性。典型编译命令：
$ 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 的非张量峰值。