diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 7cd159742..31684da30 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -761,3 +761,6 @@ 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 | +| [#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 | 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..fdb023ddd --- /dev/null +++ b/.agents/specs/sample-gen-config-and-parallel-gumbel.md @@ -0,0 +1,395 @@ +# `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::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`. + +## 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. +- [#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 + 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`) and the + `override_max_tokens` server-wide output cap derived from `max_new_tokens` + (`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 + +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. + +## 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 + 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 +`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. + +**`--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 +`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 | +|---|---|---|---| +| 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. 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..964899757 100644 --- a/include/vllm/entrypoints/openai/serving_chat.h +++ b/include/vllm/entrypoints/openai/serving_chat.h @@ -245,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 @@ -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/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/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 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; } 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..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 ''] @@ -78,6 +79,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 +314,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 @@ -439,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" @@ -631,6 +643,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 +1450,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/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..45b85633d 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,12 @@ 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::GumbelScore; +using vt::sample::kArgSentinel; // --- apply_temperature ------------------------------------------------------ __global__ void ApplyTemperatureKernel(float* logits, const float* temp, int64_t n, int64_t v, @@ -91,20 +84,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 +122,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 +161,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* 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)), val_what); + Check(cudaMalloc(&s.idx, elems * sizeof(int64_t)), idx_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 +225,15 @@ 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 val", + "argmax scratch idx"); 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,35 @@ __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 val", "random_sample scratch idx"); + 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/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()); + } +} 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); +}