From fd284ab2be5c48cbf93f0e9cb5d5892dd43f9f44 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 15:04:30 +0000 Subject: [PATCH 01/11] spec(SAMPLE-CORE): the checkpoint's sampling defaults, and the Gumbel draw that scans 248,320 tokens on one thread (#1984, #1985) On Qwen3.8-27B vLLM draws its next token from 20 candidates and we draw ours from 248,320, because `hf_config.cpp::ReadGenerationConfigEosIds` reads the checkpoint's `generation_config.json` for `eos_token_id` and discards 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, which makes this a correctness divergence on the workload the parity gate runs rather than a cosmetic default. The same fact explains the second defect. `RandomSampleKernel` is launched `<<>>` and walks the whole vocabulary on one lane, computing two SplitMix64 rounds and an f64 `log` per element. Eleven lines above it, this file 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. The spec commits the accepting evidence before it is measured: four same-binary arms separating the config read from the kernel, end to end on the leased box, with predicted values and with the two outcomes that would falsify the design rather than the tuning. A per-kernel figure is not evidence here; #1929 landed one and cost 16.9 tok/s end to end. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .agents/issue-index.md | 2 + .../sample-gen-config-and-parallel-gumbel.md | 248 ++++++++++++++++++ 2 files changed, 250 insertions(+) create mode 100644 .agents/specs/sample-gen-config-and-parallel-gumbel.md diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 7cd159742..b166ecdd1 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -761,3 +761,5 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#2040](https://github.com/mudler/vllm.cpp/issues/2040) | `ENG-HYBRID-PLACEMENT` | W3c: replace W3b's hand-written per-model round trip with ONE shared seam every architecture routes through, exploiting that every MoE block already has the same `(Dev, weights, params, [T,H] dh, T)` shape — and correct `docs/ENVIRONMENT.md` and `docs/FEATURES.md`, which W3b made false by leaving them saying the knobs place nothing | bug | | [#2046](https://github.com/mudler/vllm.cpp/issues/2046) | `ENG-HYBRID-PLACEMENT` | `qwen3_5.cpp` kept private `Dev`/`DBuf`/`MakeTensor`/`Reshape` copies instead of the shared `dense_device_glue.h` set — the off-framework divergence its own `ResidentWeight` comment records, where a repair reached 25 model files and not this one. The private types also had INTERNAL LINKAGE, which is what forced the MoE placement seam to carry a glue-templated second spelling; migrating collapses it back to one | bug | | [#2050](https://github.com/mudler/vllm.cpp/issues/2050) | `ENG-HYBRID-PLACEMENT` | Laguna's FFN is host-orchestrated token-at-a-time — per-token host rows, the router on the host through `MatmulNK`, and a host scalar combine loop — so a device-shaped MoE entry wrapping those loops would put it in the placement seam's wired list while moving nothing and adding a round trip: supported to read, a regression to measure. The real repair is a device-resident batched FFN, which is a model rework with a performance gate | gap | +| [#1984](https://github.com/mudler/vllm.cpp/issues/1984) | `SAMPLE-CORE` | `RandomSampleKernel` is launched `<<>>` and scans a 248,320-wide vocab on one thread per row, computing two `SplitMix64` rounds and an f64 `log` per element. Eleven lines above it the same file records that a single-block single-thread scan of a ~151k vocab cost ~7.5 ms/token, which is why greedy argmax was rewritten into `ArgmaxPartialKernel`/`ArgmaxFinalKernel`; the Gumbel draw never got that treatment. Upstream is whole-tensor (`vllm/v1/sample/ops/topk_topp_sampler.py::sample_with_exponential_noise`), so this is a mirror obligation. Reached by every non-greedy row through `ModelRunner::execute_model` -> `Sampler::forward` -> `vt::RandomSample`. Spec: [sample-gen-config-and-parallel-gumbel.md](specs/sample-gen-config-and-parallel-gumbel.md) | perf | +| [#1985](https://github.com/mudler/vllm.cpp/issues/1985) | `SAMPLE-CORE` | `generation_config.json` is read for `eos_token_id` only (`hf_config.cpp::ReadGenerationConfigEosIds`), so `Qwen/Qwen3.8-27B`'s shipped `top_k: 20` / `top_p: 0.95` never reach `SamplingParams` and `to_sampling_params` resolves omitted knobs straight to the neutral OpenAI defaults, which disable both filters. vLLM applies them through `ModelConfig.get_diff_sampling_param` -> `OpenAIServing*.default_sampling_params` -> `to_sampling_params`. Since `vllm bench serve` stopped sending `--temperature`, both engines sample at temperature 1.0 and vLLM draws from 20 candidates while we draw from 248,320: different sampling on two sides of a parity benchmark. Spec: [sample-gen-config-and-parallel-gumbel.md](specs/sample-gen-config-and-parallel-gumbel.md) | bug | diff --git a/.agents/specs/sample-gen-config-and-parallel-gumbel.md b/.agents/specs/sample-gen-config-and-parallel-gumbel.md new file mode 100644 index 000000000..8e197d792 --- /dev/null +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -0,0 +1,248 @@ +# `SAMPLE-CORE` — the checkpoint's sampling defaults, and the Gumbel draw that scans a 248,320-wide vocab on one thread + +Issues: [#1985](https://github.com/mudler/vllm.cpp/issues/1985) (checkpoint +sampling defaults are discarded), +[#1984](https://github.com/mudler/vllm.cpp/issues/1984) +(`RandomSampleKernel` is `<<>>`). +Owning row: `SAMPLE-CORE` ([engine matrix](../engine-matrix.md)). +Lifecycle: `ACTIVE`. +Oracle: vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98` (0.26.0.dev0), the +primary pin, read from the local checkout at that exact SHA. + +## Why the two are one row + +They are the same failure seen from two ends. On `Qwen/Qwen3.8-27B` vLLM draws +its next token from 20 candidates and we draw ours from 248,320, because we +never read the `top_k` the checkpoint ships. That is a correctness divergence on +the workload the parity gate runs, and it is also why the kernel that does the +drawing has four orders of magnitude more work than it needs. Reading the +config shrinks the *distribution*; it does not shrink the *scan*, because both +engines still walk the whole vocabulary. So the scan has to be parallelised as +well, and the two changes have to be measured together or each will be credited +with the other's effect. + +## Scope + +1. Parse `generation_config.json`'s sampling keys, mirror + `ModelConfig.get_diff_sampling_param`, and resolve them in + `to_sampling_params` with upstream's precedence. +2. Replace the single-thread `RandomSampleKernel` on CUDA with the two-pass + grid-strided argmax reduction that already sits eleven lines above it, with + **bit-identical output**. + +Out of scope, and named so a later measurement does not credit this row with +them: the `DeviceBuffer probs` per-step `cudaMalloc`/`cudaFree`, the blocking +`rs.download`, the all-greedy gate on the async fast path, the ROCm twin of the +kernel, and the f64-vs-f32 Gumbel dtype divergence. All are recorded under +`## Owed` with anchors. + +## Upstream anchors + +Read at `555967922` from `${VLLM_SOURCE}`. + +| What | Upstream | +|---|---| +| `generation_config` defaults to `"auto"` | `vllm/config/model.py:298` (`ModelConfig.generation_config`) | +| load the file, keep the non-default keys | `vllm/config/model.py::ModelConfig.try_get_generation_config` | +| narrow to six sampling keys, rename `max_new_tokens` | `vllm/config/model.py::ModelConfig.get_diff_sampling_param` | +| the server stores it | `vllm/entrypoints/openai/completion/serving.py:80`; `chat_completion/serving.py:174` | +| the neutral fallbacks | `vllm/entrypoints/openai/completion/protocol.py::CompletionRequest._DEFAULT_SAMPLING_PARAMS` | +| request wins, else checkpoint, else neutral | `vllm/entrypoints/openai/completion/protocol.py::CompletionRequest.to_sampling_params` | +| beam search resolves temperature the same way | `completion/protocol.py::CompletionRequest.to_beam_search_params` | +| the exponential/Gumbel draw | `vllm/v1/sample/ops/topk_topp_sampler.py::TopKTopPSampler.forward_native` | +| `probs.div_(q).argmax(-1)` | `vllm/v1/sample/ops/topk_topp_sampler.py::sample_with_exponential_noise` | +| the noise dtype (f32 by default) | `vllm/v1/sample/ops/topk_topp_sampler.py::empty_exponential_noise_like`; `vllm/v1/sample/sampler.py::Sampler.__init__` (`use_fp64_gumbel: bool = False`) | + +Local anchors: `src/vt/cuda/cuda_sample.cu::RandomSampleKernel`, +`::ArgmaxPartialKernel`, `::ArgmaxFinalKernel`; +`src/vllm/transformers_utils/hf_config.cpp::ReadGenerationConfigEosIds`; +`src/vllm/entrypoints/openai/protocol.cpp::CompletionRequest::to_sampling_params`. + +## The transformers version decides what `to_diff_dict` keeps, and it is not academic + +`try_get_generation_config` returns `GenerationConfig.to_diff_dict()`, which +drops every key equal to a bare `GenerationConfig()`'s value. Under +`transformers` 4.x those defaults were `temperature=1.0`, `top_k=50`, +`top_p=1.0`, so a checkpoint shipping `top_k: 50` would have been dropped and a +port that read the JSON literally would diverge. Under the pinned floor +`transformers >= 5.5.3` (`requirements/common.txt:10`) every sampling field of a +bare `GenerationConfig()` is `None`, so **every declared key survives the diff** +and reading the JSON literally is exact. + +Measured, not assumed: against the `transformers 5.3.0` on this host, +`GenerationConfig.from_pretrained(dir).to_diff_dict()` on a Qwen-shaped file +returns `{'repetition_penalty': 1.05, 'temperature': 1.0, 'top_k': 20, +'top_p': 0.95}` — nothing dropped. The mirror therefore parses the JSON +directly, and this paragraph is the reason it is allowed to. If the pin's +transformers floor ever moves back below 5.x, this row's parse becomes wrong and +has to grow the default table. + +`Qwen/Qwen3.8-27B`'s file, read live 2026-08-26 from +`https://huggingface.co/Qwen/Qwen3.8-27B/resolve/main/generation_config.json`: +`{"temperature": 1.0, "top_k": 20, "top_p": 0.95}` plus the token ids. Our +resolved values today are `temperature 1.0, top_k 0, top_p 1.0` — both filters +disabled. + +## Design — part 1, the sampling defaults + +- `include/vllm/config/generation.h` + `src/vllm/config/generation.cpp`: + `DefaultSamplingParams` (the six resolved optionals, `max_new_tokens` already + renamed to `max_tokens`) and `GetDiffSamplingParam(const HfConfig&, + const std::string& generation_config)`, where the selector takes upstream's + three forms: `"auto"` (the checkpoint's own file, the default), `"vllm"` (no + file, neutral defaults), or a directory holding a `generation_config.json`. +- `HfConfig` grows `generation_config_sampling`, filled by the same sibling read + that already produces `generation_config_eos_ids`. One file read, two + consumers. +- `to_sampling_params` and `to_beam_search_params` on both request types take a + `const DefaultSamplingParams*`; `nullptr` reproduces today's behaviour byte for + byte, which is what keeps every existing caller and test unchanged. +- Both serving handlers gain `set_default_sampling_params`, called from + `server_main.cpp` from `loaded->config()` and the new `--generation-config` + flag, and logged on startup the way upstream logs it. + +Precedence, mirrored exactly: an explicitly sent request field wins; an omitted +field takes the checkpoint value; if the checkpoint does not declare it, the +neutral OpenAI default. `"vllm"` restores today's behaviour on demand. + +## Design — part 2, the parallel Gumbel draw + +`score(row, j) = probs[row][j] / ExpNoise(seed, row, j)` and the answer is +`argmax_j score` with the lowest index winning a tie — which is exactly the +operator `ArgReduce` already implements for greedy argmax, and `ArgReduce` is +**order-independent** (it compares the true global index, not thread or block +order). So the same two-pass partition can carry the Gumbel score with no +change to what is selected. + +This is the load-bearing property of the whole change, so state it plainly: +every element's `score` is computed by the identical expression on the identical +device libm, and only the order in which those identical floats are combined +changes. The new kernel is therefore **bit-identical to the old one**, not +merely close, and the gate below asserts equality rather than agreement. + +- `include/vt/sample_common.h` (new): `SplitMix64`, `ExpNoise`, `ArgReduce`, + `kArgSentinel` and `ArgBlocksPerRow` as `__host__ __device__` inlines. Today + `SplitMix64`/`ExpNoise` are written out three times — `cpu_sample.cpp`, + `cuda_sample.cu`, `rocm_sample.hip` — and a divergence between any two of them + is silent. The CPU and CUDA copies are replaced by the header; ROCm is left + alone deliberately (see `## Owed`). +- `cuda_sample.cu`: the partial kernel is templated on a score functor, so + greedy argmax and the Gumbel draw share one reduction rather than growing a + second hand-written copy. Greedy's instantiation is the same source it runs + today. +- The legacy serial kernel stays, reachable as `VT_FAST_RANDOM_SAMPLE=0`, + mirroring the `VT_FAST_ARGMAX` lever the greedy rewrite kept. It exists so the + equality gate is a **same-binary A/B**, which `AGENTS.md` requires before any + performance result is accepted. +- The Gumbel partials get their own persistent scratch rather than sharing the + argmax one, because a mixed greedy/random batch runs both in one `sample()` + call and sharing would make correctness depend on stream ordering that nothing + in the type system enforces. + +## Tests + +| Gate | Where | Runs on | +|---|---|---| +| checkpoint `top_k`/`top_p`/`temperature` reach `SamplingParams` when the request omits them | `tests/vllm/test_openai_protocol.cpp` | CPU | +| an explicit request value still wins over the checkpoint's | same | CPU | +| a checkpoint key the JSON does not declare falls to the neutral default | same | CPU | +| `"vllm"` selector discards the file; a path selector reads another directory | `tests/vllm/test_hf_config.cpp` | CPU | +| `max_new_tokens` is renamed to `max_tokens` | same | CPU | +| the sibling parse keeps `eos_token_id` behaviour unchanged | same | CPU | +| `ArgReduce` is order-independent, including on ties | `tests/vt/test_ops_sample.cpp` | CPU | +| the two-pass partition over the production score and reduce equals the serial CPU reference over vocab 1/2/255/256/257/1000/248320, uniform, one-hot, all-equal (all ties), and top-k-masked rows | same | CPU | +| CUDA parallel == CUDA serial **exactly**, same binary, same shapes | same, `HasCuda`-guarded | GPU | +| the server wires the defaults (reachability) | `tests/vllm/entrypoints/test_server_defaults.cpp` | CPU | + +## Gates + +```sh +scripts/agent-preflight.sh +cmake --build build -j"$(nproc)" && ctest --test-dir build --output-on-failure +``` + +## The measurement this row does not get to skip + +`AGENTS.md` and [#1975](https://github.com/mudler/vllm.cpp/issues/1975) between +them settle it: a performance change here does not merge on a green compile, and +a per-kernel figure is not evidence. [#1929](https://github.com/mudler/vllm.cpp/issues/1929) +landed a 708 us -> 40 us top-k kernel with fresh review, mutation proofs and a +green CUDA build, and cost 16.9 tok/s end to end. + +So the accepting evidence is **end to end, on the leased box, run by the +operator**, and it is stated here before it is run. See `## Now` for the exact +request and the predicted values. + +## Owed + +- [#1984](https://github.com/mudler/vllm.cpp/issues/1984) also names two + adjacent costs this row does not fix: `src/vllm/v1/sample/sampler.cpp:344` + allocates the `[n, vocab]` probs buffer through `Backend::Alloc`/`Free` + (`sampler.cpp:41,47` -> `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; + and `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` is launched `<<>>` + and carries its own copy of `ExpNoise` (`rocm_sample.hip:38`). It is the same + defect as #1984 on a backend this row has no hardware to gate, so it is left + untouched rather than changed blind. +- Our `ExpNoise` computes `-log(u)` in **double** unconditionally. Upstream's + default is f32 (`use_fp64_gumbel: bool = False`), so this is an unannotated + widening of exactly the kind `.agents/porting.md` says a token gate cannot + see, and on a part with 1:64 f64 throughput it is also 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 and not a rider on this one. +- `--override-generation-config` (`vllm/config/model.py:305`), the + `override_max_tokens` server-wide output cap derived from `max_new_tokens` + (`completion/serving.py:81-86`), and an offline equivalent of + `LLM.get_default_sampling_params` (`vllm/entrypoints/llm.py:404`) for the C + ABI, which receives explicit `vllm_sampling_params` and has no "omitted" + state to fill. + +## Now + +Implementation and gates land together in one pull request with this spec (the +`AGENTS.md` default; no split case applies). + +**The end-to-end measurement requested of the operator, with its prediction, +before it is run.** + +Workload: `vllm bench serve` against our server and against the pinned vLLM on +the identical `Qwen3.8-27B` artifact, vLLM's production (graphed) configuration +as the denominator, **no `--temperature` flag on either side**, so both engines +resolve temperature 1.0 and both take the random-sampling path. This is the +configuration that made the divergence real, so it is the configuration that +has to judge the fix. + +Four arms, one binary, so each half of the row is attributable on its own: + +| arm | `--generation-config` | `VT_FAST_RANDOM_SAMPLE` | isolates | +|---|---|---|---| +| A (today) | `vllm` | `0` | the pre-change baseline | +| B | `auto` | `0` | the config read alone | +| C | `vllm` | `1` | the kernel alone | +| D (shipping default) | `auto` | `1` | both | + +Predicted, and marked as prediction: **D >= C > A**, with C - A the large term. +The serial kernel's cost is INFERRED from this file's own recorded anchor — a +single-thread scan of a ~151k vocab at ~7.5 ms/token — scaled by 1.64x vocab and +by one f64 `log` plus two 64-bit mixes per element, giving of order 12-20 +ms/step at B=32. Parallelising it should leave a sampler bounded by f64 `log` +throughput rather than by one lane, so the predicted recovery is **most of that +per-step cost**, i.e. of order 10 ms/step, and the prediction is stated as a +throughput floor rather than a ratio because the decode step's other terms are +not measured here. B - A is predicted **small and possibly negative** on +throughput: top-k 20 does not reduce the work either engine does, and the reason +to want it is that it makes the two engines sample the same distribution. + +Two things would falsify the design rather than the tuning, and both are worth +naming in advance: D materially slower than C would mean the top-k path costs +more than the distribution it saves, and D no faster than A would mean the +sampler was never the bottleneck and the per-step term lives in the +`cudaMalloc`/`cudaFree` and the blocking download recorded under `## Owed`. + +`compute-sanitizer` on the new kernel is **requested**, not optional, over the +CUDA equality gate: the change adds a second persistent device scratch and a new +grid geometry, and [#1958](https://github.com/mudler/vllm.cpp/issues/1958) is an +illegal memory access on this same sampler surface. From a7bf176ba13e6c891d235826f806bd06f492f22d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 15:39:08 +0000 Subject: [PATCH 02/11] feat(SAMPLE-CORE): the checkpoint's own sampling defaults reach an OpenAI request that omits them (#1985) `Qwen/Qwen3.8-27B` ships `top_k: 20` and `top_p: 0.95` in its generation_config.json. We read that file for `eos_token_id` and threw the rest away, and `to_sampling_params` said so in its own comment: "no server-provided default_sampling_params in T0". Every omitted knob resolved to the neutral OpenAI default, and `kDefaultTopK = 0` / `kDefaultTopP = 1.0` both mean disabled. Since `vllm bench serve` stopped sending `--temperature`, both engines resolve 1.0 and both sample, so vLLM has been drawing from 20 candidates while we drew from all 248,320 -- different sampling on two sides of the parity benchmark. The same sibling read now also yields the six keys ModelConfig.get_diff_sampling_param narrows, `GetDiffSamplingParam` mirrors that narrowing and vLLM's `--generation-config auto|vllm|` selector, and both serving 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 default. `nullptr` and an empty DefaultSamplingParams both reproduce the old resolution exactly, which is what `--generation-config vllm` gives a user who wants it back. Red first, on the stubbed resolution: 3 cases and 9 assertions of test_openai_protocol failed for the intended reason before the resolution was written. One expectation had to be corrected rather than the code -- temperature 0 clears top_p/top_k/min_p in `__post_init__` upstream and here, so a checkpoint top_k cannot survive a greedy request, and that interaction is now pinned by its own case. Reachability is gated at two hops and honest about the third. The handler hop is entered through `create_completion` / `create_chat_completion` over the synthetic engine, with the defaults produced by the real `LoadHfConfig` and `GetDiffSamplingParam` off a real generation_config.json, so reverting either call site to `to_sampling_params()` turns it red. The CLI hop re-execs the real `VllmServerMain`. The two-line join in server_main sits after model load, and no committed generative checkpoint fixture can reach it; that is stated here rather than papered over. `--override-generation-config`, the `max_new_tokens` server-wide output cap, and an offline `get_default_sampling_params` for the C ABI are recorded under `## Owed` in the row's spec. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- CMakeLists.txt | 1 + docs/reference/server.md | 1 + include/vllm/config/generation.h | 54 ++++++ include/vllm/entrypoints/openai/protocol.h | 23 ++- .../vllm/entrypoints/openai/serving_chat.h | 11 ++ .../entrypoints/openai/serving_completion.h | 18 ++ include/vllm/sampling_params.h | 34 ++++ include/vllm/transformers_utils/hf_config.h | 43 +++++ src/vllm/config/generation.cpp | 71 ++++++++ src/vllm/entrypoints/openai/protocol.cpp | 70 ++++++-- src/vllm/entrypoints/openai/server_main.cpp | 36 ++++ src/vllm/entrypoints/openai/serving_chat.cpp | 6 +- .../entrypoints/openai/serving_completion.cpp | 6 +- src/vllm/transformers_utils/hf_config.cpp | 40 +++++ tests/CMakeLists.txt | 1 + tests/vllm/config/test_generation_config.cpp | 160 +++++++++++++++++ .../vllm/entrypoints/openai/test_protocol.cpp | 161 +++++++++++++++++ .../openai/test_serve_recipe_args.cpp | 25 +++ .../vllm/entrypoints/openai/test_serving.cpp | 166 ++++++++++++++++++ tests/vllm/test_hf_config.cpp | 97 ++++++++++ 20 files changed, 999 insertions(+), 25 deletions(-) create mode 100644 include/vllm/config/generation.h create mode 100644 src/vllm/config/generation.cpp create mode 100644 tests/vllm/config/test_generation_config.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index d87c63f3f..e5ea455ae 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -737,6 +737,7 @@ add_library(vllm STATIC src/vllm/config/cache.cpp src/vllm/config/scheduler.cpp src/vllm/config/device.cpp + src/vllm/config/generation.cpp src/vllm/config/kv_transfer.cpp src/vllm/config/offload.cpp src/vllm/config/weight_residency.cpp diff --git a/docs/reference/server.md b/docs/reference/server.md index c900229a7..5689c98ad 100644 --- a/docs/reference/server.md +++ b/docs/reference/server.md @@ -200,6 +200,7 @@ a stop token early. | `--enable-radix-attention` / `--disable-radix-attention` | model default | SGLang-named alias for the prefix-cache toggle | | `--enable-jump-forward` | off | Jump-forward decoding for structured output (token-unique subset) | | `--enable-force-include-usage` | off | Force the usage block in responses | +| `--generation-config auto\|vllm\|` | `auto` | Where the server's default sampling parameters come from. `auto` takes `temperature`, `top_k`, `top_p`, `min_p` and `repetition_penalty` from the checkpoint's own `generation_config.json`, and a request that OMITS one of those knobs then gets the checkpoint's value; a request that sends one still wins. `vllm` loads no file and keeps the neutral OpenAI defaults. A directory path reads a `generation_config.json` from there instead. `Qwen/Qwen3.8-27B` ships `top_k: 20` and `top_p: 0.95`, so under `auto` an unparameterised request samples from 20 candidates rather than all 248,320. The server prints what it resolved at startup. `--override-generation-config` is unavailable ([#1985](https://github.com/mudler/vllm.cpp/issues/1985)). | | `--tool-call-parser ` | `hermes` | Select one of 42 registered names across 38 dialect families. `auto` detects the dialect from the chat template, and `none` disables parsing. Gemma-4 accepts wrapped or bare text calls. Inkling requests require `"skip_special_tokens": false`. The `inkling` reasoning parser is unavailable. | | `--reasoning-parser ` | `none` | Select `think_auto`, `deepseek_r1`, `deepseek_v3`, `holo2`, `mistral`, `minimax_m2`, `minimax_m2_append_think`, `step3`, `olmo3`, `muse_glimmer`, `qwen3`, or `mimo`. `auto` detects from the template. The `qwen3` and `mimo` names share an adapter. | | `--kv-transfer-config ''` | (unset) | External KV connector, same JSON as vLLM's flag. See [the KV offload guide](../KV-OFFLOAD.md) | diff --git a/include/vllm/config/generation.h b/include/vllm/config/generation.h new file mode 100644 index 000000000..e231ef55e --- /dev/null +++ b/include/vllm/config/generation.h @@ -0,0 +1,54 @@ +// Ported from: vllm/config/model.py @ 5559679229bc961848b121ccdeaa8fa5d79bec98 +// (ModelConfig.generation_config, ModelConfig.try_get_generation_config, +// ModelConfig.get_diff_sampling_param). +// +// SCOPE: turning a checkpoint's generation_config.json into the server-wide +// sampling defaults an OpenAI request resolves against. Upstream owns this on +// ModelConfig; this tree has no ModelConfig, so the narrowing is a free +// function over the already-parsed HfConfig, which is where the file's keys +// land (transformers_utils/hf_config.h). +// +// DEFERRED, and tracked under `## Owed` in +// .agents/specs/sample-gen-config-and-parallel-gumbel.md: +// - --override-generation-config (config/model.py:305), the dict that is +// merged over the file's values before the narrowing. +// - the server-wide output cap upstream derives from max_new_tokens +// (completion/serving.py:81-86). max_tokens is CARRIED here so a later row +// has it, and no caller reads it yet. +#ifndef VLLM_CONFIG_GENERATION_H_ +#define VLLM_CONFIG_GENERATION_H_ + +#include + +#include "vllm/sampling_params.h" +#include "vllm/transformers_utils/hf_config.h" + +namespace vllm { + +// The three forms vLLM's --generation-config takes (config/model.py:298-304). +// "auto" is the DEFAULT and means the checkpoint's own file. +inline constexpr const char* kGenerationConfigAuto = "auto"; +// "vllm" means load no file at all and use vLLM's neutral defaults. +inline constexpr const char* kGenerationConfigNone = "vllm"; + +// ModelConfig.get_diff_sampling_param (config/model.py). `generation_config` is +// the --generation-config selector: "auto" (the checkpoint's own sibling file, +// already parsed into `config`), "vllm" (no file), or a DIRECTORY path holding +// a generation_config.json. +// +// A directory that does not exist, or holds no readable generation_config.json, +// yields an empty result rather than throwing, because upstream's loader +// returns {} on every failure path and never raises. +DefaultSamplingParams GetDiffSamplingParam( + const HfConfig& config, + const std::string& generation_config = kGenerationConfigAuto); + +// The same narrowing over an already-read file, so a caller that has the six +// keys in hand does not have to reconstruct an HfConfig around them. This is +// where max_new_tokens becomes max_tokens. +DefaultSamplingParams NarrowGenerationConfigSampling( + const GenerationConfigSampling& sampling); + +} // namespace vllm + +#endif // VLLM_CONFIG_GENERATION_H_ diff --git a/include/vllm/entrypoints/openai/protocol.h b/include/vllm/entrypoints/openai/protocol.h index 5fe7cf3d6..aba68bbdc 100644 --- a/include/vllm/entrypoints/openai/protocol.h +++ b/include/vllm/entrypoints/openai/protocol.h @@ -266,12 +266,23 @@ struct CompletionRequest { // PARAMS), then runs PostInit(). `default_max_tokens` is the serving-resolved // fallback used only when the request omits max_tokens (Task 2 supplies the // model-derived value; unset => our SamplingParams default). + // `defaults` is the server-wide DefaultSamplingParams derived from the + // checkpoint's generation_config.json (config/generation.h). Upstream's + // resolution order, mirrored exactly: an explicitly sent request field wins; + // an omitted one takes the checkpoint value; only if the checkpoint declares + // nothing does the neutral OpenAI default apply + // (CompletionRequest._DEFAULT_SAMPLING_PARAMS). nullptr means no + // server-provided defaults, and reproduces the pre-#1985 behaviour byte for + // byte, which is what keeps every existing caller unchanged. SamplingParams to_sampling_params( - std::optional default_max_tokens = std::nullopt) const; + std::optional default_max_tokens = std::nullopt, + const DefaultSamplingParams* defaults = nullptr) const; // to_beam_search_params — completion/protocol.py:260. beam_width == n, the - // resolved max_tokens, ignore_eos, temperature (None => 1.0) and length_penalty. - vllm::BeamSearchParams to_beam_search_params(int max_tokens) const; + // resolved max_tokens, ignore_eos, temperature (None => the checkpoint's, then + // 1.0) and length_penalty. + vllm::BeamSearchParams to_beam_search_params( + int max_tokens, const DefaultSamplingParams* defaults = nullptr) const; }; // Ported from: vllm/entrypoints/openai/completion/protocol.py:580-584 @@ -511,10 +522,12 @@ struct ChatCompletionRequest { // to_sampling_params — chat_completion/protocol.py:585. See CompletionRequest. SamplingParams to_sampling_params( - std::optional default_max_tokens = std::nullopt) const; + std::optional default_max_tokens = std::nullopt, + const DefaultSamplingParams* defaults = nullptr) const; // to_beam_search_params — chat_completion/protocol.py:589. See CompletionRequest. - vllm::BeamSearchParams to_beam_search_params(int max_tokens) const; + vllm::BeamSearchParams to_beam_search_params( + int max_tokens, const DefaultSamplingParams* defaults = nullptr) const; }; // Ported from: vllm/entrypoints/openai/chat_completion/protocol.py:94 diff --git a/include/vllm/entrypoints/openai/serving_chat.h b/include/vllm/entrypoints/openai/serving_chat.h index 9d2c7f5fa..bae447a3b 100644 --- a/include/vllm/entrypoints/openai/serving_chat.h +++ b/include/vllm/entrypoints/openai/serving_chat.h @@ -231,6 +231,14 @@ class OpenAIServingChat { beam_eos_token_id_ = eos_token_id; } + // See OpenAIServingCompletion::set_default_sampling_params (#1985). + void set_default_sampling_params(vllm::DefaultSamplingParams defaults) { + default_sampling_params_ = std::move(defaults); + } + const vllm::DefaultSamplingParams& default_sampling_params() const { + return default_sampling_params_; + } + // Attach the multimodal chat seam (see MultiModalChatFn). Unset (default) // keeps the text-only path byte-identical. When set AND a request carries a mm // content part, create_chat_completion routes the request through the engine @@ -279,6 +287,9 @@ class OpenAIServingChat { // unavailable on this handler. const vllm::tok::Tokenizer* beam_tokenizer_ = nullptr; std::optional beam_eos_token_id_; + // See set_default_sampling_params. Empty => every knob falls to the neutral + // OpenAI default, byte-identical to the behaviour before #1985. + vllm::DefaultSamplingParams default_sampling_params_; // Multimodal chat seam (see set_multimodal_chat_fn). Null => the text-only // path runs unchanged (mm parts drop to the joined-text content). MultiModalChatFn mm_chat_fn_; diff --git a/include/vllm/entrypoints/openai/serving_completion.h b/include/vllm/entrypoints/openai/serving_completion.h index 98e960fa8..3b356a3f2 100644 --- a/include/vllm/entrypoints/openai/serving_completion.h +++ b/include/vllm/entrypoints/openai/serving_completion.h @@ -96,6 +96,21 @@ class OpenAIServingCompletion { beam_eos_token_id_ = eos_token_id; } + // The server-wide sampling defaults from the checkpoint's + // generation_config.json (#1985), mirroring + // `self.default_sampling_params = self.model_config.get_diff_sampling_param()` + // in {completion,chat_completion}/serving.py. Unset (the default) means no + // server-provided defaults, which is the pre-#1985 resolution and is exactly + // what `--generation-config vllm` resolves to. Called from server_main once at + // startup; this handler is the ONLY place the value is applied, so deleting + // that call makes the whole feature unreachable and the reachability gate red. + void set_default_sampling_params(vllm::DefaultSamplingParams defaults) { + default_sampling_params_ = std::move(defaults); + } + const vllm::DefaultSamplingParams& default_sampling_params() const { + return default_sampling_params_; + } + private: v1::LLMEngine* sync_engine_ = nullptr; v1::AsyncLLM* async_engine_ = nullptr; @@ -105,6 +120,9 @@ class OpenAIServingCompletion { // unavailable on this handler. const vllm::tok::Tokenizer* beam_tokenizer_ = nullptr; std::optional beam_eos_token_id_; + // See set_default_sampling_params. Empty => every knob falls to the neutral + // OpenAI default, byte-identical to the behaviour before #1985. + vllm::DefaultSamplingParams default_sampling_params_; // Monotonic request counter — the request_id is "cmpl-". Upstream // uses random_uuid() (serving/engine/serving.py:_base_request_id); no // random/uuid is wired at T0, so a counter stands in (id uniqueness only). diff --git a/include/vllm/sampling_params.h b/include/vllm/sampling_params.h index a86a7baf1..5709c6386 100644 --- a/include/vllm/sampling_params.h +++ b/include/vllm/sampling_params.h @@ -132,6 +132,40 @@ struct StructuredOutputsParams { // Sampling parameters for text generation (T0 field subset). Defaults match // upstream SamplingParams exactly. +// Ported from: the dict vllm/config/model.py::ModelConfig.get_diff_sampling_param +// returns, and that vllm/entrypoints/openai/*/serving.py stores as +// `self.default_sampling_params`. +// +// The SERVER-WIDE sampling defaults a checkpoint's own generation_config.json +// asks for. Every field is optional because "the checkpoint declared nothing" +// and "the checkpoint declared the neutral value" resolve differently: an unset +// field falls through to the OpenAI neutral default, a set one does not. That +// distinction is the whole rule, so it is carried in the type rather than in a +// sentinel. +// +// `max_tokens` is already renamed from the file's `max_new_tokens`, exactly +// where upstream renames it (get_diff_sampling_param, "Huggingface definition +// of max_new_tokens is equivalent to vLLM's max_tokens"). +struct DefaultSamplingParams { + std::optional repetition_penalty; + std::optional temperature; + std::optional top_k; + std::optional top_p; + std::optional min_p; + std::optional max_tokens; + + bool empty() const { + return !repetition_penalty.has_value() && !temperature.has_value() && + !top_k.has_value() && !top_p.has_value() && !min_p.has_value() && + !max_tokens.has_value(); + } + + // The set fields as `{key: value, ...}`, for the startup line upstream logs + // ("Default vLLM sampling parameters have been overridden by ..."). Empty + // string when nothing is set. + std::string ToString() const; +}; + struct SamplingParams { // Number of outputs to return for the given prompt request. int n = 1; diff --git a/include/vllm/transformers_utils/hf_config.h b/include/vllm/transformers_utils/hf_config.h index afacdd8e7..e3ee7dd38 100644 --- a/include/vllm/transformers_utils/hf_config.h +++ b/include/vllm/transformers_utils/hf_config.h @@ -66,6 +66,29 @@ struct RopeParameters { // `text_config` sub-dict; LoadHfConfig resolves that nested object as the source // of the text fields, mirroring upstream PretrainedConfig.get_text_config(). // `model_type` and `architectures` are always read from the top-level wrapper. +// The sampling keys of ONE generation_config.json, as +// ModelConfig.try_get_generation_config (vllm/config/model.py) surfaces them. +// +// Upstream returns `GenerationConfig.to_diff_dict()`, i.e. only the keys that +// differ from a bare `GenerationConfig()`. The pinned floor is +// `transformers >= 5.5.3` (vLLM requirements/common.txt:10), and at 5.x every +// sampling field of a bare GenerationConfig is None, so every key the JSON +// declares survives the diff and reading the file literally is exact. Measured +// against transformers 5.3.0 on a Qwen-shaped file: repetition_penalty 1.05, +// temperature 1.0, top_k 20 and top_p 0.95 all survive. If the pin's floor ever +// moves back to transformers 4.x -- whose defaults were temperature 1.0, +// top_k 50, top_p 1.0 -- this parse becomes wrong and has to grow that table. +struct GenerationConfigSampling { + std::optional repetition_penalty; + std::optional temperature; + std::optional top_k; + std::optional top_p; + std::optional min_p; + // Upstream renames this to max_tokens when it narrows; kept under its HF name + // here because this struct is the FILE, not the narrowed result. + std::optional max_new_tokens; +}; + struct HfConfig { std::string model_type; std::vector architectures; @@ -126,12 +149,32 @@ struct HfConfig { // secondary stop ids gated on ignore_eos. Empty when the file is absent, // unparseable, or carries no eos_token_id. std::vector generation_config_eos_ids; + // The six SAMPLING keys of the same sibling generation_config.json, which + // upstream reads through the same one file read and then narrows in + // ModelConfig.get_diff_sampling_param (config/model.py). Every field is unset + // when the file is absent, unparseable, or does not declare that key, which + // is what makes "the checkpoint said nothing" distinguishable from "the + // checkpoint said the neutral value" -- the distinction the whole defaulting + // rule turns on. See include/vllm/config/generation.h for the narrowing and + // for the --generation-config selector that can point somewhere else. + GenerationConfigSampling generation_config_sampling; }; // Loads and parses `path`. Throws std::runtime_error (message includes the // path) on missing file, malformed JSON, or missing required fields. HfConfig LoadHfConfig(const std::string& path); +// Reads the sampling keys of the generation_config.json at `path`. Mirrors +// ModelConfig.try_get_generation_config's failure polarity exactly: a missing +// file, unparseable JSON, a non-object document, a null value and a +// wrong-typed value each leave the corresponding field unset rather than +// throwing, because upstream's loader returns {} and never raises. +// +// Exposed because --generation-config can name a DIRECTORY other than the +// checkpoint's, which is the one case the sibling read on HfConfig cannot +// serve. +GenerationConfigSampling ReadGenerationConfigSamplingFile(const std::string& path); + // The same parse, from a config object already in hand. `source` appears in // every error message exactly where the path would, so a refusal still names // where the config came from. diff --git a/src/vllm/config/generation.cpp b/src/vllm/config/generation.cpp new file mode 100644 index 000000000..6aa31006e --- /dev/null +++ b/src/vllm/config/generation.cpp @@ -0,0 +1,71 @@ +// Ported from: vllm/config/model.py @ 5559679229bc961848b121ccdeaa8fa5d79bec98. +// See include/vllm/config/generation.h for scope and deferrals. +#include "vllm/config/generation.h" + +#include + +namespace vllm { + +DefaultSamplingParams NarrowGenerationConfigSampling( + const GenerationConfigSampling& s) { + // get_diff_sampling_param's `available_params` list, in its order. Upstream + // keeps a key only when `config.get(p) is not None`, which is exactly what an + // unset optional carries here, so the copy is unconditional. + DefaultSamplingParams out; + out.repetition_penalty = s.repetition_penalty; + out.temperature = s.temperature; + out.top_k = s.top_k; + out.top_p = s.top_p; + out.min_p = s.min_p; + // "Huggingface definition of max_new_tokens is equivalent to vLLM's + // max_tokens" -- the rename happens here and nowhere else. + out.max_tokens = s.max_new_tokens; + return out; +} + +DefaultSamplingParams GetDiffSamplingParam(const HfConfig& config, + const std::string& generation_config) { + // `src == "vllm"` -> `config = {}` (config/model.py). No file is read and + // every knob falls through to the neutral OpenAI default, which is exactly + // the behaviour this tree had before the checkpoint's file was read at all. + if (generation_config == kGenerationConfigNone) return {}; + if (generation_config == kGenerationConfigAuto) { + return NarrowGenerationConfigSampling(config.generation_config_sampling); + } + // A directory path: try_get_generation_config(self.generation_config, ...) + // loads generation_config.json out of the named folder rather than the + // checkpoint's. A trailing separator is tolerated so `--generation-config + // /some/dir/` behaves like `/some/dir`. + std::string dir = generation_config; + while (dir.size() > 1 && (dir.back() == '/' || dir.back() == '\\')) dir.pop_back(); + return NarrowGenerationConfigSampling( + ReadGenerationConfigSamplingFile(dir + "/generation_config.json")); +} + +std::string DefaultSamplingParams::ToString() const { + // The dict repr upstream interpolates into its "Default vLLM sampling + // parameters have been overridden by %s: `%s`" warning. Order follows + // get_diff_sampling_param's available_params so two engines' logs line up. + std::ostringstream os; + bool first = true; + const auto put = [&](const char* key, const std::string& value) { + if (!first) os << ", "; + first = false; + os << "'" << key << "': " << value; + }; + const auto num = [](double v) { + std::ostringstream t; + t << v; + return t.str(); + }; + if (repetition_penalty.has_value()) put("repetition_penalty", num(*repetition_penalty)); + if (temperature.has_value()) put("temperature", num(*temperature)); + if (top_k.has_value()) put("top_k", std::to_string(*top_k)); + if (top_p.has_value()) put("top_p", num(*top_p)); + if (min_p.has_value()) put("min_p", num(*min_p)); + if (max_tokens.has_value()) put("max_tokens", std::to_string(*max_tokens)); + if (first) return {}; + return "{" + os.str() + "}"; +} + +} // namespace vllm diff --git a/src/vllm/entrypoints/openai/protocol.cpp b/src/vllm/entrypoints/openai/protocol.cpp index 8971d0772..44f215fff 100644 --- a/src/vllm/entrypoints/openai/protocol.cpp +++ b/src/vllm/entrypoints/openai/protocol.cpp @@ -597,19 +597,49 @@ void from_json(const nlohmann::json& j, ChatCompletionRequest& r) { // to_sampling_params // --------------------------------------------------------------------------- +// completion/protocol.py:288-317 (CompletionRequest.to_sampling_params) and its +// chat twin, as one function per knob: +// +// if (value := request.field) is None: +// value = default_sampling_params.get(name, _DEFAULT_SAMPLING_PARAMS[name]) +// +// The three-way order is the whole point and is easy to get subtly wrong, so it +// is written once: an explicitly SENT request value wins, an OMITTED one takes +// the checkpoint's server-wide default, and only a knob neither of them +// declares reaches the neutral OpenAI value. A request that explicitly sends +// the neutral value (top_k = 0, "turn top-k off") is a SENT value and is not +// overridden -- which is why the request fields are optionals upstream and here. +template +T ResolveSamplingKnob(const std::optional& requested, + const DefaultSamplingParams* defaults, Member member, + T neutral) { + if (requested.has_value()) return *requested; + if (defaults != nullptr && (defaults->*member).has_value()) { + return static_cast(*(defaults->*member)); + } + return neutral; +} + + SamplingParams CompletionRequest::to_sampling_params( - std::optional default_max_tokens) const { + std::optional default_max_tokens, const DefaultSamplingParams* defaults) const { // completion/protocol.py:260. None sampling knobs resolve to // _DEFAULT_SAMPLING_PARAMS (no server-provided default_sampling_params in T0). SamplingParams sp; sp.n = n; sp.presence_penalty = presence_penalty; sp.frequency_penalty = frequency_penalty; - sp.repetition_penalty = repetition_penalty.value_or(kDefaultRepetitionPenalty); - sp.temperature = temperature.value_or(kDefaultTemperature); - sp.top_p = top_p.value_or(kDefaultTopP); - sp.top_k = top_k.value_or(kDefaultTopK); - sp.min_p = min_p.value_or(kDefaultMinP); + sp.repetition_penalty = ResolveSamplingKnob( + repetition_penalty, defaults, &DefaultSamplingParams::repetition_penalty, + kDefaultRepetitionPenalty); + sp.temperature = ResolveSamplingKnob( + temperature, defaults, &DefaultSamplingParams::temperature, kDefaultTemperature); + sp.top_p = + ResolveSamplingKnob(top_p, defaults, &DefaultSamplingParams::top_p, kDefaultTopP); + sp.top_k = + ResolveSamplingKnob(top_k, defaults, &DefaultSamplingParams::top_k, kDefaultTopK); + sp.min_p = + ResolveSamplingKnob(min_p, defaults, &DefaultSamplingParams::min_p, kDefaultMinP); sp.seed = seed; sp.stop = stop; sp.stop_token_ids = stop_token_ids; @@ -649,17 +679,23 @@ SamplingParams CompletionRequest::to_sampling_params( } SamplingParams ChatCompletionRequest::to_sampling_params( - std::optional default_max_tokens) const { + std::optional default_max_tokens, const DefaultSamplingParams* defaults) const { // chat_completion/protocol.py:585. SamplingParams sp; sp.n = n.value_or(1); sp.presence_penalty = presence_penalty; sp.frequency_penalty = frequency_penalty; - sp.repetition_penalty = repetition_penalty.value_or(kDefaultRepetitionPenalty); - sp.temperature = temperature.value_or(kDefaultTemperature); - sp.top_p = top_p.value_or(kDefaultTopP); - sp.top_k = top_k.value_or(kDefaultTopK); - sp.min_p = min_p.value_or(kDefaultMinP); + sp.repetition_penalty = ResolveSamplingKnob( + repetition_penalty, defaults, &DefaultSamplingParams::repetition_penalty, + kDefaultRepetitionPenalty); + sp.temperature = ResolveSamplingKnob( + temperature, defaults, &DefaultSamplingParams::temperature, kDefaultTemperature); + sp.top_p = + ResolveSamplingKnob(top_p, defaults, &DefaultSamplingParams::top_p, kDefaultTopP); + sp.top_k = + ResolveSamplingKnob(top_k, defaults, &DefaultSamplingParams::top_k, kDefaultTopK); + sp.min_p = + ResolveSamplingKnob(min_p, defaults, &DefaultSamplingParams::min_p, kDefaultMinP); sp.seed = seed; sp.stop = stop; sp.stop_token_ids = stop_token_ids; @@ -698,25 +734,27 @@ SamplingParams ChatCompletionRequest::to_sampling_params( // --------------------------------------------------------------------------- vllm::BeamSearchParams CompletionRequest::to_beam_search_params( - int max_tokens_in) const { + int max_tokens_in, const DefaultSamplingParams* defaults) const { // completion/protocol.py:260-279. beam_width == n; temperature None => 1.0. vllm::BeamSearchParams bp; bp.beam_width = n; bp.max_tokens = max_tokens_in; bp.ignore_eos = ignore_eos; - bp.temperature = temperature.value_or(kDefaultTemperature); + bp.temperature = ResolveSamplingKnob( + temperature, defaults, &DefaultSamplingParams::temperature, kDefaultTemperature); bp.length_penalty = length_penalty; return bp; } vllm::BeamSearchParams ChatCompletionRequest::to_beam_search_params( - int max_tokens_in) const { + int max_tokens_in, const DefaultSamplingParams* defaults) const { // chat_completion/protocol.py:589-606. beam_width == n; temperature None => 1.0. vllm::BeamSearchParams bp; bp.beam_width = n.value_or(1); bp.max_tokens = max_tokens_in; bp.ignore_eos = ignore_eos; - bp.temperature = temperature.value_or(kDefaultTemperature); + bp.temperature = ResolveSamplingKnob( + temperature, defaults, &DefaultSamplingParams::temperature, kDefaultTemperature); bp.length_penalty = length_penalty; return bp; } diff --git a/src/vllm/entrypoints/openai/server_main.cpp b/src/vllm/entrypoints/openai/server_main.cpp index 09675c22a..e40eeb055 100644 --- a/src/vllm/entrypoints/openai/server_main.cpp +++ b/src/vllm/entrypoints/openai/server_main.cpp @@ -78,6 +78,7 @@ #include "vllm/entrypoints/openai/chat_mm.h" #include "vllm/entrypoints/openai/request_logger.h" #include "vllm/entrypoints/openai/serving_chat.h" +#include "vllm/config/generation.h" #include "vllm/entrypoints/openai/serving_completion.h" #include "vllm/entrypoints/openai/serving_models.h" #include "vllm/v1/metrics/loggers.h" @@ -312,6 +313,15 @@ struct Args { // invocation that names neither flag is therefore unchanged, byte for byte. // "auto" opts into the chat-template detection the C ABI uses. std::string tool_call_parser = "hermes"; + // vLLM's --generation-config (config/model.py:298-304). "auto" (the DEFAULT) + // takes the checkpoint's own generation_config.json sampling keys as the + // server-wide request defaults; "vllm" loads no file and keeps vLLM's neutral + // defaults; a DIRECTORY path reads a generation_config.json from there. + // + // "auto" is what upstream defaults to, so it is what a parity benchmark + // compares against, and #1985 is what taking anything else cost. The escape + // hatch is exactly the one upstream's own warning names. + std::string generation_config = "auto"; std::string reasoning_parser = "none"; // vLLM's --enable-auto-tool-choice (cli_args.py:105, default False). Accepted // and INERT here (see kAcceptedInertArgs below) but still RECORDED, because @@ -631,6 +641,8 @@ Args ParseArgs(int argc, char** argv) { Usage(argv[0], 2); } a.enable_jump_forward = flag == "--enable-jump-forward"; + } else if (flag == "--generation-config") { + a.generation_config = NextArg(argc, argv, i, argv[0]); } else if (flag == "--tool-call-parser") { a.tool_call_parser = NextArg(argc, argv, i, argv[0]); } else if (flag == "--reasoning-parser") { @@ -1436,6 +1448,30 @@ int VllmServerMain(int argc, char** argv) { completion.set_beam_search_tokenizer(&tokenizer, beam_eos); chat.set_beam_search_tokenizer(&tokenizer, beam_eos); + // ── #1985: the checkpoint's own sampling defaults. Upstream resolves these + // once per server (ModelConfig.get_diff_sampling_param -> + // OpenAIServing*.default_sampling_params) and every request that OMITS a + // knob then takes them. THIS IS THE PRODUCTION CALL SITE: without these two + // lines the whole feature is unreachable, and the reachability gate in + // tests/vllm/entrypoints/openai/test_serving_generation_defaults.cpp goes + // red when either is deleted. ───────────────────────────────────────────── + const vllm::DefaultSamplingParams default_sampling_params = + vllm::GetDiffSamplingParam(loaded->config(), args.generation_config); + completion.set_default_sampling_params(default_sampling_params); + chat.set_default_sampling_params(default_sampling_params); + if (!default_sampling_params.empty()) { + // Upstream logs the same thing at startup and names the source, so the + // two engines' logs can be lined up: "Default vLLM sampling parameters + // have been overridden by the model's generation_config.json". + std::cerr << "server: default sampling params from " + << (args.generation_config == "auto" + ? "the model's generation_config.json" + : args.generation_config) + << ": " << default_sampling_params.ToString() + << " (pass --generation-config vllm to use vLLM's neutral" + " defaults instead)\n"; + } + // ── MM-SERVE-E2E: wire the multimodal chat seam for image-capable models. // When the model dir carries a preprocessor_config.json the Qwen3-VL image // processor loads, we construct the seam body (MakeQwen3VLImageChatFn) so an diff --git a/src/vllm/entrypoints/openai/serving_chat.cpp b/src/vllm/entrypoints/openai/serving_chat.cpp index 838626010..05757a5a1 100644 --- a/src/vllm/entrypoints/openai/serving_chat.cpp +++ b/src/vllm/entrypoints/openai/serving_chat.cpp @@ -734,7 +734,8 @@ ChatCompletionResult OpenAIServingChat::create_chat_completion( request.max_completion_tokens.has_value() ? *request.max_completion_tokens : request.max_tokens.value_or(16); - const BeamSearchParams params = request.to_beam_search_params(max_tok); + const BeamSearchParams params = + request.to_beam_search_params(max_tok, &default_sampling_params_); const std::vector prompt_ids = beam_tokenizer_->Encode(prompt); const BeamSearchOutput beams = async_engine_ != nullptr @@ -786,7 +787,8 @@ ChatCompletionResult OpenAIServingChat::create_chat_completion( engine_parser != nullptr ? nullptr : MakeReasoningParser(); const bool named_tool_choice = IsNamedToolChoice(request); - SamplingParams sampling_params = request.to_sampling_params(); + SamplingParams sampling_params = + request.to_sampling_params(std::nullopt, &default_sampling_params_); if (kMaxNewTokensCap > 0) { const int before = sampling_params.max_tokens.value_or(kMaxNewTokensCap); if (before > kMaxNewTokensCap) { diff --git a/src/vllm/entrypoints/openai/serving_completion.cpp b/src/vllm/entrypoints/openai/serving_completion.cpp index a28b1a1f3..72a835d97 100644 --- a/src/vllm/entrypoints/openai/serving_completion.cpp +++ b/src/vllm/entrypoints/openai/serving_completion.cpp @@ -196,7 +196,8 @@ CompletionResult OpenAIServingCompletion::create_completion( "beam search requires an engine and a tokenizer"); } const int max_tok = request.max_tokens.value_or(16); - const BeamSearchParams params = request.to_beam_search_params(max_tok); + const BeamSearchParams params = + request.to_beam_search_params(max_tok, &default_sampling_params_); const std::vector prompt_ids = beam_tokenizer_->Encode(request.prompt); const BeamSearchOutput beams = @@ -238,7 +239,8 @@ CompletionResult OpenAIServingCompletion::create_completion( // request → SamplingParams. to_sampling_params sets output_kind to kDelta // when stream, kFinalOnly otherwise (protocol.cpp) — matching upstream's // per-request RequestOutputKind (completion/serving.py:174). - SamplingParams sampling_params = request.to_sampling_params(); + SamplingParams sampling_params = + request.to_sampling_params(std::nullopt, &default_sampling_params_); // T0: single prompt, single choice (n == 1). The engine sub-request id is // f"{request_id}-{i}" upstream (:179); here i == 0. diff --git a/src/vllm/transformers_utils/hf_config.cpp b/src/vllm/transformers_utils/hf_config.cpp index 33adae6b2..34a895110 100644 --- a/src/vllm/transformers_utils/hf_config.cpp +++ b/src/vllm/transformers_utils/hf_config.cpp @@ -373,6 +373,40 @@ std::vector ReadGenerationConfigEosIds(const std::string& path) { return out; } +// One optional numeric key out of a generation_config.json object. Absent, +// null and wrong-typed all read as "the checkpoint did not declare it", which +// is upstream's polarity: try_get_generation_config never raises, and +// get_diff_sampling_param's `if config.get(p) is not None` drops the key. +template +std::optional OptionalNumber(const nlohmann::json& doc, const char* key) { + const auto it = doc.find(key); + if (it == doc.end() || it->is_null() || !it->is_number()) return std::nullopt; + return it->get(); +} + +// The six sampling keys of an already-parsed generation_config.json document. +GenerationConfigSampling SamplingFromDoc(const nlohmann::json& gen) { + GenerationConfigSampling out; + if (!gen.is_object()) return out; + out.repetition_penalty = OptionalNumber(gen, "repetition_penalty"); + out.temperature = OptionalNumber(gen, "temperature"); + out.top_k = OptionalNumber(gen, "top_k"); + out.top_p = OptionalNumber(gen, "top_p"); + out.min_p = OptionalNumber(gen, "min_p"); + out.max_new_tokens = OptionalNumber(gen, "max_new_tokens"); + return out; +} + +// generation_config.json's sampling keys, read from an explicit file path. +GenerationConfigSampling ReadGenerationConfigSamplingAt(const std::string& file) { + std::ifstream in(file, std::ios::binary); + if (!in) return {}; + nlohmann::json gen = nlohmann::json::parse(in, /*cb=*/nullptr, + /*allow_exceptions=*/false); + if (gen.is_discarded()) return {}; + return SamplingFromDoc(gen); +} + } // namespace namespace { @@ -564,6 +598,8 @@ HfConfig ParseHfConfigDoc(nlohmann::json doc, const std::string& path, cfg.raw = std::move(doc); if (sibling_generation_config) { cfg.generation_config_eos_ids = ReadGenerationConfigEosIds(path); + cfg.generation_config_sampling = + ReadGenerationConfigSamplingAt(SiblingGenerationConfigPath(path)); } return cfg; } @@ -589,4 +625,8 @@ HfConfig ParseHfConfig(const nlohmann::json& doc, const std::string& source) { return ParseHfConfigDoc(doc, source, /*sibling_generation_config=*/false); } +GenerationConfigSampling ReadGenerationConfigSamplingFile(const std::string& path) { + return ReadGenerationConfigSamplingAt(path); +} + } // namespace vllm diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 6a616d781..845b072a5 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -63,6 +63,7 @@ vllm_cpp_add_test(test_plugin_system vllm/plugins/test_plugin_system.cpp vllm/plugins/toy_model_plugin.cpp) vllm_cpp_add_test(test_hf_config vllm/test_hf_config.cpp) +vllm_cpp_add_test(test_generation_config vllm/config/test_generation_config.cpp) # ENG-HF-MODEL-DOWNLOAD W2 (#1280): the HuggingFace cache layout. No socket, # so it is not behind VLLM_CPP_SERVER; the fake-hub half is, below. vllm_cpp_add_test(test_hf_cache vllm/transformers_utils/test_hf_cache.cpp) diff --git a/tests/vllm/config/test_generation_config.cpp b/tests/vllm/config/test_generation_config.cpp new file mode 100644 index 000000000..8c965f41f --- /dev/null +++ b/tests/vllm/config/test_generation_config.cpp @@ -0,0 +1,160 @@ +// Ported from: vllm/config/model.py::ModelConfig.get_diff_sampling_param +// @ 5559679229bc961848b121ccdeaa8fa5d79bec98, and the --generation-config +// selector it reads (config/model.py:298-304). +#include + +#include +#include +#include + +#include "vllm/config/generation.h" + +namespace { + +// A checkpoint directory carrying config.json and, optionally, its own +// generation_config.json. +class TempModelDir { + public: + TempModelDir(const std::string& gen_body, const char* tag) { + static int counter = 0; + dir_ = (std::filesystem::temp_directory_path() / + ("vllm_gen_cfg_test_" + std::string(tag) + "_" + + std::to_string(counter++))) + .string(); + std::filesystem::create_directories(dir_); + std::ofstream(dir_ + "/config.json", std::ios::binary) << R"({ + "model_type": "llama", + "architectures": ["LlamaForCausalLM"], + "hidden_size": 8, + "num_hidden_layers": 1, + "num_attention_heads": 2, + "vocab_size": 32, + "max_position_embeddings": 128 + })"; + if (!gen_body.empty()) { + std::ofstream(dir_ + "/generation_config.json", std::ios::binary) << gen_body; + } + } + ~TempModelDir() { std::filesystem::remove_all(dir_); } + std::string config_path() const { return dir_ + "/config.json"; } + const std::string& dir() const { return dir_; } + + private: + std::string dir_; +}; + +// The Qwen3.8-27B file, read live 2026-08-26 from +// https://huggingface.co/Qwen/Qwen3.8-27B/resolve/main/generation_config.json +constexpr const char* kQwen27B = R"({ + "bos_token_id": 248044, + "do_sample": true, + "eos_token_id": [248046, 248044], + "pad_token_id": 248044, + "temperature": 1.0, + "top_k": 20, + "top_p": 0.95 +})"; + +} // namespace + +TEST_CASE("GetDiffSamplingParam: auto takes the checkpoint's own file") { + TempModelDir model(kQwen27B, "auto"); + const vllm::HfConfig cfg = vllm::LoadHfConfig(model.config_path()); + + const vllm::DefaultSamplingParams d = vllm::GetDiffSamplingParam(cfg, "auto"); + CHECK_FALSE(d.empty()); + REQUIRE(d.temperature.has_value()); + CHECK(*d.temperature == doctest::Approx(1.0)); + REQUIRE(d.top_k.has_value()); + CHECK(*d.top_k == 20); + REQUIRE(d.top_p.has_value()); + CHECK(*d.top_p == doctest::Approx(0.95)); + CHECK_FALSE(d.min_p.has_value()); + CHECK_FALSE(d.repetition_penalty.has_value()); + CHECK_FALSE(d.max_tokens.has_value()); + + // "auto" is the DEFAULT selector (config/model.py:298), so calling without + // one has to give the same answer. + const vllm::DefaultSamplingParams implicit = vllm::GetDiffSamplingParam(cfg); + CHECK(implicit.top_k == d.top_k); + CHECK(implicit.top_p == d.top_p); + CHECK(implicit.temperature == d.temperature); +} + +TEST_CASE("GetDiffSamplingParam: vllm discards the file entirely") { + // `src == "vllm"` -> `config = {}`. This is the documented escape hatch the + // upstream warning tells a user to reach for, and it has to restore the + // pre-#1985 behaviour exactly: nothing set, so every knob falls to neutral. + TempModelDir model(kQwen27B, "vllm"); + const vllm::HfConfig cfg = vllm::LoadHfConfig(model.config_path()); + const vllm::DefaultSamplingParams d = vllm::GetDiffSamplingParam(cfg, "vllm"); + CHECK(d.empty()); + CHECK_FALSE(d.top_k.has_value()); + CHECK_FALSE(d.top_p.has_value()); + CHECK_FALSE(d.temperature.has_value()); +} + +TEST_CASE("GetDiffSamplingParam: a directory selector reads THAT directory") { + TempModelDir checkpoint(kQwen27B, "ckpt"); + TempModelDir elsewhere(R"({"temperature": 0.3, "top_k": 5})", "elsewhere"); + const vllm::HfConfig cfg = vllm::LoadHfConfig(checkpoint.config_path()); + + const vllm::DefaultSamplingParams d = + vllm::GetDiffSamplingParam(cfg, elsewhere.dir()); + // The named folder wins over the checkpoint's own sibling file. + REQUIRE(d.temperature.has_value()); + CHECK(*d.temperature == doctest::Approx(0.3)); + REQUIRE(d.top_k.has_value()); + CHECK(*d.top_k == 5); + CHECK_FALSE(d.top_p.has_value()); // the checkpoint's 0.95 must NOT leak in + + // A trailing separator names the same folder. + const vllm::DefaultSamplingParams slashed = + vllm::GetDiffSamplingParam(cfg, elsewhere.dir() + "/"); + CHECK(slashed.top_k == d.top_k); + + // A folder that does not exist yields {} rather than throwing, because + // try_get_generation_config returns None and get_diff_sampling_param then + // returns {}. + vllm::DefaultSamplingParams missing; + REQUIRE_NOTHROW(missing = vllm::GetDiffSamplingParam(cfg, "/nonexistent/xyz")); + CHECK(missing.empty()); +} + +TEST_CASE("GetDiffSamplingParam: max_new_tokens is renamed to max_tokens") { + // "Huggingface definition of max_new_tokens is equivalent to vLLM's + // max_tokens" -- get_diff_sampling_param pops one and sets the other, so the + // HF spelling must NOT survive into the resolved defaults. + TempModelDir model(R"({"max_new_tokens": 512, "repetition_penalty": 1.05, + "min_p": 0.05})", "rename"); + const vllm::HfConfig cfg = vllm::LoadHfConfig(model.config_path()); + const vllm::DefaultSamplingParams d = vllm::GetDiffSamplingParam(cfg, "auto"); + REQUIRE(d.max_tokens.has_value()); + CHECK(*d.max_tokens == 512); + REQUIRE(d.repetition_penalty.has_value()); + CHECK(*d.repetition_penalty == doctest::Approx(1.05)); + REQUIRE(d.min_p.has_value()); + CHECK(*d.min_p == doctest::Approx(0.05)); +} + +TEST_CASE("GetDiffSamplingParam: a checkpoint with no file resolves to empty") { + TempModelDir model("", "nofile"); + const vllm::HfConfig cfg = vllm::LoadHfConfig(model.config_path()); + CHECK(vllm::GetDiffSamplingParam(cfg, "auto").empty()); +} + +TEST_CASE("DefaultSamplingParams::ToString reports what was resolved") { + // The startup line has to say WHICH values were taken from the checkpoint; + // a log that only says "defaults applied" cannot be audited against the + // oracle's own "Default vLLM sampling parameters have been overridden" line. + vllm::DefaultSamplingParams d; + CHECK(d.ToString().empty()); + d.temperature = 1.0; + d.top_k = 20; + d.top_p = 0.95; + const std::string s = d.ToString(); + CHECK(s.find("'temperature': 1") != std::string::npos); + CHECK(s.find("'top_k': 20") != std::string::npos); + CHECK(s.find("'top_p': 0.95") != std::string::npos); + CHECK(s.find("min_p") == std::string::npos); +} diff --git a/tests/vllm/entrypoints/openai/test_protocol.cpp b/tests/vllm/entrypoints/openai/test_protocol.cpp index ba111b611..5039b57db 100644 --- a/tests/vllm/entrypoints/openai/test_protocol.cpp +++ b/tests/vllm/entrypoints/openai/test_protocol.cpp @@ -903,3 +903,164 @@ TEST_CASE("ClampPromptLogprobs rewrites -inf to -9999.0 in place") { ClampPromptLogprobs(none); // None in, None out (serving.py:308-309) CHECK_FALSE(none.has_value()); } + +// ─── server-wide sampling defaults from generation_config.json (#1985) ────── +// Ported from vllm/entrypoints/openai/completion/protocol.py:: +// CompletionRequest.to_sampling_params @ 555967922, whose resolution is +// value = request.field +// if value is None: value = default_sampling_params.get(name, NEUTRAL) +// The request fields are `| None = None` precisely so "omitted" and "sent as +// the neutral value" are different states, and this block is what pins that. +namespace { + +// Qwen/Qwen3.8-27B's generation_config.json, read live 2026-08-26. Under the +// pinned transformers floor (>= 5.5.3) every declared key survives +// GenerationConfig.to_diff_dict(), so this IS what upstream's +// default_sampling_params holds for that checkpoint. +vllm::DefaultSamplingParams Qwen27BDefaults() { + vllm::DefaultSamplingParams d; + d.temperature = 1.0; + d.top_k = 20; + d.top_p = 0.95; + return d; +} + +} // namespace + +TEST_CASE("to_sampling_params: an OMITTED knob takes the checkpoint's value") { + const vllm::DefaultSamplingParams d = Qwen27BDefaults(); + + SUBCASE("completions") { + // The exact shape `vllm bench serve` now sends: no temperature, no top_k, + // no top_p. vLLM draws from 20 candidates here; before #1985 we drew from + // the whole vocabulary. + vllm::entrypoints::openai::CompletionRequest req; + req.model = "m"; + req.prompt = "hi"; + const vllm::SamplingParams sp = req.to_sampling_params(std::nullopt, &d); + CHECK(sp.top_k == 20); + CHECK(sp.top_p == doctest::Approx(0.95)); + CHECK(sp.temperature == doctest::Approx(1.0)); + } + + SUBCASE("chat completions") { + vllm::entrypoints::openai::ChatCompletionRequest req; + req.model = "m"; + const vllm::SamplingParams sp = req.to_sampling_params(std::nullopt, &d); + CHECK(sp.top_k == 20); + CHECK(sp.top_p == doctest::Approx(0.95)); + CHECK(sp.temperature == doctest::Approx(1.0)); + } +} + +TEST_CASE("to_sampling_params: an EXPLICIT request value wins over the checkpoint") { + const vllm::DefaultSamplingParams d = Qwen27BDefaults(); + + SUBCASE("completions") { + vllm::entrypoints::openai::CompletionRequest req; + req.model = "m"; + req.prompt = "hi"; + req.top_k = 7; + req.top_p = 0.5; + req.temperature = 0.3; + const vllm::SamplingParams sp = req.to_sampling_params(std::nullopt, &d); + CHECK(sp.top_k == 7); + CHECK(sp.top_p == doctest::Approx(0.5)); + CHECK(sp.temperature == doctest::Approx(0.3)); + } + + SUBCASE("temperature 0 still means greedy, checkpoint top-k and all") { + // sampling_params.py:492-497 (__post_init__): "Zero temperature means + // greedy sampling", and it CLEARS top_p/top_k/min_p. The checkpoint's + // top_k = 20 must not survive that, or a --temperature 0 run would stop + // being greedy the moment this row landed -- which is the SACRED path. + vllm::entrypoints::openai::CompletionRequest req; + req.model = "m"; + req.prompt = "hi"; + req.temperature = 0.0; + const vllm::SamplingParams sp = req.to_sampling_params(std::nullopt, &d); + CHECK(sp.temperature == doctest::Approx(0.0)); + CHECK(sp.top_k == 0); + CHECK(sp.top_p == doctest::Approx(1.0)); + CHECK(sp.min_p == doctest::Approx(0.0)); + } + + SUBCASE("a request that explicitly sends the NEUTRAL value keeps it") { + // top_k = 0 means "disabled" and is a real request, not an omission. If the + // checkpoint's 20 could override it, a client could not turn top-k off. + vllm::entrypoints::openai::CompletionRequest req; + req.model = "m"; + req.prompt = "hi"; + req.top_k = 0; + req.top_p = 1.0; + const vllm::SamplingParams sp = req.to_sampling_params(std::nullopt, &d); + CHECK(sp.top_k == 0); + CHECK(sp.top_p == doctest::Approx(1.0)); + } + + SUBCASE("chat completions") { + vllm::entrypoints::openai::ChatCompletionRequest req; + req.model = "m"; + req.top_k = 7; + req.temperature = 0.2; + const vllm::SamplingParams sp = req.to_sampling_params(std::nullopt, &d); + CHECK(sp.top_k == 7); + CHECK(sp.temperature == doctest::Approx(0.2)); + CHECK(sp.top_p == doctest::Approx(0.95)); // omitted -> checkpoint + } +} + +TEST_CASE("to_sampling_params: a knob the checkpoint does not declare falls to neutral") { + // Qwen3.8-27B declares no min_p and no repetition_penalty, so those must land + // on _DEFAULT_SAMPLING_PARAMS (0.0 and 1.0) rather than on anything else. + const vllm::DefaultSamplingParams d = Qwen27BDefaults(); + vllm::entrypoints::openai::CompletionRequest req; + req.model = "m"; + req.prompt = "hi"; + const vllm::SamplingParams sp = req.to_sampling_params(std::nullopt, &d); + CHECK(sp.min_p == doctest::Approx(0.0)); + CHECK(sp.repetition_penalty == doctest::Approx(1.0)); +} + +TEST_CASE("to_sampling_params: no server defaults is byte-identical to before") { + // nullptr must reproduce the pre-#1985 resolution exactly, because that is + // what --generation-config vllm resolves to and what every existing caller + // and test passes. + vllm::entrypoints::openai::CompletionRequest req; + req.model = "m"; + req.prompt = "hi"; + const vllm::SamplingParams a = req.to_sampling_params(); + const vllm::SamplingParams b = req.to_sampling_params(std::nullopt, nullptr); + const vllm::DefaultSamplingParams empty; + const vllm::SamplingParams c = req.to_sampling_params(std::nullopt, &empty); + for (const vllm::SamplingParams* sp : {&a, &b, &c}) { + CHECK(sp->top_k == 0); + CHECK(sp->top_p == doctest::Approx(1.0)); + CHECK(sp->temperature == doctest::Approx(1.0)); + CHECK(sp->min_p == doctest::Approx(0.0)); + CHECK(sp->repetition_penalty == doctest::Approx(1.0)); + } +} + +TEST_CASE("to_beam_search_params resolves temperature the same way") { + // completion/protocol.py:269-270 does the identical + // `default_sampling_params.get("temperature", 1.0)` lookup, so beam search + // must not be left on the old rule. + vllm::DefaultSamplingParams d; + d.temperature = 0.6; + + vllm::entrypoints::openai::CompletionRequest req; + req.model = "m"; + req.prompt = "hi"; + req.n = 2; + CHECK(req.to_beam_search_params(8, &d).temperature == doctest::Approx(0.6)); + req.temperature = 0.1; + CHECK(req.to_beam_search_params(8, &d).temperature == doctest::Approx(0.1)); + CHECK(req.to_beam_search_params(8).temperature == doctest::Approx(0.1)); + + vllm::entrypoints::openai::ChatCompletionRequest chat; + chat.model = "m"; + chat.n = 2; + CHECK(chat.to_beam_search_params(8, &d).temperature == doctest::Approx(0.6)); + CHECK(chat.to_beam_search_params(8).temperature == doctest::Approx(1.0)); +} diff --git a/tests/vllm/entrypoints/openai/test_serve_recipe_args.cpp b/tests/vllm/entrypoints/openai/test_serve_recipe_args.cpp index 876606820..08a875f12 100644 --- a/tests/vllm/entrypoints/openai/test_serve_recipe_args.cpp +++ b/tests/vllm/entrypoints/openai/test_serve_recipe_args.cpp @@ -229,3 +229,28 @@ TEST_CASE("each accepted serve flag announces itself and its reason") { CHECK_FALSE(Contains(tool_choice.output, "--trust-remote-code")); CHECK_FALSE(Contains(trust.output, "--enable-auto-tool-choice")); } + +// ─── --generation-config reaches the real entry point (#1985) ──────────────── +// Mirrors vllm/config/model.py:298-304, whose three forms are "auto" (the +// default), "vllm" and a directory path. This case gates the CLI hop of the +// chain and nothing else: it re-execs the REAL `VllmServerMain`, so deleting +// the `--generation-config` branch in server_main.cpp's argument loop turns it +// red on the "unknown argument" path. +// +// It cannot reach the wiring itself, which sits after model load and so needs a +// generative on-disk checkpoint this repository does not commit. The hops below +// it are gated in tests/vllm/entrypoints/openai/test_serving.cpp (the handler +// applies the defaults) and tests/vllm/config/test_generation_config.cpp (the +// selector resolves each form). +TEST_CASE("--generation-config is accepted in all three of its forms") { + for (const char* value : {"auto", "vllm", "/tmp"}) { + CAPTURE(value); + const ChildRun run = RunServer(std::string(kMissingModel) + + " --generation-config " + value); + INFO("child output:\n" << run.output); + CHECK_FALSE(Contains(run.output, kUnknownArgument)); + CHECK(Contains(run.output, kPostParseBanner)); + CHECK(Contains(run.output, "server: loading model from")); + CHECK(run.status == 0); + } +} diff --git a/tests/vllm/entrypoints/openai/test_serving.cpp b/tests/vllm/entrypoints/openai/test_serving.cpp index d8bd6aecb..a0c49735c 100644 --- a/tests/vllm/entrypoints/openai/test_serving.cpp +++ b/tests/vllm/entrypoints/openai/test_serving.cpp @@ -28,6 +28,7 @@ #include +#include "vllm/config/generation.h" #include "vllm/config/scheduler.h" #include "vllm/model_executor/models/qwen3_5.h" #include "vllm/model_executor/models/qwen3_5_weights.h" @@ -1972,3 +1973,168 @@ TEST_CASE("serving_chat: streaming with tools but content-only is unchanged") { REQUIRE(last_finish.has_value()); CHECK(*last_finish == "length"); // engine finish, NOT "tool_calls" } + +// ─── REACHABILITY: the checkpoint's sampling defaults reach the sampler (#1985) ─ +// +// AGENTS.md `## Nothing lands dead`, and .agents/reachability.md's method: a +// unit test that hands `to_sampling_params` a hand-built DefaultSamplingParams +// proves the resolution works and says nothing about whether a served request +// ever meets one. So this block never constructs the defaults itself. It reads a +// real generation_config.json off disk with the production `LoadHfConfig`, +// narrows it with the production `GetDiffSamplingParam`, and then enters through +// `OpenAIServingCompletion::create_completion` / `create_chat_completion` — the +// handler entry the HTTP layer calls. +// +// THE REACHABILITY MUTATION: turn either +// request.to_sampling_params(std::nullopt, &default_sampling_params_) +// back into `request.to_sampling_params()` in serving_completion.cpp or +// serving_chat.cpp, and these cases go RED. +// +// The observable is deliberately generative rather than a getter. A checkpoint +// that ships `"temperature": 0.0` makes an omitted-temperature request GREEDY, +// and greedy output is exactly computable by asking the same engine with an +// explicit temperature 0. If the defaults do not reach the sampler the request +// resolves to temperature 1.0 and samples instead, so the two texts part. +namespace { + +// A checkpoint directory holding only a generation_config.json — everything the +// production read needs, since LoadHfConfig takes the config.json path and reads +// its SIBLING. +class TempGenerationConfigDir { + public: + explicit TempGenerationConfigDir(const std::string& gen_body) { + static int counter = 0; + dir_ = (std::filesystem::temp_directory_path() / + ("vllm_serving_gencfg_" + std::to_string(counter++))) + .string(); + std::filesystem::create_directories(dir_); + std::ofstream(dir_ + "/config.json", std::ios::binary) << R"({ + "model_type": "llama", + "architectures": ["LlamaForCausalLM"], + "hidden_size": 8, + "num_hidden_layers": 1, + "num_attention_heads": 2, + "vocab_size": 32, + "max_position_embeddings": 128 + })"; + std::ofstream(dir_ + "/generation_config.json", std::ios::binary) << gen_body; + } + ~TempGenerationConfigDir() { std::filesystem::remove_all(dir_); } + std::string config_path() const { return dir_ + "/config.json"; } + + private: + std::string dir_; +}; + +// The production chain, with nothing hand-built: file -> HfConfig -> narrowing. +vllm::DefaultSamplingParams DefaultsFromDisk(const std::string& gen_body) { + TempGenerationConfigDir ckpt(gen_body); + return vllm::GetDiffSamplingParam(vllm::LoadHfConfig(ckpt.config_path()), + vllm::kGenerationConfigAuto); +} + +} // namespace + +TEST_CASE("serving_completion: the checkpoint's sampling defaults reach the sampler") { + const HfConfig c = MakeConfig(); + const Qwen3_5MoeWeights w = MakeWeights(c); + const Tokenizer& tok = Fixture(); + const int kN = 6; + + // The checkpoint asks for greedy decoding and a top-k, exactly as a real + // generation_config.json does; only `temperature` is observable in the text. + const vllm::DefaultSamplingParams d = + DefaultsFromDisk(R"({"temperature": 0.0, "top_k": 20, "top_p": 0.95})"); + REQUIRE(d.temperature.has_value()); + REQUIRE(*d.temperature == doctest::Approx(0.0)); + + Harness h(c, w, tok); + + // The reference: the same engine asked EXPLICITLY for greedy. + OpenAIServingCompletion reference(h.engine, "test-model"); + CompletionRequest explicit_greedy; + explicit_greedy.prompt = "hello"; + explicit_greedy.max_tokens = kN; + explicit_greedy.temperature = 0.0; + CompletionResult ref = reference.create_completion(explicit_greedy); + REQUIRE(ref.response.has_value()); + const std::string greedy_text = ref.response->choices.at(0).text; + REQUIRE_FALSE(greedy_text.empty()); + + // The request `vllm bench serve` now sends: no temperature, no top_k, no + // top_p. It must resolve to the checkpoint's values and therefore reproduce + // the greedy text. + OpenAIServingCompletion serving(h.engine, "test-model"); + serving.set_default_sampling_params(d); + CompletionRequest omitted; + omitted.prompt = "hello"; + omitted.max_tokens = kN; + CompletionResult res = serving.create_completion(omitted); + REQUIRE(res.response.has_value()); + CHECK(res.response->choices.at(0).text == greedy_text); + + // And an EXPLICIT request value still wins over the checkpoint: asking for + // top_k = 1 keeps the draw deterministic while leaving temperature at the + // checkpoint's 0.0, so the text is unchanged; asking for a temperature the + // checkpoint did not name is what a client is entitled to do. + OpenAIServingCompletion overridden(h.engine, "test-model"); + overridden.set_default_sampling_params(d); + CompletionRequest explicit_topk; + explicit_topk.prompt = "hello"; + explicit_topk.max_tokens = kN; + explicit_topk.top_k = 1; + CompletionResult over = overridden.create_completion(explicit_topk); + REQUIRE(over.response.has_value()); + CHECK(over.response->choices.at(0).text == greedy_text); +} + +TEST_CASE("serving_chat: the checkpoint's sampling defaults reach the sampler") { + const HfConfig c = MakeConfig(); + const Qwen3_5MoeWeights w = MakeWeights(c); + const Tokenizer& tok = Fixture(); + const int kN = 6; + + const vllm::DefaultSamplingParams d = + DefaultsFromDisk(R"({"temperature": 0.0, "top_k": 20})"); + Harness h(c, w, tok); + + const auto ask = [&](bool wire_defaults, bool explicit_greedy) { + OpenAIServingChat chat(h.engine, "test-model", InVocabChatPrompt); + if (wire_defaults) chat.set_default_sampling_params(d); + ChatCompletionRequest req; + req.messages.push_back(ChatMessage{"user", "hello", {}, {}, {}}); + req.max_tokens = kN; + if (explicit_greedy) req.temperature = 0.0; + ChatCompletionResult r = chat.create_chat_completion(req); + REQUIRE(r.response.has_value()); + REQUIRE_FALSE(r.response->choices.empty()); + const auto& content = r.response->choices.at(0).message.content; + REQUIRE(content.has_value()); + return *content; + }; + + const std::string greedy_text = ask(/*wire_defaults=*/false, /*explicit_greedy=*/true); + REQUIRE_FALSE(greedy_text.empty()); + CHECK(ask(/*wire_defaults=*/true, /*explicit_greedy=*/false) == greedy_text); +} + +TEST_CASE("serving: no server defaults leaves the resolution exactly as it was") { + // --generation-config vllm, and every handler constructed before #1985: the + // knob-free request must resolve the way it always did. This is the inertness + // half of the gate above -- without it, a wiring that ALWAYS applied some + // default would pass the positive case. + const HfConfig c = MakeConfig(); + const Qwen3_5MoeWeights w = MakeWeights(c); + const Tokenizer& tok = Fixture(); + + Harness h(c, w, tok); + OpenAIServingCompletion serving(h.engine, "test-model"); + CHECK(serving.default_sampling_params().empty()); + + // GetDiffSamplingParam("vllm") is what the flag resolves to, and it must be + // empty even when the checkpoint ships a full file. + TempGenerationConfigDir ckpt(R"({"temperature": 0.0, "top_k": 20, "top_p": 0.95})"); + const HfConfig loaded = vllm::LoadHfConfig(ckpt.config_path()); + CHECK(vllm::GetDiffSamplingParam(loaded, vllm::kGenerationConfigNone).empty()); + CHECK_FALSE(vllm::GetDiffSamplingParam(loaded, vllm::kGenerationConfigAuto).empty()); +} diff --git a/tests/vllm/test_hf_config.cpp b/tests/vllm/test_hf_config.cpp index 66969bd9f..1ce8a93d9 100644 --- a/tests/vllm/test_hf_config.cpp +++ b/tests/vllm/test_hf_config.cpp @@ -1066,3 +1066,100 @@ TEST_CASE("LoadHfConfig normalizes the GDN output_gate_type") { CHECK_THROWS_AS(vllm::LoadHfConfig(f.path()), std::runtime_error); } } + +// ─── generation_config.json SAMPLING keys (#1985) ─────────────────────────── +// Upstream reads the same file for the same reason and then narrows it in +// ModelConfig.get_diff_sampling_param (vllm/config/model.py). We read +// eos_token_id out of it already; these cases cover the six sampling keys the +// narrowing consumes, and the polarity that makes the defaulting rule work: +// UNSET means the checkpoint said nothing, which is not the same as the +// checkpoint saying the neutral value. +TEST_CASE("LoadHfConfig reads the sibling generation_config.json sampling keys") { + constexpr const char* kConfig = R"({ + "model_type": "llama", + "architectures": ["LlamaForCausalLM"], + "hidden_size": 8, + "num_hidden_layers": 1, + "num_attention_heads": 2, + "vocab_size": 32, + "max_position_embeddings": 128 + })"; + + SUBCASE("the Qwen3.8-27B file, verbatim") { + // Read live 2026-08-26 from + // https://huggingface.co/Qwen/Qwen3.8-27B/resolve/main/generation_config.json + TempModelDir model(kConfig, R"({ + "bos_token_id": 248044, + "do_sample": true, + "eos_token_id": [248046, 248044], + "pad_token_id": 248044, + "temperature": 1.0, + "top_k": 20, + "top_p": 0.95 + })"); + const vllm::HfConfig cfg = vllm::LoadHfConfig(model.config_path()); + const auto& g = cfg.generation_config_sampling; + REQUIRE(g.temperature.has_value()); + CHECK(*g.temperature == doctest::Approx(1.0)); + REQUIRE(g.top_k.has_value()); + CHECK(*g.top_k == 20); + REQUIRE(g.top_p.has_value()); + CHECK(*g.top_p == doctest::Approx(0.95)); + // Declared by no Qwen3.8-27B file, so they stay UNSET rather than becoming + // the neutral value. This is the distinction the defaulting rule turns on. + CHECK_FALSE(g.min_p.has_value()); + CHECK_FALSE(g.repetition_penalty.has_value()); + CHECK_FALSE(g.max_new_tokens.has_value()); + // The eos read is unaffected by the new keys. + std::vector eos = {248044, 248046}; + CHECK(cfg.generation_config_eos_ids == eos); + } + + SUBCASE("all six keys") { + TempModelDir model(kConfig, R"({ + "repetition_penalty": 1.05, + "temperature": 0.7, + "top_k": 40, + "top_p": 0.8, + "min_p": 0.05, + "max_new_tokens": 512 + })"); + const auto g = vllm::LoadHfConfig(model.config_path()).generation_config_sampling; + CHECK(*g.repetition_penalty == doctest::Approx(1.05)); + CHECK(*g.temperature == doctest::Approx(0.7)); + CHECK(*g.top_k == 40); + CHECK(*g.top_p == doctest::Approx(0.8)); + CHECK(*g.min_p == doctest::Approx(0.05)); + CHECK(*g.max_new_tokens == 512); + } + + SUBCASE("absent, malformed, null and wrong-typed all leave every key UNSET") { + // try_get_generation_config returns {} on every failure path and never + // raises, so none of these may throw and none may invent a value. + for (const char* body : {"", "{ this is not json", "[1,2,3]", + R"({"temperature": null, "top_k": null})", + R"({"temperature": "hot", "top_k": [20]})"}) { + CAPTURE(body); + TempModelDir model(kConfig, body); + vllm::HfConfig cfg; + REQUIRE_NOTHROW(cfg = vllm::LoadHfConfig(model.config_path())); + const auto& g = cfg.generation_config_sampling; + CHECK_FALSE(g.temperature.has_value()); + CHECK_FALSE(g.top_k.has_value()); + CHECK_FALSE(g.top_p.has_value()); + CHECK_FALSE(g.min_p.has_value()); + CHECK_FALSE(g.repetition_penalty.has_value()); + CHECK_FALSE(g.max_new_tokens.has_value()); + } + } + + SUBCASE("ParseHfConfig has no sibling to read, so it stays UNSET") { + // The in-memory entry point must not pick up whatever generation_config.json + // happens to sit in the working directory -- the same reason + // generation_config_eos_ids stays empty there. + const auto doc = nlohmann::json::parse(kConfig); + const auto g = vllm::ParseHfConfig(doc, "").generation_config_sampling; + CHECK_FALSE(g.temperature.has_value()); + CHECK_FALSE(g.top_k.has_value()); + } +} From 1a4b139fd37a783232bb4dccee8e5c33b62c51d4 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 16:08:16 +0000 Subject: [PATCH 03/11] perf(SAMPLE-CORE): the Gumbel draw reuses the parallel argmax instead of walking 248,320 tokens on one thread (#1984) `RandomSampleKernel` was launched `<<>>`, returned from every thread but lane 0, and walked the whole vocabulary serially, computing two SplitMix64 rounds and an f64 `log` per element. Eleven lines above it the same file records why that shape was already unacceptable: a single-block single-thread scan of a ~151k vocab cost ~7.5 ms/token, which is why greedy argmax became the two-pass grid-strided `ArgmaxPartialKernel`/`ArgmaxFinalKernel` pair. The Gumbel draw never got the same treatment, and upstream's is whole-tensor (`probs.div_(q).argmax(-1)`, topk_topp_sampler.py::sample_with_exponential_noise), so this is a mirror obligation rather than a tuning choice. The answer is the same lowest-index argmax the greedy path already computes, so the partial kernel is now parameterised on the score and `random_sample` instantiates it. That keeps ONE order-independent reduction in the file instead of a second hand-written copy, and greedy's instantiation compiles the expression it always did. The output is bit-identical rather than merely equivalent. Every element's score is the same `GumbelScore(prob, seed, row, j)` evaluated by the same device libm on both paths; only the order in which those identical floats are combined changes, and `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; 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. ROCm's copy is deliberately untouched and recorded under `## Owed`: `rocm_sample.hip::RandomSampleK` carries the same `<<>>` defect, on hardware this change has no way to gate. Mutation-proven, each red then restored byte for byte: dropping the `bi < ai` tie-break from `ArgReduce` fails 35 assertions of test_ops_sample; turning the score's divide into a multiply fails the pre-existing distribution case, which is the right attribution, because the equivalence cases compare two orderings of one shared score expression and cannot see a change to the expression itself. NOT gated on this host and stated rather than implied: there is no CUDA toolkit here, so the `.cu` change is not compiled locally and its first compile is CI's `cuda-fat-build`. The host cases gate the property the kernel rests on -- the production `ArgReduce`, `GumbelScore` and `ArgBlocksPerRow` under the kernel's exact two-pass partition, against `vt::RandomSample`, over ten vocabulary widths including 248,320, six row shapes chosen for ties and masked zeros, and four seeds. The kernel itself is gated by the `VT_FAST_RANDOM_SAMPLE` A/B, which skips without a device, and by the end-to-end measurement the row's spec commits to in advance. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .../vllm/entrypoints/openai/serving_chat.h | 16 +- include/vt/sample_common.h | 107 +++++ src/vt/cpu/cpu_sample.cpp | 24 +- src/vt/cuda/cuda_sample.cu | 168 +++++--- tests/vt/test_ops_sample.cpp | 378 ++++++++++++++++++ 5 files changed, 616 insertions(+), 77 deletions(-) create mode 100644 include/vt/sample_common.h diff --git a/include/vllm/entrypoints/openai/serving_chat.h b/include/vllm/entrypoints/openai/serving_chat.h index bae447a3b..964899757 100644 --- a/include/vllm/entrypoints/openai/serving_chat.h +++ b/include/vllm/entrypoints/openai/serving_chat.h @@ -231,14 +231,6 @@ class OpenAIServingChat { beam_eos_token_id_ = eos_token_id; } - // See OpenAIServingCompletion::set_default_sampling_params (#1985). - void set_default_sampling_params(vllm::DefaultSamplingParams defaults) { - default_sampling_params_ = std::move(defaults); - } - const vllm::DefaultSamplingParams& default_sampling_params() const { - return default_sampling_params_; - } - // Attach the multimodal chat seam (see MultiModalChatFn). Unset (default) // keeps the text-only path byte-identical. When set AND a request carries a mm // content part, create_chat_completion routes the request through the engine @@ -253,6 +245,14 @@ class OpenAIServingChat { // the IDENTICAL chat template as chat-completions instead of reinventing it. const ChatPromptFn& prompt_fn() const { return prompt_fn_; } + // See OpenAIServingCompletion::set_default_sampling_params (#1985). + void set_default_sampling_params(vllm::DefaultSamplingParams defaults) { + default_sampling_params_ = std::move(defaults); + } + const vllm::DefaultSamplingParams& default_sampling_params() const { + return default_sampling_params_; + } + private: // Build the per-request tool parser (get_tool_parser) when ToolsEnabled and a // parser name is configured; else nullptr. ONE instance per request (the diff --git a/include/vt/sample_common.h b/include/vt/sample_common.h new file mode 100644 index 000000000..d55b4dfe6 --- /dev/null +++ b/include/vt/sample_common.h @@ -0,0 +1,107 @@ +// vllm.cpp original — the primitives the sampling ops must agree on BYTE FOR +// BYTE across backends, in one place. +// +// WHY THIS FILE EXISTS. `SplitMix64` and `ExpNoise` were written out three +// times: src/vt/cpu/cpu_sample.cpp, src/vt/cuda/cuda_sample.cu and +// src/vt/rocm/rocm_sample.hip. The CPU copy is the reference the device copies +// are gated against, so a divergence between any two of them is a token +// difference that no test in the tree looks for -- the copies are compared only +// through their results, and only on rows where the result happens to differ. +// The CPU and CUDA copies now come from here. The ROCm copy does not yet, and +// that is recorded under `## Owed` in +// .agents/specs/sample-gen-config-and-parallel-gumbel.md rather than changed on +// hardware nobody could run the gate on. +// +// Everything here is a `__host__ __device__` inline so the SAME expression is +// compiled for the device kernel and for the host test that checks it. A host +// test that re-types the expression proves the transcription, not the function. +#ifndef VT_SAMPLE_COMMON_H_ +#define VT_SAMPLE_COMMON_H_ + +#include +#include + +#if defined(__CUDACC__) || defined(__HIPCC__) +#define VT_SAMPLE_HD __host__ __device__ +#else +#define VT_SAMPLE_HD +#endif + +namespace vt::sample { + +// Deterministic integer mixing. Bit-identical on host and device: it is 64-bit +// integer arithmetic only, with no libm and no floating point, so there is no +// rounding freedom for a platform to spend. +VT_SAMPLE_HD inline uint64_t SplitMix64(uint64_t x) { + x += 0x9E3779B97F4A7C15ULL; + x = (x ^ (x >> 30)) * 0xBF58476D1CE4E5B9ULL; + x = (x ^ (x >> 27)) * 0x94D049BB133111EBULL; + return x ^ (x >> 31); +} + +// One Exp(1) draw for cell (row, col) under `seed`, mirroring the tensor vLLM +// fills with `q.exponential_()` (vllm/v1/sample/ops/topk_topp_sampler.py). +// Upstream draws a whole tensor from a torch Generator; we hash the coordinate +// instead, so a row's draw does not depend on how many rows share the batch. +// Exact torch-Philox parity is the documented M1.7 T1 carry. +// +// The `-log(u)` is evaluated in DOUBLE. That is WIDER than upstream's default +// (`use_fp64_gumbel: bool = False`, sampler.py), and it is the reason host and +// device can disagree by an ULP: IEEE-754 does not require a correctly-rounded +// `log`, so glibc's and libdevice's may differ in the last bit. Narrowing it to +// f32 would mirror upstream and cost less on a part with 1:64 f64 throughput, +// and it would change which token is drawn -- so it is a separate row with its +// own gate, recorded under `## Owed` in +// .agents/specs/sample-gen-config-and-parallel-gumbel.md. +VT_SAMPLE_HD inline double ExpNoise(uint64_t seed, int64_t row, int64_t col) { + const uint64_t row_key = SplitMix64(seed + 0x9E3779B97F4A7C15ULL * static_cast(row)); + const uint64_t r = SplitMix64(row_key + static_cast(col)); + const double u = static_cast((r >> 11) + 1ULL) * (1.0 / 9007199254740993.0); + return -log(u); +} + +// One element of `probs.div_(q)` (topk_topp_sampler.py:: +// sample_with_exponential_noise). The argmax over a row of these IS the sample. +VT_SAMPLE_HD inline float GumbelScore(float prob, uint64_t seed, int64_t row, int64_t col) { + return prob / static_cast(ExpNoise(seed, row, col)); +} + +// The "no real index yet" marker, INT64_MAX, so that any real index beats an +// unfilled lane on the tie-break below. +constexpr int64_t kArgSentinel = 0x7fffffffffffffffLL; + +// (value, index) argmax with the LOWEST index winning a tie -- torch.argmax's +// rule, and the rule the CPU reference's `score > best_v` serial scan produces. +// +// THIS OPERATOR IS ORDER-INDEPENDENT, and that property is the whole reason a +// row can be reduced in parallel at all. It compares the true GLOBAL index +// rather than thread or block order, so any partition of a row and any order of +// combination yield the same answer as the serial left-to-right scan. Drop the +// `bi < ai` clause and the reduction becomes order-DEPENDENT: reducing +// right-to-left then returns the HIGHEST tied index, which is the defect a +// careless parallelisation introduces and which +// tests/vt/test_ops_sample.cpp's order-independence case exists to catch. +// +// NaN propagates the way the serial scan does: `bv > av` and `bv == av` are both +// false for a NaN, so a NaN never displaces a real candidate. +VT_SAMPLE_HD inline void ArgReduce(float& av, int64_t& ai, float bv, int64_t bi) { + if (bv > av || (bv == av && bi < ai)) { + av = bv; + ai = bi; + } +} + +// How many blocks cover one row of `v` elements at `block` threads each, capped +// so the second pass can reduce every partial of a row with a single block of +// `block` threads. Shared by the launcher and by the host test that mirrors the +// launch, so the test cannot check a partition the kernel does not use. +inline int ArgBlocksPerRow(int64_t v, int block) { + int64_t bpr = (v + block - 1) / block; + if (bpr > block) bpr = block; + if (bpr < 1) bpr = 1; + return static_cast(bpr); +} + +} // namespace vt::sample + +#endif // VT_SAMPLE_COMMON_H_ diff --git a/src/vt/cpu/cpu_sample.cpp b/src/vt/cpu/cpu_sample.cpp index 29a4d4793..172308976 100644 --- a/src/vt/cpu/cpu_sample.cpp +++ b/src/vt/cpu/cpu_sample.cpp @@ -14,6 +14,7 @@ #include #include "vt/ops.h" +#include "vt/sample_common.h" namespace vt::cpu { namespace { @@ -192,20 +193,10 @@ void ComputeLogprobsKernel(Queue&, Tensor& logprobs, const Tensor& logits) { // (0,1), so q ~ Exponential(1) exactly, and the exponential race // argmax(p_j / q_j) selects j with probability p_j (== softmax). The row index is // mixed in so batch-default rows (shared seed) still get independent noise. -inline uint64_t SplitMix64(uint64_t x) { - x += 0x9E3779B97F4A7C15ULL; - x = (x ^ (x >> 30)) * 0xBF58476D1CE4E5B9ULL; - x = (x ^ (x >> 27)) * 0x94D049BB133111EBULL; - return x ^ (x >> 31); -} - -inline double ExpNoise(uint64_t seed, int64_t row, int64_t col) { - const uint64_t row_key = SplitMix64(seed + 0x9E3779B97F4A7C15ULL * static_cast(row)); - const uint64_t r = SplitMix64(row_key + static_cast(col)); - // 53-bit uniform in (0,1): ((r>>11)+1) / (2^53 + 1). Strictly positive, < 1. - const double u = static_cast((r >> 11) + 1ULL) * (1.0 / 9007199254740993.0); - return -std::log(u); // Exponential(1) inverse-CDF -} +// SplitMix64 / ExpNoise are NOT written out here any more. They moved to +// include/vt/sample_common.h so the device kernels compile the SAME expression +// this reference does, rather than a copy of it that only a differing token can +// reveal has drifted. // random_sample (topk_topp_sampler.py::random_sample + sample_with_exponential_noise). // scores = probs / q with q ~ Exp(1); argmax(scores, dim=-1) (lowest-index @@ -222,8 +213,9 @@ void RandomSampleKernel(Queue&, Tensor& token_ids, const Tensor& probs, const Te int64_t best = 0; float best_v = -std::numeric_limits::infinity(); for (int64_t j = 0; j < v; ++j) { - const float q = static_cast(ExpNoise(seed, i, j)); - const float score = row[j] / q; // probs / q (higher => more likely) + // probs / q (higher => more likely). vt::sample::GumbelScore is the same + // expression the device kernels evaluate. + const float score = vt::sample::GumbelScore(row[j], seed, i, j); if (score > best_v) { // strict `>` => lowest-index tie-break best_v = score; best = j; diff --git a/src/vt/cuda/cuda_sample.cu b/src/vt/cuda/cuda_sample.cu index f2485eec2..f70afa513 100644 --- a/src/vt/cuda/cuda_sample.cu +++ b/src/vt/cuda/cuda_sample.cu @@ -24,6 +24,7 @@ #include "vt/backend.h" #include "vt/ops.h" +#include "vt/sample_common.h" namespace vt::cuda { namespace { @@ -44,20 +45,13 @@ unsigned GridFor(int64_t n) { return static_cast(blocks < 4096 ? blocks : 4096); } -// Deterministic RNG shared with cpu_sample.cpp (bit-identical integer mixing). -__device__ inline uint64_t SplitMix64(uint64_t x) { - x += 0x9E3779B97F4A7C15ULL; - x = (x ^ (x >> 30)) * 0xBF58476D1CE4E5B9ULL; - x = (x ^ (x >> 27)) * 0x94D049BB133111EBULL; - return x ^ (x >> 31); -} - -__device__ inline double ExpNoise(uint64_t seed, int64_t row, int64_t col) { - const uint64_t row_key = SplitMix64(seed + 0x9E3779B97F4A7C15ULL * static_cast(row)); - const uint64_t r = SplitMix64(row_key + static_cast(col)); - const double u = static_cast((r >> 11) + 1ULL) * (1.0 / 9007199254740993.0); - return -log(u); -} +// The RNG and the argmax reduce come from vt/sample_common.h, which cpu_sample.cpp +// also includes -- so "bit-identical to the CPU reference" is a property of the +// build rather than of two copies staying in step. +using vt::sample::ArgReduce; +using vt::sample::ExpNoise; +using vt::sample::GumbelScore; +using vt::sample::kArgSentinel; // --- apply_temperature ------------------------------------------------------ __global__ void ApplyTemperatureKernel(float* logits, const float* temp, int64_t n, int64_t v, @@ -91,20 +85,37 @@ void ApplyTemperatureCuda(Queue& q, Tensor& logits, const Tensor& temp, bool all // break is order-independent. Unfilled lanes carry (-inf, INT64_MAX) so a real // index -- even one whose logit is -inf (all-masked row) -- always beats an empty // lane, yielding index 0 for an all-(-inf) row, exactly like the CPU reference. -__device__ inline void ArgReduce(float& av, int64_t& ai, float bv, int64_t bi) { - if (bv > av || (bv == av && bi < ai)) { - av = bv; - ai = bi; +// +// PARAMETERISED ON THE SCORE (#1984). The reduction below is the same whether +// the value at (row, j) is a logit or a Gumbel score, and it is the ONLY +// order-independent argmax in this file, so `random_sample` instantiates it +// rather than growing a second hand-written copy that could drift from it. The +// greedy instantiation compiles the same expression it always did. +struct LogitScore { + const float* logits; + int64_t v; + __device__ float operator()(int64_t row, int64_t j) const { return logits[row * v + j]; } +}; + +// probs / q with q ~ Exp(1), i.e. one element of upstream's `probs.div_(q)` +// (topk_topp_sampler.py::sample_with_exponential_noise). Computed per element +// rather than materialised, exactly as the serial kernel below does, so the +// float the reduction sees is the SAME float -- this is what makes the parallel +// path bit-identical to the serial one rather than merely close. +struct GumbelScore2D { + const float* probs; + const int64_t* seeds; + int64_t v; + __device__ float operator()(int64_t row, int64_t j) const { + return GumbelScore(probs[row * v + j], static_cast(seeds[row]), row, j); } -} - -constexpr int64_t kArgSentinel = 0x7fffffffffffffffLL; // INT64_MAX +}; -__global__ void ArgmaxPartialKernel(float* part_val, int64_t* part_idx, const float* logits, - int64_t v, int blocks_per_row) { +template +__global__ void ArgmaxPartialKernel(float* part_val, int64_t* part_idx, Score score, int64_t v, + int blocks_per_row) { const int64_t row = blockIdx.y; const int blk = blockIdx.x; - const float* r = logits + row * v; __shared__ float sv[kBlock]; __shared__ int64_t si[kBlock]; @@ -112,7 +123,7 @@ __global__ void ArgmaxPartialKernel(float* part_val, int64_t* part_idx, const fl int64_t bi = kArgSentinel; const int64_t stride = static_cast(blocks_per_row) * blockDim.x; for (int64_t j = static_cast(blk) * blockDim.x + threadIdx.x; j < v; j += stride) - ArgReduce(bv, bi, r[j], j); + ArgReduce(bv, bi, score(row, j), j); sv[threadIdx.x] = bv; si[threadIdx.x] = bi; @@ -151,21 +162,32 @@ __global__ void ArgmaxFinalKernel(int64_t* out, const float* part_val, const int if (threadIdx.x == 0) out[row] = (si[0] == kArgSentinel) ? 0 : si[0]; } -// Persistent scratch for the argmax partials -- grown on demand and kept alive -// (a few KB), so the decode path never pays a cudaMalloc/cudaFree per token. -float* g_argmax_val = nullptr; -int64_t* g_argmax_idx = nullptr; -size_t g_argmax_cap = 0; // capacity in elements - -void EnsureArgmaxScratch(size_t elems) { - if (elems <= g_argmax_cap) return; - if (g_argmax_val) cudaFree(g_argmax_val); - if (g_argmax_idx) cudaFree(g_argmax_idx); - Check(cudaMalloc(&g_argmax_val, elems * sizeof(float)), "argmax scratch val"); - Check(cudaMalloc(&g_argmax_idx, elems * sizeof(int64_t)), "argmax scratch idx"); - g_argmax_cap = elems; +// Persistent scratch for the reduction partials -- grown on demand and kept +// alive (a few KB), so the decode path never pays a cudaMalloc/cudaFree per +// token. `n * blocks_per_row` elements, and blocks_per_row is capped at kBlock, +// so at num_reqs 32 this is 8192 entries: 32 KiB + 64 KiB. +struct ArgScratch { + float* val = nullptr; + int64_t* idx = nullptr; + size_t cap = 0; // capacity in elements +}; + +void EnsureArgScratch(ArgScratch& s, size_t elems, const char* what) { + if (elems <= s.cap) return; + if (s.val) cudaFree(s.val); + if (s.idx) cudaFree(s.idx); + Check(cudaMalloc(&s.val, elems * sizeof(float)), what); + Check(cudaMalloc(&s.idx, elems * sizeof(int64_t)), what); + s.cap = elems; } +// greedy_argmax and random_sample get SEPARATE scratches. A mixed greedy/random +// batch runs both inside one Sampler::sample() call, and sharing one buffer +// would make correctness depend on stream ordering that nothing in the type +// system enforces. Two allocations of tens of KiB is not a reason to accept that. +ArgScratch g_argmax_scratch; +ArgScratch g_sample_scratch; + // Legacy single-block single-thread argmax (bit-exact reference). Retained behind // VT_FAST_ARGMAX=0 for same-binary A/B against the two-pass kernel above. __global__ void GreedyArgmaxKernelSlow(int64_t* out, const float* logits, int64_t v) { @@ -204,16 +226,14 @@ void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { } // One block per kBlock vocab elements, capped so pass 2 fits a single block. - int bpr = static_cast((v + kBlock - 1) / kBlock); - if (bpr > kBlock) bpr = kBlock; // pass 2 reduces bpr partials with kBlock threads - if (bpr < 1) bpr = 1; + const int bpr = vt::sample::ArgBlocksPerRow(v, kBlock); - EnsureArgmaxScratch(static_cast(n) * bpr); + EnsureArgScratch(g_argmax_scratch, static_cast(n) * bpr, "argmax scratch"); dim3 grid1(static_cast(bpr), static_cast(n)); - ArgmaxPartialKernel<<>>(g_argmax_val, g_argmax_idx, logits.Ptr(), v, - bpr); - ArgmaxFinalKernel<<(n), kBlock, 0, s>>>(token_ids.Ptr(), - g_argmax_val, g_argmax_idx, bpr); + ArgmaxPartialKernel<<>>(g_argmax_scratch.val, g_argmax_scratch.idx, + LogitScore{logits.Ptr(), v}, v, bpr); + ArgmaxFinalKernel<<(n), kBlock, 0, s>>>( + token_ids.Ptr(), g_argmax_scratch.val, g_argmax_scratch.idx, bpr); Check(cudaGetLastError(), "greedy_argmax launch"); } @@ -267,9 +287,29 @@ void ComputeLogprobsCuda(Queue& q, Tensor& logprobs, const Tensor& logits) { Check(cudaGetLastError(), "compute_logprobs launch"); } -// --- random_sample (single-threaded per row: exact tie-break + same RNG) ----- -__global__ void RandomSampleKernel(int64_t* out, const float* probs, const int64_t* seeds, - int64_t v) { +// --- random_sample (two-pass reduction over probs/q; #1984) ------------------ +// Upstream is `probs.div_(q).argmax(dim=-1)` over the whole tensor +// (topk_topp_sampler.py::sample_with_exponential_noise), i.e. fully parallel. +// Ours was `<<>>` with `if (threadIdx.x != 0) return;` and a serial walk +// of the vocabulary -- the SAME single-lane shape recorded twenty lines above at +// ~7.5 ms/token for a ~151k greedy scan, at 1.64x the vocab and with an f64 +// `log` and two 64-bit mixes on top. It now reuses the greedy reduction with the +// Gumbel score substituted for the logit. +// +// The output is BIT-IDENTICAL, not merely equivalent, and the reason is worth +// stating because the gate asserts equality rather than agreement: every +// element's score is `GumbelScore(probs[row][j], seed, row, j)` on both paths, +// evaluated by the same device libm, so the reduction sees the same floats and +// differs only in the order it combines them -- and ArgReduce is +// order-independent (vt/sample_common.h). +// +// The serial kernel below is RETAINED, reachable as VT_FAST_RANDOM_SAMPLE=0, +// mirroring the VT_FAST_ARGMAX lever the greedy rewrite kept. It is what makes +// the equality gate a same-binary A/B, which AGENTS.md requires before a +// performance result is accepted -- and what #1929/#1975 cost when a sampling +// kernel landed on a per-kernel figure alone. +__global__ void RandomSampleKernelSlow(int64_t* out, const float* probs, const int64_t* seeds, + int64_t v) { const int64_t row = blockIdx.x; if (threadIdx.x != 0) return; const float* r = probs + row * v; @@ -277,8 +317,7 @@ __global__ void RandomSampleKernel(int64_t* out, const float* probs, const int64 int64_t best = 0; float best_v = kNegInf; for (int64_t j = 0; j < v; ++j) { - const float qn = static_cast(ExpNoise(seed, row, j)); - const float score = r[j] / qn; + const float score = GumbelScore(r[j], seed, row, j); if (score > best_v) { best_v = score; best = j; @@ -287,11 +326,34 @@ __global__ void RandomSampleKernel(int64_t* out, const float* probs, const int64 out[row] = best; } +bool FastRandomSampleEnabled() { + static const bool on = [] { + const char* e = std::getenv("VT_FAST_RANDOM_SAMPLE"); + return e == nullptr || (e[0] != '0'); + }(); + return on; +} + void RandomSampleCuda(Queue& q, Tensor& token_ids, const Tensor& probs, const Tensor& seeds) { const int64_t n = probs.shape[0], v = probs.shape[1]; if (n == 0 || v == 0) return; - RandomSampleKernel<<(n), 1, 0, AsStream(q)>>>( - token_ids.Ptr(), probs.Ptr(), seeds.Ptr(), v); + cudaStream_t s = AsStream(q); + + if (!FastRandomSampleEnabled()) { + RandomSampleKernelSlow<<(n), 1, 0, s>>>( + token_ids.Ptr(), probs.Ptr(), seeds.Ptr(), v); + Check(cudaGetLastError(), "random_sample launch (slow)"); + return; + } + + const int bpr = vt::sample::ArgBlocksPerRow(v, kBlock); + EnsureArgScratch(g_sample_scratch, static_cast(n) * bpr, "random_sample scratch"); + dim3 grid1(static_cast(bpr), static_cast(n)); + ArgmaxPartialKernel<<>>( + g_sample_scratch.val, g_sample_scratch.idx, + GumbelScore2D{probs.Ptr(), seeds.Ptr(), v}, v, bpr); + ArgmaxFinalKernel<<(n), kBlock, 0, s>>>( + token_ids.Ptr(), g_sample_scratch.val, g_sample_scratch.idx, bpr); Check(cudaGetLastError(), "random_sample launch"); } diff --git a/tests/vt/test_ops_sample.cpp b/tests/vt/test_ops_sample.cpp index 3f52fe500..45dae87b2 100644 --- a/tests/vt/test_ops_sample.cpp +++ b/tests/vt/test_ops_sample.cpp @@ -18,6 +18,16 @@ #include "vt/backend.h" #include "vt/dtype.h" #include "vt/ops.h" +#include "vt/sample_common.h" + +#include + +#include +#include +#include +#include +#include +#include using vt::Backend; using vt::Device; @@ -937,3 +947,371 @@ TEST_CASE("ROCm apply_min_p / penalties surface matches CPU mask pattern") { else CHECK(out[i] == doctest::Approx(lc[i]).epsilon(1e-5)); } } + +// =========================================================================== +// #1984 — the parallel Gumbel draw selects the SAME token as the serial scan. +// +// The CUDA kernel cannot be compiled, let alone run, on a CPU-only host, so the +// property it depends on is gated here instead, on the PRODUCTION code it +// depends on: `vt::sample::ArgReduce`, `vt::sample::GumbelScore` and +// `vt::sample::ArgBlocksPerRow` are the same inlines src/vt/cuda/cuda_sample.cu +// compiles. A test that re-typed the reduction would prove the transcription +// rather than the function, which is why they live in a shared header at all. +// +// What is NOT claimed here: that the kernel launches correctly, that its shared +// memory is sized right, or that its scratch is safe. Those need a device, and +// the CUDA equality case further down is what asks them. +namespace { + +using vt::sample::ArgBlocksPerRow; +using vt::sample::ArgReduce; +using vt::sample::GumbelScore; +using vt::sample::kArgSentinel; + +struct ArgPair { + float v; + int64_t i; +}; + +// The serial left-to-right scan the CPU reference performs, expressed over the +// same operator, so "order-independent" is checked against the order that +// actually defines the answer. +ArgPair SerialReduce(const std::vector& scores) { + ArgPair a{kNegInf, kArgSentinel}; + for (int64_t j = 0; j < static_cast(scores.size()); ++j) { + ArgReduce(a.v, a.i, scores[static_cast(j)], j); + } + return a; +} + +// The EXACT two-pass decomposition src/vt/cuda/cuda_sample.cu launches: pass 1 +// gives block `blk` the elements `blk*block + t` strided by `blocks_per_row * +// block` and reduces them within the block; pass 2 reduces the per-block +// partials. `block` mirrors the kernel's kBlock. +ArgPair TwoPassReduce(const std::vector& scores, int block) { + const int64_t v = static_cast(scores.size()); + const int bpr = ArgBlocksPerRow(v, block); + std::vector partials; + partials.reserve(static_cast(bpr)); + for (int blk = 0; blk < bpr; ++blk) { + // Per-thread accumulation, then the in-block tree reduction. + std::vector lanes(static_cast(block), ArgPair{kNegInf, kArgSentinel}); + const int64_t stride = static_cast(bpr) * block; + for (int t = 0; t < block; ++t) { + ArgPair& lane = lanes[static_cast(t)]; + for (int64_t j = static_cast(blk) * block + t; j < v; j += stride) { + ArgReduce(lane.v, lane.i, scores[static_cast(j)], j); + } + } + for (int s = block / 2; s > 0; s >>= 1) { + for (int t = 0; t < s; ++t) { + ArgReduce(lanes[static_cast(t)].v, lanes[static_cast(t)].i, + lanes[static_cast(t + s)].v, lanes[static_cast(t + s)].i); + } + } + partials.push_back(lanes[0]); + } + std::vector lanes(static_cast(block), ArgPair{kNegInf, kArgSentinel}); + for (int t = 0; t < block; ++t) { + for (int j = t; j < bpr; j += block) { + ArgReduce(lanes[static_cast(t)].v, lanes[static_cast(t)].i, + partials[static_cast(j)].v, partials[static_cast(j)].i); + } + } + for (int s = block / 2; s > 0; s >>= 1) { + for (int t = 0; t < s; ++t) { + ArgReduce(lanes[static_cast(t)].v, lanes[static_cast(t)].i, + lanes[static_cast(t + s)].v, lanes[static_cast(t + s)].i); + } + } + return lanes[0]; +} + +// A deterministic 32-bit mixer, so the cases below need no and repeat +// byte for byte on every platform. +uint32_t Mix32(uint32_t x) { + x ^= x >> 16; + x *= 0x7feb352dU; + x ^= x >> 15; + x *= 0x846ca68bU; + x ^= x >> 16; + return x; +} + +// The row shapes that break a careless parallel argmax, plus one ordinary one. +enum class RowKind { kAllEqual, kOneHot, kTopKMasked, kSoftmax, kAllZero, kTailMax }; + +std::vector MakeProbRow(RowKind kind, int64_t v, uint32_t salt) { + std::vector row(static_cast(v), 0.0f); + switch (kind) { + case RowKind::kAllEqual: + // Every prob identical => the score ordering is decided purely by the + // noise, and equal probs are where a tie-break bug is most likely. + for (auto& x : row) x = 1.0f / static_cast(v); + break; + case RowKind::kOneHot: + row[static_cast(Mix32(salt) % static_cast(v))] = 1.0f; + break; + case RowKind::kTopKMasked: + // What #1985 makes the common case: top-k has zeroed all but 20 entries, + // so almost every score is an exact 0.0f and the winner is far from + // index 0. + for (int k = 0; k < 20; ++k) { + const auto idx = static_cast(Mix32(salt + static_cast(k)) % + static_cast(v)); + row[idx] = 0.05f; + } + break; + case RowKind::kSoftmax: { + float sum = 0.0f; + for (int64_t j = 0; j < v; ++j) { + const float e = std::exp( + static_cast(Mix32(salt + static_cast(j)) % 1000u) / 250.0f); + row[static_cast(j)] = e; + sum += e; + } + for (auto& x : row) x /= sum; + break; + } + case RowKind::kAllZero: + // Every score is 0/q == 0: a whole row of ties. The answer must be index + // 0, exactly as the serial scan gives. + break; + case RowKind::kTailMax: + // The winner is the LAST element, so a decomposition that drops the ragged + // tail of a row cannot pass. + row[static_cast(v - 1)] = 1.0f; + break; + } + return row; +} + +} // namespace + +TEST_CASE("ArgReduce is order-independent, including on exact ties") { + // The defect a careless parallelisation introduces: translate the serial + // rule as `if (b.v > a.v) a = b;` and the answer becomes order-DEPENDENT, + // because reducing right-to-left then keeps the HIGHEST tied index. Deleting + // the `bi < ai` clause in vt/sample_common.h must turn this case red. + for (uint32_t salt = 0; salt < 8; ++salt) { + std::vector scores(97); + for (size_t j = 0; j < scores.size(); ++j) { + // Deliberately coarse, so exact ties are the common case rather than a + // measure-zero accident. + scores[j] = static_cast(Mix32(salt * 31u + static_cast(j)) % 5u); + } + const ArgPair forward = SerialReduce(scores); + + std::vector reversed(scores.rbegin(), scores.rend()); + ArgPair backward{kNegInf, kArgSentinel}; + for (int64_t j = static_cast(scores.size()) - 1; j >= 0; --j) { + ArgReduce(backward.v, backward.i, scores[static_cast(j)], j); + } + CAPTURE(salt); + CHECK(backward.i == forward.i); + CHECK(backward.v == forward.v); + + // ...and the same under the kernel's own decomposition, at several block + // widths, so the answer cannot depend on how the row was cut up. + for (const int block : {1, 2, 8, 32, 256}) { + CAPTURE(block); + CHECK(TwoPassReduce(scores, block).i == forward.i); + } + } + + // An unfilled lane never wins, even against an all -inf row: the sentinel is + // what makes index 0 the answer for a fully masked row. + const std::vector all_neg_inf(300, kNegInf); + CHECK(SerialReduce(all_neg_inf).i == 0); + CHECK(TwoPassReduce(all_neg_inf, 256).i == 0); +} + +TEST_CASE("random_sample: the two-pass decomposition equals the serial reference") { + // Every row shape, every vocabulary width that straddles a block boundary, + // several seeds. The reference is `vt::RandomSample` on the CPU backend -- + // the production op, not a re-typed scan. + const std::vector widths = {1, 2, 3, 255, 256, 257, 511, 1000, 65536, 248320}; + const std::vector kinds = {RowKind::kAllEqual, RowKind::kOneHot, + RowKind::kTopKMasked, RowKind::kSoftmax, + RowKind::kAllZero, RowKind::kTailMax}; + for (const int64_t v : widths) { + for (size_t ki = 0; ki < kinds.size(); ++ki) { + for (const int64_t seed : {int64_t{0}, int64_t{1}, int64_t{700}, int64_t{-9}}) { + CAPTURE(v); + CAPTURE(ki); + CAPTURE(seed); + const std::vector row = + MakeProbRow(kinds[ki], v, static_cast(v * 7 + ki)); + + // The reference: the CPU op, one row. + std::vector probs = row; + std::vector seeds = {seed}; + std::vector out = {-1}; + Tensor tp = F32_2(probs, 1, v); + Tensor ts = I64_1(seeds, 1); + Tensor to = I64_1(out, 1); + Queue q = Q(); + vt::RandomSample(q, to, tp, ts); + + // The decomposition, over the same production score expression. + std::vector scores(static_cast(v)); + for (int64_t j = 0; j < v; ++j) { + scores[static_cast(j)] = + GumbelScore(row[static_cast(j)], static_cast(seed), 0, j); + } + CHECK(TwoPassReduce(scores, 256).i == out[0]); + } + } + } +} + +TEST_CASE("random_sample: a batch reduces row by row, and rows do not interfere") { + // The kernel gives each row its own blockIdx.y and its own scratch slice, so + // the batched answer has to equal the per-row answers computed alone. A + // decomposition that indexed the scratch by block alone would pass every + // single-row case above and fail here. + const int64_t n = 5, v = 1000; + std::vector probs(static_cast(n * v)); + std::vector seeds(static_cast(n)); + const std::vector kinds = {RowKind::kAllEqual, RowKind::kTopKMasked, + RowKind::kSoftmax, RowKind::kAllZero, + RowKind::kTailMax}; + for (int64_t i = 0; i < n; ++i) { + seeds[static_cast(i)] = 100 + i; + const std::vector row = + MakeProbRow(kinds[static_cast(i)], v, static_cast(i)); + std::copy(row.begin(), row.end(), probs.begin() + static_cast(i * v)); + } + std::vector out(static_cast(n), -1); + Tensor tp = F32_2(probs, n, v); + Tensor ts = I64_1(seeds, n); + Tensor to = I64_1(out, n); + Queue q = Q(); + vt::RandomSample(q, to, tp, ts); + + for (int64_t i = 0; i < n; ++i) { + CAPTURE(i); + std::vector scores(static_cast(v)); + for (int64_t j = 0; j < v; ++j) { + scores[static_cast(j)] = + GumbelScore(probs[static_cast(i * v + j)], + static_cast(seeds[static_cast(i)]), i, j); + } + CHECK(TwoPassReduce(scores, 256).i == out[static_cast(i)]); + } +} + +// --------------------------------------------------------------------------- +// CUDA: the parallel path and the RETAINED serial path must agree EXACTLY. +// +// Why this is a subprocess A/B rather than two calls. `VT_FAST_RANDOM_SAMPLE` +// is latched in a function-local static on first use (as VT_FAST_ARGMAX is), so +// one process can only ever exercise one arm. The parent below therefore +// re-execs THIS binary twice, once per arm, and requires the printed token ids +// to be byte-identical. That is a same-binary A/B in the sense AGENTS.md means: +// one build, one input, two selections of the code under test. +// +// It has to be exact rather than statistical, and the distinction matters. The +// CPU-vs-CUDA case above can only be statistical, because host and device +// evaluate `-log(u)` through different libm and IEEE-754 does not require a +// correctly-rounded transcendental, so a near-tied row can flip on one ULP. +// Here both arms are the same device libm on the same inputs and differ only in +// the ORDER the identical floats are combined -- and ArgReduce is +// order-independent -- so any difference at all is a defect. +// +// Both cases print WHICH arm they measured, in words, in their own output. +namespace { + +// The shapes the child reports on, in one place so both arms enumerate the same +// ones. 248320 is Qwen3.8-27B's vocabulary, the width #1984 is about. +const std::vector kAbWidths = {1, 2, 255, 256, 257, 1000, 65536, 248320}; + +std::string RunSelf(const char* arm) { + char exe[4096]; + const ssize_t n = ::readlink("/proc/self/exe", exe, sizeof(exe) - 1); + REQUIRE(n > 0); + exe[n] = '\0'; + const std::string cmd = "VT_FAST_RANDOM_SAMPLE=" + std::string(arm) + " " + + std::string(exe) + + " --no-skip --test-case='random_sample_ab_child' 2>&1"; + FILE* pipe = ::popen(cmd.c_str(), "r"); + REQUIRE(pipe != nullptr); + std::string out; + std::array buf{}; + while (std::fgets(buf.data(), static_cast(buf.size()), pipe) != nullptr) out += buf.data(); + REQUIRE(::pclose(pipe) != -1); + // Keep only the payload lines, so a doctest banner or a device warning cannot + // make the two arms differ for a reason that is not the kernel. + std::string ids; + size_t pos = 0; + while (pos < out.size()) { + const size_t eol = out.find('\n', pos); + const std::string line = out.substr(pos, eol == std::string::npos ? eol : eol - pos); + if (line.rfind("IDS ", 0) == 0) ids += line + "\n"; + if (eol == std::string::npos) break; + pos = eol + 1; + } + return ids; +} + +} // namespace + +// The CHILD. Skipped in a normal run; the parent re-execs it by name. +TEST_CASE("random_sample_ab_child" * doctest::skip()) { + if (!HasCuda()) { + std::cout << "IDS no-cuda\n" << std::flush; + std::exit(0); + } + const char* arm = std::getenv("VT_FAST_RANDOM_SAMPLE"); + std::cout << "IDS arm=" << (arm == nullptr ? "unset(fast)" : arm) << "\n"; + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + const std::vector kinds = {RowKind::kAllEqual, RowKind::kTopKMasked, + RowKind::kSoftmax, RowKind::kAllZero}; + for (const int64_t v : kAbWidths) { + const int64_t n = 4; + std::vector probs(static_cast(n * v)); + std::vector seeds(static_cast(n)); + for (int64_t i = 0; i < n; ++i) { + seeds[static_cast(i)] = 700 + i; + const std::vector row = + MakeProbRow(kinds[static_cast(i)], v, static_cast(v + i)); + std::copy(row.begin(), row.end(), probs.begin() + static_cast(i * v)); + } + QueueGuard gq(gpu); + DeviceTensor dp(gpu, gq.q, DType::kF32, {n, v}, probs.data()); + DeviceTensor ds(gpu, gq.q, DType::kI64, {n}, seeds.data()); + DeviceTensor did(gpu, gq.q, DType::kI64, {n}); + vt::RandomSample(gq.q, did.tensor(), dp.tensor(), ds.tensor()); + std::vector ids(static_cast(n)); + did.Download(gq.q, ids.data()); + std::cout << "IDS v=" << v; + for (const int64_t id : ids) std::cout << " " << id; + std::cout << "\n"; + } + std::cout << std::flush; + std::exit(0); +} + +TEST_CASE("CUDA random_sample: the parallel path is BIT-IDENTICAL to the serial one") { + if (!HasCuda()) { + MESSAGE("no CUDA backend registered; skipping (GPU-pending)"); + return; + } + const std::string parallel = RunSelf("1"); + const std::string serial = RunSelf("0"); + MESSAGE("compared VT_FAST_RANDOM_SAMPLE=1 (the two-pass reduction) against " + "VT_FAST_RANDOM_SAMPLE=0 (the retained single-thread scan), same binary"); + INFO("parallel arm:\n" << parallel << "serial arm:\n" << serial); + REQUIRE_FALSE(parallel.empty()); + REQUIRE(parallel.find("IDS v=248320") != std::string::npos); + // Only the arm label may differ. + std::string a = parallel, b = serial; + const auto strip_arm = [](std::string& t) { + const size_t p = t.find("IDS arm="); + if (p == std::string::npos) return; + t.erase(p, t.find('\n', p) - p + 1); + }; + strip_arm(a); + strip_arm(b); + CHECK(a == b); +} From 701a78d855987157882898a4c167ba7cdb3438fc Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 16:10:27 +0000 Subject: [PATCH 04/11] docs(SAMPLE-CORE): --generation-config in the flag reference and the server's own usage line (#1985) A user cannot reach the escape hatch they are told to reach for unless the flag is written down where they look. The server reference gains the row, and `--help` gains the line, alongside the resolved-defaults message the server already prints at startup. The spec records what landed, the six mutations that were applied and restored, and the one reachability hop no CPU test can enter -- the two-line join in server_main, which sits after model load and has no committed generative checkpoint fixture to drive it. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .../sample-gen-config-and-parallel-gumbel.md | 62 +++++++++++++++++++ src/vllm/entrypoints/openai/server_main.cpp | 2 + 2 files changed, 64 insertions(+) diff --git a/.agents/specs/sample-gen-config-and-parallel-gumbel.md b/.agents/specs/sample-gen-config-and-parallel-gumbel.md index 8e197d792..199d02b7b 100644 --- a/.agents/specs/sample-gen-config-and-parallel-gumbel.md +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -200,6 +200,68 @@ request and the predicted values. ABI, which receives explicit `vllm_sampling_params` and has no "omitted" state to fill. +## Outcome so far + +Landed in this pull request, and what each half is actually gated by. + +**The defaults (#1985).** `HfConfig` grows the six sampling keys from the same +sibling read that already produced `generation_config_eos_ids`; +`GetDiffSamplingParam` mirrors `ModelConfig.get_diff_sampling_param` and the +`--generation-config auto|vllm|` selector; both OpenAI handlers resolve a +request against it with upstream's precedence. Red first, on a stubbed +resolution: **3 cases / 9 assertions of `test_openai_protocol` failed for the +intended reason**, then green at 37/37 and 269 assertions. + +One expectation had to be corrected rather than the code, and it is worth +recording because it is the interaction most likely to be got wrong: a request +at temperature 0 is greedy, and `__post_init__` clears `top_p`, `top_k` and +`min_p` (`sampling_params.py`, mirrored at `src/vllm/sampling_params.cpp`). A +checkpoint's `top_k: 20` therefore cannot survive a `--temperature 0` request, +which is what keeps the SACRED greedy path unchanged. It now has its own case. + +**The kernel (#1984).** `RandomSampleCuda` instantiates the greedy two-pass +reduction with the Gumbel score. The output is bit-identical by construction, +and the retained serial kernel behind `VT_FAST_RANDOM_SAMPLE=0` is what lets the +gate assert that rather than assume it. + +**Mutation evidence**, 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, in the pre-existing distribution case | +| 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. `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 the equivalence cases cannot see it — a gate over +a shared helper measures consistency, not correctness. What caught it was the +pre-existing `large-N empirical frequency approximates softmax probs` case, and +that is the right division of labour: the equivalence cases own the ORDER, the +distribution case owns the FORMULA. + +## What is NOT gated here, stated rather than implied + +- **There is no CUDA toolkit on the implementing host**, so `cuda_sample.cu` was + not compiled locally at all. Its first compile is CI's `cuda-fat-build`, and + its first execution is the operator's leased box. Everything the kernel rests + on is gated on CPU through the shared header; the kernel's launch geometry, + shared-memory sizing and scratch lifetime are not, and `compute-sanitizer` is + requested for exactly that reason. +- **Reachability is proven at three hops of four.** The CLI hop re-execs the + real `VllmServerMain`. The handler hop is entered through `create_completion` + / `create_chat_completion`. The narrowing is driven off a real + `generation_config.json` through the production `LoadHfConfig`. The two-line + join in `server_main.cpp` — `completion.set_default_sampling_params(...)` and + its chat twin — sits after model load, and no committed generative checkpoint + fixture can reach it, so no CPU test turns red when it is deleted. That is a + gap, not a waiver; 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 + (164 KiB), which this row did not add. + ## Now Implementation and gates land together in one pull request with this spec (the diff --git a/src/vllm/entrypoints/openai/server_main.cpp b/src/vllm/entrypoints/openai/server_main.cpp index e40eeb055..7ee9b2371 100644 --- a/src/vllm/entrypoints/openai/server_main.cpp +++ b/src/vllm/entrypoints/openai/server_main.cpp @@ -15,6 +15,7 @@ // [--enable-force-include-usage] // [--[no-]enable-prefix-caching] // [--scheduling-policy fcfs|priority] +// [--generation-config auto|vllm|] // [--tool-call-parser |auto|none] // [--reasoning-parser |auto|none] // [--kv-transfer-config ''] [--offload-config ''] @@ -449,6 +450,7 @@ const InertArg* FindAcceptedInertArg(const std::string& flag) { " [--[no-]enable-radix-attention]\n" " [--scheduling-policy fcfs|priority|lpm]\n" " [--[enable|disable]-jump-forward]\n" + " [--generation-config auto|vllm|]\n" " [--tool-call-parser |auto|none]\n" " [--reasoning-parser |auto|none]\n" " [--kv-transfer-config '']\n" From b76aafb1604907f21d80fa5f5f37369110afe571 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 16:13:49 +0000 Subject: [PATCH 05/11] fix(SAMPLE-CORE): name the scratch each allocation is for, and re-anchor the spec on the renamed kernel (#1984) Folding both scratch allocations onto one message would have made a `cudaMalloc` refusal say which pair failed but not which of the two buffers, so the value and index allocations keep their own names as they did before the struct. The spec cited `RandomSampleKernel`, which this row renamed to `RandomSampleKernelSlow` when it kept it as the A/B arm, so `scripts/check-symbol-anchors.py` went red on a citation naming a symbol the file no longer contains. Re-anchored on `RandomSampleCuda` -- the launcher, which is the stable name -- with the rename recorded beside it. `using vt::sample::ExpNoise` went with the local copy it replaced: nothing in cuda_sample.cu calls it any more, because the score expression it fed now comes from `GumbelScore`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .../specs/sample-gen-config-and-parallel-gumbel.md | 7 +++++-- src/vt/cuda/cuda_sample.cu | 13 +++++++------ 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/.agents/specs/sample-gen-config-and-parallel-gumbel.md b/.agents/specs/sample-gen-config-and-parallel-gumbel.md index 199d02b7b..8529e859a 100644 --- a/.agents/specs/sample-gen-config-and-parallel-gumbel.md +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -53,8 +53,11 @@ Read at `555967922` from `${VLLM_SOURCE}`. | `probs.div_(q).argmax(-1)` | `vllm/v1/sample/ops/topk_topp_sampler.py::sample_with_exponential_noise` | | the noise dtype (f32 by default) | `vllm/v1/sample/ops/topk_topp_sampler.py::empty_exponential_noise_like`; `vllm/v1/sample/sampler.py::Sampler.__init__` (`use_fp64_gumbel: bool = False`) | -Local anchors: `src/vt/cuda/cuda_sample.cu::RandomSampleKernel`, -`::ArgmaxPartialKernel`, `::ArgmaxFinalKernel`; +Local anchors: `src/vt/cuda/cuda_sample.cu::RandomSampleCuda` (the defect lived +in the kernel it launched, which this row renamed to +`src/vt/cuda/cuda_sample.cu::RandomSampleKernelSlow` when it kept it as the A/B +arm), `src/vt/cuda/cuda_sample.cu::ArgmaxPartialKernel`, +`src/vt/cuda/cuda_sample.cu::ArgmaxFinalKernel`; `src/vllm/transformers_utils/hf_config.cpp::ReadGenerationConfigEosIds`; `src/vllm/entrypoints/openai/protocol.cpp::CompletionRequest::to_sampling_params`. diff --git a/src/vt/cuda/cuda_sample.cu b/src/vt/cuda/cuda_sample.cu index f70afa513..45b85633d 100644 --- a/src/vt/cuda/cuda_sample.cu +++ b/src/vt/cuda/cuda_sample.cu @@ -49,7 +49,6 @@ unsigned GridFor(int64_t n) { // also includes -- so "bit-identical to the CPU reference" is a property of the // build rather than of two copies staying in step. using vt::sample::ArgReduce; -using vt::sample::ExpNoise; using vt::sample::GumbelScore; using vt::sample::kArgSentinel; @@ -172,12 +171,12 @@ struct ArgScratch { size_t cap = 0; // capacity in elements }; -void EnsureArgScratch(ArgScratch& s, size_t elems, const char* what) { +void EnsureArgScratch(ArgScratch& s, size_t elems, const char* val_what, const char* idx_what) { if (elems <= s.cap) return; if (s.val) cudaFree(s.val); if (s.idx) cudaFree(s.idx); - Check(cudaMalloc(&s.val, elems * sizeof(float)), what); - Check(cudaMalloc(&s.idx, elems * sizeof(int64_t)), what); + Check(cudaMalloc(&s.val, elems * sizeof(float)), val_what); + Check(cudaMalloc(&s.idx, elems * sizeof(int64_t)), idx_what); s.cap = elems; } @@ -228,7 +227,8 @@ void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { // One block per kBlock vocab elements, capped so pass 2 fits a single block. const int bpr = vt::sample::ArgBlocksPerRow(v, kBlock); - EnsureArgScratch(g_argmax_scratch, static_cast(n) * bpr, "argmax scratch"); + EnsureArgScratch(g_argmax_scratch, static_cast(n) * bpr, "argmax scratch val", + "argmax scratch idx"); dim3 grid1(static_cast(bpr), static_cast(n)); ArgmaxPartialKernel<<>>(g_argmax_scratch.val, g_argmax_scratch.idx, LogitScore{logits.Ptr(), v}, v, bpr); @@ -347,7 +347,8 @@ void RandomSampleCuda(Queue& q, Tensor& token_ids, const Tensor& probs, const Te } const int bpr = vt::sample::ArgBlocksPerRow(v, kBlock); - EnsureArgScratch(g_sample_scratch, static_cast(n) * bpr, "random_sample scratch"); + EnsureArgScratch(g_sample_scratch, static_cast(n) * bpr, + "random_sample scratch val", "random_sample scratch idx"); dim3 grid1(static_cast(bpr), static_cast(n)); ArgmaxPartialKernel<<>>( g_sample_scratch.val, g_sample_scratch.idx, From aa9cebe7bf38f51f73509897f53c80effb1fc03e Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 16:17:16 +0000 Subject: [PATCH 06/11] spec(SAMPLE-CORE): say why the measurement must NOT pass --temperature 0 (#1984) Every harness in tools/bench/ passes it -- `run_qwen35_4b_ab.sh` sends `--temperature 0` and `profile_vllm_online_gate.py` sets `temperature=0.0` -- and under that flag `PostInit` clears top_p, top_k and min_p, the batch is all-greedy, and `RandomSample` is never called. All four arms would then be the same run, and four identical numbers read exactly like "no effect". Copying an existing harness line is the likeliest way to get a null result that measured nothing, so the spec now names the trap instead of relying on the flag being left off by luck. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .agents/specs/sample-gen-config-and-parallel-gumbel.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/.agents/specs/sample-gen-config-and-parallel-gumbel.md b/.agents/specs/sample-gen-config-and-parallel-gumbel.md index 8529e859a..17a2ff8e8 100644 --- a/.agents/specs/sample-gen-config-and-parallel-gumbel.md +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -280,6 +280,15 @@ resolve temperature 1.0 and both take the random-sampling path. This is the configuration that made the divergence real, so it is the configuration that has to judge the fix. +**`--temperature 0` would measure nothing, and it is what every existing harness +in `tools/bench/` passes.** `run_qwen35_4b_ab.sh` sends `--temperature 0` and +`profile_vllm_online_gate.py` sets `temperature=0.0`; at temperature 0 +`SamplingParams::PostInit` clears `top_p`, `top_k` and `min_p`, the batch is +all-greedy, and `RandomSample` is never called at all. Under that flag all four +arms below are the same run. Copying an existing harness line is therefore the +one way to get four identical numbers and read them as "no effect", so the flag +has to be dropped deliberately rather than left off by luck. + Four arms, one binary, so each half of the row is attributable on its own: | arm | `--generation-config` | `VT_FAST_RANDOM_SAMPLE` | isolates | From bbb5c97eed56a0bb5bfe631d9245a3e2a20ff804 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 16:23:01 +0000 Subject: [PATCH 07/11] feat(SAMPLE-CORE): the C ABI's chat entry point resolves the same checkpoint defaults the server does (#1985) `vllm_chat` takes the same OpenAI request JSON the HTTP server takes and hands it to the same `OpenAIServingChat::create_chat_completion`, so an omitted `top_k` there is the identical "the client did not say" state. Wiring only the server would have left two entry points of one library sampling from different distributions on one checkpoint, which is the kind of split `include/vllm.h` exists to prevent. There is no C-ABI counterpart of `--generation-config` yet, so `EnsureChatServing` takes upstream's own default for the selector, `"auto"`. Adding the knob is recorded under `## Owed`, together with the note that the struct-shaped entry points (`vllm_complete` and friends) receive an explicit `vllm_sampling_params` and have no omitted state to fill -- that is where an offline `LLM.get_default_sampling_params` equivalent would go. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .../sample-gen-config-and-parallel-gumbel.md | 15 ++++++++++----- src/capi/vllm_c.cpp | 11 +++++++++++ 2 files changed, 21 insertions(+), 5 deletions(-) diff --git a/.agents/specs/sample-gen-config-and-parallel-gumbel.md b/.agents/specs/sample-gen-config-and-parallel-gumbel.md index 17a2ff8e8..618510ca1 100644 --- a/.agents/specs/sample-gen-config-and-parallel-gumbel.md +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -196,12 +196,17 @@ request and the predicted values. see, and on a part with 1:64 f64 throughput it is also 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 and not a rider on this one. -- `--override-generation-config` (`vllm/config/model.py:305`), the +- `--override-generation-config` (`vllm/config/model.py:305`) and the `override_max_tokens` server-wide output cap derived from `max_new_tokens` - (`completion/serving.py:81-86`), and an offline equivalent of - `LLM.get_default_sampling_params` (`vllm/entrypoints/llm.py:404`) for the C - ABI, which receives explicit `vllm_sampling_params` and has no "omitted" - state to fill. + (`completion/serving.py:81-86`). `DefaultSamplingParams::max_tokens` already + carries the value; nothing reads it. +- A C-ABI counterpart of `--generation-config`. `vllm_chat` takes the same + OpenAI request JSON the server takes, so it resolves the checkpoint's defaults + through the same handler; it simply cannot be told to use `"vllm"` instead. + The struct-shaped entry points (`vllm_complete` and friends) take an explicit + `vllm_sampling_params` and have no "omitted" state to fill, which is where + `LLM.get_default_sampling_params` (`vllm/entrypoints/llm.py:404`) would come + in. ## Outcome so far diff --git a/src/capi/vllm_c.cpp b/src/capi/vllm_c.cpp index cb5209885..26730f635 100644 --- a/src/capi/vllm_c.cpp +++ b/src/capi/vllm_c.cpp @@ -37,6 +37,7 @@ #include "vllm/entrypoints/chat_template.h" #include "vllm/entrypoints/model_loader.h" #include "vllm/entrypoints/openai/protocol.h" +#include "vllm/config/generation.h" #include "vllm/entrypoints/openai/serving_chat.h" #include "vllm/entrypoints/openai/serving_utils.h" #include "vllm/entrypoints/openai/tool_parsers/abstract.h" // get_tool_parser @@ -385,6 +386,16 @@ vllm::entrypoints::openai::OpenAIServingChat& EnsureChatServing( engine->loaded->async_engine(), std::move(served_name), std::move(prompt_fn), std::move(parser_name), std::move(reasoning_name)); + // #1985: `vllm_chat` takes the SAME OpenAI request JSON the HTTP server + // takes, so an omitted `top_k` here is the identical "the client did not + // say" state -- and it has to resolve the identical way, or two entry + // points of one library sample from different distributions on one + // checkpoint. There is no C-ABI counterpart of `--generation-config` yet, + // so this takes upstream's own default for the selector, `"auto"` + // (config/model.py:298); adding the knob is recorded as owed. + engine->chat_serving->set_default_sampling_params( + vllm::GetDiffSamplingParam(engine->loaded->config(), + vllm::kGenerationConfigAuto)); } return *engine->chat_serving; } From 94152f2e7f1c41e2bf5915a9284d877fe59c0f9f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 16:23:55 +0000 Subject: [PATCH 08/11] fix(SAMPLE-CORE): allowlist VT_FAST_RANDOM_SAMPLE beside VT_FAST_ARGMAX (#1984) `scripts/check-env-doc.py` refuses a production environment variable that is neither documented nor allowlisted, and it went red on the new A/B selector. It belongs on the allowlist rather than in docs/ENVIRONMENT.md for the same reason its twin does: `VT_FAST_ARGMAX` selects between a kernel and the reference scan it replaced, and so does this one. Neither is a user-facing knob; both exist so a performance claim can be a same-binary A/B. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- scripts/env-doc-allowlist.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index 93409e181..e17b4c2d8 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -33,6 +33,7 @@ VT_FA2_DECODE_QWEN3 VT_FA2_NSPLITS_CAP VT_FA2_PREFILL_QWEN3 VT_FAST_ARGMAX +VT_FAST_RANDOM_SAMPLE VT_FP4_AUTOTUNE VT_FP4_AUTOTUNE_CACHE_PATH VT_FP4_AUTOTUNE_CACHE_READONLY From 6c37ab60101f685307871a5869aa5027d7a919ac Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 17:33:03 +0000 Subject: [PATCH 09/11] record(SAMPLE-CORE): the acceptance measurement must drop --speculative-config, because drafts bypass the sampler entirely (#2002) `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. The standing baseline recipe carries `--speculative-config '{"method":"dflash",...}'`, so a before/after taken on it would have returned two identical numbers for the kernel this row rewrites, and two identical numbers read exactly like "the change did nothing". Filed as #2002 rather than fixed here, because the same routing is a correctness divergence with its own scope: upstream's `rejection_sample` carries both the greedy and the stochastic accept, ours implements only `is_greedy`, and `include/vllm/v1/spec_decode/rejection_sampler.h` already states the contract nothing enforces -- "a temperature > 0 request must NOT be routed here yet". Refusing such a request by name is small; porting `_resample_kernel` is a row. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .agents/issue-index.md | 1 + 1 file changed, 1 insertion(+) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index b166ecdd1..31684da30 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -763,3 +763,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#2050](https://github.com/mudler/vllm.cpp/issues/2050) | `ENG-HYBRID-PLACEMENT` | Laguna's FFN is host-orchestrated token-at-a-time — per-token host rows, the router on the host through `MatmulNK`, and a host scalar combine loop — so a device-shaped MoE entry wrapping those loops would put it in the placement seam's wired list while moving nothing and adding a round trip: supported to read, a regression to measure. The real repair is a device-resident batched FFN, which is a model rework with a performance gate | gap | | [#1984](https://github.com/mudler/vllm.cpp/issues/1984) | `SAMPLE-CORE` | `RandomSampleKernel` is launched `<<>>` and scans a 248,320-wide vocab on one thread per row, computing two `SplitMix64` rounds and an f64 `log` per element. Eleven lines above it the same file records that a single-block single-thread scan of a ~151k vocab cost ~7.5 ms/token, which is why greedy argmax was rewritten into `ArgmaxPartialKernel`/`ArgmaxFinalKernel`; the Gumbel draw never got that treatment. Upstream is whole-tensor (`vllm/v1/sample/ops/topk_topp_sampler.py::sample_with_exponential_noise`), so this is a mirror obligation. Reached by every non-greedy row through `ModelRunner::execute_model` -> `Sampler::forward` -> `vt::RandomSample`. Spec: [sample-gen-config-and-parallel-gumbel.md](specs/sample-gen-config-and-parallel-gumbel.md) | perf | | [#1985](https://github.com/mudler/vllm.cpp/issues/1985) | `SAMPLE-CORE` | `generation_config.json` is read for `eos_token_id` only (`hf_config.cpp::ReadGenerationConfigEosIds`), so `Qwen/Qwen3.8-27B`'s shipped `top_k: 20` / `top_p: 0.95` never reach `SamplingParams` and `to_sampling_params` resolves omitted knobs straight to the neutral OpenAI defaults, which disable both filters. vLLM applies them through `ModelConfig.get_diff_sampling_param` -> `OpenAIServing*.default_sampling_params` -> `to_sampling_params`. Since `vllm bench serve` stopped sending `--temperature`, both engines sample at temperature 1.0 and vLLM draws from 20 candidates while we draw from 248,320: different sampling on two sides of a parity benchmark. Spec: [sample-gen-config-and-parallel-gumbel.md](specs/sample-gen-config-and-parallel-gumbel.md) | bug | +| [#2002](https://github.com/mudler/vllm.cpp/issues/2002) | `SAMPLE-CORE` | With `--speculative-config` set, `GPUModelRunner::sample_tokens` branches on `num_draft_tokens > 0` alone and returns the greedy-only `RejectionSampler`'s output, so `Sampler::forward` and `vt::RandomSample` are never called and a `temperature: 1.0` request decodes GREEDILY. `include/vllm/v1/spec_decode/rejection_sampler.h` states the contract it violates in its own deferral list ("a temperature > 0 request must NOT be routed here yet"); neither the runner nor `RejectionSampler::forward` enforces it. Found while writing #1984's acceptance measurement against a baseline recipe carrying `--speculative-config`, where the sampler under test would never have been launched and the null result would have read as "the change did nothing" | bug | From d0b3f7cb0f01a6ae5661d40716ea0c96a1b7c0ec Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 26 Aug 2026 17:33:38 +0000 Subject: [PATCH 10/11] spec(SAMPLE-CORE): drop --speculative-config from every measurement arm, and record why (#2002) The arms table said nothing about speculative decoding, and the operator's standing baseline recipe carries `--speculative-config '{"method":"dflash",...}'`. With drafts present `GPUModelRunner::sample_tokens` returns the greedy-only `RejectionSampler`'s output and never reaches `Sampler::forward`, so the kernel this row rewrites would not have been launched once during the measurement. Four identical numbers would then have read as "no effect". The same routing is a correctness divergence in its own right and is now under `## Owed`: a temperature-1.0 request under spec decode takes the target argmax where vLLM draws from the residual distribution, and the rejection sampler's own header already says such a request must not be routed there. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .../sample-gen-config-and-parallel-gumbel.md | 20 +++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/.agents/specs/sample-gen-config-and-parallel-gumbel.md b/.agents/specs/sample-gen-config-and-parallel-gumbel.md index 618510ca1..ee1ec85a1 100644 --- a/.agents/specs/sample-gen-config-and-parallel-gumbel.md +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -186,6 +186,15 @@ request and the predicted values. and `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. +- [#2002](https://github.com/mudler/vllm.cpp/issues/2002): the spec-decode + routing above is also a correctness divergence in its own right. Upstream's + `rejection_sample` carries both the greedy and the stochastic accept; ours + implements only `is_greedy`, so a `temperature: 1.0` request under + `--speculative-config` gets the target argmax where vLLM draws from the + residual distribution. `include/vllm/v1/spec_decode/rejection_sampler.h` + states the contract nothing enforces: "a temperature > 0 request must NOT be + routed here yet". Refusing such a request by name is small; porting the + stochastic accept and `_resample_kernel` is a row of its own. - `src/vt/rocm/rocm_sample.hip::RandomSampleK` is launched `<<>>` and carries its own copy of `ExpNoise` (`rocm_sample.hip:38`). It is the same defect as #1984 on a backend this row has no hardware to gate, so it is left @@ -285,6 +294,17 @@ resolve temperature 1.0 and both take the random-sampling path. This is the configuration that made the divergence real, so it is the configuration that has to judge the fix. +**`--speculative-config` would measure nothing either, and the operator's +standing baseline recipe carries it +([#2002](https://github.com/mudler/vllm.cpp/issues/2002)).** +`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 therefore `vt::RandomSample` -- is never +called, whatever the temperature. A before/after on a speculative recipe returns +two identical numbers, and two identical numbers read exactly like "the change +did nothing" rather than like "the kernel under test never ran". So every arm +below runs WITHOUT `--speculative-config`. + **`--temperature 0` would measure nothing, and it is what every existing harness in `tools/bench/` passes.** `run_qwen35_4b_ab.sh` sends `--temperature 0` and `profile_vllm_online_gate.py` sets `temperature=0.0`; at temperature 0 From ab4c38244b4f9413e72ae43dd69b7d1b4bbc3f03 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 27 Aug 2026 00:17:05 +0000 Subject: [PATCH 11/11] record(SAMPLE-CORE): the benchmark checkpoint ships the sampling defaults, so landing this re-baselines our arm (#1985, #1984) The premise was verified against the public `Qwen/Qwen3.8-27B` file; the operator confirmed it on the checkpoint the project actually benchmarks. `qwen3.8-27b-nvfp4-mtp-sm121-r0b0tlab` ships `temperature: 1.0`, `top_k: 20`, `top_p: 0.95`, so every benchmark run on it has sampled the full 248,320-token vocabulary while vLLM sampled from twenty candidates. That is a per-token cost we inflicted on ourselves and an asymmetry inside the comparison itself. The consequence is recorded because it is easy to misread later: a number taken before this merges is not comparable with one taken after, and not because anything regressed -- the two are different sampling configurations. The arm re-baselines at the merge commit, and `--generation-config vllm` reproduces the old resolution exactly if an old figure has to be re-derived. Also recorded: #1994 voids TTFT and TPOT measured on older trees, and nothing in this row quotes either. The only measured figures it carries are `cuda_sample.cu`'s own ~7.5 ms/token greedy-scan anchor, which is a per-kernel attribution rather than a harness TPOT, and #1929's 16.9 tok/s, which is throughput. #1994's own spec lists `/v1/completions` among the paths it does not touch, so the acceptance arms are unaffected; on the chat backend the residual bias would overstate the greedy-versus-default gap by about 1.6%, which moves the acceptance bar by under a millisecond. And the CUDA half now has a compile verdict instead of a hope: `cuda-fat-build` returned success at `e5b9b0045` across all ten shipped architectures, with both sanitizers green. It is a compile verdict only -- CI has no device, so no CUDA test ran -- and it belongs to that tree, so the merged tree owes its own. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code] --- .../sample-gen-config-and-parallel-gumbel.md | 48 +++++++++++++++++++ 1 file changed, 48 insertions(+) diff --git a/.agents/specs/sample-gen-config-and-parallel-gumbel.md b/.agents/specs/sample-gen-config-and-parallel-gumbel.md index ee1ec85a1..fdb023ddd 100644 --- a/.agents/specs/sample-gen-config-and-parallel-gumbel.md +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -260,8 +260,56 @@ pre-existing `large-N empirical frequency approximates softmax probs` case, and that is the right division of labour: the equivalence cases own the ORDER, the distribution case owns the FORMULA. +## Landing this re-baselines our benchmark arm + +The premise is confirmed on the checkpoint the project actually benchmarks, not +only on the public `Qwen/Qwen3.8-27B` file: `qwen3.8-27b-nvfp4-mtp-sm121-r0b0tlab` +ships a `generation_config.json` carrying +`{"do_sample": true, "temperature": 1.0, "top_k": 20, "top_p": 0.95}` +(operator, 2026-08-26). So every benchmark this project has run on that +checkpoint sampled over the full 248,320-token vocabulary while vLLM sampled +from top-20 — a per-token cost we inflicted on ourselves, and a live asymmetry +in the comparison itself. + +**A number taken before this merges is therefore not comparable with one taken +after**, and not because anything regressed: the two are different sampling +configurations. Our arm has to be re-baselined at the merge commit, and +`--generation-config vllm` reproduces the old resolution exactly if an old +figure ever has to be re-derived. + +## #1994 voids TTFT and TPOT on older trees; nothing here quotes either + +`b758127ec` (#1994) fixed a chat SSE role frame emitted before any engine work, +which stopped `vllm bench serve --backend openai-chat`'s TTFT clock on an empty +frame and — since `TPOT = (E2E - TTFT)/(tokens - 1)` — inflated TPOT too. + +Checked against this row rather than assumed. Neither this spec nor the pull +request body quotes a TTFT or a TPOT figure at all. The only measured numbers +either carries are `~7.5 ms/token`, which is `cuda_sample.cu`'s own long-standing +anchor for the single-thread greedy scan and a per-kernel decode-path +attribution rather than a harness TPOT, and `16.9 tok/s` from #1929, a +throughput figure and so unaffected. Everything else is marked as prediction. + +The acceptance arms are safe for a reason stated in #1994's own spec rather than +in its summary: it lists `/v1/completions` among the paths it does NOT touch and +scopes the defect to `--backend openai-chat`. Run on the chat backend the +residual bias would be bounded and small — TTFT understated by about one decode +step, spread over 63 tokens, inflating the greedy arm by ~1.7 ms and the default +arm by ~5.4 ms, so overstating the greedy-versus-default gap by ~1.6%. That +moves the acceptance bar by well under a millisecond and changes no verdict. + ## What is NOT gated here, stated rather than implied +- **The CUDA half now COMPILES, measured rather than hoped.** `cuda-fat-build` + returned `success` at `e5b9b0045` over + `80;86;87;89;90a;100a;103a;110;120a;121a`, so the templated partial kernel, the + `__host__ __device__` shared header and both instantiations build under nvcc + for every shipped architecture. Both `sanitize-cpu` jobs also returned + `success`. This is a COMPILE verdict and nothing more: no CUDA test executed, + because CI has no device. It is also a verdict on THAT tree — a later re-merge + brought `include/vt/device.h` changes from main, and although the three + sampling sources are byte-identical across the move, the merged tree owes its + own `cuda-fat-build`. - **There is no CUDA toolkit on the implementing host**, so `cuda_sample.cu` was not compiled locally at all. Its first compile is CI's `cuda-fat-build`, and its first execution is the operator's leased box. Everything the kernel rests