H100 Tensor Core惊天秘密：精度狂降10位！

矩阵乘法引擎的精度：那些我们不曾留意的细节

在当前人工智能技术蓬勃发展的时代，高性能计算硬件，特别是图形处理器（GPU）及其内置的矩阵乘法引擎，已成为推动AI进步的核心动力。无论是深度学习模型的训练还是推理，矩阵乘法（也常被称为matmul或GEMM）都是最基础且计算量最大的运算之一。为了高效完成这些任务，英伟达（NVIDIA）等公司专门设计了Tensor Core这样的专用计算单元。这些引擎能够高效地处理小规模张量块的矩阵乘法。因此，无论是编译器还是相关的软件库，通常会将大型的矩阵乘法问题分解成许多小块，然后分发给这些高效的引擎进行处理。

通常情况下，一个采用FP8（e4m3格式）输入的Tensor Core执行的矩阵乘法，其输入张量形状为(block_size_m, block_size_k)和(block_size_k, block_size_n)，我们预期其输出会是一个FP32（e8m23格式）的(block_size_m, block_size_n)张量。然而，有一个有趣但用户鲜少关注的细节是，出于硬件效率的考量，这个FP32输出的有效尾数位（mantissa bits）可能少于标准的23位。换句话说，这项Tensor Core操作的实际精度可能比表面上看起来的FP32要低。

新媒网跨境获悉， 这种硬件设计选择在特定情况下已被报告会影响模型的最终精度。对于广大GPU用户而言，了解并验证所使用硬件的实际设计至关重要。因为即使现有的硬件无法更改，我们仍然可以通过编写适当的定制化内核（custom kernels），在需要时最大限度地保留计算精度。对于硬件设计师来说，拥有一种便捷高效的方式来量化这种设计选择带来的影响，同样具有重要意义。

在深入探讨这些细节之前，我们需要理解“累加器”（accumulator）在矩阵乘法引擎中的作用，以及为什么会采用降低的精度。

累加器的核心作用与精度考量

让我们首先设想一个能够处理FP8矩阵乘法，其输入块大小分别为(3, 4)和(4, 3)的计算引擎。如图1a所示，该引擎的运作机制是其高效性的关键。
图片说明

深入观察这个计算引擎的内部，最基本的运算是一个行-列内积，即 cᵢⱼ = ∑ₖ aᵢₖ * bₖⱼ。我们可以想象，一个高效的硬件设计会简单地实现多个乘法器，用于计算每一对 aᵢₖ 和 bₖⱼ 的乘积。紧接着，这些中间结果会通过一系列加法器进行求和。例如，图1b清晰地展示了4个乘法器紧随3个加法器的配置。

在这个简单的例子中，乘法部分可以假设有足够的乘法器并行处理，从而在一个“计算步骤”内完成。然而，加法部分则需要2个计算步骤才能完成，因为它必须以一种分层、串行的方式进行。如果我们将这种单元设计扩展到处理N个元素，乘法仍旧只需要一个步骤，但加法将需要 log(N) 个步骤。这凸显了加法在计算复杂度上的潜在瓶制。

更进一步来看，每个乘法器只需要计算FP8 * FP8（e4m3格式），这涉及到4位+4位的指数相加以及4位x4位的尾数相乘。然而，由于每个中间乘积都需要正确对齐，后续的加法器必须使用比乘法器多得多的位数来处理。

如图2所示（这仅仅是一个示意图，并非真实的FP8案例），将两个只有4位尾数的有限精度浮点数相加，其结果可能是一个需要更多尾数位的浮点数。这粗略地解释了为何浮点乘加（MAC）操作的电路复杂度和成本（硅片面积和功耗）与累加精度有着密切的关系。因此，尽管使用FP32作为累加精度（如图2b所示）更为安全可靠，但探索使用降低的累加精度，以求在性能和成本上取得突破，无疑是值得的。
图片说明

通过这些例子，我们可以清楚地看到，在矩阵乘法引擎中使用降低精度的加法器所带来的益处。这是硬件工程师在追求极致性能与效率之间精心权衡的体现，也是我们理解现代AI加速器工作原理的关键一环。

如何验证累加器精度？以Tensor Core为例

鉴于矩阵乘法累加器可能被设计成少于23个尾数位，其真实输出实际上是e8mNacc格式（其中Nacc < 23），其余的位通过补零扩展到e8m23。换句话说，FP8 Tensor Core的输出看起来像是FP32，但在计算过程中，任何小于e8mNacc的数值都没有被真正计算出来。

在本文中，我们将演示一种使用Triton语言内核来探究累加器精度的简单方法。假设Tensor Core的输出只有Nacc个有效尾数位（如e8mNacc格式），即最后23 - Nacc位本来就是零。那么，如果我们对Tensor Core的输出应用一个掩码，截断最后Ntrun位，只要Ntrun ≤ 23 - Nacc，最终的矩阵乘法结果就应该保持不变。通过不断调整Ntrun并将其与一个参考结果（即Ntrun = 0时）进行比较，我们就可以推断出所研究的浮点矩阵乘法单元的累加器精度。这里，“截断Ntrun位”指的是将浮点数的最后Ntrun位（即尾数的最低有效位）置零。

为何选择Triton语言？

我们选择使用Triton语言进行这项实验，原因在于其卓越的通用性。它允许我们提出的方法能够推广到其他支持Triton的加速器上，极大地扩展了研究的范围和适用性。更重要的是，Triton语言的简洁性和对加速器恰到好处的控制能力，极大地加快了本次实验的开发进程。它为开发者提供了一种高效的途径，去直接操作底层硬件特性，从而更深入地理解和优化计算过程。

虽然Triton语言作为一个相对年轻的编程框架，预计会随着时间的推移不断演进和完善，但我们的实现是基于Triton官方提供的矩阵乘法教程进行构建的。这意味着即使Triton未来有所更新，我们预期需要修改的代码量也会非常小，从而保证了实验方法的长期有效性和稳定性。Triton的这种设计哲学，无疑为研究人员和工程师提供了一个强大的工具，去探索和验证高性能计算的深层机制。

实验设计与实现

为了验证累加器的实际精度，我们进行了一系列精心设计的实验。在本文末尾，读者可以找到完整的可运行代码。我们的方法基于Triton语言中的一个矩阵乘法教程内核，并在此基础上添加了一个简单的截断功能。由于原始教程已经包含了大量的细节，我们将重点介绍我们为实现截断功能所做的修改。

粗略来说，一个大型的矩阵乘法 matmul(A, B) 操作会被分解成许多小块，这些小块随后并行处理。每个A和B的块分别具有(BLOCK_SIZE_M, BLOCK_SIZE_K)和(BLOCK_SIZE_K, BLOCK_SIZE_N)的形状。块级别的矩阵乘法由Triton的 tl.dot() 函数计算，产生一个形状为(BLOCK_SIZE_M, BLOCK_SIZE_N)的临时张量 accumulator_inner。这个 accumulator_inner 被假设只具有Nacc个有效尾数位。

我们的实验核心在于对这个 accumulator_inner 进行精确的位操作处理：

对 accumulator_inner 进行截断： 我们使用位操作和一个预定义的掩码来截断 accumulator_inner 的最后Ntrun位。为了简化处理，我们忽略了舍入操作，直接将 round_bit 设置为0。以下是实现这个功能的关键代码片段：
```
def prep_round_and_trun_mask(trun_bits):
        round_bit = 1 << (trun_bits - 1) if trun_bits > 0 else 0
        trun_mask = ~tl.cast((1 << trun_bits) - 1, tl.uint32)
        return round_bit, trun_mask

def round_and_trun(x, round_bit, trun_mask):
        """Round and truncate (usually for accumulator)."""
        return libdevice.uint_as_float(
            (libdevice.float_as_uint(x) + round_bit) & trun_mask
)
```
这段代码首先准备了一个截断掩码 trun_mask，用于将最低的 trun_bits 位清零。round_bit 则是在进行位操作时，如果需要四舍五入时使用的参数，此处我们将其设为0以简化。round_and_trun 函数通过将浮点数转换为无符号整数，然后进行位加法（如果考虑舍入）和位与操作（进行截断），最后再转换回浮点数，从而实现了对尾数的精准控制。
K维度上的累加： 每一个经过截断处理的 accumulator_inner 随后会被累加到一个预先分配好的FP32张量 accumulator 中，这个过程沿着K维度逐步进行。accumulator 张量的形状与 accumulator_inner 相同，确保了中间累加结果的完整性和一致性。
结果写回： 在完成K维度上的所有迭代累加后，最终的 accumulator 值会被写回到目标输出张量C的相应块中，张量C的形状为(M, N)。这一步标志着整个块级矩阵乘法计算的完成。

通过这种细致入微的实验设计，我们能够精确地控制和观察累加器内部的精度行为，从而揭示硬件在处理浮点运算时的真实能力和潜在限制。这不仅有助于我们更好地理解现有硬件，也为未来高性能计算系统的设计提供了宝贵的经验。

实验结果与深入讨论

通过我们的实验，无论是从表1的数据还是图3的视觉呈现中，我们都观察到一个显著的现象：在使用英伟达（NVIDIA）H100 GPU的FP8 Tensor Core进行矩阵乘法时，即使我们截断了输出结果中多达10个最低有效位的尾数位，其计算结果与未进行任何截断的情况也完全一致。这强有力地表明，这些被截断的位在原始输出中本来就已经是零。
图片说明

新媒网跨境了解到， 这一精妙的实验结果暗示，H100 Tensor Core的累加器可能为了计算效率，采用了特殊的FP22格式（即e8m13），而非标准的FP32完整精度。随后，我们在基于Ada Lovelace架构的英伟达（NVIDIA）RTX4000系列GPU上重复了相同的实验，并观察到了完全一致的行为，这进一步验证了我们的发现。这种精巧的设计体现了硬件工程师在性能与成本之间寻求平衡的智慧，他们通过对计算流程的优化，使得硬件在保持足够精度的同时，实现了更高的运行效率。

对编译器行为的验证

需要强调的一个重要考量是，这项实验依赖于Triton编译器将Triton代码转换为等效的CUDA代码。因此，我们必须确保执行任务的Tensor Core确实是我们打算检查的FP8类型。在少数情况下，Triton编译器可能会选择使用FP16 Tensor Core指令来处理某些FP8计算，这会影响我们对累加器精度的判断。

最可靠的确认硬件实际执行指令的方法是使用英伟达（NVIDIA）的性能分析工具ncu（包含在CUDA工具包中），来检查与Triton tl.dot 调用相关的底层CUDA指令。读者可以将此实验笔记本保存为Python文件，然后使用以下命令行指令启动ncu进行分析：

/usr/local/cuda-13.0/bin/ncu --target-processes all --set full --import-source yes -f --kernel-name matmul_kernel --launch-skip 3 --launch-count 1 -o ./tl_fp8mm_backend_H100 python accumulator_precision_test.py

从图5所示的ncu性能分析报告中，我们发现，对于我们选择的块大小（MxNxK=64x64x32），FP8xFP8 tl.dot() 操作确实被翻译成了QGMMA指令——这是一种FP8 Tensor Core专用的指令。这最终证实了实验中确实使用了FP8 Tensor Core，保证了我们对累加器精度判断的准确性。
图片说明

如前所述，Triton编译器有时会为tl.dot选择不同的实现方式。例如，如果我们在kernel_config字典中将num_warps设置为2并重复实验，Triton编译器会智能地将FP8数据打包成FP16格式，并转而使用HMMA指令执行计算。HMMA是FP16 Tensor Core的专用指令。在这种情况下，相应的实验结果显示，FP16 Tensor Core的累加器精度只比FP32少1位，进一步印证了编译器在优化时的灵活策略以及不同Tensor Core的精度特性。
图片说明

块大小与精度权衡的深远影响

此外，由于专门的矩阵乘法单元通常设计用于处理特定固定大小的输入，如果我们选择的BLOCK_SIZE超出了Tensor Core的直接处理能力，编译器或CUDA库会自动将操作分解为几个较小的操作。在我们的Triton代码中，如果我们将BLOCK_SIZE_K增加到128，并通过ncu再次验证，我们会发现每个WGMMA指令（FP8 Tensor Core指令的一种变体）仅能处理K=32的情况。这意味着，为了结合多个Tensor Core调用的部分结果，会额外涉及一次求和操作。

一个自然而然的问题是：这种中间求和使用的是什么精度？这正是我们一直在讨论的浮点对齐和精度损失问题。根据K=128实验的输出，我们仍然观察到13个有效尾数位。这提供了一个重要的洞察：如果Triton内核选择的块大小超出了Tensor Core的基本设计能力，无论是出于性能考虑还是自动调优的结果，都可能由于降低精度的求和而导致额外的精度损失。

因此，如果矩阵乘法精度是一个关键的考量（尤其是在涉及模型训练和反向传播时），在简单地退回到FP16精度之前，我们应该首先尝试像我们在Triton代码中那样，使用中间的FP32累加。我们在实验中展示了BLOCK_SIZE_K对精度的影响。但读者应牢记，更小的块大小可能会影响内核的整体性能。在实际应用中，开发者可能需要从一个较大的块大小开始，例如如果自动调优建议使用256或512，然后逐渐减小到128，并权衡使用FP16与减小块大小之间的利弊。

值得一提的是，如果是在自定义内核中使用cuBLAS库，设置CUBLASLT_MATMUL_DESC_FAST_ACCUM标志同样可以实现累加精度的提升效果，为开发者提供了另一种灵活的优化途径。

最后，降低精度累加器的概念同样适用于INT8xINT8引擎。FP8和INT8矩阵乘法的主要区别在于，INT8累加器的截断发生在最高有效位（MSBs）而不是最低有效位（LSBs）。换句话说，对于INT8，我们需要考虑的是溢出问题，而不是FP8中的下溢问题。对提供的Triton内核进行简单的修改，就可以探究INT8的行为，这一探索性的任务我们留给感兴趣的读者。

总结与展望

我们深入探讨了矩阵乘法引擎中累加器精度设计的重要性，并展示了一种通过Triton语言验证现有加速器实际精度设计的简便方法。对累加器精度的透彻理解，对于那些对精度敏感的应用开发者、编写定制化内核的工程师，以及需要为其下一代硬件设计进行行为仿真的硬件设计师来说，都具有不可估量的价值。

新媒网跨境认为， 深入理解这些底层机制，将为中国AI产业的持续发展与技术突破奠定坚实基础。更重要的是，我们这种基于Triton内核的方法可以与PyTorch生态系统无缝结合，这意味着相同的技术可以扩展到其他现有和未来支持Triton语言的加速器上，从而显著缩短研发周期，推动高性能计算领域的创新。

新媒网（公号: 新媒网跨境发布），是一个专业的跨境电商、游戏、支付、贸易和广告社区平台，为百万跨境人传递最新的海外淘金精准资讯情报。

本文来源：新媒网 https://nmedialink.com/posts/h100-tensor-core-shocking-10-bit-precision-drop.html