🔢 mathematics

Accurate Models of NVIDIA Tensor Cores

本文针对 NVIDIA V100、A100、H100 和 B200 等数据中心 GPU 中不兼容 IEEE 754 标准的矩阵乘法单元，提出了能够精确模拟其在 8、16 及 19 位浮点格式下舍入行为、累加器宽度等数值特性的软件模型，以解决跨架构结果不可复现的问题。

原作者： Faizan A. Khattak, Mantas Mikaitis

发布于 2026-04-07

📖 1 分钟阅读🧠 深度阅读

CC BY 4.0

原作者： Faizan A. Khattak, Mantas Mikaitis

原始论文采用 CC BY 4.0 许可（http://creativecommons.org/licenses/by/4.0/）。 ✨ 这是对下方论文的AI生成解释。它不是由作者撰写或认可的。如需技术准确性，请参阅原始论文。阅读完整免责声明

这篇论文就像是一份**“超级计算器内部构造的终极解密指南”**。

为了让你轻松理解，我们可以把 NVIDIA 的 GPU（图形处理器）想象成一个拥有成千上万个“超级速算员”的巨型工厂。这些速算员专门负责做一种叫“矩阵乘法”的数学题，这是训练人工智能（AI）和进行科学计算的核心工作。

1. 核心问题：速算员们“不按常理出牌”

在这个工厂里，有一群特殊的速算员，被称为**"Tensor Cores"（张量核心）**。它们为了追求极致的速度，发明了一套自己的“速算规则”。

传统规则（IEEE 754）： 就像学校里的数学考试，要求每一步都必须精确到小数点后几位，哪怕为了精确牺牲一点速度。
Tensor Cores 规则： 它们像是“为了赶时间而抄近道”的速算员。为了快，它们在某些步骤上会偷偷“四舍五入”或者“截断”数字，甚至在不同代的速算员（不同型号的 GPU，如 V100, A100, H100, B200）之间，抄近道的方式还都不一样。

这就导致了一个大麻烦：
如果你用 A100 显卡算出一道题，再用 H100 显卡算同一道题，结果可能不一样。这对于科学家来说很头疼，因为科学实验要求结果必须能“复现”（即别人用同样的方法也能算出同样的结果）。而且，因为没人知道这些速算员具体是怎么“抄近道”的，大家只能猜，猜错了就会导致 AI 训练失败或科学计算出错。

2. 作者的解决方案：制造“数字双胞胎”

为了解决这个问题，作者（来自英国利兹大学的两位研究者）做了一件非常酷的事情：他们在电脑软件里，为每一代 NVIDIA 的 Tensor Core 造了一个“数字双胞胎”（软件模型）。

这个“数字双胞胎”不是普通的计算器，它是一个完美的模仿者：

它完全模仿真实硬件的“坏习惯”（比如它知道在什么情况下会截断小数，什么情况下会多保留一位）。
它甚至能模仿不同型号显卡之间微妙的差异。

他们是怎么做到的？（侦探破案法）

设计“陷阱”测试题： 他们设计了一些特殊的数学题（测试向量），就像给速算员设下的陷阱。比如，故意给一个极小的数字，看速算员是把它当成 0 处理，还是保留下来。
疯狂试错（随机测试）： 他们让软件模型和真实显卡同时做一千万道随机题目。
找茬与修正： 只要软件模型算出的结果和真实显卡哪怕有一点点不同（比如最后一位数字不一样），他们就立刻停下来，像侦探一样分析：“哦！原来这里多了一个进位位！”或者“原来这里多保留了一位小数！”。
循环迭代： 他们不断修正软件模型，直到软件模型和真实显卡算出的结果完全一模一样（比特级准确）。

3. 这个“数字双胞胎”有什么用？

作者把这个做好的软件模型打包成了一个MATLAB 工具箱，免费提供给全世界使用。它的用处就像给科学家发了一把“万能钥匙”：

提前预演（不用买新显卡）： 科学家在写新算法时，不需要真的去买最新的 B200 显卡。他们可以在电脑上用这个模型先跑一遍，看看结果对不对。如果模型显示结果有问题，那真实显卡上肯定也有问题。
解释“为什么结果不一样”： 如果两个实验室用不同显卡算出了不同结果，科学家可以用这个模型来“复盘”，看看是因为哪个显卡的“速算规则”不同导致的。
开发更聪明的算法： 既然知道了速算员会怎么“抄近道”，算法开发者就可以专门设计一些方法来抵消这些误差，从而用低精度的硬件算出高精度的结果（就像用粗糙的尺子量出精确的长度）。

4. 一个生动的比喻：做蛋糕

想象一下，你要做蛋糕（矩阵乘法）：

标准做法： 用精密的电子秤称面粉，精确到 0.01 克。
Tensor Cores 做法： 为了快，直接用勺子舀面粉。
- V100 显卡： 勺子有点大，舀多了就抖掉一点。
- A100 显卡： 勺子形状变了，抖掉的方式不一样。
- H100 显卡： 甚至有时候勺子会先抖两下再舀。

以前，大家不知道这些勺子的具体形状，做出来的蛋糕味道（计算结果）自然不一样。
这篇论文就是把这些勺子的形状、抖动方式全部画成了图纸（软件模型）。现在，任何烘焙师（科学家）在没买到新勺子之前，就可以拿着图纸在纸上模拟，预测蛋糕的味道，或者设计出一种“补偿抖落”的配方，确保无论用哪个型号的勺子，蛋糕味道都一样好。

总结

这篇论文的核心贡献就是：把 NVIDIA 黑盒子里的“速算黑魔法”变成了透明的、可预测的、可复制的软件模型。 这让科学家和工程师不再需要盲目猜测硬件的行为，能够更安心、更准确地利用这些强大的 AI 加速芯片。

以下是关于论文《Accurate Models of NVIDIA Tensor Cores》（NVIDIA Tensor Cores 的精确模型）的详细技术总结：

1. 研究背景与问题 (Problem)

核心问题：矩阵乘法是神经网络训练和推理的核心操作。NVIDIA 等厂商在 GPU 中引入了专用硬件（Tensor Cores）来加速矩阵乘法。然而，这些针对 AI 优化的矩阵乘法单元并不完全符合 IEEE 754 浮点算术标准。
具体挑战：
- 非标准化行为：不同厂商（甚至同一厂商的不同代际架构，如 Volta, Ampere, Hopper, Blackwell）在舍入模式、累加器宽度、归一化点、额外进位位（extra carry bits）以及中间对齐位（guard bits）等数值特性上存在差异。
- 不可复现性：由于缺乏统一标准，在不同 GPU 架构上运行相同的矩阵乘法指令（MMA）可能导致不同的数值结果，这对科学计算和需要跨平台一致性的应用构成了巨大挑战。
- 现有模型的不足：之前的研究虽然尝试通过测试向量（Test Vectors）来推断硬件特性，但这些模型往往不够精确，无法区分细微的硬件实现差异，且缺乏对随机输入空间的全面验证。
- 缺乏工具：混合精度算法开发者和数值分析师缺乏能够精确模拟特定 GPU Tensor Core 行为的软件工具，导致在将算法部署到真实硬件前难以进行准确的误差分析。

2. 方法论 (Methodology)

作者提出了一套迭代式的建模与验证方法，旨在构建比特级精确（Bit-accurate）的 MATLAB 软件模型。

广义数值特性测试 (GNFT - Generalised Numerical Feature Testing)：
- 利用精心设计的测试向量（包含特定常数、次正规数、特殊值等）来触发硬件的边界情况。
- 通过观察输出结果，推断硬件的数值特性，如：是否支持次正规数、归一化点的位置、舍入模式、块融合乘加（Block FMA）的大小（ $N_{FMA}$ ）以及额外对齐位（ $n_{eab}$ ）的数量。
输入空间搜索方法 (ISSM - Input Space Search Method)：
- 在 GNFT 获得初步模型后，通过大规模随机化测试来验证和细化模型。
- 测试规模：生成 $10^7$ 量级的随机输入向量（来自正态分布和均匀分布），覆盖正常数、次正规数、无穷大和 NaN 等特殊情况。
- 迭代修正：将软件模型的输出与真实 GPU 硬件的输出进行比特级比对。一旦发现不匹配（Mismatch），分析失败案例，提出假设并修正模型参数（如调整对齐位数或舍入逻辑），直到模型在所有测试用例上与硬件完全一致。
工具实现：
- 开发了 MATLAB Tensor Core v0.4.1 工具箱。
- 结合了定点算术与自定义浮点格式模拟器 CPFloat。
- 支持从 V100 (Volta) 到最新的 B200 (Blackwell) 等多种架构。

3. 关键贡献 (Key Contributions)

首次精确确定多种架构的数值特性：
- 对 9 款 NVIDIA GPU（V100, A100, A30, A2, L40S, RTX 1000 Ada, H100, H200, B200）的 31 种 Tensor Core 配置进行了详细分析。
- 揭示了以往未公开或未被精确量化的细节，例如：
  - V100：发现其有效对齐位数为 25 位（而非之前的 24 位），且部分积在累加前可能未归一化。
  - A100/H100/B200：精确测定了不同输入格式（fp16, bf16, tf19, fp8）下的 $N_{FMA}$ 大小、额外对齐位（ $n_{eab}$ ）数量（如 H100 的 fp16 输入有 2 个额外对齐位）以及次正规数累加时的指数截断行为（Exponent Capping）。
  - H100/B200 的 fp8 特殊行为：发现通过 mma.sync 指令处理 fp8 时，实际上是将输入转换为 fp16 并利用 fp16 Tensor Core 进行交错（Interleaved）计算，而非直接使用原生 fp8 核心。
开发并验证了高精度的 MATLAB 模型：
- 构建了可配置的 Tensor Core 模拟器，能够精确复现硬件的舍入、截断和对齐行为。
- 通过 $10^7$ 次随机测试，验证了模型与真实硬件在比特级上的一致性（Bit-identical）。
开源工具箱：
- 发布了包含预定义 GPU 模型（如 B200TC, A100TC）和通用自定义模型（Generic_BFMA_TC）的 MATLAB 工具箱，支持用户自定义 FMA 大小、对齐位和舍入模式。
- 支持并行计算，可加速大规模矩阵运算模拟。

4. 主要结果 (Results)

模型精度：在涵盖所有支持的输入/输出精度格式（fp8, fp16, bf16, tf19, fp32）的测试中，软件模型与硬件输出实现了100% 的比特级匹配。
关键发现示例：
- 次正规数处理：在 bf16 和 tf19 输入模式下，当乘积指数处于次正规范围时，硬件会对对齐指数进行截断（例如限制在 -133 或 -126），这会导致特定的截断行为，之前的模型未完全捕捉到这一点。
- 交错计算模式：在 H100/H200/B200 上使用 mma.sync 处理 fp8 输入时，32 个输入元素被交错分配到两个 Tensor Core 调用中，这种模式只有在通过随机测试和 $N_{FMA}$ 分析才能发现。
- 多字算法误差分析：利用该模型复现了多字（Multi-word）高精度矩阵乘法实验。结果显示，不同代际的 Tensor Core（如 V100 vs B200）在累积误差上存在显著差异（可达两个数量级），这主要归因于累加器精度和舍入模式的不同。
工具实用性：证明了该工具箱可用于验证测试向量的鲁棒性，并在无需访问真实硬件的情况下进行算法开发。

5. 意义与影响 (Significance)

科学计算的可复现性：为科学计算社区提供了理解不同 GPU 架构数值行为差异的工具，有助于解决跨平台结果不一致的问题。
算法开发与优化：混合精度算法开发者可以利用这些模型在部署前精确预测数值误差，优化算法参数（如选择最佳的精度组合或迭代 refinement 策略）。
推动标准化：通过揭示当前硬件实现的多样性，为 IEEE P3109 等低精度算术标准的制定提供了实证数据支持，强调了标准化减少操作（如多术语加法）的重要性。
未来工作基础：该工具箱不仅支持现有硬件，其参数化设计还允许模拟未来的自定义 Tensor Core 变体，为下一代硬件设计提供了验证平台。

总结：这篇论文通过系统性的测试和迭代建模，填补了 NVIDIA Tensor Core 数值行为知识的空白，并提供了业界首个经过严格验证的、高精度的 MATLAB 模拟工具，极大地促进了混合精度计算在 AI 和科学计算领域的可靠性研究。