（原标题：NVIDIA Tensor Core 的演变：从 Volta 到 Blackwell）
公众号记得加星标??，第一时间看推送不会错过。
来源：内容编译自semianalysis。
在人工智能和深度学习领域，GPU 计算能力的提升速度远超摩尔定律，年复一年地持续实现着“黄氏定律”般显著的性能提升。推动这一进步的核心技术正是 Tensor Core。
尽管 Tensor Core 无疑是现代人工智能和机器学习的基石，但即使是许多经验丰富的从业者，对其也仍未有深入的理解。GPU 架构以及基于该架构的编程模型的快速发展，使得机器学习研究人员和科学家越来越难以跟上 Tensor Core 的最新变化并理解这些变化的影响。
在本问中，我们将介绍主流数据中心 GPU 的核心特性，首先解释性能工程的重要基本原理。然后，我们将追溯 Nvidia Tensor Core 架构和编程模型的演变，并重点阐述其演变背后的动机。我们的最终目标是提供资源，帮助理解 Nvidia 的 GPU 架构，并直观地了解其架构的演变。只有在解释完每个架构之后，我们才能解释 Blackwell 张量核心及其全新内存层次结构的精妙之处。
需要强调的是，扎实的计算机架构理解能力是理解本文诸多讲解和讨论的先决条件。本文将简要介绍 CUDA 编程，供读者复习，而非解释 GPU 架构的基本概念。相反，我们将以 Tensor Core 的前沿知识为基础，通过详尽的讲解，将目前较为零散的知识转化为易于理解的结构化见解，从而拓展读者对这项前沿技术的理解。
性能优先原则
阿姆达尔定律（Amdahl’s Law）
对于固定的问题规模，阿姆达尔定律规定了通过增加计算资源进行并行化可以获得的最大加速比。具体而言，扩展计算资源只会缩短并行部分的执行时间，因此性能提升受限于串行部分的执行时间。为了量化这一点，最大性能提升如下：
其中 S 是并行工作的执行时间，p 是可并行化工作的加速比。在并行部分完全并行化的理想情况下，加速比 p 可以是处理单元的数量。
强扩展和弱扩展
（Strong and Weak Scaling）
强扩展和弱扩展描述了针对不同问题设置扩展计算资源的性能提升。强扩展是指扩展计算资源以解决固定规模的问题，而阿姆达尔定律则量化了强扩展的加速比。另一方面，弱扩展是指扩展计算资源以在恒定时间内解决更大的问题。例如，使用 4 倍的计算资源，在相同的时间内处理 4 倍大的图像。
强扩展和弱扩展在不同规模的问题上意味着不同的性能提升。强扩展可以为所有规模的问题提供加速，而弱扩展仅在使用更多计算来解决更大问题时才保证性能提升。
数据移动是大忌
数据移动是一种罪过，因为就运行时间和扩展性而言，计算成本低廉，而数据移动成本高昂。数据移动速度从根本上来说更慢，因为现代 DRAM 单元的运行速度为数十纳秒，而晶体管的开关速度为亚纳秒。就扩展性而言，虽然自 2000 年代以来计算速度的提升有所放缓，但内存速度的提升也更慢，从而形成了“内存墙”。
Tensor Core 架构演进
张量核心生成概述
在本节中，我们将介绍使用 Tensor Core 的主要 Nvidia GPU 架构，即 Tesla V100 GPU、A100 Tensor Core GPU、H100 Tensor Core GPU 以及 Blackwell GPU。我们还添加了一个 Tensor Core 之前的章节，作为 CUDA 编程模型的复习。我们将简要介绍与理解 Tensor Core 相关的主要特性和变化，并将详细信息保留到其他来源，我们会在每个小节中提供链接。
预张量核心
一、PTX编程模型
并行线程执行 (PTX：Parallel Thread Execution) 是跨 GPU 代抽象的虚拟指令集。PTX 程序描述了一个核函数，该函数由大量 GPU 线程执行，这些线程在 GPU 的硬件执行单元（即 CUDA 核心）上执行。线程被组织为网格，每个网格由协作线程阵列 ( CTA )组成。PTX 线程可以访问来自多个状态空间的数据，这些状态空间是具有不同特性的内存存储区域。具体而言，线程具有每个线程的寄存器，CTA 内的线程具有共享内存，并且所有线程都可以访问全局内存。
二、PTX 机器模型
GPU 架构围绕流多处理器 ( SM )阵列构建。SM 由标量处理核心、多线程指令单元和片上共享内存组成。SM 将每个线程映射到一个标量处理核心（也称为 CUDA 核心），而多线程指令单元则以 32 个并行线程组（称为Warp）的形式管理线程。
在指令发出时，指令单元选择一个 Warp，并向 Warp 中的线程发出指令。这种执行方法称为单指令多线程 ( SIMT )。与单指令多数据 ( SIMD ) 类似，SIMT 使用一条指令控制多个处理单元，但与 SIMD 不同的是，SIMT 指定的是单线程行为，而不是向量宽度。
三、流式汇编器
流式汇编器 (SASS：Streaming Assembler ) 是 PTX 虚拟化所基于的特定于架构的指令集。有关更多信息，请参阅CUDA 二进制实用程序文档。遗憾的是，由于 NVIDIA 向竞争对手隐藏了其架构 ISA 的细节，SASS 的文档并不完善。
Volta
NVIDIA 为何添加 Tensor Core
随着深度学习变得越来越突出，业界注意到 ML 工作负载需要硬件加速。2015 年初，Google 部署了 TPUv1 来加速其内部 ML 工作负载，2017 年，Nvidia 推出了用于矩阵数学的专用硬件。虽然 GPU 由于其简单的硬件流水线在发出指令时会消耗少量能量（~30pJ），但简单的浮点运算消耗的HFMA 能量甚至更少，仅为 1.5pJ。这使得指令所需的功耗是浮点运算本身的 20 倍。因此，执行大量浮点运算进行矩阵乘法是功耗低的。为了摊销指令开销，我们需要使用每个指令可以执行更多计算的复杂指令。为此，Nvidia 设计了半精度矩阵乘法和累加（HMMA）指令，这是一条执行半精度矩阵乘法的专用指令。执行该指令的相应专用硬件是 Tensor Core，它于 2017 年在 Volta 架构的 Tesla V100 GPU 中推出。Volta 张量核心是在 Volta 架构开发的后期添加的，仅在流片前几个月，这证明了 Nvidia 对其架构的调整速度有多快。
MMA 指导概述
给定一个矩阵，乘法和累加 (MMA) 指令计算 D = A * B + C：
A 是 M×K 矩阵
B 是 K×N 矩阵
C 和 D 是 M×N 矩阵
我们将矩阵形状表示为mMnNkK或 MxNxK。
为了执行完整的计算，我们首先将矩阵 A、B 和 C 从共享内存加载到线程寄存器，以便每个线程保存矩阵的片段。其次，我们执行 MMA 指令，该指令从线程寄存器读取矩阵，在 Tensor Core 上执行计算，并将结果存储到线程寄存器。最后，我们将结果从线程寄存器存储回共享内存。完整的计算由多个线程共同执行，这意味着每个步骤都需要协作线程之间的同步。
第一代 Tensor Core – Warp-scoped MMA
Tesla V100 GPU 的 SM 包含 8 个 Tensor Core，分为两部分。每个 Tensor Core 每周期能够计算相当于 4x4x4 矩阵乘法的运算，相当于每个 SM 每周期 1024 FLOP。
NVIDIA 设计了 PTX 指令 mma，以针对较低级别的HMMA指令。在 Volta 架构上，MMA 指令执行 8x8x4 矩阵乘法，由 8 个线程组成的四对 (quadpair) 通过共同保存输入和输出矩阵参与运算。其中，T0 表示线程 0，T0, T1, T2, T3 和 T16, T17, T18, T19 表示线程组，这两个线程组组成一个四对 (quadpair)。
在数据类型方面，Volta Tensor Core 支持 FP16 输入和 FP32 累积，这与 NVIDIA 的混合精度训练技术相呼应。该技术表明，可以在不损失模型精度的情况下以较低的精度训练模型。
Turing
Turing 架构包含第二代 Tensor Core，这是 Volta Tensor Core 的增强版，增加了对 INT8 和 INT4 精度的支持。Turing Tensor Core 支持全新的 Warp-Level 同步 MMA，我们将在下一节中讨论。Turing Tensor Core 还支持深度学习超级采样 (DLSS)，标志着 NVIDIA 开始将深度学习应用于游戏图形。
Ampere
异步数据复制
NVIDIA 在 Ampere 架构中引入了异步数据复制，这是一种以异步方式将数据直接从全局内存复制到共享内存的方法。要在 Volta 架构上将数据从全局内存加载到共享内存，线程必须先将数据从全局内存加载到寄存器，然后再将其存储到共享内存。然而，MMA 指令的寄存器使用率很高，并且必须与数据加载操作共享寄存器文件，这会导致寄存器压力过大，并浪费内存带宽来复制数据进出 RF。
异步数据复制通过从全局内存 (DRAM) 获取数据并将其直接存储到共享内存（可选 L1 访问）来缓解此问题，从而释放更多寄存器用于 MMA 指令。数据加载和计算可以异步进行，这从编程模型的角度来看更加困难，但可以提高性能。
此功能通过 PTX 指令线程级异步复制 cp.async 实现（文档）。对应的 SASS 是 LDGSTS，即异步全局到共享内存复制。具体的同步方法是异步组和基于 mbarrier 的完成机制。
第三代 Tensor Core – 曲速级同步
MMA（Warp-level Synchronous MMA）
Ampere 每个 SM 有 4 个 Tensor Core，每个 Tensor Core 每周期能够执行 512 FLOP，每个 SM 每周期总计 2048 Dense FLOP，性能是 Volta 的两倍。
Volta 需要 8 个线程组成的四对才能参与 MMA 运算，而 Ampere 则需要 32 个线程的完整 Warp。采用 Warp 宽度的 MMA 指令简化了线程布局，并降低了 Ampere 的 RF 压力。例如，以下是 16x8x16 形状的混合精度浮点的线程和数据布局：
NVIDIAldmatrix在 Ampere 中引入了增强型矢量化加载操作。与类似mma，ldmatrix它也是 Warp 范围的，这意味着一组 Warp 线程共同加载一个矩阵。与发出多个加载指令相比，这减少了地址生成寄存器的使用，从而降低了寄存器压力。有关更多信息，请参阅CUDA 文档。
ldmatrix将数据加载到寄存器中，其布局与 Tensor Core 的数据布局相匹配。与 Volta 的交错模式相比（参见“Tensor Core 编程：使用 CUTLASS 实现原生 Tensor Core”），更简单的线程和数据布局极大地提升了编程的人体工程学。观看 GTC 演讲“开发 CUDA 内核以在 NVIDIA A100 上将 Tensor Core 推向极限”。
Ampere MMA 采用 Brain 浮点格式 (BF16)，该格式已成为半精度数据类型的事实标准。BF16 提供与 FP32 相同的 8 位指数范围，但尾数为 7 位，从而能够以一半的存储成本实现 FP32 级别的动态范围。BF16 还消除了混合精度训练中损失缩放的需要。
Hopper
线程块集群
随着 SM 数量的增加，单个 SM 与整个 GPU 之间的大小差异也随之增大。为了在 CTA（映射到 SM）和网格（映射到整个 GPU）之间提供更精细的控制粒度，NVIDIA 在 Hopper 上添加了一个新的线程层次结构级别——线程块集群，它映射到物理上位于同一图形处理集群 (GPC) 中的一组 SM。线程块集群也称为协作网格阵列 (CGA)，在 CUDA 文档中简称为集群。
线程块集群中的 CTA 保证在同一 GPC 内的各个 SM 上协同调度，并且默认每个 SM 分配一个 CTA。这些 SM 的共享内存分区构成分布式共享内存 (DSMEM)。线程可以通过专用的 SM 到 SM 网络（无需经过 L2 缓存）以低延迟访问其他 SM 的共享内存。通过将 GPC 硬件执行单元暴露给编程模型，程序员可以减少数据移动并提高数据局部性。
张量记忆加速器
为了提高数据获取效率，NVIDIA 为每个 Hopper SM 添加了张量内存加速器 (TMA)。TMA 是一个专用硬件单元，可加速全局内存和共享内存之间的大量异步数据传输（批量异步复制）。
CTA 中的单个线程可以启动 TMA 复制操作。TMA 释放线程来执行其他独立工作，处理地址生成并提供额外的优势，例如越界处理。在 PTX 中，相应的指令是cp.async.bulk，详情请参阅CUDA 文档中的 部分。
然而，对于小型请求，由于地址生成开销，TMA 加载的延迟比常规异步数据复制更高。因此，NVIDIA 建议程序员使用 TMA 进行大型数据复制，以分摊开销。例如，在 LLM 推理中，TMA 不适用于以小块加载键值缓存的工作负载，但当每个块是 16 字节的倍数时，TMA 效果良好。
TMA 还支持一种称为多播的数据加载模式，该模式将数据从全局内存加载到线程块集群中多个 SM 的共享内存中，由多播掩码指定。多播加载不是发出多个全局内存加载请求，将同一段数据加载到多个 SM 中，而是一次性完成。具体来说，线程块集群中的多个 CTA 将部分数据加载到其对应的 SMEM 中，并通过 DSMEM 共享数据。这减少了二级缓存流量，进而减少了 HBM 流量。
第四代 Tensor Core – Warpgroup
级异步 MMA
NVIDIA 通过 Hopper 引入了一种新型的 MMA，即 Warpgroup 级 MMA（wgmma）。wgmma它是 Warpgroup 级的，这意味着由 4 个 Warp 组成的 Warpgroup 共同执行 MMA 操作。wgmma支持更广泛的形状。例如，混合精度 MMA 支持m64nNk16，其中 N 可以是 8 的倍数，范围从 8 到 256。wgmma.mma_async降低到一组新的 SASS：GMMA。在另一个示例中，半精度wgmma指令降低到HGMMA。
虽然 Warpgroup 中的所有线程都会将输出矩阵保存在其寄存器中，但 Hopper Tensor Core 可以直接从共享内存（而非寄存器）加载操作数，从而节省寄存器空间和带宽。具体来说，操作数矩阵 A 可以驻留在寄存器或共享内存中，而操作数矩阵 B 只能通过共享内存访问。
在wgmma数据类型方面，Hopper 引入了 8 位浮点数据类型（E4M3 和 E5M2），并采用 FP32 累加。实际上，累加路径采用 22 位定点格式（13 位尾数加符号位和指数位），与真正的 32 位累加相比，动态范围有所限制。由于张量核心精度降低，为了避免影响训练精度，每次 N_c 次累加都必须在 CUDA 核心中进行。。这种降低精度的累加提高了效率，但却以牺牲精度为代价。
Blackwell
巨大的寄存器压力并没有让 Hopper 放松，这促使Tensor Memory (TMEM)应运而生，这是一种专门用于 Tensor Core 运算的新型内存。在每个 SM 上，TMEM 拥有 128 行（通道）和 512 列 4 字节单元，总计 256 KB，这也是 SM 上寄存器文件的大小。
TMEM 的内存访问模式受到限制。具体来说，访问整个 TMEM 需要一个 WarpGroup，并且 WarpGroup 中的每个 Warp 只能访问一组特定的通道。通过限制内存访问模式，硬件设计人员可以减少访问端口的数量，从而节省芯片空间。另一方面，这种设计也意味着结语操作也需要一个 WarpGroup 来运行。与共享内存不同，程序员必须明确管理 TMEM，包括分配、释放以及将数据复制到 TMEM 和从 TMEM 复制数据。
CTA Pair
如果线程块集群中的两个 CTA 在其线程块集群中的 CTA 排序最后一位不同（例如 0 和 1、4 和 5），则它们会形成一个CTA 对。一个 CTA 对映射到一个纹理处理集群 (TPC)，该集群由两个 SM 组成，并与其他 TPC 组合形成一个 GPC。当 Blackwell Tensor Core 运算以 CTA 对粒度执行时，这两个 CTA 能够共享输入操作数。这种共享可以降低 SMEM 的容量和带宽需求。
Tensor Core 第五代 MMA
Tensor Core 第五代 MMA 指令（tcgen05.mma在 PTX 中）已完全不再使用寄存器来保存矩阵。操作数现在驻留在共享内存和 Tensor Memory 中。
具体来说，假设 MMA 计算 D = A * B + D：不使用线程寄存器可以消除复杂的数据布局，并释放线程寄存器空间用于其他工作，例如尾声操作。与wgmma使用 Warpgroup 启动 MMA 操作不同，tcgen05.mma它具有单线程语义，这意味着单个线程可以启动 MMA 操作。这消除了 Warp 发出 MMA 操作的作用。
一个值得注意的 MMA 变体是 MMA.2SM，它使用两个 SM 共同执行 MMA 操作。
MMA.2SM 以 CTA 对级粒度执行，并且由于其tcgen05.mma具有单线程语义，因此 CTA 对中领导者 CTA 中的单个线程将启动 MMA.2SM。这里我们展示了数据路径组织布局 A。布局 A 显示，与 1SM 版本（布局 D）相比，MMA.2SM 将 M 维度增加了一倍，因此这两个 SM 加载不同的矩阵 A 和 D 分块。此外，MMA.2SM 拆分了矩阵 B，将加载的数据量减半。
矩阵 B 在两个 SM 之间共享，这意味着 B0 和 B1 块需要通过 DSMEM 进行通信。虽然 DSMEM 和 SMEM 之间存在带宽差异，但由于我们加载的是较小的块，因此对协调的影响很小。即便如此，我们怀疑在 Blackwell 上，TPC 中 SM 之间的通信带宽高于 DSMEM，因此 MMA.2SM 会利用这一点来实现更好的性能。
第五代 Tensor Core 除了可以执行常规矩阵乘法之外，还可以执行卷积。tcgen05.mma它支持权重平稳模式，并带有一个收集器缓冲区，用于缓存矩阵 B 以供复用。更多信息，请参阅CUDA 文档和相应的权重平稳 MMA 指令。
在支持的数据类型方面，Blackwell 支持微缩放浮点格式 (MXFP)，包括 MXFP8、MXFP6 和 MXFP4。详情请参阅本文。Blackwell 还支持 NVIDIA 自己的 NVFP4 格式，该格式以比 MXFP4 更精确而闻名。这可能是因为 NVFP4 的块大小更小、缩放因子数据格式不同以及采用了两级量化方法。
对于 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 倍的用于加载权重的内存带宽。
遗憾的是，2:4 结构化稀疏 GEMM 内核与 Hopper 上的密集 GEMM 内核相比，无法达到接近 2 倍的加速。这是因为在保持模型准确率的同时进行结构化剪枝存在困难、cuSPARSELt 内核未经优化以及 TDP 限制。除了中国人工智能实验室和少数西方实验性研究 论文外，大多数人工智能实验室在生产推理中忽略了 2:4 结构化稀疏性，而专注于量化和提炼。Meta 正在 Llama 上进行实验，但在很多情况下，这也是一条死路。
此外，目前缺乏使用 2:4 FP8 结构化稀疏性或 4:8 FP4 结构化稀疏性实现性能提升，同时保持零精度损失的封闭式或开放式模型，而且用于结构化剪枝的资源普遍匮乏。我们建议 NVIDIA 停止在主题演讲和市场推广材料中提及Jensen 数学结构化稀疏性失败案例，除非他们开始持续展示能够利用结构化剪枝进行推理的 SOTA 开放模型。一个好的开端是尝试在 DeepSeek 上进行结构化稀疏性测试，并证明其性能可以叠加到其他技术之上，例如 NVFP4 的蒸馏和量化。
NVIDIA 在第五代 Tensor Core 中，为 NVFP4 数据类型引入了成对 4:8 结构化稀疏性。在此方案中，每八个元素被分组为四对连续元素，其中恰好两对必须包含非零值，而其余两对则被修剪为零。由于 NVFP4 是子字节数据类型，我们认为这一约束促使 NVIDIA 采用了成对 4:8 模式。虽然 4:8 稀疏性可能看起来比之前的 2:4 模式更宽松，但增加的成对要求意味着，对于寻求在修剪的同时保持模型精度的机器学习工程师来说，在实践中，它并不是一个更宽松的约束。
张量核心尺寸增加
历代以来，NVIDIA 扩展 Tensor Core 大小的步伐远大于增加 Tensor Core 数量。NVIDIA 选择扩展张量核心大小而非核心数量，是因为它更符合矩阵乘法的性能特征。具体而言，当扩展问题规模时，矩阵乘法计算量呈立方增长，而数据移动量呈平方增长，这意味着运算强度呈线性增长。O(n) 的运算强度，加上数据移动量比计算量更昂贵的事实，刺激了张量核心大小的增加。
然而，增加核心大小和核心数量都会以牺牲量化效应为代价。具体来说，核心数量过多会受到“图块量化效应”的影响，而核心大小过大则会导致“波量化效应”。当工作单元数量不能被工作器数量完全整除时，就会出现“波量化效应”，导致在处理最终较小批量工作时利用率下降。增加张量核心大小本质上就是增加工作单元大小，导致小型矩阵的利用率低下。
运算强度的线性增长也推动了 MMA 形状的增加。更大的 MMA 形状可以增强操作数共享粒度。具体来说，启动更少的较大块可以提高数据重用率，从而节省内存占用以及 RF 和 SMEM 的带宽。对于 Blackwell 之前的架构，这导致执行 MMA 操作所需的线程数量不断增加，从 8 个线程的四对 (Volta) 到 32 个线程的 Warp (Ampere)，再到 128 个线程的 Warpgroup (Hopper)。
增加内存大小
共享内存几乎每一代都在增加，而寄存器文件的大小却保持不变。这是因为 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，我们可以利用这种设计来提升 TMEM 的能效。例如，在一个朴素的分块矩阵乘法中，为了计算一个分块，矩阵 D 的分块会被访问 2Kt 次（Kt 次读取和 Kt 次写入。Kt：K 维度上的分块数量），而矩阵 A 的分块和矩阵 B 的分块仅会被访问一次。
MMA教学的异步性
其中的“H”UTCHMMA,HGMMA,HMMA 代表半精度，因为它是16位格式；“Q”代表QGMMA,UTCQMMA四分之一精度（8位），因为8位是全精度（32位）的四分之一。“O”代表“八进制”，即FP4中32位的八分之一UTCOMMA。
MMA 指令看似从同步跳到了异步。实际上，由于需要重叠，MMA 指令在 SASS 级别逐渐变为异步LDSM指示。
在 SASS 级别，MMA 操作涉及执行一条LDSM指令，将矩阵块从共享内存加载到寄存器文件，然后执行两条HMMA执行 MMA 指令。在执行过程中，两个HMMA指令异步发出，并通过硬件互锁阻止寄存器的使用。由于硬件互锁不允许重叠的 LDSM 指令，因此顺序执行一条LDSM或两条HMMA指令在指令发布流水线中会产生一个小气泡。然而，随着 Tensor Core 的速度越来越快，这个气泡会造成不可忽略的性能损失，因此需要为 MMA 引入异步完成机制。
Hopper 支持异步完成机制提交和隔离wgmma。HGMMA发出指令时，没有硬件互锁来保护寄存器的使用。相反，编译器会调度LDSM下一个 MMA 指令，并使用FENCE保持下一个HGMMA等待的指令。使用 Blackwell，MMA 操作是完全异步的。加载到 Tensor Memory（tcgen05.ld / tcgen05.st / tcgen05.cp）的指令全部是显式异步的。
数据类型精度降低
每一代 NVIDIA Tensor Core 中，NVIDIA 都会不断添加低精度数据类型，从 16 位到 4 位。这是因为深度学习工作负载对低精度的容忍度极高。对于推理尤其如此，推理可以使用比训练更低的精度。低精度更节能，占用更少的硅片面积，并实现更高的计算吞吐量。在新一代 Tensor Core 中，我们还看到 NVIDIA 移除了 FP64 支持，以便在硅片面积和功耗预算允许的情况下优先处理低精度数据类型。
有趣的是，优先级排序也影响了整数数据类型的支持。自 Hopper 以来，INT4 数据类型已被弃用，而在 Blackwell Ultra 上，我们发现 INT8 计算吞吐量有所下降。这是由于低精度整数数据类型的普及延迟造成的。尽管 Turing 支持 INT8 和 INT4，但直到 4 年后，新的推理量化方法才能够利用 INT4 的紧凑性来服务于 LLM。那时，NVIDIA 已经在 Hopper 上弃用了 INT4 wgmma。
https://semianalysis.com/2025/06/23/nvidia-tensor-core-evolution-from-volta-to-blackwell/#ampere
*免责声明：本文由作者原创。文章内容系作者个人观点，半导体行业观察转载仅为了传达一种不同的观点，不代表半导体行业观察对该观点赞同或支持，如果有任何异议，欢迎联系半导体行业观察。
今天是《半导体行业观察》为您分享的第4074期内容，欢迎关注。
加星标??第一时间看推送，小号防走丢
求推荐
股票配资在线提示：文章来自网络，不代表本站观点。
