Research Paper · 2026 · GB203 / SM_120

Dissecting the SM_120 Microarchitecture

A Cycle-Level Characterization of NVIDIA Blackwell Consumer GPUs
—— 21 个测试类别 · 879 项 trace · 96.1% 覆盖率 · RTX PRO 5000 72GB

NVIDIA Blackwell · GB203
SM_120 · CC 12.0
CUDA 13.0 / 13.3 · PTX ISA 9.3
2.58 GHz (持续无 throttle)
96 MB · 1176 GB/s GDDR7
975TFLOPS
CUTLASS NVFP4 GEMM 峰值
29 / 23cyc
TC 统一延迟 / 吞吐
6.9TF/W
MX-FP4 能效
63.1%
FlashAttn ALU 占比
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

目录 / Contents

13 章 + 3 附录
§ 01引言 / Introduction § 02相关工作 / Related Work § 03测试平台与方法学 § 04SM_120 架构概览 § 05执行管线与指令发射 § 06指令延迟特性化 § 07存储层级 § 08TensorCore § 09TMA 与异步数据移动 § 10Warp 调度器与控制位 § 11系统级特性与应用 § 12讨论 / Discussion § 13结论 / Conclusion REF参考文献 APP附录 A · B · C
§ 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

相关工作

定位
Positioning

2.1 GPU 微架构 microbenchmark 方法论

系统性 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 交叉验证。

2.2 NVIDIA 架构演进特性化

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

  • Jarmusch et al. [1]：对同一芯片 GB203 (RTX 5080) 进行了初步特性化，但仅报告"4 管线"模型，未识别 V-pipe 独立性，未涉及 TC co-issue / FlashAttention / 功耗模型。
  • Jarmusch et al. [2]：后续工作转向 B200 (SM_100)，提供 tcgen05 + TM 数据，与本文形成互补。
  • Huerta et al. [3]：提供调度器分析方法学（CGGTY 阈值、控制位逆向）。本文扩展其方法到 SM_120 并纠正 scoreboard 深度模型（6 → ≥12）。

2.3 TensorCore 与低精度计算

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

2.4 功耗建模

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

§ 03
Methodology

测试平台与方法学

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 测量方法

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
FEATURE SM_120 · GB203 RTX PRO 5000 / RTX 5080 SM_100 · GB100/200 B200 (Datacenter) SM_90 · GH100 H100 (Hopper) TC API mma.sync only mma.sync + wgmma + tcgen05 mma.sync + wgmma Tensor Memory ✗ 不支持 ✓ 256 KB/CTA ✗ FP6 / FP4 dense MMA ✓ 588 TFLOPS ✓ ✗ MX block-scale MMA ✓ 1175 TFLOPS ✓ ✗ TMA Multicast ✗ ~10000× 退化 ✓ ✓ Cluster Max Size 12 (= GPC SM) 16 16 L2 Cache 96 MB 126 MB 50 MB DRAM Type GDDR7, 372 cyc HBM3e, ~180 cyc HBM3, ~200 cyc setmaxnreg (CUDA 13.3) ✓ inc 61 / dec 50 cyc ✓ ✗ (Hopper 之后) FP64 : FP32 Ratio 1 : 39 (pro) 1 : 2 (datacenter) 1 : 2 SM 数量 (典型) 110 148 132
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 级递归拓扑
GB203 — NVIDIA Blackwell · 110 SM · 96 MB L2 · GDDR7 384-bit Compute Capability 12.0 · 2.58 GHz Boost · 350 W TDP L2 — 9× GPC (GRAPHICS PROCESSING CLUSTER) GPC 0 12 SM · 6 TPC GPC 1 12 SM GPC 2 12 SM GPC 3 12 SM GPC 4 ★ 详情下展开 GPC 5 12 SM GPC 6 12 SM GPC 7 12 SM GPC 8 ~2 SM L3 — GPC 4 内部：6× TPC (TEXTURE PROCESSING CLUSTER) TPC 0 2 SM TPC 1 2 SM TPC 2 ★ SM N · SM N+1 TPC 3 2 SM TPC 4 2 SM TPC 5 2 SM L4 — SM 内部：4× SUB-CORE (WARP SCHEDULER PARTITION) Streaming Multiprocessor (SM) SUB-CORE 0 P0 · INT/FP32 (FMA) └ V-pipe (SIMD) P1 · FP64 (×2 EU) P2 · SFU/MUFU P3 · LSU P4 · TensorCore SUB-CORE 1 warp → warpid % 4 == 1 独立 Warp Scheduler 独立 Scoreboard ≥ 12 独立 Operand Collector 独立 RF (4 bank) peak 0.5 inst/cyc SUB-CORE 2 warp → warpid % 4 == 2 CGGTY 阈值 = 5 warps ≤4 warp: 无 latency hiding ≥5 warp: 6× 加速 → 每 sub-core ≥ 2 warps SUB-CORE 3 warp → warpid % 4 == 3 round-robin 调度 warp 切换 = 0 cycle issue 宽度: 2/cycle no triple-issue VERIFIED BY: FFMA dep-chain warp scaling · %warpid readback · per-pipe backpressure
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）
FRONTEND EXECUTION BACKEND Warp Scheduler round-robin 12 warps/sub-core Dispatch 2 ports / cycle no triple-issue Operand Collector 4 bank RF 透明 bank conflict Scoreboard depth ≥ 12 stall counter linear P0 unified INT / FP32 Pipe FFMA · IADD3 · IMAD · LOP3 · DP4A latency 2.22-4.22 cyc · throughput 0.41 inst/cyc V-pipe ★ NEW SIMD / Min-Max sub-unit VIADD.16x2 · VIADD.U8x4 · FMNMX · VIMNMX +61% peak vs IADD3 · 共享 dispatch / 独立 backend P1 FP64 Pipe DFMA · DMNMX · DMMA 2 EU · 非流水线 P2 SFU / MUFU RCP · SIN · COS · EX2 · LG2 4 SFU/SM · ★ FA bottleneck P3 LSU (Load/Store) LDG · STG · LDS · STS · ATOM · LDMATRIX P4 Tensor Pipe HMMA · QMMA · OMMA · IMMA · QMMA.SF 29 / 23 cyc 统一 SM-LEVEL PEAK IPC (32 WARPS): FFMA = 2.07 IADD3 = 2.0 SFU = 0.11 SMEM = 0.5 MMA = 0.25 CO-ISSUE PROPERTIES FP64 ⊥ FP32 100% overlap P1 完全独立于 P0 (MIX/MAX = 1.002) SFU ⊥ FP32 100% overlap 独立 dispatch slot, ratio = 1.000 INT + FP32 共享 P0 前端 独立 backend → 混合 +19% vs 纯 IADD3 TC + ALU ≈ free 8× IADD3 +1 cyc · 4× FFMA +0 cyc SFU + ALU + LDS = 序列化 +8-10% overhead → 仅支持 2-way co-issue 单 warp issue: 0.55 inst/cyc ≈ 1.81 cyc/inst, dispatch round-robin CGGTY 5-warp threshold 4→5 warps: 268→45 cyc/iter (6× speedup)
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 比率）
INT FP32 FP64 SFU MEM TC V-pipe INT (IADD3) FP32 (FFMA) FP64 (DFMA) SFU (MUFU) MEM (LDG) TC (HMMA) — 1.00 1.00 1.00 1.00 0.80 0.85 1.00 — 1.002 1.000 1.00 0.95 0.85 1.00 1.002 — 0.85 0.90 0.78 0.95 1.02 1.00 0.85 — 0.55 0.70 0.85 1.00 1.00 0.90 0.55 — 0.95 0.85 0.80 0.96 0.78 0.70 0.95 — 0.85 CO-ISSUE EFFICIENCY 0.5 0.85 1.0+ (full overlap) ★ FP64⊥FP32, SFU⊥FP32 = 100% overlap ★ SFU+MEM = 0.55 (序列化)
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）
1000 100 10 1 CYCLES (LOG SCALE) 2.22 IADD3 P0 4.0 VIADD.U8x4 V-pipe ★ 4.22 IMAD / DP4A P0 / LOP3 4.25 IADD.64 native 5.25 VIADD.16x2 V-pipe 9.13 IMAD.HI 2-pass 10.06 POPC throttled 16.25 mul.lo.u64 3×IMAD ~60 div.u32 MUFU+IMAD ~419 div.u64 30+IMAD ▎ 单管线最快 ▎ V-pipe 子单元 ▎ 多 stage 流水 ▎ 软件仿真序列 从 2.22 cyc 到 419 cyc，190× 跨度
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 新增）

§ 07
Memory Hierarchy

存储层级: 4 级延迟阶梯

7.1 延迟阶梯：L1 → L2 → DRAM

Fig. 6 SM_120 存储层级 4 级延迟阶梯（实测 cycles @ 2.58 GHz）
400 cyc 300 cyc 200 cyc 100 cyc 0 33 L1 / SMEM 128 KB SRAM (shared bank) ~13 ns 79 L2 Near 本地 GPC slice ~31 ns 180 L2 Far 跨 NoC 远端 slice (96 MB total) ~70 ns 372 DRAM (GDDR7) 72 GB · 1176 GB/s 实测 ~144 ns +46 +101 +192 每跨一级延迟近乎翻倍 L2 = 96 MB · +92% vs Hopper
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 种精度共享同一物理通路
12 PRECISION INPUTS FP16 / BF16 TF32 FP8 (e4m3/e5m2) FP6 dense FP4 dense MX-FP4 ★ NV-FP4 ★ MX-FP6/8 INT8 dense INT8 sparse 2:4 FP16 sparse 2:4 MX-FP4 sparse UNIFIED TC PIPELINE 29 cycle latency 23 cycle throughput systolic array · same physical path FP64 EXCEPTION (DIFFERENT PATH) DMMA · 175 cyc latency · 449 cyc throughput · 0.99 TFLOPS EFFECTIVE TFLOPS FP16 / BF16 293 / 290 TF TF32 145 TF FP8 588 TF FP6 / FP4 dense 588 TF MX-FP4 (k=64) 1175 TF ★ NV-FP4 1175 TF ★ Sparse INT8 (k=64) 1176 TF ★ Dense INT8 588 TF FP64 (DMMA) 0.99 TF MICROARCHITECTURAL MODEL · 物理机制 同一 systolic array → systolic 深度固定 → 延迟固定 Multiplier 配置随精度切换：FP16 → 16-bit · FP8 → 8-bit (2× 密度) · FP4 → 4-bit (4× 密度) 同一物理通路在 23 cycle 内处理更多 OPs（FP4 = 4× FP16） ★ 开发者结论：精度选择由 存储带宽 决定，而非 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

8.4 TC Co-issue 与功率

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× 差距）
CYCLES TMA Load 2D (1024 B) DRAM → L2 → SMEM 全路径 620 cyc DRAM access address translation tiling + write TMA Load 1D (1024 B) 无 swizzle 开销 488 cyc TMA Store 2D (1024 B) SMEM → L2 buffered (writeback async) 33 cyc ★ 19× faster than load 设计启示 epilogue 优先使用 TMA store · prologue 应隐藏 load 延迟于 prefetch pipeline · pipeline depth ≥ 32 可饱和 1176 GB/s
TMA Load 需完整 DRAM → L2 → SMEM 路径，而 Store 仅需 SMEM → L2 缓冲写（writeback 异步完成）。这种不对称应指导 GEMM epilogue 设计。

9.2 mbarrier 时序

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 新增

§ 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
bit 0 bit 127 bit 64 bit 41 (ctrl start) bit 57 (ctrl end) Lower 64-bit · Opcode + Operands opcode · src/dst regs · immediate · predicate extra operand bits [40:0] of upper 17-bit CONTROL CODE bits [57:41] of upper reserved 17-BIT CONTROL CODE BREAKDOWN Stall Counter 4 bits · [44:41] · 0-15 stall < latency → data corruption Yield 1 bit · [45] 反逻辑: 0=yield 阻断 latency hide Write SB 3 bits · [48:46] SB# to allocate 实测 ≥ 12 个 SB Read SB 3 bits · [51:49] SB# to release Wait Barrier Mask 6 bits · [57:52] SB bitmap to wait on 最多 6 个并行依赖 Reuse no-op Blackwell 不再使用
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。

10.3 CGGTY 5-Warp 阈值

Fig. 10 CGGTY 调度阈值：4 → 5 warp 出现 6× 加速跳跃
0 100 200 300 400 cyc/iter 1 2 3 4 5 6 7 8 10 12 14 16 20 24 28 32 Active Warps per Sub-core × 4 68 136 204 268 45 46 47 48 ★ CGGTY 阈值 6× speedup at 5 warps ≤ 4 warps: 无 latency hiding（每 sub-core 1 warp） ≥ 5 warps: 调度器交织（每 sub-core ≥ 2 warps）
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

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）
PRECISION MEASURED PEAK · TFLOPS EFFICIENCY 0 300 600 900 1200 TF CUTLASS 4.5 NVFP4 block-scale 975.1 ★ 80.7% CUTLASS 4.5 MXFP8 sparse 2:4 664.9 — cuBLAS INT8 607.0 — cuBLAS FP8 E4M3 574.9 95.0% cuBLAS FP16 283.6 97.8% cuBLAS BF16 280.4 96.7% cuBLAS TF32 130.3 89.9% cuBLAS FP32 52.5 ~100% 1 PFLOPS
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 稳态）
63.1% Softmax ALU 96.8 cyc/tile 36.9% 2× MMA 56.7 cyc SOFTMAX SUB-COMPONENT BREAKDOWN 5× expf (MUFU.EX2) ~100 cyc · SFU bottleneck ★ rowmax reduction ~28 cyc rowsum reduction ~28 cyc rescale (FFMA) ~12 cyc ★ OPTIMIZATION DIRECTIONS 1. fast math / polynomial approximation 替代 expf 2. multi-warp overlap (MMA / softmax 交错) 3. SM_100 TM 路径在 FA 场景中优势显著（释放 regs 给 SFU） 4. Tensor Memory 结构性消除 SFU + RF 竞争 → TC 有 63% 时间在等待 ALU 完成 softmax
在 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 数与总功耗）
0 W 50 100 150 200 222 W (peak) 80 W 0 SM (idle) L2/NoC/MC 102 W 22 SM (1 GPC) +22 W 135 W 55 SM (50%) 190 W 110 SM all-on +110 W compute 222 W ★ TC saturated 63% of TDP 350 W TDP nameplate includes PCIe/HW 基础设施 80W (clk/L2/NoC/MC) per-SM 增量 ~1.0 W/SM TC 满载额外 ~30W TDP 上限
基础设施功耗（时钟域 / L2 / NoC / MC）固定 80W，占满载的 36-45%。增加每个 SM 仅 +1W 计算功耗——优化策略应优先减少活跃 GPC 数而非降低 per-SM 利用率。
Fig. 14 每精度能效（TFLOPS/W）：MX-FP4 比 FP32 高 28 倍
8 6 4 2 0 TF/W 6.90 ★ Sparse INT8 28× FP32 6.80 ★ MX-FP4 28× FP32 3.45 Dense INT8 14× FP32 3.32 FP8 14× FP32 1.73 FP16 7× FP32 0.24 FP32 FMA baseline 0.005 FP64 ratio 1:39 PEAK MAP 1175 TOPS Sparse INT8 1175 TF MX-FP4 588 TF FP8 / Dense 290 TF FP16
MX-FP4 在相同 23 cyc 内处理 4× FP16 OPs，但 TC 功耗（开关功耗由物理结构决定）几乎不变。吞吐 4× 而功耗不变 → 能效 4×。

11.4 DVFS 与最优能效频率

Fig. 15 DVFS：频率-吞吐线性 vs 频率-功耗超线性
TFLOPS POWER (W) 600 1200 1500 1800 2200 2580 MHz SM Clock Frequency ★ Best TFLOPS/W 1500-1800 MHz TFLOPS · 线性 (R² > 0.99) Power · P ∝ f^1.5-2.0 (V²f) boost: +73% TFLOPS but +127% power
频率-吞吐完美线性，但功率超线性（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 Cache96 MB (统一)100 MB ×2 (双 die)单 die 简化制程
DRAMGDDR7 72 GB / 384-bitHBM3e 192 GB工作站领域需中容量低成本
NVLink无 (PCIe Gen5)NVLink 5 (1.8 TB/s)单卡定位，无多卡训练
SM 数量110 SM (smid 0–109)132+132 (双 die)单 die 高密度布局
FP4 Peak975 TFLOPS~5000 TFLOPS5× 差距，主要来自 SM 数
TDP350 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 ∝ f1.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 cyc28–33 cyc持平
DRAM 延迟372 cyc (GDDR7)~430 cyc (HBM3)改善 14%
L2 容量96 MB50 MB+92%
SFU (rcp)44 cyc~28 cyc退化 57%
mma.sync FP1629/23 cyc~16/8 cyc退化 ~3×
Max warps/SM4864−25%
FP64 : FP321: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 方法学局限性

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

§ 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. [1]Jarmusch, J., et al. Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks. arXiv:2507.10789, 2025.
  2. [2]Jarmusch, J., et al. Microbenchmarking NVIDIA's Blackwell Architecture. arXiv:2512.02189v3, 2026.
  3. [3]Huerta, S., et al. Analyzing Modern NVIDIA GPU cores. arXiv:2503.20481, 2025.
  4. [4]Wong, H., et al. Demystifying GPU Microarchitecture through Microbenchmarking. IEEE ISPASS, 2010.
  5. [5]Mei, X. and Chu, X. Dissecting GPU Memory Hierarchy through Microbenchmarking. IEEE TPDS, 2017.
  6. [6]Volkov, V. Understanding Latency Hiding on GPUs. PhD Thesis, UC Berkeley, 2016.
  7. [7]Jia, Z., et al. Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking. arXiv:1804.06826, 2018.
  8. [8]Jia, Z., et al. Dissecting the Ampere GPU Architecture through Microbenchmarking. GPU Technology Conference, 2022.
  9. [9]Raihan, M. A., et al. Modeling Deep Learning Accelerator Enabled GPUs. IEEE ISPASS, 2019.
  10. [10]Yan, D., et al. Demystifying Tensor Cores to Optimize Half-Precision Matrix Multiply. IEEE IPDPS, 2020.
  11. [11]Dao, T., et al. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. NeurIPS, 2022.
  12. [12]Dao, T. FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning. ICLR, 2024.
  13. [13]Thakkar, V., et al. CUTLASS: CUDA Templates for Linear Algebra Subroutines and Solvers. NVIDIA, 2023.
  14. [14]Leng, J., et al. GPUWattch: Enabling Energy Optimizations in GPGPUs. ISCA, 2013.
  15. [15]Kandiah, V., et al. AccelWattch: A Power Modeling Framework for Modern GPUs. MICRO, 2021.
  16. [16]NVIDIA Corporation. NVIDIA Blackwell Architecture Technical Brief. 2024.
  17. [17]NVIDIA Corporation. CUDA C++ Programming Guide. v13.0, 2026.
  18. [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/decSec 12 / 06inc ~61 cyc, dec ~50 cyccompute_120a + CUDA 13.3
SIMD 8×4Sec 4 / 154.0 cyc, 1.0 inst/cyc, V-pipecompute_120a + CUDA 13.3
cvt.pack (I2IP)Sec 6 / 084 cyc, 2 ops/cyc @ ILP=8compute_120a + CUDA 13.3
FP8 cvt (e4m3)Sec 6 / 09pack 8 cyc, unpack 4 cyccompute_120a + CUDA 13.3
MX narrow cvtSec 6 / 108.22 cyc round-tripcompute_120a + CUDA 13.3
tensormap.replaceSec 1026-45 cyc/opcompute_120a + CUDA 13.3
L2::cache_hintSec 10UTMACCTL + cp.async.bulkcompute_120a + CUDA 13.3
CUTLASS 4.5 GeForceSec 18 / 05NVFP4 975 T, CLC + warp specCUDA 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 的非张量峰值。