feat(SAMPLE-CORE): the checkpoint's sampling defaults, and the Gumbel draw that scanned 248,320 tokens on one thread (#1984, #1985)

[localai-bot]

On Qwen/Qwen3.8-27B vLLM draws its next token from 20 candidates and we draw ours from 248,320, because we read the checkpoint's generation_config.json for eos_token_id and discard every sampling key in it. vllm bench serve no longer sends --temperature, so both engines resolve 1.0 and both take the random-sampling path: that is a correctness divergence on the workload the parity gate runs, not a cosmetic default.

The same fact explains the second defect. RandomSampleKernel was launched <<<n, 1>>> and walked the whole vocabulary on one lane, computing two SplitMix64 rounds and an f64 log per element. Eleven lines above it, cuda_sample.cu already records that a single-block single-thread scan of a ~151k vocab cost ~7.5 ms/token, and that is why greedy argmax was rewritten into the two-pass grid-strided reduction sitting there now. The Gumbel draw never got the same treatment, and its answer is the same lowest-index argmax that reduction already computes.

What changed

The defaults (#1985). The sibling read that already produced generation_config_eos_ids now also yields the six keys ModelConfig.get_diff_sampling_param narrows. GetDiffSamplingParam mirrors that narrowing and vLLM's --generation-config auto|vllm|<dir> selector. Both OpenAI handlers resolve a request against it with upstream's exact precedence: an explicitly sent field wins, an omitted field takes the checkpoint's value, and only a knob neither declares reaches the neutral OpenAI default. nullptr and an empty DefaultSamplingParams both reproduce the old resolution, which is what --generation-config vllm gives a user who wants it back.

The kernel (#1984). ArgmaxPartialKernel is parameterised on the score, and random_sample instantiates it rather than growing a second hand-written reduction. The output is bit-identical, not merely equivalent: every element's score is the same GumbelScore(prob, seed, row, j) on the same device libm, and only the order in which those identical floats are combined changes, while ArgReduce compares the true global index rather than thread or block order. The serial kernel is retained behind VT_FAST_RANDOM_SAMPLE=0, mirroring VT_FAST_ARGMAX, so the equality gate is a same-binary A/B rather than a claim.

SplitMix64, ExpNoise, ArgReduce and the launch partition moved to include/vt/sample_common.h, which cpu_sample.cpp and cuda_sample.cu both include. They had been written out three times, and the CPU copy is the reference the device copies are gated against, so a drift between any two of them was a token difference nothing looked for.

Red before green

test_openai_protocol failed 3 cases / 9 assertions for the intended reason against a stubbed resolution, then passed 37/37 over 269 assertions. One expectation had to be corrected rather than the code: temperature 0 is greedy and __post_init__ clears top_p, top_k and min_p upstream and here, so a checkpoint's top_k: 20 cannot survive a --temperature 0 request. That interaction now has its own case, because it is what keeps the SACRED greedy path unchanged.

Six mutations, each applied, measured, and restored byte for byte:

mutation	gate	result
drop bi < ai from ArgReduce	test_ops_sample	RED, 35 assertions
prob * q instead of prob / q in GumbelScore	test_ops_sample	RED, 2 assertions
delete the sibling sampling read in hf_config.cpp	test_generation_config / test_hf_config / test_openai_serving	RED 3 / RED 1 / RED 3
revert to_sampling_params(..., &default_sampling_params_) in the completion handler	test_openai_serving	RED (reachability)
the same in the chat handler	test_openai_serving	RED (reachability)
delete the --generation-config argument branch	test_serve_recipe_args	RED, 12 assertions

The second row is the honest one, and it is why the table names where each mutation landed. GumbelScore is shared by the CPU reference and by the decomposition the equivalence cases drive, so a change to the score moves both sides together and those cases cannot see it. What caught it was the pre-existing large-N empirical frequency approximates softmax probs case. That is the right division of labour rather than a gap: the equivalence cases own the ORDER, the distribution case owns the FORMULA.

What is not gated here

There is no CUDA toolkit on the implementing host, so cuda_sample.cu was not compiled locally. Its first compile is this pull request's cuda-fat-build, and its first execution is the operator's leased box. The property the kernel rests on is gated on CPU against the production inlines — ArgReduce, GumbelScore and ArgBlocksPerRow under the kernel's exact two-pass partition, compared with vt::RandomSample over ten vocabulary widths including 248,320, six row shapes chosen for ties and masked zeros, and four seeds. The kernel's launch geometry, shared-memory sizing and scratch lifetime are not, and compute-sanitizer is requested for exactly that reason: this change adds a second persistent device scratch and a new grid geometry, and #1958 is an illegal memory access on this same sampler surface.

Reachability is proven at three hops of four. The CLI hop re-execs the real VllmServerMain. The handler hop enters through create_completion / create_chat_completion with the defaults produced by the real LoadHfConfig and GetDiffSamplingParam off a real generation_config.json. The two-line join in server_main.cpp sits after model load and no committed generative checkpoint fixture can reach it, so no CPU test turns red when it is deleted. Stated as a gap, not waived: the smallest thing that would close it is a tiny on-disk causal-LM fixture of the tests/vllm/models/fixtures/llama_embed_e2e kind, which this row did not add.

The measurement this does not merge without

AGENTS.md and #1975 settle it between them: a performance change here does not merge on a green compile, and a per-kernel figure is not evidence. #1929 landed a 708 us to 40 us top-k kernel with fresh review, mutation proofs and a green CUDA build, and cost 16.9 tok/s end to end.

The request is therefore stated before it is run, and it is in ## Now of the row's spec. vllm bench serve against our server and the pinned vLLM on the identical Qwen3.8-27B artifact, vLLM's graphed production configuration as the denominator, no --temperature on either side, in four same-binary arms so each half of the row is attributable on its own:

arm	--generation-config	VT_FAST_RANDOM_SAMPLE	isolates
A	vllm	0	the pre-change baseline
B	auto	0	the config read alone
C	vllm	1	the kernel alone
D	auto	1	the shipping default

--speculative-config would measure nothing either, and the standing baseline recipe carries it (#2002, filed from this row). GPUModelRunner::sample_tokens branches on exec_state_.step.num_draft_tokens > 0 alone and returns the greedy-only RejectionSampler's output, so with drafts present Sampler::forward — and with it vt::RandomSample — is never called, whatever the temperature. Every arm above therefore runs without it. The same routing is a correctness divergence in its own right: at temperature 1.0 under spec decode we emit the target argmax where vLLM draws from the residual distribution, and include/vllm/v1/spec_decode/rejection_sampler.h already states the contract nothing enforces.

--temperature 0 would measure nothing, and it is what every harness in tools/bench/ passes today (run_qwen35_4b_ab.sh sends --temperature 0, profile_vllm_online_gate.py sets temperature=0.0). At temperature 0 PostInit clears top_p, top_k and min_p, the batch is all-greedy, and RandomSample is never called, so all four arms are the same run and four identical numbers read exactly like "no effect".

Predicted, and marked as prediction: D >= C > A, with C - A the large term, of order 10 ms/step recovered at B=32. That figure is INFERRED from this file's own recorded ~7.5 ms/151k anchor scaled by vocabulary and by the extra per-element transcendental, not measured. B - A is predicted small and possibly negative on throughput; the reason to want it is that it makes the two engines sample the same distribution.

Two outcomes would falsify the design rather than the tuning, and both are named in advance: D materially slower than C means the top-k path costs more than the distribution it saves; D no faster than A means the sampler was never the bottleneck and the per-step term lives in the cudaMalloc/cudaFree and the blocking download recorded under ## Owed.

Landing this re-baselines every benchmark number we have

Do not compare a number taken before this merges with one taken after. The benchmark checkpoint qwen3.8-27b-nvfp4-mtp-sm121-r0b0tlab ships a generation_config.json containing {"do_sample": true, "temperature": 1.0, "top_k": 20, "top_p": 0.95}. Today we ignore it and sample over the full 248,320-token vocabulary while vLLM samples from top-20 — so our arm has been paying a per-token cost we inflict on ourselves, on a live benchmark asymmetry, and the tokens we emit are drawn from a different distribution than the oracle's. After this lands, an unparameterised request on that checkpoint samples top-20 / top-p-0.95 like vLLM's.

That is the correct behaviour and it is the point of #1985, but it moves the numbers. Our arm must be re-baselined at the merge commit, and any figure quoted across the boundary is comparing two different sampling configurations rather than two engines. --generation-config vllm restores the old resolution exactly if an old number ever needs reproducing.

What #1994 does and does not invalidate here

Nothing in this pull request or in the row's spec quotes a TTFT or a TPOT figure, so there is nothing to void. Checked rather than assumed: the only measured numbers either document carries are ~7.5 ms/token — cuda_sample.cu's own long-standing anchor for the single-thread greedy scan, a per-kernel decode-path attribution rather than a harness TPOT — and 16.9 tok/s, an end-to-end throughput figure from #1929. Throughput figures are unaffected by #1994. Everything else here is explicitly marked as prediction.

The acceptance measurement in ## Now is likewise safe, and the reason is in #1994's own spec rather than in its summary: it lists /v1/completions under what it does not touch, and scopes the defect to vllm bench serve --backend openai-chat. Should any arm be run on the chat backend instead, the residual bias is bounded and small — TTFT was understated by roughly one decode step, and TPOT = (E2E - TTFT)/(tokens - 1) spreads that over 63 tokens, so the greedy arm inflates by about 1.7 ms and the default arm by about 5.4 ms, overstating the greedy-versus-default gap by about 1.6%. That shifts the acceptance bar by well under a millisecond and changes no verdict.

Owed, with anchors

• src/vllm/v1/sample/sampler.cpp:344 allocates the [n, vocab] probs buffer through Backend::Alloc/Free (sampler.cpp:41,47 to src/vt/cuda/cuda_backend.cu:80-85) — a 30.3 MiB raw cudaMalloc plus a device-synchronising cudaFree every decode step at B=32.
• src/vllm/v1/sample/sampler.cpp:358 rs.download(...) is a blocking Synchronize every step, because the zero-sync device-resident path at sampler.cpp:456-461 requires sm.all_greedy and is structurally unreachable at temperature 1.0.
• src/vt/rocm/rocm_sample.hip::RandomSampleK carries the identical <<<n, 1>>> defect and its own copy of ExpNoise (rocm_sample.hip:38). Left untouched rather than changed on hardware this row cannot gate.
• ExpNoise computes -log(u) in double, where upstream's default is f32 (use_fp64_gumbel: bool = False). That is an unannotated widening of the kind .agents/porting.md says a token gate cannot see, and on a part with 1:64 f64 throughput it is the dominant remaining cost of the parallel kernel. Narrowing it changes which token is drawn, so it is a separate row with its own gate.
• --override-generation-config, the max_new_tokens server-wide output cap, and an offline LLM.get_default_sampling_params equivalent for the C ABI.

Closes #1984
Closes #1985

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: claude-code:claude-opus-5-1m [Claude Code]

[localai-bot]

PRE-CHECK RESULT — the premise holds, and it was understated by 12x

This PR asked for a pre-check before anything else: on a tree without the fix, speculation off, does the server default temperature cost 12-20 ms/token more than --temperature 0? And it set the falsifier itself: under 3 ms refutes the premise and this must not merge on a speed claim.

Run on dgx:gpu0 (GB10) idle under an rc lease, binary build12 = main @ 993b62e4e (pre-fix), artifact-gated. --num-blocks 3744 --max-num-seqs 4 --max-model-len 8192 --no-enable-prefix-caching, no --speculative-config, 852-token prompt, 64 output, warm request first, two samples per arm:

path	ms/token
--temperature 0 → device GreedyArgmax (sampler.cpp:456-461)	105.08, 104.89
server default 1.0 → vt::RandomSample (cuda_sample.cu:271-296)	341.12, 340.75

Difference: ~236 ms/token. Predicted 12-20. Falsifier 3. The premise is confirmed by roughly an order of magnitude more than forecast, and the run-to-run spread is under 0.4 ms in both arms.

In throughput terms the default path costs 3.25x — 2.93 tok/s against 9.53.

What the 236 ms is

RandomSampleKernel is launched <<<n, 1>>> with if (threadIdx.x != 0) return; — one thread per row, serially scanning all 248,320 vocabulary entries, evaluating two SplitMix64 rounds and an f64 log per element. 236 ms over 248,320 iterations is ~950 ns per element, which is the right order for one GPU thread doing that work.

The tree already recorded this shape as catastrophic eleven lines above the kernel: cuda_sample.cu:84-88 notes a single-thread serial scan of a ~151k vocab cost ~7.5 ms/token and is why greedy argmax was rewritten. The random path never got the same treatment. The measured 236 ms against that 7.5 ms anchor is consistent once the f64 log and the two 64-bit mixes are added.

Upstream is fully parallel — q.exponential_() then argmax(probs/q) (vllm/v1/sample/ops/topk_topp_sampler.py:199-204,426-433). This is a mirror obligation, not an optimisation.

Scope of the impact, stated precisely

This does not affect the #1574 DFlash2 ladder. As this PR established, GPUModelRunner::sample_tokens (runner.cpp:2347) branches on num_draft_tokens > 0 and returns the greedy-only RejectionSampler's output, so vt::RandomSample is never called with a speculative config at any temperature. Our spec-on c=1 row (24.70 out tok/s, TPOT 37.90 ms) is unaffected.

It affects every non-speculative request at the default temperature, which is the common serving configuration. vllm bench serve no longer sets temperature (vllm/benchmarks/serve.py:2073-2079), so any benchmark of this engine without a draft has been paying 3.25x.

It also compounds the correctness divergence in #2002: at temperature 1.0 under speculation we emit the target argmax where vLLM draws from the residual distribution.

What is still owed before merge

The A-vs-D measurement on the integrated tree, at the same recipe and predicted by this PR at 10-20 ms/token faster than A, recovering >=80% of the gap above, with its falsifier: do not merge on that evidence if D is less than 8 ms/token faster, or if D is slower than A on any arm by more than 1%. Given the gap is 236 ms rather than the 12-20 assumed, the "recover >=80%" criterion is the one to hold this to, not the absolute figure.

compute-sanitizer on the new kernel is also still owed, as this PR requested.

[localai-bot]

Success criterion, agreed and recorded BEFORE the run

I offered "recover >=80% of the measured gap". The author asked to be held to a tighter bar and gave a checkable reason, so the criterion is:

bar	residual allowed	arm C must reach
pass	<= 11.8 ms	<= 116.8 ms/token (>=95% recovery)
hard fail	> 23.6 ms	> 128.6 ms/token (<90% recovery)

Floor 104.99 ms/token (measured, --temperature 0), pre-fix 340.94 (measured), gap 235.95.

The reason 95% is defensible rather than optimistic: the temperature-0 arm already runs this exact reduction. GreedyArgmax at 104.99 ms/token is ArgmaxPartialKernel/ArgmaxFinalKernel over the same 248,320-wide row, so that reduction demonstrably does not cost tens of milliseconds. The fix makes the random path run the same reduction with a Gumbel score substituted for the raw logit; what remains on top is ComputeProbs (one block, three passes), the 993 KB probs cudaMalloc/cudaFree and the blocking download — the two items already listed under ## Owed.

Three arms, not two — the earlier framing was wrong

I had planned an A-vs-D comparison. That would have misattributed a cost, because D is not "A with a faster sampler": it also switches on the checkpoint's top_k=20 / top_p=0.95, adding an ApplyTopKTopP pass that arm A never ran (pre-fix, both filters resolved to disabled).

arm	configuration	judged against
A	pre-fix binary, server default temperature	the 340.94 ms/token baseline
C	--generation-config vllm, VT_FAST_RANDOM_SAMPLE=1	the >=95% bar above — this is the kernel measured alone, and it is the throughput headline
D	full parity, checkpoint defaults active	within 10 ms/token of C

D exceeding C by more than 10 ms/token does not block this row — D is the parity-correct configuration and matching upstream's distribution is what it is for — but it would open a follow-up on the sort-free pivot search at n=1.

Also being collected on the same runs

The residual becomes measurable for the first time. At 341 ms/token the serial scan was ~69% of the step, so the probs cudaMalloc/cudaFree (src/vllm/v1/sample/sampler.cpp:344, feeding a device-synchronising cudaFree) and the blocking rs.download (:358) were hidden underneath it. If C lands at 112 against the 105 floor, that ~7 ms is those two, and it deserves a row rather than being rounded away.

Token identity, which is the correctness half and is free to collect. A and C at the same seed must emit identical token ids — that is the bit-identical claim. A and D must emit different ids, because D samples from 20 candidates where A sampled from 248,320; identical ids there would mean the checkpoint defaults are not reaching the sampler at all.

Why the original estimate missed by 12x, recorded because the reasoning generalises

The 12-20 ms/token forecast scaled this file's own 7.5 ms/151k-vocab anchor, which is a load-bound scan, and assumed the RNG and log would partly overlap it. On a single lane they overlap nothing: it is a fully serial dependent chain through a software f64 log with emulated 64-bit multiplies. The measured 236 ms over 248,320 elements is ~0.95 us per element on one thread, which is what that chain costs. The post-fix estimate is a throughput calculation over 65,536 threads and is far better conditioned — which is precisely why the author named a falsifier they could actually fail.