diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md
index 1260fde25..828131428 100644
--- a/.agents/benchmark-record.md
+++ b/.agents/benchmark-record.md
@@ -19,6 +19,71 @@ from relative link targets repointed for this file's location.
 
 # Benchmarks
 
+## A2-Q2b — NemotronH host re-expansion, attributed per tensor group, per decode token (2026-08-19, `row/A2-Q2b-lmhead-nvfp4`, #810)
+
+**Why this was measured at all.** The row was dispatched on the claim that
+NemotronH decode re-expands ~1.24e9 elements (~2.49 GB) per token on the host,
+and that `lm_head` (guessed at 131072 x 4096 = 537e6 elements, ~43%) is a large
+share of it. That was arithmetic, not measurement, and the brief said so. This
+entry replaces it with a count.
+
+**Instrument.** A shape-keyed tally at `NemotronHOwned::DenseBf16`
+(`nemotron_h.cpp:395`), the SINGLE seam every host dequant in this architecture
+passes through, dumped per forward call from `ForwardNemotronHForCausalLM`.
+Shape is an unambiguous key on this checkpoint, so no call-site labelling was
+needed. Driven by `examples/nemotron_h_gen` against the real 21 GB checkpoint
+(revision `29f2d1746d8f41e316523194b19018707749b1b1`) on `mudler-ubuntu-box`,
+CPU host arm, engine load 302.4 s, peak RSS 19.2 GiB. The instrument was scratch
+and was reverted; the run logs are the evidence.
+
+**The decode step (T=1, `top_k` 6, 23 MoE layers, 369 dequant calls):**
+
+| group | shape | calls | elements | per call | % |
+|---|---|---|---|---|---|
+| routed expert `up_proj` | `[1856, 2688]` | 138 | 688 472 064 | 4 988 928 | 22.36% |
+| routed expert `down_proj` | `[2688, 1856]` | 138 | 688 472 064 | 4 988 928 | 22.36% |
+| shared expert `down_proj` | `[2688, 3712]` | 23 | 229 490 688 | 9 977 856 | 7.45% |
+| shared expert `up_proj` | `[3712, 2688]` | 23 | 229 490 688 | 9 977 856 | 7.45% |
+| `lm_head` | `[131072, 2688]` | 1 | 352 321 536 | 352 321 536 | 11.44% |
+| mamba `out_proj` (FP8) | `[2688, 4096]` | 23 | 253 231 104 | 11 010 048 | 8.23% |
+| mamba `in_proj` (FP8) | `[10304, 2688]` | 23 | 637 034 496 | 27 697 152 | 20.69% |
+| **TOTAL** | | **369** | **3 078 512 640** | | **100%** |
+
+`138 == 6 x 23` exactly, which is what identifies this as the decode shape
+rather than a prefill aggregate. The T=5 prefill of the same run reported 1115
+calls / 6 800 252 928 elements, and is recorded here only so the two are not
+confused: the per-token figure is the T=1 one.
+
+**Three findings.**
+
+1. The dispatching estimate was wrong in both numbers, in the same direction.
+   `hidden_size` is 2688, not 4096, so `lm_head` is 352 321 536 elements, and
+   its share of the population the brief meant is 28.35%, not 43%.
+2. **The 1.24e9 figure names a REGIME, not a total.** It is not the host arm's
+   3 078 512 640. It is `mamba + lm_head` = 1 242 587 136 = 2.485 GB, the
+   residue after A2-Q2a moved the MoE arm to the device, matching to four
+   significant figures. Anyone quoting 1.24e9 is quoting the post-A2-Q2a state.
+3. **`lm_head` is the LAST one.** Against the A2-D1 three-leg discriminator on
+   `dgx:gpu0` (`/workspace/a2d1-discriminate/20260819T200231Z`: device mamba ON
+   1.554 s/token, 108.2x vs vLLM, GPU busy 10.18%; OFF 10.319 s/token, 718.1x,
+   GPU busy 7.86%), the mamba arm is worth 6.64x on its own. `lm_head` is on the
+   host in all three legs. Once the mamba arm lands, `lm_head` is 352 321 536 of
+   352 321 536, i.e. 100% of the host re-expansion left in a decode step.
+
+**Also load-bearing, and not visible in the percentages:** `lm_head` is the
+largest SINGLE re-expansion in the model by 12.7x (352.3e6 elements in ONE call
+against 27.7e6 for mamba `in_proj`), so it allocates a 704.6 MB transient bf16
+buffer once per step. On a unified-memory box where `gpu_memory_utilization`
+does not bound host RAM and the kernel reboots rather than OOM-kills, that
+single transient is the one that matters.
+
+**What this does NOT measure.** Time. The element counts are portable; the
+seconds are not, and the box that produced these counts is not a gate host. No
+speed claim is made from this entry, and the device `lm_head` arm's own numeric
+gate is PENDING a `dgx:gpu0` window.
+
+---
+
 ## ENG-EXPERT-STREAM-DEVICE W0a — a GB10 kernel CAN dereference the host slot arena (2026-08-19, `row/ENG-EXPERT-STREAM-DEVICE-W0`, #1124)
 
 **W0a is the probe the whole W0 mechanism rests on, and it answered
diff --git a/.agents/issue-index.md b/.agents/issue-index.md
index 13cd6cdb8..2f11019cd 100644
--- a/.agents/issue-index.md
+++ b/.agents/issue-index.md
@@ -474,6 +474,8 @@ 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 |
+| [#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 |
+| [#1421](https://github.com/mudler/vllm.cpp/issues/1421) | `ENG-RECORD-ANCHOR-RATCHET` | Two NVFP4 W4A16 upstream anchors are WRONG AT THE PIN `5559679229bc961848b121ccdeaa8fa5d79bec98`, verified against a local checkout of `vllm-project/vllm` at that exact commit (`origin` is upstream, not a fork). The quoted `elif linear_backend == "auto" and use_a16: force_kernel = MarlinNvFp4LinearKernel` is at `kernels/linear/__init__.py:922-924`, not `:879-881`; `:879-881` is the tail of `init_wfp8_a16_linear_kernel` (`:848`), an FP8 helper sharing the `a16` vocabulary, and `init_nvfp4_linear_kernel` does not start until `:883`. And `packed_modules_mapping` is at `models/qwen3.py:275-278`, not `:271-274`, which is the `class Qwen3ForCausalLM(...)` line. Both are the worst shape of anchor rot: the cited line EXISTS and holds plausible neighbouring code, so a reader who checks the anchor is confirmed in a wrong belief. Nine code and test sites over SEVEN files (`dense_nvfp4_gemm.h:18,627,783`, `nvfp4.h:10`, `gguf_keep_quant.h:201`, `dense_weight_loaders.h:674`, `minimax_h3_device.cpp:153`, `test_qwen3_32b_nvfp4a16_paged_engine.cpp:63`, `test_qwen3_32b_nvfp4a16_load.cpp:19` -- count the parenthetical, the earlier `six files` did not) plus EIGHT record sites over FIVE record files, `parity-ledger.md:599,915,916` carrying three of them and `quantization-matrix.md:124`, `specs/sweep-qwen3-32b-nvfp4a16.md:86,465`, `specs/minimax-h3.md:332` and `completed/state-events/2026-08/STATE-20260806T141900-001.md:14` the rest. PRE-EXISTING and NOT fixed in flow: it is not small or confined, it spans paths several rows own, it would turn `parity-ledger.md` into a lock, and it needs the whole citation set RE-DERIVED at the pin rather than a sed -- range-checking cannot find it, because every one of these is in range. Found while repairing fresh-review findings on [#1415](https://github.com/mudler/vllm.cpp/pull/1415) (A2-Q2b, [#810](https://github.com/mudler/vllm.cpp/issues/810)), which relocates the `:783` citation from `:725` and deliberately leaves its VALUE unchanged so this issue owns every site uniformly instead of half of them looking reviewed. Owned by row `ENG-RECORD-ANCHOR-RATCHET`, the anchor-rot ratchet row | bug |
 | [#1327](https://github.com/mudler/vllm.cpp/issues/1327) | `SPEC-DFLASH2` | `.agents/specs/dflash2-spec-decode.md` `## Upstream chain` said the three output scalars `input_embedding_scale`, `output_multiplier` and `final_logit_softcapping` are "ABSENT from this config" and that "no published checkpoint exercises them, so the port implements them and gates them synthetically". That was measured on `z-lab/Qwen3.8-27B-DFlash2` alone. `z-lab/Muse-Glimmer-30B-DFlash2` — the SECOND published DFlash2 checkpoint, `config.json` sha256 `cb684d6f688a22619a63ea1debe7d30c139c195bf3141fd86a763763ab34b5d9`, read 2026-08-19 — sets `output_multiplier` to `0.19611613513818404` and `final_logit_softcapping` to `20.0`, and ships `block_size` 16 against the 27B's 8, hidden 6656 (416 groups, a 1664-wide `kernel_projection`) and `rope_theta` 500000.0. Both scalars are applied to the candidate VALUES in `compute_candidates` BEFORE the selector scores them, so a wrong value reorders the top-K and moves acceptance without raising — the `is_causal` failure class one layer up, and the class no token gate here can see. A port reading all three with `.get(key, default)` would pass every gate built from the 27B draft and be measuring the default path. The same file also falsifies `## Scope`'s exclusion of "a second DFlash2 target family": upstream registers ONE architecture and both checkpoints declare `model_type` `qwen3`, so what the second adds is values rather than a class. FIXED IN FLOW by `SPEC-DFLASH2` W2, which is the wave that had to read both configs anyway: `## Scope` drops the exclusion, `## Upstream chain` records both values with their source, `## Gates` G1 now requires BOTH published block shapes (upstream's own reference test parametrises 5 and 8 and never reaches 16), and `## Risks/decisions` D9 binds W3 to gate the scalars against the checkpoint that sets them | gap |
 | [#1413](https://github.com/mudler/vllm.cpp/issues/1413) | `LTX25-DEVICE-RESIDENCY` | The LTX-2.5 render phase table [#1408](https://github.com/mudler/vllm.cpp/pull/1408) lands for [#1010](https://github.com/mudler/vllm.cpp/issues/1010) is written by the SUCCESS PATH ONLY — `WritePhaseLog` sits after `im.trace.completed = true` at `src/vllm/multimodal/ltx2_video.cpp:4655-4658 @ 4f3c24380` — so a render that is killed, aborted by a lease governor, or still running writes no table at all. That is the population the campaign actually has: [#1375](https://github.com/mudler/vllm.cpp/issues/1375) is `ABORT[92] PROJECTED OVERRUN`, `child exit=-15`, 0 frames; [`ltx25-decode-speed.md`](specs/ltx25-decode-speed.md) rung 1 is `EXIT=137`, 0 frames; rung 2 is `EXIT=1`, 0 frames. **And nothing is emitted WHILE a render runs**: `PhaseLog::Open`/`Close` print nothing and `VLLM_RENDER_PHASE_LOG_STDERR` fires inside `WriteJson`, i.e. on the success path again, so between `ltx2-gen: family=...` and `wrote N frames` a 2.5-hour render is silent and working is byte-identical to hung. **The unit that costs the wall has no counter**: `denoise` is one leaf covering ~2.7 h, against #1375's measured ~162 s per DiT forward at 60 structural forwards (30 steps x 2 CFG legs — `cfg_scale != 1.0` forces the unconditional branch at `ltx2_pipeline.cpp:521-523`). External sampling is NOT the fallback and was tried: #1375 records `phase=OTHER` throughout because `eu-stack` unwinds ZERO frames inside the `rc` worker container, and its own text names an in-process phase marker as the way to attribute the 162 s. Owed: a stderr line on every phase open and close, so the LAST LINE PRINTED names the phase in flight, plus one tick per DiT forward carrying phase, step `k/N`, cumulative forward index, elapsed and seconds-since-previous — on the shipped default, because the failure happened on default settings and `VT_H3_PROGRESS` (`minimax_h3.cpp:776-793`) is the opt-in shape that is exactly why no LTX-2.5 run has one. Precondition for attributing #1375, which caps render resolution. Stage W0-live in [`ltx25-device-residency.md`](specs/ltx25-device-residency.md) | feature |
 | [#1320](https://github.com/mudler/vllm.cpp/issues/1320) | `ENG-EXPERT-STREAM-DEVICE` | `VT_CPU_QUANT_REPACK` rewrites a Q8_0 weight into the `block_q8_0x4` i8mm interleave at load, and only the CPU `MatmulBTKernel` understands that layout — the CUDA quant dot reads plain `block_q8_0` and never consults `Tensor::repacked`, so a repacked weight uploaded to a device produces WRONG TOKENS rather than a crash. Its sibling `elem_kn_repack` has BOTH a CPU-platform gate in the loader policy (`gguf_keep_quant.cpp`, `CurrentPlatform().device_type() == kCPU`) and a refusal in `ResidentWeight`; `quant_repack` had NEITHER, because it rides `vt::cpu::QuantRepackActive()`, a HOST-CPU Arm i8mm probe that says nothing about where the weight will execute. An aarch64 box doing `--device cuda` — which is exactly `dgx:gpu0`, a Grace-Blackwell GB10 — satisfies it. Currently silent rather than firing, and that is measured: eligibility is Q8_0 with `n % 4 == 0 && k % 32 == 0`, `Qwen3.8-2.4T-A95B UD-Q1_0` has exactly one Q8_0 tensor at 0.01% of parameters, and the instrumented load of that checkpoint recorded `quant_repack = 0`. Found while implementing W0f ([#1299](https://github.com/mudler/vllm.cpp/issues/1299)), which changes the same function and whose G0-CORRECT gate compares CUDA tokens against CPU tokens — precisely the instrument that would have reported this as a W0f defect. FIXED IN FLOW as the tripwire its sibling already has, placed beside it so it covers the staging branch AND W0f's aliasing branch (where the bytes live was never the question; the kernel that reads them is), gated red-first on both arms and mutation-proven. `VT_CPU_QUANT_REPACK=0` is the operator's way past it. NOT fixed here, deliberately: adding the missing CPU-platform gate to `p.quant_repack` itself is a loader-policy semantics change owned by `QUANT-GGUF-KEEPQ-LOADER`, carried under `## Owed` in [`expert-stream-device-slots.md`](specs/expert-stream-device-slots.md) | bug |
diff --git a/.agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md b/.agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md
index 4344d63b7..71cd32889 100644
--- a/.agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md
+++ b/.agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md
@@ -5,7 +5,8 @@
 ([#517](https://github.com/mudler/vllm.cpp/issues/517)).
 **Sibling / predecessor:** [`nemotron-h-a2q2-nvfp4-moe-lmhead.md`](nemotron-h-a2q2-nvfp4-moe-lmhead.md)
 (A2-Q2a — the MoE arm, landed on the synthetic gate).
-**State:** spec only. No product code.
+**State:** IMPLEMENTED on `row/A2-Q2b-lmhead-nvfp4`; the GPU legs of §3 are pending a
+`dgx:gpu0` window. See `## Now`.
 
 ---
 
@@ -36,7 +37,7 @@ bounded (A2-Q2a §13.6.1) and cannot stand in for what this unit measures.
 | In A2-Q2b | Out |
 |---|---|
 | the per-block numeric gate on the REAL checkpoint: every one of the 23 MoE layers against `trace.mixer[l]` | anything A2-P owns (paging, carried state, batching, the G-SAFE narrowing) |
-| `lm_head` through the NVFP4 dense route, with the §4.3 residency decision applied as A2-Q2a applied it | the FP8 mamba arm — A2-Q1 |
+| `lm_head` through the NVFP4 dense route, with the `## 5. Owed` residency decision (#984) applied as A2-Q2a applied it | the FP8 mamba arm — A2-Q1 |
 | hybrid-vs-host token identity, and the disclosure that A2-R's attributability property ENDS when `lm_head` moves | fixing [#984](https://github.com/mudler/vllm.cpp/issues/984) or [#962](https://github.com/mudler/vllm.cpp/issues/962) |
 | the §5.3 mutations A2-Q2a left owed (Q2-M3 … Q2-M7) | any throughput number, on any axis |
 
@@ -100,13 +101,26 @@ the box reboots rather than OOM-killing a process).
 
 So:
 
-1. Take `$GPU_LOCK` with a **blocking** `flock` and wait. Never race.
-2. **Check `free -g` headroom INSIDE the locked region**, not before acquiring
-   it. A blocking flock says the previous holder released; it never says the box
-   recovered. Abort loudly below a stated floor (A2-Q2a used 60 GB).
+1. **Claim `dgx:gpu0` through the fleet controller, and run the work as the
+   lease's payload**: `rc run -d dgx:gpu0 --max-runtime <N>h -- <cmd>`, with
+   `scripts/nemotron-h-a2q2b-gpu-gate.sh` as `<cmd>`. Never `ssh` to the box,
+   and never `rc hold` an idle one. **Do not take `$GPU_LOCK` here.** `dgx:gpu0`
+   is a fleet device, so `AGENTS.md` makes the lease the required path, and the
+   file mutex is the wrong instrument for it: the fleet cannot see that mutex,
+   so a `flock` taken over `ssh` does not exclude a concurrent `rc` holder and
+   the controller keeps reporting the box free while somebody is on it. That is
+   not hypothetical — on 2026-08-17 exactly this pair of non-excluding mutexes
+   cost `.agents/specs/minimax-music3.md` §13.10 a whole speed axis, which is
+   the [#777](https://github.com/mudler/vllm.cpp/issues/777) failure again. The
+   gate script correspondingly takes **no** mutex of its own; the lease is what
+   serialises this row's window.
+2. **Check `free -g` headroom INSIDE the lease**, not before claiming it. A
+   granted lease says the previous holder released; it never says the box
+   recovered. Abort loudly below a stated floor (A2-Q2a used 60 GB), which is
+   what `scripts/nemotron-h-a2q2b-gpu-gate.sh` PRECONDITION 1 does.
 3. **Sample `free -g` on a loop for the whole load and record the PEAK**, not the
    final value. The figure that matters is transient.
-4. Build `-j 4`. One log per run. Never hang holding the lock.
+4. Build `-j 4`. One log per run. Never hang holding the lease.
 5. Verify the configure log reads `ENABLED for [121a]` — a `DISABLED` line or a
    `[121]` **voids** the result rather than failing it.
 
@@ -138,6 +152,453 @@ and an 8h19m outage ended by a human power cycle. Size the plan for that.
   `b_q_weight` and nothing at all about `b_scales`**, so a stride defect is
   silent at the op boundary and only the numeric gate can see it.
 
+
+- [#1410](https://github.com/mudler/vllm.cpp/issues/1410) — `check-runner-routing-consistency.py`
+  cannot resolve a cross-TU free-function device forward, so NemotronH still
+  classifies HOST although `NemotronHPagedForward` now assigns both
+  `fl.device_tensor` and `fl.device_storage`. The allowlist entry is NARROWED
+  to that instrument limit rather than removed. Not fixed in this flow because
+  it changes checker semantics, which `AGENTS.md` routes to its own row.
+- **The 23-layer MoE per-block sweep of §3 and the `lm_head` real-checkpoint
+  numeric leg** run only on `dgx:gpu0` and are PENDING a window, not waived.
+  What this change lands is the implementation, the CPU-buildable arms, and the
+  synthetic device gate that A2-Q2a's preamble argues for. `## Now` records the
+  exact state of each leg, so a `PENDING` here is a scheduled measurement with
+  a named blocker rather than a gate nobody ran.
+- **Q2-M3 … Q2-M7** (A2-Q2a's owed mutations) are unchanged by this row and
+  stay owed. They gate the MoE arm, not `lm_head`.
+
+---
+
 ## 6. Now
 
-Claimable once A2-Q2a lands. Its blocker is a GPU window, not a design question.
+**Implemented.** The device `lm_head` arm lands on `row/A2-Q2b-lmhead-nvfp4`.
+
+### What was measured BEFORE anything was built
+
+The row's premise — that host re-expansion dominates NemotronH decode and that
+`lm_head` is a large share of it — was ARITHMETIC when this row was dispatched.
+It is now a measurement, taken at the single dequant seam
+(`NemotronHOwned::DenseBf16`) on the REAL 21 GB checkpoint through the
+production ABI driver (`examples/nemotron_h_gen`), one decode step, T=1,
+top_k=6, 23 MoE layers:
+
+| group | shape | calls | elements | per call | % |
+|---|---|---|---|---|---|
+| routed expert `up_proj` | `[1856, 2688]` | 138 | 688 472 064 | 4 988 928 | 22.36% |
+| routed expert `down_proj` | `[2688, 1856]` | 138 | 688 472 064 | 4 988 928 | 22.36% |
+| shared expert `down_proj` | `[2688, 3712]` | 23 | 229 490 688 | 9 977 856 | 7.45% |
+| shared expert `up_proj` | `[3712, 2688]` | 23 | 229 490 688 | 9 977 856 | 7.45% |
+| **`lm_head`** | **`[131072, 2688]`** | **1** | **352 321 536** | **352 321 536** | **11.44%** |
+| mamba `out_proj` (FP8) | `[2688, 4096]` | 23 | 253 231 104 | 11 010 048 | 8.23% |
+| mamba `in_proj` (FP8) | `[10304, 2688]` | 23 | 637 034 496 | 27 697 152 | 20.69% |
+| **TOTAL** | | **369** | **3 078 512 640** | | **100%** |
+
+`138 == 6 * 23` exactly, which is what confirms this is the decode shape and not
+a prefill aggregate.
+
+Three things follow, and the third is why the row proceeded:
+
+1. **The dispatching estimate was wrong in both of its numbers, in the same
+   direction.** It put `lm_head` at 131072 x 4096 = 537e6 elements and at ~43%
+   of the population. `hidden_size` is 2688, not 4096: the true count is
+   352 321 536, and the true share of that population is 28.35%.
+2. **The "~1.24e9 elements / ~2.49 GB per token" figure is real and now has a
+   name.** It is not the host arm's total (3.079e9 / 6.157 GB). It is exactly
+   `mamba + lm_head` = 1 242 587 136 elements = 2.485 GB — the residue AFTER
+   A2-Q2a moved the MoE arm to the device. It matches to four significant
+   figures, which is what identifies which regime the number describes.
+3. **`lm_head` is the LAST one.** Against the three-leg discriminator on
+   `dgx:gpu0` (`/workspace/a2d1-discriminate/20260819T200231Z`: device mamba ON
+   1.554 s/token and 108.2x vs vLLM, OFF 10.319 s/token and 718.1x), the mamba
+   arm is worth 6.64x and is in flight on A2-D1. `lm_head` is on the HOST in
+   every one of those three legs. Once the mamba arm lands, `lm_head` is
+   352 321 536 of 352 321 536 — 100% of the host re-expansion left in a decode
+   step. It is also the largest SINGLE re-expansion in the model by 12.7x
+   (352.3e6 in one call against 27.7e6 for mamba `in_proj`), so its 704.6 MB
+   transient bf16 buffer is the allocation that matters most on a
+   unified-memory box that reboots rather than OOM-kills.
+
+The refutation is therefore narrow and the conclusion survives: the estimate's
+share was wrong, the direction was right, and the case is STRONGER after the
+discriminator than before it.
+
+### Leg status
+
+| Leg | State |
+|---|---|
+| seam extension (caller-owned `MarlinDenseResident`) | DONE, built on BOTH arms (see below) |
+| device `lm_head` arm + `DeviceLmHeadEligible` | DONE, built on BOTH arms |
+| production wiring in `NemotronHPagedForward` -> device `ForwardLogits` | DONE |
+| host arm retained as the gate's operand + the non-NVFP4 fallback | DONE |
+| routing-allowlist entry narrowed, [#1410](https://github.com/mudler/vllm.cpp/issues/1410) filed | DONE |
+| the CPU-reachable half (`n_out`, the shared `final_normed` download) gated through `GPUModelRunner` | DONE and RUN — `test_nemotron_h_paged_forward.cpp` §12, red-first below |
+| synthetic device `lm_head` numeric gate (measured band, asserted counts) | COMPILES; runs on CUDA only. **NEVER RUN** |
+| `nvcc` build of the Marlin KERNEL + any execution of the device arm | PENDING a `dgx:gpu0` window |
+| real-checkpoint `lm_head` numeric leg + token identity | PENDING a `dgx:gpu0` window |
+| reachability deletion mutation | PENDING the same window |
+| 23-layer MoE per-block sweep (§3) | PENDING; owed above |
+
+### What "built" means here, corrected
+
+The first submission of this row said "the CPU build compiles the `#else` arms
+only, so it does not compile the Marlin path at all", and recorded the CUDA
+build of the Marlin arm as blocked on a GPU window. **That was wrong, and it
+overstated the blocker.** `include/vt/cuda/marlin_repack.h` includes only
+`<cstdint>`, `<cstddef>` and `<vector>`, so the HOST side of the Marlin arm
+needs no CUDA toolkit at all. Measured on this box, which has no `nvcc`:
+
+```
+c++ -std=c++20 -I include -I src -isystem third_party -DVT_MARLIN_NVFP4=1     -Wall -Wextra -Werror -c -o nhd_marlin.o     src/vllm/model_executor/models/nemotron_h_device.cpp
+-> rc 0, 0 errors, 0 warnings, a 1 268 552-byte object
+```
+
+**That byte figure is anchored to `72d736867`**, gcc 13.3.0 (Ubuntu
+13.3.0-6ubuntu2~24.04.1) on a box with no `nvcc`, and compilation is
+deterministic here: two runs of the identical command produced the identical
+`sha256 42d670b6...`. The anchor is the point. Measured with that one command
+at seven commits of this branch, the size MOVES with the code:
+
+| commit | object bytes |
+|---|---|
+| `806b263e7`, `8fa900a62`, `1c62d9974`, `29b1128e3` (the reviewed head) | 1 269 696 |
+| `fedf78d86` (the `-Werror` declaration repair) | 1 272 808 |
+| `bff2b7b2f` (the F4 predicate repair) onwards, through `18f3da188` and `72d736867` | 1 268 552 |
+
+So the 1 269 696 this section carried until now was CORRECT when it was written
+and correct at the head a fresh reviewer read; it went stale two commits later
+and nothing said so, because the evidence block named no commit. An unanchored
+byte count is a measurement nobody can reproduce or falsify, which is why the
+number now travels with the SHA it was taken at. Nothing after `72d736867`
+touches this translation unit or its headers.
+
+The size is incidental either way; what the block asserts is `rc 0` with zero
+errors and zero warnings on a toolchain with no CUDA, and that holds at every
+one of those seven commits.
+
+Both arms are therefore compiled, and both are compiled `-Werror`. What
+genuinely needs `nvcc` is the Marlin **kernel** and every **execution** of the
+device path. Landing without those is acceptable and is what the PENDING rows
+above record; claiming the host arm could not be compiled was not.
+
+### A gate this row was already failing, and cannot repair in place
+
+`scripts/check-doc-checkpoint.py` is red on this branch, and it was red at the
+reviewed head `29b1128e` with the identical two errors, measured with
+`--base 96ed8346f --head 29b1128e3`:
+
+```
+ERROR: commit 1c62d9974: changed user_usage but did not update docs/USAGE.md
+ERROR: commit 8fa900a62: changed .agents/benchmark-record.md: measurement
+       recorded but did not update docs/STATUS.md
+```
+
+The PR description claimed this gate had one real failure and that it was
+repaired. It was not, and `.github/workflows/ci.yml:519` runs the same
+`--base/--head` invocation, so the lane is red for this reason independently of
+[#1371](https://github.com/mudler/vllm.cpp/issues/1371).
+
+The second error names a real gap and it is now closed: this row moved a
+lifecycle state and recorded a measurement, and `docs/STATUS.md` said neither.
+It does now.
+
+The first error, and the historical form of the second, cannot be closed by a
+later commit. `check_doc_checkpoint` iterates `commits_in_range` and judges each
+commit on its own contents, so the obligation belongs to `1c62d9974` and
+`8fa900a62` and nothing appended afterwards discharges it. Repairing those two
+would mean rewriting commits that are already the reviewed base, which resets
+the pull request's continuous-integration approval and moves the head a fresh
+reviewer was asked to look at. That is a scheduling decision rather than a
+technical one, so it is recorded here as owed and raised for the operator rather
+than taken unilaterally by a repair pass.
+
+### The gate, as it actually ran
+
+Every number below is from a run on this box, `-DCMAKE_BUILD_TYPE=RelWithDebInfo`,
+`-Wall -Wextra -Werror`, no CUDA toolkit.
+
+**Read the `overlay` column before any other one.** This tree cannot run
+`GPUModelRunner` at all — [#1371](https://github.com/mudler/vllm.cpp/issues/1371)
+throws at its construction — so every green below was taken with
+[#1392](https://github.com/mudler/vllm.cpp/pull/1392)'s production fix applied
+to the working tree, never committed here and reverted byte-for-byte afterwards.
+Without it the §12 case does not merely score differently, it never reaches its
+first assertion. Saying "a run on this tree" without that clause made the §12
+row false of the tree it named.
+
+| binary | overlay | `run_rc` | cases | assertions | verdict |
+|---|---|---|---|---|---|
+| `test_nemotron_h_moe_device` | none | 0 | 4, **4 passed** | 4 | `SUCCESS!` — and **all four SKIP**, both A2-Q2b cases included |
+| `test_nemotron_h_paged_forward`, whole binary | **none** | 1 | 13, 2 passed, **11 failed** | 18 | `FAILURE!` — [#1371](https://github.com/mudler/vllm.cpp/issues/1371), not this row |
+| `test_nemotron_h_paged_forward`, whole binary | **#1392** | 0 | 13, **13 passed** | 3269 | `SUCCESS!` |
+| the new §12 case alone | **none** | **1** | 1, **0 passed, 1 failed**, 12 skipped | **0** | **`FAILURE!` — throws #1371 before the first assertion** |
+| the new §12 case alone | **#1392** | 0 | 1, 1 passed, 12 skipped | 13 | `SUCCESS!` |
+
+The un-overlaid §12 row is the baseline the two mutations below are measured
+against, and it is a red already. That is exactly why they are reported with the
+overlay: a mutation cannot be shown to turn a case red when the case is red
+without it. Re-derived at `72d736867`: `-tc=` the §12 case on the clean tree
+gives `run_rc=1`, `test cases: 1 | 0 passed | 1 failed | 12 skipped`,
+`assertions: 0`, and `ERROR: test case THREW exception: No valid attention
+backend for device type 0 from {FLASH_ATTN: [head_size not supported]}`. Note
+`assertions: 0` — the shape this project has read as a pass before.
+
+**The moe_device row is the honest reading of the synthetic gate, and it is not
+a pass.** The binary builds and exits 0, but all four cases take the
+`TryCudaQueue` skip on a GPU-less box, so the 4 assertions are the skip notices
+themselves. The numeric gate examined nothing. That is why the table above says
+`NEVER RUN` and not `green`.
+
+**The paged-forward red is [#1371](https://github.com/mudler/vllm.cpp/issues/1371)
+and belongs to nobody here.** All 11 failing cases throw the identical
+`No valid attention backend for device type 0 from {FLASH_ATTN: [head_size not
+supported]}` at `GPUModelRunner` construction, 10 of them cases this row never
+touched. Overlaying [#1392](https://github.com/mudler/vllm.cpp/pull/1392)'s
+production fix in the working tree — never committed here, and reverted
+afterwards — turns the same binary green at 13/13 and 3269 assertions. Note the
+shape of the red: doctest reported `assertions: 18 | 18 passed | 0 failed` while
+11 cases were throwing, so the assertion line alone would have read as a pass.
+
+### The red-first, on the cases that can run
+
+Each mutation was applied to a scratch copy of
+`src/vllm/model_executor/models/nemotron_h_device.cpp`, verified to have applied
+by `git diff --stat`, built (a mutation that fails to build proves nothing), run,
+and restored to an identical sha256 (`abf6e21f...`).
+
+**Every row here was run with [#1392](https://github.com/mudler/vllm.cpp/pull/1392)'s
+production fix overlaid**, for the reason the table above gives: un-overlaid, the
+§12 case is already `run_rc=1` on the unmutated tree, so an M1 or M2 red measured
+there would prove nothing at all. The overlay is what makes the vehicle able to
+report a green, and only then can a mutation take it away.
+
+| mutation | overlay | applied | `compile_rc` | `run_rc` | verdict |
+|---|---|---|---|---|---|
+| — (unmutated control) | **#1392** | — | 0 | 0 | **GREEN**, `SUCCESS!`, 13 assertions |
+| — (unmutated control) | **none** | — | 0 | **1** | already **RED** — #1371, and the reason the rest of this table is overlaid |
+| **M1** — `n_out` -> `R` at the host projection | **#1392** | 1 ins / 1 del | 0 | **1** | **RED**, `FAILURE!` |
+| **M2** — never fill `trace->final_normed` | **#1392** | 1 ins / 2 del | 0 | **1** | **RED**, `FAILURE!` |
+| **M4** — restore the exact pre-repair two-download shape | **#1392** | 3 ins / 8 del | 0 | 0 | **GREEN — reported, not hidden** |
+
+Re-derived independently at `72d736867`, on a clean tree with the same overlay
+applied and reverted: the unmutated control is `run_rc=0`, `1 passed`,
+`13 assertions`, `SUCCESS!`; M1 gives `compile_rc=0`, `run_rc=1`, and
+`ERROR: test case THREW exception: vt: NemotronH lm_head: gathered row count
+does not match hidden_size at nemotron_h.cpp:1029`; M2 gives `compile_rc=0`,
+`run_rc=1`, `assertions: 8 | 7 passed | 1 failed` on `REQUIRE( 0 == 288 )`. The
+M2 edit was written as one deletion rather than the recorded `1 ins / 2 del`
+reshape, so the applied-lines cell is the original author's shape and the
+verdict is reproduced. `M4`'s cell needs no separate attestation: un-overlaid,
+every run of this case is red, so a GREEN is only reachable with the overlay in
+place and its own verdict entails the column.
+
+M1 is the red the `n_out` rename exists for: with the request count substituted
+the returned row count is 1 where the gather asked for 3. M2 arms the
+trace-operand assertion. Both showed the trap this project has been bitten by
+before — doctest printed `assertions: 2 | 2 passed | 0 failed` on M1 while the
+case was failing, because a `REQUIRE` throws rather than counting.
+
+**M4 stays green and that is a finding, not a gap to paper over.** The duplicate
+`DownloadF32` the first submission introduced copies the same unchanged buffer
+twice and produces identical bytes, so nothing observable from outside the
+function can distinguish it. It is repaired structurally — one call is reached on
+any path — and no assertion here pretends to catch it.
+
+The device arm's own red-first does not exist on a CPU box and is not claimed.
+`MarlinW4A16Selects` is false on a CPU queue, so the device branch, the
+`fallback_gemms` assertion and the reachability deletion mutation are all
+unreachable here and stay PENDING a `dgx:gpu0` window.
+
+### The fresh review, and what it found
+
+A fresh reviewer returned FINDINGS on the first submission. The two blocking
+ones are recorded here because both are about EVIDENCE, and evidence is what a
+spec is for.
+
+1. **`tests/vllm/models/test_nemotron_h_moe_device.cpp` had never compiled.**
+   `NemotronHHostWeights` was used unqualified and was missing from the file's
+   using-block, so the file failed `-Wall -Wextra -Werror` with 9 errors on a
+   plain CPU build — the PR head at `rc 1`, the merge-base version of the same
+   file at `rc 0` on the same command. **The red-first result claimed for the
+   synthetic numeric gate therefore did not exist and could not have existed.**
+   The declaration is repaired and the file now compiles at `rc 0` on both
+   arms, but the gate itself is still CUDA-only and still has never executed;
+   the table above says `NEVER RUN` rather than restating a red nobody saw.
+   `vllm_cpp_add_test(test_nemotron_h_moe_device)` is deliberately registered
+   with NO CUDA guard, and that is what surfaced this: a case that skips at run
+   time still has to parse and type-check on every CPU build.
+2. **The production source asserted a protection that does not exist.** A
+   comment at the device branch claimed the allowlist entry was removed and
+   that "the routing checker, not a comment, is what now holds this branch in
+   place". All three parts were false: the entry is narrowed and still present,
+   and deleting the entire `if (DeviceLmHeadEligible(...)) { ... }` block leaves
+   `scripts/check-runner-routing-consistency.py` at `rc 0` with byte-identical
+   output ("3 host-logits off-framework (3 allowlisted)"), reproduced on the
+   repaired tree with the file restored byte-for-byte afterwards (identical
+   sha256). **Nothing automated holds that branch.** The checker is not widened
+   to make it — that changes checker semantics, and [#1410](https://github.com/mudler/vllm.cpp/issues/1410)
+   owns it with its own red-before. The comment now says so, and the
+   reachability deletion mutation stays PENDING rather than claimed.
+
+The reviewer also found a real defect that a token gate structurally cannot
+see. `DeviceLmHeadEligible` restated the shared dispatcher's selection clauses
+and dropped `MarlinW4A16Enabled()`, so under an explicit `VT_NVFP4_MARLIN=0`
+the predicate said eligible while `MatmulNvfp4W4A16D` took its naive
+redundant-dequant arm — computing the SAME logits while re-uploading the whole
+`[131072, 2688]` operand on every decode step, because `LmHeadNvfp4View` hands
+out a stack temporary that `ResidentNvfp4`'s weight-keyed cache can never hit.
+The repair is structural rather than a patched clause: the three clauses now
+live once, in `dense_nvfp4::MarlinW4A16Selects`, and both the dispatcher and
+the model call it, so they cannot drift. `DeviceLmHeadD` additionally refuses
+BY NAME if that same predicate is false on the operand it is about to hand over,
+which catches an eligibility answer taken against a different queue or dtype.
+Deliberately the PREDICATE and not the seam's `fallback_gemms` counter:
+`MutableW4A16Stats()` is a plain non-atomic process-wide static, so a counter
+window in production would refuse a correct run whenever anything else took a
+fallback GEMM concurrently, and a false refusal is worse than the silence it
+replaces. The counter is the right instrument in a test, where
+single-threadedness is a property of the harness, and the synthetic gate asserts
+it there. It is demonstrably armed rather than assumed, because
+`tests/vllm/models/test_qwen3_forward.cpp:559` already asserts on CPU that it
+reaches exactly `5 * num_hidden_layers` when the dispatcher does fall back.
+
+**And the counter is not the only process-wide static on this route.** A row
+whose whole thesis is declining to inherit the sibling's process-static defect
+([#984](https://github.com/mudler/vllm.cpp/issues/984), the address-keyed Marlin
+repack cache) should say what it DOES inherit, so:
+`dense_nvfp4::DenseMarlinWorkspace` (`dense_nvfp4_gemm.h:506`) is a
+process-static **device** allocation — `static void* ws` behind a `static
+std::mutex`, sized from `MarlinDeviceSms` and keyed on nothing, not even the
+device index — and every caller of `MatmulNvfp4MarlinD` (`:539`) and
+`GateUpFusedMarlinD` (`:700`) shares the one buffer. It already has three
+consumers (the shared dense route in `dense_attn_block.h`, MiniMax-H3's
+`minimax_h3_device.cpp`, and the compressed-tensors NVFP4 scheme in
+`schemes/nvfp4.h`); NemotronH's `lm_head` becomes a fourth. It is PRE-EXISTING
+and NOT addressed here: it arrived at `80d1da096`, and the function body is
+byte-identical at this row's merge base and at its head
+(`sha256 bf4685f4...` over the whole definition, both sides). Nothing in this
+row's diff touches it. It is called out because the asymmetry — refusing one
+shared static by name while silently taking another — is the kind of thing a
+reader is entitled to see stated rather than to discover. The
+behavioural red for this class needs CUDA and is PENDING with the rest.
+
+### The DSR ratchet, which no review round ever saw run
+
+`device-leakage` had not COMPLETED on this row through three review passes, so
+its verdict was never an input to any of them. It completed after the third and
+failed: `vt_ifdef` **35 against a baseline of 32**, `rc 1`. Three
+`#ifdef VT_MARLIN_NVFP4` sites had been added to the device-agnostic shared
+layer, and the ratchet exists precisely to stop that accreting.
+
+The three, located by running `scripts/check-device-leakage.py --report` at the
+failing head `7a3909187` and diffing the per-file table against `9ecaf1bb3`
+rather than by reading the diff for guards:
+
+| # | site at `7a3909187` | what the guard decided | resolution |
+|---|---|---|---|
+| 1 | `include/vllm/model_executor/models/dense_nvfp4_gemm.h:768` — `MarlinW4A16Selects` | **nothing.** `MarlinW4A16Enabled()` is declared above the guarded region and `vt::OpRegistered` is the op table's own availability answer | **removed** |
+| 2 | `src/vllm/model_executor/models/nemotron_h_device.cpp:883` — around `LmHeadNvfp4View` | **nothing.** The function names no symbol the Marlin build adds | **removed** |
+| 3 | `src/vllm/model_executor/models/nemotron_h_device.cpp:985` — `DeviceLmHeadD`'s body | `ResidentIn` and a complete `dense_nvfp4::MarlinDenseResident`, neither of which EXISTS without the guarded arena region | **`DSR-ALLOW(A2-Q2b)`** |
+
+**(1) is the case the checker's own message describes.** The build flag and the
+registration are the same condition, not two: `CMakeLists.txt`'s single
+`if(VLLM_CPP_MARLIN)` block adds `src/vt/cuda/cuda_moe_marlin.cu` — whose
+file-scope `Registrar` holds the tree's only
+`RegisterOp(OpId::kMoeGroupedGemmNvfp4Marlin, …)` — and defines
+`VT_MARLIN_NVFP4=1`, in that same block. A build without the macro therefore
+registers nothing and the query already resolves false on exactly the builds the
+`#ifdef` excluded. This is the call `nemotron_h_device.cpp`'s `moe_on_device`
+selection had already made, in a comment that says so.
+
+**(2) was measured, not reasoned.** The claim is that the function's external
+linkage at `namespace vllm` scope is what makes an unused definition harmless in
+a build where its only call site is compiled out. Adding `static` to that
+definition — the mutation that removes exactly that property — turns the same
+CPU compile RED, `error: 'vllm::Nvfp4Weight vllm::LmHeadNvfp4View(...)' defined
+but not used [-Werror=unused-function]`, `rc 1`; the file was restored to an
+identical sha256 (`9719ea70…`) afterwards.
+
+**(3) is TYPES-not-behaviour, and takes the checker's documented escape hatch
+rather than a baseline change.** `AGENTS.md` forbids making a red gate green by
+widening an assertion, and a baseline bump is that. `DSR-ALLOW` is not: the site
+is excluded from the count but COUNTED AND PRINTED on every run, so the
+exemption is visible in CI output. It is the same class, and carries the same
+stated reason, as the five sibling guards A2-Q2a and A2-P already hold in this
+file. The SELECTION for this arm is a runtime op-table query
+(`DeviceLmHeadEligible` → `MarlinW4A16Selects`, which now carries no guard);
+only the call site needs the build guard, and its `#else` refuses by name.
+
+**Measured on this tree**, a CPU build with `VT_MARLIN_NVFP4` absent from
+`build/compile_commands.json` (positive control: 1020 `VLLM_CPP` hits in the
+same file, so the grep is not silently wrong) — which is the configuration that
+exercises both removals, since it is the arm the deleted `#else` branches used
+to serve.
+
+| what | before (`7a3909187`) | after | `rc` |
+|---|---|---|---|
+| `check-device-leakage.py` `vt_ifdef` | 35 | **32** | 1 → **0** |
+| `DSR-ALLOW` exemptions in force | 20 | **21** | — |
+| `scripts/device-leakage-baseline.json` | 32 | **32, untouched** | — |
+| per-file table vs `origin/main` | +1 header, +2 model TU | **identical** | — |
+| `nemotron_h_device.cpp`, `nemotron_h.cpp`, `qwen3_5.cpp` at `-Wall -Wextra -Werror` | — | compile | **0** |
+
+Each of the three repairs is load-bearing, proven by reverting it alone in a
+scratch worktree and re-running the gate. Every mutation was verified applied by
+`git diff --stat` and restored to an identical sha256, with the unmutated
+control green immediately before and after.
+
+| mutation | applied | `vt_ifdef` | `rc` | verdict |
+|---|---|---|---|---|
+| — (control) | — | 32 | 0 | `ratchet holds` |
+| **M-B** — restore the guard on `MarlinW4A16Selects` | 4 ins | **33** | **1** | **RED**, `DSR REGRESSION` |
+| **M-C** — restore the guard around `LmHeadNvfp4View` | 2 ins | **33** | **1** | **RED**, `DSR REGRESSION` |
+| **M-D** — delete the `DSR-ALLOW(A2-Q2b)` line | 1 del | **33** | **1** | **RED**, `DSR REGRESSION` |
+| — (control, after restore) | — | 32 | 0 | `ratchet holds` |
+
+The two pre-existing allowlist entries the report also prints —
+`deepseek_v4_device.cpp [kcuda] x8` and `platform.cpp [dev_cast] x1` — are
+byte-identical at `9ecaf1bb3` and here. This change moves one bucket and
+nothing else.
+
+**The #1392 overlay is retired.** The two evidence tables above in this section
+carry an `overlay` column because this tree could not construct `GPUModelRunner`
+at all ([#1371](https://github.com/mudler/vllm.cpp/issues/1371)) and every green
+was taken with [#1392](https://github.com/mudler/vllm.cpp/pull/1392)'s fix
+applied to the working tree and reverted. #1392 has since landed on `main` and
+this branch has merged it, so the tree now carries the fix as a committed
+object. The historical rows keep their `overlay` cells, because they describe
+the tree they were measured on and rewriting them would make them false.
+
+### The full CPU suite, and a red that arrived from `main` mid-repair
+
+`origin/main` moved four times while this repair was gated, and the third sync
+brought `4712dac40` (`VT-ACT-ROUND-POLARITY`, [#1322](https://github.com/mudler/vllm.cpp/issues/1322)
+via [#1347](https://github.com/mudler/vllm.cpp/pull/1347)). Both `ctest` runs
+below are on this box, `-Wall -Wextra -Werror`, no CUDA toolkit, and neither
+carries an overlay of any kind — #1392 is a committed object here now.
+
+| tree | base | `ctest` | result |
+|---|---|---|---|
+| this row + `9ecaf1bb3` | before `4712dac40` | 567 | **`100% tests passed, 0 tests failed out of 567`**, `rc 0` |
+| this row + `01854663c` | after `4712dac40` | 569 | `rc 8`, **4 failed**: `test_minimax_music3_ar`, `test_ltx2_text_encoder`, `test_muse_glimmer_text`, `test_muse_glimmer_text_fallback` |
+
+The four are [#1458](https://github.com/mudler/vllm.cpp/issues/1458), filed from
+another flow before this control was run, and they are inherited rather than
+caused. That is proven in both directions on the same tree, each mutation
+verified applied by `git diff --stat` and restored to an identical sha256:
+
+| control | change | `compile_rc` | `run_rc` (music3 / ltx2 / glimmer) | verdict |
+|---|---|---|---|---|
+| **A** | revert THIS row's two source files to their pre-repair `7a3909187` content | 0 | **1 / 1 / 1** | still RED — **not this row's** |
+| **B** | revert `src/vt/cpu/cpu_ops.cpp` alone to `4712dac40^` | 0 | **0 / 0 / 0**, `37/37`, `27/27`, `24/24` | GREEN — **`4712dac40` is the cause** |
+
+Control A is the one that answers the attribution question and it answers it
+alone; control B is here because naming a cause is more useful to the next
+reader than clearing oneself. Neither repairs anything: #1458 needs
+`VT-ACT-ROUND-POLARITY` to decide whether the kernel or four never-re-derived
+bf16 error floors are wrong, which is that row's oracle work and not a small,
+clear, in-flow fix.
+
+`test_nemotron_h_paged_forward` — the test all three of this row's red CI jobs
+failed on — is green on both trees: `13 | 13 passed | 0 failed`,
+`assertions: 3269 | 3269 passed | 0 failed`, `Status: SUCCESS!`, `run_rc 0`.
diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md
index 933b949c4..fcac92ebb 100644
--- a/docs/BENCHMARKS.md
+++ b/docs/BENCHMARKS.md
@@ -13,6 +13,7 @@
 | **Container images (ACTIVE; arm64 cuda verified on GB10 + Orin 2026-08-11)** | `ENG-RELEASE-CONTAINERS` ([spec](../.agents/specs/container-images.md)) | cpu amd64 783 MB; cuda arm64 **1.71 GB**. GB10 `sm_121a`: `/health`+`/version`+SIGTERM on `--gpus all`. Orin `sm_87` (Tegra): Qwen3-0.6B **generates**, GPU **GR3D 95-97%** | n/a |
 | **Developer/row protocol** | Contribution entry point; `ENG-NOW-DERIVED` #374 @`dbd0d51c` | Entry-point gates retained. #374 W1-W5 DONE; benchmark/runtime/parity `VOID`; row specs now carry `## Now` | n/a |
 | **NemotronH paged forward** (`MODEL-NEMOTRON-H-ABI-A2P`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | **No speed number, by the unit's own rule** ([spec](../.agents/specs/nemotron-h-a2p-paged-forward.md) §5) | **A3 host gate PASSES 96/96 `STRICT PASS`.** GB10 read 4/24; cause and fix [#1157](https://github.com/mudler/vllm.cpp/issues/1157), sm_121a re-run pending a lease | CPU gate 12/12. Load 209.0 s, peak RSS 20 142 392 KB |
+| **NemotronH host re-expansion / decode token** (`A2-Q2b`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | Real Nemotron-3.5-Lightning-30B NVFP4 (rev `29f2d174`), T=1 decode step, counted at the `DenseBf16` seam via `examples/nemotron_h_gen` | Full host arm **3 078 512 640 elements (6.157 GB bf16)**, 369 calls. `lm_head` **352 321 536 (11.44%)**, **28.35%** of the post-A2-Q2a residue, **100%** once the mamba arm lands | Per-group table in [the record](../.agents/benchmark-record.md). Attribution only, no speed number. Device `lm_head` gate PENDING a `dgx:gpu0` window |
 | **Record-anchor ratchet** (`ENG-RECORD-ANCHOR-RATCHET`, #632) | **No number owed:** a record checker. At `8daa67b39`, **832 of 867** in-scope citations (**96.0%**) were already parsed; no symbol test and no report ran. Rot **38** (32 stale, 6 broken), **32 in range** |
 | **LoRA runtime W2** (`LORA-RUNTIME`, #278) | **No number owed:** correctness-only; a grid PENDS the W7 model gate |
 | **ARCH audit: ABI is text-only** | 4 capabilities (H3 video, Laguna, Kimi-Linear, DeepSeek-V4) reachable only from `examples/`, none registry-backed. No gate asks whether a CONSUMER can reach a capability. Documentation only |
diff --git a/docs/FEATURES.md b/docs/FEATURES.md
index 3289ebc35..26ff27229 100644
--- a/docs/FEATURES.md
+++ b/docs/FEATURES.md
@@ -148,7 +148,7 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks.
 | `LagunaForCausalLM` | poolside/Laguna-S-2.1-NVFP4, GGUF-Q4_K, Laguna-XS | byte-exact near-tie (distributional vs vLLM) | vLLM parity+ 1.03x, default on, via the `laguna-gen` CLI; the registered engine forward VT_CHECKs non-bf16 (`ARCH-ONE-SURFACE` fold) |
 | `KimiLinearForCausalLM` | Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE) | **Folded onto the shared paged runner (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; vs golden 122/128 (the intrinsic near-tie profile); FA2 paged MLA default-ON; SACRED post-fold green** | Served via `vllm_engine_load` + `vllm_complete_tokens` (ABI v13); server 19.0 tok/s wall vs vLLM ~21 (~0.90×), speed residual open |
 | `KimiK3ForConditionalGeneration` | Kimi-K3 (2.8T MoE) | scaffold: registry+config+enumeration gated, forward refuses | HW-infeasible (~1.56 TB); no run |
-| `NemotronHForCausalLM` | Nemotron-3.5-Lightning-30B-A3B-NVFP4 (`nvidia` @`29f2d174`) | config+enumeration+KV-shape gated; hybrid forward COMPUTES; loader materializes 18487/18487 as SHIPPED | **PAGED (#810 A2-P): K/V go into the runner's pages and conv+SSM rows are carried at the metadata's state indices.** G-SAFE narrows to `num_reqs <= 1`. Owed: device `lm_head` (A2-Q2b), FP8 mamba (A2-Q1), MTP, GGUF |
+| `NemotronHForCausalLM` | Nemotron-3.5-Lightning-30B-A3B-NVFP4 (`nvidia` @`29f2d174`) | config+enumeration+KV-shape gated; hybrid forward COMPUTES; loader materializes 18487/18487 as SHIPPED | **PAGED (#810 A2-P): K/V go to the runner's pages; conv+SSM rows carry at the metadata's state indices.** G-SAFE: `num_reqs <= 1`. Device `lm_head` (A2-Q2b), UNMEASURED. Owed: FP8 mamba (A2-Q1), MTP, GGUF |
 | `MuseGlimmerForCausalLM` | real tensors, **bf16 depth 4/52 only**: 5 prefill argmax positions match a torch transcription of vllm#51655 and HF. GGUF full depth generates coherently (#347, #359) but is **NOT token-exact** | text forward + loader vs an fp32 reference, per-mechanism property tests, scaffold 11/11, GGUF gate 17/17. An ABSENT config key now takes the architecture's constant (#412): GGUF post-norms ran at 1e-5, not 1e-8 | no vLLM denominator (pin cannot load it); SECONDARY llama.cpp, same GGUF, GB10 CPU: prefill tie **0.997x**, decode 0.232x, RSS 1.92x (#333) |
 | `MuseGlimmerForConditionalGeneration` | vision: **no reference run of any kind**; enumeration gated vs the released 30B index (1436/1436). Image/video need bf16 safetensors: `mmproj-kquant.gguf` is refused by name | perception encoder loaded and wired, so an image or video prompt runs; `perception_emb_norm` now armed by default (#405). Reachability plus placeholder scatter only, no image or video correctness | not measurable; anchored to open vllm#51655 |
 | `LlamaModel` | landed tiny synthetic embedding fixture (engine path == direct pooler path, identical vectors; f64 LAST+normalize reference); real checkpoint (e5-mistral class) is a NAMED residual | pooling/embed only, text paths refuse by task; `vllm_embed` + `/v1/embeddings` | n/a (CPU correctness-grade embeddings) |
diff --git a/docs/USAGE.md b/docs/USAGE.md
index a3047fc51..f4afb6dc2 100644
--- a/docs/USAGE.md
+++ b/docs/USAGE.md
@@ -227,7 +227,7 @@ lists other published arms when they have not been used as a gated checkpoint.
 | Model or component | File | Size | Repository and revision | Quantized SHA-256 | Supported arms | Refused arms or missing part |
 |---|---|---|---|---|---|---|
 | DSpark for Qwen3.8-27B | `model.safetensors` | 2,718,576,122 bytes | `RadixArk/Qwen3.8-27B-DSpark` @ `85ef153be924f17ce4bf62726954eeaa4a73e854` | n/a (non-quantized) | Qwen3 DSpark routing | Token-exact decode gate is pending |
-| Nemotron-3.5-Lightning-30B | `model-000{01..52}-of-00052.safetensors` | 21,583,809,748 bytes total | `nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` @ `29f2d1746d8f41e316523194b19018707749b1b1` | `672c8bda10fdec0256e0819e112d2aa3a936cc3e5d311a05fd3ff773ca9a44b9` (first shard) | Device bf16, GQA, and NVFP4 experts; host FP8 Mamba2 and NVFP4 head | GGUF, MTP, and batched decode |
+| Nemotron-3.5-Lightning-30B | `model-000{01..52}-of-00052.safetensors` | 21,583,809,748 bytes total | `nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` @ `29f2d1746d8f41e316523194b19018707749b1b1` | `672c8bda10fdec0256e0819e112d2aa3a936cc3e5d311a05fd3ff773ca9a44b9` (first shard) | Device bf16, GQA, NVFP4 experts, and the NVFP4 head (A2-Q2b, unmeasured); host FP8 Mamba2 | GGUF, MTP, and batched decode |
 | MiniMax-H3 FL2VA | `MiniMax-H3-FL2VA-Q4_K_M.gguf` | 19,864,208,160 bytes | `realrebelai/MiniMax-H3_GGUFs` @ `daf03b4ca652cce16dfd4fcf91e79c52ffa5c1e7` | `5e8fa6e960d5fbd547390ceec63fcead275435d8f3bd2466a8a2cbd8c2e361e3` | Q4_K_M `t2va` and `fl2va`, verified end to end | `ref2va` requires the REF2VA partition |
 | MiniMax-H3 REF2VA | `MiniMax-H3-REF2VA-Q4_K_M.gguf` | 19,864,208,064 bytes | `realrebelai/MiniMax-H3_GGUFs` @ `daf03b4ca652cce16dfd4fcf91e79c52ffa5c1e7` | `17925612821ea3037ffaf5f7f9789f5460e87025385bd45e9ec6c7d536684d56` | Q4_K_M `ref2va`, verified end to end | `t2va` and `fl2va` require the FL2VA partition |
 | MiniMax-H3 encoder | `qwen3vl-32B-MiniMax-H3-Q4_K_M.gguf` | 14,576,977,888 bytes | `realrebelai/MiniMax-H3_GGUFs` @ `daf03b4ca652cce16dfd4fcf91e79c52ffa5c1e7` | `1bf75e038c5895b97b6ea16cc1e3d32076254b06ec3df10657650d86dc82279e` | Q4_K_M text and multimodal conditioning | No separate refused arm recorded |
diff --git a/docs/models/nemotron-3-5-lightning.md b/docs/models/nemotron-3-5-lightning.md
index aff529fb6..b98a9557e 100644
--- a/docs/models/nemotron-3-5-lightning.md
+++ b/docs/models/nemotron-3-5-lightning.md
@@ -58,7 +58,7 @@ comparison can see, so it is written down here instead.
 | bf16 layers, norms, the 6 GQA attention blocks | **device** |
 | MoE experts, NVFP4 W4A16 g16 | **device** (Marlin arena) |
 | FP8 W8A8 static Mamba2 input projections | **host** — the device arm is owed, [#940](https://github.com/mudler/vllm.cpp/issues/940) |
-| `lm_head`, NVFP4 W4A16 g16 | **host** — it refuses a non-CPU queue by name, so the forward's last step is a host projection and the model still returns host logits. Owed to A2-Q2b, [#810](https://github.com/mudler/vllm.cpp/issues/810) |
+| `lm_head`, NVFP4 W4A16 g16 | **device** on the paged forward (A2-Q2b, [#810](https://github.com/mudler/vllm.cpp/issues/810)) — it runs through the shared NVFP4 W4A16 route `dense_nvfp4::MatmulNvfp4W4A16D`, so `NemotronHPagedForward` returns device-resident logits to the on-GPU sampler instead of a host projection. **The device arm's own numeric gate has NOT run yet** (it needs a `dgx:gpu0` window); until it does, treat this row as "implemented and unmeasured", which is what the spec's `## Now` records. The host projection is retained and is what serves a build without the Marlin NVFP4 GEMM, a non-NVFP4 `lm_head`, an explicit `VT_NVFP4_MARLIN=0`, and the numeric gate's reference side. `VT_NVFP4_MARLIN=0` is the same-binary A/B escape hatch for the Marlin NVFP4 W4A16 GEMM, and it now routes this arm to the host projection in EVERY build: the model's eligibility test and the shared dispatcher call one predicate, `dense_nvfp4::MarlinW4A16Selects`, rather than restate its clauses, and that predicate asks the op/provider table whether the Marlin arm is realized for the device instead of asking the preprocessor |
 
 And the arms that are **refused by name** rather than substituted:
 
@@ -68,6 +68,27 @@ And the arms that are **refused by name** rather than substituted:
 | the MTP draft head | deferred by name at load (W5) |
 | batched decode (`num_reqs > 1`) | refused at the forward. One request's KV pages and one request's recurrent state are carried per step; a multi-request step would be decoded as ONE concatenated causal sequence and would return plausible wrong tokens instead of failing. Owed to A2-B, [#810](https://github.com/mudler/vllm.cpp/issues/810) |
 
+### What the host arms cost, per decode token
+
+Counted on the real checkpoint at the single dequant seam, one decode step
+(T=1, `top_k` 6, 23 MoE layers), because "where the arithmetic happens" is only
+half the answer — the other half is how much of it there is:
+
+| group | elements re-expanded per token | share |
+|---|---|---|
+| routed experts (6 of 128, x 23 layers) | 1 376 944 128 | 44.7% |
+| shared experts (x 23) | 458 981 376 | 14.9% |
+| Mamba2 FP8 `in_proj` + `out_proj` (x 23) | 890 265 600 | 28.9% |
+| `lm_head` | 352 321 536 | 11.4% |
+| **total** | **3 078 512 640 (6.157 GB at bf16)** | |
+
+Only the groups marked **host** above are actually paid at run time. With the
+MoE arm on the device that is 1 242 587 136 elements (2.485 GB) per token, and
+with `lm_head` on the device too it is the Mamba2 FP8 pair alone. `lm_head` is
+the largest SINGLE re-expansion in the model by 12.7x — one call, a 704.6 MB
+transient bf16 buffer — which is why it matters more than its 11.4% share
+suggests on a unified-memory box that reboots rather than OOM-kills.
+
 ## What has NOT been measured
 
 **No token gate result exists for this checkpoint yet.** The example above is the
diff --git a/docs/reference/model-loading.md b/docs/reference/model-loading.md
index 5401f2545..eda8a0a72 100644
--- a/docs/reference/model-loading.md
+++ b/docs/reference/model-loading.md
@@ -185,7 +185,7 @@ tokens quietly.
 | Architecture | Why it refuses |
 |---|---|
 | `KimiK3ForConditionalGeneration` | Needs ~1.56 TB (MXFP4); no host here can run it |
-| `NemotronHForCausalLM` | Single-request paged decode runs. Batched decode refuses when `num_reqs > 1`. The token gate remains pending, `lm_head` and FP8 Mamba2 projections run on the host, and GGUF is unavailable. See the [Nemotron-3.5-Lightning model recipe](../models/nemotron-3-5-lightning.md) and [benchmark state](../BENCHMARKS.md). |
+| `NemotronHForCausalLM` | Single-request paged decode runs. Batched decode refuses when `num_reqs > 1`. The token gate remains pending, `lm_head` reaches the device on the paged forward (A2-Q2b, implemented and unmeasured) while the FP8 Mamba2 projections still run on the host, and GGUF is unavailable. See the [Nemotron-3.5-Lightning model recipe](../models/nemotron-3-5-lightning.md) and [benchmark state](../BENCHMARKS.md). |
 
 This is a deliberate state, not a bug: registering the architecture is what lets
 the config parse and weight-name mapping be tested before the forward exists.
diff --git a/include/vllm/model_executor/models/dense_nvfp4_gemm.h b/include/vllm/model_executor/models/dense_nvfp4_gemm.h
index 6a21cccad..fb2c2fdf4 100644
--- a/include/vllm/model_executor/models/dense_nvfp4_gemm.h
+++ b/include/vllm/model_executor/models/dense_nvfp4_gemm.h
@@ -77,6 +77,31 @@ using vt::Backend;
 using vt::DType;
 using vt::Tensor;
 
+// ── A2-Q2b (#810): the CALLER-OWNED resident, so a consumer can opt OUT of the
+// address-keyed cache below without hand-rolling a parallel Marlin path.
+//
+// `MarlinDenseResidentFor` (below) keys its `static` map on the ADDRESS of the
+// `Nvfp4Weight`. That is issue #984: destroy one engine and build another in
+// the same process and the allocator can hand the new weights the old address,
+// so the new weights inherit an entry marked `ready` whose device pointers
+// belong to the previous engine. `Nvfp4Weight` already carries the
+// `resident_marlin` `ResidentSlot` that #237 added to fix exactly this, and
+// this header does not read it.
+//
+// A2-Q2b does not fix #984 — that needs the two-engine red-before #984 asks
+// for, and it is out of this row's scope. It declines to INHERIT it: NemotronH
+// owns its lm_head resident in a `ResidentSlot` (the same shape A2-Q2a used for
+// the MoE arena) and hands it in here. Threading the resident through is
+// additive — every existing caller passes nothing and keeps the address-keyed
+// default byte-for-byte — so this EXTENDS the seam rather than forking it,
+// which is what AGENTS.md `## Shared seams` asks for when the seam cannot
+// represent the behaviour a consumer needs.
+//
+// Declared unconditionally (defined only under VT_MARLIN_NVFP4) so the
+// dispatcher's signature does not change with the build flag; a pointer to an
+// incomplete type is all the non-Marlin build ever needs.
+struct MarlinDenseResident;
+
 // VT_NVFP4_MARLIN (default ON): the vendored Marlin NVFP4 W4A16 GEMM is the
 // validated path (35B gate +22%, token-for-token vs the pinned oracle). Only an
 // explicit VT_NVFP4_MARLIN=0 opts back out to the naive redundant-dequant kernel
@@ -504,9 +529,11 @@ inline void* DenseMarlinWorkspace(Dev d, int* out_sms) {
 
 // y[M,N] = x[M,K] bf16 @ dequant(w).T via the single-expert Marlin W4A16 GEMM.
 inline DBuf MatmulNvfp4MarlinD(Dev d, const Tensor& x, const Nvfp4Weight& w,
-                               DType out_dtype) {
+                               DType out_dtype, MarlinDenseResident* resident = nullptr) {
   const int64_t M = x.shape[0], K = x.shape[1], N = w.n;
-  MarlinDenseResident& mr = MarlinDenseResidentFor(&w);
+  // A2-Q2b: a caller that owns the resident's lifetime supplies it; everyone
+  // else keeps the address-keyed cache, unchanged (#984).
+  MarlinDenseResident& mr = resident != nullptr ? *resident : MarlinDenseResidentFor(&w);
   if (!mr.ready) BuildMarlinDenseResident(d, w, mr);
   int sms = 0;
   void* ws = DenseMarlinWorkspace(d, &sms);  // zeroed once; kernel self-resets
@@ -721,6 +748,47 @@ inline DBuf GateUpFusedMarlinD(Dev d, const Tensor& x, const Nvfp4Weight& gw,
 #endif  // VT_MARLIN_NVFP4
 
 // --- The W4A16 dispatcher ---------------------------------------------------
+
+// ── A2-Q2b (#810): the dispatcher's OWN gate, as ONE predicate ───────────────
+//
+// `MatmulNvfp4W4A16D` below selects Marlin on THREE clauses. A caller that has
+// to know the answer BEFORE it calls — because the fallback is not merely
+// slower for it but pathological — must ask this function rather than restate
+// the clauses, and this exists because a restatement is exactly what went
+// wrong. A2-Q2b's first submission restated two of the three in NemotronH's
+// `DeviceLmHeadEligible` and dropped `MarlinW4A16Enabled()`, so under an
+// explicit `VT_NVFP4_MARLIN=0` the model said "eligible", the dispatcher fell
+// to the naive redundant-dequant arm, and because the caller hands in a
+// TRANSIENT `Nvfp4Weight` view the `ResidentNvfp4` cache below could never hit
+// — re-uploading the whole `[vocab, hidden]` operand on every decode step.
+//
+// A token gate cannot see any of that: the fallback computes the SAME value.
+// One expression with two call sites is what makes the two unable to drift.
+inline bool MarlinW4A16Selects(Dev d, DType act_dtype) {
+  // NO `#ifdef VT_MARLIN_NVFP4` HERE, and the absence is the design. This site
+  // only SELECTS a path, and every term it reads exists in every build:
+  // `MarlinW4A16Enabled()` is declared above the guarded region, and
+  // `vt::OpRegistered` IS the op/provider table's own answer to "is the Marlin
+  // arm realized for this device".
+  //
+  // The build flag and the registration are the SAME condition, so the guard
+  // would decide nothing the query does not already decide. `CMakeLists.txt`'s
+  // one `if(VLLM_CPP_MARLIN)` block adds `src/vt/cuda/cuda_moe_marlin.cu` —
+  // whose file-scope `Registrar` holds the tree's only
+  // `RegisterOp(OpId::kMoeGroupedGemmNvfp4Marlin, …)` — and defines
+  // `VT_MARLIN_NVFP4=1`, in that same block. A build without the macro
+  // therefore registers nothing, and this resolves false on exactly the builds
+  // a `#ifdef` would have excluded. Asking the table rather than the
+  // preprocessor is what `scripts/check-device-leakage.py` asks of a selection,
+  // and it is the call `nemotron_h_device.cpp`'s `moe_on_device` already made.
+  //
+  // The device clause is likewise an OP-AVAILABILITY question, not a
+  // `== kCUDA` one (registered only for kCUDA today, so this is byte-identical
+  // on the production build — accelerator-seam audit class A, work row S4).
+  return vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, d.q.device.type) &&
+         MarlinW4A16Enabled() && act_dtype == DType::kBF16;
+}
+
 // y[M,N] = x[M,K] @ dequant(w).T for an NVFP4 W4A16 weight. Mirrors vLLM's
 // forced-Marlin selection for `use_a16` (__init__.py:879-881): on CUDA with a
 // BF16 activation take Marlin; otherwise fall back to the naive
@@ -728,19 +796,18 @@ inline DBuf GateUpFusedMarlinD(Dev d, const Tensor& x, const Nvfp4Weight& gw,
 // reference). `w` MUST be W4A16 (alpha == 0) — a true-W4A4 weight belongs to
 // qwen3_5.cpp's private fp4-activation path and is rejected here.
 inline DBuf MatmulNvfp4W4A16D(Dev d, const Tensor& x, const Nvfp4Weight& w,
-                              DType out_dtype) {
+                              DType out_dtype, MarlinDenseResident* resident = nullptr) {
   const int64_t M = x.shape[0], K = x.shape[1], N = w.n;
   VT_CHECK(!w.IsTrueW4A4(),
            "dense_nvfp4: true-W4A4 weight routed into the W4A16 dispatcher");
 #ifdef VT_MARLIN_NVFP4
   // Marlin requires a bf16 activation (vLLM's a16 path is bf16/fp16 too). The
-  // device gate is an OP-AVAILABILITY question, not a "== kCUDA" question: ask
-  // the vt::OpProvider table whether the Marlin NVFP4 grouped-GEMM is realized
-  // for this device (registered only for kCUDA today, so this is byte-identical
-  // on the production build — accelerator-seam audit class A, work row S4).
-  if (vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, d.q.device.type) &&
-      MarlinW4A16Enabled() && x.dtype == DType::kBF16)
-    return MatmulNvfp4MarlinD(d, x, w, out_dtype);
+  // three clauses live in `MarlinW4A16Selects` above so a caller can ask the
+  // SAME question in advance instead of restating them (A2-Q2b).
+  if (MarlinW4A16Selects(d, x.dtype))
+    return MatmulNvfp4MarlinD(d, x, w, out_dtype, resident);
+#else
+  (void)resident;
 #endif
   ++MutableW4A16Stats().fallback_gemms;
   DBuf dout(d, out_dtype, {M, N});
diff --git a/scripts/nemotron-h-a2q2b-gpu-gate.sh b/scripts/nemotron-h-a2q2b-gpu-gate.sh
new file mode 100755
index 000000000..76d4aceaa
--- /dev/null
+++ b/scripts/nemotron-h-a2q2b-gpu-gate.sh
@@ -0,0 +1,173 @@
+#!/usr/bin/env bash
+# A2-Q2b (#810) dgx:gpu0 gate runner — the device `lm_head` on NVFP4 W4A16 g16.
+#
+# .agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md.
+#
+# ── WHY THIS FILE EXISTS RATHER THAN A COMMAND LINE ────────────────────────
+# The spec's §4 records that this row has lost FOUR GB10 windows to environment
+# rather than to code — `121` instead of `121a`, an unconstrained build, a
+# CUTLASS fetch with no egress, and an 8h19m outage. Every one of those is a
+# precondition that can be checked before the expensive part starts, so they are
+# checked here, once, and a violation VOIDS the run loudly instead of producing
+# a number nobody can use.
+#
+# Env:
+#   SRC        checkout to build (required)
+#   BUILD      build directory (required; put it on /tmp, NOT /workspace — CIFS
+#              holds no symlink and the link step fails there)
+#   ARCH       CUDA arch, must be 121a on GB10 (default 121a)
+#   LOG_ROOT   where to write the run log (required)
+#   CKPT       Nemotron-3.5-Lightning NVFP4 snapshot (default from CHECKPOINT_ROOT)
+set -uo pipefail
+
+SRC="${SRC:?set SRC}"
+BUILD="${BUILD:?set BUILD}"
+ARCH="${ARCH:-121a}"
+LOG_ROOT="${LOG_ROOT:?set LOG_ROOT}"
+CKPT="${CKPT:-${CHECKPOINT_ROOT:-/usr/local/nas_share/checkpoints}/nemotron-3.5-lightning-30b-nvfp4}"
+mkdir -p "$LOG_ROOT"
+LOG="$LOG_ROOT/a2q2b-$(date -u +%Y%m%dT%H%M%SZ).log"
+rc=0
+
+say() { echo "=== $* ===" | tee -a "$LOG"; }
+
+# `run` reports the COMMAND's exit status, never the pipeline's. `cmd | tail`
+# reports tail's status, which is 0 essentially always — a whole gate series has
+# read green that way in this tree.
+run() {
+  local name="$1"; shift
+  say "$name"
+  "$@" >>"$LOG" 2>&1
+  local status=$?
+  echo "### $name -> exit $status" | tee -a "$LOG"
+  [ "$status" -ne 0 ] && rc=1
+  return 0
+}
+
+say "A2Q2B GATE start $(date -u +%FT%TZ) host=$(hostname) src=$SRC arch=$ARCH"
+(cd "$SRC" && git rev-parse HEAD && git status --short) | tee -a "$LOG"
+
+# ── PRECONDITION 1: host headroom, sampled INSIDE the run, not before it ─────
+free -g | tee -a "$LOG"
+AVAIL=$(free -g | awk '/^Mem:/{print $7}')
+if [ "${AVAIL:-0}" -lt 60 ]; then
+  say "VOID: only ${AVAIL}G available, floor is 60G (spec §4.2). The box reboots \
+rather than OOM-killing at gpu_memory_utilization; this is not a failure of the code."
+  exit 3
+fi
+
+# ── the toolchain: the worker provisions ITSELF ──────────────────────────────
+# No CUDA toolkit is preinstalled on an rc worker (AGENTS.md), and the first
+# submission of this gate VOIDed in 60 seconds on exactly that. The recipe below
+# is the one row A2-D1 proved on this box, reused verbatim rather than
+# re-derived: `sbsa` (not x86) for GB10, and the 13-0 package set.
+say "toolchain"
+export DEBIAN_FRONTEND=noninteractive
+apt-get update -qq >>"$LOG" 2>&1
+apt-get install -y -qq git cmake ninja-build g++ curl ca-certificates python3 python3-dev >>"$LOG" 2>&1
+if ! command -v nvcc >/dev/null 2>&1 && [ ! -x /usr/local/cuda/bin/nvcc ]; then
+  curl -fsSL -o /tmp/cuda-keyring.deb \
+    https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2404/sbsa/cuda-keyring_1.1-1_all.deb >>"$LOG" 2>&1
+  dpkg -i /tmp/cuda-keyring.deb >>"$LOG" 2>&1
+  apt-get update -qq >>"$LOG" 2>&1
+  apt-get install -y -qq cuda-nvcc-13-0 cuda-cudart-dev-13-0 libcublas-dev-13-0 \
+     cuda-nvrtc-dev-13-0 cuda-nvtx-13-0 cuda-profiler-api-13-0 libcurand-dev-13-0 >>"$LOG" 2>&1
+fi
+export PATH=/usr/local/cuda/bin:$PATH
+nvcc --version >>"$LOG" 2>&1 || { say "VOID: no nvcc after the install step"; exit 3; }
+
+# ── PRECONDITION 2: the toolkit is proved by a LINK, not by --version ────────
+# A2-D1's check, kept because `nvcc --version` succeeds on an install that
+# cannot link, and this row has already lost windows to environment.
+cat > /tmp/probe_a2q2b.cu <<'PROBE'
+#include <cublasLt.h>
+#include <cstdio>
+int main() {
+  cublasLtHandle_t h = nullptr;
+  const auto s = cublasLtCreate(&h);
+  std::printf("cublasLtCreate=%d\n", static_cast<int>(s));
+  return s == CUBLAS_STATUS_SUCCESS ? 0 : 1;
+}
+PROBE
+if ! nvcc -arch="sm_$ARCH" /tmp/probe_a2q2b.cu -o /tmp/probe_a2q2b -lcublasLt >>"$LOG" 2>&1; then
+  say "VOID: the CUDA toolkit does not link cublasLt at sm_$ARCH"
+  exit 3
+fi
+/tmp/probe_a2q2b >>"$LOG" 2>&1
+
+# ── CUTLASS: staged, never fetched. The HOST has no egress to github.com, and
+# a CUTLASS fetch with no egress is one of the four windows this row has lost.
+CUTLASS="${CUTLASS:-/root/cutlass-v4.5.0}"
+if [ ! -f "$CUTLASS/include/cutlass/cutlass.h" ]; then
+  if [ -f /workspace/cutlass-v4.5.0.tar.gz ]; then
+    mkdir -p "$CUTLASS" && tar xzf /workspace/cutlass-v4.5.0.tar.gz -C "$CUTLASS" --strip-components=1 >>"$LOG" 2>&1
+  else
+    git clone --depth 1 --branch v4.5.0 https://github.com/NVIDIA/cutlass "$CUTLASS" >>"$LOG" 2>&1
+  fi
+fi
+[ -f "$CUTLASS/include/cutlass/cutlass.h" ] || { say "VOID: no CUTLASS headers at $CUTLASS"; exit 3; }
+
+# ── the build ────────────────────────────────────────────────────────────────
+# -j 4: unconstrained parallelism has OOM-REBOOTED this box (AGENTS.md).
+say "configure"
+CFGLOG="$LOG_ROOT/configure-$(date -u +%H%M%SZ).log"
+cmake -S "$SRC" -B "$BUILD" -G Ninja \
+  -DCMAKE_BUILD_TYPE=Release \
+  -DVLLM_CPP_CUDA=ON \
+  -DVLLM_CPP_CUDA_ARCHITECTURES="$ARCH" \
+  -DVLLM_CPP_CUTLASS_DIR="$CUTLASS" > "$CFGLOG" 2>&1
+CFG=$?
+cat "$CFGLOG" >>"$LOG"
+echo "### configure -> exit $CFG" | tee -a "$LOG"
+[ "$CFG" -ne 0 ] && { say "VOID: configure failed"; exit 3; }
+
+# ── PRECONDITION 3: the arch the configure log actually RESOLVED ─────────────
+# A `[121]` instead of `[121a]`, or a DISABLED marlin cell, VOIDS the run rather
+# than failing it: the result would describe a different kernel set than the one
+# this row is about. Both are spec §4.5 conditions.
+grep -E "CUDA feature .*(ENABLED|DISABLED)" "$CFGLOG" | tee -a "$LOG"
+if ! grep -qE "marlin-nvfp4: ENABLED for \[[^]]*$ARCH" "$CFGLOG"; then
+  say "VOID: marlin-nvfp4 is not ENABLED for [$ARCH]; there is no NVFP4 GEMM to gate"
+  grep -iE 'marlin' "$CFGLOG" | tee -a "$LOG"
+  exit 3
+fi
+
+say "build (-j 4)"
+cmake --build "$BUILD" -j 4 --target test_nemotron_h_moe_device nemotron-h-gen >>"$LOG" 2>&1
+B=$?
+echo "### build -> exit $B" | tee -a "$LOG"
+[ "$B" -ne 0 ] && { say "BUILD FAILED — this IS a code result, not a void"; exit 1; }
+
+# ── LEG 1: the synthetic device gate ─────────────────────────────────────────
+# `-tc` filters are NOT used: doctest splits them on commas, and a comma in a
+# case name yields `0 cases ran` + `SUCCESS!`. The whole TU runs, and the case
+# count is asserted below.
+run "LEG1 synthetic device MoE + lm_head" "$BUILD/tests/test_nemotron_h_moe_device" -s
+
+# A run that executed NO cases prints SUCCESS and exits 0. Assert a non-zero
+# case count from the summary line rather than trusting the exit status.
+if ! grep -qE 'test cases:[[:space:]]+[1-9]' "$LOG"; then
+  say "VOID: the device TU reported no executed test cases"
+  rc=1
+fi
+grep -E 'test cases:|assertions:|Status:' "$LOG" | tail -6 | tee -a "$LOG"
+
+# ── LEG 2: the REAL checkpoint, through the PRODUCTION ABI ───────────────────
+# This is the leg that enters through the production entry point:
+# `vllm_engine_load` -> the paged runner -> ForwardNemotronHForCausalLM ->
+# NemotronHPagedForward -> the device lm_head branch. A unit test that calls
+# NemotronHDeviceLmHead directly proves the projection; only this proves the
+# capability is REACHED.
+if [ -d "$CKPT" ]; then
+  run "LEG2 real-checkpoint token identity (device lm_head)" \
+    "$BUILD/examples/nemotron-h-gen" \
+      --model "$CKPT" \
+      --golden "$SRC/tests/parity/goldens/nemotron_35_lightning_greedy/oracle.json" \
+      --steps 8 --prompts 3 --max-model-len 2048
+else
+  say "SKIP LEG2: no checkpoint at $CKPT — this is a SKIP with a named reason, never a pass"
+  rc=1
+fi
+
+say "A2Q2B GATE overall exit: $rc  log: $LOG"
+exit "$rc"
diff --git a/scripts/runner-routing-allowlist.txt b/scripts/runner-routing-allowlist.txt
index c25a2b799..03a9f4743 100644
--- a/scripts/runner-routing-allowlist.txt
+++ b/scripts/runner-routing-allowlist.txt
@@ -23,4 +23,4 @@
 # --- KNOWN OFF-FRAMEWORK: private resident/graph decode, pending framework-routing ---
 laguna        # LagunaModel::ForwardDevice is a HOST stub (returns HostLogits); resident/graph decode fully off-framework (private f32-host-vector KV, bespoke DecodeAttnGqaKernel, host RoPE, host logit download) -> pending framework-routing (see AGENTS.md decode/runtime seam)
 qwen3_vl      # ForwardQwen3VLForConditionalGeneration returns HostLogits + private VLGenerateCore host ArgMax loop (MM-ENGINE-FORWARD partial seam) bypasses the runner -> pending framework-routing (see AGENTS.md decode/runtime seam)
-nemotron_h    # NARROWED AGAIN by A2-P (#810, .agents/specs/nemotron-h-a2p-paged-forward.md §3.5). What this entry USED to say -- "the forward is the HOST reference; it holds K and V for the whole prompt and pages nothing" -- is no longer true. ForwardNemotronHForCausalLM now selects NemotronHPagedForward whenever the runner supplies paged KV and recurrent state: K/V go into the runner's pages at attn_meta.slot_mapping, attention reads back out of them, and the conv/SSM rows are gathered from and scattered to the runner's GdnStateCache at the metadata's state indices. Of the checker's three clauses, (b) no private host generate loop and (c) bf16-resident activations are both satisfied, and EXACTLY ONE is not: (a) on-device logits. `lm_head` is NVFP4 W4A16 g16 on the released checkpoint and refuses on a non-CPU queue (nemotron_h.cpp:1031-1034), so the forward's last step is still the HOST projection and it returns HostLogits. A2-Q2b ports the device NVFP4 lm_head arm and A2-Q2b is what removes this entry. Deleting it while the forward still returns HostLogits reds the checker, and widening the allowlist to satisfy the checker is the defect the checker exists to stop -> pending A2-Q2b (see AGENTS.md decode/runtime seam)
+nemotron_h    # NARROWED AGAIN by A2-Q2b (#810, .agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md). The clause this entry existed for IS NOW MET: `lm_head` runs on the device through dense_nvfp4::MatmulNvfp4W4A16D and NemotronHPagedForward returns a device-resident ForwardLogits (fl.device_tensor / fl.device_storage = dlogits.ReleaseShared()) on the NVFP4 W4A16 checkpoint, which is every released one. The entry stays for ONE reason and it is an INSTRUMENT limit, not a routing one: check-runner-routing-consistency.py resolves a delegate across translation units only for the `Class::ForwardDevice` shape (_DELEGATE, :142) or a helper defined in the registry TU itself (classify_with_helpers, :246). NemotronH's device forward is a FREE function in another TU (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 -- classifies the model HOST. The checker already builds the cross-TU `free_fn_file` map it would need (used by invariant (b)), so this is a resolution hole rather than a design limit. NOT fixed in this flow: it CHANGES CHECKER SEMANTICS, which AGENTS.md `## Changing the rules or a checker` routes to the normal row, spec and fresh-review path with its own red-before. Tracked as #1410, owned by this row, listed under `## 5. Owed` in the A2-Q2b spec. REACHABILITY IS NOT PROVEN BY THIS ENTRY AND IS NOT YET PROVEN AT ALL: the deletion mutation reachability.md asks for needs CUDA, because every case that can enter the device branch is gated on `dense_nvfp4::MarlinW4A16Selects`, which is false on a CPU queue. It is recorded PENDING a `dgx:gpu0` window under `## 6. Now` in the spec, not claimed. Deleting the device branch in nemotron_h_device.cpp leaves THIS checker at rc=0 with byte-identical output (measured), so nothing here holds that branch in either direction; #1410 is what would make the checker able to see it. A seam checker never proved reachability and was never meant to.
diff --git a/src/vllm/model_executor/models/nemotron_h_device.cpp b/src/vllm/model_executor/models/nemotron_h_device.cpp
index 4455998d9..f02d86f53 100644
--- a/src/vllm/model_executor/models/nemotron_h_device.cpp
+++ b/src/vllm/model_executor/models/nemotron_h_device.cpp
@@ -110,15 +110,26 @@
 
 // DSR-ALLOW(A2-Q2a): TYPES, not behaviour -- vt/cuda/marlin_repack.h is BUILT ONLY under VT_MARLIN_NVFP4 (its own header note), so this include cannot be resolved by a runtime query. This is the platform leg for the CUDA Marlin repack.
 #ifdef VT_MARLIN_NVFP4
-// A2-Q2a: the load-time NVFP4 -> Marlin repack primitives. DELIBERATELY NOT
-// `dense_nvfp4_gemm.h`: that header's `MarlinDenseResidentFor` (:379) keys its
-// repack cache on the WEIGHT'S ADDRESS, which is issue #984, and including it
-// here would put the unsafe accessor one unqualified call away from a reviewer's
-// eye. These four functions are the same primitives it and qwen3_5.cpp both
-// drive, with no cache attached.
+// A2-Q2a: the load-time NVFP4 -> Marlin repack primitives, which the MoE arena
+// drives directly with no cache attached.
 #include "vt/cuda/marlin_repack.h"
 #endif
 
+// A2-Q2b: the SHARED NVFP4 W4A16 dense route, for `lm_head`.
+//
+// ★ THIS INCLUDE'S OWN HISTORY IS THE REVIEW NOTE. A2-Q2a deliberately did NOT
+// include this header, and said so here, because `MarlinDenseResidentFor`
+// (dense_nvfp4_gemm.h:379) keys its repack cache on the WEIGHT'S ADDRESS —
+// issue #984 — and including it put that unsafe accessor one unqualified call
+// away from a reviewer's eye. That reason was correct and it has been ANSWERED
+// rather than overruled: the seam now accepts a CALLER-OWNED resident, and the
+// only entry point this file calls (`dense_nvfp4::MatmulNvfp4W4A16D`) is called
+// with one, out of the `lm_head_marlin` slot the weights own. Nothing in this
+// translation unit calls `MarlinDenseResidentFor`, and a reviewer should check
+// that claim rather than take it — it is one grep, and it is the whole
+// difference between inheriting #984 and declining it.
+#include "vllm/model_executor/models/dense_nvfp4_gemm.h"
+
 namespace vllm {
 namespace {
 
@@ -833,6 +844,203 @@ std::vector<float> NemotronHMoeBlockDeviceHostIO(const NemotronHMoeWeights& w,
 #endif
 }
 
+
+// ═══ A2-Q2b: `lm_head` on the DEVICE, through the shared NVFP4 dense route ═══
+//
+// .agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md, issue #810.
+//
+// ── WHY THIS IS THE LAST HOST RE-EXPANSION, MEASURED ────────────────────────
+//
+// `lm_head` is [131072, 2688] NVFP4 W4A16 g16. On the HOST arm every decode
+// step widens it whole through `NemotronHOwned::DenseBf16` — 352 321 536
+// elements, a 704.6 MB transient bf16 buffer, for ONE GEMM against ONE row.
+//
+// Counted at the single dequant seam on the real checkpoint, one decode step
+// (T=1, top_k=6, 23 MoE layers), the host arm re-expands 3 078 512 640
+// elements. `lm_head` is 11.44% of that, and it is 28.35% of the 1 242 587 136
+// that remain once A2-Q2a moved the MoE arm to the device. It is 352 321 536
+// of 352 321 536 — everything left — once the mamba arm follows. It is also
+// the largest SINGLE re-expansion in the model by 12.7x: 352.3e6 elements in
+// one call, against 27.7e6 for the next largest (mamba `in_proj`). On a
+// unified-memory box that transient is the allocation that matters.
+//
+// ── THE RESIDENCY DECISION (spec `## 5. Owed`), TAKEN EXPLICITLY ────────────
+//
+// The spec's `## 5. Owed` required this row to CHOOSE rather than default —
+// "A2-Q2a routed around it by never calling either `MarlinDenseResidentFor`;
+// `lm_head` must do the same or say why not" — because that accessor keys its
+// repack cache on the weight's ADDRESS (issue #984) and NemotronH is exactly
+// the second-consumer condition an address key cannot survive.
+//
+// Chosen: route through the shared seam `dense_nvfp4::MatmulNvfp4W4A16D`, and
+// hand it a resident this model OWNS, in the `lm_head_marlin` `ResidentSlot`
+// on the weights. That is the property `qwen3_5.cpp` gets by hand-rolling its
+// own Marlin path, reached WITHOUT hand-rolling one, because
+// the seam was extended to accept a caller-owned resident rather than forked.
+// A2-Q2a made the same call for the MoE arena; this keeps the two arms
+// consistent, and it leaves #984 exactly as it was for every other caller
+// (which is where its own two-engine red-before belongs).
+
+// A transient `Nvfp4Weight` VIEW over the host `lm_head` payload.
+//
+// UNGUARDED, unlike the arena helpers above, because it names nothing that the
+// Marlin build adds: `Nvfp4Weight` comes from qwen3_5_weights.h and
+// `OwnedBytes` from the loader, both unconditional, and no `vt::cuda::` symbol
+// appears below. A `#ifdef` here would have been a device-specific reference
+// paying for nothing — the one thing `scripts/check-device-leakage.py` counts.
+// It has external linkage at `namespace vllm` scope, so a build without
+// VT_MARLIN_NVFP4, where its single call site below is compiled out, emits no
+// unused-function diagnostic.
+//
+// It BORROWS: `Nvfp4Weight` wants `OwnedTensor`s and `NemotronHOwned` already
+// holds the packed codes and the group scales in exactly the layout
+// `BuildMarlinDenseResident` reads ([N, K/2] and [N, K/16], torch [out, in]).
+// Copying would add 373 MB of host residency to a load whose report is pinned
+// to the byte (`rep.host_bytes == 18888922112`,
+// test_nemotron_h_loader.cpp:310) and whose whole design is that a quantized
+// weight KEEPS its quantized memory format.
+//
+// The keep-alive is a no-op deleter, and that is correct rather than lazy:
+// `OwnedBytes::Borrow` rejects a null owner because a borrowed view must not
+// outlive its bytes, and here the bytes are the caller's `NemotronHHostWeights`,
+// which outlive this stack temporary by construction — the value is consumed by
+// the repack inside this same call and never stored.
+Nvfp4Weight LmHeadNvfp4View(const NemotronHOwned& w, int64_t V, int64_t H) {
+  VT_CHECK(w.form == NemotronHWeightForm::kNvfp4W4A16G16,
+           "NemotronH device lm_head: this arm is NVFP4 W4A16 g16 only");
+  VT_CHECK(w.shape.size() == 2 && w.shape[0] == V && w.shape[1] == H,
+           "NemotronH device lm_head: weight shape is not [vocab_size, hidden_size]");
+  VT_CHECK(w.bytes.size() == static_cast<size_t>(V) * static_cast<size_t>(H) / 2,
+           "NemotronH device lm_head: packed payload is not [vocab_size, hidden_size/2]");
+  VT_CHECK(w.scale.size() == static_cast<size_t>(V) * static_cast<size_t>(H) / 16,
+           "NemotronH device lm_head: group scales are not [vocab_size, hidden_size/16]");
+  // A keep-alive that owns nothing: see the note above.
+  static const std::shared_ptr<const void> kNoopOwner(reinterpret_cast<const void*>(1),
+                                                      [](const void*) {});
+  Nvfp4Weight nw;
+  nw.packed.bytes = OwnedBytes::Borrow(w.bytes.data(), w.bytes.size(), kNoopOwner);
+  nw.packed.dtype = DType::kI8;
+  nw.packed.rank = 2;
+  nw.packed.shape[0] = V;
+  nw.packed.shape[1] = H / 2;
+  nw.scale.bytes = OwnedBytes::Borrow(w.scale.data(), w.scale.size(), kNoopOwner);
+  nw.scale.dtype = DType::kI8;
+  nw.scale.rank = 2;
+  nw.scale.shape[0] = V;
+  nw.scale.shape[1] = H / 16;
+  nw.scale2 = w.global_scale;
+  nw.n = V;
+  nw.k = H;
+  nw.group_size = 16;
+  nw.is_mxfp4 = false;
+  // W4A16, never true-W4A4: the released checkpoint quantizes no activation, so
+  // `alpha` stays 0 and `IsTrueW4A4()` is false. The seam asserts this too; the
+  // fields are left at their defaults deliberately rather than set to 0 twice.
+  return nw;
+}
+
+// True when the DEVICE `lm_head` arm can serve this weight on this queue.
+//
+// TWO clauses, and the split between them is the point:
+//   * `dense_nvfp4::MarlinW4A16Selects` is the shared dispatcher's OWN gate,
+//     called rather than restated. It carries VT_MARLIN_NVFP4, the op table's
+//     realization for THIS device (an availability question, not a `== kCUDA`
+//     one), `MarlinW4A16Enabled()` (the VT_NVFP4_MARLIN A/B escape hatch), and
+//     Marlin's bf16 a/c contract.
+//   * the weight is actually NVFP4 — the one clause that is THIS model's and
+//     not the dispatcher's (a dense `lm_head` is a different arm, and a
+//     synthetic fixture ships one).
+//
+// ★ IT WAS A RESTATEMENT AND THE RESTATEMENT WAS WRONG. The first submission
+// listed the dispatcher's clauses by hand and dropped `MarlinW4A16Enabled()`,
+// so under an explicit `VT_NVFP4_MARLIN=0` this predicate said eligible, the
+// seam fell to the naive redundant-dequant arm, and `LmHeadNvfp4View`'s
+// transient weight made `ResidentNvfp4`'s cache miss every time — re-uploading
+// the whole [131072, 2688] operand on EVERY decode step. The value is
+// unchanged, so a token gate sees nothing. Calling the seam's own predicate is
+// what makes the two unable to disagree.
+//
+// Anything else falls through to `NemotronHHostLmHead`, which refuses BY NAME
+// on a non-CPU queue rather than computing on the wrong operand.
+bool DeviceLmHeadEligible(Dev d, const NemotronHHostWeights& host, DType adt) {
+  return host.lm_head.form == NemotronHWeightForm::kNvfp4W4A16G16 &&
+         dense_nvfp4::MarlinW4A16Selects(d, adt);
+}
+
+// A device-side row gather, the same helper every other model's logits path
+// carries verbatim (opt.cpp:229, qwen3.cpp, gemma3.cpp:283). It exists here
+// because the host arm's gather was a `std::memcpy` out of a DOWNLOADED f32
+// copy of `final_normed`; the device arm must never make that download, which
+// is the whole point of moving the projection.
+void GatherRowsD(Dev d, void* dst, const Tensor& src, const std::vector<int64_t>& idx,
+                 int64_t row_elems) {
+  const size_t rb = static_cast<size_t>(row_elems) * vt::SizeOf(src.dtype);
+  auto* dp = static_cast<char*>(dst);
+  const auto* sp = static_cast<const char*>(src.data);
+  for (size_t s = 0; s < idx.size(); ++s)
+    d.b.Copy(d.q, dp + s * rb, sp + static_cast<size_t>(idx[s]) * rb, rb);
+}
+
+// y[R, vocab] f32 = gathered[R, hidden] bf16 @ dequant(lm_head).T, on the device.
+//
+// `out_dtype` is f32 and that is the CARRIER, not a widened compute: Marlin's
+// output is bf16 (`c_type=kBFloat16`) and the seam upcasts it, so this is the
+// same value the bf16 GEMM produced. f32 is what `ForwardLogits` and the
+// on-device sampler take, exactly as every other model's device logits do
+// (opt.cpp:308, minicpm3.cpp:303).
+DBuf DeviceLmHeadD(Dev d, const NemotronHHostWeights& host,
+                   const NemotronHParams& params, const Tensor& gathered) {
+// DSR-ALLOW(A2-Q2b): TYPES, not behaviour -- this body names two symbols that do not EXIST without the guarded arena region above: the `ResidentIn` template, which is defined INSIDE that region, and `dense_nvfp4::MarlinDenseResident`, which dense_nvfp4_gemm.h DECLARES unconditionally but DEFINES only under VT_MARLIN_NVFP4, so the reference `ResidentIn` returns cannot bind to an incomplete type. The SELECTION is already a runtime op-table query: `DeviceLmHeadEligible` calls `dense_nvfp4::MarlinW4A16Selects`, which carries NO guard, and only this call site needs one. It has an #else that refuses BY NAME, so a build without Marlin reports the missing arm rather than silently computing on the host. Same shape and same reason as the `NemotronHMoeBlockDeviceD` body and the `moe_on_device` dispatch branches.
+#ifdef VT_MARLIN_NVFP4
+  const int64_t V = params.vocab_size;
+  const int64_t H = params.hidden_size;
+  VT_CHECK(gathered.rank == 2 && gathered.shape[1] == H,
+           "NemotronH device lm_head: gathered rows are not [R, hidden_size]");
+  VT_CHECK(gathered.shape[0] > 0, "NemotronH device lm_head: no rows requested");
+  const Nvfp4Weight nw = LmHeadNvfp4View(host.lm_head, V, H);
+  // SLOT-KEYED, never the header's address-keyed static (#984, and the `## 5.
+  // Owed` bullet in the spec that asks this arm to route around it or say why
+  // not).
+  dense_nvfp4::MarlinDenseResident& mr =
+      ResidentIn<dense_nvfp4::MarlinDenseResident>(host.lm_head_marlin);
+  // ★ A FALLBACK HERE IS A DEFECT, NOT A SLOWER ANSWER, so it refuses by name.
+  //
+  // `DeviceLmHeadEligible` asked `MarlinW4A16Selects` and was told yes. Asking
+  // the SAME predicate again on the operand actually about to be handed over is
+  // what catches an eligibility answer taken against a different dtype or a
+  // different queue. If it were ever false here the dispatcher would take its
+  // naive redundant-dequant arm, and the consequence is invisible to every
+  // value-based gate: that arm computes the SAME logits, while re-uploading
+  // 198.18 MB — the whole [131072, 2688] packed operand plus its group scales —
+  // on EVERY call, because `LmHeadNvfp4View` hands out a stack temporary and
+  // `ResidentNvfp4` caches on `w.d_packed`, a member of the weight it was given.
+  //
+  // The predicate and not the seam's `fallback_gemms` counter, deliberately.
+  // `MutableW4A16Stats()` is a plain non-atomic static shared by every consumer
+  // in the process, so a counter window would refuse a correct run whenever
+  // anything else took a fallback GEMM concurrently — a false refusal is worse
+  // than the silence it replaces. The counter is the right instrument in a test,
+  // where single-threadedness is a property of the harness, and that is where
+  // `tests/vllm/models/test_nemotron_h_moe_device.cpp` asserts it.
+  VT_CHECK(dense_nvfp4::MarlinW4A16Selects(d, gathered.dtype),
+           "NemotronH device lm_head: the shared NVFP4 W4A16 dispatcher would take "
+           "its redundant-dequant fallback on this queue and dtype although the "
+           "device arm was selected — the fallback re-uploads the whole lm_head on "
+           "every call, so this refuses rather than computing the right answer the "
+           "wrong way");
+  return dense_nvfp4::MatmulNvfp4W4A16D(d, gathered, nw, DType::kF32, &mr);
+#else
+  (void)d;
+  (void)host;
+  (void)params;
+  (void)gathered;
+  VT_CHECK(false,
+           "NemotronH device lm_head: this build has no Marlin NVFP4 GEMM "
+           "(VT_MARLIN_NVFP4 is off), so the device lm_head arm is not compiled in");
+  return DBuf(d, DType::kF32, {1, 1});
+#endif
+}
+
 // ─── the hybrid forward ─────────────────────────────────────────────────────
 
 std::vector<float> NemotronHDeviceForward(const NemotronHHostWeights& host,
@@ -1001,11 +1209,25 @@ std::vector<float> NemotronHDeviceForward(const NemotronHHostWeights& host,
   const std::vector<float> fvec = DownloadF32(d, final_normed, adt, T * H);
   if (trace != nullptr && trace->capture) trace->final_normed = fvec;
 
-  // --- lm_head, on the HOST: it is NVFP4 W4A16 g16 on the released
-  // checkpoint. Both arms therefore end in the IDENTICAL host projection, which
-  // is what makes A2-R's token gate attributable: a token difference can only
-  // have come from the 6 device attention blocks and the device residual
-  // stream, never from the output projection.
+  // --- lm_head, on the HOST, and A2-Q2b LEAVES IT THERE DELIBERATELY.
+  //
+  // A2-Q2b moves `lm_head` to the device in `NemotronHPagedForward` — the
+  // production path — and not here. Three reasons, and the first is the one
+  // that matters:
+  //
+  //   * THIS SEAM IS THE GATE'S OPERAND. A2-Q2b's numeric gate compares the
+  //     device projection against the host projection on the SAME gathered
+  //     rows. Moving both arms to the device deletes the reference the
+  //     comparison is against, which is the same mistake A2-Q2a's §4.1 refused
+  //     when it supplemented the weights instead of replacing them.
+  //   * A2-R's token gate is attributable BECAUSE both arms end in the
+  //     identical host projection: a token difference can only have come from
+  //     the 6 device attention blocks and the device residual stream. That
+  //     property is worth keeping on the seam that still has it. The spec
+  //     required this row to DISCLOSE that the property ends where `lm_head`
+  //     moves — it ends in the paged forward, and it survives here.
+  //   * `NemotronHDeviceForward` has no production caller at all (:651), so
+  //     leaving it costs no user-visible behaviour.
   std::vector<int64_t> want;
   if (logits_indices.empty()) {
     want.resize(static_cast<size_t>(T));
@@ -1752,13 +1974,27 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host,
     Tensor ot = final_normed.t();
     AddRmsNorm(d, ot, xt, wt, rt, nargs, params.layer_norm_epsilon);
   }
-  const std::vector<float> fvec = DownloadF32(d, final_normed, adt, T * H);
-  if (trace != nullptr && trace->capture) trace->final_normed = fvec;
+  // ── ONE download of `final_normed`, whichever arm and whichever trace ──────
+  //
+  // Two consumers want these bytes and NEITHER always runs: the trace (the
+  // numeric gate's operand, and nothing else reads it) and the HOST projection
+  // below. Before A2-Q2b the download was unconditional because the host
+  // projection always needed it. Making it conditional is the point of this row
+  // — the device arm must not pay a T*H device-to-host copy per step — but the
+  // first submission wrote the condition as an EXTRA download rather than as a
+  // shared one, so a traced host step copied the same buffer twice.
+  //
+  // `have_fvec` rather than `fvec.empty()`: emptiness is a property of T*H and
+  // would silently become "download again" if a caller ever asked for zero rows.
+  std::vector<float> fvec;
+  bool have_fvec = false;
+  if (trace != nullptr && trace->capture) {
+    fvec = DownloadF32(d, final_normed, adt, T * H);
+    trace->final_normed = fvec;
+    have_fvec = true;
+  }
 
-  // The gather-before-lm_head rows, then the HOST projection. `lm_head` is
-  // NVFP4 W4A16 g16 and its device arm is A2-Q2b's, so this forward returns
-  // HOST logits and `scripts/runner-routing-allowlist.txt` is NARROWED rather
-  // than removed (spec §3.5). An EMPTY `logits_indices` is the runner's
+  // The gather-before-lm_head rows. An EMPTY `logits_indices` is the runner's
   // VT_LOGITS_GATHER=0 path and means "every row", which is also what the two
   // non-paged seams mean by it — so this branch serves both gather settings and
   // no runner step can escape the paged path on that flag.
@@ -1773,15 +2009,116 @@ ForwardLogits NemotronHPagedForward(const NemotronHHostWeights& host,
       want.push_back(idx);
     }
   }
+  // `n_out`, not `R`: `R` is already this function's REQUEST count (:1673) and
+  // these are the gathered LOGIT rows, which are not the same number.
+  const int64_t n_out = static_cast<int64_t>(want.size());
+
+  // ── A2-Q2b: THE DEVICE PROJECTION, AND THE PRODUCTION CALL SITE ───────────
+  //
+  // This is the line `scripts/runner-routing-allowlist.txt` was NARROWED for —
+  // narrowed, not removed; the entry is still in that file and its text says
+  // why.
+  //
+  // ★ NOTHING AUTOMATED HOLDS THIS BRANCH IN PLACE, AND THAT IS MEASURED.
+  // Delete this whole `if` block and `scripts/check-runner-routing-consistency.py`
+  // still exits 0 with byte-identical output ("3 host-logits off-framework
+  // (3 allowlisted)"), because the allowlist entry that would have to be
+  // removed first is exactly what keeps the model off the checker's red list —
+  // and the entry cannot be removed while #1410 makes the checker misclassify
+  // a cross-TU free-function device forward as HOST. So the checker is not a
+  // guard here in either direction. An earlier draft of this comment claimed
+  // the opposite; it was wrong, and it was wrong in the direction that stops a
+  // reader looking for the real evidence.
+  //
+  // What DOES hold it is the reachability deletion mutation
+  // (`.agents/reachability.md`): delete this block, run the device `lm_head`
+  // gate, and it must go red. That mutation needs CUDA — every case that can
+  // reach this branch is `TryCudaQueue`-gated, because `MarlinW4A16Selects` is
+  // false on a CPU queue — so it is recorded PENDING a `dgx:gpu0` window in the
+  // spec's `## 6. Now`, not claimed. #1410 owns making the checker able to see
+  // this at all.
+  if (DeviceLmHeadEligible(d, host, adt)) {
+    DBuf grows(d, adt, {n_out, H});
+    GatherRowsD(d, grows.ptr(), final_normed.t(), want, H);
+    DBuf dlogits = DeviceLmHeadD(d, host, params, grows.t());
+    ForwardLogits fl;
+    fl.rows = n_out;
+    fl.vocab = params.vocab_size;
+    fl.device_tensor = dlogits.t();
+    // The pool block's lifetime moves into a shared_ptr whose deleter returns
+    // it to the DevicePool — the shared `WrapDeviceLogits` carrier every
+    // device-logits model uses (opt.cpp:315, minicpm3.cpp:303).
+    fl.device_storage = dlogits.ReleaseShared();
+    return fl;
+  }
+
+  // The HOST projection, kept below the fold exactly as the host reference
+  // forward is: it is the operand A2-Q2b's numeric gate compares against, and
+  // it is what serves a build with no Marlin NVFP4 GEMM and a checkpoint whose
+  // `lm_head` is dense. It refuses BY NAME on a non-CPU queue.
+  if (!have_fvec) fvec = DownloadF32(d, final_normed, adt, T * H);
   std::vector<float> gathered(want.size() * static_cast<size_t>(H));
   for (size_t r = 0; r < want.size(); ++r) {
     std::memcpy(gathered.data() + r * static_cast<size_t>(H),
                 fvec.data() + static_cast<size_t>(want[r]) * static_cast<size_t>(H),
                 static_cast<size_t>(H) * sizeof(float));
   }
-  return HostLogits(NemotronHHostLmHead(host, params, gathered,
-                                        static_cast<int64_t>(want.size()), hq),
+  return HostLogits(NemotronHHostLmHead(host, params, gathered, n_out, hq),
                     params.vocab_size);
 }
 
+// ─── A2-Q2b: the PUBLIC per-block entry points for `lm_head` ────────────────
+//
+// The twins of `NemotronHHostLmHead`, argument for argument, so the numeric
+// gate compares like with like instead of against a second implementation of
+// the same projection written inside the test. That symmetry is the whole
+// point: this file's header says the per-block entry points "are public so the
+// gate can compare a BLOCK's activations against an independent reference".
+//
+// THEY ARE NOT THE REACHABILITY PROOF, and nothing here should be read as one.
+// Production reaches the device projection through
+// `ModelRegistry::Forward` -> `ForwardNemotronHForCausalLM` ->
+// `NemotronHPagedForward`, at the `DeviceLmHeadEligible` branch. These two
+// localize a failure to the projection; the real-checkpoint gate is what
+// enters through the production path (reachability.md, "Does a test enter
+// through it?").
+bool NemotronHDeviceLmHeadEligible(const NemotronHHostWeights& host, vt::DType adt,
+                                   vt::Queue& dev_queue) {
+  Dev d{vt::GetBackend(dev_queue.device.type), dev_queue};
+  return DeviceLmHeadEligible(d, host, adt);
+}
+
+std::vector<float> NemotronHDeviceLmHead(const NemotronHHostWeights& host,
+                                         const NemotronHParams& params,
+                                         const std::vector<float>& gathered_normed,
+                                         int64_t num_rows, vt::Queue& dev_queue) {
+  const DType adt = host.act_dtype;
+  const int64_t H = params.hidden_size;
+  const int64_t V = params.vocab_size;
+  VT_CHECK(num_rows > 0, "NemotronH device lm_head: no rows requested");
+  VT_CHECK(static_cast<int64_t>(gathered_normed.size()) == num_rows * H,
+           "NemotronH device lm_head: gathered row count does not match hidden_size");
+  Dev d{vt::GetBackend(dev_queue.device.type), dev_queue};
+  VT_CHECK(DeviceLmHeadEligible(d, host, adt),
+           "NemotronH device lm_head: this arm needs an NVFP4 W4A16 g16 `lm_head`, a "
+           "bf16 activation dtype, and a device whose op table realizes the Marlin "
+           "NVFP4 grouped GEMM — the host arm is NemotronHHostLmHead");
+  // Upload the gathered rows at the MODEL dtype, not f32: the operand Marlin
+  // takes is bf16, and widening here would move twice the bytes for a value the
+  // kernel narrows again on the way in.
+  DBuf rows(d, adt, {num_rows, H});
+  {
+    // PackF32 lives in nemotron_h.cpp; the conversion is one loop and doing it
+    // here avoids widening this file's dependency surface for two call sites.
+    std::vector<uint16_t> bf(static_cast<size_t>(num_rows * H));
+    for (size_t i = 0; i < bf.size(); ++i) bf[i] = vt::F32ToBF16(gathered_normed[i]);
+    d.b.Copy(d.q, rows.ptr(), bf.data(), bf.size() * sizeof(uint16_t));
+    d.b.Synchronize(d.q);
+  }
+  DBuf dlogits = DeviceLmHeadD(d, host, params, rows.t());
+  std::vector<float> out(static_cast<size_t>(num_rows * V));
+  dlogits.Download(d, out.data());
+  return out;
+}
+
 }  // namespace vllm
diff --git a/src/vllm/model_executor/models/nemotron_h_forward.h b/src/vllm/model_executor/models/nemotron_h_forward.h
index 25360fa43..80b9192ec 100644
--- a/src/vllm/model_executor/models/nemotron_h_forward.h
+++ b/src/vllm/model_executor/models/nemotron_h_forward.h
@@ -360,6 +360,22 @@ struct NemotronHHostWeights {
   // meaningful — both arms end in the identical host projection, so any token
   // difference is attributable to the 6 device attention blocks.
   NemotronHOwned lm_head;  // [vocab_size, hidden_size]
+  // ── A2-Q2b (#810): the DEVICE lm_head's Marlin resident ────────────────────
+  //
+  // The repacked NVFP4 operands for `lm_head`, built once on first device
+  // forward and reused on every step. It lives HERE, on the weights, for the
+  // reason issue #237 added `ResidentSlot` at all: state keyed on the ADDRESS
+  // of a weight outlives the engine that built it, and a second engine in the
+  // same process can inherit device pointers the first one freed. That is
+  // issue #984, which the shared `dense_nvfp4_gemm.h` cache still has; A2-Q2a
+  // declined to inherit it for the MoE arena (`moe_marlin` above) and this
+  // slot is the same decision for `lm_head`, taken explicitly per that spec's
+  // §4.3 rather than by default.
+  //
+  // Opaque `shared_ptr<void>`: the resident type is a CUDA-path detail of
+  // nemotron_h_device.cpp, and holding it here would drag Marlin's headers
+  // into every consumer of this one.
+  ResidentSlot lm_head_marlin;
   // False until a loader materializes the enumerated tensors. The forward
   // refuses by name on false rather than computing on zeros.
   bool materialized = false;
@@ -526,6 +542,18 @@ std::vector<float> NemotronHMoeBlockDeviceHostIO(const NemotronHMoeWeights& w,
 // the identical projection, a token difference between them is attributable to
 // the device attention blocks and the device residual stream alone. A second
 // copy of this projection would quietly destroy that property.
+// A2-Q2b: the DEVICE twins of `NemotronHHostLmHead` below, argument for
+// argument. `...Eligible` is the same predicate the production paged forward
+// branches on, exposed so a gate can assert it DISCRIMINATES rather than
+// assuming it; a predicate that is constantly false disables the whole arm and
+// still passes every "the host arm was used" assertion.
+bool NemotronHDeviceLmHeadEligible(const NemotronHHostWeights& host, vt::DType adt,
+                                   vt::Queue& dev_queue);
+std::vector<float> NemotronHDeviceLmHead(const NemotronHHostWeights& host,
+                                         const NemotronHParams& params,
+                                         const std::vector<float>& gathered_normed,
+                                         int64_t num_rows, vt::Queue& dev_queue);
+
 std::vector<float> NemotronHHostLmHead(const NemotronHHostWeights& host,
                                        const NemotronHParams& params,
                                        const std::vector<float>& gathered_normed,
diff --git a/tests/vllm/models/test_nemotron_h_moe_device.cpp b/tests/vllm/models/test_nemotron_h_moe_device.cpp
index 181f070cb..7912de375 100644
--- a/tests/vllm/models/test_nemotron_h_moe_device.cpp
+++ b/tests/vllm/models/test_nemotron_h_moe_device.cpp
@@ -39,6 +39,7 @@
 #include <string>
 #include <vector>
 
+#include "vllm/model_executor/models/dense_nvfp4_gemm.h"
 #include "vllm/model_executor/models/nemotron_h_forward.h"
 #include "vt/backend.h"
 #include "vt/dtype.h"
@@ -47,6 +48,12 @@ namespace {
 
 using vllm::NemotronHBlock;
 using vllm::NemotronHExpertWeights;
+// A2-Q2b: the lm_head cases below name this type UNQUALIFIED. It was missing
+// from this block on the first submission and the file never compiled on any
+// platform, which is why `vllm_cpp_add_test(test_nemotron_h_moe_device)` is
+// registered with NO CUDA guard: the CPU build is what makes a case that
+// SKIPS at run time still have to parse and type-check.
+using vllm::NemotronHHostWeights;
 using vllm::NemotronHMoeWeights;
 using vllm::NemotronHOwned;
 using vllm::NemotronHParams;
@@ -343,3 +350,172 @@ TEST_CASE("NemotronH A2-Q2a: the device MoE arm refuses a dense expert rather th
   const std::vector<float> x = SynthVec(static_cast<size_t>(T * H), 78, 0.5F);
   CHECK_THROWS(vllm::NemotronHMoeBlockDeviceHostIO(w, p, x, T, dt, dq));
 }
+
+// ═══ A2-Q2b (#810): the DEVICE `lm_head` arm ════════════════════════════════
+//
+// The cheap local arm in front of the real-checkpoint gate, for exactly the
+// reason the preamble of this file gives for the MoE arm: three GB10 windows
+// on this row were VOID for reasons a synthetic case catches for free.
+//
+// ── THE GEOMETRY IS NOT ARBITRARY (same rule as the MoE fixture) ────────────
+// Marlin refuses a shape it has no thread config for: `is_valid_config` needs
+// `prob_k % thread_k == 0 && prob_n % thread_n == 0` over {128,128,256},
+// {64,128,128}, {128,64,128}. `lm_head` is [V, H] consumed as K=H, N=V, so
+// H=256 / V=512 resolves on {128,128,256} (256%128==0, 512%128==0). The real
+// checkpoint's H=2688 / V=131072 also resolves, on more than one config —
+// which is precisely why the real-checkpoint arm is still owed and this case
+// does not stand in for it (spec §2: "the other Marlin thread configs").
+//
+// H is also divisible by 16, so the group-scale grid is [V, H/16] exactly.
+namespace {
+
+NemotronHParams LmHeadParams() {
+  NemotronHParams p = MoeParams();
+  p.hidden_size = 256;
+  p.vocab_size = 512;
+  return p;
+}
+
+// The host weights this arm needs, and NOTHING else: `NemotronHDeviceLmHead`
+// reads `lm_head` and `lm_head_marlin`, and the host reference reads
+// `lm_head` and `act_dtype`. Building a whole model here would test the model.
+NemotronHHostWeights LmHeadWeights(const NemotronHParams& p, DType dt, bool nvfp4) {
+  NemotronHHostWeights h;
+  h.act_dtype = dt;
+  h.lm_head = nvfp4 ? MakeNvfp4(p.vocab_size, p.hidden_size, 77, dt)
+                    : OwnF32(SynthVec(static_cast<size_t>(p.vocab_size * p.hidden_size), 5, 0.2F),
+                             dt, {p.vocab_size, p.hidden_size});
+  h.materialized = true;
+  return h;
+}
+
+}  // namespace
+
+// NO COMMA IN THIS NAME (doctest `-tc` splits filters on commas -> `0 cases ran`
+// + `SUCCESS!`; the case-count assertion below is the other half of that guard).
+TEST_CASE("NemotronH A2-Q2b: the device lm_head matches the host projection on the same rows") {
+  Queue dq{Device{DeviceType::kCPU, 0}, nullptr};
+  if (!TryCudaQueue(&dq)) {
+    NoteDeviceSkip("device lm_head vs host projection");
+    return;
+  }
+  const NemotronHParams p = LmHeadParams();
+  Queue hq{Device{DeviceType::kCPU, 0}, nullptr};
+  const DType dt = DType::kBF16;  // Marlin's a/c contract (ops.cpp:879)
+  const int64_t H = p.hidden_size;
+  const int64_t V = p.vocab_size;
+  const NemotronHHostWeights host = LmHeadWeights(p, dt, /*nvfp4=*/true);
+
+  // ★ T == 1 IS THE DECODE SHAPE, and it is the one this projection spends its
+  // whole life in: `lm_head` runs once per step over the GATHERED rows, and a
+  // decode step gathers exactly one. The width loop keeps a prefill shape
+  // beside it so a config that resolves only at one M cannot hide.
+  for (int64_t R : {static_cast<int64_t>(1), static_cast<int64_t>(3)}) {
+    CAPTURE(R);
+    const std::vector<float> rows = SynthVec(static_cast<size_t>(R * H), 11, 0.5F);
+
+    // The HOST reference, through the SAME entry point the production fallback
+    // takes — not a hand-rolled dequant-and-multiply, which would be a second
+    // implementation for this case to agree with.
+    const std::vector<float> want = vllm::NemotronHHostLmHead(host, p, rows, R, hq);
+    REQUIRE(want.size() == static_cast<size_t>(R * V));
+
+    // ── ★ THE ONE DEFECT CLASS A VALUE COMPARISON CANNOT SEE ────────────────
+    //
+    // Everything below this line compares NUMBERS, and the shared dispatcher's
+    // naive redundant-dequant fallback computes the SAME numbers as Marlin. So
+    // if `DeviceLmHeadEligible` and `dense_nvfp4::MarlinW4A16Selects` ever
+    // disagree, every assertion in this case still passes while the arm
+    // re-uploads the whole [vocab, hidden] operand on every single call —
+    // `LmHeadNvfp4View` hands out a stack temporary, so `ResidentNvfp4`'s
+    // weight-keyed cache can never hit. That is the shape of the defect this
+    // row shipped in its first submission (the predicate restated the
+    // dispatcher's clauses and dropped `MarlinW4A16Enabled()`).
+    //
+    // The counter is the seam's own, and it is DEMONSTRABLY ARMED rather than
+    // assumed: `test_qwen3_forward.cpp:497` asserts on CPU that
+    // `fallback_gemms` reaches exactly `5 * num_hidden_layers` when the
+    // dispatcher does fall back. An absent hook reads exactly like an armed one.
+    vllm::dense_nvfp4::ResetW4A16Stats();
+    const std::vector<float> got = vllm::NemotronHDeviceLmHead(host, p, rows, R, dq);
+    const vllm::dense_nvfp4::Nvfp4W4A16Stats st = vllm::dense_nvfp4::GetW4A16Stats();
+    MESSAGE("R=" << R << " W4A16 counters: marlin_gemms=" << st.marlin_gemms
+                 << " fallback_gemms=" << st.fallback_gemms
+                 << " dense_gemms=" << st.dense_gemms);
+    CHECK(st.fallback_gemms == 0);
+    CHECK(st.marlin_gemms + st.dense_gemms >= 1);
+
+    // ★ THE COUNT IS ASSERTED AGAINST THE GEOMETRY, never against either
+    // buffer's own size — a buffer that agrees with itself proves nothing, and
+    // a maximum over zero elements is 0.0, which is also exactly what a
+    // bit-exact comparison prints (this file's MaxRel note, and A2-Q2a's first
+    // GB10 run, which printed `worst relative deviation: 0` from a mute loop).
+    REQUIRE(got.size() == static_cast<size_t>(R * V));
+    int64_t examined = 0;
+    const double dev = MaxRel(got, want, &examined);
+    REQUIRE(examined == R * V);
+
+    // ── THE BAND IS MEASURED IN THE CASE, AND THE GUARD IS A PROPERTY ───────
+    // (spec §3, and A2-Q2a's §5.2, both paid for on this row: a hard-coded
+    // bf16 band of 3e-2 once sat ABOVE a 2.11e-2 defect and accepted a wrong
+    // answer.)
+    //
+    // `separation` is what this comparison could possibly resolve: the worst
+    // deviation between the reference and a DELIBERATELY WRONG answer built by
+    // the same arithmetic — here the reference with its rows rotated by one,
+    // which is a real failure mode of a gathered projection (an off-by-one row
+    // gather) and not a synthetic perturbation.
+    std::vector<float> rotated(want.size());
+    for (int64_t r = 0; r < R; ++r)
+      for (int64_t c = 0; c < V; ++c)
+        rotated[static_cast<size_t>(r * V + c)] =
+            want[static_cast<size_t>(((r + 1) % R) * V + c)];
+    int64_t sep_examined = 0;
+    const double separation = MaxRel(rotated, want, &sep_examined);
+    // Agreement, separation and the guard MUST report the same count or the
+    // band between them is fiction (spec §3).
+    REQUIRE(sep_examined == examined);
+
+    if (R > 1) {
+      // `separation > 0` is REQUIRED, not assumed: `separation / 2` with a
+      // separation of 0 is a band of 0, and `<` against it rejects even exact
+      // agreement. A2-Q2a's first band collapsed on exactly this and failed on
+      // the BEST possible outcome.
+      REQUIRE(separation > 0.0);
+      MESSAGE("R=" << R << " examined=" << examined << " deviation=" << dev
+                   << " separation=" << separation << " band=" << (separation / 2.0));
+      // Strict `<`, never `<=`: `<=` admits a band of 0.
+      CHECK(dev < separation / 2.0);
+    } else {
+      // R==1 cannot be row-rotated, so there is no measured separation to build
+      // a band from and this width reports agreement WITHOUT a band rather than
+      // borrowing R=3's. Said out loud, because a case that quietly compares
+      // nothing at the decode width is the one failure this file exists to stop.
+      MESSAGE("R=1 examined=" << examined << " deviation=" << dev
+                              << " (no row-rotation separation exists at R=1)");
+      CHECK(examined == V);
+    }
+  }
+}
+
+TEST_CASE("NemotronH A2-Q2b: the device lm_head refuses a DENSE weight rather than reading garbage") {
+  Queue dq{Device{DeviceType::kCPU, 0}, nullptr};
+  if (!TryCudaQueue(&dq)) {
+    NoteDeviceSkip("device lm_head dense refusal");
+    return;
+  }
+  const NemotronHParams p = LmHeadParams();
+  // A DENSE `lm_head` is a different arm. The eligibility predicate must route
+  // it to the host projection rather than handing dense bf16 elements to a
+  // kernel that reads them as packed nibbles — plausible garbage, finite, and
+  // completely invisible to a token comparison.
+  const NemotronHHostWeights host = LmHeadWeights(p, DType::kBF16, /*nvfp4=*/false);
+  CHECK_FALSE(vllm::NemotronHDeviceLmHeadEligible(host, DType::kBF16, dq));
+  // ...and the NVFP4 one IS eligible on the same queue, so the predicate is
+  // discriminating rather than constantly false (which would pass the line
+  // above while disabling the whole arm).
+  const NemotronHHostWeights q = LmHeadWeights(p, DType::kBF16, /*nvfp4=*/true);
+  CHECK(vllm::NemotronHDeviceLmHeadEligible(q, DType::kBF16, dq));
+  // An f32 activation is not Marlin's contract either.
+  CHECK_FALSE(vllm::NemotronHDeviceLmHeadEligible(q, DType::kF32, dq));
+}
diff --git a/tests/vllm/models/test_nemotron_h_paged_forward.cpp b/tests/vllm/models/test_nemotron_h_paged_forward.cpp
index ca07832f4..d203f58bc 100644
--- a/tests/vllm/models/test_nemotron_h_paged_forward.cpp
+++ b/tests/vllm/models/test_nemotron_h_paged_forward.cpp
@@ -1508,3 +1508,141 @@ TEST_CASE("NemotronH paged: the recurrent state index comes from the block table
   REQUIRE(gd.non_spec_state_indices_tensor->size() == 1);
   CHECK((*gd.non_spec_state_indices_tensor)[0] == assigned);
 }
+
+// ═══════════════════════════════════════════════════════════════════════════
+// 12. A2-Q2b (#810) — THE CPU-REACHABLE HALF OF THE DEVICE `lm_head` CHANGE.
+//
+//     A2-Q2b's device arm is `MarlinW4A16Selects`-gated and therefore CUDA-only:
+//     no case in the tree can enter it on this box, and its reachability
+//     deletion mutation is PENDING a `dgx:gpu0` window. But that row ALSO
+//     refactored the part of `NemotronHPagedForward` that every CPU step runs —
+//     the gathered-row count and the download that feeds both the trace and the
+//     host projection — and this file is the vehicle that already drives that
+//     function through a real `GPUModelRunner`. So the refactor is gated HERE,
+//     rather than by a hand-built `NemotronHHostWeights` in a unit test, which
+//     `AGENTS.md` §"Nothing lands dead" is explicit does not count.
+//
+//     WHAT THIS CASE CANNOT SEE, said out loud: the redundant SECOND download
+//     of `final_normed` that A2-Q2b's first submission introduced is not
+//     observable from outside the function — both copies are of the same
+//     unchanged buffer and produce the same bytes. It is repaired structurally
+//     (one `DownloadF32` call reached on any path) and no assertion here
+//     pretends to catch it. What IS asserted is the property that made the
+//     duplicate possible: the trace's copy and the projection's operand must be
+//     the SAME bytes.
+// ═══════════════════════════════════════════════════════════════════════════
+
+// NO COMMA IN THIS NAME: doctest `-tc` splits filters on commas, so a comma
+// here turns a targeted run into `0 cases ran` + `SUCCESS!`.
+TEST_CASE("NemotronH paged: the logit rows are counted from logits_indices and never from num_reqs") {
+  setenv("VT_KV_CACHE_F32", "1", 1);
+  Fixture fx("float32");
+  const vllm::ModelRegistration& reg = fx.model->registration();
+  KVCacheConfig kv = reg.factory->make_kv_cache(fx.cfg, kBlockSize, kNumBlocks);
+  GPUModelRunner runner(fx.cfg, *fx.model, kv, Q(), /*max_num_reqs=*/2, kMaxModelLen,
+                        /*max_num_batched_tokens=*/64);
+
+  const std::vector<int32_t> prompt{1, 7, 3, 9, 2, 14, 5, 11, 0, 6, 8, 4};
+  const int64_t T = static_cast<int64_t>(prompt.size());
+
+  // One prefill through the runner, purely to BUILD the step metadata and the
+  // block table. Nothing below fabricates a `ModelForwardInput`.
+  std::vector<NewRequestData> reqs;
+  reqs.push_back(MakeNewReq("R0", prompt, {0, 1}, /*state_slot=*/0));
+  std::map<std::string, int> sched;
+  sched["R0"] = static_cast<int>(prompt.size());
+  SchedulerOutput s1 = NewStep(std::move(reqs), std::move(sched));
+  CHECK_FALSE(runner.execute_model(s1).has_value());
+  (void)runner.sample_tokens(std::nullopt);
+
+  auto& pages = const_cast<std::vector<vllm::GdnStateCache>&>(runner.gdn_state());
+  // The recurrent pages are advanced IN PLACE, so a replayed prefill would carry
+  // the prompt twice. Both runs below start from zeroed state (same surgery, and
+  // same justification, as the per-layer case above).
+  auto zero_recurrent = [&pages]() {
+    for (vllm::GdnStateCache& c : pages)
+      for (vt::Tensor* p : {&c.conv_state, &c.ssm_state}) {
+        int64_t n = 1;
+        for (int r = 0; r < p->rank; ++r) n *= p->shape[r];
+        std::memset(p->data, 0, static_cast<size_t>(n) * vt::SizeOf(p->dtype));
+      }
+  };
+
+  vt::Queue hq = Q();
+  std::vector<int32_t> positions(static_cast<size_t>(T), 0);
+
+  // ★ THREE gathered rows against ONE request. `num_reqs` is 1 and
+  // `logits_indices.size()` is 3, so the two numbers CANNOT be confused for
+  // each other by accident here — which is the point. A2-Q2b named the gathered
+  // count `n_out` precisely because the surrounding function already binds `R`
+  // to the REQUEST count, and this asserts that the returned row count follows
+  // the gather rather than the batch.
+  const std::vector<int32_t> logits_indices{1, 4, static_cast<int32_t>(T - 1)};
+  const int64_t kWantRows = static_cast<int64_t>(logits_indices.size());
+  REQUIRE(kWantRows != 1);  // or this case cannot tell the two apart
+
+  auto run = [&](bool capture, NemotronHTrace* tr) {
+    zero_recurrent();
+    vllm::ModelForwardInput input{
+        .token_ids = prompt,
+        .positions = positions,
+        .attn_meta = runner.last_attn_meta(),
+        .gdn_meta = runner.last_gdn_meta(),
+        .attn_kv = const_cast<std::vector<vllm::PagedKvCache>&>(runner.attn_kv()),
+        .gdn_state = pages,
+        .config = fx.cfg,
+        .queue = hq,
+        .logits_indices = logits_indices,
+        .num_reqs = 1};
+    if (tr != nullptr) tr->capture = capture;
+    return vllm::NemotronHPagedForward(fx.host, fx.params, input, tr);
+  };
+
+  NemotronHTrace tr;
+  const vllm::ForwardLogits traced = run(/*capture=*/true, &tr);
+
+  // ── (a) THE ARM. On a CPU queue the device projection is not selectable, so
+  //        this step takes the HOST arm — stated as an assertion because every
+  //        other claim below is about the host arm's operands.
+  CHECK_FALSE(vllm::NemotronHDeviceLmHeadEligible(fx.host, fx.host.act_dtype, hq));
+  CHECK_FALSE(traced.on_device());
+
+  // ── (b) THE ROW COUNT, from the gather and not from the batch.
+  CHECK(traced.rows == kWantRows);
+  CHECK(traced.vocab == kVocab);
+  REQUIRE(traced.host.size() == static_cast<size_t>(kWantRows) * kVocab);
+
+  // ── (c) THE TRACE'S COPY AND THE PROJECTION'S OPERAND ARE THE SAME BYTES.
+  //        `final_normed` is downloaded ONCE and serves both consumers. Feeding
+  //        the trace's copy back through the very entry point the production
+  //        fallback calls must reproduce the returned logits EXACTLY — not to a
+  //        band, bit for bit, because it is the same function over the same f32
+  //        values. A second, independently taken download would still pass this;
+  //        a trace taken from a DIFFERENT buffer, or at a different point in the
+  //        graph, would not.
+  REQUIRE(tr.final_normed.size() == static_cast<size_t>(T) * kHidden);
+  std::vector<float> gathered(static_cast<size_t>(kWantRows) * kHidden);
+  for (size_t r = 0; r < logits_indices.size(); ++r)
+    std::memcpy(gathered.data() + r * kHidden,
+                tr.final_normed.data() + static_cast<size_t>(logits_indices[r]) * kHidden,
+                static_cast<size_t>(kHidden) * sizeof(float));
+  const std::vector<float> from_trace =
+      vllm::NemotronHHostLmHead(fx.host, fx.params, gathered, kWantRows, hq);
+  REQUIRE(from_trace.size() == traced.host.size());
+  CHECK(std::memcmp(from_trace.data(), traced.host.data(),
+                    from_trace.size() * sizeof(float)) == 0);
+
+  // ── (d) THE DOWNLOAD IS CONDITIONAL AND THE ANSWER IS NOT. A2-Q2b put the
+  //        trace download behind `capture` so the device arm pays no T*H
+  //        device-to-host copy per step. The host arm must be unaffected: same
+  //        logits, bit for bit, with the trace off.
+  const vllm::ForwardLogits untraced = run(/*capture=*/false, nullptr);
+  CHECK(untraced.rows == kWantRows);
+  REQUIRE(untraced.host.size() == traced.host.size());
+  CHECK(std::memcmp(untraced.host.data(), traced.host.data(),
+                    untraced.host.size() * sizeof(float)) == 0);
+  MESSAGE("A2-Q2b CPU arm: " << kWantRows << " gathered rows x " << kVocab
+                             << " vocab, trace operand byte-identical to the "
+                                "projection operand");
+  unsetenv("VT_KV_CACHE_F32");
+}