Back to blog

Blackwell GPU 架构综述

NVIDIA Blackwell GPU 架构的综合技术综述——涵盖第五代张量核心(tcgen05、FP4/FP6)、张量内存 TMEM、硬件解压引擎 DE、重构内存层次(L1/L2/HBM3e/GDDR7)、双芯 B200 互连,以及与 Hopper/H200 的微基准对比与调优策略。综合自 Jarmusch et al. 两篇微基准论文。

Blackwell GPU 架构综述

本文是对 NVIDIA Blackwell GPU 架构的综合技术综述,综合自同一研究组(Jarmusch et al., University of Delaware)的两篇微基准论文。它把「消费级 vs 数据中心」「Blackwell vs Hopper/H200」的实测发现整合为统一的架构画像。

来源论文(均已单独分析):

基础背景:若不熟悉 GPU 内存层次、CUDA/SM/warp、张量核心与 Roofline 模型,建议先读 GPU 架构基础:内存层次、CUDA 与张量核心。


一、核心定位

NVIDIA Blackwell 面向 exascale 仿真、AI 训练/推理、高吞吐科学计算。相较 Hopper/H200,四大架构革新:

  1. 第五代张量核心(含 FP4/FP6 与 tcgen05 warp 级 MMA)
  2. 张量内存 TMEM(每 SM 专属 256KB)
  3. 硬件解压引擎 DE(Decompression Engine)
  4. 统一大 L2 缓存 + warp 级调度创新

二、微架构核心模块

2.1 第五代张量核心流水线

  • 新 PTX 指令 tcgen05.mma 取代 Hopper 的 wgmma。
  • warp 级、单线程 MMA:每线程独立发射 MMA,消除 warp 宽指令屏障。
  • 指令延迟近乎恒定(~11 周期),跨 tile 尺寸与数值精度基本一致。
  • 支持格式:FP64、TF32、BF16、FP16、FP8、FP6、FP4、INT8 → 映射到 SASS 指令 DMMA / HMMA / QMMA / OMMA / IMMA。
  • warp 级派发使指令延迟较 Hopper warp-group 派发降低 2.9–11.6×。

2.2 张量内存(TMEM)

  • 每 SM 专属 256KB(512 列 × 128 lane × 32bit),经 tcgen05.cp / .ld / .st 访问。
  • 读带宽 16 TB/s per SM,写带宽 8 TB/s,与传统 shared memory 叠加。
  • 显著降低缓存缺失延迟:降 58%(≈420 周期 vs H200 ≈1000 周期)。

2.3 硬件解压引擎(DE)

  • 原生加速 LZ4、Snappy、Zstandard、GZIP、Cascaded、Bitcomp、ANS。
  • 100MB 块亚毫秒延迟,吞吐达 539 GB/s。
  • 相比软件解压 50–200× 加速;吞吐随压缩比变化,饱和 batch 下并行良好。

2.4 双芯设计与互连(B200)

  • B200 = 两颗相同 die(各 74 SM,共 148 SM),8 GPC,每 die 4 个 L2 分区、8 个 HBM3e 栈。
  • 统一 192GB 地址空间;die 间 NV-HBI ≈1.5 TB/s 相干互连。

三、内存层次(vs Hopper/H200)

层级BlackwellHopper/H200
L1 / Shared(可配)128KB(较 Hopper 减半)256KB
L264–65MB 单块/分区(B200: 8×8MB 私有分区)50MB(2×25MB)
GlobalRTX 5080: 16GB GDDR7 / B200: 192GB HBM3e—
Cache line128 字节—
替换策略(实测推断)L1 近似 LRU、L2 伪 LRU(未公开)—

延迟(实测周期)

子系统BlackwellHopper/H200
L1 命中≈22–40≈30–40
L2 命中≈128–358≈273(Hopper)/ ≈128(H200)
TMEM 缺失≈420—
DRAM 访问≈877(5080)/ ≈4200(B200)≈659(Hopper)/ ≈1000(H200)

持续 DRAM 带宽(微基准)

方向BlackwellHopper/H200
读8.2(5080)、8.4(B200)15.8(Hopper)、4.38(H200)
写1.6(5080)2.2(Hopper)

有效带宽由缓存缺失率 MRMR 决定:

Beff=Bpeak×(1−MR),MR=NmissNhit+NmissB_{\mathrm{eff}} = B_{\mathrm{peak}} \times (1 - MR), \quad MR = \frac{N_{\text{miss}}}{N_{\text{hit}} + N_{\text{miss}}}

⚠️ 带宽回退:消费级 RTX 5080 读带宽 8.2 TB/s 显著低于 Hopper 15.8 TB/s;但 B200 双芯设计达 ~8.4 TB/s 片外带宽,STREAM 微基准实测 7.48 TB/s(1.71× H200)。


四、SM 执行与调度

  • Warp 调度器:每 SM 4 个,每周期发 2 条独立指令;峰值 8 标量 op(INT32/FP32)或 4 张量 op。
  • ALU:统一 INT32/FP32 簇,每 lane 每周期 4-way fused-MAC;真实依赖延迟 INT32/FP32 = 4 周期,FP64 ~32–64 周期(模拟),MMA(FP4) ~1.21 周期/tile。
  • 张量核心吞吐:FP4 在 ILP=6、25 warps 时达 ~11 TFLOP/s per SM。
  • 性能模型:

ThroughputSM=fcore×IPC×Wactive\mathrm{Throughput}_{\mathrm{SM}} = f_{\mathrm{core}} \times \mathrm{IPC} \times W_{\mathrm{active}}

其中 fcoref_{core} 为 SM 时钟(2.4–2.65 GHz),WactiveW_{active} 为活跃 warp(≤64);全局 T=#SM×ThroughputSMT = \#SM \times \mathrm{Throughput}_{SM}。

  • 资源争用:warp 增多时因寄存器文件与 shared-L1 争用,指令发射率下降。
  • 最佳配置:延迟隐藏与 ILP 在 32–48 warps/SM 且 ILP≥4 时最优。

代表性延迟/吞吐表:

操作真实延迟(周期)完成延迟(周期)吞吐(ops/cycle/SM)
INT32 (mad)16.97—0.25
FP32 (fma)7.97—0.50
FP64~37.537–64(流水)0.054
MMA (FP4, tile)~1.21见张量核心表—

五、张量核心与混合精度

FP4(E2M1)/FP6(E3M2)为 Blackwell 新增硬件数据通路,Hopper/H200 缺失。

精度Blackwell TFLOPSHopper/H200 TFLOPS加速
FP6444.834.01.32×
FP32481.2378.41.27×
FP83851.43026.91.27×
FP65134.8N/A新增
FP47702.5N/A新增
INT83927.13088.41.27×

(@2.4 GHz;峰值 TFLOPS 随子字并行度线性增长)

端到端 Transformer 训练:Blackwell 混合精度吞吐 1.56×、能效 +42% vs H200。


六、与 Hopper/H200 对比要点

维度发现
内存层次L2 较 Hopper +30%,但 L1 减半;B200 的 TMEM + 8 路 L2 分区降低跨流量与缺失延迟
带宽RTX 5080 带宽回退(8.2 vs 15.8 TB/s);B200 双芯 ~8.4 TB/s,STREAM 1.71× H200
GEMM 回退8192³ 大 GEMM Blackwell 反而慢(0.233 vs Hopper 0.887 TFLOP/s)、功耗更高(114W vs 68W)——归因 cuBLASLt 驱动/编译器不成熟
推理能效Transformer 推理各精度功耗均低于 Hopper(FP8: 45.1W vs 57–60W)
解压引擎真实负载 50–200× 超软件解压

七、工作负载优化与调优策略

  • 稠密 GEMM:operand tile 取 64×64 最大化 TMEM 带宽;用 tcgen05.cp 在 MMA 计算中预取 tile;中间累加驻留 TMEM 以减少 L2/global 流量。
  • 稀疏 GEMM / SpMV:块以压缩形式(RLE/Bitcomp)存 HBM,DE 即时解压后流入 TMEM 归约。
  • Transformer 推理:权重量化至 FP8/FP4(E4M3/E2M1),激活 FP16;用 DE 加载量化权重入 TMEM,注意力与中间结果驻留 TMEM。
  • 混合精度训练:BF16/FP16 累加更稳(纯 FP32 有 ~50% 性能损失);TPC 配对协同调度 CTA 共享 TMEM 数据;小层偏 L1、大层偏 TMEM。
  • 线程块调度:round-robin 映射直至 SM 资源耗尽;最佳延迟隐藏与 ILP 在 32–48 warps/SM、ILP≥4。

八、总结、局限与展望

核心影响:第五代张量核心、TMEM、DE、统一 L2 与双芯设计,改变了内存受限内核设计以及混合/超低精度推理格局。Blackwell 在混合精度 AI、能效、解压加速上均有可量化提升,兼顾轻 warp 场景与大规模 exascale 部署。

局限:

  • RTX 5080 的 DRAM 带宽回退与 FP8 GEMM 性能回退,凸显 cuBLASLt 编译器/运行时成熟度的影响。
  • L1 减半(vs Hopper)须靠 TMEM 感知算法弥补。

展望:未来工作负载优化将越来越依赖最大化 TMEM 利用与硬件 DE 集成,以抵消 L1/L2 约束、在低精度下最大化吞吐。


九、参考资源