fix(#1312): hold the NemotronH device MoE router pair resident instead of re-uploading it per token

[localai-bot]

Closes #1312. Files #1341 and #1346, both fixed-in-flow or owned by this spec.

The question the issue asked first, answered first

Synchronize is a real drain on this backend, not a no-op. src/vt/cuda/cuda_backend.cu:110-112 implements it as Check(cudaStreamSynchronize(AsStream(q)), "cudaStreamSynchronize"), which blocks the calling host thread until every operation previously submitted to that stream has completed. The issue's refutation clause -- "refuted if the sync count is already near zero, i.e. if Synchronize is a no-op on this backend" -- therefore fails, and #1312 is serialisation as well as bandwidth. The no-op is the Backend base implementation (include/vt/backend.h:42), which suits the synchronous CPU backend; Vulkan overrides it with a batch flush (src/vt/vulkan/vulkan_backend.cpp:110). Neither runs this arm.

What was wrong

NemotronHMoeBlockDevice called UploadOwned for mixer.gate.weight and for mixer.gate.e_score_correction_bias on every forward, and UploadOwned ended in d.b.Synchronize(d.q). At the released checkpoint's 23 MoE layers a decode token paid 46 host->device copies and 46 full pipeline drains for two tensors that do not change after load. vLLM has no counterpart to mirror, because it holds the router as an nn.Parameter that already lives on the device.

What changed

Both tensors now sit in two slots on NemotronHMoeMarlinResident, filled on first use and read in place after that. The seam is the one this same function already uses for the 16.5 GB Marlin arena, and the one include/vllm/model_executor/models/dense_attn_block.h:177 states for every other resident weight in the tree. No new caching mechanism is introduced.

The refusal keeps its polarity, and that is the part to review. UploadOwned's four VT_CHECKs are now RequireOwned, and ResidentOwned runs them BEFORE the residency test. A mis-shaped or absent router weight still refuses by name on every forward rather than on the first one only. The first-use copy keeps its drain: one per layer per process is not what this row is about, and removing it is a separate argument.

DenseMarlinE1 stops uploading std::vector<float> ones(M). That vector was handed to DBuf's constructor, whose copy is async, and then died at the end of the function with the copy possibly still in flight -- the use-after-free shape UploadAs documents at :183-189, uncovered. It is fixed by removal, not by adding a drain: the kernel dereferences topk_weights_ptr only under if (mul_topk_weights) (marlin_template.h:518 is the only load, :1879 the only read, :1797 and :1880-1884 the other branches on the same flag), and this call site passes false. The tensor is allocated and left unwritten.

One stale comment corrected in flow (#1346), because it sits directly above the changed lines. The lazy-arena justification claimed this block "has no production caller at all (ForwardNemotronHForCausalLM still routes to the host reference)". A2-P landed the paged fold, so ForwardNemotronHForCausalLM -> NemotronHPagedForward (nemotron_h_registry.cpp:200-202) -> here is one, and the cited line numbers point at an unrelated refusal. The conclusion still stands on its other clause and the comment now says which.

Not fixed, filed instead: #1341. DenseMarlinE1 still routes a dense GEMV through the grouped-MoE alignment machinery. Nobody has measured what that costs at M = 1, and the obvious alternative is reached through an accessor A2-Q2a deliberately avoids on account of #984, so it is not a one-line swap. It is filed rather than left unremarked.

Evidence

All of it from ONE rc run lease on dgx:gpu0 (job 75f0121b-ed8a-42cc-a639-e00bfefafb9c, 2026-08-19), one build tree, RED and GREEN from the same configure.

The box and the stack, re-derived in the container rather than inherited. rc-worker-4b8lj, Ubuntu 24.04.4, aarch64, 20 cores, 119 GB, NVIDIA GB10, GPU at 0% on entry. nvcc release 13.0, V13.0.88 (asserted 13.x, not matched by a wildcard). CUTLASS v4.5.0 at e406c186f. Configure: -G Ninja -DCMAKE_BUILD_TYPE=Release -DVLLM_CPP_CUDA=ON -DVLLM_CPP_CUDA_ARCHITECTURES=121a -DVLLM_CPP_TRITON=ON -DVLLM_CPP_CUTLASS_DIR=..., CONFIGURE_RC=0. The feature lines say this is the production stack and not a degraded one:

-- CUDA target architectures: 121a
--   CUDA feature fp4-mma: ENABLED for [121a]
--   CUDA feature cutlass-nvfp4: ENABLED for [121a]
--   CUDA feature cutlass-fp8: ENABLED for [121a]
--   CUDA feature marlin-nvfp4: ENABLED for [121a]
--   CUDA feature fa2: ENABLED for [121a]
-- CUTLASS found at ...; enabling sm120a NVFP4 cutlass GEMM
-- Marlin NVFP4 W4A16 MoE GEMM enabled (vendored) for [121a]
-- Triton AOT: ... <- sm_121a as vt_aot_sm_121a_..._default

The three DISABLED lines are cutlass-nvfp4-sm100, scaledmm-c3x-sm90 and scaledmm-c3x-sm100, each "no requested arch in [121a] provides it". They are other architectures' features, not a lost lever on this one.

The RED, at 9d55ef92f

[doctest] test cases:  3 |  2 passed | 1 failed | 0 skipped
[doctest] assertions: 47 | 42 passed | 5 failed |
[doctest] Status: FAILURE!
RED_EXIT=1

That one case alone, RED_TC_EXIT=1, 1 | 0 passed | 1 failed | 2 skipped, 18 | 13 passed | 5 failed:

MESSAGE: forward 1: gate uploads 1, bias uploads 1, host syncs 5
MESSAGE: forward 2: gate uploads 1, bias uploads 1, host syncs 4
MESSAGE: forward 3: gate uploads 1, bias uploads 1, host syncs 4
ERROR: CHECK( gate_uploads[1] == 0 ) is NOT correct!   values: CHECK( 1 == 0 )
ERROR: CHECK( gate_uploads[2] == 0 ) is NOT correct!   values: CHECK( 1 == 0 )
ERROR: CHECK( bias_uploads[1] == 0 ) is NOT correct!   values: CHECK( 1 == 0 )
ERROR: CHECK( bias_uploads[2] == 0 ) is NOT correct!   values: CHECK( 1 == 0 )
ERROR: CHECK( syncs[2] == 2 ) is NOT correct!          values: CHECK( 4 == 2 )

The GREEN, at e551bfa17, same build tree

[doctest] test cases:  3 |  3 passed | 0 failed | 0 skipped
[doctest] assertions: 47 | 47 passed | 0 failed |
[doctest] Status: SUCCESS!
GREEN_EXIT=0

That one case alone, GREEN_TC_EXIT=0, 1 | 1 passed | 0 failed | 2 skipped, 18 | 18 passed | 0 failed:

MESSAGE: forward 1: gate uploads 1, bias uploads 1, host syncs 5
MESSAGE: forward 2: gate uploads 0, bias uploads 0, host syncs 2
MESSAGE: forward 3: gate uploads 0, bias uploads 0, host syncs 2
MESSAGE: forward 2 vs forward 1: worst relative deviation 0 over 128 elements
MESSAGE: forward 3 vs forward 1: worst relative deviation 0 over 128 elements

The deviation is 0 over an element count asserted against the geometry, REQUIRE( n == T * H ), 128 == 128, on both comparisons. The answer did not move.

Counts per MoE block per forward, before and after

	before	after
gate uploads, steady state	1	0
bias uploads, steady state	1	0
host drains charged to the router	2	0
host drains charged to the test seam	2	2
total, through the host-in/host-out seam	4	2

The two that remain are UploadAs of the activation and DownloadF32 of the result, both artefacts of the test seam. NemotronHPagedForward keeps the carry on the device and has neither. At the released checkpoint's 23 MoE layers, per decode token: 46 copies and 46 drains, to 0 and 0.

The mutation

Put the per-forward upload back in a scratch copy, one line, slot = nullptr; immediately before the residency test. This is the check that the case detects the defect rather than merely passing beside it:

MUTATION_APPLIED
MUTATE_RC=0
 src/vllm/model_executor/models/nemotron_h_device.cpp | 1 +
 1 file changed, 1 insertion(+)
BUILD_MUT_RC=0                      <- it BUILT, so this is not a build failure reading as a pass
0                                   <- lines matching "error" in the mutation build log
MUT_EXIT=1
[doctest] test cases:  1 |  0 passed | 1 failed | 2 skipped
[doctest] assertions: 18 | 13 passed | 5 failed |

git diff --stat reports the insertion, so the mutation applied; the build returned 0, so the red is the test's and not the compiler's; and the tree restored to an empty git diff and an empty git status --porcelain afterwards.

The A3 end-to-end token gate, at GREEN, on the released checkpoint

[nemotron-h] libvllm 0.0.3+cuda (ABI 21, header 21)
[nemotron-h] engine loaded in 254.5s
[nemotron-h] prompt 0: prompt_tokens=5  generated=32 golden_width=32 compared=32 matched=32 wall=335.87s
[nemotron-h] prompt 1: prompt_tokens=8  generated=32 golden_width=32 compared=32 matched=32 wall=323.42s
[nemotron-h] prompt 2: prompt_tokens=13 generated=32 golden_width=32 compared=32 matched=32 wall=322.71s

[nemotron-h] TOKEN MATCH: 96/96 over 3 prompt(s) (full rows=3, short rows=0, mode=decode)
[nemotron-h] STRICT PASS
A3_EXIT=0

Golden revision 29f2d1746d8f41e316523194b19018707749b1b1, three entries of width 32, against oracle vllm=0.23.1rc1.dev1511+g555967922. Through include/vllm.h alone.

No speed claim is made, and here is why one would be wrong

No A/B was run, so nothing here is a throughput result. What the numbers do settle is that #1312 was never going to be the decode gap. The gate above decodes 32 tokens in 335.87 s, which is 10.50 s per token. The tree's own figure for a cudaStreamSynchronize is ~66 us (deepseek_v4.cpp:203, and that is a DeepSeek-V4 comment, not a measurement taken on this arm), so 46 drains is of order 3 ms, about 0.03% of the step. Removing them is correct and it is free, but it cannot be what a 6.31% decode occupancy is made of.

That is not a ceiling and nothing here says the gap is closed. The next traceable step is the one this arm makes obvious rather than the one this row owns: the 23 Mamba2 blocks still compute their FP8 projections on the HOST every decode token, which is #940 / #1289 / #1311, and that is where a 10.50 s token has to be attributed before anything else on this model is worth measuring.

The local gate, and the two reds that are not this branch's

scripts/agent-ready.py on the merged tree reports two failures, and both were checked
against origin/main rather than assumed:

• test_check_gate_commands fails identically at origin/main 601b576c6: 44 tests,
  8 failures, and the failing-test sets on the two trees diff clean. It arrived with
  main's own ENG-CUDAGRAPH-BREAK commit, which made ENG-CUDAGRAPH-BREAK runnable while
  RUNNABLE_BASELINE still expects it absent. Inherited, subtracted, not chased.
• test_cpu_x86_llamacpp_floor fails the way
  #618 records it: the contended leg exits
  4 (NO_QUIET_WINDOW) instead of 2 (GIVING_UP), printing busy=104% load=25.20 at a
  box loadavg of 27.31. Environmental, and this change touches no CPU path.

Also inherited and worth naming so it is not read as this branch's:
#1371, where main's 369ea7fd4
head_size % 8 == 0 FA guard turns test_nemotron_h_paged_forward red.

Earlier in this row a preflight run reported ten record suites red; every one of them was a
tempfile.mkdtemp failure with the root filesystem at 100%, and all passed bare once space
returned. That is now indexed as
#1353 on main.

How this is tested

A numeric gate cannot see this defect: the answer is identical either way, which is why the existing equivalence case stayed green while it was true. The observable is a count, taken over more than one forward, at T = 1, which is the decode width the claim is about.

The new case installs a DELEGATING vt::Backend over the real CUDA one -- the arm calls real Marlin kernels, so a backend that owned its own memory could not run the block and the counts would be counts of a simulation -- and counts Copy BY SOURCE HOST POINTER and Synchronize per forward. Counting by pointer rather than by byte width means no unrelated 4096-byte transfer can be read as the router's. It asserts that forward 1 uploads each tensor once, that forwards 2 and 3 upload neither, that the steady-state drain count is 2 with each of the two named, and that all three forwards returned the same answer over an element count asserted against the geometry.

The synthetic NVFP4 fixture is the only operand that reaches this arm without a 21 GiB checkpoint, and its entry point is a host-in/host-out seam. Production reachability of the changed lines is carried by the A3 end-to-end gate, which enters through include/vllm.h alone. Both are run. Neither is sufficient alone.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]