fix(SAMPLER-HOST-LOGITS-DEVICE-TAG): the on-device sampler was handed a host address, and the CUDA arm dereferences it

.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<float>` 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<float>()` 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 |

.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<float>` 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<float>(), v)`
+* `src/vt/cuda/cuda_sample.cu::ArgmaxPartialKernel` — `ArgmaxPartialKernel<<<...>>>(..., logits.Ptr<float>(), 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`.