Dissecting the SM_120 Microarchitecture

NVIDIA's Blackwell consumer GPU (GB203/SM_120) features a unified TensorCore pipeline where all 12 non-FP64 precision formats share identical 29-cycle latency and 23-cycle throughput, reducing precision selection to a pure memory bandwidth decision. A comprehensive cycle-level characterization of the RTX PRO 5000 72GB revealed a "5 pipeline + V-pipe" architecture correcting prior "4 pipeline" models, achieved 975 TFLOPS (80.7% efficiency) using CUTLASS 4.5 NVFP4, and established the first per-SM power decomposition model showing MX-FP4 delivers 6.9 TFLOPS/W — 28 times FP32 efficiency. The study also found FlashAttention on SM_120 is ALU-bound with softmax consuming 63.1% of cycles versus TensorCore's 36.9%, shifting the bottleneck to SFU units.

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