Dynamic batching, kernel switching, and the off-policy trap in RLVR

Setting temperature=0 and a fixed seed should make an LLM deterministic. It does not, and the standard “GPU hardware is inherently nondeterministic” explanation that the field carried for most of 2018 to 2024 is wrong. The actual cause is software: vLLM-style dynamic batching changes the batch size at runtime, the autotuner switches to a different reduction kernel, and the resulting reorder of floating-point summation produces bitwise different logits across runs. The drift is small per token, but autoregressive decoding amplifies it, and in RLHF/RLVR pipelines it cascades into a hidden off-policy trap where rollout and update use different effective policies on identical weights and reward collapses without warning. Built on Thinking Machines Lab’s Defeating Nondeterminism in LLM Inference (He, 2025), which laid out the diagnostic spine and supplied the figures attributed inline below.

Floating-point and the condition number of summation

Modern LLM training and inference run almost entirely in mixed precision, with BF16 (Brain Float, 1 sign / 8 exponent / 7 mantissa) the default for activations and parameters on Ampere-and-later GPUs. BF16 keeps the dynamic range of FP32 but slashes its precision: the unit roundoff is $\mathbf{u} = 2^{-7} \approx 8 \times 10^{-3}$, two orders of magnitude coarser than FP32’s $2^{-23}$.

Floating-point addition is non-associative at any finite precision, $(a + b) + c \ne a + (b + c)$ in general. Two different reduction orders therefore sit on different rounding paths and produce two different bitwise outputs, with the gap controlled jointly by the precision $\mathbf{u}$, the sum length $n$, and the cancellation among the summands.

The condition-number bound, in formal form

The natural quantity for tracking how much the cancellation hurts in a sum $S = \sum_{i=1}^{n} x_i$ is the condition number of summation

$\kappa_{\mathrm{sum}} = \frac{\sum_{i=1}^{n} |x_i|}{|\sum_{i=1}^{n} x_i|},$

which measures how much the magnitudes cancel. A standard worst-case bound for the accumulated error then reads

$\frac{|S - \hat{S}|}{|S|} \;\lesssim\; \underbrace{n}_{\text{seq len}} \cdot \underbrace{\mathbf{u}}_{\text{precision}} \cdot \underbrace{\kappa_{\mathrm{sum}}}_{\text{data-dependent}}$

(Higham, 2002). For an LLM the three factors are uncomfortable at once: $n$ runs into the thousands inside attention sums; $\mathbf{u}$ is BF16’s $8 \times 10^{-3}$; $\kappa_{\mathrm{sum}}$ exceeds $1$ wherever signed quantities cancel. The bound itself is invariant under reordering, but the realized error inside the bound is not; two reduction orders sit on different points inside the envelope, neither equal to the other.

That is the lever every later section pulls.

The puzzle: temperature 0 still drifts

Two recent measurements make the empirical situation concrete.

Atil et al. (2025) ran the same prompt ten times with temperature=0 and a fixed seed across a benchmark suite. On a college-math task they observed a max-min accuracy gap of up to 15% across the ten “identical” runs.

Yuan et al. (2025) localized one mechanism. In DeepSeek-R1, a long-context BF16 reasoning model, changing the batch size (eight versus thirty-two requests) is enough to flip individual tokens. A single early flip cascades through autoregressive decoding into a 9% accuracy drop and response-length variation of up to 9000 tokens. Long generation has the texture of a chaotic system: microscopic numerical drift in the early logits amplifies into macroscopic divergence of the trajectory.

A single-prompt experiment by Thinking Machines Lab makes the drift visible up close. Sampling 1000 completions of “Tell me about Richard Feynman” at temperature=0 from Qwen3-235B-A22B-Instruct produces 80 unique outputs, the most common appearing 78 times. The completions agree on the first 102 tokens; at the 103rd, 992 of them continue with “Queens, New York” and eight with “New York City”. Argmax decoding is sampler-deterministic by construction, so the drift lives in the logits themselves.

Where does the drift come from when seed, sampler, weights, prompt, and hardware are all fixed?

What is not the cause: concurrency and floating point

The story the field told itself for most of 2018 to 2024 is that GPU hardware is intrinsically nondeterministic because thread concurrency reorders accumulations and atomicAdd introduces races. Tri Dao, the FlashAttention author, gave a clean version of that diagnosis in 2022 (Dao, 2022, GH Issue #16). Asked about a numerical mismatch between FlashAttention and the reference attention in fp16, he attributed it to “the difference in the order of operations”, demonstrated with a one-line reductio that (x * 3.0) / 3.0 differs from x by the same magnitude.

That diagnosis is correct for what it covers and wrong as an explanation of LLM inference nondeterminism. Modern LLM forward kernels (matmul, attention) almost never use global atomicAdd; their summation order is fixed by the static thread-block layout, not by which thread happens to finish first. At a fixed batch size, the same input on the same hardware gives bitwise identical results across runs. The forward pass is run-to-run deterministic.

How the kernels achieve run-to-run determinism

The forward kernels use deterministic tree reductions inside a CUDA thread block: threads share fast L1 / shared memory and synchronize with __syncthreads(), so the in-block reduction tree is determined by the static block layout rather than by arrival order. The property is fragile in exactly one way: the block layout depends on input shape, so changing the batch size silently changes the tree.

The variability users observe is at a different level. The inference server is presented different effective inputs on different runs (because concurrent traffic varies the batch dimension), and the kernels respond differently to different effective inputs. What the kernels lack is not run-to-run determinism but invariance in the right property, and that property turns out to be batch size.

The actual mechanism

Floating-point summation is not bitwise-invariant under reordering. In production LLM inference, the autotuner reorders the sum based on batch size, and serving traffic varies the batch size. That coupling is the entire mechanism. The rest of this section gives the two concrete instances.

Load fluctuation. vLLM serves multiple concurrent requests by dynamically grouping them into batches; the number of in-flight requests $B$ varies at runtime with traffic, and so does the batch dimension presented to each kernel. The vLLM FAQ (vLLM Contributors, 2024) admits the effect and recommends switching to fp32, preferring fp16 over bf16, or fixing request seeds. All three attack the precision factor and leave the batch-invariance problem untouched.

Kernel non-invariance. Kernel libraries (cuBLAS, cuDNN, FlashAttention) ship multiple implementations of the same op, and an autotuner picks the fastest one for the input shape. The choice depends on $B$. When $B$ changes the chosen kernel changes, the reduction tree changes, and the bitwise output changes with it.

Concretely, take a reduction along the inner dimension $K$ inside a matmul, i.e., a long sum of $K$ products. At large $B$ the autotuner gives each core its own rows to sum, accumulating left-to-right. At small $B$ that strategy leaves cores idle, so the autotuner switches to Split-K: split the inner sum into chunks computed in parallel by different cores, then add the partial sums. Same input, different parenthesization, different bitwise output. The lever is not dynamic batching as such; it is that the kernel choice depends on a property the system does not pin. Dynamic batching is the trigger LLM inference happens to expose.

The two reduction trees, in summation form

• Large batch (high $B$). Data-parallel: each core handles multiple rows sequentially; summation order is $((x_1 + x_2) + x_3) + \cdots$.
• Small batch (low $B$). Split-K: the reduction dimension is split into chunks processed by different cores, then partial sums are added; summation order becomes $(x_1 + \cdots + x_k) + (x_{k+1} + \cdots + x_n)$.

Attention has a sharper version. The per-token reduction order in vLLM’s decoding attention kernel depends on how many other tokens from the same sequence are being simultaneously processed. The reason is that the kernel reduces over the KV cache separately from the current K/V values, and the boundary between the two moves with traffic.

Where the FlashAttention reduction breaks invariance

FlashAttention tiles attention into blocks and runs an online softmax that incrementally updates the local max $m$ and partial denominator $\ell$ (Dao et al., 2022). In decoding, vLLM’s Triton attention kernel reduces over the KV cache separately from the current K/V values. The correctness requirement is that the reduction order for a given token must not depend on how many other tokens from its sequence are being simultaneously processed. The default kernel violates that whenever a request mixes prefill and decode tokens, which is exactly when serving sees mixed traffic.

The mechanism is structural, not a kernel bug. Autotuners do their job; dynamic batching is a throughput optimization; the coupling between them is what generates the nondeterminism.

Why TPU + JAX sidesteps this

TPU systolic arrays demand fixed lock-step data flow, and JAX enforces that contract from above by running everything through jit, which compiles a program for a specific input shape. Practical pipelines therefore force shapes to be static. Google’s JetStream serving engine makes the contract concrete in code: jetstream/engine/token_utils.py defines DEFAULT_PREFILL_BUCKETS = [16, 32, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384, 32768]; an incoming sequence is rounded up by take_nearest_length to the next bucket, padded with pad_id tokens, and dispatched to a JAX function warmed up once per bucket via compile_prefill(length) (Google, 2024). The autotuner has nothing to autotune over. Determinism falls out of the contract; GPU + PyTorch + vLLM keeps shapes dynamic for higher throughput and pays the corresponding determinism cost.

The downstream consequence: silent off-policy in RLVR

Modern RLHF/RLVR pipelines split the workload across two engines for efficiency (Yao et al., 2025). The rollout runs in vLLM with small dynamic batches optimized for token throughput. The update runs in FSDP with large fixed batches optimized for compute efficiency. Both run on the same weights $\theta$.

By the previous section, the two engines run on different kernels, so identical $\theta$ produces different forward arithmetic, different logits, different action probabilities. Hence

$\pi_{\mathrm{rollout}}(a \mid s) \ne \pi_{\mathrm{train}}(a \mid s),$

with the difference small per token and large per trajectory. PPO and its descendants assume on-policy data; without correction the policy gradient estimate is biased and the KL between rollout and train policies spikes during training.

A Bigmath RLVR experiment (Thinking Machines Lab, 2025) with the policy initialized from Qwen-2.5-VL-Instruct 8B and max rollout length 4096 compares three configurations:

• No off-policy correction. $\mathrm{KL}(\pi_{\mathrm{rollout}} \,\|\, \pi_{\mathrm{train}})$ stays small until around step 318, then spikes; mean reward collapses in the same window.
• With truncated importance sampling. KL stays around $10^{-3}$ with occasional spikes; training proceeds smoothly.
• With batch-invariant kernels enabled. $\pi_{\mathrm{rollout}}$ is bitwise identical to $\pi_{\mathrm{train}}$, KL is flat at zero, and training is stable without any importance-weighting correction.

Same algorithm, same hyperparameters, same data; the only difference is whether the rollout kernel depends on batch size. This is the failure that gets misdiagnosed as a learning bug when it is a numerics bug.

Mitigations

Two routes, one mathematical and one systems-level.

Mathematical patch (Truncated Importance Sampling). Treat the mismatch as a deliberate off-policy setup and correct the gradient by importance sampling with the ratio truncated for variance control. This works while the rollout-train mismatch is small. When it is large (e.g., INT8-quantized rollout against BF16 train), the importance ratio variance explodes faster than truncation can damp, and learning fails anyway.

The TIS gradient estimator

$\nabla J(\theta) \;\approx\; \mathbb{E}_{a \sim \pi_{\mathrm{rollout}}}\!\left[\, \min\!\left( \frac{\pi_{\mathrm{train}}(a \mid s)}{\pi_{\mathrm{rollout}}(a \mid s)},\; C \right) \cdot A(s, a) \cdot \nabla \log \pi_{\mathrm{train}}(a \mid s) \,\right].$

The truncation cap $C$ bounds the per-sample weight: uncapped IS is unbiased but high-variance; clipping at $C$ trades a small bias for tractable variance.

Systems fix (batch-invariant kernels). Modify the three non-pointwise kernels in vLLM, namely RMSNorm, matrix multiplication, and attention, so that the forward computation does not depend on batch size. The TML benchmark (Qwen-3-8B serving 1000 sequences with output length 90 to 110 tokens) reports 42 seconds against vLLM’s default 26 seconds, about 1.6× with an optimized attention kernel and 2.1× without. Benefit: $\pi_{\mathrm{rollout}}$ is bitwise identical to $\pi_{\mathrm{train}}$ and the off-policy correction becomes unnecessary.

Per-kernel batch-invariance recipe

One fix per non-pointwise reduction:

• RMSNorm. Use a single data-parallel reduction strategy, sized for the smallest expected batch.
• Matrix multiplication. Use fixed tile sizes with no dynamic Split-K.
• Attention. Update the KV cache and page table before the attention kernel runs, so keys and values are laid out consistently regardless of how many tokens are being processed. Adopt a fixed split-size strategy for FlashDecoding: fix the size of each split rather than the number of splits, so the reduction tree for a given token does not depend on the total sequence length.

A fraction of throughput in exchange for keeping PPO on-policy.

What it amounts to

The mental model “GPU hardware is nondeterministic” misroutes debugging twice. It points at concurrency primitives that, for forward inference, are already deterministic; and it hides failures (15%-accuracy benchmark spreads, RLVR reward collapse at step 318) that look algorithmic but are numerical.

The decomposition this post has been using quietly throughout is worth stating explicitly:

System-level nondeterminism = (kernel non-invariance to a property) × (system-level fluctuation of that property).

For LLM inference the property is batch size. Two intervention points follow: pin the property in the system, as in the JAX-on-TPU bucketing route, or remove the kernel’s dependence on it via batch-invariant kernels. When training falls over without an obvious bug, the diagnostic move is to ask which property the forward depends on that you did not think you were varying. The candidate list is short: input shape, sharding layout, request order, autotune cache state. That is where autotuners hide.

References

1. Higham, N. J. (2002). Accuracy and Stability of Numerical Algorithms (2nd ed.). Society for Industrial and Applied Mathematics.
2. Dao, T., Fu, D. Y., Ermon, S., Rudra, A., and Ré, C. (2022). FlashAttention: fast and memory-efficient exact attention with IO-awareness. arXiv:2205.14135.
3. Dao, T. (2022). Numerical stability concerns (Issue #16). https://github.com/Dao-AILab/flash-attention/issues/16.
4. Google (2024). JetStream: a throughput- and memory-optimized engine for LLM inference on XLA devices. https://github.com/AI-Hypercomputer/JetStream.
5. vLLM Contributors (2024). Frequently asked questions. vLLM documentation.
6. Atil, B., Aykent, S., Chittams, A., Fu, L., Passonneau, R. J., Radcliffe, E., Rajagopal, G. R., Sloan, A., Tudrej, T., Ture, F., Wu, Z., Xu, L., and Baldwin, B. (2025). Non-determinism of “deterministic” LLM settings. arXiv:2408.04667.
7. He, H., and Thinking Machines Lab (2025). Defeating nondeterminism in LLM inference. Thinking Machines Lab: Connectionism.
8. Yuan, J., Li, H., Ding, X., Xie, W., Li, Y.-J., Zhao, W., Wan, K., Shi, J., Hu, X., and Liu, Z. (2025). Understanding and mitigating numerical sources of nondeterminism in LLM inference. arXiv:2506.09501.
9. Yao, F., Liu, L., Zhang, D., Dong, C., Shang, J., and Gao, J. (2025). Your efficient RL framework secretly brings you off-policy RL training. Notion blog.