gpu, gpu-gates, model, refset: Nemotron 3 Nano 4B on the card and bloomery-serve --model nemotron

[midagedev]

gpu, gpu-gates, model, refset: Nemotron 3 Nano 4B on the card and bloomery-serve --model nemotron

On main after #236 and #235. The group-5 flash entry and the Q5_0 embed branch are #235's as merged; the body calls
enqueue_pass::<GROUP_5>(…, false). relu_sqr, which #235 dropped until a body called it (7852eac), returns here
with its gate case and spill rows (a revert of that commit; unsafe pin 1787 -> 1789). Head keeps no q8_1 form for a
Q8_0 lm_head (its width need not divide 256), on top of main's Head::enqueue_over.

What it adds.

• gpu::arch::nemotron_h: a dense whole-card decode body over the common owners: the site gemvs, the per-head
  plane append (flash::enqueue_kv_append_planes, lifted out of dspark's dense draft, which now calls it),
  recur::RecStore for the Mamba-2 state, the group-5 flash entry (now seg_scalar), relu², the Q5_0 embed rows,
  and the common Head. The body holds only the layer kinds, the in-projection views (z 0, xBC 7680, dt 17408, as
  mainline's mamba-base.cpp:197-202) and the residual order. Program::NemotronHBody and its coverage rows.
• bloomery-serve --model nemotron: one slot; a prompt is one step per token; a prefix is kept only where it
  stands (the Mamba state cannot step back), so no draft and no prefix reuse past the live position. The context
  comes from the common ladder over nemotron_h::place (an attention layer's K/V rows at the context, a Mamba-2
  layer's ctx-free state from runtime::stores::mamba2_bytes, which the body allocates by too).
• Commonized: bind::SeatFlags and Place::machine_untiered lifted out of the mimo2 seat, which calls them; the
  per-head append shared with dspark; model::arch::names::blk is the body's layer names.
• jinja: a comment ends at its first #} (Jinja2's lexer); the Nemotron template's comment holds an apostrophe.

What I ran

Head: see the SHA in the push request. Pre-review CLEAN at 2b2b19d.

• gate-gpu-nemotron-e2e (3090, the real file), every clause PASS: structure (mamba2_step 21/21, relu_sqr
  17/17, gqa_flash_seg_g5 4/4, kv_append_pos_buf 64/64, embed_rows_q5_0 1/1, no group-8 entry, 380 nodes),
  eager_equals_replay, reset_rewinds, state_carries, fault_word_clean, top1_is_mainlines,
  logits_near_mainlines. Against mainline's last-position logits (floor: cosine ≥ 0.999, top-5 ids equal as a set):

  Prompt	top-1	max |Δ|	cosine
  code5	16473	0.2406	0.999816
  code29	5803	0.2306	0.999886
  prose64	4317	0.2777	0.999858
  prose512	6911	0.3773	0.999685

  prose64's ranks 3 and 4 swap (body 15.147 vs 15.141; mainline 15.173 vs 15.042); order inside the top-5 is printed,
  not gated.

• FAIL-first, each a reverted mutant: layer 0's A NaN (fault named), the attention as group 8 (structure red),
  a Mamba layer skipped, relu² as relu, the embed row off by one (both mainline clauses red on all four prompts,
  cosine 0.62–0.90); the seat's load ctx halved, prefill skipping the first id, keep granting a shorter cut
  (gate-gpu-nemotron-serve red); the plan's Mamba state × ctx (two unit tests red).

• gate-gpu-nemotron-serve 7/7 PASS; gate-nemotron-meta, gate-gpu-mamba2, gate-gpu-head,
  gate-gpu-qwen3moe-types, gate-gpu-qwen3moe-flash, gate-gpu-qwen3moe-experts, gate-ptx-spill,
  gate-gpu-qwen35-dspark, gate-gpu-qwen35moe-e2e, gate-gpu-qwen4exp-e2e, gate-jinja, gate-serve,
  gate-decision-clef, gate-decision-lev, gate-ds41-bind, gate-gpu-clef-serve, gate-gpu-glm5next-serve,
  gate-gpu-lev-serve, gate-gpu-qwen38-serve, gate-q38-serve-rules, gate-placement, gate-refset rc 0 on
  da1cc8eb; gate-gpu-qwen3-serve PASS on the A6000 (s613; on the 3090 the whole loads do not fit and the seat
  places them, which is that gate's card fact, not this change); gate-gpu-ds41-chat PASS (s688),
  weekly-gpu-mimo2-serve rc 0, weekly-gpu-ds41-serve exit 0 (s689) on 5e77a2c8's line. Jobs whose client lost
  the reply (rc 84) were read back with boxq.py attach: s640, s641, s651, s652, s688, s689.

• After the rebase onto main (conflicts only in accepts.rs program arms, unsafe-ratchet.txt counts and
  gate-paths.tsv rows): check-unsafe ok, check-comments ok, check-recipes table ok; on 272d0beb
  check, gate-gpu-qwen3moe-experts, gate-ptx-spill, gate-gpu-qwen3moe-flash, gate-gpu-mamba2, gate-placement,
  gate-nemotron-meta, gate-gpu-nemotron-e2e, gate-gpu-nemotron-serve, fmt-check, lint (41) all rc 0; after the
  last folds (the shared e2e::argmax, generate::trained_ctx in the serve gate), gate-gpu-nemotron-e2e,
  gate-gpu-nemotron-serve, fmt-check and lint (41) rc 0 on ea2bffd1.

• Restacked on refset: Nemotron 3 Nano last-position logits, read from one reference set #236's 07aa96b (the reader keeps refset: Nemotron 3 Nano last-position logits, read from one reference set #236's id checks in ids; the reader test reads refset: Nemotron 3 Nano last-position logits, read from one reference set #236's synthetic
  ids): gate-gpu-nemotron-e2e (every clause PASS, the same Δ and cosine as above), gate-refset, fmt-check,
  lint (41) rc 0 on 10b2e5e; cargo test -p bloomery-refset 208 passed.

• Review fold (216e9681, Shingler on 3a71fa2): the nemotron type pins (coverage::TYPE_PINS: q5_0 token table;
  q8_0, q6_K or q4_K head; the body's four site types for projections and the FFN), so --plan refuses a file the
  body would refuse at load; the body takes names::blk; plan is the one-slot plan (no plan_with_slots); the
  keep rule sits in the seat; Hparams copies dropped; the seat records GpuModel::ctx_max, the context the model
  holds. FAIL-first: the new coverage test red with q8_0 added to the table's pin; gate-gpu-nemotron-serve red
  (load_and_listen: FAIL, job s834) with the seat opening the model at half the planned context. Rebased onto main 73c307c (gpu: geglu_tanh, the GELU-gated unit over f32 rows #237's geglu_tanh and
  relu_sqr side by side in elem.rs; unsafe pin 1788 -> 1790); gate-paths rows filled by recipes.py check --fill.

• Review fold 2 (Shingler on 7448bb0): the nemotron pins read every tensor of their roles, F32 included (the
  projections and the head by name), so an F32 table, head or projection is an item at --plan; a Mamba-2 layer
  whose file holds no ssm_conv1d.bias (the reader takes it as optional, as mainline does) reads one derived zero
  vector resident from load; the seat's context is GpuModel::ctx_max alone; PlaceError keeps the default
  PlanRefusal::short. FAIL-first: the coverage test's F32 rows red with the pins back on matrices only. Not
  gated: a bias-free file on the card; every Nemotron file at hand holds the bias, so the zero path has no file to
  run, and the bias-present path is the one the gates below run.
  Machine (box, the gates' host), nvidia-smi --query-gpu=name,driver_version --format=csv,noheader, 2026-10-15:

  NVIDIA GeForce RTX 3090, 615.71.09
  NVIDIA RTX A6000, 615.71.09

Head 2b2b19d on main 22c6540, through the box queue (BLOOMERY_BOX_QUEUE=1 tools/gate-batch.sh check lint fmt-check gate-nemotron-meta gate-gpu-nemotron-serve gate-gpu-nemotron-e2e, each item just <recipe>; the card items with
BLOOMERY_BOX_ENV='BLOOMERY_GATE_CARD=3090' or '…=a6000'):

$ BLOOMERY_BOX_ENV='BLOOMERY_GATE_CARD=3090' just gate-gpu-nemotron-e2e
structure: layers mamba=21 attn=4 mlp=17, graph_nodes=380 host=0, entries (got/want) mamba2_step=21/21 relu_sqr=17/17 gqa_flash_seg_g5=4/4 …
eager_equals_replay: walk16 16 tokens, ids [1046, 6013, 1046, 5963], logits bit-identical true, finite true PASS
reset_rewinds: walk16 ids equal true, logits bit-identical true PASS
state_carries: walk16 last token 119798 stepped twice from the same reset: ids [1010] vs [1010] equal true; … PASS
top1_is_mainlines: code5 16473 / code29 5803 / prose64 4317 / prose512 6911, each the mainline's PASS
logits_near_mainlines: cosine 0.999816 / 0.999886 / 0.999858 / 0.999685 (floor 0.999), top-5 same ids true PASS
fault_word_clean: word None, poisoned None PASS
gpu-gate.sh: gate_nemotron_e2e ran 13 s (exit 0)

$ BLOOMERY_BOX_ENV='BLOOMERY_GATE_CARD=a6000' just gate-gpu-nemotron-serve
check plan_only_names_its_default: PASS
check parallel_above_one_is_refused: PASS
check load_and_listen: PASS
check chat_is_200: PASS
check edit_resend_is_a_fresh_run: PASS
check extension_keeps_the_held_sequence: PASS
check completion_ids: PASS
gpu-gate.sh: gate_nemotron_serve ran 14 s (exit 0)

$ just gate-nemotron-meta      # test result: ok. 1 passed
$ just check                   # rc 0
$ just fmt-check               # rc 0
$ just lint                    # rc 0, 41 `^warning:` lines, the ratchet's 41
$ just contrib-check --tier t0 # DONE tier=t0 total=83 red=0 (check-arch, check-recipes included)

• Review fold 4 (Shingler on a25cbe0): the nemotron pins for the Mamba-2 parameters, the normalization gains and
  output_norm (f32 and bf16, what f32_buf reads; a q8_0 one is an item at --plan); the serve gate's three
  request clauses in one loop over check (no ok: &mut bool wrapper, R24); the weights test keeps no blk.
  literal outside arch/ (check-arch). FAIL-first: the coverage test's q8_0 rows red with q8_0 added to the
  parameter pin.
• Review fold 3 (Shingler on 242bc22): Weights::load refuses a file tensor with a derived. name on both its
  filing paths (load_named, load_segments), so a lookup of a derived name, such as the zero conv bias, finds
  only what insert_derived filed; gate-paths' nemotron_h row names the six gates coverage.rs now reaches it
  through. FAIL-first: gate-gpu-lib's a_file_tensor_never_takes_a_derived_name red with the refusal off.
  Rebased onto main 22c6540 (unsafe pin 1814 -> 1816 for relu_sqr beside the head-512 entries).
• On a25cbe0 (main 22c6540): check, lint (41), fmt-check, gate-gpu-lib, gate-ptx-spill,
  gate-nemotron-meta, gate-gpu-nemotron-e2e (cosine 0.999816 / 0.999886 / 0.999858 / 0.999685),
  gate-gpu-nemotron-serve rc 0.
• On 242bc22: check, lint (41), fmt-check, gate-nemotron-meta, gate-gpu-nemotron-e2e (cosine 0.999816 /
  0.999886 / 0.999858 / 0.999685, as above), gate-gpu-nemotron-serve rc 0.
• On 7448bb0: gate-gpu-nemotron-e2e rc 0 on the 3090 (s848; cosine 0.999816 / 0.999886 / 0.999858 / 0.999685,
  as above), gate-gpu-nemotron-serve rc 0 on the A6000 (s847, 7/7 PASS); check, gate-nemotron-meta,
  gate-placement, gate-ptx-spill rc 0 on its tree less one unused test import. gate-refset is red on main
  73c307c the same way (the mimo2 draft unit test, eight EmbeddingGemma sets not on the box), not here.
• fmt-check rc 0; lint 41 ^warning: lines, the ratchet's 41.

PTX scan and draft gate: base 22c6540 vs head a25cbe0

Shingler's non-blocking note: the shared FlashKernels::enqueue_kv_append_planes (flash.rs:1947) changes the DSpark dense draft's enqueue order. Before, each key head appended its K and V inside the head loop. Now, after the loop, each pass appends every head's K and then every head's V. Each head reads its own kr[h] / vb[h] and writes its own cache rows on one stream, and the launch count is the same. So the prediction is identical PTX entries and identical draft bits.

PTX entries (tools/narrow-scan.sh 22c65407 a25cbe0f, --no-jit, the scan set covering every kernel carrier):

bin (features)	base entries	head entries	added	removed	changed
bloomery-serve-ds41 (deepseek41,vision)	324	325	relu_sqr	none	0
gate_kquant (gpu)	245	246	relu_sqr	none	0
gate_swap (gpu)	242	243	relu_sqr	none	0

narrow-scan refuses the pair by name, as the add class requires: added relu_sqr in a pair with a --no-jit scan (both) — rescan without --no-jit: the add and occupancy classes need the JIT columns. So gate_kquant was also scanned with the JIT on the RTX 3090 (just ptx-scan gate_kquant --features gpu, jobs s881 base and s882 head, jit-card=NVIDIA_GeForce_RTX_3090 jit-cc=8.6 jit-cuda=13.4, ptxas 13.3.73, sm_86):

• 245 base entries: every column equal, the JIT columns included. Every ptx-scan-md5 (method decl1) equal.

• The one added entry:

  entry     reqntid depot ld.local st.local fma cvt.f16 regs smem spill blk/SM(static) jit_regs jit_local ld.global.scalar ld.global.v2 ld.global.v4
  relu_sqr  256     no    0        0        0   0       10   0    0     6              10       0         1                0            0
  relu_sqr  md5 1950d4df3f21ced0ee6c421e54b69dbb (39)
  

Draft gate (just gate-gpu-qwen35-dspark, real tier, RTX 3090):

• Base 22c6540 (s868) and head a25cbe0 (s873) both exit 0, with no FAIL line.
• The gate's 16 output lines are equal between the two runs except for timing fields: the loader time, the append time (316.057 µs vs 314.906 µs at R=3, 157.39 ms vs 157.74 ms at R=2100) and the host reference time. Every value, ratio and verdict is the same, including "the captured block pass replays the eager bits", 2 appends at R=3 and 487 at R=2100, and a largest Q8_0 ratio of 0.056.

Cards (nvidia-smi on the box):

GPU 0: NVIDIA GeForce RTX 3090 (UUID: GPU-307fa0f6-daae-24e5-6fd3-cd50620de6b1)
GPU 1: NVIDIA RTX A6000 (UUID: GPU-8c129fa6-7382-35a5-2464-9ff01d99fcd4)
0, NVIDIA GeForce RTX 3090, 615.71.09, 00000000:41:00.0, 24576 MiB
1, NVIDIA RTX A6000, 615.71.09, 00000000:61:00.0, 49140 MiB

Fold 4 (2b2b19d, on top of a25cbe0) changes no kernel and no launch: coverage pins, the serve gate's clause loop and a weights unit test.

Not here: multi-slot (Slots), a prompt batch path (pp ≈ decode rate), and the Q5_0 wide GEMM.

🤖 Generated with Claude Code

[github-actions]

Verdict: CLEAN: no failure path survived refutation and no proven duplicate, unneeded element or unpinned gate change is in the diff.

Non-blocking

⚠ ids_of leaks temp files on failure: dump-nemotron.sh:51 makes raw with a bare mktemp, outside $WORK. If llama-tokenize fails under set -e, $raw, $raw.in and $raw.bracket stay in /tmp. Use mktemp -p "$WORK" so the EXIT trap removes them.

Optimization headroom

• G per-launch cost: step.rs:377-379 q, k and v each quantize the same s.xn into the same act. That is 2 redundant quantize launches per attention layer per token when the three sites share a type and width.

Checked

• Duplicates: about 40 new items Grepped, none duplicated. Dryer: 10 pairs read, all in place.rs. The PlaceError glue repeats in other models, but a shared owner would touch six models and add more elements than it removes, so it is not a finding. Acts here and in qwen3moe/wide.rs only share a name.
• Complexity: crapper's flagged functions are new gate, seat and loader code. Head::with_norm and Head::enqueue grew by an Option and a match arm. drive::run and check_with grew by one arm each. No #[kernel] body is split.
• Silent failure: 9 sites followed (unwrap_or(0) on ff, longest and found in body.rs; unwrap_or_default on the log in the serve gate; .ok() in refresh; the others in gate code). Each is either refused by name earlier or has no consumer when it is empty. A row word past u32 is a named error in refresh. A position past the context is refused by GpuModel::step (model.rs:1671).
• Tests: mutants checked and each turns an assertion red. n <= pos in keep_where_standing fails (6, 7). A ring row count off by one fails mamba2_sizes. Dropping the #} branch fails the jinja test. relu_sqr with no finite guard fails the NaN bit compare. relu_sqr with >= is an equivalent mutant.
• Gates: pins moved with a dated PIN(2026-10-15) and reason: the jinja closer test, the spec.rs and nemotron_meta coverage clauses, the accepts.rs card-leg assert, the recipes.py set, the unsafe-ratchet 1787→1789 and the clef MODELS table. The new gates carry FAIL-first mutants in the PR body.
• Escape catalog: C1 to C4 checked on the new scripts, none matched. No pipe or substitution hides an exit code, no unquoted heredoc, no pgrep or pkill.

Notes

• gpu.json evidence is false: the body has gate results but no nvidia-smi line and no commit hashes ("see the SHA in the push request").
• The body asks to review only the commits after #236. diff.patch holds the whole branch, and this review covers all of it.

_{Card check: needs-gpu. This pull request changes crates/gpu-gates/Cargo.toml, crates/gpu-gates/src/bin/bloomery_serve.rs, crates/gpu-gates/src/bin/gate_clef_serve.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

Verdict: MUST_FIX: Planning accepts unsupported formats; duplicated and unneeded elements remain, and the context gate misses its claimed mutant.

Must fix

1. Reject unsupported tensor formats

crates/model/src/arch/nemotron_h/place.rs:102 An otherwise supported file with a Q8_0 embedding passes --plan, then fails Body::load at body.rs:95, which requires Q5_0. Add these coverage::TYPE_PINS rows before callers rely on the plan:

TypePin {
    program: Program::NemotronHBody, role: Role::TokenEmbedding,
    matrices: false, names: None, what: "embedding",
    reads: &[GgmlType::Q5_0],
},
TypePin {
    program: Program::NemotronHBody, role: Role::Head,
    matrices: false, names: Some(|n| n == "output.weight"),
    what: "output head",
    reads: &[GgmlType::Q8_0, GgmlType::Q4_K, GgmlType::Q6_K],
},
TypePin {
    program: Program::NemotronHBody, role: Role::Attention,
    matrices: false,
    names: Some(|n| [
        ".ssm_in.weight", ".ssm_out.weight", ".attn_q.weight",
        ".attn_k.weight", ".attn_v.weight", ".attn_output.weight",
    ].iter().any(|&s| n.ends_with(s))),
    what: "projection",
    reads: &[GgmlType::Q5_0, GgmlType::Q8_0, GgmlType::Q4_K, GgmlType::Q6_K],
},
TypePin {
    program: Program::NemotronHBody, role: Role::DenseFfn,
    matrices: false, names: None, what: "dense FFN",
    reads: &[GgmlType::Q5_0, GgmlType::Q8_0, GgmlType::Q4_K, GgmlType::Q6_K],
},

2. Reuse the tensor-name owner

crates/gpu/src/arch/nemotron_h/body.rs:275 duplicates crates/model/src/arch/names.rs:7 (R14). Share the existing owner before callers multiply. Make it public:

pub fn blk(l: usize, stem: &str) -> String {
    format!("blk.{l}.{stem}")
}

Import model::arch::names::blk as layer_name in the body and delete its layer_name.

3. Remove unused slot planning

crates/model/src/arch/nemotron_h/place.rs:145 Only plan and a test call plan_with_slots; serving refuses multiple sequences. Delete it, seq_terms, the slot-only test and imports before they gain callers. Replace plan with:

pub fn plan<'a>(
    &'a self, machine: &'a Machine, ctx_max: u64, levers: &PlanLevers,
) -> Result<Plan<'a>, PlaceError> {
    if ctx_max == 0 {
        return Err(PlaceError::CtxZero);
    }
    if let Some(card) = machine.tiers.first() {
        return Err(PlaceError::Tier { card: card.name.clone(), tier: 0 });
    }
    if machine.cards.len() != 1 {
        return Err(PlaceError::Cards(machine.cards.len()));
    }
    checked(placement::plan_routed(
        &self.model, machine, ctx_max, &self.kv, levers,
        routed_rule(Program::NemotronHBody),
    )?).map_err(PlaceError::Broken)
}

4. Inline the single-user prefix rule

crates/gpu-gates/src/bind.rs:627 Only Nemotron calls this helper; MiMo calls Session::keep_query. Delete the helper, import and unit test before it gains callers. The serve gate covers both outcomes. Use:

fn keep(&self, n: usize) -> (usize, Option<String>) {
    if n == self.pos() {
        (n, None)
    } else {
        (0, Some(NO_STEP_BACK.to_owned()))
    }
}

5. Remove redundant hyperparameter fields

crates/gpu/src/arch/nemotron_h/body.rs:44 retains Hparams only for epsilon, already in Shape. PlanInputs::hp (place.rs:36) supplies only ModelTensors' layer count. Delete both fields and initializers before their copies diverge. Replace inputs.hp.n_layer with inputs.model.layers. Use:

fn head_eps(&self) -> f32 {
    self.sh.eps
}

6. Read the held context in the gate

crates/gpu-gates/src/bin/gate_nemotron_serve.rs:337 compares records that echo a.ctx. Mutant: pass a.ctx / 2 to NemotronHModel::open in serve_seats/nemotron.rs:131. The records stay unchanged and short prompts fit, so the gate stays green despite its stated contract. This is a kind B gate defect. Add a documented GpuModel accessor:

/// The context limit held by the resident model.
pub fn ctx_max(&self) -> usize {
    self.ctx_max
}

In Nemotron::open, set let ctx = m.ctx_max();, record .u("ctx", ctx) and return Ok(Nemotron { m, ctx }). The mutant then fails.

Optimization headroom

• G. Per-launch and per-sync fixed cost: crates/gpu/src/arch/nemotron_h/step.rs:377, quantize s.xn once for Q/K/V sites sharing a form and width, using existing Acts buffers: up to two fewer quantization launches per attention layer per token [derived].

Checked

Evidence: traced all six findings and their callers; cited lines are added by the diff. The naming fix reuses an existing owner.
Tests: statically traced the context-halving mutant; no GPU gates run. Test removals above need dated reasons.
Scope: full review; no ledger lessons in force; escape patterns C1–C4 do not apply to these findings.

Notes

The PR text lacks the nvidia-smi line and complete gate command output.

_{Card check: needs-gpu. This pull request changes crates/gpu-gates/Cargo.toml, crates/gpu-gates/src/bin/bloomery_serve.rs, crates/gpu-gates/src/bin/gate_clef_serve.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

Verdict: MUST_FIX: planning accepts files the new loader cannot run, and two elements repeat existing owners.

Must fix

1. Apply the type pins to F32 projections

crates/model/src/arch/coverage.rs:801 weight_formats excludes F32 from matrix. An F32 token table, head or projection passes --plan, then fails loading. The test misses F32.
Fix admission before callers rely on these plans. Replace the four Nemotron pins with:

TypePin {
    program: Program::NemotronHBody,
    role: Role::TokenEmbedding,
    matrices: false,
    names: None,
    what: "token embedding (the card reads q5_0 rows)",
    reads: &[GgmlType::Q5_0],
},
TypePin {
    program: Program::NemotronHBody,
    role: Role::Head,
    matrices: false,
    names: Some(|name| name == "output.weight"),
    what: "output head (the head reads q6_K, q4_K and q8_0)",
    reads: &[GgmlType::Q6_K, GgmlType::Q4_K, GgmlType::Q8_0],
},
TypePin {
    program: Program::NemotronHBody,
    role: Role::Attention,
    matrices: false,
    names: Some(|name| {
        use super::nemotron_h::names as n;
        [n::SSM_IN, n::SSM_OUT, n::ATTN_Q, n::ATTN_K,
         n::ATTN_V, n::ATTN_OUTPUT]
            .into_iter().any(|stem| name.ends_with(stem))
    }),
    what: "Mamba-2 and attention projections (the body reads q5_0, q8_0, q4_K and q6_K)",
    reads: NEMOTRON_SITES,
},
TypePin {
    program: Program::NemotronHBody,
    role: Role::DenseFfn,
    matrices: false,
    names: None,
    what: "dense FFN (the body reads q5_0, q8_0, q4_K and q6_K)",
    reads: NEMOTRON_SITES,
},

Add F32 rejection assertions for these four roles.

2. Honor the optional convolution bias

crates/gpu/src/arch/nemotron_h/body.rs:242 unconditionally resolves ssm_conv1d.bias. The reader marks it optional and its existing the_conv_bias_may_be_absent test accepts its absence. Such a file passes planning, uploads its weights, then fails instead of using zero bias.
Resolve the absent bias once at load. Let Body::load and Names::w take mutable weights, make kinds mutable and Names::layer take &mut self. Add this method to Names:

fn conv_bias(
    &mut self, l: usize, len: usize, stream: &CudaStream,
) -> Result<String, GpuError> {
    if self.file.find(&layer_name(l, names::SSM_CONV_BIAS)).is_some() {
        return self.vector(l, names::SSM_CONV_BIAS, len);
    }
    let name = "derived.nemotron_h.zero_conv_bias".to_owned();
    if self.w.get(&name).is_none() {
        self.w.insert_derived(name.clone(), crate::weights::DevWeight::F32 {
            w: DeviceTensor::zeroed(stream, 1, len)?,
            k: len,
        })?;
    }
    Ok(name)
}

Replace the initializer with:

conv_b: self.conv_bias(l, s.conv_dim, stream)?,

Gate a bias-free file against the equivalent explicit-zero bias.

3. Read context from the model

crates/gpu-gates/src/bin/shared/serve_seats/nemotron.rs:120 duplicates m.ctx_max() for one consumer (R14).
Remove it now so the serving limit has one owner. Use:

struct Nemotron {
    m: NemotronHModel,
}

Initialize with Ok(Nemotron { m }) and replace the getter:

fn ctx_max(&self) -> usize {
    self.m.ctx_max()
}

4. Use the inherited refusal method

crates/model/src/arch/nemotron_h/place.rs:85 repeats PlanRefusal::short exactly. The override adds no behavior.
Delete the override and its doc comment, and remove Short from the import. The inherited method supplies the same behavior.

Non-blocking

⚠ PTX evidence: crates/gpu/src/dspark/dense.rs:830, the append extraction is the move class, retaining 4 × n_kv append launches per layer per prompt chunk [derived]. Attach before/after ptx-scan output showing identical existing entries and only the restored relu_sqr addition.

Optimization headroom

• G. Per-launch and per-sync fixed cost: crates/gpu/src/arch/nemotron_h/step.rs:377, when Q/K/V are all Q5_0 or all K-quants, each quantizes the same input. Reusing the existing activation buffer's quantized input removes two quantization launches per attention layer per token [derived].

Checked

Duplicates: verified the context copy and redundant override against their existing owners.
Complexity: diagnostics inspected; no finding rests on complexity alone.
Behavior: traced coverage admission through planning to body/head loading, including missing-bias lookup.
Tests: the coverage test omits F32; the reader test accepts absent bias. Validation was static; no gates or mutants run.
Gates: PTX comparison requested above. Escape catalog: no matching pattern. Ledger: no lessons in force.

Notes

Card runs lack the nvidia-smi card/driver line and exact commands with output.

_{Card check: needs-gpu. This pull request changes crates/gpu-gates/Cargo.toml, crates/gpu-gates/src/bin/bloomery_serve.rs, crates/gpu-gates/src/bin/gate_clef_serve.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

Apply this pin to F32 token tables too: they pass planning, but the loader accepts only Q5_0.

Suggested change

	matrices: true,
	matrices: false,

[github-actions]

Since the last review:

• F32 admission: fixed at crates/model/src/arch/coverage.rs:801, with rejection assertions for all four roles.
• Optional bias: fixed at crates/gpu/src/arch/nemotron_h/body.rs:215; the new cache collision is finding 1 below.
• Context copy: fixed at crates/gpu-gates/src/bin/shared/serve_seats/nemotron.rs:154.
• Refusal override: fixed at crates/model/src/arch/nemotron_h/place.rs:68; the inherited method applies.

Verdict: MUST_FIX: the new zero-bias cache can silently use nonzero file data.

Must fix

1. Refuse a file tensor that occupies the zero-bias name

crates/gpu/src/arch/nemotron_h/body.rs:224: open a bias-free GGUF containing a finite, nonzero F32 vector named derived.nemotron_h.zero_conv_bias, of conv_dim values. Public NemotronHModel::open uploads globals through crates/gpu/src/model.rs:638. This lookup then skips allocation and insert_derived's collision guard. Mamba uses the file vector as the missing bias and runs with nonzero bias without an error. Serving's classifier rejects the extra name; direct model opening does not call it.

The fix reuses the existing insertion guard and preserves the shared zero vector. Replace the cache insertion block with:

if self.file.find(&name).is_some() || self.w.get(&name).is_none() {
    let w = DeviceTensor::zeroed(stream, 1, len)?;
    self.w
        .insert_derived(name.clone(), DevWeight::F32 { w, k: len })?;
}

This preserves zero-vector reuse and lets the existing owner reject the occupied name. Add a direct-open regression with the colliding vector; removing the file-name arm must make its error assertion fail.

Checked

Duplicates: the new constant and helper have no second definition in the GPU crate. The context copy and redundant refusal override are removed.
Complexity: diagnostics inspected; this delta changes Body::load and Names::layer for mutable weights and a load-time bias call. The new branch belongs to conv_bias.
Silent failure: 3 skip sites followed. Missing bias gets zeros; the cache collision consumes file values; the test filter only excludes F32 format items.
Tests: static mutation check: restoring matrices=true at coverage.rs:801, 809, 817 or 837 breaks the corresponding F32 assertion; dropping SSM_IN at coverage.rs:821 breaks its assertion. Changing RopeMode::None at spec.rs:63 breaks an_attention_layer_needs_no_rope. No assertion covers the new bias fallback. No Rust tests or GPU gates executed here.
Performance: the allocation and name resolution occur once at load.
Gates: no numeric threshold moved; the need-test exemption has a dated PIN and reason. Escape catalog: no matching pattern.

Notes

The author reports no bias-free card run because available files contain bias. Card evidence still lacks the nvidia-smi card/driver line and exact gate commands with output.

_{Card check: needs-gpu. This pull request changes crates/gpu-gates/Cargo.toml, crates/gpu-gates/src/bin/bloomery_serve.rs, crates/gpu-gates/src/bin/gate_clef_serve.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

A direct model open uploads a file tensor at this derived name; force that collision through insert_derived instead of using its nonzero values as the missing bias.

Suggested change

	if self.w.get(&name).is_none() {
	if self.file.find(&name).is_some() || self.w.get(&name).is_none() {

[github-actions]

Verdict: MUST_FIX: planning accepts tensor types the new body cannot load, and a new check helper violates R24.

Must fix

1. Pin the parameter and normalization types before planning

crates/model/src/arch/coverage.rs:814 The new pins omit normalization gains and Mamba parameters. Change blk.0.attn_norm.weight to Q8_0 with its valid [3136] shape: header validation, coverage and placement accept it, so --plan succeeds. Loading then fails at crates/gpu/src/arch/nemotron_h/body.rs:235 because f32_buf requires F32 residency.
Reject this before uploading the model; diagnosing a refusal after allocation costs more than checking the header.

Insert these entries into TYPE_PINS. Extend the existing coverage test with Q8_0 gains and biases refused, and F32/BF16 accepted.

    TypePin {
        program: Program::NemotronHBody,
        role: Role::Attention,
        matrices: false,
        names: Some(|n| {
            use super::nemotron_h::names as nm;
            [
                nm::ATTN_NORM, nm::SSM_A, nm::SSM_D,
                nm::SSM_DT_BIAS, nm::SSM_CONV,
                nm::SSM_CONV_BIAS, nm::SSM_NORM,
            ]
            .into_iter()
            .any(|stem| n.ends_with(stem))
        }),
        what: "Mamba-2 parameters and normalization gains",
        reads: &[GgmlType::F32, GgmlType::BF16],
    },
    TypePin {
        program: Program::NemotronHBody,
        role: Role::FfnNorm,
        matrices: false,
        names: None,
        what: "normalization gains",
        reads: &[GgmlType::F32, GgmlType::BF16],
    },
    TypePin {
        program: Program::NemotronHBody,
        role: Role::Head,
        matrices: false,
        names: Some(|n| n == "output_norm.weight"),
        what: "output normalization gain",
        reads: &[GgmlType::F32, GgmlType::BF16],
    },

2. Remove the new accumulator wrapper

crates/gpu-gates/src/bin/gate_nemotron_serve.rs:386 R24 explicitly forbids an ok: &mut bool accumulator in a new check helper. clause adds this interface around the existing serve_client::check.
Remove the wrapper now, before another gate copies it. Delete clause and replace its three calls in served with:

        for (name, run) in [
            (
                "chat_is_200",
                chat_is_those_ids(
                    &url, &d, messages, N, (&turn_answer, &turn_body),
                ),
            ),
            (
                "edit_resend_is_a_fresh_run",
                edit_resend(&url, &d, &prose),
            ),
            (
                "extension_keeps_the_held_sequence",
                extension(&url, &d, &prose, &prose_answer),
            ),
        ] {
            let pass = run.unwrap_or_else(|e| {
                println!("{name}: {e}");
                false
            });
            check(ok, name, pass);
        }

Non-blocking

⚠ Refactoring proof: crates/gpu/src/flash.rs:1947 The shared append extraction changes the draft's enqueue order while retaining its launch count. The move/split class requires unchanged existing PTX entries and preserved graph structure. Attach a base/head scan comparison showing only the declared relu_sqr addition beside the draft gate result; retaining it now avoids another verification run.

Optimization headroom

• G, Per-launch and per-sync fixed cost: crates/gpu/src/arch/nemotron_h/step.rs:388 Batch heads in the common append owner. Two launches per attention layer instead of 2 × n_kv remove 56 launches/token for this file's four attention layers and eight KV heads [derived].

Checked

Duplicates: 179 new declarations searched; all six dryer pairs read. They are small adapters over different error types or the existing coverage owner; sharing them adds machinery.
Complexity: gate plan_only_names_its_default/served and seat run orchestrate checks/setup; SeatFlags::parse/Names::layer add match arms; Scratch::new/Body::load allocate/load; Head's two functions add type branches; server run/coverage::check_with add enum arms.
Silent failure: 11 default/conversion/result-discard sites followed. Missing log text fails the context assertion; conversion errors become named errors; other defaults represent absence or unreachable conversion failure.
Tests: static mutation analysis covers dropped layers/reset, wrong group/embedding/ReLU, shorter keep, half context, context-scaled state, restored comment quote scanning and removed type/name refusals. Missing parameter-type coverage is finding 1. Gates were not run here.
Gates: removed contract clauses have dated reasons; no tolerance relaxed.
Escape catalog: C1–C4 do not match.

Notes

GPU evidence is incomplete: the PR gives gate summaries but omits the nvidia-smi card/driver line and complete commands with output. gpu.json records false.

_{Card check: needs-gpu. This pull request changes crates/gpu-gates/Cargo.toml, crates/gpu-gates/src/bin/bloomery_serve.rs, crates/gpu-gates/src/bin/gate_clef_serve.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

Verdict: MUST_FIX: byte accounting is incomplete, and new code violates R14 and R24.

Must fix

1. Count the Mamba scratch and projection ids

crates/gpu/src/arch/nemotron_h/step.rs:191 excludes Mamba2Scratch's part and done; crates/gpu/src/arch/nemotron_h/body.rs:362 excludes KGemvKernels::one. A successful seat load passes these totals through GpuModel::resident_bytes() into its load record, undercounting Nano 4B by 3,904 bytes [derived]. Count these buffers before using load records as memory evidence.

Add inside impl Mamba2Scratch:

pub(crate) fn device_bytes(&self) -> usize {
    self.part.num_bytes() + self.done.num_bytes()
}

Add inside impl KGemvKernels:

pub(crate) fn device_bytes(&self) -> usize {
    self.one.num_bytes()
}

Append + self.mamba.device_bytes() to Scratch::bytes. Replace the final expression in Body::resident_bytes with:

layers + self.s.bytes() + self.acts.bytes() + self.r.num_bytes()
    + self.k.kgemv.device_bytes()

2. Leave zero-context validation with the loader

crates/gpu/src/arch/nemotron_h/body.rs:77 duplicates crates/gpu/src/model.rs:627. The loader rejects this same context before calling the private body loader. Delete the unreachable branch to keep one owner for the bound (R14).

--- a/crates/gpu/src/arch/nemotron_h/body.rs
+++ b/crates/gpu/src/arch/nemotron_h/body.rs
@@ -77,3 +76,0 @@
-        if ctx == 0 {
-            return Err(GpuError::shape(WHAT, "a cache of 0 rows"));
-        }

3. Return gate status instead of mutating the caller's status

crates/gpu-gates/src/bin/gate_nemotron_serve.rs:166 and lines 247, 309 and 431 add ok: &mut bool check parameters, forbidden by R24. Return status before adding clauses. This preserves checks; apply both diffs with git apply --unidiff-zero.

--- a/crates/gpu-gates/src/bin/gate_nemotron_serve.rs
+++ b/crates/gpu-gates/src/bin/gate_nemotron_serve.rs
@@ -166,2 +166,2 @@
-        ok: &mut bool,
-    ) -> Result<Option<u64>, GateError> {
+    ) -> Result<(Option<u64>, bool), GateError> {
+        let mut ok = true;
@@ -236 +236 @@
-            ok,
+            &mut ok,
@@ -240 +240 @@
-        Ok(ctx_line(&text).and_then(|l| field(l, "ctx")))
+        Ok((ctx_line(&text).and_then(|l| field(l, "ctx")), ok))
@@ -247,2 +247,2 @@
-        ok: &mut bool,
-    ) -> Result<(), GateError> {
+    ) -> Result<bool, GateError> {
+        let mut ok = true;
@@ -261 +261 @@
-            ok,
+            &mut ok,
@@ -265 +265 @@
-        Ok(())
+        Ok(ok)
@@ -309,2 +309,2 @@
-        ok: &mut bool,
-    ) -> Result<Option<Answers>, GateError> {
+    ) -> Result<(Option<Answers>, bool), GateError> {
+        let mut ok = true;
@@ -322,2 +322,2 @@
-                check(ok, "load_and_listen", false);
-                return Ok(None);
+                check(&mut ok, "load_and_listen", false);
+                return Ok((None, ok));
@@ -348 +348 @@
-            ok,
+            &mut ok,
@@ -376 +376 @@
-            check(ok, name, pass);
+            check(&mut ok, name, pass);
@@ -379 +379 @@
-        Ok(Some(Answers {
+        Ok((Some(Answers {
@@ -384 +384 @@
-        }))
+        }), ok))
@@ -431 +431,2 @@
-    fn completion_ids(file: &Path, a: &Answers, ok: &mut bool) -> Result<(), GateError> {
+    fn completion_ids(file: &Path, a: &Answers) -> Result<bool, GateError> {
+        let mut ok = true;
@@ -441,2 +442,2 @@
-        check(ok, "completion_ids", prose && turn);
-        Ok(())
+        check(&mut ok, "completion_ids", prose && turn);
+        Ok(ok)
@@ -457,3 +458,5 @@
-        let default_ctx = plan_only_names_its_default(&dir, &file, &mut ok)?;
-        parallel_above_one_is_refused(&dir, &file, &mut ok)?;
-        let answers = served(&dir, &file, default_ctx, layers, &mut ok)?;
+        let (default_ctx, pass) = plan_only_names_its_default(&dir, &file)?;
+        ok &= pass;
+        ok &= parallel_above_one_is_refused(&dir, &file)?;
+        let (answers, pass) = served(&dir, &file, default_ctx, layers)?;
+        ok &= pass;
@@ -461 +464 @@
-            completion_ids(&file, &a, &mut ok)?;
+            ok &= completion_ids(&file, &a)?;

Optimization headroom

• G, per-launch and per-sync fixed cost: crates/gpu/src/arch/nemotron_h/step.rs:377 can reuse the existing activation buffer when Q/K/V read the same quantized format: two fewer quantization launches per such attention layer per token [derived].

Checked

Traced omitted allocations into the serve load record. Confirmed the body's sole caller rejects zero context and R24 covers all four new accumulator parameters. Both proposed patches match the head source. No escape-catalog pattern applies to these findings.

Notes

Static review; GPU gates were not rerun.

_{Card check: gpu-evidence. This pull request changes crates/gpu-gates/Cargo.toml, crates/gpu-gates/src/bin/bloomery_serve.rs, crates/gpu-gates/src/bin/gate_clef_serve.rs … and carries a gate run on the author's card, which counts as verified.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}