Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks

文章标题：通过微基准测试剖析NVIDIA Blackwell架构
作者/机构：Aaron Jarmusch, Nathan Graddon, Sunita Chandrasekaran (University of Delaware)

A1 主要贡献

本文旨在通过微基准测试对NVIDIA Blackwell GPU架构进行微架构层面的分析，并与前代Hopper架构进行比较。研究的核心问题在于，面对层出不穷的新型GPU架构，如何确定哪种架构最适合特定的工作负载。尽管已有多种性能分析工具和方法，但由于现代商业GPU缺乏公开文档，深入的微架构分析仍然受限。因此，本研究通过自行设计的微基unchmarks，揭示Blackwell架构的关键子系统，如内存层次结构、流式多处理器（SM）执行管线和第五代张量核心（Tensor Cores），为应用开发者、编译器编写者和性能工程师提供优化工作负载的实践性见解。

核心贡献如下：

• 开发新型微基准测试：创建了一套新的微基准测试，用于评估NVIDIA Blackwell架构的关键组件，并与前代Hopper架构的GPU进行比较。
• 深入分析SM与内存层次结构：对流式多处理器（SM）子单元（包括支持FP4/FP6的下一代张量核心、统一的INT32/FP32核心和FP64执行核心）以及内存层次结构进行了深度分析。
• 提供性能优化指南：为软件和应用程序开发者提供性能指导，帮助他们有效利用新硬件的优势。
• 探索新型低精度数据类型：研究了Blackwell张量核心中新支持的低精度浮点数据类型（FP4和FP6）的行为及其性能影响。

A3 架构概述

GB203与GH100的设计哲学差异。尽管GB203基于GH100，但它们代表了两种不同的设计理念。GH100（NVIDIA Hopper架构GPU芯片）专为大规模人工智能和科学计算工作负载优化，而GB203（Blackwell架构的一部分）则是一款注重能效、面向消费者的GPU，旨在支持游戏、渲染和小型推理任务。两者都实现了相似的CUDA核心编程模型，但在指令集、执行单元数量、内存层次结构和资源调度方面存在差异。

1. 流式多处理器（SM）与执行管线。两种架构的核心都是流式多处理器（SM），负责warp调度、指令发射和执行。虽然高层SM设计相似，但存在一些差异，如表I所示。Blackwell改进了warp调度，降低了发散工作负载的warp分派延迟。而GH100则为训练规模的工作负载提供了更高的执行吞吐量和更大的片上缓冲。
2. 缓存层次结构。两种GPU的一个显著区别在于缓存和内存子系统。表II展示了详细对比，GB203通过增加L2带宽来弥补其较小的L1缓存。GH100的HBM内存支持更大的批量大小和工作集容量，这对于训练大型模型很有帮助。
3. 张量核心与AI加速。张量核心是为加速矩阵运算而设计的专用单元。Hopper引入了支持FP8工作负载的第四代张量核心，而Blackwell则升级到第五代，扩展了对FP4/FP6格式的支持，同时保持了对FP8工作负载的支持。
4. 指令集与软件兼容性。Hopper架构通过PTX和CUDA 11.8支持wgmma和FP8算术等张量核心指令，并与传统的mma和CUDA C++指令保持向后兼容。CUDA 12.8和PTX 8.7则扩展了对Blackwell第五代张量核心指令的支持，包括tcgen05以及针对FP4和FP6格式的增强操作数类型。虽然Hopper的wgmma指令与Blackwell不兼容，但可以使用新的tcgen05指令进行warp-group计算。

总结与测试平台。总而言之，GH100是一个拥有庞大内存资源的训练优化架构，而GB203则专注于在散热限制内最大化能效。尽管存在差异，两者共享的微架构相似性使得可以采用统一的基准测试策略。本文后续章节将通过针对性的微基准测试和分析，深入剖析Hopper H100 PCIe（GH100芯片）和Blackwell GeForce RTX 5080（GB203芯片）GPU的各项特性，并重点分析Blackwell架构，同时与前代架构进行对比。

TABLE I GH100和GB203 GPU上的执行单元。

A2 方法细节

计算管线

GPU上的大部分计算由计算管线完成，其核心是流式多处理器（SM）。SM负责分配计算资源、调度指令和移动数据。为完成此任务，每个SM包含四个执行单元分区或子核心，用于整数、浮点和张量操作。这些子核心的布局和硬件实现可能提升或阻碍应用性能。了解这些资源可以帮助开发者改进应用并充分利用GPU。为了对GB203和GH100进行全面评估，本文采用了以下指标和分析方法：

1. 延迟（Latency）：指一条指令产生可被后续指令使用的结果所需的时钟周期数。我们测量了两种延迟：
   • 真实延迟（True Latency）：在串行化、相关的指令链中测得，反映了无并行或重叠时的数据就绪延迟。
   • 完成延迟（Completion Latency）：在一组允许重叠和并行的独立指令中测得。
     两种延迟均以每个指令的时钟周期数（#clock cycles/#instruction）报告。
2. 吞吐量/带宽（Throughput/Bandwidth）：根据测量的运行时间和指令计数，衡量每个SM每时钟周期完成的指令数。

微基准测试实现方法。我们的微基准测试使用PTX指令和CUDA编写。为了防止编译时优化，我们没有在CUDA文件中使用内联指令，而是在独立的PTX内核文件中编写PTX内核，这些内核在运行时由编译后的CUDA文件执行。为了确认PTX指令在运行时未被优化，我们审查了每个PTX内核文件生成的SASS代码，以确保指令未被优化。

TABLE III GH100 vs GB203 延迟结果（真实延迟/完成延迟）。

A. 时钟开销

时钟周期测量方法。所有时钟周期均使用%clock64进行测量，这是一个预定义的只读特殊寄存器，读取时返回时钟周期的计数值【索引26，Parallel Thread Execution (PTX) ISA, Release 8.8, 2025】。例如，清单1中的代码在指令执行前后获取计数器值，然后将两者相减得到测量的时钟周期。

.entry my_kernel(param_u64_1) {
    .reg .b64 start, end, a;
    mov.u64 start, %clock64; // 开始计时
    mad.lo.s32 %r0, %r0, 1, 1; // 目标指令
    mov.u64 end, %clock64; // 结束计时
    sub.s64 a, end, start; // 计算差值
    st.global.u64 [%rd0], a; // 存储结果
    ret;
}

测量开销分析。在GB203上，如果两个计时寄存器之间没有指令，相减的值为1，而GH100为2。此外，当包裹一组混合指令（即混合工作负载）时，该值取决于具体的指令，这有助于揭示指令的工作流程。

B. INT和FP32执行单元

架构演进与统一核心。在Volta和Ampere等早期架构中，INT32和FP32指令通过独立的执行管线发布，当工作负载由单一类型主导时，常导致利用率不佳【索引7，Dissecting the NVIDIA volta GPU architecture via microbenchmarking, 2018】。在GB203中，NVIDIA引入了可同时处理INT32和FP32指令的统一执行单元，从而能够根据指令混合进行动态调度，这可能减少空闲周期并提高混合工作负载的吞吐量。然而，这些统一核心在任一时钟周期内只能作为INT32或FP32操作使用，这可能在混合INT32/FP32工作负载中产生冲突。

实验设计与结果。为了精确测试这些统一核心的功能，我们使用PTX指令fma和mad分别执行标准的FP32和INT32算术和整数操作。我们比较了三组内核的结果：（A）纯INT32，（B）纯FP32，和（C）混合INT32/FP32指令，以代表利用两种操作的混合工作负载。每个工作负载执行1024次并取平均值，以提供对这些核心的无噪声理解。

性能分析。我们的测量结果显示，在两种GPU上，纯INT32和FP32工作负载的真实延迟均为四个周期。在纯内核的完成延迟方面，GH100稍好一些。尽管GH100对INT32和FP32操作使用独立的执行管线，但我们观察到，在执行混合指令序列时，GH100的性能差于GB203（见表III）。这表明Blackwell上引入的统一INT32/FP32执行核心拥有更高效的计算管线。值得一提的是，GB203在首次运行纯内核和混合内核时，真实延迟和完成延迟都较高，这一结果未包含在表III中。其他研究也曾因缓存预热前的缓存未命中而排除了类似结果【索引27，A performance model for gpus with caches, 2015】。然而，在Hopper上并未出现这种情况。

结论。Blackwell架构在混合工作负载的延迟方面表现出改进，而Hopper在纯指令工作负载方面表现更佳。

C. FP64执行单元

FP64单元的重要性与架构配置。双精度（FP64）执行单元对于需要高数值精度的科学模拟等工作负载至关重要。尽管人工智能和图形应用越来越依赖低精度格式（如FP16、BF16、FP8），FP64仍然是HPC和研究领域的关键特性。GH100和GB203都包含一组专用的FP64执行单元，与INT32/FP32单元物理分离，这使得调度器可以独立发布FP64指令。GB203芯片每个SM有两个FP64执行单元，而GH100则有64个（见表I）。

性能测试与分析。我们的微基准测试显示了预期的结果，表III显示，对于1024条FP64相关指令，GH100的延迟低于GB203。当只执行两条相关指令时，GB203的延迟降至37.5个周期。通常情况下，运行更多指令可以隐藏延迟，但在这种情况下，由于只有两个执行单元，这表明这两个单元仅用于类型和指令支持，而计算可能意在通过其他精度（如使用FP32执行单元或张量核心）来模拟。

对开发者的启示。这些发现对于希望在数据中心级和消费级GPU上实现可移植性能的用户尤其重要，因为了解FP64的瓶颈可以为精度权衡或算法设计提供信息。

D. Warp调度器行为和发射模型

实验设计。为了评估warp调度敏感性和在依赖链下的延迟处理能力，我们实现了一个串行化的依赖指令基准测试，其中每个线程在寄存器中为每个执行单元执行一连串依赖的算术操作。图2和图3显示了当每个线程的依赖指令数从1增加到1024（图中显示1到64的结果）时的结果。为了公平比较总周期和吞吐量，我们通过调整每个链长度的循环迭代次数来控制执行的总指令数。

吞吐量分析。如图3所示，在前1-9条指令内，吞吐量稳步增长，之后管线在不同架构间出现差异，并且所有INT32、FP32和FP64工作负载的吞吐量都逐渐提升。短依赖链的低吞吐量是由于每个线程内部的指令级并行（ILP）不足，这限制了GPU隐藏指令延迟的能力。只有几条依赖指令时，线程很快会停顿，调度器无法充分利用执行单元，导致性能较低。随着链长度增加，调度器能更好地重叠执行并隐藏延迟，从而提高吞- 吐量。

架构差异对比。值得注意的是，与GH100较为不规则的爬升相比，GB203的吞吐量增长更平滑、更一致。然而，GB203在FP64指令上的吞吐量低于其INT32或FP32的对应指令。同样，图2显示，在8条指令之前，GH100实现了更低的总周期数，表明在短依赖链下，其延迟隐藏能力比GB203更有效。在8条指令之后，两种架构的总周期数都急剧下降，这可能是由于指令调度预热或管线效应。然而，与吞吐量类似，随着指令数的增加，GH100的总周期数波动更大。

调度策略推断。GH100在少量依赖指令下稍好的吞吐量可能是由于更深的指令缓冲和更激进的warp调度器，它能在压力较大时容忍指令依赖。然而，激进的调度也可能在指令数较多时引入不稳定性。另一方面，GB203更平稳的进展表明其发射策略更为保守。不过，其在少量依赖指令下较高的总周期数仍需进一步分析。

结论。虽然两种GPU的整体性能相当，但GH100能更好地容忍短的延迟受限指令序列，而GB203则为更规则的高ILP内核进行了优化。

第五代张量核心

张量核心简介。张量核心自Volta架构引入以来，是专门用于加速矩阵乘法的单元，这是深度学习和科学计算中的基础操作。为了充分利用这些张量核心，我们使用包含矩阵乘法操作的内核，例如矩阵乘法累加（MMA）。为了观察GB203和GH100的行为，我们定制的PTX微基准测试测量了执行延迟、吞吐量和操作数暂存行为。本节评估了GB203和GH100张量核心在不同指令级并行度（ILP）和warp数量下的能力。

A. 指令集和支持的数据类型

两代张量核心的对比。第五代GB203和第四代GH100的张量核心均支持具有不同数据类型、操作数处理和性能调优的指令。表IV比较了第四代和第五代张量核心支持的指令和数据类型精度。Blackwell的第五代张量核心引入了新的数据类型（FP4和FP6），并通过新的SASS级指令（如OMMA、QMMA）在CUDA中实现，反映了硬件对低精度格式的支持。而Hopper则支持wgmma指令，该指令支持warp-group异步矩阵操作，但缺乏对FP4和FP6的支持。

B. 可变MMA和基于Tile的指令

指令支持现状。尽管GB203不支持wgmma指令，且GB203的tcgen05指令尚未实现，但NVIDIA在两种GPU中都实现了带有相应数据类型的mma指令，可供分析。

MMA操作详解。矩阵乘法累加（MMA）操作支持GEMM和深度学习工作负载的矩阵计算。每个mma指令指定一个tile形状，表示为M×N×K，这决定了每个warp或每个threadgroup处理的矩阵片段的维度。例如，公式1中的指令使用16×32（M×K）和32×8（K×N）的输入计算一个16×8（M×N）的输出tile。

mma.sync.aligned.m16n8k32.f32.f16.f16.f32

支持的形状与精度。还支持多种其他tile形状，如m8n8k16或m16n8k64，这些形状支持更细的粒度或每次指令发射时更大的操作数复用。除了tile形状，mma指令还支持多种输入/输出精度，包括但不限于FP4、FP8、FP16和FP32。这些数据类型在公式1中编码，其中f16和f32表示FP16输入，FP32累加和输出。

Blackwell新指令集。在为Blackwell发布的最新指令集CUDA 12.9中，必须在PTX指令上明确指定.kind::f8f6f4后缀才能在GB203上使用FP6或FP4的mma操作。尝试在GH100上使用这些格式或不带kind说明符都会导致PTX错误。表V显示了我们在微基准测试中跨精度格式测试的矩阵形状和在mma.sync.aligned.kind::f8f6f4之后使用的PTX指令。

TABLE V

第四代和第五代张量核心上支持的、并用MMA指令测试的数据类型（D-TYPES）的比较。E8M0仅用于块内缩放指数，因此未进行测试【索引26，Parallel Thread Execution (PTX) ISA, Release 8.8, 2025】。

PTX到SASS的转换。通过我们的实验，PTX级别的mma.sync指令被翻译成OMMA、QMMA或HMMA SASS指令。通过观察GH100上生成的SASS指令，我们发现每个mma.sync对每种数据类型都使用了HMMA指令。对于Blackwell，CUDA Binary Utilities 12.9文档【索引28，CUDA Binary Utilities - Instruction Set Reference, 2025】指出，QMMA用于整个warp的FP8矩阵乘法累加，而OMMA用于FP4。我们的微基准测试证实，两种FP8输入格式以及两种FP6输入格式都使用了新的QMMA指令。虽然FP4输入的mma本应使用OMMA SASS指令，但我们观察到实际使用的是QMMA指令。然而，当使用带有FP8 ue8m0作为缩放格式的块缩放时，在SASS代码中观察到了OMMA。这表明在当前软件中，QMMA是FP4输入的备用方案。

总结。总而言之，由于NVIDIA刚开始开发这些功能，软件支持有限，我们希望提供对这些管线可用性的当前理解和深入分析。

C. 精度权衡

低精度格式的优势与权衡。低精度格式被用来减少内存占用并提高吞吐量，尤其是在推理工作负载中。前一节提到的支持的数据类型（FP4、FP6、FP8）都有不同的格式，并且都属于低精度。这些数据类型格式通过调整指数位和尾数位的数量来减少表示浮点数所用的比特数，因此在动态范围和精度之间存在权衡。

性能功耗分析。以往的研究致力于理解这些低精度格式的准确性，本节我们将分析这两种架构上这些低精度格式的每瓦性能和功耗。

实验结果。值得注意的是，表VI总结了我们在GB203和GH100上的微基准测试结果。GH100缺乏对FP4和FP6格式的原生支持，但在两种FP8格式下功耗稍高（约55W），而GB203在相同格式下的最大功耗为46W。功耗通常随着精度降低而减少，FP4的功耗最低，为16.75W，而FP6和FP8格式的功耗分别超过39W和46W。

结论。这些结果表明Blackwell在低精度下的架构效率，并展示了在mma张量核心工作负载中数值表达能力与能耗之间的权衡。

D. Warp扩展和共享内存访问

实验设计。通过我们实现的低精度输入mma微基准测试，我们改变ILP和warp数量来考察指令映射，并比较GB203和GH100的warp调度行为。

ILP扩展能力。对于每种精度格式，在逐渐减少的warp数量下，能够达到持续吞吐量的最大ILP水平是：GH100为ILP=5，有29个活动warp；GB203为ILP=6，有25个活动warp。这意味着与GH100较低的ILP扩展能力相比，Blackwell能够为每个线程发布更多独立的mma指令。

延迟与吞吐量分析。当ILP=1且warp=1时，从指令发射到数据可用的时钟周期数即为完成延迟。对于GB203，所有精度格式的完成延迟均为1.21094个周期。GH100的完成延迟为1.65625个周期，这表明在各自的架构上，所有低精度格式的mma指令使用相同的执行管线。同样，GB203在所有低精度格式上的吞吐量都高于GH100，在6个ILP和32个活动warp时达到峰值超过11 TFLOP/s。这表明增加ILP能显著提升低warp数量下的吞吐量，证实了当并发性有限时，Blackwell的warp调度器能有效利用线程内并行性。

总体趋势。我们对每种低精度格式的ILP延迟和吞吐量进行了平均，以呈现Blackwell和Hopper的趋势。如图4所示，GB203在每种精度格式上的吞吐量都有所提高。同样，图5展示了延迟随格式变化的扩展情况。GB203持续保持较低的延迟，尤其是在FP4和FP6上，而GH100随着warp数量的增加，延迟出现阶梯式增长，这是更深但不够灵活的调度队列的标志。这表明GH100需要更多的在途warp来饱和执行单元，而GB203在warp较少但独立指令较多的情况下表现更佳。综合来看，这些结果表明Blackwell的warp调度器针对低精度、高ILP、控制流清晰的工作负载进行了优化，而Hopper则依赖于大规模并发和更深的缓冲来在不那么规则的条件下维持性能。

架构权衡。这一比较表明，Blackwell针对更高的每线程指令吞吐量进行了优化，而两种架构无论数据格式如何，都具有相似的warp调度能力，反映了它们在张量核心微架构上的不同权衡。我们的方法可作为评估未来架构上张量核心性能的参考框架，并突出了在warp级别上精度、吞吐量和执行行为之间的关键权衡。

内存子系统

内存瓶颈的重要性。GPU性能越来越受到内存子系统行为的限制，而非原始计算吞吐量。高效利用内存层次结构，包括共享内存、各级缓存和全局内存，对于实现架构效率至关重要。虽然GH100和GB203都采用了相似的内存布局，但它们在延迟、带宽和容量方面表现出明显的权衡。本节通过微基准测试方法对内存子系统进行比较评估，测量其延迟、饱和行为以及对访问模式的敏感性。

A. 内存层次结构概述

研究范围。本研究专注于设备级内存访问，不包括主机-设备传输性能，后者受系统互连（如PCIe与NVLink）的严重影响。GPU内存访问模式针对全局内存、共享内存和硬件管理的缓存层（L2、L1、L0指令缓存），以及寄存器文件。

延迟测量方法。为了分离延迟特性，我们采用指针追逐（pointer-chase）微基准测试，进行随机串行化的内存访问。图6展示了GB203和GH100在数据大小增加时的延迟（以周期为单位）。

缓存区域划分。图中显示了三个缓存区域：
1) L1缓存：范围从0到（约128KB或约256KB）
2) L2缓存：范围从L1结束到（约30MB或约60MB）
3) 全局内存：超出L2缓存范围

延迟的突增与缓存边界相对应，这与架构规格（表II）一致。

B. 共享内存和L1缓存行为

统一内存空间。现代NVIDIA GPU将共享内存和L1缓存组合在每个SM的统一内存空间中。为了评估GB203和GH100芯片上这种统一设计的性能和特性，我们开发了微基准测试来测量访问延迟趋势、bank冲突敏感性和warp扩展行为。

L1缓存延迟。从图6的指针追逐基准测试中可以看出，两种GPU在L1缓存区域的延迟几乎相同，稳定在30-40个周期，表明硬件管理的数据路径中的命中延迟相似。尽管架构不同，这表明L1访问路径经过了良好优化。

容量差异与配置。然而，缓存容量差异显著。GH100每个SM拥有高达256 KB的L1/共享内存组合，而GB203则减少到128 KB/SM。GH100和GB203都向软件开放了这部分内存的可配置部分作为共享内存。通过使用cudaFuncSetAttribute和cudaFuncAttributeMaxDynamicSharedMemorySize属性进行动态分配，我们确定了可配置的共享内存限制在GH100上约为227 KB/SM，在GB203上约为99 KB/SM。不进行动态分配时，两种架构的默认静态共享内存限制均为48 KB/SM。

访问行为微基准测试。为了探究访问行为，我们设计了两个微基准测试。对于共享内存，我们访问一个静态声明的共享数组，步长和warp数可配置。类似地，对于L1缓存，我们通过restrict float* gmem访问全局内存，工作集设计为适合L1缓存容量，并通过跨步加载引发冲突。两个基准测试都扫描了1到32个warp，步长为1和4，进行了32次内存访问。每个测试重复1024次，并记录中位延迟。

共享内存延迟分析。图7显示了共享内存延迟随warp数量增加的变化。对于两种步长，GB203在低warp数量（1-4个warp）时表现出较低的延迟，表明在轻负载下有更优化的路径。然而，GH100在较高warp压力（6-32个warp）下表现更佳，这可能归因于其更大的共享内存容量，但也可能来自更强大的bank冲突缓解机制。

L1缓存延迟分析。在步长为4时，GH100在较高warp级别保持了更平滑的扩展和更低的延迟，表明对访问倾斜的容忍度更高。相比之下，GB203在步长为4时延迟增长更陡峭，这可能是由于bank争用及其较小内存分区的饱和。图8显示了随着更多warp访问L1缓存的延迟趋势。虽然两种架构在首次访问时延迟都有所增加，但GB203在步长为1的情况下，从2到11个warp时延迟略低。与共享内存相比，L1延迟在不同warp数下的变化更为平缓，尤其是在GH100上，这可能是由于空间局部性和L1对冲突的更高容忍度。然而，在GB203上步长为4时，延迟增长更急剧，表明即使L1对步长的敏感性低于共享内存，访问倾斜仍然影响性能，尤其是在较小的内存分区容易饱和的情况下。

总结。总体而言，共享内存延迟对warp数量和访问步长高度敏感，尤其是在GB203上，bank冲突的扩展更为剧烈。相比之下，L1缓存表现出更高的延迟，但随着warp增加，其弹性更好，共享内存和L1缓存的延迟在32个warp时趋于一致。GH100更大的统一内存和平滑的warp扩展使其在具有密集复用的高线程内核中具有优势。而GB203则改进了低延迟访问路径和在少量warp下的冲突解决能力，这可能是通过多端口bank或warp感知调度等微架构增强实现的。这些结果表明，在warp扩展内核设计方面有所改进，尤其对GB203而言，但该芯片仍受限于内存分区的大小。

C. L2缓存

架构设计差异。GH100和GB203 GPU中的L2缓存架构反映了两种不同的设计，特别是在处理分区和负载下的扩展方式上。L2缓存位于全局内存和SM之间，是最大的片上内存块。在GH100中，L2缓存被分为两个独立的分区，每个分区服务于一部分GPC（图形处理集群）。这种分区设计支持更好的数据局部性和缓存访问的并行性。相比之下，GB203采用由所有GPC共享的单片L2缓存。这种统一的方法简化了全局内存路由和一致性，并可以改善小型或面向图形的工作负载的空间局部性。然而，当许多SM同时发出未缓存或流式内存访问时，也可能导致更大的争用。

延迟测量对比。延迟测量突显了这些差异。对于标准的L2命中，GB203表现出约358个周期的固定延迟，而GH100则实现了约273个周期的较低延迟。GH100的延迟优势可能源于其分区设计，该设计通过将访问分布在两个单元上来减少争用。然而，随着内存需求的增长，GH100的优势减弱：当两个分区都饱和时，对于31 MB到45 MB的内存大小，延迟增加到约508个周期。相比之下，GB203在其内存占用范围内能更长时间地维持其基线延迟，这归功于其更大的总L2容量（65 MB，而GH100为50 MB）。

Warp级并发性能。为了理解这些架构在warp级并发下的性能，我们开发了一个微基准测试，每个线程发出1024次全局内存加载/存储操作，并跟踪每个warp的周期计时。这个设置使我们能够评估L2吞吐量随warp数量增加的扩展情况。

结果分析。图9显示，在低warp数量（1-4）时，GH100始终提供更好的性能，平均每个warp的周期时间约为43.5k，而GB203为49k。这种差异不仅反映了GH100更快的L2延迟，还反映了其更深的warp调度器管线和更高效的缓冲。在8-16个warp的范围内，GH100保持其优势且性能下降最小，而GB203开始显示饱和迹象，在16个warp时达到约66k周期。这表明随着并发内存压力的增长，GB203的单个L2接口成为瓶颈。有趣的是，在高warp数量（16-32）时，GB203追赶上来，并在20个warp时最终略微超过GH100。在32个warp时，GB203以每个warp约128.4k周期的速度完成基准测试，而GH100为约128.9k。这种转变反映了GB203在极端负载下更高的聚合L2带宽，这可能是其更大缓存大小和减少的分区开销的结果。虽然GH100在低到中等并发水平下提供了一致的性能和确定性的warp调度，但在全压下，受其分区仲裁的限制，它达到了吞吐量上限。

架构适用性。这些趋势表明，GH100更适合对延迟敏感且在中等并发下运行的动态工作负载，这得益于其激进的warp调度和分区缓存布局。另一方面，GB203在充分利用下提供卓越的性能，使其更适合大规模、带宽受限的应用，如深度学习推理或密集矩阵运算。总之，GH100的分区L2架构针对高并发和计算密集的服务器级工作负载进行了优化。GB203的统一L2设计简化了硬件复杂性，并倾向于混合计算-图形用例。在为内存受限内核中的特定性能目标进行调优时，无论是优先考虑延迟、吞吐量还是数据局部性，都必须考虑这些架构上的权衡。

D. 全局内存

带宽测试。我们使用一系列持续传输基准测试将分析扩展到全局内存带宽。如图10所示，GH100实现了15.8 TB/s的峰值读取带宽，远高于GB203的8.2 TB/s。两者的写入带宽都较低，分别为2.2 TB/s（GH100）和1.6 TB/s（GB203），这表明架构设计倾向于读取密集型工作负载，可能原因在于较窄的回写路径或不够积极的写合并策略。

延迟分析。图6中观察到的延迟趋势表明，全局内存访问在71 MB（GB203）和55 MB（GH100）之后开始，相应的延迟分别为约876.7个周期和约658.7个周期。GH100卓越的延迟性能归功于其使用的HBM2e内存，该内存比GB203的GDDR7提供更高的带宽和更低的延迟。

Blackwell的潜在问题。尽管Blackwell架构在内存调度和子系统设计上引入了显著的增强，但这些变化可能导致在不规则或延迟敏感的工作负载中一致性降低。

A4 实验环境

• 硬件配置:
  • GPU:
    • NVIDIA Blackwell GeForce RTX 5080 (GB203芯片)
    • NVIDIA Hopper H100 PCIe (GH100芯片)
  • 互连: 未明确指定，但提到了PCIe和NVLink作为影响主机-设备传输的因素。
• 软件配置:
  • 核心库: CUDA, PTX。
  • CUDA/PTX版本:
    • Hopper (GH100): CUDA 11.8。
    • Blackwell (GB203): CUDA 12.8, PTX 8.7, CUDA 12.9。
  • 工具:
    • NVIDIA cuBLASLt API (用于D-GEMM案例研究)。
    • NVIDIA TensorRT 10.0 (用于Transformer推理案例研究)。
    • nvidia-smi (用于测量功耗)。
• 模型与数据集:
  • 模型: GPT-neox (用于Transformer推理案例研究)。
  • 数据类型: 实验中测试了多种数据类型，包括INT32, FP32, FP64, FP16, BF16, FP8 (nv_fp8_e4m3), FP6, FP4。
  • 工作负载: 微基准测试使用不同长度的依赖指令链（1到1024）和不同数量的Warp（1到32）。D-GEMM案例研究使用了不同规模的矩阵（边长从1024到8192）。

A4 实验结果

计算管线性能

• INT32/FP32单元：在纯INT32和FP32工作负载中，GH100和GB203的真实延迟均为4个周期。但在混合指令工作负载中，采用统一执行单元的GB203性能优于使用独立管线的GH100，表明其调度效率更高（见表III）。
• FP64单元：GH100拥有64个FP64单元，而GB203仅有2个，因此GH100在FP64计算上延迟显著更低。GB203的FP64性能表明其FP64单元主要用于功能支持，而非高性能计算（见表III）。
• Warp调度器：在短依赖链（1-8条指令）下，GH100能更有效地隐藏延迟。随着依赖链变长，GB203的吞吐量增长更平滑和稳定，而GH100则表现出更多波动。这表明GH100的调度器更激进，适合短延迟受限序列，而GB203则优化了规则的高ILP内核（见图2和图3）。

张量核心性能

• 低精度支持：Blackwell (GB203) 引入了对FP4和FP6的硬件支持。实验证实，这些新格式的mma指令在GB203上通过新的QMMA SASS指令执行（见表IV和V）。
• 功耗与效率：GB203在低精度格式下展现出卓越的能效。其FP4功耗仅为16.75W，远低于FP8的46W。相比之下，GH100在FP8格式下功耗约为55W，且不支持FP4/FP6。这表明Blackwell在能耗和数值表达能力之间提供了新的权衡（见表VI）。
• 吞吐量与延迟：GB203在所有测试的低精度格式上均实现了比GH100更高的吞吐量和更低的延迟。GB203能以更少的活动Warp和更高的ILP达到峰值性能，表明其Warp调度器能更有效地利用线程内并行性（见图4和图5）。

内存子系统性能

• L1缓存与共享内存：GB203在轻负载（1-4个Warp）下L1和共享内存访问延迟更低。但在高Warp压力下，拥有更大容量（256KB vs 128KB）的GH100表现更佳，尤其是在存在bank冲突的跨步访问模式下（见图7和图8）。
• L2缓存：GH100的分区L2设计在低到中等并发下提供了更低的延迟（273周期 vs 358周期）。然而，GB203的单片式、更大容量（65MB vs 50MB）的L2在极端并发（>20个Warp）下展现出更高的总带宽，最终性能反超GH100（见图9）。
• 全局内存：得益于HBM2e内存，GH100的全局内存带宽和延迟远优于使用GDDR7的GB203。GH100的峰值读取带宽为15.8 TB/s，而GB203为8.2 TB/s（见图10）。

案例研究

• D-GEMM (FP8)：在FP8密集矩阵乘法中，Hopper (GH100) 的性能全面优于Blackwell (GB203)。Hopper的运行时间更短、吞吐量更高（在8192³矩阵下为0.887 TFLOP/s，是Blackwell的近4倍），且功耗更稳定（约60W）。Blackwell不仅性能较低，且功耗随矩阵尺寸增大而急剧上升，峰值超过114W，表明当前软件栈或编译器启发式未能充分利用其硬件潜力（见图11、图12和表VII）。
• Transformer推理 (GPT-neox)：在使用TensorRT进行推理时，Blackwell展现出更好的功耗调节能力。随着精度的降低（从FP16到FP8），Blackwell的功耗从58.8W降至45W，而Hopper则稳定在57-60W。这表明Blackwell在可调的推理工作负载中具有更好的能效潜力（见表VIII）。

TABLE VII GH100和GB203 GPU上的D-GEMM吞吐量

A5 结论

本研究通过精心设计的微基准测试，对NVIDIA的Blackwell架构（GB203芯片）进行了详细的实验性分析。通过将其微架构特性与Hopper架构（GH100芯片）进行对比，我们揭示了Blackwell在内存层次结构、SM执行管线以及第五代张量核心方面的进步。我们的分析强调了该硬件对FP4和FP6等低精度格式的增强支持，并揭示了它们在功耗和性能效率方面的实际影响。本研究中提出的指导方针和观察结果提供了一种微架构层面的理解，旨在帮助开发者优化软件，从而更有效地利用硬件，实现AI和HPC工作负载的更高效部署。

引用文献汇总

• [7] Z. Jia, M. Maggioni, B. Staiger, and D. P. Scarpazza, “Dissecting the NVIDIA volta GPU architecture via microbenchmarking,” 2018.

  • 引用位置: 方法细节 - INT和FP32执行单元
  • 引用描述: 该文献被引用以说明在Volta和Ampere等早期架构中，INT32和FP32指令通过独立的执行管线发布，这在某些情况下会导致利用率不佳。

• [26] NVIDIA Corporation, Parallel Thread Execution (PTX) ISA, Release 8.8, 2025.

  • 引用位置: 方法细节 - 时钟开销；方法细节 - 可变MMA和基于Tile的指令
  • 引用描述: 第一次引用是为了说明%clock64寄存器的定义和功能，它是一个预定义的只读特殊寄存器，用于读取时钟周期计数值。第二次引用是为了指出在MMA指令的测试中，E8M0数据类型因其仅用于块内指数缩放的特定用途而未被测试。

• [27] T. T. Dao, J. Kim, S. Seo, B. Egger, and J. Lee, “A performance model for gpus with caches,” 2015.

  • 引用位置: 方法细节 - INT和FP32执行单元
  • 引用描述: 该文献被引用以支持排除GB203首次运行时出现的较高延迟数据的做法，因为其他研究也曾因缓存预热前的缓存未命中而排除了类似结果。

• [28] NVIDIA Corporation, CUDA Binary Utilities - Instruction Set Reference, 2025.

  • 引用位置: 方法细节 - 可变MMA和基于Tile的指令
  • 引用描述: 引用该文档以说明在Blackwell架构中，PTX mma指令到SASS指令的转换规则，即QMMA用于FP8矩阵乘累加，OMMA用于FP4。