diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 506af525e..90bb28f94 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -474,6 +474,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1394](https://github.com/mudler/vllm.cpp/issues/1394) | `ENG-CUDAGRAPH-BREAK` | The CPU paged attention reads `btab[r * bt_row + (j / block_size) * bt_col]` for every `j < seq_lens[r]` without checking that the block table has that many columns, so a caller with a short table gets an out-of-bounds read, a plausible block index out of it, and attention over the WRONG page — silently. `tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp`'s `SpecAttnMeta` supplies one: hardcoded `block_table_num_cols = 1` against shape C's `seq_lens = 24` at `block_size = 16`. Found while fixing [#1380](https://github.com/mudler/vllm.cpp/issues/1380), whose `DevicePool` change moved the bytes after the table and turned the same read into a SIGSEGV on `thor:gpu0` (`gdb` at `src/vt/cpu/cpu_paged_attn.cpp:224` under `FullAttnBlockPaged`). PRE-EXISTING: the case passes at `origin/main` only because the read landed on bytes that decoded to an in-range index. FIXED IN FLOW in both halves — the kernel refuses a short table with one compare per request outside the token loop, and the helper sizes its table for the sequence length it declares. Owned by row `ENG-CUDAGRAPH-BREAK` | bug | | [#1387](https://github.com/mudler/vllm.cpp/issues/1387) | `ENG-EXPERT-STREAM-DEVICE` | `docs/FEATURES.md`'s routed-expert-streaming row still read "CPU keep-quant towers only" after W0c made the seam take the slot arm on `is_cpu()` OR `host_memory_is_device_addressable()` and W0d made the load-time fit refusal drop those towers from its bound. `AGENTS.md` routes a feature-surface change to that page, and the change that moved the surface did not write it. `scripts/check-doc-checkpoint.py` said so, on commit `939755f99` of `row/ENG-EXPERT-STREAM-DEVICE-W0`: a measurement was appended to `.agents/benchmark-record.md` with no `docs/FEATURES.md` edit beside it. The PAGE is FIXED IN FLOW while repairing the fresh review of [#1377](https://github.com/mudler/vllm.cpp/pull/1377): the row now names both the device arm ([#1124](https://github.com/mudler/vllm.cpp/issues/1124)) and the residency condition ([#1378](https://github.com/mudler/vllm.cpp/issues/1378)), and the observability detail it displaced to stay inside the 220-character cell budget is stated at `docs/USAGE.md:4598-4620`. The GATE is NOT fixed and needs a decision: the checker walks a range one COMMIT at a time, so once a commit is published on a branch that may not be force-pushed, no later commit can make it green, and `scripts/agent-preflight.sh` keeps reporting `doc-checkpoint range` red on this branch until it merges. The squashed commit that lands on `main` carries both paths and passes. Changing the walk is checker semantics and needs its own row, spec and red-first evidence per `AGENTS.md` "Changing the rules or a checker", so it is not folded in here. Spec [`expert-stream-device-slots.md`](specs/expert-stream-device-slots.md) | bug | | [#1386](https://github.com/mudler/vllm.cpp/issues/1386) | — | `tools/bench/gpu_clock_state.py`'s `QUERY_FIELDS` collects nine fields and **none of them is thermal and none is electrical**, so the driver's own `SwThermalSlowdown` label can never be checked against a die reading on any window this helper has ever recorded. The measured consequence is that the nine windows of 2026-08-19 cannot distinguish a load transition from a thermal excursion. The concrete evidence is `clock-c1-r1.samples.json` in `/mnt/nas_share/rc/q38bf16/out/bench-20260819T035148Z/`: ours c1 r1 dips five times on the same period at the same `utilization.gpu = 96` — 48.83 s / 2177 MHz, 80.60 s / 2320 MHz, 109.28 s / 2210 MHz, 137.98 s / 2359 MHz, 166.07 s / 2268 MHz — and **two of those five carry `0x0000000000000000`**, no throttle bit at all (2210 and 2359), while three carry `0x20`. The 2210 MHz unlabelled dip is deeper than two of the three labelled ones, so the driver labels comparable excursions inconsistently and the bit alone cannot decide it. What would settle it: add `temperature.gpu` and `power.draw` to `QUERY_FIELDS`. That changes the clock-record schema, so it owes its own row and spec. Split out of [#1354](https://github.com/mudler/vllm.cpp/issues/1354) and owed under `## Owed` in [lease-clock-pinning.md](specs/lease-clock-pinning.md) | gap | +| [#1313](https://github.com/mudler/vllm.cpp/issues/1313) | `SAMPLER-HOST-LOGITS-DEVICE-TAG` | The `[rows, vocab]` logits tensor the ON-DEVICE sampler runs on was built from a host `std::vector` pointer stamped with the queue device, at FOUR sites (three in `src/vllm/v1/worker/gpu/runner.cpp::assemble_sample_logits`, one in `src/vllm/v1/worker/gpu/runner.cpp::collect_prompt_logprobs`) and not the one the issue names: the (A') `VT_GPU_SAMPLE=0` download-then-sample A/B leg, the reported host-logits leg, the (B) `VT_LOGITS_GATHER=0` re-gather leg, and the `collect_prompt_logprobs` prompt-row slice. CONFIRMED rather than refuted: `vt::GreedyArgmax` does NOT copy its input — `GreedyArgmaxCuda` hands `logits.Ptr()` straight to a kernel on both arms (`src/vt/cuda/cuda_sample.cu::GreedyArgmaxKernelSlow` slow, `src/vt/cuda/cuda_sample.cu::ArgmaxPartialKernel` fast), and so do `ApplyTemperature` / `ApplyTopKTopP` / `ComputeProbs` / `ComputeLogprobs` through the same gate (`src/vt/ops.cpp::CheckSamplingLogits`), so the address IS dereferenced on device. Invisible on GB10 because `CudaBackend::UnifiedMemory()` is `pageable_memory_access && integrated` (`src/vt/cuda/cuda_backend.cu::UnifiedMemory`), true there, and the driver services the host pointer through ATS at host-memory latency — which is why the NemotronH A3 gate reads `96/96 STRICT PASS` over it. A discrete GPU reports `integrated == false` and the same address is illegal. Affects exactly the three models that return `ForwardLogits.host`, i.e. the `scripts/runner-routing-allowlist.txt` set: `nemotron_h`, `laguna`, `qwen3_vl`. FIXED by routing all four sites through a new `HostBufferStaging` beside `DeviceScratch` in `include/vllm/v1/sample/device_scratch.h`, which applies the residency contract that header already states (unified wraps in place 0-copy, discrete allocs and copies up) with one grow-only allocation instead of an alloc/free per step, because this sits on the per-token decode path. The unified branch is the replaced expression unchanged, so GB10 does not move. The issue's suggested `Tensor::Contiguous` residency assertion was REJECTED and the reason recorded in the spec: two registered backends legitimately stamp a host-dereferenceable pointer with a non-unified device — Tenstorrent, whose `Alloc` returns `aligned_alloc` host memory by design (`src/vt/tenstorrent/tenstorrent_backend.cpp::Alloc`, with `src/vt/tenstorrent/tenstorrent_backend.cpp::UnifiedMemory` returning false), and a discrete Vulkan device, whose buffers are HOST_VISIBLE and persistently mapped (`include/vt/backend.h::DeviceMemoryIsHostAddressable`) — and both register a `kGreedyArgmax` provider, so a blanket assert false-fires on both. Spec [sampler-host-logits-device-tag.md](specs/sampler-host-logits-device-tag.md) | bug | | [#1074](https://github.com/mudler/vllm.cpp/issues/1074) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | The NemotronH model-matrix row described `main` as of 2026-08-12 and stayed `INVENTORIED` while A2-R, A2-P, A2-Q2a and the A3 driver landed on top of it. Reconciled against `main` at `b626be75a`, every claim re-verified rather than inherited: the non-gated `relu²` MoE EXISTS (`4d0c399e1`, `vt::MoeRelu2` called from `nemotron_h.cpp:354`), ModelOpt `MIXED_PRECISION` per-module loading EXISTS (`1bc5ef82c`), the MTP head is STILL OWED (W5, 270 tensors deferred by name), and the `KERNEL-SSM-MAMBA` block is FALSE — [#496](https://github.com/mudler/vllm.cpp/issues/496) W1 landed the host arm at `47960a009` and W2 the CUDA arm at `43a6c5518`, in `src/vt/cuda/cuda_mamba2_ssd.cuh`, a `.cuh` included by `cuda_gdn.cu` rather than a translation unit of its own, which is why a `src/vt/*mamba*` FILE GLOB reports absence; `nemotron_h.cpp:597` calls that op today. Row moved `INVENTORIED` -> `PARTIAL` with the rollup, the checklist entry and the projections it owes. `PARTIAL` and not `ACTIVE` deliberately: `check-agent-record.py` requires an `ACTIVE` row to name a `CLAIM-*` row a claim source carries, no file under `.agents/claims/` claims this row, and authoring one for another session's in-flight work would be a fabricated record. NO end-to-end token gate has passed and no throughput, latency or memory number is claimed; the A3 gate stays PENDING and [#1157](https://github.com/mudler/vllm.cpp/issues/1157) is the open decode divergence | bug | | [#1080](https://github.com/mudler/vllm.cpp/issues/1080) | — | `scripts/check-doc-checkpoint.py:153` matches a spec's live-position section with `^##\s+Now\s*$`, and specs in this tree write it as `## N. Now`, so `spec_now_errors` reports "has no `## Now` section" about a section that is present and current. It fires only when a row moves lifecycle state. `nemotron-h-model.md` was one of them and is repaired in flow by [#1074](https://github.com/mudler/vllm.cpp/issues/1074), which is the change that made it the spec a moving row links. Re-measured at `b626be75a` after that repair: 15 specs still write the numbered spelling — `gate-audit-branch-evidence`, `ltx25-a2v-audio-input`, `ltx25-image-conditioning`, `ltx25-t2a-one-stage`, `ltx25-token-append`, `ltx2-device-staged-view-uaf`, `mamba2-ssd`, `nas-mount-path`, `nemotron-h-a2p-paged-forward`, `nemotron-h-a2q1-fp8-mamba`, `nemotron-h-a2q2b-realckpt-lmhead`, `nemotron-h-a2q2-nvfp4-moe-lmhead`, `nemotron-h-abi-e2e`, `offload-docs-refusal`, `registry-downcast-sweep`. The population GREW from the twelve this was filed against, which is the argument for the checker-semantics fix over a rename sweep. Either close needs its own spec and a red-before test, so neither rides in a records reconcile. `tests/scripts/` covers `NOW_SECTION` nowhere, which is how the mismatch survived. Listed under `## Owed` in [`nemotron-h-model.md`](specs/nemotron-h-model.md) | bug | | [#1410](https://github.com/mudler/vllm.cpp/issues/1410) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | `scripts/check-runner-routing-consistency.py` resolves the registry hook's delegate across translation units only for the `Class::ForwardDevice` shape (`_DELEGATE`, `:142-145`) or a helper defined in the registry TU itself (`classify_with_helpers`, `:246-271`). NemotronH's production device forward is a FREE function in another TU — `ForwardNemotronHForCausalLM` (`nemotron_h_registry.cpp`) calls `NemotronHPagedForward` (`nemotron_h_device.cpp`) — so the hop finds nothing and the hook's own host-reference fall-through, which A2-P deliberately keeps below the paged fold as the numeric gate's operand, classifies the model HOST. Measured on `row/A2-Q2b-lmhead-nvfp4` after A2-Q2b put `lm_head` on the device: `NemotronHPagedForward` assigns both `fl.device_tensor` and `fl.device_storage = dlogits.ReleaseShared()`, which IS `_DEVICE_SEAM` (`:125-128`), and the checker still names the model. A FALSE RED, so the safe direction — but it holds an allowlist entry open for a clause that is MET, and the allowlist is what a reader trusts to know what is still unrouted; it is latent the other way for any future model whose device logits come from a cross-TU free function. The checker already builds the `free_fn_file` map the fix needs (invariant (b) uses it). NOT fixed in flow: it CHANGES CHECKER SEMANTICS, which `AGENTS.md` `## Changing the rules or a checker` routes to the normal row, spec and fresh-review path, and widening a classification to turn a red gate green is exactly the move that section slows down — it needs its own red-before in `tests/scripts/test_check_runner_routing_consistency.py`. Owned by row `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm`, listed under `## 5. Owed` in [`nemotron-h-a2q2b-realckpt-lmhead.md`](specs/nemotron-h-a2q2b-realckpt-lmhead.md) | bug | diff --git a/.agents/specs/sampler-host-logits-device-tag.md b/.agents/specs/sampler-host-logits-device-tag.md new file mode 100644 index 000000000..baebb2ca0 --- /dev/null +++ b/.agents/specs/sampler-host-logits-device-tag.md @@ -0,0 +1,142 @@ +# SAMPLER-HOST-LOGITS-DEVICE-TAG — a host pointer stamped with the CUDA device + +Issue: [#1313](https://github.com/mudler/vllm.cpp/issues/1313) +Row: SAMPLER-HOST-LOGITS-DEVICE-TAG +State: ACTIVE + +## Scope + +`src/vllm/v1/worker/gpu/runner.cpp` builds the `[rows, vocab]` logits tensor the +on-device sampler runs on from a host `std::vector` pointer, and stamps +`queue_.device` on it. On a unified-memory device the address is valid. On a +discrete GPU it is a host address handed to a kernel. + +In scope: every site in `runner.cpp` that views a host logits buffer as a tensor +on the queue device, and the shared staging seam they route through. + +Out of scope: removing the host-logits path itself (A2-Q2b for `nemotron_h`, the +framework-routing rows for `laguna` and `qwen3_vl`). This row makes the existing +path correct on every backend; it does not change which models take it. + +## The defect, grounded + +`vt::GreedyArgmax` does NOT copy its input before touching it. The CUDA arm +passes the caller's pointer straight to a kernel: + +* `src/vt/cuda/cuda_sample.cu::GreedyArgmaxKernelSlow` — `GreedyArgmaxKernelSlow<<<...>>>(..., logits.Ptr(), v)` +* `src/vt/cuda/cuda_sample.cu::ArgmaxPartialKernel` — `ArgmaxPartialKernel<<<...>>>(..., logits.Ptr(), v, bpr)` + +So the pointer IS dereferenced on device. This is a latent illegal access on a +discrete GPU, not a mislabelling. + +The same is true of the rest of the sampling pipeline the tensor feeds +(`ApplyTemperature`, `ApplyTopKTopP`, `ComputeProbs`, `ComputeLogprobs` — +`src/vt/ops.cpp::CheckSamplingLogits`), and of `apply_grammar_bitmask`. + +## Affected sites (four, not one) + +| Site | Path | Host buffer | +|---|---|---| +| `src/vllm/v1/worker/gpu/runner.cpp::assemble_sample_logits` | (A') `VT_GPU_SAMPLE=0` download-then-sample A/B | `sampled_logits` | +| `src/vllm/v1/worker/gpu/runner.cpp::assemble_sample_logits` | host logits, rows already gathered (the reported one) | `fl.host` | +| `src/vllm/v1/worker/gpu/runner.cpp::assemble_sample_logits` | (B) `VT_LOGITS_GATHER=0` host re-gather | `sampled_logits` | +| `src/vllm/v1/worker/gpu/runner.cpp::collect_prompt_logprobs` | `collect_prompt_logprobs` prompt-row slice | `fl.host + offset` | + +Affected models are the three that return `ForwardLogits.host`, i.e. exactly the +`scripts/runner-routing-allowlist.txt` entries: `nemotron_h`, `laguna`, +`qwen3_vl`. + +## Why it has not bitten + +`vt::CudaBackend::UnifiedMemory()` is `caps.pageable_memory_access && +caps.integrated` (`src/vt/cuda/cuda_backend.cu::UnifiedMemory`). GB10 is integrated, so the +driver services an ordinary host pointer through ATS and the NemotronH A3 gate +reads `96/96 STRICT PASS` at host-memory latency. A discrete GPU reports +`integrated == false` and the same address is illegal. + +## Design + +The repository already states this contract, in +`include/vllm/v1/sample/device_scratch.h`: + +> unified-memory backends (CPU, GB10) wrap the host buffer in place (0-copy); +> discrete backends alloc device memory and copy the host buffer up + +`DeviceScratch` applies it to the small derived sampling tensors. The logits +tensor — the largest one, and the one every sampling kernel reads — bypasses it. + +Add `HostBufferStaging` beside `DeviceScratch`: the same residency contract, with +one grow-only device allocation reused across steps instead of an alloc/free per +construction, because this sits on the per-token decode path. Route all four +sites through it. + +Two staging members, not one: `collect_prompt_logprobs` runs while the assembled +sample-logits tensor is still live (`src/vllm/v1/worker/gpu/runner.cpp::sample_tokens`), so one shared buffer +would invalidate it. + +### Alternatives rejected + +**The sampler refuses a host pointer carrying a device tag.** Rejected. It +converts a repairable path into a refusal: the three models would stop decoding +on a discrete GPU rather than start working. It also cannot tell a GB10-valid +host pointer from a discrete-invalid one without a per-backend residency probe, +which is the same new seam the repair needs anyway — so it costs what the repair +costs and delivers less. + +**Build the tensor with the host device and dispatch accordingly.** Rejected. +`CheckSamplingLogits` (`src/vt/ops.cpp::CheckSamplingLogits`) requires `logits.device == +q.device` for every sampling op, so a host-device logits tensor needs a host +queue and the whole pipeline re-dispatched to CPU. That regresses GB10 from +on-device sampling to host sampling, and `apply_grammar_bitmask`, the random +path, `compute_prompt_logprobs` and `sample_tokens_async`'s device id buffer all +read the device off the same tensor. + +**A residency assertion in `vt::Tensor::Contiguous`.** Rejected as stated in the +issue, and this is worth recording because it looks free. Two registered +backends legitimately stamp a host-dereferenceable pointer with a non-unified +device: Tenstorrent, whose `Alloc` returns `aligned_alloc` host memory by design +(`src/vt/tenstorrent/tenstorrent_backend.cpp::Alloc`, with +`src/vt/tenstorrent/tenstorrent_backend.cpp::UnifiedMemory` returning false), and a discrete Vulkan device, whose buffers are HOST_VISIBLE and +persistently mapped (`include/vt/backend.h::DeviceMemoryIsHostAddressable`). Both register a +`kGreedyArgmax` provider. A blanket assert would false-fire on both. Making it +not false-fire needs a new per-backend residency virtual that every backend has +to get right, on the hottest constructor in the tree, for a debug-only net — +against a repair that removes the defect outright on the four sites that have it. + +## Tests + +`tests/vllm/v1/sample/test_host_buffer_staging.cpp`, using the fake-backend idiom +already established in `tests/vt/test_reference_tier.cpp` (a `Backend` over host +memory registered on the otherwise-unused `kXPU` slot, one unified instance and +one discrete instance). No GPU required. + +1. DISCRETE — the staged tensor's `data` is NOT the host pointer, and it carries + the host bytes. This is the red. +2. UNIFIED — the staged tensor's `data` IS the host pointer, byte-for-byte the + zero-copy wrap the GB10 path has today. This makes the "A3 stays green" + claim executable rather than asserted. +3. Grow-only reuse — a second smaller stage reuses the same allocation; a larger + one grows it. No per-step alloc churn on the decode path. +4. Two independent staging buffers do not alias, which is what + `collect_prompt_logprobs` overlapping the live sample-logits tensor requires. + +## Gates + +* `test_host_buffer_staging` — new, focused. +* `test_sampler`, `test_runner`, `test_sampling_metadata`, `test_logits_processors` — the seam's neighbours. +* `scripts/agent-preflight.sh`. +* NemotronH A3 (GB10) — the working path this must not break. + +## Stop conditions + +Stop and report if the unified arm is not byte-identical to the wrap it +replaces: the whole safety argument for GB10 is that the unified branch is the +same expression. + +## Owed + +Nothing. #1313 is fixed in this change. + +## Now + +ACTIVE — repair landed on `fix/sampler-host-logits-device-tag`. diff --git a/docs/USAGE.md b/docs/USAGE.md index e89cd2baf..734625d33 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -227,6 +227,20 @@ supported. a family can move -- for MiniMax-Music3 the language model, the RVQ depth decoder and the flow-matching transformer -- and there is no per-stage switch and no environment variable that turns one of them on by itself. +- `nemotron_h`, `laguna` and `qwen3_vl` finish their forward on the host and + hand the runner a host logits buffer, while the sampler itself runs on device + (`scripts/runner-routing-allowlist.txt` lists them and names what removes each + entry). On a unified-memory device — GB10 and other integrated CUDA devices, + integrated Vulkan, and CPU — those logits are sampled where they are. On a + discrete GPU they are staged into device memory once per step, into a buffer + that is reused and only ever grows, so you pay one host-to-device transfer of + `rows x vocab x 4` bytes per step on these three models and on no others. + Before [#1313](https://github.com/mudler/vllm.cpp/issues/1313) the host + address was handed to the sampling kernel directly, which is valid only on a + unified-memory device; an illegal-address abort during sampling on a discrete + GPU on an older build was this. The discrete arm is gated at the seam + (`tests/vllm/v1/sample/test_host_buffer_staging.cpp`) and has no hardware run + behind it, because every GPU on the project's fleet reports unified memory. - `tokenizer: merge token "..." at merge rank N ... is not in the vocabulary` means the tokenizer file names a merge whose left token, right token, or joined result is missing from its own vocabulary. Both `tokenizer.json` and a diff --git a/include/vllm/v1/sample/device_scratch.h b/include/vllm/v1/sample/device_scratch.h index d6fd0ee43..403436ccc 100644 --- a/include/vllm/v1/sample/device_scratch.h +++ b/include/vllm/v1/sample/device_scratch.h @@ -58,6 +58,76 @@ class DeviceScratch { vt::Tensor tensor_; }; +// Stages a HOST buffer so a DEVICE kernel may dereference it, reusing ONE +// grow-only device allocation across steps. +// +// Same residency contract as DeviceScratch above, and for the same reason — +// but for the biggest tensor in the sampling path rather than the small derived +// ones, so it must not alloc/free per step. The runner assembles the +// [rows, vocab] logits the on-device sampler runs on from a host buffer +// whenever the forward returned ForwardLogits.host (nemotron_h, laguna, +// qwen3_vl — scripts/runner-routing-allowlist.txt), and every sampling op +// dereferences that pointer ON DEVICE: vt::GreedyArgmax's CUDA arm passes it +// straight to a kernel (src/vt/cuda/cuda_sample.cu:199,207), and so do +// ApplyTemperature / ApplyTopKTopP / ComputeProbs / ComputeLogprobs. +// +// On a unified-memory backend a host address is a valid device address, so the +// wrap is in place and free — this is the GB10 path, and it is byte-for-byte +// the expression it replaces. On a DISCRETE backend it is not, and the bytes +// must be copied up first (#1313). +class HostBufferStaging { + public: + HostBufferStaging() = default; + ~HostBufferStaging() { release(); } + HostBufferStaging(const HostBufferStaging&) = delete; + HostBufferStaging& operator=(const HostBufferStaging&) = delete; + + // A [shape] view of `host`, dereferenceable by `device`. The returned Tensor + // is valid until the next Stage() on THIS object, or its destruction. + vt::Tensor Stage(vt::Device device, vt::Queue& q, const void* host, vt::DType dtype, + std::initializer_list shape) { + vt::Backend& b = vt::GetBackend(device.type); + if (b.UnifiedMemory()) { + // Host and device share one address space (CPU, GB10, integrated Vulkan): + // point straight at the host buffer. const_cast is safe on exactly the + // terms DeviceScratch states — and note the sampling ops DO mutate this + // one (temperature / top-k / top-p / the grammar bitmask), in place, + // which is what the host path has always done. This branch is the + // expression it replaces, character for character, so GB10 does not move. + return vt::Tensor::Contiguous(const_cast(host), dtype, device, shape); + } + + int64_t numel = 1; + for (int64_t s : shape) numel *= s; + const size_t bytes = static_cast(numel) * vt::SizeOf(dtype); + + // A backend swap invalidates the allocation: free it against the backend + // that made it, never against the new one. + if (backend_ != nullptr && backend_ != &b) release(); + backend_ = &b; + if (bytes > capacity_) { + if (owned_ != nullptr) backend_->Free(owned_); + capacity_ = bytes == 0 ? 1 : bytes; + owned_ = backend_->Alloc(capacity_); + } + if (bytes != 0) backend_->Copy(q, owned_, host, bytes); + return vt::Tensor::Contiguous(owned_, dtype, device, shape); + } + + // Drop the staging allocation. Only meaningful on a discrete backend. + void release() { + if (owned_ != nullptr) backend_->Free(owned_); + owned_ = nullptr; + capacity_ = 0; + backend_ = nullptr; + } + + private: + vt::Backend* backend_ = nullptr; + void* owned_ = nullptr; + size_t capacity_ = 0; +}; + } // namespace vllm::v1 #endif // VLLM_V1_SAMPLE_DEVICE_SCRATCH_H_ diff --git a/include/vllm/v1/worker/gpu/runner.h b/include/vllm/v1/worker/gpu/runner.h index 1d992a0a5..f6bfe58fc 100644 --- a/include/vllm/v1/worker/gpu/runner.h +++ b/include/vllm/v1/worker/gpu/runner.h @@ -76,6 +76,7 @@ #include "vllm/v1/attention/backends/gdn_attn.h" #include "vllm/v1/engine/types.h" #include "vllm/v1/kv_cache_interface.h" +#include "vllm/v1/sample/device_scratch.h" // HostBufferStaging (#1313) #include "vllm/v1/sample/sampler.h" #include "vllm/v1/worker/gpu/async_output.h" #include "vllm/v1/worker/gpu/input_batch.h" @@ -530,6 +531,19 @@ class GPUModelRunner final : public ModelRunnerBase { std::unique_ptr draft_model_; std::vector draft_attn_kv_; vt::Queue queue_; + // #1313 — the [rows, vocab] logits the ON-DEVICE sampler runs on are assembled + // from a HOST buffer whenever the forward returned ForwardLogits.host + // (nemotron_h, laguna, qwen3_vl). Those bytes must be reachable by the queue + // device before any sampling kernel dereferences them: free on a unified + // backend (GB10 wraps in place, unchanged), a grow-only H2D copy on a discrete + // one, where the old host address was simply illegal. + // + // TWO buffers, not one: collect_prompt_logprobs stages the prompt rows while + // the assembled sample-logits tensor is still live (sample_tokens calls it + // between assemble_sample_logits and the sampler), so a shared buffer would + // invalidate the tensor the sampler is about to read. + HostBufferStaging sample_logits_staging_; + HostBufferStaging prompt_logprobs_staging_; InputBatch input_batch_; Sampler sampler_; // ARCH-ONE-SURFACE ROW 6 (mirror of gpu/model_runner.py:368-369 diff --git a/src/vllm/v1/worker/gpu/runner.cpp b/src/vllm/v1/worker/gpu/runner.cpp index bb8041a41..697cfc460 100644 --- a/src/vllm/v1/worker/gpu/runner.cpp +++ b/src/vllm/v1/worker/gpu/runner.cpp @@ -1727,13 +1727,13 @@ vt::Tensor GPUModelRunner::assemble_sample_logits( b.Copy(queue_, sampled_logits.data(), fl.device_tensor.data, sampled_logits.size() * sizeof(float)); b.Synchronize(queue_); - logits = vt::Tensor::Contiguous( - sampled_logits.data(), vt::DType::kF32, queue_.device, + logits = sample_logits_staging_.Stage( + queue_.device, queue_, sampled_logits.data(), vt::DType::kF32, {static_cast(num_logits), vocab}); } } else if (fl.rows == num_logits) { - logits = vt::Tensor::Contiguous( - fl.host.data(), vt::DType::kF32, queue_.device, + logits = sample_logits_staging_.Stage( + queue_.device, queue_, fl.host.data(), vt::DType::kF32, {static_cast(num_logits), vocab}); } else { // (B) VT_LOGITS_GATHER=0: re-gather the logits rows from full [T,vocab] host. @@ -1747,8 +1747,8 @@ vt::Tensor GPUModelRunner::assemble_sample_logits( static_cast(row) * static_cast(vocab), static_cast(vocab) * sizeof(float)); } - logits = vt::Tensor::Contiguous( - sampled_logits.data(), vt::DType::kF32, queue_.device, + logits = sample_logits_staging_.Stage( + queue_.device, queue_, sampled_logits.data(), vt::DType::kF32, {static_cast(num_logits), vocab}); } @@ -1841,10 +1841,11 @@ void GPUModelRunner::collect_prompt_logprobs( if (fl.on_device()) { view = fl.device_tensor.Slice(0, first, first + r.num_rows); } else { - view = vt::Tensor::Contiguous( + view = prompt_logprobs_staging_.Stage( + queue_.device, queue_, fl.host.data() + static_cast(first) * static_cast(vocab), - vt::DType::kF32, queue_.device, + vt::DType::kF32, {static_cast(r.num_rows), vocab}); } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 930c6a964..ea1398999 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1334,6 +1334,7 @@ vllm_cpp_add_test(test_input_batch vllm/v1/worker/test_input_batch.cpp) vllm_cpp_add_test(test_sampling_metadata vllm/v1/sample/test_metadata.cpp) vllm_cpp_add_test(test_logits_processors vllm/v1/sample/test_logits_processors.cpp) vllm_cpp_add_test(test_sampler vllm/v1/sample/test_sampler.cpp) +vllm_cpp_add_test(test_host_buffer_staging vllm/v1/sample/test_host_buffer_staging.cpp) vllm_cpp_add_test(test_prepare_inputs vllm/v1/worker/test_prepare_inputs.cpp) vllm_cpp_add_test(test_combine_tokens vllm/v1/worker/test_combine_tokens.cpp) vllm_cpp_add_test(test_async_output vllm/v1/worker/test_async_output.cpp) diff --git a/tests/vllm/v1/sample/test_host_buffer_staging.cpp b/tests/vllm/v1/sample/test_host_buffer_staging.cpp new file mode 100644 index 000000000..dade444dc --- /dev/null +++ b/tests/vllm/v1/sample/test_host_buffer_staging.cpp @@ -0,0 +1,215 @@ +// vllm.cpp original (vt-runtime residency; vLLM has no mirror — torch owns +// residency there via .to(device)). +// +// The gate for #1313: the [rows, vocab] logits tensor the ON-DEVICE sampler runs +// on was built from a HOST std::vector pointer stamped with the queue device +// (src/vllm/v1/worker/gpu/runner.cpp, four sites). vt::GreedyArgmax does NOT copy +// its input — its CUDA arm hands the caller's pointer straight to a kernel +// (src/vt/cuda/cuda_sample.cu:199 and :207) — so on a DISCRETE GPU that is an +// illegal address. On GB10 it survives because CudaBackend::UnifiedMemory() is +// `pageable_memory_access && integrated` (src/vt/cuda/cuda_backend.cu:363), true +// there, and the driver services the host pointer through ATS. +// +// That is why this test uses FAKE backends rather than hardware: every backend +// registered on this box reports UnifiedMemory() == true, so the defect is +// invisible here by construction. The idiom — a Backend over ordinary host +// memory registered on the otherwise-unused kXPU slot, one unified instance and +// one discrete instance — is the one tests/vt/test_reference_tier.cpp already +// uses to gate the reference tier's discrete/unified split without a GPU. +// +// This file is its own executable (tests/CMakeLists.txt: one add_executable per +// test), so the kXPU registration cannot leak into another test binary. +#include "vllm/v1/sample/device_scratch.h" + +#include + +#include +#include +#include +#include + +#include "vt/backend.h" +#include "vt/dtype.h" +#include "vt/tensor.h" + +namespace { + +using vllm::v1::HostBufferStaging; +using vt::Backend; +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::Queue; + +// A Backend over ordinary host memory. `unified` is the only property under +// test; everything else is a plain allocator so a staged copy actually happens +// and can be read back. Alloc/Free are counted so the grow-only contract is +// observable rather than assumed. +class FakeBackend final : public Backend { + public: + explicit FakeBackend(bool unified) : unified_(unified) {} + void* Alloc(size_t bytes) override { + ++allocs; + return std::malloc(bytes == 0 ? 1 : bytes); + } + void Free(void* p) override { + if (p != nullptr) ++frees; + std::free(p); + } + void Memset(Queue&, void* p, int v, size_t bytes) override { std::memset(p, v, bytes); } + void Copy(Queue&, void* dst, const void* src, size_t bytes) override { + ++copies; + std::memcpy(dst, src, bytes); + } + Queue CreateQueue() override { return Queue{Device{DeviceType::kXPU, 0}, nullptr}; } + bool UnifiedMemory() const override { return unified_; } + + int allocs = 0; + int frees = 0; + int copies = 0; + + private: + bool unified_; +}; + +FakeBackend& Unified() { + static FakeBackend b(true); + return b; +} +FakeBackend& Discrete() { + static FakeBackend b(false); + return b; +} + +constexpr Device kXpu{DeviceType::kXPU, 0}; + +// A [rows, vocab] host logits buffer, exactly the shape the runner assembles. +std::vector HostLogits(int64_t rows, int64_t vocab) { + std::vector v(static_cast(rows * vocab)); + for (size_t i = 0; i < v.size(); ++i) v[i] = static_cast(i) * 0.5f - 3.0f; + return v; +} + +} // namespace + +// --------------------------------------------------------------------------- +// 1. THE DEFECT (#1313) — a DISCRETE device must never be handed a host address. +// --------------------------------------------------------------------------- +TEST_CASE("host staging: a discrete device gets DEVICE memory, not the host pointer") { + vt::RegisterBackend(DeviceType::kXPU, &Discrete()); + Queue q = Discrete().CreateQueue(); + + constexpr int64_t kRows = 3, kVocab = 64; + std::vector host = HostLogits(kRows, kVocab); + + HostBufferStaging staging; + vt::Tensor t = staging.Stage(kXpu, q, host.data(), DType::kF32, {kRows, kVocab}); + + // The whole point: what a kernel will dereference is NOT the std::vector. + CHECK(t.data != static_cast(host.data())); + // ...and it carries the bytes, so staging is a copy and not just a different + // address. A guard that only checked the pointer would pass on an empty one. + REQUIRE(t.data != nullptr); + const auto* staged = static_cast(t.data); + for (size_t i = 0; i < host.size(); ++i) { + CHECK(staged[i] == doctest::Approx(host[i])); + } + // The tensor still reports the device the sampler dispatches on. + CHECK(t.device.type == DeviceType::kXPU); + CHECK(t.rank == 2); + CHECK(t.shape[0] == kRows); + CHECK(t.shape[1] == kVocab); + CHECK(t.IsContiguous()); +} + +// --------------------------------------------------------------------------- +// 2. THE WORKING PATH — a UNIFIED device keeps the zero-copy wrap byte-for-byte. +// --------------------------------------------------------------------------- +// This is the executable form of "the NemotronH A3 gate stays 96/96". GB10 is a +// unified backend, so the repair must not move it off the in-place wrap it has +// today: no copy, no allocation, the same address. +TEST_CASE("host staging: a unified device wraps IN PLACE, with no copy and no alloc") { + vt::RegisterBackend(DeviceType::kXPU, &Unified()); + Queue q = Unified().CreateQueue(); + + constexpr int64_t kRows = 3, kVocab = 64; + std::vector host = HostLogits(kRows, kVocab); + + const int allocs_before = Unified().allocs; + const int copies_before = Unified().copies; + + HostBufferStaging staging; + vt::Tensor t = staging.Stage(kXpu, q, host.data(), DType::kF32, {kRows, kVocab}); + + CHECK(t.data == static_cast(host.data())); + CHECK(Unified().allocs == allocs_before); + CHECK(Unified().copies == copies_before); + + // In-place means the sampler's mutation is visible on the host buffer, which + // is what apply_temperature / the grammar bitmask have always done here. + static_cast(t.data)[0] = 42.0f; + CHECK(host[0] == doctest::Approx(42.0f)); +} + +// --------------------------------------------------------------------------- +// 3. GROW-ONLY — no per-step allocation on the decode path. +// --------------------------------------------------------------------------- +TEST_CASE("host staging: the discrete allocation is reused across steps and grows once") { + vt::RegisterBackend(DeviceType::kXPU, &Discrete()); + Queue q = Discrete().CreateQueue(); + + std::vector big = HostLogits(4, 64); + std::vector small = HostLogits(1, 64); + + HostBufferStaging staging; + const int allocs_before = Discrete().allocs; + + vt::Tensor a = staging.Stage(kXpu, q, big.data(), DType::kF32, {4, 64}); + const int after_first = Discrete().allocs; + CHECK(after_first == allocs_before + 1); + void* first_ptr = a.data; + + // A SMALLER step must not allocate again — this is the per-token decode step. + vt::Tensor b = staging.Stage(kXpu, q, small.data(), DType::kF32, {1, 64}); + CHECK(Discrete().allocs == after_first); + CHECK(b.data == first_ptr); + const auto* staged = static_cast(b.data); + for (size_t i = 0; i < small.size(); ++i) CHECK(staged[i] == doctest::Approx(small[i])); + + // A LARGER one grows exactly once, and frees the old block rather than leaking. + const int frees_before = Discrete().frees; + std::vector bigger = HostLogits(8, 64); + vt::Tensor c = staging.Stage(kXpu, q, bigger.data(), DType::kF32, {8, 64}); + CHECK(Discrete().allocs == after_first + 1); + CHECK(Discrete().frees == frees_before + 1); + const auto* staged_c = static_cast(c.data); + for (size_t i = 0; i < bigger.size(); ++i) CHECK(staged_c[i] == doctest::Approx(bigger[i])); +} + +// --------------------------------------------------------------------------- +// 4. TWO BUFFERS DO NOT ALIAS. +// --------------------------------------------------------------------------- +// collect_prompt_logprobs stages the prompt rows while the assembled sample +// logits are still live (runner.cpp:2083-2090), so the runner holds two staging +// objects. One shared buffer would invalidate the tensor the sampler is about to +// read — this pins that they are independent. +TEST_CASE("host staging: two staging buffers are independent") { + vt::RegisterBackend(DeviceType::kXPU, &Discrete()); + Queue q = Discrete().CreateQueue(); + + std::vector sample_rows = HostLogits(2, 32); + std::vector prompt_rows(static_cast(2 * 32), 7.0f); + + HostBufferStaging sample_staging; + HostBufferStaging prompt_staging; + + vt::Tensor s = sample_staging.Stage(kXpu, q, sample_rows.data(), DType::kF32, {2, 32}); + vt::Tensor p = prompt_staging.Stage(kXpu, q, prompt_rows.data(), DType::kF32, {2, 32}); + + CHECK(s.data != p.data); + // Staging the prompt rows must not have disturbed the sample rows. + const auto* sd = static_cast(s.data); + for (size_t i = 0; i < sample_rows.size(); ++i) CHECK(sd[i] == doctest::Approx(sample_rows[i])); + const auto* pd = static_cast(p.data); + for (size_t i = 0; i < prompt_rows.size(); ++i) CHECK(pd[i] == doctest::Approx(7.0f)); +}