Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

D ETENT: BATCH -I NVARIANT LLM S ERVING AND RL ROLLOUTS AT N EAR -S TOCK S PEED

A BSTRACT Reproducible evaluation, verifiable inference and on-policy reinforcement learning (RL) need an LLM engine to give a request the same output bits whichever requests share its batch. Batch-invariant kernels exist, but engines still select each kernel and its parameters from the batch. Two kernels that compute the same sum can round it differently. The engines’ deterministic modes remove the choice by keeping one conservative kernel per operator and pay for it in latency. Many fast kernels are batch-invariant by themselves, so it is enough to fix the few batch-dependent values that kernel selection reads. D ETENT expresses kernel selection as a program that the engine executes. It finds the fewest values in this program that must be fixed and tunes them for the workload before deployment. On multi-GPU deployments of Llama-3.1-70B and Qwen3-30B-A3B in vLLM and SGLang, D ETENT is batch-invariant. Tuned per load, it serves at 1.04–1.17× the latency per output token of Stock (each engine with its standard settings), against 1.67–10.3× for the deterministic modes. In SGLang it matches the GSM8K accuracy of SGLang’s deterministic mode. D ETENT covers multiple deployment combinations, including multiple engines, GPU generations, tensor-parallel degrees, and dense and MoE models. Seeded RL rollouts of models up to Llama-3.1-70B replay bit for bit under other batches and cache splits. A trainer on D ETENT’s kernels and values matches every rollout logprob bit for bit at 1.03× Stock’s rollout time per token. BF16 GRPO then stays stable where Stock’s run collapses.

SGLang (Zheng et al., 2024) removes the choice in its deterministic mode (SGLang-Det; The SGLang Team, 2025). vLLM (Kwon et al., 2023) does the same in its batchinvariant mode (vLLM-Det; The vLLM Team, 2026). Both 1 Independent Researcher, Boston University, Boston, MA, USA Google, Mountain View, CA, USA 3 RadixArk, Palo Alto, CA, USA 4 Princeton University, Princeton, NJ, USA 5 Georgia Institute of Technology, Atlanta, GA, USA.

0

ShareGPT

arXiv

arXiv

(closed loop)

ShareGPT

arXiv

1.06

1.67

3.76

10.31

3.42

9.75

1×

1.04

1.17

1.09

2.72

SGLang-Det vLLM-Det

1.06

3.48

2.56

(b)

1.13

1.08

6.01

4.86

Detent in SGLang Detent in vLLM

1.13

1.16

1.12

2.02

2×

2.83

(a) 3×

1.14

A request is batch-invariant when its output bits do not depend on the other requests in its batch. Engines with their standard settings, which we call Stock, are not. Changing only the load changes Stock SGLang’s answer on 33 of GSM8K’s 1,319 problems (§8.4). A trainer that recomputes its rollouts matches only 1.0–47.7% of their tokens bit for bit. This mismatch alone can collapse RL training (Zhong et al., 2026; Figure 12a,c). Evaluations are not repeatable (Yuan et al., 2025), and third parties cannot verify inference by re-execution (Alves et al., 2026). Prior work makes the reduction order inside each kernel independent of the batch (He & Thinking Machines Lab, 2025; Zhang et al., 2026). The engine, however, still selects each kernel and its parameters from batch-dependent quantities such as the batch’s token count. Two kernels that compute the same sum can round it differently.

2

Stock (1×) (not batch-invariant)

I NTRODUCTION

1.09

1

p50 latency per output token over the same engine’s Stock

arXiv:2611.10611v1 [cs.LG] 7 Nov 2026

Peixu Xin 1 Runchuan Feng 1 Jiazhen Guo 2 Kan Wu 3 Ming Yin 4 5

arXiv

(closed loop)

Figure 1. ShareGPT replays multi-turn chats. Their prompts are short and mostly hit the prefix cache, so most of the work is decode. arXiv sends 8k-token documents with Poisson arrivals and long prefills. Both run at loads up to Stock’s peak. Closed loop sends the arXiv documents one or eight at a time, without queueing (§8.1). (a) Qwen3-30B-A3B; (b) Llama-3.1-70B. Bars: geometric mean over a workload’s loads, with D ETENT tuned per load (§6).

modes keep one conservative kernel per operator: unsplit attention, NCCL’s tree all-reduce and one hard-coded MoE configuration. They are batch-invariant, but their serving latency per output token is 1.59–6.34× (SGLang-Det) and 2.83–10.3× (vLLM-Det) Stock’s per load (Figures 1, 7 and 9). Our observation is that many fast kernels are batch-invariant

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

by themselves. For those kernels the batch reaches a request’s bits only through kernel selection, so we only need to fix the batch-dependent values that kernel selection reads. Writing BI for batch invariance, kernel | {z BI} =⇒ end-to-end BI, | {z BI} ∧ selection

user-supplied

(1)

D ETENT

provided each request’s output is mathematically a function of its own input (§3). We present D ETENT. It keeps the fast kernels and makes kernel selection batch-invariant. For each deployment and engine version, D ETENT decides which values in kernel selection to fix. We call them pins. For each workload, it picks the pins’ values and a few deployment settings, because the fastest choice depends on the traffic. These values form a commitment tuple, fixed while serving. Both decisions are hard to make by hand. The source does not show which kernels agree bit for bit, and there are too many configurations to measure (§2). D ETENT automates both decisions and integrates into existing engines such as vLLM and SGLang (Figure 4, §7). We write kernel selection in SelIR (Selection IR), an executable DSL. The engine runs the user’s selection program as its selector (§4). A solver finds the fewest values in that program to pin so that no batch-dependent input reaches the selection. It also finds the values each pin may legally take (§5). A tuner picks the commitment tuple for the workload (§6). The user supplies batch-invariant kernels. On multi-GPU deployments of Llama-3.1-70B and Qwen330B-A3B in vLLM and SGLang, D ETENT is batch-invariant as long as the commitment tuple stays the same. Tuned per workload, D ETENT serves within 12% of Stock’s p50 latency at the geometric mean and is 2.4–5.7× faster than the engines’ deterministic modes (§8.3). In SGLang it is as accurate on GSM8K as SGLang-Det (§8.4). On multiple types of workloads, one tuple for all loads of a type costs 1.1% (Llama-3.1-70B) and 2.7% (Qwen3-30B-A3B) more. The tuple of another type can take 4.17× as long. With seeded sampling committed, a trainer on the serving kernels and tuple matches every rollout logprob bit for bit on multiple models. Rollout takes 1.03× Stock’s time per token, and a whole RL step up to 1.18×. Seeded rollouts of Qwen3-30B-A3B and Llama-3.1-70B on four GPUs also replay bit for bit under other batches and cache splits, and they match their rescoring. On the BF16 recipe of Qi et al. (2025), Stock’s GRPO (Shao et al., 2024) run collapses while D ETENT’s run stays stable (§8.6). We make the following contributions: 1. We give a sufficient condition for end-to-end batch invariance. With batch-invariant kernels, any SelIR selection is batch-invariant on the deployment and engine version

solved for once the values that separate it from the batch are pinned to legal values (§3, §5). 2. We design SelIR, a selection language with a JIT cache. The solver analyzes its programs and the engine runs them (§4). 3. We build a solver that finds the fewest such pins and their legal values. A tuner commits a tuple per workload in minutes, and the tuple serves at near-stock speed (§5, §6). 4. We implement D ETENT in vLLM and SGLang. We evaluate it across both engines, two GPU generations, dense and MoE models, and tensor-parallel (TP) degrees (§7, §8). 5. We show that a trainer integrated with D ETENT achieves on-policy RL training and matches RL rollouts bit for bit at near-stock speed (§8.6).

2

BACKGROUND AND M OTIVATION

Observation 1: selection reads the batch. Floating-point addition is not associative, so a kernel’s bits depend on its reduction order. A request’s output is batch-invariant when its bits do not depend on the batch it runs in (He & Thinking Machines Lab, 2025). The batch gets in by several routes (Figure 2). Kernel selection picks each kernel and its parameters. Each pick fixes how a request’s values are grouped and added. FlashAttention-3 (FA3; Shah et al., 2024) sets its split count (Figure 2b) and tile width (Figure 3a) from the batch. The fused MoE takes the row of its tuned configuration table nearest the batch’s token count, and the all-reduce picks an algorithm by the bytes the batch sends (Figure 3a). Across steps, the prefix cache and the scheduler divide a request’s tokens between prefill and decode. If a request prefills 100 tokens and then decodes 50, its keys are grouped differently than if it prefilled all 150 (Figure 2a). An RL rollout prefills and then decodes, while training prefills everything. Batch-invariant kernels fix the reduction order once these choices are made, but not the choices themselves (Figure 2c). Observation 2: agreement must be measured. SGLangDet and vLLM-Det fix these choices. SGLang-Det’s decode step then takes 1.36–2.00× Stock’s on five multi-GPU deployments (Figure 3d). Keeping faster kernels instead requires knowing which of them agree bit for bit. Only measurement shows that (Figure 3b). At head dimension 64, two tile widths (Figure 3b) disagree on 31 of 32 logprobs. On Qwen3-30B-A3B, by contrast, the 18 tuned MoE configurations agree bit for bit with the one SGLang-Det hard-codes. Nothing in the source tells these two cases apart, and the answer changes with the deployment.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed (a)

(b)

RL rollout: prefill, then decode prefill 100

alone: 4 splits

decode 50

0 1 SM1

̸= prefill 150

in a busy batch: 1 split

0

3 2 SM2

1

0–7

0–5

0–3 combine 0–1 2–3

6–7

5 4 SM3

6 7 SM4

SM1 3 4

2

5

6

(c)

kernel selection fixed by D ETENT

4–5

batch

̸=

batch-invariant kernel supplied by the user

7

RL training: one prefill

output bits

0–7

Figure 2. The batch regroups a request’s additions, so its bits change. (a) Where prefill ends. (b) Kernel selection sets the split count; each split sums its KV blocks, and a combine adds the sums in order. (c) D ETENT fixes selection; the user supplies kernels. (a)

FA3 tile (hd 64) 1

longest query (tokens) 128-wide 256 4K total KV (tokens) 1 256K 4M batch tokens

192-wide

16 FA3 splits (1K ctx) 8 4 3 2 1K 16K MoE row (1 of 18) 1

16

All-reduce (AR, TP=4) one-shot 16 KB 256 KB

256 two-shot 4 MB

4K message bytes NCCL 64 MB

(b)

(c)

(d)

FA3, hd 64: 128- vs 192-wide tile 31 of 32 logprobs change bits

MoE, 18 rows: tuned vs fixed

1M 737,280

10K

The source shows neither

1K

SGLang-Det

FA3

splits by batch

one split

332,464

MoE

18-row table

one config

AR

3 algorithms

NCCL tree

GEMM cuBLAS by M

DeepGEMM

100K

18 of 18 bit-identical

Stock

measuring these once takes about 4.8 days

18,472 3,280 template compiled product kernels

SGLang-Det’s decode step takes with splits

hd128, BF16

1.36–2.00× Stock’s

on five multi-GPU deployments

Figure 3. Kernel selection reads the batch, and only measurement shows which kernels agree. (a) Where four choices switch (log axes; H100; tile on Llama-3.2-1B, MoE row on Qwen3-30B-A3B, splits and all-reduce on Llama-3.1-70B under CUDA graphs). (b) Agreement is measured. (c) Too many FA3 configurations on H100 to measure (log scale). (d) Fixing choices is slow.

Observation 3: too many configurations to measure. Measuring each of FA3’s 18,472 configurations once would take about 4.8 days of a four-GPU node (Figure 3c; Appendix B). D ETENT fixes what kernel selection reads and never relies on two kernels agreeing (Observations 1 and 2). It finds what to fix by analyzing the selection program (Observation 3).

3

OVERVIEW

The condition.

such as FA3’s split count. An assignment gives every pin a value. It is legal if every precondition holds on the inputs the solver judges. A knob is a deployment setting that no input determines, such as the all-reduce workspace (push or pull) or the padding of the matrix multiplication (GEMM). The tuner then picks a commitment tuple: a legal value for every pin and a setting for every knob. For a workload, the commitment tuple is fixed before deployment and does not change while serving, so the batch’s composition never affects kernel selection.

A request’s output is y = F (x, s; M, D, T ; B),

(2)

with prompt x, sampling seed s, model code and weights M , deployment D (engine and its build, GPU, parallelism, . . . ), commitment tuple T and batch state B (co-running requests, row, prefix-cache hit, . . . ). We call it batch-invariant when F (x, s; M, D, T ; B) = F (x, s; M, D, T ; B ′ ) for all B, B ′ . CoRun states a similar contract (Zhao et al., 2026). Kernel batch invariance: each candidate kernel, given its parameters, reduces a request’s values in an order that depends only on that request. Selection batch invariance: the values a kernel selection returns, its selection points, depend only on M , D, T and the request. Like prior work, we assume each request’s output is mathematically a function of its own input. A kernel family is the set of kernels one SelIR program selects among. A kernel with its parameters fixed, such as one FA3 tile width and split count, counts as one kernel. For each family, the solver returns pins. A pin is a value of the program that D ETENT replaces with a constant,

The pipeline. The solver finds the pins and their legal values in a SelIR program once per deployment and engine version (§5). The tuner picks a commitment tuple per workload (§6; Figure 4). During serving, the same SelIR program selects kernels with the tuple’s values pinned. The pins block every path from the batch to a selection point (§5), so selection batch invariance holds by construction. Kernel batch invariance is the user’s responsibility (§7). An RL trainer that runs the same kernels under the same tuple differs from the server only in the batch, so it matches the rollout’s logprobs. Admitting only seeded sampling also fixes the sampled tokens (§8.6).

4

S EL IR

Each kernel call site uses a SelIR program in place of the engine’s selection code. Because the solver analyzes this same program, pinning a value in SelIR also pins it in the running engine.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

written by the user SelIR programs (selection, knobs, require, assume) deployment facts

1 Compile

2 Tune

3 Serve

minimum vertex cut of the optimized program: the pins

tuner: profiles a workload close to the real one and picks

patched engine: calls the compiled selector for each kernel family (FA3, MoE, all-reduce) batch-dependent inputs

VRA: the pins’ legal values, found by running the program

JIT cache: the pin-marked AST and the legal values; a hit skips solving

commitment tuple: a legal value for every pin, a setting for every knob

compiled selector: invokes the kernel with the pins as constants; each new input is checked once against require/assume (fail: exit) pass batch-invariant kernels

JIT cache: the executable selector, the pin-marked AST pinned by the tuple

batch-invariant result

Figure 4. D ETENT isolates kernel selection from the batch.

The language. SelIR is a Python-subset DSL with Python’s semantics. Its programs are pure functions over records and scalars, with conditionals and reductions over static collections. Batch-dependent inputs are marked Runtime, or Unpinnable if they cannot be pinned. Each value the program returns is a selection point. knob declares a deployment setting (§3). require and assume together state when the program’s selection is valid. The user never names pins or their candidate values. A program may transcribe an engine’s selection or state a new policy for the user’s kernels. Our programs transcribe an engine’s selection, then adapt its heuristics to the kernels D ETENT selects among. For example, the FA3 program adds a split stride (§7). The engine’s heuristic makes no such choice. Both engines we evaluate run these programs. Selections for other kernels and engines are written the same way. Optimization and lowering. Most selection inputs are known before serving: the model’s head dimension and head counts, the data type, the GPU architecture, the degree of parallelism, the page size, the maximum context, etc. A few change at run time, such as the batch’s token count and KV lengths. We therefore partially evaluate each program for its deployment (Jones et al., 1993). Compiler optimizations lower it to a single dependence graph. Its edges capture every dependence on the batch, including control dependences. The graph’s values are simplified as far as the deployment facts allow. With constant folding alone, FA3’s cut has 12 pins instead of 3 (§8.2). 4.1

Compiling and Caching Selections

As with TorchDynamo’s code cache (Ansel et al., 2024), TVM’s tuning records (Chen et al., 2018) and TensorRT’s engines (NVIDIA Corporation, 2026), we compile a selection once and reuse it (Figure 4). A solve is stored on disk under a hash of everything that can change its answer. A hit in this JIT cache skips solving (§5). For FA3, a hit takes about 20 ms and a cold solve of one committed tuple about 5 s. The compiled selector has the pins as constants. While serving, each new combination of Runtime inputs

9 batch-dependent inputs Unpinnable

require, assume

num_tokens

row key

8,076 values deployment facts and a knob

14 via the row key

3 pins + 23 not batch-dependent

align, combine

26 selection points

Figure 5. The solver separates kernel selection from the batch with a few pins. Left, the whole selection of FA3 on Llama3.1-70B; right, the MoE’s on Qwen3-30B-A3B. In FA3, red inputs are Runtime or computed at launch, purple ones other Unpinnable inputs, each value is tinted by the inputs it derives from, and gray values derive only from deployment facts and a knob.

is checked once against every require and assume. The compiled selector runs at every call. Unlike a failed shape guard in torch.compile, a new batch never triggers a recompile.

5

S OLVER

The solver decides which values of a SelIR program to pin and which values each pin may take. It runs once per deployment and engine version. The cut. After optimization, a selection is a dependence graph from its inputs to its selection points. For FA3 on Llama-3.1-70B, the graph has 8,076 values and 26 selection points, 3 of them batch-dependent (Figure 5). The batch must not reach any selection point. Pinning a value turns it into a constant, which no path crosses. The solver therefore pins a vertex cut, a set of values that blocks every path from a batch-dependent input to a selection point. Pinning every

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed family

candidates

legal

measured

FA3

7,168

2,976

48

MoE

76

36

20

AR

204

3

3

GEMM

5

5

5

one family at a time (sum)

3,020

76

all families jointly (product)

1,607,040

. . . any other family written in SelIR

Figure 6. The tuner searches one kernel family at a time. A case on Qwen3-30B-A3B: candidates multiply a family’s pin and knob values, VRA keeps the legal ones, and the tuner measured 76 of the 3,020 legal assignments.

Q legal tuples are L = f Lf . We make three assumptions. (i) Kernels of different families run back to back on one stream without overlap. (ii) Because the recorded workload fixes batching and outputs, a tuple changes only which kernels each family selects. (iii) The time tf (c, af ) of a call c depends only on the call and its family’s assignment. Then, with C the time of all other work, the workload’s time under a is X X T (a) = C + Tf (af ), Tf (af ) = n · tf (c, af ). f

(c,n)∈Wf

(3) Hence
arg min T (a) = arg min Tf (af ) f . af ∈Lf

a∈L

batch-dependent selection point always gives such a cut, but each pin adds a dimension to the tuner’s search space. The solver instead takes a weighted minimum cut via maximum flow (Edmonds & Karp, 1972). The weights favor the fewest pins and, among cuts of equal size, pins near the selection points. On FA3 the cut falls on batch-dependent selection points. On the MoE it falls inside the program, on the row key (the token count of the tuned configuration row it uses). This pin fixes 14 of the MoE’s 16 batch-dependent selection points. Legal values. Pin Value Range Analysis (VRA) finds the values each pin may take. A pin’s candidates are the values it can take in the program. VRA runs the program with an assignment of candidates substituted. It judges the assignment legal if every require and assume holds at every input point it tries. It tries all points when the input domains are small and a fixed sample otherwise. While serving, the selector checks these conditions on each new input. If one fails, the engine stops with an error and runs no other kernel.

6

T UNER

The fastest tuple depends on the workload, for example on its prefix-hit rate, concurrency and prompt lengths (§8.3). The tuner picks a commitment tuple for each workload. Profile. For each workload, we tune a commitment tuple once using a profiling workload with similar arrivals. The target workload then runs with the fixed tuple. Tuning first serves the profiling workload with the default tuple. It gives each pin the legal value the unpinned program takes most often and each knob its default. For each kernel family f , the tuner records in Wf the family’s distinct calls c with their batch-dependent quantities and counts n. Additivity. A tuple a = (af )f gives each family f a legal assignment af ∈ Lf . Since legality is judged per family, the

∗ Proof: the per-family minimizers a tuple a∗ ∈ L. P af form ∗ Every a ∈ L has T (a) ≥P C + f Tf (af ) = T (a∗ ). The tuner Q therefore searches f |Lf | assignments instead of f |Lf |. On Qwen3-30B-A3B, it searches 3,020 instead of 1,607,040 (Figure 6). If the sum holds only within a relative error δ, i.e. |TP (a) − T̂ (a)| ≤ δ T (a) for every legal a with T̂ (a) = C + f Tf (af ), then for every legal a ∗

(a ) (a) 1+δ T (a∗ ) ≤ T̂1−δ ≤ T̂1−δ ≤ 1−δ T (a).

On Qwen3-30B-A3B, the largest interaction between families was δ = 1.04% of the measured GPU time per output token. The per-family choice is then within 2.1% of the best tuple. The assumptions fail if an engine overlaps two families (e.g., kernel fusions) or if faster kernels change how arriving requests are batched. The user must therefore declare kernels that overlap as one family. Search. The tuner searches each family’s legal grid and measures only a small part of it. The grid’s axes are the family’s pins and the knobs that only its selection reads. Each candidate is a repin that swaps the pins’ constants and regenerates the selector without lowering or solving again. On Llama-3.1-70B’s exhaustive grids, which were not used to design the tuner, the tuner measures 8.0% of the candidates and ends 0.04% from the optimum on average (0.75% at worst).

7

I MPLEMENTATION

D ETENT’s solver with VRA, selector runtime, tuner and engine integrations are implemented in Python. So are the SelIR programs for FA3, fused MoE, all-reduce and GEMM. The all-reduce kernels and FA3 modifications are implemented in CUDA. In each engine, a patch routes kernel selection at the FA3, MoE and all-reduce call sites through the selector. The rest of D ETENT runs unchanged. We run SGLang (a 2026 development version) and vLLM 0.26.0. For RL, a second patch outside SelIR commits seeded sampling. An RL trainer uses the serving kernels.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed Table 1. Arms (§8.1). BI: batch-invariant. Stock: each engine as shipped, deterministic mode off. SGLang-Det, vLLM-Det: the engines’ deterministic modes (Appendix A). MoE: Qwen3-30BA3B only. No arm uses FlashInfer’s all-reduce fusion; Appendix C ablates it and the matrix-multiplication and all-reduce kernels. Arm

BI Kernel selection

Matrix mult. Attention split

All-reduce

Stock

× stock heuristics

cuBLAS

SGLang-Det vLLM-Det

✓ one kernel per operator ✓ SelIR, pinned

DeepGEMM none (Triton in vLLM) DeepGEMM fixed (tuned pad) stride or none

custom, by nearest the batch’s the batch bytes NCCL tree fixed

D ETENT

by the batch

ours, pinned

MoE row key

pinned

Kernels. D ETENT uses the same kernels in both engines: DeepGEMM (Zhao et al., 2025) for matrix multiplication, Triton fused-MoE kernels, our all-reduce kernels and a rebuilt FA3. The rebuilt FA3 keeps split-KV and follows the fixed split-size principle (He & Thinking Machines Lab, 2025), so a request’s split boundaries do not depend on the batch. Our all-reduce kernels sum each element in a fixed rank order. The tuner also picks two knobs like pins: FA3’s split stride and the GEMM padding.

8

E VALUATION

We ask what solving and tuning cost (§8.2), how close D ETENT’s batch-invariant serving comes to Stock in vLLM and SGLang (§8.3), whether it keeps accuracy (§8.4), how it compares with LLM-42 (§8.5), whether a trainer matches its rollouts (§8.6) and how far the solver carries to other deployments (§8.7). We ablate the solver’s optimizations (§8.2), the tuner (§8.3, §8.7) and our kernels (Appendix C). 8.1

decode. W-B sends 200 arXiv documents (Cohan et al., 2018) of about 8k tokens with Poisson arrivals. It mixes long prefills with decode under load. W-B-c sends the first 64 of them in a closed loop to measure latency without queueing. W-A and W-B can queue. Running each of these three workloads at several loads gives nine workloads. A1–A4 run W-A at 0.25–1.0λmax , B1–B3 run W-B at 0.33– 1.0λmax , and C1, C2 run W-B-c at concurrency 1 and 8. Here λmax is Stock’s maximum offline throughput in that engine. In vLLM we serve A4, B3, C1 and C2. On Qwen330B-A3B, C1 and C2 there send all 200 documents. W-B and W-B-c run with a 32k-token context limit. A workload tuple that is legal only at 32k serves all nine workloads at 32k. Every Llama-3.1-70B arm runs GSM8K at 32k. Tuples and metrics. D ETENT serves three kinds of tuple: the default tuple, a tuned tuple for each workload A1–C2 (§6), and a workload tuple for each of W-A, W-B and W-B-c. Each stays fixed while serving. A workload tuple has the lowest mean cost over that workload’s loads in an exhaustive table that predicts every legal tuple’s GPU time on the nine workloads. We report the p50, p90 and p99 of a request’s latency per output token as the median of three runs. A ratio is a latency over the same engine’s Stock, at the geometric mean over the workloads A1–C2 that the engine serves. We test batch invariance with a battery of 26 cells. Each cell runs a request at temperature 0 in a reference batch and again in the same or a changed batch. It compares the bits of the request’s first 16 output logprobs. This also catches a change in the logits that leaves the tokens alone. Two cells rerun the same batch to check reproducibility, fourteen change the co-batched requests, four move the request to another row, and six change how much of its prompt hits the prefix cache. Nine further strong cells place the request next to 1 to 32 short, long or mixed requests and move it within the batch.

Setup

Deployments and arms. The main deployments run Llama-3.1-70B-Instruct (Grattafiori et al., 2024) and Qwen330B-A3B (MoE; Yang et al., 2025). Both have head dimension 128 and are served by vLLM and SGLang in bf16 with tensor parallelism (TP) over four H100 GPUs. D ETENT runs unchanged in both engines. Only the deployment facts in its specs change (in vLLM, 16-token pages, its CUDAgraph capture sizes and its batch’s range). The solver finds the same pins. In each engine, all arms (Table 1) run the same version (§7) with FA3, CUDA graphs, prefix caching, chunked prefill and greedy decoding. Workloads. W-A replays 600 requests from about 160 ShareGPT conversations (anon8231489123, 2023) turn by turn. Its prompts are short, and about 64% of their tokens repeat the previous turn. W-A exercises prefix caching and

8.2

Solving and Tuning Costs

Ablating the solver’s optimizations. On the main deployments the solver finds 3+3+0 pins (attention + all-reduce + MoE) on Llama-3.1-70B and 3+3+3 on Qwen3-30B-A3B. Both engines get the same pins. Each selection with pins needs a different part of the analysis. With constant folding alone, FA3’s cut has 12 pins instead of 3. Without a cut, pinning every value the MoE’s batch-dependent selection points read takes 10 pins instead of 3. Each extra pin multiplies the tuner’s search space by its number of candidates. Without the analysis, FA3’s selection has 227 candidate assignments (Appendix B). Without one of the simplifications, all but one of the all-reduce’s legal assignments bypass the selection’s own tuned table. The GEMM selection has no cut. The batch reaches neither of its two selection points, the kernel and the padding.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

Solving and tuning time. Once per deployment, judging FA3’s 7,168 candidate assignments takes 130–180 s on 32 cores. Repinning takes about 4 ms. With a five-minute budget per family, tuning one workload took 605–1,148 s on Qwen3-30B-A3B and 690–1,006 s on Llama-3.1-70B. This includes worker startup. For some workloads we also tuned with 10- or 30-minute budgets and kept the fastest tuple. 8.3

Invariance and Serving

Invariance. D ETENT, SGLang-Det and vLLM-Det are batch-invariant on both models. Stock fails 9 (Llama-3.170B) and 5 (Qwen3-30B-A3B) of 26 cells in SGLang, and 19 and 2 of the 24 and 23 cells that apply in vLLM. With 16-token pages, the request’s short prompt gets no prefix hit in the remaining cache-split cells. D ETENT and the deterministic modes pass all nine strong cells. Stock fails 4 to 7 of them. Every KV-splitting tuple we served and ran the battery on also passes it. Latency. Tuned per workload, D ETENT’s p50 is 1.10× Stock’s on Llama-3.1-70B and 1.11× on Qwen3-30B-A3B in SGLang (Figure 7). In vLLM it is 1.08× and 1.11× (Figure 9). Across workloads its p50 is 1.01–1.22× and its geometric-mean p99 1.12–1.22×. At the geometric mean, SGLang-Det is 2.4–2.6× slower than D ETENT at p50 and 2.8–3.0× at p99. vLLM-Det is 3.4–5.7× slower at p50 and 3.7–9.1× at p99. On one workload their p50 reaches 6.1× and 9.2× D ETENT’s. Both values come from runs whose queue keeps growing. We also serve each arm at 0.67 to 1.0 of its own maximum throughput so that its queue does not grow. vLLM-Det on Llama-3.1-70B’s W-A and SGLangDet on Qwen3-30B-A3B’s W-B then still take 2.0–2.3× D ETENT’s p50. Ablating the tuner. Without tuning, the default tuple is up to 103% (Llama-3.1-70B) and 1–233% (Qwen3-30B-A3B) slower than the tuned tuple in SGLang (Figure 7a,b). In vLLM the gaps are up to 18% and 72%. At p99 they reach 120% and 273% in SGLang and 31% and 91% in vLLM. The tuples choose from the same kernels. Tuning alone therefore accounts for 1.34× (Llama-3.1-70B) and 1.67× (Qwen3-30B-A3B) of SGLang-Det’s gap at the geometric mean, and 1.13× and 1.31× of vLLM-Det’s. Two things explain part of the remaining gap to Stock. First, the rebuilt FA3 adds 3 µs per call, and its best split stride is within 1.06× of Stock’s split. Second, one tuple serves long prefills and small decode batches. Ablating our kernels. Swapped into Stock, D ETENT’s GEMM and all-reduce kernels add 0.6% (Qwen3-30B-A3B) and 2.8% (Llama-3.1-70B) GPU time per output token. So they are not why D ETENT is close to Stock. With Stock’s FlashInfer all-reduce fusion on, D ETENT’s decode step takes

0.996–1.103× Stock’s. Turning the fusion off does not weaken Stock (Appendix C). Workload tuples. Stock adapts its selection to every batch. A commitment tuple forbids this. We adapt per workload instead. Figure 8 shows what one tuple for all loads of WA, W-B or W-B-c costs. On its own workload, the tuple is within 2.7% (Qwen3-30B-A3B) and 1.1% (Llama-3.170B) of the tuned tuples at the geometric mean. Another workload’s tuple can be much slower. On Qwen3-30BA3B, the W-A tuple at B2 takes 4.17× as long as the tuned tuple and raises the median time to first token from 238 ms to 4.75 s. On Llama-3.1-70B that cell takes 1.15×. Offdiagonal cells are what traffic drift costs. If traffic moves and the tuple stays, latency moves to that cell. 8.4

Accuracy

Table 2. GSM8K accuracy (%) on the 1,319 test problems, four comparisons. At G2 and G3, SGLang-Det’s and D ETENT’s outputs are identical, so the two load points form one comparison and only Stock has two accuracies (G2, G3). ∆ is D ETENT minus SGLangDet in percentage points with its 95% interval, and p the p-value of the exact McNemar test. Wilson 95% intervals on the accuracies are about ±1.3 percentage points. Point

Stock

SGLang-Det D ETENT

∆ [95%]

p

Qwen3-30B-A3B G1 93.71 G2, G3 93.93, 94.01

93.63 93.63

93.48 94.16

−0.15 [−0.92, +0.66] 0.85 +0.53 [−0.15, +1.02] 0.14

Llama-3.1-70B G1 94.69 G2, G3 94.92, 94.62

94.69 94.69

94.77 94.84

+0.08 [−0.39, +0.50] 1.00 +0.15 [−0.36, +0.57] 0.75

D ETENT runs different kernels from SGLang-Det, so its outputs can differ from SGLang-Det’s. They should be as accurate and should not depend on the load. We serve GSM8K’s 1,319 test problems (Cobbe et al., 2021) five-shot and greedily in SGLang at three load points: concurrency 8 (G1), concurrency 64 (G2) and all problems at once (G3, an overload point). Each problem meets different batches at the three points. An engine whose output depends on the batch can answer it differently. We tune D ETENT on the training split at each point. Its G1 tuple splits the KV, and its G2 and G3 tuples do not. vLLM runs the same D ETENT, so we did not repeat these runs in vLLM. D ETENT is as accurate as SGLang-Det (Table 2, Figure 10a). D ETENT’s accuracy minus SGLang-Det’s is −0.15 to +0.53 percentage points (95% intervals on the difference within [−0.92, +1.02]). The exact McNemar test (McNemar, 1947) gives p ≥ 0.14. The intervals on the accuracies in Figure 10a are 95% Wilson intervals (Wilson, 1927). On the 1,209 corrected problems of GSM8K-Platinum (Vendrow et al., 2025), accuracy is 96.11–

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed Stock

Detent, default

(a) 6× 5× 4× 3× 2× 1×

latency / Stock

1.5× 1×

(c)

SGLang-Det

p50

p99 2×

5×

4×

4×

3×

2×

1.5×

3× 2×

1×

W-A

p90

(b)

3×

2×

1.4×

Detent, tuned

14.1×

2×

1×

W-B

W-B-c

1.3×

1×

1×

(d)

W-A

W-B

W-B-c

1.2× 1.2× 1.1× 1.0×

1.0× A1

A2

A3

A4

B1

B2

B3

C1

A1

C2

A2

A3

A4

B1

B2

B3

C1

C2

W-A tuple

(a) A1 1.02

A2 1.01

A3 1.00

A4 0.99

B1 1.83

B2 4.17

B3 2.34

C1 1.03

C2 1.26

(b) A1 0.97

A2 1.02

A3 1.02

A4 1.03

B1 1.03

B2 1.15

B3 1.22

C1 1.00

C2 1.05

W-B tuple

1.24

1.19

1.05

1.07

1.00

1.02

1.17

1.19

1.04

1.03

1.04

1.03

1.03

1.03

1.00

1.01

1.08

1.00

W-B-c tuple

1.12

1.09

0.96

1.08

0.95

1.19

1.37

1.02

1.01

0.97

1.00

1.01

1.10

1.01

1.05

1.09

1.01

1.02

4× 3× 2× 1.5× 1×

p50 / tuned

Figure 7. Latency per output token over Stock’s in SGLang, median of three runs. (a, c) Qwen3-30B-A3B, (b, d) Llama-3.1-70B. A1–A4 (W-A), B1–B3 (W-B): loads up to Stock’s peak. C1, C2 (W-B-c): one and eight in flight. (a, b) p50, SGLang-Det also p99, each workload on its own scale. (c, d) Tuned D ETENT at p50, p90 and p99: 1.01–1.22× at p50, 1.04–1.34× at p90 and p99.

Figure 8. Rows: the workload tuples of W-A, W-B and W-B-c. Columns: the loads. Cells: the row’s p50 over the p50 of the column’s tuned tuple. (a) Qwen3-30B-A3B, (b) Llama-3.1-70B. Outlined: a tuple at its own workload’s loads. On another workload’s loads a tuple can take up to 4.17× (Qwen3-30B-A3B) and 1.22× (Llama-3.1-70B) as long.

96.77% on Qwen3-30B-A3B and 96.86–97.11% on Llama3.1-70B. There D ETENT stays within [−0.93, +0.96] points of SGLang-Det (McNemar’s p 0.18–1).

Only Stock’s answers depend on the load (Figure 10b). Across the three points, Stock’s text changes on 790 (Qwen330B-A3B) and 484 (Llama-3.1-70B) of the 1,319 problems, and its answer on 33 and 16. SGLang-Det’s text never changes. D ETENT serves G2 and G3 with different tuples but gives identical text, so SGLang-Det and D ETENT each form one comparison across G2 and G3. D ETENT’s G1 text can differ from its G2 text because its G1 tuple splits the KV and so sums in a different order. Its answers follow its tuple, not the load. At G1 and G2, D ETENT’s p50 latency per output token is 1.006–1.057× Stock’s, and SGLang-Det’s is 1.38–1.76× D ETENT’s. At G3, D ETENT’s throughput is 0.945× and 0.934× Stock’s. So D ETENT gains its speed without losing accuracy. 8.5

Table 3. Online: ported LLM-42’s median end-to-end latency over D ETENT’s on Llama-3.1-8B, at four request rates and six deterministic shares. Deterministic

12 QPS

14 QPS

16 QPS

18 QPS

2% 5% 10% 20% 50% 100%

1.155 1.187 1.212 1.265 2.023 3.946

1.161 1.191 1.221 1.302 2.408 4.322

1.185 1.212 1.242 1.379 2.842 4.756

1.198 1.236 1.289 1.530 3.249 5.147

Table 4. Offline: D ETENT’s output throughput over ported LLM42’s on Llama-3.1-8B, at six deterministic shares. Deterministic 2% 5% 10% 20% 50% 100%

ShareGPT

arXiv

1.098 1.224 1.357 1.471 1.727 2.147

0.975 0.982 0.987 0.996 1.063 1.217

Comparison with LLM-42

LLM-42 (Gond et al., 2026a) makes a chosen share of requests deterministic without changing the kernels. It verifies fast-path tokens and rolls back mismatches. We port it to our SGLang base and reproduce its single-GPU setting on one H100 SXM. That setting runs Llama-3.1-8B on ShareGPT at 12 to 18 requests per second (QPS) with 2 to 100% of requests deterministic. LLM-42 used an H100 PCIe. D E -

TENT ’s tuner judged a tie on this workload and committed the default tuple. With prefix caching on, D ETENT is batch-

invariant while Stock fails 14 of the 26 cells. Online. With 4,096 requests and three independent servers per point, D ETENT’s median end-to-end latency is 3.5– 5.1% above Stock’s, while SGLang-Det’s is 39–49% above.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed Stock Detent, tuned

Detent, default SGLang-Det / vLLM-Det

(a) 2.3×

p50

Table 5. Seeded rollouts. 512/512: bit-identical replays, rescorings. Ratios: D ETENT’s throughput over Stock’s and SGLang-Det’s.

(b) 3.4×

2.6×

2.5×

2.8×

6.0×

3.4×

3.5×

2×

D ETENT

Model, H100s

SGLang-Det

BI Replay, rescoring / Stock / SGLang-Det 1× 9.2×

p99

6.4×

latency / Stock

Llama-3.1-8B, 1 yes Qwen3-30B-A3B, 4 yes Llama-3.1-70B, 4 yes

1.15× 1.27× 1.35×

BI, mixed no no no

2×

A4

B3

C1

C2

(c)

p50

0.917× 0.837× 0.872×

4×

1× A4

B3

C1

C2

3.7×

3.8×

(d)

3.2×

3.0×

9.8×

10.3×

4.5×

5.2×

49.2×

18.4×

2× 1×

p99

512/512 512/512 512/512

8.6

On-Policy RL Rollouts and Training

For RL, we commit seeded temperature sampling and refuse any other sampling. D ETENT’s trainer runs the serving kernels under the same tuple (§3). As in TBIK (Zhang et al., 2026), we compare it with Stock and with SGLang-Det plus a transformers trainer on batch-invariant operators (BIO). Stock trains with transformers.

4× 2× 1× A4

B3

C1

C2

A4

B3

C1

C2

Figure 9. Latency per output token over Stock’s at the loads vLLM serves, p50 and p99 (rows). SGLang (left), vLLM (right). (a, b) Qwen3-30B-A3B, (c, d) Llama-3.1-70B.

Ported LLM-42 is slower than D ETENT at every rate and deterministic share (Figure 11a,b and Table 3). At 2%, LLM-42 makes only 2% of requests deterministic and D E TENT makes all of them. Even so, LLM-42’s median latency there is 1.15–1.20× D ETENT’s. No deterministic arm’s output changed at the token level: 0 of 45,184 comparisons for D ETENT and 0 of 87,916 for LLM-42’s deterministic requests. Offline. D ETENT’s throughput is 0.964× (ShareGPT) and 0.952× (arXiv) Stock’s. It is 1.16× and 1.11× SGLangDet’s. On ShareGPT, it is 1.10–2.15× ported LLM-42’s (Table 4). On arXiv, LLM-42 at 20% deterministic or less has 0.4–2.6% higher throughput than D ETENT. Fidelity of the port. On the official fork’s online grid (12 and 18 requests per second at 2, 10 and 100%), the port is within 1.1% of official LLM-42. The pre-registered offline check failed by 44%. It divided each stack’s LLM-42 throughput by that of the stack’s own SGLang deterministic mode. That mode runs 1.76× as fast on our SGLang base as on the official fork’s older one. We corrected the check after seeing the data. It now compares LLM-42 throughput directly and passes within its 10% tolerance. The port’s throughput is 0.987× official LLM-42’s.

Seeded rollouts. Table 5 gives D ETENT’s seeded rollouts on Llama-3.1-8B (one H100) and on Qwen3-30B-A3B and Llama-3.1-70B (four H100s each). Its seeded sampling is batch-invariant. All 512 GSM8K (Cobbe et al., 2021) rollouts at temperature 1 replay bit for bit under other batches and cache splits and match their rescoring. At 1,024 output tokens, its throughput is 0.84–0.92× Stock’s and 1.15– 1.35× SGLang-Det’s. SGLang-Det’s seeded rollouts change in all 20 cells where one co-batched request is greedy or uses top-p. They hold when every request samples with temperature. Trainer agreement. We use TBIK’s on-policy protocol on Llama-3.1-8B (GSM8K, full fine-tuning, 100 steps) but lower the learning rate from TBIK’s 10−5 to 10−6 . D E TENT ’s trainer matches the rollout’s bits on every token of every step, so KL is 0. Stock matches 1.0–47.7% of the tokens and SGLang-Det with BIO 9.9–63.9% (Figure 12a). Per generated token, D ETENT’s rollout takes 1.03× Stock’s time and its whole step stays within 4% of Stock’s, while SGLang-Det with BIO takes 1.47× (Figure 12b). Its median step takes 1.09× Stock’s because it generates 11% more tokens. BF16 GRPO stability. In the sanity test of Qi et al. (2025), BF16 GRPO on DeepSeek-R1-Distill-Qwen-1.5B (Guo et al., 2025) collapses. We run that recipe in D ETENT with one GPU serving and one training. Stock’s reward falls from 0.733 (step 299) to 0.436 (step 500), and its AIME 2024 (Mathematical Association of America, 2024) Pass@1 from 29.2% to 8.5% (Figure 12c,d). D ETENT keeps KL at 0 through step 600, and its reward reaches 0.843. It takes 1.18× Stock’s time per token, most of it in training.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed Stock

SGLang-Det

(a)

Detent, tuned

(b)

∆ vs SGLang-Det

−0.15

G1

33

[−0.92, +0.66]

Qwen3-30B-A3B

G2

G3

G2, G3

Qwen3-30B-A3B 0

+0.53

0

[−0.15, +1.02]

16

+0.08

G1

[−0.39, +0.50]

Llama-3.1-70B

G2

G3

G2, G3

92

Llama-3.1-70B 0

+0.15

0

[−0.36, +0.57]

93

94

95

96

accuracy (%), 95% interval

0

20

40

problems

Figure 10. D ETENT keeps GSM8K accuracy. (a) Accuracy with 95% Wilson intervals. Right: D ETENT minus SGLang-Det with its 95% interval. (b) Answers that change across G1–G3 (D ETENT: G2 vs. G3, as its G1 tuple splits the KV). LLM-42, 10%

SGLang-Det

2 1.5

1.00

Detent (all) Stock

1 12

14

16

request rate (QPS)

18

2

5

10

20

50

0

2 10 100

100

600

698

690

400 200 0

2 10 100

LLM-42 (%)

LLM-42 deterministic share (%)

681

800

560

7.2k

9.7k

8.8k

10.1k

5k

offline output tokens/s

1.25

3

10k

4.5k

1.50

5

(d)

8.4k

8

1.75

LLM-42, 100%

(c)

715

LLM-42, 2%

611

Detent

(b)

offline output tokens/s

2.00

SGLang-Det

median E2E (s)

median E2E latency (s)

Stock

(a)

LLM-42 (%)

Figure 11. D ETENT makes every request deterministic at 3.5–5.1% above Stock’s median latency. LLM-42 costs more even when only 2% of requests are deterministic. Llama-3.1-8B on one H100, ShareGPT online, median of three servers per point. (a) Online, median end-to-end latency against request rate. LLM-42 at 100% (4.6–6.9 s) is off this scale. (b) Online at 18 requests per second, LLM-42 as its deterministic share grows. (c, d) Offline output throughput on ShareGPT and arXiv, one run per configuration. LLM-42 runs on its port to our SGLang base. D ETENT’s tuner judged a tie on this workload and committed the default tuple. Stock

50 25 0

1×

50 RL step

(b) trainer 40

1.47 1.00

0.96

1.00

100

0.8

1.18

0.5× 0

0

(b) rollout

Pass@1 (%)

75

Detent

(c)

1.5×

Detent: 100% every step

time / token

identical (%)

100

SGLang-Det + BIO

(b)

reward

(a)

0.6 0.4

8B (TBIK)

1.5B (sanity)

(d)

30 20 10 0

0

200

RL step

400

600

0

200

RL step

400

600

Figure 12. D ETENT’s trainer matches the rollout’s bits at near-stock cost. Its GRPO run does not collapse. (a) Bit-identical tokens and (b) time per token over Stock’s, TBIK’s protocol on Llama-3.1-8B. Sanity test on DeepSeek-R1-Distill-Qwen-1.5B, (c) training reward (20-step mean), (d) AIME 2024.

8.7

Ablation Across Deployments

Figure 13 adds eight deployments solved from specs alone. They include Llama-3.1-70B at TP=8, Phi-3.5-mini (Abdin et al., 2024), gemma-7b-it (Gemma Team et al., 2024) and Llama-3.1-8B in both engines. With the main deployments and a B200, they span two engines, two GPU generations,

tensor-parallel degrees 1, 2, 4 and 8, head dimensions 64 to 256, and dense and MoE models. Each ablates one part of the main deployments. D ETENT is batch-invariant on all eight, while Stock changes outputs in 8 to 16 cells. On these eight, §8.3’s tuner ablation repeats. The default tuple, which never splits the KV, takes 0.99–1.14× Stock’s latency on WA and 1.14–1.65× on W-B-c. Tuning each workload brings

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

0 engine GPUs head dim pins Stock fails

Llama-3.1 70B SGLang 8×H100 128 3+3+0 15/26

Qwen3 14B

SGLang 4×H100 128 3+3+0

8/26

Qwen3 8B

SGLang 2×H100 128 3+3+0

8/26

Llama-3.2 1B SGLang 1×H100 64 8+0+0 9/26

16/26

Llama-3.1 8B

SGLang

14/26

1×H100 128 3+0+0

vLLM

10/24

1.60

2.34 gemma 7b-it SGLang 1×H100 256 8+0+0 16/26

1.07 1.09 1.29

1.01 1.03 1.03 1.30 1.16 1.43

0.97 1.02 0.99 1.14

0.99 1.01 1.00 1.14 1.19 1.33

0.99 1.02 1.00 1.20 1.19 1.39 Phi-3.5 mini

SGLang 1×H100 96 2+0+0

2.61

SGLang-Det / vLLM-Det (light: W-A, dark: W-B-c)

2.84 3.07

Detent, default

1.03 1.08 1.12 1.65 1.38 1.65

1×

Detent, tuned

1.00 1.03 1.04 1.31 1.57 1.83

2×

1.09 1.13 1.10 1.43 1.79 2.10

3×

1.07 1.09 1.14 1.48 1.82 2.00

latency / Stock

Stock (1×)

Llama-3.1 8B SGLang 1×B200 128 2+0+0 16/26

Figure 13. On eight more deployments and a B200, solved from specs alone, D ETENT is batch-invariant. Bars: closed-loop median at W-A and W-B-c. Pins: attention + all-reduce + MoE. Stock fails: battery cells where Stock’s bits change.

this to 0.97–1.09× and 1.01–1.13×, which is 1.14–3.01× faster than the deterministic modes. On a B200 with Llama3.1-8B, D ETENT and SGLang-Det pass 26 cells while Stock fails 16. There, the default tuple takes 1.29× (W-A) and 2.34× (W-B-c) Stock’s latency. The tuned tuples take 1.07× and 1.09×, 1.49× and 2.38× faster than SGLang-Det.

9

R ELATED W ORK

Batch-invariant kernels. D ETENT takes batch-invariant kernels as given and makes the selection among them batchinvariant. Thinking Machines wrote batch-invariant kernels with one matrix-multiplication configuration and a fixed attention split size (He & Thinking Machines Lab, 2025). SGLang’s deterministic mode builds on them and runs FA3 with one split (The SGLang Team, 2025). So does vLLM’s batch-invariant mode (The vLLM Team, 2026). TBIK fixes one reduction tree, so outputs also match across TP degrees (Zhang et al., 2026).

Controlling what kernels see. LLM-42 decodes on the fast path, replays tokens under a fixed reduction schedule and rolls back mismatches (Gond et al., 2026a; §8.5). CoRun prefills requests alone and pads decode batches to a fixed row count, with split-KV off (Zhao et al., 2026). Both control the batches that stock kernels see. D ETENT instead fixes what kernel selection reads. This lets it keep split-KV and reuse prefixes across turns.

RL rollout and training. Rollout and training engines disagree on token probabilities (Yao et al., 2025). Qi et al. (2025) switch to FP16 to narrow the gap. Thinking Machines, Miles, VeXact and TBIK make the two bit-identical with batch-invariant kernels (He & Thinking Machines Lab, 2025; RadixArk Team, 2025; Zhong et al., 2026; Zhang et al., 2026). Our trainer gets the same result from pinned kernel selection at near-stock speed.

Reproducibility and verification. Hawkeye replays GPU matrix multiplication on a CPU (Badash et al., 2026), and EigenAI audits inference by re-execution (Alves et al., 2026). TOPLOC and DiFR verify inference despite nondeterminism (Ong et al., 2025; Karvonen et al., 2025), and Vosti verifies a single-GPU engine (Qin et al., 2026). D E TENT makes multi-GPU engines like vLLM and SGLang batch-invariant. Kernel selection and autotuning. Engines adapt kernels to each step’s batch (Kwon et al., 2023; Zheng et al., 2024; Ye et al., 2025; Gond et al., 2026b). Triton’s autotuner remeasures when its arguments change (Tillet et al., 2019; The Triton Developers, 2026). D ETENT tunes its pins once per workload before serving. Program analysis. The cut enforces the information-flow policy (Denning, 1976) that batch state must not reach kernel selection. Quantitative information flow (McCamant & Ernst, 2008) and Swift (Chong et al., 2007) also use maximum flows. D ETENT uses partial evaluation (Jones et al., 1993), slicing (Weiser, 1984) and value range analysis (Harrison, 1977). VRA applies it to pins.

10

C ONCLUSION

D ETENT makes LLM serving engines such as vLLM and SGLang batch-invariant while keeping their fast kernels. It pins the few values through which the batch reaches kernel selection and tunes them per workload, so latency per output token stays at 1.04–1.17× Stock’s. The engines’ deterministic modes take 1.67–10.3×. A trainer on the same kernels and tuple matches every rollout logprob bit for bit at near-stock speed, keeping KL at 0. BF16 GRPO runs stay stable with D ETENT where Stock’s collapses. Because selection then depends only on the deployment, the tuple and the request, anyone with the same deployment and tuple can re-execute a request and check its output bit for bit.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

Limits and next steps. • Scope. Outputs match only within one deployment, model and commitment tuple. Other TP degrees also need TPinvariant kernels like TBIK’s. • Writing SelIR. Users need to write each selection in SelIR. We plan LLM-assisted transcription and native-code selectors. • Traffic drift. Speed depends on how well the tuned workloads match traffic. Re-tuning must switch tuples where no batch mixes them. • Quantization. Dynamic per-tensor FP8 activation scales span the batch. Batch-invariant FP8 serving needs pertoken or static scales. • Agent RL. Agentic rollouts grow long contexts and end on-policy steps with few requests, where split-KV matters most. D ETENT is one of the few batch-invariant systems that keep it, so agent RL is our next target.

R EFERENCES Abdin, M., Aneja, J., Awadalla, H., Awadallah, A., Awan, A. A., Bach, N., Bahree, A., Bakhtiari, A., Bao, J., Behl, H., Benhaim, A., Bilenko, M., Bjorck, J., Bubeck, S., Cai, M., Cai, Q., Chaudhary, V., Chen, D., Chen, D., Chen, W., Chen, Y.-C., Chen, Y.-L., Cheng, H., Chopra, P., Dai, X., Dixon, M., Eldan, R., Fragoso, V., Gao, J., Gao, M., Gao, M., Garg, A., Del Giorno, A., Goswami, A., Gunasekar, S., Haider, E., Hao, J., Hewett, R. J., Hu, W., Huynh, J., Iter, D., Jacobs, S. A., Javaheripi, M., Jin, X., Karampatziakis, N., Kauffmann, P., Khademi, M., Kim, D., Kim, Y. J., Kurilenko, L., Lee, J. R., Lee, Y. T., Li, Y., Li, Y., Liang, C., Liden, L., Lin, X., Lin, Z., Liu, C., Liu, L., Liu, M., Liu, W., Liu, X., Luo, C., Madan, P., Mahmoudzadeh, A., Majercak, D., Mazzola, M., Mendes, C. C. T., Mitra, A., Modi, H., Nguyen, A., Norick, B., Patra, B., Perez-Becker, D., Portet, T., Pryzant, R., Qin, H., Radmilac, M., Ren, L., de Rosa, G., Rosset, C., Roy, S., Ruwase, O., Saarikivi, O., Saied, A., Salim, A., Santacroce, M., Shah, S., Shang, N., Sharma, H., Shen, Y., Shukla, S., Song, X., Tanaka, M., Tupini, A., Vaddamanu, P., Wang, C., Wang, G., Wang, L., Wang, S., Wang, X., Wang, Y., Ward, R., Wen, W., Witte, P., Wu, H., Wu, X., Wyatt, M., Xiao, B., Xu, C., Xu, J., Xu, W., Xue, J., Yadav, S., Yang, F., Yang, J., Yang, Y., Yang, Z., Yu, D., Yuan, L., Zhang, C., Zhang, C., Zhang, J., Zhang, L. L., Zhang, Y., Zhang, Y., Zhang, Y., and Zhou, X. Phi-3 technical report: A highly capable language model locally on your phone. arXiv preprint arXiv:2404.14219, 2024. Alves, D. R., Patankar, V., Pereira, M., Stephens, J., Vaziri, N., and Kannan, S. EigenAI: Deterministic inference, verifiable results. arXiv preprint arXiv:2602.00182, 2026.

anon8231489123. ShareGPT Vicuna unfiltered. Hugging Face dataset anon8231489123/ShareGPT Vicuna unfiltered, revision 192ab21, 2023. Ansel, J., Yang, E., He, H., Gimelshein, N., Jain, A., Voznesensky, M., Bao, B., Bell, P., Berard, D., Burovski, E., Chauhan, G., Chourdia, A., Constable, W., Desmaison, A., DeVito, Z., Ellison, E., Feng, W., Gong, J., Gschwind, M., Hirsh, B., Huang, S., Kalambarkar, K., Kirsch, L., Lazos, M., Lezcano, M., Liang, Y., Liang, J., Lu, Y., Luk, C. K., Maher, B., Pan, Y., Puhrsch, C., Reso, M., Saroufim, M., Siraichi, M. Y., Suk, H., Suo, M., Tillet, P., Wang, E., Wang, X., Wen, W., Zhang, S., Zhao, X., Zhou, K., Zou, R., Mathews, A., Chanan, G., Wu, P., and Chintala, S. PyTorch 2: Faster machine learning through dynamic Python bytecode transformation and graph compilation. In Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 2 (ASPLOS), pp. 929–947. ACM, 2024. doi: 10.1145/3620665.3640366. Badash, E., Boneh, D., Komargodski, I., and Srivastava, M. Hawkeye: Reproducing GPU-level non-determinism. In Proceedings of Machine Learning and Systems (MLSys), volume 8, pp. 1830–1842. MLSys, 2026. Chen, T., Moreau, T., Jiang, Z., Zheng, L., Yan, E., Shen, H., Cowan, M., Wang, L., Hu, Y., Ceze, L., Guestrin, C., and Krishnamurthy, A. TVM: An automated end-to-end optimizing compiler for deep learning. In 13th USENIX Symposium on Operating Systems Design and Implementation (OSDI), pp. 578–594. USENIX Association, 2018. Chong, S., Liu, J., Myers, A. C., Qi, X., Vikram, K., Zheng, L., and Zheng, X. Secure web applications via automatic partitioning. In Proceedings of the 21st ACM Symposium on Operating Systems Principles (SOSP), pp. 31–44. ACM, 2007. Cobbe, K., Kosaraju, V., Bavarian, M., Chen, M., Jun, H., Kaiser, L., Plappert, M., Tworek, J., Hilton, J., Nakano, R., Hesse, C., and Schulman, J. Training verifiers to solve math word problems. arXiv preprint arXiv:2110.14168, 2021. Cohan, A., Dernoncourt, F., Kim, D. S., Bui, T., Kim, S., Chang, W., and Goharian, N. A discourse-aware attention model for abstractive summarization of long documents. In Proceedings of the 2018 Conference of the North American Chapter of the Association for Computational Linguistics: Human Language Technologies, Volume 2 (Short Papers) (NAACL-HLT), pp. 615–621. Association for Computational Linguistics, 2018. Denning, D. E. A lattice model of secure information flow. Communications of the ACM, 19(5):236–243, 1976.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

Edmonds, J. and Karp, R. M. Theoretical improvements in algorithmic efficiency for network flow problems. Journal of the ACM, 19(2):248–264, 1972. Gemma Team, Mesnard, T., Hardin, C., Dadashi, R., Bhupatiraju, S., Pathak, S., Sifre, L., Rivière, M., Kale, M. S., Love, J., Tafti, P., Hussenot, L., Sessa, P. G., Chowdhery, A., Roberts, A., Barua, A., Botev, A., Castro-Ros, A., Slone, A., Héliou, A., Tacchetti, A., Bulanova, A., Paterson, A., Tsai, B., Shahriari, B., Le Lan, C., ChoquetteChoo, C. A., Crepy, C., Cer, D., Ippolito, D., Reid, D., Buchatskaya, E., Ni, E., Noland, E., Yan, G., Tucker, G., Muraru, G.-C., Rozhdestvenskiy, G., Michalewski, H., Tenney, I., Grishchenko, I., Austin, J., Keeling, J., Labanowski, J., Lespiau, J.-B., Stanway, J., Brennan, J., Chen, J., Ferret, J., Chiu, J., Mao-Jones, J., Lee, K., Yu, K., Millican, K., Sjoesund, L. L., Lee, L., Dixon, L., Reid, M., Mikuła, M., Wirth, M., Sharman, M., Chinaev, N., Thain, N., Bachem, O., Chang, O., Wahltinez, O., Bailey, P., Michel, P., Yotov, P., Chaabouni, R., Comanescu, R., Jana, R., Anil, R., McIlroy, R., Liu, R., Mullins, R., Smith, S. L., Borgeaud, S., Girgin, S., Douglas, S., Pandya, S., Shakeri, S., De, S., Klimenko, T., Hennigan, T., Feinberg, V., Stokowiec, W., Chen, Y.-h., Ahmed, Z., Gong, Z., Warkentin, T., Peran, L., Giang, M., Farabet, C., Vinyals, O., Dean, J., Kavukcuoglu, K., Hassabis, D., Ghahramani, Z., Eck, D., Barral, J., Pereira, F., Collins, E., Joulin, A., Fiedel, N., Senter, E., Andreev, A., and Kenealy, K. Gemma: Open models based on Gemini research and technology. arXiv preprint arXiv:2403.08295, 2024. Gond, R., Kamath, A. K., Ramjee, R., and Panwar, A. LLM42: Enabling determinism in LLM inference with verified speculation. In Proceedings of the ACM SIGOPS 32nd Symposium on Operating Systems Principles (SOSP), pp. 916–931. ACM, 2026a. doi: 10.1145/3830418.3843901. arXiv:2601.17768. Gond, R., Kwatra, N., and Ramjee, R. TokenWeave: Efficient compute-communication overlap for distributed LLM inference. In Proceedings of Machine Learning and Systems (MLSys), volume 8, pp. 36–56. MLSys, 2026b. Grattafiori, A., Dubey, A., Jauhri, A., Pandey, A., Kadian, A., Al-Dahle, A., Letman, A., Mathur, A., Schelten, A., Vaughan, A., Yang, A., Fan, A., Goyal, A., Hartshorn, A., Yang, A., Mitra, A., Sravankumar, A., Korenev, A., Hinsvark, A., Rao, A., Zhang, A., Rodriguez, A., Gregerson, A., Spataru, A., Roziere, B., Biron, B., Tang, B., Chern, B., Caucheteux, C., Nayak, C., Bi, C., Marra, C., McConnell, C., Keller, C., Touret, C., Wu, C., Wong, C., Ferrer, C. C., Nikolaidis, C., Allonsius, D., Song, D., Pintz, D., Livshits, D., Wyatt, D., Esiobu, D., Choudhary, D., Mahajan, D., Garcia-Olano, D., Perino, D., Hupkes, D., Lakomkin, E., AlBadawy, E., Lobanova, E., Dinan,

E., Smith, E. M., Radenovic, F., Guzmán, F., Zhang, F., Synnaeve, G., Lee, G., Anderson, G. L., Thattai, G., Nail, G., Mialon, G., Pang, G., Cucurell, G., Nguyen, H., Korevaar, H., Xu, H., Touvron, H., Zarov, I., Ibarra, I. A., Kloumann, I., Misra, I., Evtimov, I., Zhang, J., Copet, J., Lee, J., Geffert, J., Vranes, J., Park, J., Mahadeokar, J., Shah, J., van der Linde, J., Billock, J., Hong, J., Lee, J., Fu, J., Chi, J., Huang, J., Liu, J., Wang, J., Yu, J., Bitton, J., Spisak, J., Park, J., Rocca, J., Johnstun, J., Saxe, J., Jia, J., Alwala, K. V., Prasad, K., Upasani, K., Plawiak, K., Li, K., Heafield, K., Stone, K., El-Arini, K., Iyer, K., Malik, K., Chiu, K., Bhalla, K., Lakhotia, K., Rantala-Yeary, L., van der Maaten, L., Chen, L., Tan, L., Jenkins, L., Martin, L., Madaan, L., Malo, L., Blecher, L., Landzaat, L., de Oliveira, L., Muzzi, M., Pasupuleti, M., Singh, M., Paluri, M., Kardas, M., Tsimpoukelli, M., Oldham, M., Rita, M., Pavlova, M., Kambadur, M., Lewis, M., Si, M., Singh, M. K., Hassan, M., Goyal, N., Torabi, N., Bashlykov, N., Bogoychev, N., Chatterji, N., Zhang, N., Duchenne, O., Çelebi, O., Alrassy, P., Zhang, P., Li, P., Vasic, P., Weng, P., Bhargava, P., Dubal, P., Krishnan, P., Koura, P. S., Xu, P., He, Q., Dong, Q., Srinivasan, R., Ganapathy, R., Calderer, R., Cabral, R. S., Stojnic, R., Raileanu, R., Maheswari, R., Girdhar, R., Patel, R., Sauvestre, R., Polidoro, R., Sumbaly, R., Taylor, R., Silva, R., Hou, R., Wang, R., Hosseini, S., Chennabasappa, S., Singh, S., Bell, S., Kim, S. S., Edunov, S., Nie, S., Narang, S., Raparthy, S., Shen, S., Wan, S., Bhosale, S., Zhang, S., Vandenhende, S., Batra, S., Whitman, S., Sootla, S., Collot, S., Gururangan, S., Borodinsky, S., Herman, T., Fowler, T., Sheasha, T., Georgiou, T., Scialom, T., Speckbacher, T., Mihaylov, T., Xiao, T., Karn, U., Goswami, V., Gupta, V., Ramanathan, V., Kerkez, V., Gonguet, V., Do, V., Vogeti, V., Albiero, V., Petrovic, V., Chu, W., Xiong, W., Fu, W., Meers, W., Martinet, X., Wang, X., Wang, X., Tan, X. E., Xia, X., Xie, X., Jia, X., Wang, X., Goldschlag, Y., Gaur, Y., Babaei, Y., Wen, Y., Song, Y., Zhang, Y., Li, Y., Mao, Y., Coudert, Z. D., Yan, Z., Chen, Z., Papakipos, Z., Singh, A., Srivastava, A., Jain, A., Kelsey, A., Shajnfeld, A., Gangidi, A., Victoria, A., Goldstand, A., Menon, A., Sharma, A., Boesenberg, A., Baevski, A., Feinstein, A., Kallet, A., Sangani, A., Teo, A., Yunus, A., Lupu, A., Alvarado, A., Caples, A., Gu, A., Ho, A., Poulton, A., Ryan, A., Ramchandani, A., Dong, A., Franco, A., Goyal, A., Saraf, A., Chowdhury, A., Gabriel, A., Bharambe, A., Eisenman, A., Yazdan, A., James, B., Maurer, B., Leonhardi, B., Huang, B., Loyd, B., De Paola, B., Paranjape, B., Liu, B., Wu, B., Ni, B., Hancock, B., Wasti, B., Spence, B., Stojkovic, B., Gamido, B., Montalvo, B., Parker, C., Burton, C., Mejia, C., Liu, C., Wang, C., Kim, C., Zhou, C., Hu, C., Chu, C.H., Cai, C., Tindal, C., Feichtenhofer, C., Gao, C., Civin, D., Beaty, D., Kreymer, D., Li, D., Adkins, D., Xu, D., Testuggine, D., David, D., Parikh, D., Liskovich, D., Foss,

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

D., Wang, D., Le, D., Holland, D., Dowling, E., Jamil, E., Montgomery, E., Presani, E., Hahn, E., Wood, E., Le, E.-T., Brinkman, E., Arcaute, E., Dunbar, E., Smothers, E., Sun, F., Kreuk, F., Tian, F., Kokkinos, F., Ozgenel, F., Caggioni, F., Kanayet, F., Seide, F., Florez, G. M., Schwarz, G., Badeer, G., Swee, G., Halpern, G., Herman, G., Sizov, G., Zhang, G., Lakshminarayanan, G., Inan, H., Shojanazeri, H., Zou, H., Wang, H., Zha, H., Habeeb, H., Rudolph, H., Suk, H., Aspegren, H., Goldman, H., Zhan, H., Damlaj, I., Molybog, I., Tufanov, I., Leontiadis, I., Veliche, I.-E., Gat, I., Weissman, J., Geboski, J., Kohli, J., Lam, J., Asher, J., Gaya, J.-B., Marcus, J., Tang, J., Chan, J., Zhen, J., Reizenstein, J., Teboul, J., Zhong, J., Jin, J., Yang, J., Cummings, J., Carvill, J., Shepard, J., McPhie, J., Torres, J., Ginsburg, J., Wang, J., Wu, K., U, K. H., Saxena, K., Khandelwal, K., Zand, K., Matosich, K., Veeraraghavan, K., Michelena, K., Li, K., Jagadeesh, K., Huang, K., Chawla, K., Huang, K., Chen, L., Garg, L., A, L., Silva, L., Bell, L., Zhang, L., Guo, L., Yu, L., Moshkovich, L., Wehrstedt, L., Khabsa, M., Avalani, M., Bhatt, M., Mankus, M., Hasson, M., Lennie, M., Reso, M., Groshev, M., Naumov, M., Lathi, M., Keneally, M., Liu, M., Seltzer, M. L., Valko, M., Restrepo, M., Patel, M., Vyatskov, M., Samvelyan, M., Clark, M., Macey, M., Wang, M., Hermoso, M. J., Metanat, M., Rastegari, M., Bansal, M., Santhanam, N., Parks, N., White, N., Bawa, N., Singhal, N., Egebo, N., Usunier, N., Mehta, N., Laptev, N. P., Dong, N., Cheng, N., Chernoguz, O., Hart, O., Salpekar, O., Kalinli, O., Kent, P., Parekh, P., Saab, P., Balaji, P., Rittner, P., Bontrager, P., Roux, P., Dollar, P., Zvyagina, P., Ratanchandani, P., Yuvraj, P., Liang, Q., Alao, R., Rodriguez, R., Ayub, R., Murthy, R., Nayani, R., Mitra, R., Parthasarathy, R., Li, R., Hogan, R., Battey, R., Wang, R., Howes, R., Rinott, R., Mehta, S., Siby, S., Bondu, S. J., Datta, S., Chugh, S., Hunt, S., Dhillon, S., Sidorov, S., Pan, S., Mahajan, S., Verma, S., Yamamoto, S., Ramaswamy, S., Lindsay, S., Lindsay, S., Feng, S., Lin, S., Zha, S. C., Patil, S., Shankar, S., Zhang, S., Zhang, S., Wang, S., Agarwal, S., Sajuyigbe, S., Chintala, S., Max, S., Chen, S., Kehoe, S., Satterfield, S., Govindaprasad, S., Gupta, S., Deng, S., Cho, S., Virk, S., Subramanian, S., Choudhury, S., Goldman, S., Remez, T., Glaser, T., Best, T., Koehler, T., Robinson, T., Li, T., Zhang, T., Matthews, T., Chou, T., Shaked, T., Vontimitta, V., Ajayi, V., Montanez, V., Mohan, V., Kumar, V. S., Mangla, V., Ionescu, V., Poenaru, V., Mihailescu, V. T., Ivanov, V., Li, W., Wang, W., Jiang, W., Bouaziz, W., Constable, W., Tang, X., Wu, X., Wang, X., Wu, X., Gao, X., Kleinman, Y., Chen, Y., Hu, Y., Jia, Y., Qi, Y., Li, Y., Zhang, Y., Zhang, Y., Adi, Y., Nam, Y., Wang, Y., Zhao, Y., Hao, Y., Qian, Y., Li, Y., He, Y., Rait, Z., DeVito, Z., Rosnbrick, Z., Wen, Z., Yang, Z., Zhao, Z., and Ma, Z. The Llama 3 herd of models. arXiv preprint arXiv:2407.21783, 2024.

Guo, D., Yang, D., Zhang, H., Song, J., Wang, P., Zhu, Q., Xu, R., Zhang, R., Ma, S., Bi, X., Zhang, X., Yu, X., Wu, Y., Wu, Z. F., Gou, Z., Shao, Z., Li, Z., Gao, Z., Liu, A., Xue, B., Wang, B., Wu, B., Feng, B., Lu, C., Zhao, C., Deng, C., Ruan, C., Dai, D., Chen, D., Ji, D., Li, E., Lin, F., Dai, F., Luo, F., Hao, G., Chen, G., Li, G., Zhang, H., Xu, H., Ding, H., Gao, H., Qu, H., Li, H., Guo, J., Li, J., Chen, J., Yuan, J., Tu, J., Qiu, J., Li, J., Cai, J. L., Ni, J., Liang, J., Chen, J., Dong, K., Hu, K., You, K., Gao, K., Guan, K., Huang, K., Yu, K., Wang, L., Zhang, L., Zhao, L., Wang, L., Zhang, L., Xu, L., Xia, L., Zhang, M., Zhang, M., Tang, M., Zhou, M., Li, M., Wang, M., Li, M., Tian, N., Huang, P., Zhang, P., Wang, Q., Chen, Q., Du, Q., Ge, R., Zhang, R., Pan, R., Wang, R., Chen, R. J., Jin, R. L., Chen, R., Lu, S., Zhou, S., Chen, S., Ye, S., Wang, S., Yu, S., Zhou, S., Pan, S., Li, S. S., Zhou, S., Wu, S., Yun, T., Pei, T., Sun, T., Wang, T., Zeng, W., Liu, W., Liang, W., Gao, W., Yu, W., Zhang, W., Xiao, W. L., An, W., Liu, X., Wang, X., Chen, X., Nie, X., Cheng, X., Liu, X., Xie, X., Liu, X., Yang, X., Li, X., Su, X., Lin, X., Li, X. Q., Jin, X., Shen, X., Chen, X., Sun, X., Wang, X., Song, X., Zhou, X., Wang, X., Shan, X., Li, Y. K., Wang, Y. Q., Wei, Y. X., Zhang, Y., Xu, Y., Li, Y., Zhao, Y., Sun, Y., Wang, Y., Yu, Y., Zhang, Y., Shi, Y., Xiong, Y., He, Y., Piao, Y., Wang, Y., Tan, Y., Ma, Y., Liu, Y., Guo, Y., Ou, Y., Wang, Y., Gong, Y., Zou, Y., He, Y., Xiong, Y., Luo, Y., You, Y., Liu, Y., Zhou, Y., Zhu, Y. X., Huang, Y., Li, Y., Zheng, Y., Zhu, Y., Ma, Y., Tang, Y., Zha, Y., Yan, Y., Ren, Z. Z., Ren, Z., Sha, Z., Fu, Z., Xu, Z., Xie, Z., Zhang, Z., Hao, Z., Ma, Z., Yan, Z., Wu, Z., Gu, Z., Zhu, Z., Liu, Z., Li, Z., Xie, Z., Song, Z., Pan, Z., Huang, Z., Xu, Z., Zhang, Z., and Zhang, Z. DeepSeek-R1 incentivizes reasoning in LLMs through reinforcement learning. Nature, 645(8081):633– 638, 2025. doi: 10.1038/s41586-025-09422-z. Harrison, W. H. Compiler analysis of the value ranges for variables. IEEE Transactions on Software Engineering, SE-3(3):243–250, 1977. doi: 10.1109/TSE.1977.231133. He, H. and Thinking Machines Lab. Defeating nondeterminism in LLM inference. Thinking Machines Lab: Connectionism, 2025. Published September 10, 2025. https://doi.org/10.64434/tml.20250910. Jones, N. D., Gomard, C. K., and Sestoft, P. Partial Evaluation and Automatic Program Generation. Prentice Hall, 1993. Karvonen, A., Reuter, D., Rinberg, R., Marks, L., GarrigaAlonso, A., and Warr, K. DiFR: Inference verification despite nondeterminism. arXiv preprint arXiv:2511.20621, 2025. Kwon, W., Li, Z., Zhuang, S., Sheng, Y., Zheng, L., Yu, C. H., Gonzalez, J. E., Zhang, H., and Stoica, I. Ef-

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

ficient memory management for large language model serving with PagedAttention. In Proceedings of the 29th Symposium on Operating Systems Principles (SOSP), pp. 611–626. ACM, 2023. Mathematical Association of America. American Invitational Mathematics Examination (AIME) 2024. AIME I and AIME II, 30 problems, 2024. https://maa. org/maa-invitational-competitions/, accessed October 23, 2026. McCamant, S. and Ernst, M. D. Quantitative information flow as network flow capacity. In Proceedings of the 29th ACM SIGPLAN Conference on Programming Language Design and Implementation (PLDI), pp. 193–205. ACM, 2008.

Shao, Z., Wang, P., Zhu, Q., Xu, R., Song, J., Bi, X., Zhang, H., Zhang, M., Li, Y. K., Wu, Y., and Guo, D. DeepSeekMath: Pushing the limits of mathematical reasoning in open language models. arXiv preprint arXiv:2402.03300, 2024. Tillet, P., Kung, H. T., and Cox, D. Triton: An intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages (MAPL), pp. 10–19. ACM, 2019. The Triton Developers. triton.autotune. Triton documentation, 2026. https://triton-lang.org/ main/python-api/generated/triton.auto tune.html, accessed October 13, 2026.

McNemar, Q. Note on the sampling error of the difference between correlated proportions or percentages. Psychometrika, 12(2):153–157, 1947.

Vendrow, J., Vendrow, E., Beery, S., and Madry, A. Do large language model benchmarks test reliability? arXiv preprint arXiv:2502.03461, 2025.

NVIDIA Corporation. NVIDIA TensorRT documentation. Sections “TensorRT Capabilities” (version 11.4.0) and “Optimizing Builder Performance” (version 10.16.1), 2026. https://docs.nvidia.com/deeplear ning/tensorrt/, accessed October 13, 2026.

The vLLM Team. Batch invariance. vLLM documentation (beta feature), 2026. https://docs.vllm.ai/e n/latest/features/batch_invariance/, accessed October 13, 2026.

Ong, J. M., Di Ferrante, M., Pazdera, A., Garner, R., Jaghouar, S., Basra, M., Ryabinin, M., and Hagemann, J. TOPLOC: A locality sensitive hashing scheme for trustless verifiable inference. In Proceedings of the 42nd International Conference on Machine Learning (ICML), volume 267 of Proceedings of Machine Learning Research, pp. 47196–47211. PMLR, 2025. Qi, P., Liu, Z., Zhou, X., Pang, T., Du, C., Lee, W. S., and Lin, M. Defeating the training-inference mismatch via FP16. arXiv preprint arXiv:2510.26788, 2025. Qin, J., Du, A., Zhang, D., Lentz, M., and Zhuo, D. Vosti: Specifying, implementing, and verifying deterministic LLM inference. arXiv preprint arXiv:2609.38981, 2026. RadixArk Team. Introducing Miles: RL framework to fire up large-scale MoE training. LMSYS Org blog, 2025. Published November 19, 2025. https://lmsys.or g/blog/2025-11-19-miles/. The SGLang Team. Towards deterministic inference in SGLang and reproducible RL training. LMSYS Org blog, 2025. Published September 22, 2025. https: //lmsys.org/blog/2025-09-22-sglang-d eterministic/. Shah, J., Bikshandi, G., Zhang, Y., Thakkar, V., Ramani, P., and Dao, T. FlashAttention-3: Fast and accurate attention with asynchrony and low-precision. In Advances in Neural Information Processing Systems (NeurIPS), volume 37, pp. 68658–68685. Curran Associates, Inc., 2024.

Weiser, M. Program slicing. IEEE Transactions on Software Engineering, SE-10(4):352–357, 1984. Wilson, E. B. Probable inference, the law of succession, and statistical inference. Journal of the American Statistical Association, 22(158):209–212, 1927. Yang, A., Li, A., Yang, B., Zhang, B., Hui, B., Zheng, B., Yu, B., Gao, C., Huang, C., Lv, C., Zheng, C., Liu, D., Zhou, F., Huang, F., Hu, F., Ge, H., Wei, H., Lin, H., Tang, J., Yang, J., Tu, J., Zhang, J., Yang, J., Yang, J., Zhou, J., Zhou, J., Lin, J., Dang, K., Bao, K., Yang, K., Yu, L., Deng, L., Li, M., Xue, M., Li, M., Zhang, P., Wang, P., Zhu, Q., Men, R., Gao, R., Liu, S., Luo, S., Li, T., Tang, T., Yin, W., Ren, X., Wang, X., Zhang, X., Ren, X., Fan, Y., Su, Y., Zhang, Y., Zhang, Y., Wan, Y., Liu, Y., Wang, Z., Cui, Z., Zhang, Z., Zhou, Z., and Qiu, Z. Qwen3 technical report. arXiv preprint arXiv:2505.09388, 2025. Yao, F., Liu, L., Zhang, D., Dong, C., Shang, J., and Gao, J. On the rollout-training mismatch in modern RL systems. In NeurIPS 2025 Workshop on Efficient Reasoning, 2025. https://openreview.net/forum?id=8MHq vb4lK9. Ye, Z., Chen, L., Lai, R., Lin, W., Zhang, Y., Wang, S., Chen, T., Kasikci, B., Grover, V., Krishnamurthy, A., and Ceze, L. FlashInfer: Efficient and customizable attention engine for LLM inference serving. In Proceedings of Machine Learning and Systems (MLSys), volume 7. MLSys, 2025. https://proceedings.mlsys.org/paper_

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

files/paper/2025/file/dbf02b21d77409 a2db30e56866a8ab3a-Paper-Conference. pdf. Yuan, J., Li, H., Ding, X., Xie, W., Li, Y.-J., Zhao, W., Wan, K., Shi, J., Hu, X., and Liu, Z. Understanding and mitigating numerical sources of nondeterminism in LLM inference. In Advances in Neural Information Processing Systems (NeurIPS), volume 38, pp. 188045–188077. Curran Associates, Inc., 2025. doi: 10.52202/085713-5653. Zhang, Z., Ding, X., Yuan, J., Liu, R., Mao, H., Xing, J., and Liu, Z. Deterministic inference across tensor parallel sizes that eliminates training-inference mismatch. In Proceedings of the 43rd International Conference on Machine Learning (ICML), volume 306 of Proceedings of Machine Learning Research, pp. 155026–155042. PMLR, 2026. Zhao, C., Xu, Z., Zhao, L., Li, J., Xu, C., Xu, A., Liu, S., Zhou, K., and Yu, K. DeepGEMM: Clean and efficient BLAS kernel library on GPU. GitHub repository, 2025. https://github.com/deepseek-ai/DeepGE MM. Zhao, S., Yang, J., Chen, Q., Hu, J., Zheng, J., Chen, G., and Chen, X. CoRun: Padding is simple and efficient for deterministic LLM inference. arXiv preprint arXiv:2608.14376, 2026. Zheng, L., Yin, L., Xie, Z., Sun, C., Huang, J., Yu, C. H., Cao, S., Kozyrakis, C., Stoica, I., Gonzalez, J. E., Barrett, C., and Sheng, Y. SGLang: Efficient execution of structured language model programs. In Advances in Neural Information Processing Systems (NeurIPS), volume 37, pp. 62557–62583. Curran Associates, Inc., 2024. Zhong, T., Ling, N., Pi, Y., Wei, Z., Yu, T., Fox, G., Wu, P., and Yu, X. Diagnosing training inference mismatch in LLM reinforcement learning via a zero-mismatch reference. arXiv preprint arXiv:2605.14220, 2026.

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed

A

T HE D ETERMINISTIC M ODES ’ S WITCHES

We turn on SGLang’s deterministic mode with --enabledeterministic-inference and vLLM’s batchinvariant mode with VLLM BATCH INVARIANT=1. Table 6 lists what each mode changes in the versions we run. Figure 3d also shows SGLang’s changes. Both modes keep one conservative kernel per operator.

B

C OUNTING THE S ELECTION S PACE

These are the counts behind Figure 3c, for the FA3 and SGLang versions we evaluate. FA3’s run flash fwd has 19 template parameters, 13 of them Boolean. On sm90, the product of the numbers of values they can take is 737,280. A full build compiles 4,720 forward kernels, 3,280 of them for sm90. A split kernel takes any split count from 2 to 256. On sm90 this gives 332,464 kernel configurations, 18,472 of them at head dimension 128 in bf16. One tuning run on Qwen3-30B-A3B at TP=4 spent 22.4 s per measured point on average. At that rate, measuring the 18,472 configurations at head dimension 128 in bf16 once takes 4.8 days. Without the analysis, FA3’s selection has 227 candidate assignments. SGLang’s fused-MoE tuning grid has 1,920 configurations for each of 18 token counts. One tuned table takes 34,560 runs.

C

A BLATION OF S TOCK ’ S K ERNELS

Stock and D ETENT differ in more than kernel selection. They use different matrix-multiplication and all-reduce kernels. Every arm runs with FlashInfer’s all-reduce fusion off (Table 1). This appendix ablates both differences. It checks that neither weakens Stock, that is, that D ETENT does not come close to Stock by giving itself faster kernels or by taking a feature away from Stock. Matrix-multiplication and all-reduce kernels. We record each workload’s GEMM and all-reduce calls in SGLang while profiling. We time them once on Stock’s kernels and once on D ETENT’s (Table 7, Figure 14). D ETENT’s kernels are the slower ones. They add 0.6% (Qwen3-30B-A3B) and 2.8% (Llama-3.1-70B) to Stock’s GPU time per output token. This is more than twice the measurement noise, so Stock keeps the faster kernels. The GEMM row runs DeepGEMM unpadded, so it weighs the two libraries and not D ETENT’s padding knob. Read with the padding each tuple fixes, the two kernels add 0.6% and 2.1%. D ETENT’s are still the slower ones. In most cases, most of the difference is in the all-reduce and comes from fixing one algorithm per tuple. When both run the one-shot push, D ETENT and Stock take the same time (1.003×). When both run the twoshot pull, D ETENT’s own two-shot kernel takes 1.10–1.21× Stock’s.

FlashInfer’s all-reduce fusion. The fusion does little for Stock, so turning it off does not weaken Stock. On Qwen330B-A3B, Stock without the fusion is 1.3–2.1% faster at a decode step with one request and 11–13% faster on a 1,447-token prefill. From 8 to 256 requests, the difference on this model is within run-to-run noise. With the fusion on in Stock, D ETENT’s decode step takes 0.996–1.103× Stock’s on both models (Table 7, bottom).

Detent: Batch-Invariant LLM Serving and RL Rollouts at Near-Stock Speed Table 6. What the engines’ deterministic modes fix. Operator

Stock

SGLang-Det

vLLM-Det

FA3

split count from a batch-wide heuristic

one split

All-reduce (TP>1)

NCCL tree with fixed channels

Fused MoE

custom, symmetric-memory or NCCL kernels by message size tuned table, nearest token count

GEMM

cuBLAS, kernel picked by shape

batch-invariant mm, addmm on DeepGEMM

Other row-local operators

native kernels

Sampling

default backend

batch-invariant log softmax, mean, rms norm PyTorch backend

one split; ahead-of-time scheduler and cascade attention off NCCL tree, one channel; custom and symmetric-memory all-reduce off one hard-coded configuration; kernels without batch-invariant support refused Triton persistent matmul for linear layers; cuBLAS without split-K batch-invariant log softmax, softmax, mean, rms norm unchanged

GEMM: DeepGEMM vs cuBLAS

added GPU time per output token

(a)

W-A

one hard-coded configuration

all-reduce: ours (pinned) vs Stock’s choice W-B

W-B-c

(b)

sum: unpadded (left), padded (right)

W-A

W-B

W-B-c

+10% +7.5% +5% +2.5% 0 A1

A2

A3

A4

B1

B2

B3

C1

C2

A1

A2

A3

A4

B1

B2

B3

C1

C2

Figure 14. What D ETENT’s GEMM and all-reduce kernels add to Stock’s GPU time per output token at each workload, split by family. Left bar: the GEMM read unpadded. Right bar: with the padding that workload’s tuple fixes. The padding leaves the all-reduce untouched. A1 and A4 in (a) and A1 and C1 in (b) are served unpadded, so their bars are identical. (a) Qwen3-30B-A3B, +0.6% and +0.6% on average; (b) Llama-3.1-70B, +2.8% and +2.1%. Diamonds: the sums, with twice the measurement noise of the unpadded reading.

Table 7. Ablation of Stock’s kernels. Top: D ETENT’s kernels over Stock’s in GPU time (geometric mean and range over the workloads), and the GPU time they add to Stock’s GPU time per output token (mean and range). Middle: the all-reduce row broken down by the algorithm the tuple fixes and by the algorithm Stock picks for each message (in parentheses: share of D ETENT’s extra allreduce time). Bottom: D ETENT’s decode step over Stock’s when Stock runs FlashInfer’s all-reduce fusion. Measure

Qwen3-30B-A3B

Llama-3.1-70B

Kernel swap: D ETENT’s kernel over Stock’s GEMM: DeepGEMM / cuBLAS 0.99 (0.93–1.05) 1.02 (0.99–1.15) with the tuple’s padding 1.00 (0.93–1.06) 1.01 (0.99–1.06) All-reduce: ours / Stock’s pick 1.10 (0.97–1.29) 1.16 (1.04–1.37) Added to Stock’s GPU time +0.6% (−0.3 to +1.6) +2.8% (+0.1 to +10.3) with the tuple’s padding +0.6% (−0.3 to +1.3) +2.1% (+0.5 to +5.5) Twice the noise 0.5% 1.4% All-reduce by the tuple’s workspace and Stock’s pick push, Stock one-shot push 1.003 (1%) push, Stock two-shot pull 1.54 (15%) push, Stock NCCL 2.03 (61%) pull, Stock one-shot push 1.68 (39%) pull, Stock two-shot pull 1.10 (2%) pull, Stock NCCL 0.94 (−18%)

1.003 (0%) 1.40 (2%) 2.03 (1%) 1.72 (84%) 1.21 (6%) 1.01 (6%)

Stock with FlashInfer’s all-reduce fusion: decode step, D ETENT / Stock 1 request per step 1.103 1.009 32 requests per step 0.996 1.046

Related documents

• When Fine-Tuning LLMs Meets Data Privacy: An Empirical Study of Federated Learning in LLM-Based Program Repair#303019
• KernelGenBench: Can LLMs and Agents Write Efficient Kernels Across Operator Sources and Hardware Platforms?#987764
• Blink: CPU-Free LLM Inference by Delegating the Serving Stack to GPU and SmartNIC#2544
• Adaptive Context Parallelism for Production LLM Serving#660891
• Irminsul: MLA-Native Position-Independent Caching for Agentic LLM Serving#168276