NVIDIA Tensor Core 的演变：从 Volta 到 Blackwell

https://semianalysis.com/2025/06/23/nvidia-tensor-core-evolution-from-volta-to-blackwell/

在我们去年 AI Scaling Laws article from late last year中，我们探讨了多层 AI 扩展定律如何持续推动 AI 行业向前发展，使得模型能力的增长速度超过了摩尔定律，并且单位 token 成本也相应地迅速降低。这些扩展定律由训练和推理的优化与创新推动，但计算能力的提升超越摩尔定律也在其中发挥了关键作用。

在 AI 扩展定律文章中，我们重新审视了围绕计算扩展的长期争论，回顾了 2000 年代末 Dennard 缩放定律的终结，以及 2010 年代末经典摩尔定律每晶体管成本下降速度的结束。尽管如此，计算能力仍在以快速的步伐不断改进，接力棒传递给了诸如先进封装、3D 堆叠、新型晶体管和专用架构（如 GPU）等其他技术。

来源：Nvidia

在 AI 和深度学习领域，GPU 的计算能力提升速度超过了摩尔定律，每年都能持续提供令人惊叹的“黄氏定律”性能改进。推动这一改进的核心技术就是 Tensor Core。

尽管 Tensor Core 毫无疑问是现代 AI 和机器学习基础的基石，但即使是该领域的许多经验丰富的从业者，对其也了解甚少。GPU 架构和运行在此架构上的编程模型的快速演变，使得机器学习研究人员和科学家越来越难以跟上 Tensor Core 的最新变化，并理解这些变化的影响。

来源：SemiAnalysis, HC2023-K2: Hardware for Deep Learning

在本报告中，我们将介绍主要数据中心 GPU 的核心特性，首先解释性能工程的重要基本原理。然后，我们将追溯 NVIDIA Tensor Core 架构和编程模型的演变，突出其背后的动机。我们的最终目标是提供一个理解 NVIDIA GPU 架构的资源，并对其架构演变提供直观的见解。只有在解释了每种架构之后，我们才能解释 Blackwell Tensor Core 的美妙之处以及其新的内存层次结构。

重要的是，我们需要说明，掌握计算机架构是能够理解本文中许多解释和讨论的先决条件，本文将提供一个关于 CUDA 编程的简短部分作为复习，而不是解释 GPU 架构的基础概念。相反，我们将在 Tensor Core 知识的前沿基础上进行拓展，通过详细解释，将目前仅在小范围内流传的知识转化为易于获取、有条理的见解。

正如大学会开设 101 课程以及 4000 级课程一样，SemiAnalysis 的不同文章将迎合不同层次的理解能力，以及不同职业和专业领域的读者。

我们感谢以下合作者：

• Jay Shah，Colfax Research：提供了出色的 CUTLASS 教程，并在多次会议中仔细检查技术细节

• Ben Spector，斯坦福 Hazy Research：提供了关于编程模型变化的深刻见解以及写作建议

• Tri Dao，普林斯顿和 Together AI：审阅了草稿并给出了详细反馈

• Neil Movva，Together AI：审阅了草稿并提供了关于 GPU 内核编写的见解

• Charles Frye，Modal：提供了教育性的 GPU 术语表并对草稿进行了总体审阅

• Simon Guo，斯坦福 Hazy Research：绘制了封面图像

• NVIDIA：分享了 Tensor Core 设计的进展。团队包括：

• Ian Buck，CUDA 的发明者
• Jonah Alben，GPU 架构和工程负责人许多

• 其他 GPU 专家

性能基本原理

Amdahl 定律

对于固定问题规模，Amdahl 定律规定了通过增加计算资源进行并行化所能获得的最大加速比。具体而言，增加计算资源仅能减少并行部分的执行时间，因此性能改进受到串行部分的限制。为了量化这一点，最大性能改进为：

其中 (S) 是并行工作执行时间，(p) 是可并行化工作的加速比。在理想情况下，如果并行部分能够完美并行化，加速比 (p) 可以是处理单元的数量。

强扩展与弱扩展

强扩展和弱扩展描述了在不同问题设置下，增加计算资源所带来的性能改进。强扩展是指增加计算资源来解决固定规模的问题，Amdahl 定律量化了强扩展的加速比。另一方面，弱扩展是指增加计算资源来解决更大规模的问题，同时保持恒定的时间。例如，使用 4 倍的计算资源在相同时间内处理一个 4 倍大的图像。我们推荐阅读这篇博客文章以获取更详细的解释。

来源：SemiAnalysis, Performance and Scalability – SCENET Summer School

强扩展和弱扩展意味着不同问题规模下的性能改进有所不同。强扩展为所有问题规模提供加速，而弱扩展仅在我们使用更多计算资源解决更大问题时保证性能改进。

来源：SemiAnalysis

数据移动是原罪

数据移动之所以是“原罪”，是因为在运行时间和扩展性方面，计算是廉价的，而数据移动是昂贵的。数据移动本质上更慢，因为现代 DRAM 单元的操作时间以数十纳秒计，而晶体管的切换速度仅为亚纳秒级。在扩展性方面，尽管自 2000 年代以来计算速度的提升已经放缓，但内存速度的提升更慢，从而形成了“内存墙”。

Tensor Core 架构演变

Tensor Core 代际概述

在本节中，我们将介绍使用 Tensor Core 的主要 NVIDIA GPU 架构，即 Tesla V100 GPU、A100 Tensor Core GPU、H100 Tensor Core GPU 以及 Blackwell GPU。我们还增加了一个“Tensor Core 之前”的部分，以便复习 CUDA 编程模型。我们将简要介绍与理解 Tensor Core 相关的主要特性和变化，并将详细内容留给其他来源，我们在每个子部分中都链接了这些来源。

Tensor Core 之前

PTX 编程模型

并行线程执行（PTX）是一种虚拟指令集，它抽象了 GPU 代际之间的差异。PTX 程序描述了一个内核函数，该函数由大量 GPU 线程执行，这些线程在 GPU 的硬件执行单元（即 CUDA 核心）上运行。线程被组织成网格，而网格由协作线程数组（CTA）组成。PTX 线程可以访问多个状态空间的数据，这些状态空间是具有不同特性的内存存储区域。具体而言，线程拥有每个线程的寄存器，CTA 内的线程共享共享内存，所有线程都可以访问全局内存。更多信息请参阅 CUDA 文档的这一部分。

来源：SemiAnalysis

PTX 机器模型

GPU 架构围绕一系列流式多处理器（SM）构建。SM 由标量处理核心、多线程指令单元和片上共享内存组成。SM 将每个线程映射到标量处理核心（也称为 CUDA 核心），多线程指令单元以 32 个并行线程为一组的线程束来管理线程。

在指令发布时，指令单元选择一个线程束，并向该线程束的线程发布指令。这种执行方式称为单指令多线程（SIMT）。与单指令多数据（SIMD）类似，SIMT 使用单条指令控制多个处理单元，但与 SIMD 不同的是，SIMT 指定单个线程的行为，而不是向量宽度。更多信息请参阅 CUDA 文档的这一部分。

PTX 机器模型。来源：SemiAnalysis, PTX ISA 文档 - 图 4

流式汇编器

流式汇编器（SASS）是 PTX 虚拟化的特定架构指令集。更多信息请参阅 CUDA 二进制工具文档。遗憾的是，由于 NVIDIA 隐瞒其架构 ISA 细节以防止竞争对手获取信息，SASS 并未得到很好的记录。

Volta

NVIDIA 添加 Tensor Core 的原因

随着深度学习变得越来越重要，业界注意到机器学习工作负载需要硬件加速。2015 年初，谷歌设计了用于加速矩阵乘法的 TPUv1，大约在同一时期，Nvidia 也开始开发用于矩阵数学的专用硬件。尽管 GPU 在发布指令时消耗的能量较少（约 30pJ），因为它们的硬件流水线简单，但简单的浮点运算（如 HFMA）消耗的能量更少，仅为 1.5pJ。这使得指令与浮点运算本身的能耗比高达 20 倍。因此，执行大量浮点运算进行矩阵乘法是能源效率低下的。为了分摊指令开销，我们需要使用复杂的指令，这些指令可以在每条指令中执行更多的计算。为此，Nvidia 设计了 半精度矩阵乘法和累加（HMMA）指令，这是一种执行半精度矩阵乘法的专用指令。执行此指令的相应专用硬件是 Tensor Core，于 2017 年在 Volta 架构的 Tesla V100 GPU 中首次推出。Volta Tensor Core 在 Volta 架构的开发过程中非常晚才加入，仅在流片前几个月，这证明了 Nvidia 在架构上快速转向的能力。

来源：深度学习硬件趋势：专用指令分摊开销

MMA 指令概述

给定一个矩阵，乘法和累加（MMA）指令计算 (D = A \times B + C)：

我们将矩阵形状表示为 mMnNkK 或 。

为了执行完整的计算，我们首先将矩阵 (A)、(B) 和 (C) 从共享内存加载到线程寄存器中，以便每个线程都持有矩阵的片段。其次，我们执行 MMA 指令，该指令从线程寄存器读取矩阵，在 Tensor Core 上执行计算，并将结果存储回线程寄存器。最后，我们将结果从线程寄存器存储回共享内存。整个计算由多个线程共同完成，这意味着每个步骤都需要协作线程之间的同步。

来源：SemiAnalysis

第一代 Tensor Core - 线程束级 MMA

Tesla V100 GPU 的一个 SM 包含 8 个 Tensor Core，以两个一组的形式分组。每个 Tensor Core 每个时钟周期能够计算相当于 (4 \times 4 \times 4) 矩阵乘法的运算量，总共每个 SM 每个时钟周期可达到 1024 次浮点运算。

来源：Volta Tensor Core 培训

NVIDIA 设计了 PTX 指令 mma 来针对底层的 HMMA 指令。在 Volta 架构中，一个 MMA 指令执行 (8 \times 8 \times 4) 矩阵乘法，一个四对线程（共 8 个线程）参与操作，共同持有输入和输出矩阵。这里 T0 表示线程 0，[T0, T1, T2, T3] 和 [T16, T17, T18, T19] 是线程组，这两个线程组形成一个四对。

来源：SemiAnalysis。由 CUTLASS 可视化器生成

在数据类型方面，Volta Tensor Core 支持 FP16 输入和 FP32 累加，这与 NVIDIA 的混合精度训练技术相对应。这种技术表明，可以在不损失模型精度的情况下以较低精度训练模型。

要完全理解 MMA 的布局，请参考 Citadel 的微基准测试论文《通过微基准测试剖析 NVIDIA Volta GPU 架构》。要查看 Volta Tensor Core MMA 的交错布局模式，请阅读幻灯片《使用 CUTLASS 编程 Tensor Core：原生 Tensor Core》。最后，关于 Volta 架构的其他信息，请参考 NVIDIA Tesla V100 GPU 架构白皮书。

Turing

Turing 架构包括第二代 Tensor Core，这是 Volta Tensor Core 的增强版本，增加了对 INT8 和 INT4 精度的支持。Turing Tensor Core 支持一种新的线程束级同步 MMA，我们将在下一节中讨论。Turing Tensor Core 还启用了深度学习超级采样（DLSS），标志着 NVIDIA 开始将深度学习应用于游戏图形。感兴趣的读者可以参考 NVIDIA 的博客文章《深入解析 NVIDIA Turing 架构》以及 Turing 架构白皮书。

Ampere

异步数据拷贝

借助 Ampere，NVIDIA 引入了异步数据拷贝，这是一种以异步方式直接从全局内存拷贝数据到共享内存的方法。在 Volta 上，要将数据从全局内存加载到共享内存，线程必须首先将数据从全局内存加载到寄存器，然后再存储到共享内存。然而，MMA 指令占用大量寄存器，并且必须与数据加载操作共享寄存器文件，导致寄存器压力过高，并浪费了用于数据拷贝进出寄存器文件的内存带宽。

异步数据拷贝通过从全局内存（DRAM）获取数据并直接将其存储到共享内存（可选地访问 L1）来缓解这一问题，从而为 MMA 指令释放更多寄存器。数据加载和计算可以异步进行，这在编程模型方面更具挑战性，但从性能角度来看，能够解锁更高的性能。

此功能以 PTX 指令线程级异步拷贝 cp.async（文档）的形式实现。相应的 SASS 是 LDGSTS，即异步全局到共享内存拷贝。确切的同步方法是基于异步组和 mbarrier 的完成机制，详细信息请参阅此处。

来源：NVIDIA A100 Tensor Core GPU 架构白皮书

第三代 Tensor Core - 线程束级同步 MMA

Ampere 每个 SM 配备了 4 个 Tensor Core，每个 Tensor Core 每个时钟周期能够执行 512 次浮点运算，总共每个 SM 每个时钟周期可达到 2048 次密集浮点运算，是 Volta 的两倍。

虽然 Volta 需要一个四对 8 个线程参与 MMA 操作，但 Ampere 需要一个完整的线程束（32 个线程）。将 MMA 指令扩展到整个线程束，简化了 Ampere 的线程布局并降低了寄存器文件的压力。例如，以下是形状为 16x8x16 的混合精度浮点数的线程和数据布局：

来源：SemiAnalysis。由 CUTLASS 可视化器生成

NVIDIA 在 Ampere 中引入了 ldmatrix，这是一种增强型向量化加载操作。与 mma 类似，ldmatrix 是线程束级的，这意味着一个线程束的线程共同加载一个矩阵。与发布多个加载指令相比，这减少了地址生成寄存器的使用，从而降低了寄存器压力。更多信息请参阅 CUDA 文档。

ldmatrix 以与 Tensor Core 数据布局相匹配的方式将数据加载到寄存器中。与 Volta 的交错模式（见《使用 CUTLASS 编程 Tensor Core：原生 Tensor Core》）相比，更简单的线程和数据布局极大地提高了编程的便捷性。观看 GTC 演讲《开发 NVIDIA A100 上的 CUDA 内核，将 Tensor Core 推向绝对极限》以了解更多关于 Ampere 的内存加载如何与 Tensor Core 协调一致。

Ampere MMA 引入了脑浮点格式（BF16），它已成为半精度数据类型的事实标准。BF16 提供了与 FP32 相同的 8 位指数范围，但只有 7 位尾数，允许在存储成本减半的情况下达到 FP32 级别的动态范围。BF16 还消除了混合精度训练中对损失缩放的需求。

Hopper

线程块簇

随着 SM 数量的增加，SM 与整个 GPU 之间的尺寸差异变得越来越大。为了在 CTA（映射到 SM）和网格（映射到整个 GPU）之间提供更细粒度的控制，在 Hopper 上，NVIDIA 增加了一个新的线程层次结构级别，即线程块簇，它映射到物理上位于同一图形处理簇（GPC）中的一组 SM。线程块簇也称为协作网格数组（CGA），在 CUDA 文档中称为簇。

线程块簇中的 CTA 默认情况下保证在同一个 GPC 的 SM 上共同调度，并且每个 SM 分配一个 CTA。这些 SM 的共享内存分区形成了一个分布式共享内存（DSMEM）。线程可以通过专用的 SM 到 SM 网络（无需经过 L2 缓存）以低延迟访问另一个 SM 的共享内存。通过将 GPC 硬件执行单元暴露给编程模型，程序员可以减少数据移动并提高数据局部性。

来源：GTC 演讲《深入 NVIDIA Hopper 架构》

Tensor 内存加速器

为了提高数据获取效率，NVIDIA 在每个 Hopper SM 中添加了 Tensor 内存加速器（TMA）。TMA 是一个专用硬件单元，用于加速全局内存和共享内存之间大量数据的异步传输（批量异步拷贝）。

一个 CTA 中的单个线程可以发起一个 TMA 拷贝操作。TMA 释放了线程，使其可以执行其他独立工作，处理地址生成，并提供诸如越界处理等额外好处。在 PTX 中，相应的指令是 cp.async.bulk，详细信息请参阅 CUDA 文档的这一部分。

然而，对于小请求，TMA 加载的延迟高于常规异步数据拷贝，因为存在地址生成开销。因此，NVIDIA 建议程序员使用 TMA 进行大数据拷贝，以分摊开销。例如，在 LLM 推理中，TMA 不适合以小块加载 KV 缓存的工作负载，但如果每个块是 16 字节的倍数，则效果良好。关于这一点的更具体示例，请参阅 SGLang 前缀缓存、论文《FlashInfer》第 3.2.1 节、论文《Hardware-Efficient Attention for Fast Decoding》第 4.2 节以及 ThunderKittens MLA 解码。

TMA 还支持一种名为多播的数据加载模式，TMA 从全局内存将数据加载到线程块簇中多个 SM 的共享内存中，由一个多播掩码指定。与向多个 SM 发起多个全局内存加载以将相同的数据加载到多个 SM 中相比，多播只需一次加载即可完成。具体来说，线程块簇中的多个 CTA 将数据的一部分加载到它们各自的 SMEM 中，并通过 DSMEM 共享数据。这减少了 L2 缓存流量，进而减少了 HBM 流量。我们建议阅读 Jay Shah 的 TMA 教程以获取更多详细信息。

来源：SemiAnalysis, GTC 演讲《在 Hopper Tensor Core 上开发最佳 CUDA 内核》

第四代 Tensor Core - 战争组级异步 MMA

NVIDIA 在 Hopper 中引入了一种新的 MMA 类型，即战争组级 MMA（wgmma）。wgmma 是战争组级的，这意味着一个战争组（4 个线程束）共同执行一个 MMA 操作。wgmma 支持更广泛的形状范围。例如，混合精度 MMA 支持 m64nNk16，其中 N 可以是从 8 到 256 的 8 的倍数。wgmma.mma_async 降低到一组新的 SASS：GMMA。在另一个例子中，半精度 wgmma 指令降低到 HGMMA。有关 MMA 形状和数据类型的详细信息，请参阅 CUDA 文档的这一部分。

虽然战争组中的所有线程共同在寄存器中持有输出矩阵，但 Hopper Tensor Core 可以直接从共享内存加载操作数，而不是寄存器，从而节省寄存器空间和带宽。具体来说，操作数矩阵 A 可以位于寄存器或共享内存中，而操作数矩阵 B 只能通过共享内存访问。有关 wgmma 的完成机制、SMEM 布局等详细信息，请参阅 CUDA 文档的 wgmma 部分。

来源：SemiAnalysis

对于 wgmma 数据类型，Hopper 引入了 8 位浮点数据类型（E4M3 和 E5M2），并带有 FP32 累加。在实践中，累加路径被实现为 22 位定点格式（13 位尾数加上符号和指数位），与真正的 32 位累加相比，动态范围有限。由于 Tensor Core 精度降低，每 N_c 次累加必须在 CUDA 核心中进行，以防止限制训练精度（见 3.2.2）。这种降低精度的累加提高了效率，但以精度为代价。

有关 Hopper 架构的更多信息，请参阅以下内容：

• GTC 演讲：《深入 NVIDIA Hopper 架构》

• NVIDIA 博客文章概述：《深入解析 NVIDIA Hopper 架构》

• 白皮书：《NVIDIA H100 Tensor Core GPU 架构》

• 微基准测试：《基准测试和剖析 NVIDIA Hopper GPU 架构》

• 微基准测试：《通过微基准测试和多级分析剖析 NVIDIA Hopper 架构》

有关如何为 Hopper GPU 编程的示例，请参阅：

• GTC 演讲：《优化 Hopper 架构的应用程序》

• CUTLASS 演讲：《在 Hopper Tensor Core 上开发最佳 CUDA 内核》

• Colfax 博客文章：《CUTLASS 教程：在 NVIDIA Hopper GPU 上使用 WGMMA 进行快速矩阵乘法》

Blackwell

Tensor 内存

Hopper 上极端的寄存器压力促使了 Tensor 内存（TMEM） 的诞生，这是一种专为 Tensor Core 操作设计的新内存。在每个 SM 上，TMEM 拥有 128 行（通道）和 512 列 4 字节单元，总共 256 KB，这与 SM 上的寄存器文件大小相同。

TMEM 的内存访问模式受到限制。具体来说，需要一个战争组来访问整个 TMEM，而战争组中的每个线程束只能访问特定的一组通道。通过限制内存访问模式，硬件设计人员可以减少访问端口的数量，从而节省芯片空间。另一方面，这种设计也意味着尾声操作需要一个战争组来执行。与共享内存不同，程序员必须明确管理 TMEM，包括分配、释放以及数据的进出拷贝。

来源：GTC 演讲《使用 CUTLASS 编程 Blackwell Tensor Core》

CTA 对

如果线程块簇中的两个 CTA 的 CTA 排名在其线程块簇中相差最后一位（例如，0 和 1，4 和 5），则这两个 CTA 形成一个 CTA 对。CTA 对映射到一个纹理处理簇（TPC），TPC 由两个 SM 组成，并与其他 TPC 结合形成 GPC。当 Blackwell Tensor Core 操作以 CTA 对粒度执行时，这两个 CTA 能够共享输入操作数。这种共享减少了对 SMEM 容量和带宽的需求。

第五代 Tensor Core MMA

第五代 Tensor Core MMA 指令（在 PTX 中为 tcgen05.mma）完全摒弃了使用寄存器来存储矩阵。操作数现在位于共享内存和 Tensor 内存中。

具体来说，假设 MMA 计算 (D = A \times B + D)：不使用线程寄存器消除了复杂的数据布局，并为其他工作（如尾声操作）释放了线程寄存器空间。与使用战争组发起 MMA 操作的 wgmma 不同，tcgen05.mma 具有单线程语义，这意味着一个单线程发起 MMA 操作。这移除了线程束在发布 MMA 方面的角色。

来源：SemiAnalysis

一个值得注意的 MMA 变体是 MMA.2SM，它使用 2 个 SM 共同执行一个 MMA 操作。MMA.2SM 在 CTA 对粒度级别执行，并且由于 tcgen05.mma 具有单线程语义，CTA 对中的领导者 CTA 中的一个线程发起 MMA.2SM。这里我们展示了数据路径组织布局 A。布局 A 显示，与 1SM 版本（布局 D）相比，MMA.2SM 将 M 维度翻倍，因此两个 SM 加载不同的矩阵 A 和 D 瓦片。此外，MMA.2SM 将矩阵 B 分割，将加载的数据量减半。

来源：SemiAnalysis, GTC 演讲《使用 CUTLASS 编程 Blackwell Tensor Core》

矩阵 B 在两个 SM 之间共享，这意味着需要在 DSMEM 中通信 B0 和 B1 瓦片。尽管 DSMEM 和 SMEM 之间存在带宽差异，但由于我们加载的是较小的瓦片，因此对协调的影响最小。也就是说，我们怀疑在 Blackwell 上，TPC 中 SM 之间的通信带宽高于 DSMEM 的，因此 MMA.2SM 利用这一点来实现更好的性能。

第五代 Tensor Core 还可以执行卷积，而不仅仅是通用矩阵乘法。tcgen05.mma 支持权重静止模式，并带有收集器缓冲区，用于缓存矩阵 B 以供重用。更多信息，请参阅 CUDA 文档以及相应的权重静止 MMA 指令。

在支持的数据类型方面，Blackwell 支持微缩放浮点格式（MXFP），包括 MXFP8、MXFP6 和 MXFP4。详情请参阅此论文。Blackwell 还支持 NVIDIA 自己的 NVFP4 格式，该格式以其比 MXFP4 更高的精度而闻名。这可能是因为它的块大小较小、不同的缩放因子数据格式以及两级量化方法（见此 GitHub 问题）。详情请参阅此论文，了解数据格式比较。

在 Blackwell 上，由于 FP8 和 FP6 的理论吞吐量相同，我们相信它们在 Tensor Core 中共享物理电路。相比之下，CDNA4 的 FP6 吞吐量是 FP8 的两倍，因为它们的 FP6 单元与 FP4 共享数据路径。我们相信 UDNA 将改为让 FP6 单元与 FP8 共享。

附注：结构化稀疏性

Ampere 特点是 2:4 结构化稀疏性，理论上可以将 Tensor Core 吞吐量翻倍。它通过修剪权重矩阵实现，使得每 4 个元素中有 2 个为零。在这种格式中，矩阵通过移除零元素进行压缩，并且额外的元数据索引矩阵记录它们的位置，大致将内存使用量和带宽减半。

根据中国工程师破解的这篇微基准测试论文，Ampere 的结构化稀疏性可以在指令级别为大型形状 MMA 操作实现 2 倍加速。它还表明，在 Hopper 中，结构化稀疏性 wgmma 指令可以达到 2 倍加速，并在加载权重时节省高达 2 倍的内存带宽。

不幸的是，与 Hopper 上的密集对应物相比，2:4 结构化稀疏性 GEMMs 内核无法接近 2 倍加速。这是由于在保持模型精度的同时进行结构化修剪存在困难，cuSPARSELt 内核未经过优化，以及 TDP 限制。除了中国 AI 实验室和少数西方实验性研究论文外，大多数 AI 实验室在生产推理中忽略了 2:4 结构化稀疏性，转而关注量化和蒸馏。Meta 正在 Llama 中尝试使用它，但在许多情况下，这是一条死胡同。

此外，缺乏闭源或开源模型在保持零精度损失的同时，通过 2:4 FP8 结构化稀疏性或 4:8 FP4 结构化稀疏性显示出性能提升，并且普遍缺乏专门用于结构化修剪的资源。我们建议 NVIDIA 应该停止在主题演讲和营销材料中使用 Jensen 数学结构化稀疏性 flops，除非他们开始持续展示能够利用结构化修剪进行推理的 SOTA 开源模型。一个好的第一步是在 DeepSeek 上进行结构化稀疏性，并展示性能可以与其他技术（如蒸馏和量化）如 NVFP4 叠加。

来源：NVIDIA

在第五代 Tensor Core 中，NVIDIA 为 NVFP4 数据类型引入了成对的 4:8 结构化稀疏性。在这种模式下，每 8 个元素被分组为 4 个连续的对，其中恰好有 2 个对包含非零值，而其余的 2 个对被修剪为零。由于 NVFP4 是一种子字节数据类型，我们相信这一约束促使 NVIDIA 采用了成对的 4:8 模式。尽管 4:8 稀疏性看起来比之前的 2:4 模式更宽松，但由于增加了成对的要求，实际上对于在修剪时寻求保持模型精度的 ML 工程师来说，它并不是一个更宽松的约束。

来源：NVIDIA

Tensor Core 尺寸增大

来源：SemiAnalysis, NVIDIA

在各代产品中，NVIDIA 比增加 Tensor Core 数量更激进地增加了 Tensor Core 的尺寸。NVIDIA 选择增加 Tensor Core 尺寸而不是核心数量，是因为这更适合矩阵乘法的性能特性。具体来说，当扩展问题规模时，矩阵乘法计算量呈立方增长，但数据移动呈二次方增长，这意味着算术强度呈线性增长。O(n) 算术强度，加上数据移动比计算更昂贵的事实，促使了 Tensor Core 尺寸的增加。

来源：SemiAnalysis, NVIDIA

然而，无论是增加核心尺寸还是核心数量，都会以量化效应为代价。具体来说，拥有大量核心会受到瓦片量化效应的影响，而拥有较大核心尺寸会导致波浪量化效应。当工作单元数量不能被工作器数量整除时，会发生波浪量化效应，导致在处理最后一批较小的工作时利用率下降。增加 Tensor Core 尺寸本质上是增加工作单元尺寸，从而导致小矩阵的利用率较低。

来源：SemiAnalysis

算术强度的线性增长也促使了 MMA 形状的增加。拥有更大的 MMA 形状增强了操作数共享粒度。具体来说，启动较少的较大瓦片会增加数据重用，节省寄存器文件和共享内存的内存占用和带宽。对于 Blackwell 之前的架构，这导致了执行 MMA 操作的线程数量的增加，从 Volta 的四对 8 个线程，到 Ampere 的 32 个线程的线程束，再到 Hopper 的 128 个线程的战争组。

内存尺寸增大

来源：SemiAnalysis, NVIDIA

几乎每一代，共享内存都有所增加，而寄存器文件大小保持不变。原因是 Tensor Core 吞吐量的增加需要更深的暂存缓冲区。

由于 Tensor Core 消耗数据的速度远远快于全局内存加载的速度，我们使用暂存内存来缓冲数据，以便内存加载可以领先于 MMA 操作。Tensor Core 吞吐量每代翻倍，但全局内存加载延迟并未减少，实际上反而增加了。因此，我们需要增加暂存内存大小以缓冲更多数据。 为了实现这一点，NVIDIA 选择了共享内存作为 Tensor Core 的暂存内存，这解释了为什么共享内存增加了，但寄存器文件大小保持不变。

然而，Blackwell 的共享内存大小并没有从 Hopper 增加。这是因为 tcgen05 MMA 可以利用 2 个 SM，因此每个 SM 的共享内存只需要加载一半的操作数。因此，Blackwell 的共享内存大小实际上翻倍了。

NVIDIA 的暂存内存选择也解释了为什么操作数位置逐渐从寄存器转移到共享内存。也就是说，NVIDIA 在 Blackwell 上增加了 TMEM 以支持增加的 Tensor Core 吞吐量。由于 TMEM 更接近 Tensor Core，它可以更节能。此外，拥有单独的内存增加了总内存带宽，以饱和 Tensor Core。

在所有操作数中，矩阵 D 始终保留在 TMEM 中。我们可以利用这种设计的优势，因为矩阵 D 比矩阵 A 和 B 更频繁地被访问。例如，在计算一个瓦片的简单瓦片矩阵乘法中，矩阵 D 瓦片被访问 (2K_t) 次（(K_t) 次读取和 (K_t) 次写入。(K_t)：沿 K 维度的瓦片数量），而矩阵 A 瓦片和矩阵 B 瓦片仅被访问一次。

来源：SemiAnalysis, NVIDIA

MMA 指令的异步性

来源：SemiAnalysis, NVIDIA

UTCHMMA、HGMMA、HMMA 中的 “H” 代表半精度，因为它是 16 位格式，而 QGMMA、UTCQMMA 中的 “Q” 代表四分之一精度（8 位），因为 8 位是全精度（32 位）的四分之一。“O” 代表 “八进制”，意味着是 32 位的八分之一，因为 UTCOMMA 是 FP4。

MMA 指令似乎从同步跳转到了异步。实际上，由于需要重叠 LDSM 指令，MMA 指令在 SASS 级别逐渐变得异步。

在 SASS 级别，MMA 操作涉及执行一条 LDSM 指令，将矩阵瓦片从共享内存加载到寄存器文件中，然后执行两条 HMMA 指令来执行 MMA。在执行过程中，两条 HMMA 指令是异步发布的，并且通过硬件互锁阻塞寄存器使用。由于硬件互锁不允许重叠 LDSM 指令，因此依次执行一条 LDSM 和两条 HMMA 指令会在指令发布管道中产生一个小气泡。然而，Tensor Core 变得如此之快，以至于这个气泡会导致不可忽视的性能损失，这促使了 MMA 的异步完成机制。

Hopper 支持 wgmma 的异步完成机制提交和围栏。当发布 HGMMA 指令时，没有硬件互锁来保护寄存器使用。相反，编译器为下一个 MMA 调度 LDSM，并使用 FENCE 指令让下一个 HGMMA 等待。在 Blackwell 上，MMA 操作完全异步。加载到 Tensor 内存的指令（tcgen05.ld / tcgen05.st / tcgen05.cp）都是明确的异步指令。

来源：SemiAnalysis

数据类型精度降低

来源：SemiAnalysis, NVIDIA

在 NVIDIA Tensor Core 的每一代产品中，NVIDIA 继续增加低精度数据类型，从 16 位到 4 位。这是因为深度学习工作负载对低精度极为容忍。在推理中尤其如此，甚至可以使用比训练时更低的精度。低精度更节能，占用的硅片面积更小，并且能够实现更高的计算吞吐量。在新一代产品中，我们也看到 NVIDIA 为了优先考虑低精度数据类型，在硅片面积和功耗预算下，取消了对 FP64 的支持。

有趣的是，这种优先级也影响了整数数据类型的支持。自 Hopper 以来，INT4 数据类型已被弃用，在 Blackwell Ultra 上，我们看到 INT8 的计算吞吐量降低。这是由于低精度整数数据类型受欢迎的时间延迟。尽管 Turing 支持 INT8 和 INT4，但直到 4 年后，新的推理量化方法才开始利用 INT4 的紧凑性来为 LLM 提供服务。到那时，NVIDIA 已经在 Hopper wgmma 上弃用了 INT4。

以下内容略

编程模型的演化

强大的扩展性和单一 CTA 占用率

异步执行

文章来自于微信公众号“human five”。