Ask an inference engineer why two temperature-0 runs of the same prompt return different text, and you'll usually get the same answer: floating-point addition isn't associative, GPUs run thousands of threads concurrently, thread completion order varies, therefore results vary. It sounds airtight. It's also, for the LLM forward pass, mostly wrong — and the real cause means your output depends on what other users happen to be sending the server at the same moment.

The definitive treatment here is Horace He's September 2025 post for Thinking Machines Lab, "Defeating Nondeterminism in LLM Inference", which comes with runnable kernels. The findings are worth internalizing because they change where you look when reproducibility matters.

The usual suspect has an alibi

Floating-point non-associativity is real. (0.1 + 1e20) - 1e20 evaluates to 0.0, while 0.1 + (1e20 - 1e20) gives 0.1 — whenever you add numbers of very different magnitudes, the order of operations decides which low-order bits survive. The post demonstrates that summing the same list of floats in shuffled orders produced 102 distinct results.

But non-associativity only causes run-to-run nondeterminism if something actually changes the reduction order between runs. The classic mechanism is atomic adds: many threads accumulating into one memory location, in whatever order they finish. Here's the thing — a typical LLM forward pass contains essentially no atomic adds. (The well-known offender, FlashAttention's backward pass, is a training concern, not an inference one.) Run torch.mm on the same 2048×2048 bfloat16 matrices a thousand times and you get bitwise-identical results every time.

The post resolves the apparent paradox with four statements that are simultaneously true:

  1. Some GPU kernels are nondeterministic.
  2. Every kernel in an LLM forward pass is deterministic.
  3. An inference server's forward pass can be deterministic.
  4. From the user's point of view, the endpoint is nondeterministic.

The missing variable in statement 4 is the batch.

Your output depends on who else is online

Modern inference servers batch concurrent requests together, and the batch composition at any instant is a function of server load. The forward pass is a deterministic function of (your tokens, everyone else's tokens sharing your batch) — but you only control the first argument.

That would be harmless if kernels were batch-invariant: if the result computed for your row didn't depend on how many other rows ride along. They aren't. The post's minimal demo: take a of shape (2048, 4096) and b of shape (4096, 4096), then compare torch.mm(a[:1], b) against torch.mm(a, b)[:1]. Same row, same weights, mathematically identical. The results differ — by a maximum of about 1669.25 in that bfloat16 example. Each call is perfectly repeatable; they just disagree with each other.

Why? Performance engineering. A kernel tuned for a large batch parallelizes across batch elements and reduces within each row in one pass. At small batch sizes that strategy leaves most of the GPU idle, so libraries switch strategies — and switching strategies means changing the reduction order, which changes the floating-point result:

  • RMSNorm: the data-parallel version (one batch element per core) is batch-invariant, but small batches push libraries toward split reductions, which aren't.
  • Matmul: Split-K partitions the reduction dimension across blocks, breaking invariance. Worse, small batches can select different tensor-core instructions entirely — a wgmma.mma_async.sync.aligned.m64n128k16 instruction has its own internal reduction order. The fix — compile one kernel configuration and use it at every shape — costs about 20% versus cuBLAS in the post's measurement.
  • Attention: FlashAttention-2's data-parallel prefill is invariant, but chunked prefill and prefix caching change how a given token's reduction is sliced. Decode needs split-KV (FlashDecoding) because query length is tiny, and the standard schedules — a fixed number of splits, or FlashInfer's load-balanced split sizing — both make the split boundaries depend on batch state. The batch-invariant alternative is a fixed split size: a KV length of 1000 always becomes three 256-element splits plus one 232-element split, regardless of load.

What it looks like in practice

The post's headline experiment: 1,000 completions of "Tell me about Richard Feynman" at temperature 0, 1,000 tokens each, against Qwen3-235B-A22B-Instruct-2507. Result: 80 unique completions, the most frequent appearing just 78 times. Every completion was identical through token 102; at token 103, after "Feynman was born on May 11, 1918, in," 992 runs wrote "Queens, New York" and 8 wrote "New York City." One bit of accumulated numeric drift crossed an argmax boundary, and from there the trajectories never reconverged.

With batch-invariant kernels swapped in, all 1,000 completions were bitwise identical.

The cost, measured on one GPU with Qwen3-8B generating 1,000 sequences of 90–110 tokens: default vLLM took 26 seconds, the unoptimized deterministic build 55 seconds, and 42 seconds after improving the attention kernel — with the caveat that the FlexAttention integration wasn't heavily optimized yet. Determinism isn't free, but it's a constant factor, not an order of magnitude.

Why RL training teams care most

For chat serving, 80-way output variance is an annoyance. For reinforcement learning it's a correctness bug. If your sampler (inference engine) and trainer (training framework) produce numerically different logprobs for the same tokens, your "on-policy" RL is silently off-policy — you're training on samples from a slightly different policy than the one you're updating. In the post's RLVR experiment (initialized from Qwen 2.5-VL 8B), running without any off-policy correction made reward collapse mid-training, with a loss spike around step 318 coinciding with a KL divergence spike between sampler and trainer. Importance weighting patches it, holding KL around 0.001. With bitwise-identical sampler and trainer — true on-policy — the KL is exactly 0 and the correction term becomes unnecessary.

What you can do today

  • The companion repo, batch_invariant_ops (MIT), swaps batch-invariant implementations of torch.mm, torch.addmm, torch.log_softmax, and torch.mean into existing code via torch.Library, behind a context manager:
python
from batch_invariant_ops import set_batch_invariant_mode

with set_batch_invariant_mode():
    output = model(input_tensor)
  • vLLM's reproducibility documentation is explicit that vLLM "does not guarantee reproducible results by default, for the sake of performance." For offline inference you can set VLLM_ENABLE_V1_MULTIPROCESSING=0 to make scheduling deterministic (with the side effect that vLLM mutates your process's random state), or enable batch invariance. For an online server, batch invariance is the only option — you can't control what the scheduler batches together. And either way, reproducibility holds only on the same hardware with the same vLLM version.

The durable lesson: when a system is "nondeterministic," ask what the output is actually a function of before blaming the hardware. The kernels were deterministic all along. The input just included a hidden argument — everyone else's requests — that nobody told you about.

References

determinismbatch-invariancevllmgpu-kernelsinference