译文-击败 LLM 推理中的非确定性
date
Sep 14, 2025
slug
defeating-nondeterminism-in-llm-inference-translate
status
Published
tags
LLM
RL
翻译
summary
type
Post
‣
可重现性是科学进步的基石。然而，从大型语言模型中获得可重现的结果却异常困难。
例如，您可能观察到多次向 ChatGPT 询问相同问题会提供不同的结果。这本身并不令人惊讶，因为从语言模型获取结果涉及"采样"，这是一个将语言模型的输出转换为概率分布并概率性地选择一个标记（token）的过程。
可能更令人惊讶的是，即使我们将温度（temperature）调整到 0¹（这意味着大型语言模型（LLM）始终选择最高概率的标记，这被称为贪婪采样（greedy sampling）），从而使采样在理论上成为确定性的，LLM API 在实践中仍然不是确定性的（参见以往的讨论这里、这里或这里）。即使在您自己的硬件上使用诸如 vLLM 或 SGLang 等开源推理库运行推理，采样仍然不是确定性的（参见这里或这里）。
但为什么 LLM 推理引擎不是确定性的呢？一个常见的假设是浮点非结合性和并发执行的某种组合导致基于哪个并发核心首先完成的非确定性。我们将其称为 LLM 推理非确定性的"并发 + 浮点"假设。例如，最近的一篇 arXiv 预印本写道：
GPU 中的浮点运算表现出非结合性，意味着 ，这是由于有限精度和舍入误差造成的。这个性质直接影响 Transformer 架构中注意力分数和 logits 的计算，其中跨多个线程的并行操作可以根据执行顺序产生不同的结果。
您也可以发现其他人重复"并发 + 浮点"假设，比如这里（"存在速度权衡，为了使端点快速，使用了 GPU，它们进行并行[非确定性]计算。任何现代 GPU 神经网络计算都会受到这些影响。"），或者这里（"因为 GPU 是高度并行化的，加法或乘法的顺序在每次执行时可能不同，这可能会级联到输出中的小差异。"）。
虽然这个假设并非完全错误，但它没有揭示完整的图景。例如，即使在 GPU 上，在相同数据上重复运行相同的矩阵乘法总是会提供按位相等的结果。我们确实在使用浮点数。我们的 GPU 确实有很多并发性。为什么我们在这个测试中看不到非确定性？
为了理解 LLM 推理非确定性的真正原因，我们必须深入研究。
不幸的是，即使定义 LLM 推理确定性的含义也很困难。可能令人困惑的是，以下陈述都同时为真：
- GPU 上的一些核（kernel）是非确定性的。
- 然而，语言模型前向传播中使用的所有核都是确定性的。
- 此外，LLM 推理服务器（如 vLLM）的前向传播也可以被声称是确定性的。
- 尽管如此，从任何使用推理服务器的人的角度来看，结果都是非确定性的。
在这篇文章中，我们将解释为什么"并发 + 浮点"假设错过了要点，揭开 LLM 推理非确定性背后的真正罪魁祸首，并解释如何击败非确定性并在 LLM 推理中获得真正可重现的结果。
原罪：浮点非结合性
在谈论非确定性之前，解释为什么会有数值差异是有用的。毕竟，我们通常将机器学习模型视为遵循交换律或结合律等结构规则的数学函数。难道不应该有一个我们的机器学习库应该为我们提供的"数学上正确"的结果吗？
罪魁祸首是浮点非结合性。也就是说，对于浮点数：
具有讽刺意味的是，打破结合性正是使浮点数有用的原因。
浮点数之所以有用，是因为它们允许"动态"精度级别。为了解释的目的，我们将使用十进制（而不是二进制），其中浮点数的格式为 。我们还将为尾数使用 3 位数字，为指数使用 1 位数字。
例如，对于值 3450，我们可以将其精确表示为 。我们也可以表示更小的值，如 0.486 为 。通过这种方式，浮点允许我们表示非常小以及非常大的值。在科学中，我们可以说浮点允许我们维护恒定数量的"有效数字"。
如果您将两个具有相同指数的浮点数相加，它看起来类似于整数加法。例如，123（）+ 456（）得到 579（）。
但是当我们将两个具有不同指数的浮点数相加时会发生什么，比如 1230 和 23.4？在这种情况下，确切的结果是 1253.4。然而，我们一次只能保持 3 位数的精度。因此，浮点加法将丢弃最后 2 位数字并获得值 （或 1250）。
但此时，我们已经破坏了信息。请注意，每当我们将两个具有不同"尺度"（即不同指数）的浮点数相加时，这种情况就会发生。而将不同指数的浮点数相加是经常发生的。事实上，如果我们能保证永远不需要不同的指数，我们就可以只使用整数！
换句话说，每当我们以不同的顺序将浮点数相加时，我们就可能得到完全不同的结果。举一个极端的例子，根据顺序的不同，对这个数组求和有 102 种不同的结果。
尽管这是非相同输出的根本原因，但它没有直接回答非确定性来自哪里。它没有帮助我们理解为什么浮点值会以不同的顺序相加、何时发生这种情况以及如何避免它。
答案在于核的实现方式。
为什么核不总是以相同的顺序添加数字？
如上所述，核为什么以不同顺序添加数字的一个常见解释是"并发 + 浮点"假设。该假设指出，如果并发线程完成的顺序是非确定性的，并且累积顺序取决于并发线程完成的顺序（比如原子加法（atomic add）），那么我们的累积顺序也将是非确定性的。
令人困惑的是，虽然这可能导致非确定性核，但并发性（和原子加法）最终在 LLM 推理非确定性中完全不涉及！为了解释真正的罪魁祸首是什么，让我们首先理解为什么现代 GPU 核很少需要原子加法。
何时需要原子加法？
通常 GPU 跨许多"核心"（即 SM）并发启动程序。由于核心之间没有固有的同步，如果核心需要相互通信，这就构成了挑战。例如，如果所有核心都必须累积到同一个元素，您可以使用"原子加法"（有时称为"取并加"）。原子加法是"非确定性的"——结果累积的顺序纯粹取决于哪个核心首先完成。
具体来说，想象您正在用 100 个核心归约一个 100 元素向量（例如
torch.sum()）。虽然您可以并行加载所有 100 个元素，但我们最终必须归约到单个元素。完成此操作的一种方法是使用某种"原子加法"原语，其中硬件保证所有加法都将被处理，但不保证顺序。
这通常是人们所说的"非确定性"的含义——您用完全相同的输入执行相同的核两次，但得到不同的结果。这被称为运行间非确定性（run-to-run nondeterminism），即您用完全相同的依赖项运行相同的 Python 脚本两次，但得到不同的结果。
虽然并发原子加法确实使核非确定性，但对于绝大多数核来说，原子加法是不必要的。事实上，在 LLM 的典型前向传播中，通常没有单个原子加法存在。
考虑到并行化归约可以从原子加法中受益，这可能令人惊讶。原子加法最终不需要的原因有两个主要原因。
- 沿着"批次"维度通常有足够的并行性，因此我们不需要沿着归约维度并行化。例如，假设我们不是归约单个 100 维向量，而是并行归约 500 个向量。在这种情况下，我们可以在每个核心中归约一个完整的向量，并允许每个核心处理不同的向量。
- 随着时间的推移，大多数神经网络库都采用了各种策略来在不牺牲性能的情况下实现确定性。例如，我们可以执行"分割"（或树）归约，其中我们将 100 元素归约分成五个 20 元素归约（因此实现五路并行性）。然后，为了组合其余的五个元素，我们可以执行单独的"清理"归约（它不是并行化的，但对足够少的元素进行操作以保持便宜）或利用信号量（它确保每个并发线程块将以确定性顺序累积）²。
由于这两个因素，避免原子加法对于绝大多数神经网络操作来说是微不足道的性能损失。
仍然有几个常见操作在避免原子时会有显著的性能损失。例如，PyTorch 中的
scatter_add（a[b] += c）。然而，在 LLM 中唯一常用的是 FlashAttention 反向传播³。然而，LLM 的前向传播不涉及需要原子加法的操作。因此，LLM 中的前向传播实际上是"运行间确定性的"。
维基百科写道："确定性算法是一种给定特定输入，总是产生相同输出的算法。"在这种情况下，给定完全相同的输入（即推理服务器正在处理的确切请求），前向传播总是产生完全相同的输出。
然而，前向传播本身是"确定性的"并不足以确保包含它的系统是确定性的。例如，如果我们请求的输出依赖于并行用户请求（例如批量标准化（batch-norm））怎么办？由于每个单独的请求无法知道并行请求将是什么，从它们的角度来看，我们的整体 LLM 推理也是非确定性的！
事实证明，我们请求的输出确实依赖于并行用户请求。不是因为我们以某种方式在批次间泄露信息——而是因为我们的前向传播缺乏"批次不变性"，导致我们请求的输出依赖于我们前向传播的批次大小。
批次不变性和"确定性"
为了解释批次不变性，让我们简化系统并仅查看矩阵乘法（matmul）。您可以假设所有矩阵乘法实现都是"运行间确定性的"⁴。然而，它们不是"批次不变的"。换句话说，当批次大小改变时，批次中的每个元素都可能得到不同的结果。
从数学角度来看，这是一个相当不寻常的特性。矩阵乘法应该沿着批次中的每个元素"独立"——批次中的其他元素或批次大小都不应该影响批次中特定元素的计算结果。
然而，正如我们可以经验性地观察到的，这并不成立。
请注意，这是"运行间确定性的"。如果您多次运行脚本，它将确定性地返回相同的结果⁵。
然而，当非批次不变核被用作更大推理系统的一部分时，系统可能变得非确定性。当您向推理端点进行查询时，服务器承受的负载量从用户的角度来看实际上是"非确定性的"。负载决定了核运行的批次大小，从而改变每个单独请求的最终结果！
如果您将核不变的某种属性（即批次大小）与该属性的非确定性（即服务器承受的负载）结合，您就得到了一个非确定性系统。
换句话说，几乎所有 LLM 推理端点非确定性的主要原因是负载（因此批次大小）非确定性地变化！这种非确定性并非 GPU 独有——从 CPU 或 TPU 提供的 LLM 推理端点也将有这种非确定性来源。
所以，如果我们想在推理服务器中避免非确定性，我们必须在核中实现批次不变性。为了理解如何实现这一点，让我们首先看看为什么核首先没有批次不变性。
如何使核批次不变？
为了使 Transformer 实现批次不变，我们必须使每个核批次不变。幸运的是，我们可以假设每个逐点操作都是批次不变的⁶。因此，我们只需要担心涉及归约的 3 个操作——RMSNorm、矩阵乘法和注意力⁷。
方便的是，这些操作也按照实现批次不变性难度的递增级别排序。每一个都需要一些额外的考虑来以合理的性能实现批次不变性。让我们首先谈谈 RMSNorm。
批次不变 RMSNorm
批次不变性的要求是，无论核的批次大小如何，每个元素的归约顺序都必须固定。请注意，这并不意味着我们必须始终使用相同的归约策略。例如，如果我们改变要归约的元素数量，即使我们的归约策略发生变化，我们仍然可以是批次不变的⁸。
因此，只有当我们的批次大小影响归约策略时，我们才会打破批次不变性。
让我们看看 RMSNorm 的标准并行策略。一般来说，并行算法受益于最小化跨核心的通信。为了这次讨论的目的，您可以假设当我们提到"核心"时，我们指的是 SM。更具体地说，这里重要的特性是我们的核启动的线程块数量大于 SM 的数量。因此，我们可以开始的一个策略是将每个批次元素分配给一个核心，如图5所示。
增加我们的批次大小不会影响我们的归约策略；如果批次大小为 200 为我们的核提供足够的并行性，那么批次大小为 2000 绝对会提供足够的并行性。
另一方面，减少批次大小可能会带来挑战。因为我们将每个批次元素分配给一个核心，减少我们的批次大小最终将导致核心数量多于批次元素，使一些核心处于空闲状态。
在遇到这种情况时，一个好的核工程师会采用前面部分提到的解决方案之一（原子加法或分割归约），保持良好的并行性，从而获得良好的性能。不幸的是，这改变了归约策略，阻止了这个核成为批次不变的。
最简单的解决方案是简单地完全忽略这些情况。这并非完全不合理——小批次大小意味着核可能很快就会执行完毕，因此速度下降可能不会是灾难性的。
如果我们不得不优化这个用例，一种方法是始终使用即使对于非常小的批次大小也有足够并行性的归约策略。这种归约策略对于更大的批次大小会导致过多的并行性，但允许我们在整个大小范围内实现不错的（但不是峰值）性能。
批次不变矩阵乘法
从核心上讲，您可以将矩阵乘法简单地视为逐点操作后跟归约。然后，如果我们通过将输出分块为 tile 来并行化我们的矩阵乘法，我们就有了一个类似的"数据并行"核策略，将每个归约保持在一个核心内。
也类似于 RMSNorm，我们的"批次"维度（M 和 N）可能变得太小，迫使我们沿着归约维度（K）分割。尽管有两个"批次"维度，矩阵乘法也要求我们每个核心有更多的"工作"以便有效地利用张量核心。例如，如果您有一个 [1024, K] x [K, 1024] 的矩阵乘法和标准的 2D tile 大小 [128, 128]，数据并行策略只能将这个矩阵乘法分割成 64 个核心，不足以饱和 GPU。
矩阵乘法中沿着归约维度分割被称为 Split-K 矩阵乘法。就像 RMSNorm 一样，使用这种策略会破坏批次不变性。
矩阵乘法的另一个有趣的并行策略是 stream-k。Stream-k 很有趣，因为它比典型的矩阵乘法具有更少的不变性。如所讨论的，大多数矩阵乘法库不是批次不变的，但它们至少是您可以称之为批次位置不变的（即改变元素在批次内的位置不影响数值）。然而，stream-k 也不是批次位置不变的！其核心洞察是您可以通过对不同的输出 tile 以不同方式沿 k 分割来获得更清洁的负载平衡，但利用这一点也使我们的核不是批次位置不变的。
矩阵乘法有一个额外的复杂性——张量核心指令。尽管对于归约我们可以简单地一次操作一行，有效的矩阵乘法核必须一次操作整个"tile"。
每个张量核心指令（比如
wgmma.mma_async.sync.aligned.m64n128k16）在内部可能有不同的归约顺序。使用不同张量核心指令的一个原因可能是批次大小非常小。例如，如果我们使用操作长度为 256 的 tile 的张量核心 PTX 指令，但批次大小只有 32，我们就浪费了几乎所有的计算！在批次大小为 1 时，最快的核通常根本不使用张量核心。
因此，确保矩阵乘法批次不变性的最简单方法是编译一个核配置并将其用于所有形状。尽管我们会失去一些性能，但这在 LLM 推理中通常不是灾难性的。特别是，当 M 和 N 都很小时最需要 split-k，幸运的是在我们的情况下，N（即模型维度）通常相当大！⁹
批次不变注意力
在获得矩阵乘法的批次不变性之后，注意力引入了两个额外的复杂性——恰当地说，因为它包含两个矩阵乘法。
- 与只沿特征维度归约的 RMSNorm 和矩阵乘法相反，我们现在沿特征维度和序列维度归约。
- 由于上述原因，注意力必须处理影响序列处理方式的各种推理优化（分块预填充（chunked prefill）、前缀缓存（prefix caching）等）。
因此，为了在 LLM 推理中实现确定性，我们的数值必须既不变于一次处理多少请求，也不变于每个请求在推理引擎中如何被切片。
让我们首先走过注意力的标准并行策略，这首先在 FlashAttention2 中引入。类似于 RMSNorm 和矩阵乘法，默认策略是"数据并行"策略。由于我们沿着键/值张量归约，数据并行策略只能沿着查询张量并行化。
例如，根据推理引擎的选择，一个序列可能在几个部分中处理（如在分块预填充中）或者一次全部处理（如果预填充没有被拆分）。为了实现"批次不变性"，给定标记的归约顺序不依赖于来自其序列的多少其他标记正在同时处理是必要的。如果您将 KV 缓存中的 K/V 值与正在处理的当前标记中的 K/V 值分开归约（如在 vLLM 的 Triton 注意力核中），这无法实现。例如，在处理序列中的第 1000 个查询标记时，无论 KV 缓存中有 0 个标记（预填充）还是 999 个标记（解码），归约顺序都必须相同。
为了解决这个问题，我们可以在注意力内核本身之前更新KV缓存和页表，确保无论处理多少个token，我们的键和值始终以一致的方式布局。
有了这个额外的细节（以及前面部分提到的所有内容，例如一致的 tile 大小），我们就能够实现一个批处理不变的注意力实现！
然而，这里有一个重要的问题。与矩阵乘法不同，我们在 LLM 推理中看到的注意力形状通常确实需要一个拆分规约的核函数，通常称为 Split-KV 或 FlashDecoding。这是因为如果我们不沿着规约维度并行化，我们只能沿着批次维度、头维度和“查询长度”维度并行化。在注意力机制的解码阶段，查询长度非常小，因此除非我们有非常大的批次大小，否则我们通常无法使 GPU 饱和。
不幸的是，像处理 RMSNorm 和 Matmuls 那样忽略这种情况并不容易。例如，如果你有一个非常长的 KV 缓存，尽管只处理一个请求，注意力核函数也可能花费很长时间。
此外，通常用于注意力机制的拆分规约策略也对批次不变性构成了挑战。例如，FlashInfer 的“平衡调度算法”选择了能够使 GPU 所有核心饱和的最大拆分大小，从而使规约策略不是“批次不变的”。然而，与 RMSNorm/Matmuls 不同，无论批次大小如何，选择一个固定的拆分数量是不够的。
相反，为了实现批次不变性，我们必须采用“固定拆分大小”的策略。换句话说，我们不是固定拆分的数量，而是固定每次拆分的大小，然后最终得到一个可变数量的拆分。通过这种方式，我们可以保证无论我们正在处理多少个 token ，我们总是执行相同的规约顺序。
这需要一些我们代码发布中未包含的内部 FlexAttention 更改。我们将在不久的将来将它们上游化！
实现
我们通过利用 vLLM 的 FlexAttention 后端以及
torch.Library，提供了一个在 vLLM 之上进行确定性推理的演示。通过 torch.Library，我们能够以非侵入性的方式替换掉大多数相关的 PyTorch 操作符。你可以在 thinking-machines-lab/batch-invariant-ops 找到“批次不变”核函数库，以及在“确定性”模式下运行的 vLLM 示例。实验
补全结果的不确定性有多大？
我们使用 Qwen/Qwen3-235B-A22B-Instruct-2507，在温度为 0 的情况下，使用提示“Tell me about Richard Feynman”（非思考模式）采样 1000 个补全，每个生成 1000 个 token 。令人惊讶的是，我们生成了 80 个独特的补全，其中最常见的出现了 78 次。
观察补全的不同之处，我们发现补全内容在前 102 个 token 上实际上是完全相同的！补全内容出现分歧的第一个实例发生在第 103 个 token 。所有补全都生成了序列“Feynman was born on May 11, 1918, in”。然而，其中 992 个补全接着生成了“Queens, New York”，而 8 个补全生成了“New York City”。
另一方面，当我们启用批次不变的核函数时，我们所有的 1000 个补全都是完全相同的。这是我们从采样器中数学上所期望的，但如果没有我们的批次不变的核函数，我们是无法实现确定性结果的。
性能
我们尚未在优化批次不变核函数的性能方面投入大量精力。但是，让我们运行一些实验来验证我们的性能仍然可用。
我们将设置一个带有一个 GPU 的 API 服务器，运行 Qwen-3-8B，并请求 1000 个序列，输出长度在 90 到 110 之间。
配置 | 时间（秒） |
vLLM 默认 | 26 |
未优化的确定性 vLLM | 55 |
+ 改进的注意力核函数 | 42 |
大部分的性能下降来自于 vLLM 中的 FlexAttention 集成尚未经过深度优化。尽管如此，我们看到性能并非灾难性的。
真正的在线策略强化学习
正如研究人员指出的，训练和推理之间不同的数值，隐含地将我们的在线策略（on-policy）强化学习（RL）变成了离线策略（off-policy）强化学习。
当然，如果我们甚至无法从两个相同的推理请求中获得逐位相同的结果，那么在训练和推理之间获得逐位相同的结果是不可能的。然后，确定性推理使我们能够修改我们的训练栈，以在采样和训练之间获得逐位相同的结果，从而实现真正的同策略强化学习。
我们在 Bigmath 上的 RLVR 设置中进行了实验，其中 RL 策略从 Qwen 2.5-VL instruct 8B 初始化，最大 rollout 长度为 4096。
如果我们不进行离策略校正（即重要性加权）进行训练，我们的奖励会在训练中途崩溃，而添加一个离策略校正项则可以使训练顺利进行。但是，如果我们在采样器和训练器之间实现了逐位相同的结果，我们就是完全的同策略（即 0 KL 散度），并且也可以顺利训练。
我们还可以绘制采样器和训练器之间对数概率（logprobs）的 KL 散度，其中所有 3 次运行都表现出显著不同的行为。当使用重要性加权运行时，它保持在 0.001 左右，并有偶尔的峰值。然而，不使用重要性加权运行最终会导致 KL 散度在奖励崩溃的大约同一时间出现峰值。当然，当运行“真正的同策略强化学习”时，我们的 KL 散度保持在 0 的平坦线上，表明训练策略和采样策略之间没有分歧。
结论
现代软件系统包含多层抽象。在机器学习中，当我们遇到不确定性和细微的数值差异时，我们常常倾向于将它们掩盖过去。毕竟，我们的系统已经是“概率性”的了，多一点不确定性又有什么关系呢？在失败的单元测试中提高 atol/rtol 的值有什么问题呢？训练器和采样器之间的对数概率差异可能不是一个真正的错误，对吧？
我们拒绝这种失败主义。只要稍加努力，我们就能理解不确定性的根源，甚至解决它们！我们希望这篇博文能为社区提供一个关于如何解决推理系统中不确定性的坚实理解，并激励其他人去全面了解他们的系统。
脚注：
¹ 这意味着大型语言模型（LLM）始终选择最高概率的标记，这被称为贪婪采样（greedy sampling）。
² 信号量策略可以在这里找到描述。
³ 有趣的事实：您知道广泛使用的 FlashAttention 反向传播的 Triton 实现实际上在算法上与 Tri Dao 的 FlashAttention-2 论文不同吗？标准的 Triton 实现在反向传播中进行额外的重计算，避免原子但成本增加 40% FLOPs！
⁴ 这并不完全正确，但大多数常见的矩阵乘法实现确实具有这个特性。
⁵ 它不是"硬件/软件版本不变的"——您的 GPU/PyTorch 版本可能返回不同的值，但它应该确定性地返回相同的值。
⁶ 尽管这对于 PyTorch 中的所有核都是真的，但它并非本质上真实。例如，CPU 上的一些核实现会在数组的某些部分使用向量化内在函数，在其他部分使用非向量化内在函数，这些内在函数不一定总是具有按位相同的数值。
⁷ 与并行性相关的归约超出了本讨论的范围，但相同的原理适用。一个可能有用的事实是，在 Blackwell 以及带有 CUDA 12.8+ 的 Hopper 上，NVLink-Sharp 交换机内归约是确定性的。与许多事情一样，这个信息可以在 NCCL 的 github issues 上找到。
⁸ The Quack 博客文章有一些很好的例子，展示了您可以执行的各种归约策略的层次结构（例如线程归约、warp 归约、块归约、集群归约）。
⁹ 可以在这里找到有关 LLM 推理中矩阵乘法形状的有趣分析。意力核本身之前更新 KV 缓存和页表，确保我们的键和值始终一致地布局，无论正在处理多少标记。
有了这个额外的细节（以及前面部分提到的所有内容，如一致的 tile 大小），我们能够实现批次不变的注意力实现！
然而，这里有一个重大问题。与矩阵乘法不同，我们在 LLM 推理中看到的注意力形状通常确实需要分割归约核，通常称为 Split-KV 或 FlashDecoding。这是因为如果我们不沿着归约并行化，我们只能沿着批次维度、头维度和"查询长度"维度并行化。在注