NVIDIA Blackwell架构微基准深度解析：FP4/FP6赋能LLM推理2.5倍加速，36.3TFLOPS FP64重塑科学计算

关键词：Blackwell、GPU、 Microbenchmark 、5th-generation Tensor Core 、 TMEM

本文工作量化了张量内存（TMEM）对矩阵密集型负载的影响，评估了硬件解压缩引擎（DE）的吞吐量及最优使用方式，通过新的tcgen05 PTX 指令分析了第五代张量核心的执行特性。

此外，还评估了 FP4 与 FP6 精度的权衡关系，在大型语言模型（LLM）推理、科学计算核函数及混合精度训练等多样化负载中对 Blackwell 进行了基准测试，并为面向该下一代架构的开发者提炼了可落地的性能优化指南。

• Microbenchmarking NVIDIA’s Blackwell Architecture: An in-depth Architectural Analysis
• https://arxiv.org/pdf/2512.02189
• 本文提出了一套新开发的开源微基准测试套件（ 因双盲评审要求，目前暂无法公开代码 ）
• 1.7 万字，阅读 70 分钟，播客 23 分钟

本文针对 NVIDIA Blackwell（B200）GPU 架构开展了系统的微基准测试与深度分析 ，旨在量化其核心创新的性能提升，为开发者优化工作负载提供实践指导。随着 AI 与 HPC 对硬件效率、扩展性的需求激增，B200 引入第 5 代张量核心、张量内存（TMEM）、硬件解压引擎（DE）及双芯片设计等关键创新， 但相关性能量化方法滞后于硬件发展。

研究团队开发了开源微基准套件 ，从内存子系统、张量核心流水线、浮点精度（FP32 至 FP4）等维度，对比 B200 与 H200 的性能表现 。测试结果显示：

• B200 的张量核心增强使混合精度吞吐量较 H200 提升 1.56 倍，能效优化 42%；
• 缓存未命中场景下内存访问延迟降低 58%，重塑了算法设计逻辑。
• TMEM 作为张量运算专用内存，实现 16TB/s 读取带宽，有效减少内存瓶颈；
• DE 支持 7 种压缩格式，Bitcomp 格式输出吞吐量达 462.4GB/s，显著提升数据处理效率。

在实际负载测试中，B200 在 LLM 推理（FP4 精度下 Mistral-7B 吞吐量提升 2.5 倍 ）、科学计算（FP64 稠密矩阵乘法性能达 36.3TFLOPS ）、神经网络训练（ResNet-50 训练速度提升 1.54 倍 ）等场景均展现显著优势。

本文为 B200 提供了首个详细的微基准表征，明确了各创新组件的最优使用场景，为 AI 与 HPC 应用的性能优化及未来 GPU 设计提供了关键参考。

本文目录

• 关键问题
  • 问题一：TMEM 的引入是否真正解决了内存墙问题，还是仅仅将瓶颈转移到了其他子系统？
  • 问题二：FP4/FP6 等超低精度格式在实际 AI 推理中的可用性究竟如何？精度损失是否被低估？
  • 问题三：Blackwell 的“专用化”架构演进，是否在制造一种新型的“碎片化”与“锁定”风险？
• 一、引言
• 二、相关工作
• 三、BLACKWELL 架构
  • A. Blackwell 架构概述
• 四、PTX 微基准测试方法
  • A. 针对 Blackwell 特有特性的新型基准测试设计
• 五、内存子系统
  • 5.1 张量内存（TMEM）
  • 5.2 解压缩引擎（DE）
• 六、GPU 核心微架构
  • 6.1 第五代张量核心
  • 6.2 扩展精度支持：FP4 与 FP6
• 七、性能分析与案例研究
  • 7.1 实验方法
  • 7.2 大型语言模型（LLM）推理
  • 7.3 科学计算负载
  • 7.4 混合精度训练：端到端训练性能
• 八、讨论
  • 8.1 架构权衡
  • 8.2 软件生态
  • 8.3 部署建议
• 九、结论
• 参考文献

关键问题

问题一：TMEM 的引入是否真正解决了内存墙问题，还是仅仅将瓶颈转移到了其他子系统？

TMEM（Tensor Memory）作为 Blackwell 架构的专用内存子系统，旨在减少张量计算中的内存瓶颈。但这是否真正解决了“内存墙”问题，还是将瓶颈转移到了其他部分（如编程复杂性、资源竞争或非张量工作负载）？

TMEM 在特定场景下显著改善了内存性能，但并未完全解决内存墙问题，且可能引入了新的瓶颈和复杂性：

1. 延迟与带宽改进 ：TMEM 在缓存未命中情况下将内存访问延迟降低了58%（从 Hopper 的1000周期降至420周期），并提供每SM 16 TB/s的读取带宽和8 TB/s的写入带宽。在FP8数据上，TMEM 实现了8 TB/s的带宽，是传统 ld.global 路径（3.8 TB/s）的2.1倍提升。这表明 TMEM 确实缓解了张量工作负载中的内存瓶颈。
2. 软件管理与复杂性 ：作者指出，TMEM 需要全新的指令序列（如 tcgen05.cp、tcgen05.ld 和 tcgen05.sr），这增加了编程复杂性和编译器优化负担。传统的数据移动指令无法与 TMEM 交互，开发者必须显式管理数据放置和移动，这可能将瓶颈从硬件转移到软件层面。
3. 容量与访问模式限制 ：TMEM 容量为每SM 256KB，且针对64×64元素块（如FP8精度下4KB）优化。作者提到，小于32×32的块会因内存接口未充分利用而带宽降至峰值的45%，大于128×128的块则会引入多阶段传输和流水线停顿。这意味着不规则或小规模工作负载可能无法充分受益于 TMEM，瓶颈可能转移到计算效率或调度上。
4. 资源竞争与通用性 ：作者【未】直接讨论 TMEM 与其他子系统（如SMEM、L1/L2缓存）的资源竞争，但指出 TMEM 带宽与L1/SMEM带宽是叠加而非竞争关系 。然而，TMEM 专用于张量核心操作，对于非张量工作负载（如标量运算或图形渲染），其优势可能不明显，甚至可能因芯片面积分配而间接影响通用计算资源。

结论：TMEM 显著提升了张量密集型任务的内存性能，但并未完全解决内存墙问题，而是将部分瓶颈转移到了软件复杂性、访问模式优化和专用硬件适应性上。 对于非张量或混合工作负载，TMEM 的收益可能有限，且需要额外的开发成本。

问题二：FP4/FP6 等超低精度格式在实际 AI 推理中的可用性究竟如何？精度损失是否被低估？

论文展示了 FP4 和 FP6 格式在理论上能带来显著的吞吐量提升和内存节省，但这并不意味着它们已准备好广泛应用于实际生产环境。文中报告的“可接受”精度损失在更广泛、更复杂的任务中可能面临挑战：

1. 显著的性能提升与明确的精度代价：对于 Mistral-7B 模型，FP8 和 FP4 分别带来了 1.73 倍和 2.5 倍的推理吞吐量提升。然而，这种性能增益伴随着明确的模型质量下降：FP8 导致困惑度增加 1.9%-2.4%，而 FP4 则导致更大幅度的 7.7%-9.1% 的困惑度增加。对于需要复杂逻辑或事实一致性的精度敏感型应用，这种损失可能是不可接受的。
2. 应用场景的特定性与“可接受”的定义模糊：在标准语言建模困惑度指标上，FP4 在稀疏混合专家模型（Mixtral-8x7B）上显示出更大的收益（2.69倍提升），但这并未证明其在所有任务上的鲁棒性。“可接受精度损失”是一个高度依赖于具体应用的定义。论文未探讨 FP4/FP6 在需要数学推理、代码生成或多轮对话一致性等更严苛任务上的表现，而这些正是实际部署中可能暴露低精度弱点的场景。
3. 硬件支持超前于软件生态与算法成熟度：虽然 FP6 硬件已存在，但缺乏成熟的软件工具链支持。此外，FP4/FP6 需要“逐层精度选择”，这引入了巨大的超参数调优和算法工程复杂性。作者也指出，FP32 累加器会将 FP16 计算的吞吐量减半，这迫使开发者在数值稳定性和性能之间做出艰难取舍。这表明，可用性不仅受限于硬件，更受限于整个软件栈和算法优化技术的成熟度。
4. FP6 的定位：有希望的折中，但非终极解决方案：文中介绍 FP6（e3m2 格式）在动态范围和内存节省（相比 FP8）之间提供了折中。然而，其性能（TFLOPS）介于 FP8 和 FP4 之间。它更像是一个针对特定精度需求的专用工具，而非一个通用解决方案。其价值取决于能否在特定模型上找到优于 FP8（精度更高）和 FP4（性能更好）的“甜点”，而这需要大量的每模型、每任务调优。

结论：FP4/FP6 格式在实际 AI 推理中的可用性目前是有限且有条件的。 它们为追求极致吞吐量和能效、且对一定精度损失有容忍度的特定推理场景（如某些检索增强生成或已知对量化不敏感的模型）提供了强大工具。然而，论文中的数据并未低估精度损失，反而清晰地揭示了其代价。真正的挑战在于：
* 精度损失的普适性：报告的平均困惑度增长可能掩盖了某些关键能力（如推理、规划）的更严重退化。
* 系统复杂性：为充分利用它们所需的硬件感知算法设计和细致的逐层调优，极大地提高了部署门槛。
因此，这些超低精度格式目前更适用于由专家主导、针对已知工作负载的优化，而非面向所有 AI 应用的、“开箱即用”的通用解决方案。它们的广泛采用有待于更自动化的量化工具链和更坚实的算法研究来减轻精度损失。

问题三：Blackwell 的“专用化”架构演进，是否在制造一种新型的“碎片化”与“锁定”风险？

Blackwell 通过引入 TMEM、解压引擎（DE）、FP4/FP6 支持等专用组件，针对 AI 和 HPC 工作负载进行优化。这确实在提升性能的同时，潜在地引入了碎片化和锁定风险：

1. 专用硬件组件：文中详细描述了 TMEM、DE 和第五代张量核心（支持 FP4/FP6）。这些组件针对特定任务（如张量乘法、数据解压）优化，可能挤占了通用计算资源。例如，晶体管数量增至 2080 亿，但未讨论这是否减少了通用核心的面积或影响了传统工作负载的性能。
2. 软件生态依赖：CUDA 13.0 提供了对 TMEM 和 CTA 配对的初步支持，但框架集成仍在进行中。FP6 硬件已支持，但缺乏软件工具链。此外，FP4/FP6 需要逐层精度选择，这增加了开发复杂性。这种对 NVIDIA 专有工具链（如 PTX 指令 tcgen05）的依赖，可能抑制跨平台编译器（如 SYCL、OpenCL）的发展，加剧生态锁定。
3. 通用性牺牲：论文比较了不同工作负载的性能，显示 Blackwell 在 LLM 推理和训练中表现优异，但对于科学计算（如 FP64 DGEMM），其性能提升主要来自 TMEM 和内存带宽优化，而非通用核心改进。这暗示专用化可能以通用计算性能的平衡为代价。例如，在较小工作集下，B200 的内存带宽效率低于 H200，表明其优化更偏向大规模张量任务。
4. 能效与可持续性：Blackwell 在训练中实现了 42% 的能效提升，但这是基于特定工作负载的测量。未考虑制造更复杂芯片的全生命周期能耗（如制造、散热）。专用化路径可能导致硬件迭代更频繁，增加电子废弃物和环境成本。

结论：Blackwell 的专用化架构在提升目标工作负载性能的同时，确实带来了碎片化和锁定风险。 它加深了对 NVIDIA 专有硬件和软件栈的依赖，可能削弱通用计算能力，并挑战生态多样性和长期可持续性。开发者需权衡性能增益与移植成本，而行业需关注避免过度依赖单一厂商的技术路线。

一、引言

人工智能（AI）与高性能计算（HPC）已发展成为数据密集型学科，持续对硬件的效率、可扩展性及精度提出挑战。
* 如今，大型语言模型（LLMs）的参数规模已突破千亿，处理的上下文窗口可涵盖数百万个 token；
* 与此同时，多物理场模拟与气候模拟需实现万亿次浮点运算每秒（TFLOPS）的持续性能——这些需求推动着 GPU 设计向“兼顾大规模并行与架构适应性”的方向演进。

在该规模下，现代加速器需平衡多方面需求：既要维持密集张量负载的算术吞吐量，又要最小化片上与片外内存延迟，同时还需提供能有效支持混合精度计算的硬件原语。

日益增长的需求暴露了当前 GPU 架构的若干局限性，尤其体现在内存层级结构、精度灵活性及延迟敏感型任务调度方面。因此，加速器的持续架构创新对于推动“吞吐量优化的训练任务”与“时间关键型推理任务”的发展至关重要。NVIDIA 的 Blackwell 架构便是为应对部分此类挑战而设计的架构，标志着 GPU 代际演进的重要里程碑。

作为 Hopper 架构的直接继任者，Blackwell 架构在计算流水线、内存层级结构及张量处理子系统等多个方面对 NVIDIA GPU 设计进行了改进。
* Blackwell 引入了支持 FP4 与 FP6 精度执行的第五代张量核心，为大规模训练任务提供了精度与性能的权衡空间。
* 此外，新增的张量内存（TMEM）子系统作为张量数据移动的专用片上内存，可在矩阵密集型操作中减少对共享内存（SMEM）和每个流式多处理器（SM）寄存器文件（RF）的依赖。
* 其次，NVIDIA 还集成了硬件解压缩引擎（DE），并重新设计了指令流水线以支持对压缩模型权重的访问。
* 除了原始计算能力的增强，Blackwell 还优化了线程与线程块（CTA）调度模型，以充分利用 SM 间通信与内存并发能力。

鉴于 Blackwell 架构引入了诸多变革，对其微架构与新指令进行分析至关重要，这将帮助应用开发者与科研人员充分挖掘现代及未来 GPU 的性能潜力。

本文提出了一套新开发的开源微基准测试套件（因双盲评审要求，目前暂无法公开代码）。该套件基于并行线程执行（PTX，NVIDIA GPU 的中间汇编语言）与 CUDA 实现，旨在对 NVIDIA Blackwell GPU 进行全面架构分析。

该套件重点关注 Blackwell 与 Hopper 架构的差异化创新，通过在计算密集型与内存密集型负载下进行压力测试，系统评估其性能表现，进而揭示对并行计算应用的影响。

本文的核心贡献如下：

• 构建针对性微基准测试，对 NVIDIA Blackwell B200 的关键组件进行表征分析；据我们所知，本文是首个对该下一代 GPU 进行详细微基准测试表征的研究。
• 量化张量内存（TMEM）对矩阵密集型负载的影响，及其在缓解张量计算中内存瓶颈的作用。
• 评估解压缩引擎（DE）在不同格式下的吞吐量，并确定其最优使用方式。
• 通过新的 tcgen05 PTX 指令分析第五代张量核心的执行过程，研究其性能影响。
• 评估混合精度张量操作中 FP4/FP6 精度的性能与精度权衡，量化“精度-性能”折衷关系。
• 在大型语言模型（LLM）推理/训练、科学计算核函数及混合精度负载场景下对 Blackwell 进行基准测试，以展示其实际应用价值与性能提升。
• 为开发者提供基于 Blackwell 架构进行优化的可落地性能指导方针。

本文其余部分结构如下：

• 第二节详细阐述我们在现有 GPU 微基准测试领域的贡献；
• 第三节概述 Blackwell B200 架构；
• 第四节详细介绍用于系统表征 Blackwell 微架构的微基准测试方法；
• 第五节分析内存子系统；
• 第六节阐述张量核心流水线；
• 第七节对关键负载进行性能分析；
• 最后，第八节讨论研究意义与权衡关系。

二、相关工作

在高性能计算（HPC）研究领域，理解 GPU 性能长期以来都是核心关注点。多年来，已有多项研究通过微基准测试及其他方法剖析架构层级，对 GPU 微架构进行细粒度分析。

| 架构 | 研究重点 |
| :— | :— |
| Tesla, Fermi | 内存与缓存行为 |
| Kepler, Pascal, Maxwell | 线程束调度与指令延迟 |
| Turing 到 Hopper | 混合精度与张量核心性能，针对矩阵乘加（MMA）指令、分块大小（tile size）及数据布局开发了相应基准测试 |
| 近期研究 | 指令级并行（ILP）及高寄存器压力下的流水线动态特性 |

除微基准测试外，研究人员还构建了用于表征 GPU 性能的框架。

• 应用程序剖析：可收集运行时指标，但存在开销较大且架构可见性有限的问题。
• 屋顶线模型（Roofline models）：可提供“吞吐量-计算强度”关系图，但会过度简化瓶颈问题且无法捕捉动态内存行为。
• 缓存停顿预测：可通过访问模式估算流水线延迟，但无法涵盖现代 GPU 的复杂特性（如缓存旁路、线程束调度及内存合并）。
• 分析模型：基于 Hong 与 Kim 研究成果开发的 Accel-Sim 与 GCoM 等模型，虽能为 GPU 性能提供有价值的洞察，但均未对 Blackwell 架构特有的张量内存（TMEM）或解压缩引擎（DE）进行建模——这是因为准确模拟这些组件所需的详细架构信息尚未公开。

因此，若缺乏对这些组件的系统性理解，研究界将难以获取性能建模、负载优化及数据中心部署所需的 AI 推理负载准确模拟所需的关键数据。

三、BLACKWELL 架构

本节将介绍基于 NVIDIA Blackwell 架构的数据中心级 GPU——B200 的架构，并详细阐述其与前代设计的差异。

A. Blackwell 架构概述

B200 GPU 标志着 GPU 架构理念的重大进步。在此之前，从 Tesla 到 Hopper 的各代 GPU 均以最大化大规模模型训练的浮点运算每秒（FLOPS）为核心目标；与之不同，Blackwell 架构更注重训练后任务与推理任务的效率，在内存与计算组织方面均实现了变革性改进。

图 1：通过 NV-HBI 互联的 NVIDIA Blackwell GPU 双芯片设计

单颗 B200 GPU 采用双芯片配置：两颗 GPU 芯片共包含 2080 亿个晶体管，148 个流式多处理器分布在 8 个图形处理集群（GPCs，包含多个 SM 的硬件单元）中，配备 4 个 L2 缓存分区（数量为 Hopper 架构的 2 倍），并集成 8 个 HBM3e 内存堆叠。

尽管两颗芯片在物理上分离，但通过 NVIDIA 高带宽接口（NV-HBI，用于连接双芯片的高速互联技术）实现了统一，对软件而言呈现为一个连贯的单一设备，拥有 192GB 的统一 HBM3e 内存空间。

在每个 SM 内部，Blackwell 引入的第五代张量核心打破了前代架构（Volta、Ampere、Hopper）的“线程束同步范式”。前代架构要求线程束内的所有 32 个线程必须先同步，再通过 mma.sync 或 wgmma 指令执行矩阵乘加（MMA）操作。这种“锁步”模式降低了调度灵活性，尤其在处理长度可变的依赖链时会产生空闲周期。

Blackwell 用 tcgen05.mma（单线程指令）取代了线程束同步的 MMA 指令。如今，每个线程可独立发起 MMA 操作，消除了线程束级同步，实现了张量操作的真正单线程调度。操作数现在可从共享内存（SMEM）及新的内存通路——张量内存（TMEM）中获取。每个 SM 配备的 TMEM 可为张量核心提供内存访问能力，其分配、数据移动及释放需通过 tcgen PTX 指令集由软件显式管理，这使得编译器工具链能精确控制分块（tile）的局部性与流量模式。

独立 MMA 调度的灵活性减少了空闲周期，并为编译器带来了优化空间，但同时也引发了关于新性能限制的疑问：依赖关系下的指令延迟、张量核心使用的并发性及流水线饱和情况。这些信息在厂商文档中尚未公开，本文将通过系统性表征进行探索。

在数值支持方面，Blackwell 的张量核心引入了原生 FP4 与 FP6，用于量化推理任务，进一步提升了 AI 负载的内存与计算效率。架构创新还延伸至线程块层级：通过“CTA 对执行”（两个相邻序号的协作线程数组共享操作数）减少冗余数据移动。每个 CTA 对映射到一个张量处理集群（TPC），并利用专用的 TPC 内通信网络实现高效的操作数共享。

为进一步拓展功能，Blackwell 的张量核心还原生支持“权重固定数据流”（weight-stationary dataflows）的卷积算子——通过收集器缓冲区缓存并复用矩阵 B（权重张量）操作数，从而优化可从操作数局部性中获益的卷积核函数。此外，针对模型与数据规模不断增长的问题，Blackwell 引入了基于硬件的解压缩引擎（DE），将解压缩任务从通用 SM 中卸载。该子系统支持多种算法，可将模型权重与大型数据库表以压缩形式存储在 HBM3e 中，并在内存访问过程中透明地完成解压缩。

尽管部分架构细节已公开，但关键的微架构信息（如指令延迟、流水线深度、缓存交互及饱和特性）仍不明确。本文通过 PTX 微基准测试实验（第五、六节）对这些信息进行系统性研究，以填补 AI 与 HPC 性能相关的知识空白。

图 2：tcgen05、wgmma 及 Volta/Ampere 架构的张量核心指令流水线

四、PTX 微基准测试方法

本文采用基于 NVIDIA 并行线程执行（PTX，GPU 的中间汇编语言）的微基准测试方法，对 Blackwell 的微架构特性进行表征。

尽管此前的 GPU 表征研究[4], [7], [14]已建立了 PTX 级基准测试的基础原则，但本文通过创新技术拓展了这些方法，专门用于剖析 Blackwell 此前未被表征的组件，包括第五代张量核心的 FP4/FP6 精度模式、解压缩引擎（DE）及改进后的缓存层级结构。

该方法利用 PTX 对特定架构的寄存器与内存操作进行显式控制。PTX 代码会被编译为流式汇编语言（SASS，即 GPU 机器码）并由 GPU 执行。通过记录 PTX 到 SASS 的转换过程，并结合观测性能进行验证，本文确保微基准测试能准确隔离并测量特定的微架构行为。

A. 针对 Blackwell 特有特性的新型基准测试设计

1. 张量内存（TMEM）

与前代架构中“MMA 操作完全依赖共享内存（SMEM）、动态共享内存（DSMEM）及寄存器文件（RF）”不同，Blackwell 引入的 TMEM 是专为张量操作设计的专用片上内存。这一变革催生了新的数据移动范式，也带来了尚未被探索的性能优化机遇。理解 TMEM 的性能特性至关重要，原因如下：

• 首先，传统的数据移动指令（包括 wmma.load、ldmatrix、ld.shared 及 cp.async）无法与 TMEM 交互，开发者必须采用全新的指令序列（tcgen05.ld、tcgen05.st、tcgen05.cp）；
• 其次，这一新内存层级的性能影响尚未被表征，应用开发者缺乏“何时及如何有效利用 TMEM”的指导。

本文通过首个全面的实证分析，填补了 TMEM 性能特性及其对实际计算核函数影响的研究空白。研究方法包含三种评估策略：

1. 通过指针追踪基准测试（pointer-chase benchmarks，通过创建依赖的内存访问来避免流水线重叠，从而揭示各内存层级的基础访问开销）对比传统共享内存与 TMEM 的内存访问延迟，建立性能基准。
2. 系统对比新型 TMEM 数据移动指令（tcgen05.*系列）与前代指令在不同访问模式下的表现。
3. 通过改变操作数大小与访问步长，确定带宽饱和点，并测量不同配置下的单次访问延迟——这一过程可揭示新指令集的能力与局限性。

2. 解压缩引擎表征

为系统表征 B200 的硬件解压缩引擎（DE），本文开发了定制化微基准测试套件，在受控测试条件下针对七种压缩格式（LZ4、Snappy、Zstandard、GZIP、Cascaded、Bitcomp、ANS）进行测试。测试使用 100MB 数据集，测量每种支持格式的端到端解压缩吞吐量：输入吞吐量定义为从 GPU 内存读取压缩数据的速率，输出吞吐量定义为生成解压缩数据的速率，延迟则包含内存传输在内的完整解压缩时间。为确保热稳定性与缓存稳定性，每个测量值均为“100 次预热迭代后 1000 次迭代”的平均值。

本文生成了具有不同熵值的合成数据集：随机数据（不可压缩，压缩比 1.00 倍）、混合字母数字数据（压缩比 1.98 倍）、重复模式数据（压缩比 15.02 倍）及零填充缓冲区（压缩比 245.45 倍）。所有实验均采用 LZ4 格式，以排除“格式特定优化”对“压缩比影响”的干扰——数据先在 CPU 上通过参考实现进行压缩，再加载到 GPU 内存供 DE 处理，从而将压缩算法的影响与硬件解压缩特性分离。

本文通过系统性改变块大小（32KB、64KB、128KB、256KB）与批处理并发度（1–1024 个并发操作），确定最优并行度水平。

• 峰值吞吐量定义为“效率下降前的最大可持续带宽”；
• 流水线深度定义为“维持约 85%效率（效率定义为‘单次操作吞吐量/峰值单次操作吞吐量’）的并发操作数”；
• 饱和点定义为“额外并发度仅带来约 5%边际吞吐量提升”的临界点。

该方法揭示了 NVIDIA 未公开的硬件资源限制与内存带宽约束。

3. 张量核心表征

本文开发了定制化 GPU 核函数，利用 Blackwell 新引入的张量核心指令集（tcgen05）执行形如 D = A × B + C 的 MMA 操作。通过改变指令类型、矩阵分块形状及操作数布局，测量延迟与吞吐量，以表征执行流水线行为。能效分析通过对比计算吞吐量与板级功耗，确定不同精度模式与分块配置下的能效最优工作点。

为测试“单线程指令改进”的效果，本文在 FP16 输入输出精度下，发起多种 MMA 分块形状的指令。首先隔离指令以测量单指令延迟（SI-LAT），然后对比不同线程束（warp）规模下的表现（需考虑 wgmma 指令仅能在“线程束组”层级发起——每个线程束组包含 4 个异步线程束）。测试中 TMEM 的使用量控制在最低限度（仅用于累加器），标记为“SS”（仅累加器使用 TMEM）。

4. 扩展精度表征

与此前聚焦 FP8、FP16 及 INT8 张量操作的研究[8], [11]不同，本文利用支持 e2m1（FP4）、e3m2（FP6）及 e2m3（FP6）编码格式的 tcgen05 PTX 操作码，开发了首个针对 Blackwell FP4 与 FP6 MMA 指令的系统性基准测试。该方法的核心贡献是“依赖链方法”——可隔离这些超低精度操作的真实指令延迟。

5. 负载基准测试

为评估上述各独立特性及 B200 的整体性能，本文开发了集成负载，可同时触发多种架构创新特性。具体设计如下：

• 选择 Mistral 模型家族[23]作为大型语言模型（LLM）测试对象，原因包括：
  1. Mistral-7B 采用典型的密集型解码器架构，性能与更大规模模型具有可比性；
  2. Mixtral-8x7B 的混合专家（MoE）架构可触发不同数据流模式，对 Blackwell 的内存层级结构形成压力；
  3. Mistral 家族开源可获取，确保实验可复现。该家族涵盖从密集型（Mistral-7B）到稀疏 MoE 型（Mixtral-8x7B、Mixtral-8x22B）的架构多样性，可全面覆盖现代 LLM 部署场景。
• 开发基于 FP64 的定制化矩阵乘法核函数，以测量科学计算负载的实际性能。此外，运行 STREAM Triad 基准测试[24]以测量内存带宽，使用真实世界数据进行稀疏矩阵向量乘法（SpMV）测试以基准化解压缩引擎（DE）性能。
• 采用 ResNet50[25]与 GPT-1.3B[26]的混合精度训练任务，测量端到端训练性能。

本文上述 PTX 微基准测试方法，提供了现有 B200 特性模拟框架中缺失的实证性能数据。通过隔离 TMEM、解压缩引擎及扩展精度张量核心的“单独影响”与“组合影响”，本文为研究人员、HPC 从业者及 AI 框架开发者（针对新兴 GPU 架构的内存密集型与计算密集型负载）提供了可落地的洞察。

五、内存子系统

在本章中，我们通过微基准测试方法对包含张量内存（TMEM）和解压缩引擎（DE）在内的内存子系统进行对比评估，测试内容包括延迟、饱和行为及对访问模式的敏感性。

与前代架构相比，这些新特性可能从根本上改变数据移动模式和内存带宽利用率。

5.1 张量内存（TMEM）

张量内存是每个流式多处理器专用的 256KB 片上内存，仅用于张量核心操作。其结构为 512 列 ×128 通道的 32 位单元二维数组，采用通道-列寻址方式。张量内存将张量核心存储与寄存器分离，使得矩阵中间结果可在跨线程束组中保留，同时减少对全局内存或共享内存的依赖。

我们的延迟表征结果显示，在缓存未命中场景下，张量内存的端到端内存访问延迟为 420 个时钟周期，相比 Hopper 架构 1000 个时钟周期的全局内存延迟降低了 58%。

这一改进源于张量内存的专用仲裁逻辑——该逻辑可绕过传统内存层级中固有的 L2 缓存分区竞争问题。更关键的是，每个 SM 的张量内存可提供 16TB/s 的读取带宽和 8TB/s 的写入带宽，且该带宽与 L1 缓存/共享内存带宽呈叠加关系，而非竞争同一资源。在我们针对 FP8 数据的背靠背矩阵乘加操作微基准测试中，张量内存可持续维持 8TB/s 的内存带宽，与 HBM3e 的峰值性能持平，相比传统 ld.global 路径（因需经过 L1/L2 缓存遍历开销而稳定在 3.8TB/s）实现了 2.1 倍的性能提升。

向张量内存的迁移需要全新的指令序列，因为传统数据移动指令无法与该内存层级交互。在 Hopper 架构中，张量操作的标准流水线依赖 cp.async.bulk.tensor.2d 指令实现从全局内存到共享内存的异步 2D 分块复制，随后通过 ldmatrix 或 wmma.load 指令将操作数暂存到寄存器中，再执行 MMA 操作。这些指令构成了一套成熟的流水线：全局内存复制引擎填充共享内存分块，屏障同步保证数据可用性，最后通过显式加载指令将操作数传输到寄存器文件。

在 Blackwell 架构中，tcgen05 指令家族取代了上述整个序列。其中，tcgen05.cp 指令负责张量数据与张量内存之间的异步传输；tcgen05.ld 和 tcgen05.st 指令提供张量内存与寄存器或共享内存之间的专用加载/存储操作，支持对数据放置的细粒度控制。重要的是，Hopper 架构要求 A、B 操作数矩阵必须先经过共享内存才能被张量核心使用，而 tcgen05 指令允许 MMA 指令直接从共享内存或张量内存读取操作数，且会将累加器结果直接写入张量内存，形成非对称但效率更高的数据流。

我们对不同操作数大小和访问步长的指令级分析揭示了关键性能特征：

张量内存在 64×64 元素分块（FP8 精度下为 4KB）时实现最优效率，这与每个 SM 的 256KB 共享内存容量相匹配，且能充分利用 1024 位内存接口宽度。

这与 Hopper 架构 32×32 的最优分块大小形成显著差异，意味着迁移到 Blackwell 架构的核函数需要进行算法调整。

• 小于 32×32 元素的分块会导致宽内存接口利用不足，仅能达到峰值带宽的 45%；
• 而大于 128×128 元素的分块会触发多阶段传输，引发流水线停顿，导致有效吞吐量降低 30%。

这些带宽饱和点为核函数优化提供了明确指导：

• 矩阵乘法核函数应将计算分解为 64×64 分块以最大化张量内存利用率；
• 链式操作如 Transformer 注意力机制中的查询矩阵与键矩阵的转置相乘，随后执行 softmax 和值矩阵乘法，应将中间结果保存在张量内存中，以利用 16TB/s 的读取带宽提升后续操作性能。

传统 Hopper 架构的操作存在串行依赖链：全局内存读取 → L2 缓存遍历 → 共享内存写入 → 屏障等待 → 寄存器加载 → 最终 MMA 执行。每个阶段都会引入延迟和带宽竞争。而 Blackwell 架构的张量内存配备独立的地址生成单元，可将张量分块直接预取到暂存缓冲区，使 tcgen05.mma 指令能实现数据移动与计算的重叠。当一个线程束组在某一张量内存分块上执行张量操作时，复制引擎会异步填充下一个分块，实现接近理想的双缓冲。这种流水线效率在 MMA 操作存在生产者-消费者关系的负载中尤为明显。

对于形如的链式矩阵乘法，将中间结果保存在张量内存中，相比 Hopper 架构“将中间结果写回全局内存”的方式，在 SM 完全利用的情况下，每秒可减少约 12TB 的数据移动。

能效测量结果揭示了张量内存使用中的细微权衡：对于将 D 矩阵累加器部署在张量内存而非传统共享内存的核函数，在处理大型矩阵（2048×2048）时，我们观察到在计算吞吐量相当的情况下，整个 GPU 板卡的功耗降低了 15%。这一能效提升源于中间结果保存在片上，减少了 L2 缓存抖动和 DRAM 流量。然而，对于工作集可完全放入 L1 缓存的小规模问题，强制分配张量内存会因额外的复制操作引入轻微开销，导致功耗增加 3%-5%。

这些测量结果给出了明确指导：

• 对于具有大型工作集的多阶段张量流水线，应优先使用张量内存；
• 而对于小规模矩阵的单次操作，传统内存层级仍是最优选择。

5.2 解压缩引擎

NVIDIA Blackwell B200 GPU 引入了专用硬件解压缩引擎，相比前代 H100 架构（仅支持软件解压缩）实现了显著的架构升级。

该子系统原生支持多种主流压缩格式，可加速 AI 和 HPC 负载中关键的数据加载与预处理过程。解压缩速率直接决定批处理延迟、GPU 利用率及系统整体吞吐量。

表 I 格式特异性性能分析。输入吞吐量衡量压缩数据处理速率；输出吞吐量衡量解压缩数据生成速率；延迟表示 100MB 数据集的端到端解压缩时间

为表征 B200 解压缩引擎的性能，我们开发了一套定制化微基准测试套件，在受控条件下针对多种压缩格式，测试不同数据大小、压缩比和内存带宽场景下的性能。该设计可对比硬件解压缩与基于软件的 GPU/CPU 解压缩在速率、延迟及与计算重叠能力方面的差异。通过对多种压缩格式的系统性基准测试，我们发现 B200 解压缩引擎硬件存在显著的格式特异性优化。

表 I 展示了各支持格式的综合性能指标，结果显示：不同压缩算法的吞吐量差异显著（42-462 GB/s），这表明解压缩引擎存在专用优化路径。最值得注意的是，Bitcomp 格式实现了 462.4 GB/s 的极高输出吞吐量，且延迟仅为 0.227ms，这可能得益于其针对数值数据的整数特异性优化。

所有测试格式的解压缩延迟均低于 1 毫秒（0.227-1.251ms），这表明无论格式复杂度如何，解压缩引擎都能维持稳定的低延迟性能——即使是最古老的 GZIP 算法，也能实现亚毫秒级响应时间。这种通用低延迟能力使解压缩引擎适用于交互式应用和实时数据流场景。

• Zstandard (zstd) 格式在各类数据中表现出均衡性能：输入吞吐量 77.5 GB/s，输出吞吐量 154.9 GB/s，是通用负载的最优选择。
• Snappy 格式以牺牲峰值吞吐量为代价，优先保证超低延迟（0.894ms），适用于响应时间至关重要的实时应用。
• GZIP 作为较旧算法，仍能维持合理性能（输入 42.0 GB/s，输出 83.8 GB/s），并兼容遗留系统和标准化数据格式。

表 II 压缩比敏感性分析（揭示压缩有效性与输入带宽的反比关系）。所有测量基于 LZ4 格式和 100MB 数据集。

表 II 数据表明，解压缩引擎的性能瓶颈本质上是压缩输入带宽，而非解压缩计算能力。
* 不可压缩的随机数据（压缩比 1.00 倍）实现了 173.23 GB/s 的峰值输入吞吐量。
* 而高压缩比的零填充数据（245.45 倍）输入吞吐量仅为 0.85 GB/s。

这种反比关系表明：对于高压缩比数据，硬件需执行更复杂的解压缩操作，每个输入字节消耗的周期比例显著增加。

尽管输入处理速率差异极大，但所有数据模式的输出吞吐量均保持在 160-220 GB/s 的稳定范围，其中重复模式数据的峰值输出吞吐量达 219.80 GB/s。这种稳定性表明解压缩引擎内部存在约 200 GB/s 的解压缩吞吐量上限，说明其架构设计优先保证持续输出带宽，而非最大化输入处理速率。不同压缩比下输出性能的稳定性证明，解压缩引擎是“解压缩吞吐量受限”设计，而非“输入带宽受限”设计。

表 II 压缩比敏感性分析（揭示压缩有效性与输入带宽的反比关系）。所有测量基于 LZ4 格式和 100MB 数据集。

表 II 中的延迟数据显示，所有数据模式的延迟均保持在较低水平（100MB 数据集为 0.477-0.660ms），这表明无论压缩复杂度如何，硬件都能维持可预测的响应时间。这种时间稳定性对实时应用至关重要（此类应用中可预测性能比峰值吞吐量更重要），也 暗示解压缩引擎采用了复杂的负载均衡机制 以维持稳定服务水平。

表 III 不同块大小的流水线深度特性。流水线深度指维持>85%效率（效率定义为“单操作吞吐量/峰值单操作吞吐量”）的并发操作数。

表 III 的流水线深度分析揭示了块大小与最优并发度之间的明确反比关系：块大小从 32KB 增至 256KB 时，流水线深度从 16 个并发操作降至 4 个。这一模式表明，更大的块大小会导致每个操作消耗更多硬件资源或内存带宽，因此系统需降低并行度以维持最优效率。不同块大小下 85%的效率阈值保持一致，这暗示解压缩引擎采用了复杂的资源管理机制，可避免资源竞争导致的性能下降。

峰值吞吐量的缩放特性体现了解压缩引擎出色的并行处理能力：总吞吐量随块大小增加而提升，从 32KB 小块的 53.8 GB/s 增至 256KB 大块的 151.6 GB/s。这一缩放行为表明，尽管块大小增加会导致最优并发度降低，但单次操作处理的数据量增加足以补偿并行度的损失。相比串行处理，最大加速比达 40-70 倍，这证明 解压缩引擎实现了真正的硬件级并行，可根据负载特性同时处理数十至数百个解压缩流。

饱和点分析显示不同块大小存在不同缩放行为：
* 多数配置在并发操作数达到 1024 时仍能提升吞吐量（尽管超过流水线深度阈值后效率显著下降）。
* 而 128KB 块的饱和点出现在 256 个并发操作，这表明 更大块会引发内存带宽或资源压力，限制其向更高并发度的有效缩放。

这一行为为应用开发者提供了明确指导：需根据块大小选择最优批处理策略，以最大化硬件利用率并维持高效率。

单次请求性能随块大小变化显著：小块（32KB）为 0.75 GB/s，大块（256KB）为 3.21 GB/s，这表明解压缩引擎的基准性能随数据粒度增加而提升。然而，硬件的真正优势体现在批处理中——256KB 块并发处理时，总吞吐量可达 151.6 GB/s。这种缩放行为表明，应用需仔细平衡块大小与并发度，以实现最优性能。

在各块大小的流水线深度阈值以内，硬件效率保持在 90%以上，为最优工作点提供了明确指导。超过流水线深度后，并发度超过硬件最优利用率，效率开始下降——尽管大批量处理仍能提升吞吐量，但效率大幅降低。这一效率曲线表明，应用应根据特定块大小的流水线深度设定目标并发度，以同时最大化吞吐量和资源利用效率。

块大小、流水线深度与峰值吞吐量之间的关系揭示了解压缩引擎架构的根本内存带宽限制：随着块大小增加，流水线深度降低但单次操作吞吐量提升，这表明硬件主要受内存带宽限制，而非计算能力限制。这一架构特性解释了为何解压缩引擎优先保证持续解压缩吞吐量而非输入处理速率——因为内存带宽的稳定利用是性能的主要决定因素。

基于实证表征结果，最优利用策略需根据应用需求和数据特性灵活调整：
* 处理大量小文件的应用：应采用 32KB 块大小和 16 个并发操作，在维持高效率（>90%）的同时最大化总带宽。该配置可实现 53.8 GB/s 的总吞吐量，且单次操作延迟极低，适用于实时数据摄入流水线。
* 处理大文件的应用：应采用 256KB 块大小和 4 个并发操作，以实现 151.6 GB/s 的最大单次操作吞吐量。尽管该配置支持的并发流较少，但更高的单次操作带宽足以补偿并行度损失，适用于处理大型数据集或文件的应用。
* 科学计算负载：可进一步利用 Bitcomp 格式优化性能——该格式对数值数据处理的输出吞吐量达 462.4 GB/s，且延迟仅 0.227ms。

这些发现使开发者能根据特定负载需求，通过选择合适的格式、块大小和并发度，最大化解压缩引擎的利用率。此前受 CPU 限制的解压缩操作，如今已转变为可维持 100+ GB/s 吞吐量的 GPU 加速流水线，这从根本上改变了数据密集型计算的成本效益——不仅支持压缩数据流的实时处理，还大幅缩短了 AI、HPC 和分析领域应用的洞察获取时间。因此，B200 的解压缩引擎为硬件加速数据处理建立了新的性能基准，这一设计或将影响未来专用计算加速器的架构发展方向。

六、GPU 核心微架构

本章将阐述 Blackwell GPU 核心微架构的研究发现，重点介绍张量核心微架构（特别是 tcgen05 指令）、线程块对调度及扩展精度支持。

6.1 第五代张量核心

此前研究表明，张量核心的 PTX 指令会根据操作数精度编译为一系列 SASS 指令（如 HMMA、HGMMA、QGMMA、IGMMA 或 BGMMA）[8,16]。

表 IV Blackwell 张量核心的 SASS 指令映射（与前代 SASS 指令对比）。wgmma 指令仅支持 Hopper 架构；OMMA 指令为 Blackwell 新增，专为八字节浮点格式设计。

在表 IV 中，我们观察到 tcgen05.mma PTX 指令会针对不同精度编译为相应的 SASS 指令，且包含新增指令。分析显示，发起 tcgen05 指令时，会针对每种精度编译为对应的 SASS 操作（详见表 IV）。

表 V Hopper 架构线程束组级 wgmma 与 Blackwell 架构线程束级 tcgen05.mma 张量核心操作的单指令延迟对比。所有测量采用 FP16 输入/输出；张量内存仅用于累加器。

Blackwell 架构引入了 tcgen05.mma PTX 指令，该指令会根据操作数精度编译为不同的 SASS 指令。这与 Hopper 架构统一的 wgmma 指令不同，可在硬件层面实现精度特异性优化。

架构设计上的关键变化是：从 Hopper 架构的“线程束组级（wgmma，128 个线程）执行”转变为 Blackwell 架构的“线程束级（tcgen05.mma，32 个线程）执行”。基准测试结果揭示了这一设计选择对延迟的影响。

表 V Hopper 架构线程束组级 wgmma 与 Blackwell 架构线程束级 tcgen05.mma 张量核心操作的单指令延迟（SI-LAT）对比。所有测量采用 FP16（16 位浮点）输入/输出；张量内存（TMEM）仅用于累加器。

测量结果显示，Blackwell 的单指令延迟比 Hopper 低 2.9-11.6 倍。关键在于，Blackwell 的延迟在不同分块大小下几乎保持恒定（11.0-11.4 个时钟周期），而 Hopper 的延迟随分块宽度线性增加。这证实 Blackwell 采用了完全不同的流水线架构：分块大小仅影响吞吐量，不影响延迟。这表明其采用空间阵列设计（硬件并行处理不同分块区域），而非 Hopper 的时间流水线设计（按时间顺序处理分块的不同阶段）。

此外，线程束级粒度支持更细粒度的调度，减少了同步开销。
* 在 Hopper 架构中，每次 wgmma 操作都需要 4 个线程束同步。
* 而 Blackwell 消除了这一要求。

性能分析显示，在“张量核心利用率受数据可用性（而非计算能力）限制”的内存受限核函数中，这一改进可将调度器停顿（因资源或数据未就绪导致的执行停顿）减少 18%-23%。

表 VI 各支持精度下的张量核心综合性能表征。

进一步扩展分析了各支持精度的性能，表 VI 结果显示：尽管 FP64（44.8 TFLOPS）与 FP4（7702.5 TFLOPS）的吞吐量差异达 177 倍，但延迟差异仅为 1.27 倍（11.2-14.2 个时钟周期）。这证实吞吐量的提升源于并行度增加（更宽的数据通路），而非流水线深度增加。因此，Blackwell 的设计优先级是“在所有精度下维持稳定的低延迟”，确保无论量化水平如何，性能都具有可预测性。

对比“FP16 输入+FP16 累加器”与“FP16 输入+FP32 累加器”的场景，发现一个关键瓶颈：FP32 累加会使吞吐量减半（从 1929.2 TFLOPS 降至 964.6 TFLOPS）。这表明限制吞吐量的是累加器数据通路，而非乘法单元。这意味着存在性能权衡：需高数值精度的应用必须牺牲 50% 性能，而采用 FP16 累加器的推理负载可实现最大吞吐量。

INT8（3927.1 TOPS）的性能略高于 FP8（3851.4 TFLOPS），而 FP4（7702.5 TFLOPS）的性能又优于 FP8。这一优势表明，整数与浮点操作共享相同的执行单元，且整数格式所需的控制逻辑略简单。此外，多数精度的延迟相近，这证实各精度的流水线结构相似，吞吐量提升主要源于并行度增加。

6.2 扩展精度支持：FP4 与 FP6

Blackwell 最显著的改进之一是对 FP4 和 FP6 数据类型的原生硬件支持。

表 VII 各精度下的张量核心吞吐量（单位：TFLOPS/TOPS）。

FP4 格式采用 e2m1 编码（1 个符号位、2 个指数位、1 个尾数位）。Blackwell 支持两种 FP4 变体：MXFP4（微缩放浮点）和 NVIDIA 自研的 NVFP4。
* MXFP4 通过将数据划分为 32 个元素的块，每个块使用 E8M0 格式（8 个指数位、0 个尾数位）的缩放因子，增强了低精度训练的性能。
* 而 NVFP4 则将数据划分为 16 个元素的块，缩放因子采用 e4m3 格式（4 个指数位、3 个尾数位），支持更细粒度的缩放。

尽管 FP4 的精度极低，看似不切实际，但近期量化研究表明，对于推理负载，FP4 可维持可接受的精度。Blackwell 的 FP4 支持包含硬件反量化逻辑，可在矩阵乘法过程中将 FP4 值转换为更高精度（通常为 FP8 或 FP16），从而在实现 FP4 存储和带宽节省的同时，维持计算精度。

另一方面，FP6 格式提供了中间权衡：采用 1 个符号位、3 个指数位、2 个尾数位的编码。与 FP4 相比，FP6 的动态范围（可表示的数值范围）显著更宽；同时，与 FP8 相比，仍能实现 1.33 倍的内存和带宽节省。

表 VII 各精度下的张量核心吞吐量（单位：TFLOPS/TOPS）。

测量结果验证了架构规格：在矩阵操作中，Blackwell 的 FP8 模式吞吐量达 3851.4 TFLOPS，占理论峰值的 96.3%；FP4 操作的吞吐量达 7702 TFLOPS，同样占理论峰值的 96.3%。所有精度下的实际吞吐量均达到理论峰值的 96%-99%，这表明张量核心并非性能瓶颈，真正的瓶颈在于内存带宽和核函数启动开销。

在第七章中，将分析张量核心、张量内存（TMEM）和解压缩引擎（DE）在不同实际负载中的应用效果。

七、性能分析与案例研究

本章对 GPU 在三类关键负载中的性能进行综合实证评估：大型语言模型（LLM）推理、科学计算应用及混合精度神经网络训练。通过与 H200 基准架构对比，量化了 NVIDIA B200 架构创新带来的性能收益。

7.1 实验方法

所有报告的指标均为“10 次预热迭代后 100 次迭代”的平均值。延迟测量包含中位数、95 百分位（P95）和 99 百分位（P99），以捕捉尾部行为特征。

能耗通过 NVIDIA 管理库（NVML API）监测，采样间隔为 10ms，以实现高分辨率功耗分析。

7.2 大型语言模型（LLM）推理

7.2.1 精度模式的影响

评估了四种量化方案对推理吞吐量和模型质量的影响：FP16、FP8（E4M3 编码，张量级动态缩放）和 FP4（E2M1 编码，仅权重量化，采用 NVFP4 16 元素块，激活值为 FP8）。所有实验采用标准化配置：批处理大小 32，序列长度 2048。

表 VIII 不同精度模式下的 LLM 推理性能（批处理大小 32，序列长度 2048）。

研究发现，低精度格式相比 FP16 基准可实现性能提升。

• 对于 Mistral-7B 模型，FP8 和 FP4 分别实现了 1.73 倍和 2.5 倍的吞吐量提升。虽然这些收益未达到理论带宽提升（FP8 为 2 倍，FP4 为 4 倍），但仍代表了实际负载中的显著加速。性能提升源于内存流量减少和缓存局部性改善——随着精度降低，L2 缓存命中率从 68% 提升至 84%。 此外，精度降低还会使负载从“内存受限”转变为“计算吞吐量受限”：B200 的带宽利用率从 FP16 的 67.3% 降至 FP4 的 47.6%，这表明低精度格式能更充分地利用可用计算资源，而非受限于内存子系统性能瓶颈。

• 稀疏混合专家模型从量化中获得的收益比密集型模型更显著。对于 FP4 量化，Mixtral-8x7B 的吞吐量提升达 2.69 倍（76,900 tok/s vs FP16 基准的 28,600 tok/s），高于密集型 Mistral-7B 的 2.50 倍。这一额外收益源于量化优化了专家权重缓存，并减少了专家路由机制的开销。

在两种架构均支持的精度模式下，B200 相比 H200 始终保持性能优势：FP16 和 FP8 模式下，B200 的吞吐量比 H200 高 1.57-1.59 倍。这一缩放因子反映了多方面改进的综合贡献：SM 数量增加（1.09 倍）、张量核心效率提升（1.27 倍）及有效内存带宽改善（1.23 倍）。

最后需注意，量化虽带来显著性能收益，但会导致可测量但通常可接受的模型质量下降：FP8 的困惑度仅增加 1.9%-2.4%（各模型间），而 FP4 的困惑度增加更显著但仍具实用性（7.7%-9.1%）。

7.2.2 批处理大小敏感性

为理解批处理大小与推理延迟的关系，我们以 FP8 精度对 Mixtral-8x7B 模型进行了全面分析，测试不同批处理大小下的性能，结果如表 IX 所示。分析揭示了推理流水线中的不同运行模式。

表 IX 延迟与批处理大小的关系（Mixtral-8x7B 模型，FP8 精度，2048 个 token）

• B200 相比 H200 的性能提升在低批处理大小时达 1.48-1.52 倍，高于高批处理大小时的 1.44 倍。这一性能优势很可能源于“自动流水线重构”——将处理阶段从 18-20 个减少至 8-10 个，实现了低于 20ms 的延迟。在高批处理大小时，系统优先优化最大吞吐量而非单请求延迟，性能比稳定在 1.44 倍：尽管单请求延迟增加，但系统总吞吐量持续提升，在批处理大小 32 左右达到峰值效率。
• 此外，B200 的性能稳定性更优：P99 与中位数延迟比为 1.12-1.14，而 H200 为 1.23-1.38。尾部行为的改善对要求稳定响应时间的生产环境至关重要。

7.3 科学计算负载

7.3.1 FP64 性能

科学计算应用的计算特征与深度学习负载存在根本差异：需高精度算术运算、稳定内存带宽及不规则访问模式。我们采用双精度浮点（FP64）评估密集矩阵乘法性能——该操作对需数值准确性的科学模拟至关重要。

表 XII FP64 精度下的 DGEMM 性能

表 XII 展示了不同矩阵维度下的结果：

• B200 在大型矩阵上的 DGEMM 性能达 36.3 TFLOPS，占其 40 TFLOPS 理论峰值的 80.7%；
• 而 H200 的性能为 18.9 TFLOPS，仅占其 34 TFLOPS 理论峰值的 55.6%。

B200 的效率额外提升 45%（1.92/1.32=1.45 倍），这源于两方面改进：

1. 张量内存支持的累加操作减少了 L2 缓存流量。
2. 优化的内存访问合并模式更充分地利用了可用内存带宽。

7.3.2 稳定内存带宽

内存密集型科学应用需稳定的高带宽数据移动能力。我们采用 STREAM Triad 基准测试，在不同工作集大小下测量可实现内存带宽，结果如表 XIII 所示。

表 XIII STREAM Triad 基准测试的内存带宽

结果显示：对于可放入 GPU 缓存层级的小型数组，H200 的绝对带宽虽低，但效率更高（60-60.6%），而 B200 为 51.7-51.8%——这表明 H200 的设计更针对小型工作集优化。然而，当工作集超过缓存容量时，两种架构的内存带宽利用率均超过 90%，表现优异。B200 的 1.71 倍带宽加速比与原始带宽比（8.0/4.8）高度吻合，这表明内存受限代码的性能随可用带宽线性缩放。

7.3.3 稀疏操作

有限元方法和图负载中的不规则模式，对“为规则执行优化”的 GPU 构成挑战。我们采用稀疏矩阵向量乘法，结合解压缩特性进行测试，结果如表 XIV 所示。

表 XIV B200 上结合硬件解压缩的 SpMV 性能

解压缩引擎在稀疏矩阵上实现了稳定的 3.16 倍加速比：游程编码对稀疏行指针数组实现了 8.2 倍压缩比。专用解压缩硬件引入的延迟开销低于 5%，同时为“指针密集型负载”减少了 35% 的内存流量。

7.4 混合精度训练：端到端训练性能

我们针对不同模型架构进行了全面的训练基准测试，以评估架构改进在实际训练场景中的实际影响，结果如表 XI 所示。

表 XI 端到端训练性能

训练性能实现了稳定的 1.54-1.56 倍加速比，这一收益可分解为：

• SM 数量增加 1.09 倍
• 线程块对调度 1.27 倍
• 张量内存优化 1.26 倍

尽管 GPT 训练的功耗增加了 14%，但其能效提升了 42%。

八、讨论

表 X 全面汇总了所有评估负载类别的性能改进，重点标注了每项性能提升对应的具体架构特性。

表 X 各负载性能汇总

8.1 架构权衡

张量内存、双模张量核心和解压缩引擎的引入，使晶体管数量从 H200 的 1800 亿增加到 B200 的 2080 亿，但相应带来了 1.5-3.9 倍的性能提升。每个 SM 配备的 256KB 张量内存占 SM 内存的 10%，其命中率达 61-82%， 这验证了该内存大小的设计合理性。

8.2 软件生态

CUDA 13.0 提供了对张量内存和线程块对的初步支持，但框架集成仍在进行中。尽管 FP6 的硬件支持已存在，但相应的软件工具链仍不完善。FP4/FP6 的使用需要“逐层精度选择”——FP4 导致的 8.2% 困惑度下降为平均值：部分网络层可容忍 FP4 精度，而其他层则需 FP8 以维持模型质量。

8.3 部署建议

• LLM 推理：B200 相比 H200 提供 1.8-3.9 倍的性能优势；FP4 精度对 70B 规模模型具有实用性（可平衡性能与精度）。
• 训练任务：1.54-1.56 倍的性能提升支持 33% 的批处理大小增加（更大批处理可加速训练收敛）。
• HPC 领域：FP64 性能 1.92 倍的提升，使 B200 在科学计算中具备竞争力。

九、结论

NVIDIA B200 GPU 标志着 GPU 架构的重大变革。本文首次基于微基准测试套件，对 NVIDIA Blackwell B200 GPU 进行了详细表征分析，深入揭示了其架构创新与性能表现。

具体而言，我们量化了张量内存（TMEM）对矩阵密集型负载的影响，评估了硬件解压缩引擎（DE）的吞吐量及最优使用方式，并通过新的 tcgen05 PTX 指令分析了第五代张量核心的执行特性。

此外，我们还评估了 FP4 与 FP6 精度的权衡关系，在大型语言模型（LLM）推理、科学计算核函数及混合精度训练等多样化负载中对 Blackwell 进行了基准测试，并为面向该下一代架构的开发者提炼了可落地的性能优化指南。

参考文献

关注“鲸栖”小程序，掌握最新AI资讯