[midagedev]

Metal이 Qwen3.8 27B의 MLX MTP sidecar를 본 모델 옆에 열고 draft step 하나를 돌립니다(Metal MTP의 R3; --draft mtp 연결은 공용 계층 정리 뒤 R4에서 합니다).

What and why

generate_metal --mtp SIDECAR opens an MLX MTP sidecar beside the main model. The sidecar is one draft layer: the embedding and hidden norms, fc, a dense attention layer and the shared head. It is a 65th layer outside the plan, with its own buffers and KV cache, so a session without it keeps its plain program.

Session::draft_step runs embed, mtp_join, fc, the layer and the head in one command buffer. --mtp-check holds its logits to tools/mlx-logits.py mtp-dump's reference. This is step R3 of Metal MTP. The decode loop does not call the draft yet. R4 plugs it into the shared draft layer (runtime Verify/Draft, #1135) once that layer is lifted.

Release notes: internal — the Metal backend can load a Qwen3.8 MTP sidecar and check its draft step; decoding does not use it yet.

Code shape

• No second copy. The sidecar reuses the session's mat-vec, attention and head encoders. mtp_join is the one new kernel. The draft policy is not copied from crates/app; R4 uses the lifted common owner.
• Model code holds only architecture facts. The sidecar shape is Qwen3.8's MTP layer. A sidecar of another model is refused by name (hidden size, vocab, tensor set, quant, shape).
• Dead code is gone.
• No silent failure. A NaN or infinite join input raises the norm fault site (clause mtp_join_nan). Bad geometry is refused by name (mtp_join_refused).
• Comments state what is true now.
• No relaxed gate. New clauses only. The draft band is 10 · Δ · the row's largest logit. Δ = 2.8e-4 comes from the BF16 rounding model the e2e gate's bands use, and the derivation is in mtp_check.rs's module doc.
• cargo fmt, just mac-metal-lint clean.

Duplication found or removed: none.

What I ran

just affected origin/main selects 0 box gates. Every changed file is in crates/metal, plus tools/mlx-logits.py, docs/mac.md and Cargo.lock.

Second Mac (M1 Pro, macOS 26), on this branch's head rebased on main 0d6f89c:

test result: ok. 167 passed; 0 failed   (cargo test -p bloomery-metal --lib)
clause mtp_join: PASS (pre and post hidden rows at 5120 and 5632 columns are the composed norms' bits)
clause mtp_join_nan: PASS / clause mtp_join_refused: PASS ; gate_metal_qwen35moe_kernels: PASS (14 instances)
mutant mtp_join_skips_output_norm: FAIL (join post[5120] is -0.09669225, expected -0.06389876)
mutant mtp_join_swaps_halves:      FAIL (join pre[0] is -0.13744682, expected -0.4810251)
mutant mtp_join_swaps_gains:       FAIL (join pre[0] is -0.21071179, expected -0.4810251)
PASS mtp draft logits, hidden pre:  32 positions, worst |metal − reference| 0.0035
PASS mtp draft pick, hidden pre:    29 of 29 clear positions are the reference's pick
PASS mtp draft logits, hidden post: 32 positions, worst |metal − reference| 0.0048
PASS mtp draft pick, hidden post:   28 of 28
INFO mtp tapped: the draft's pick is the session's next pick at 22/32 (pre) and 19/32 (post), equal to the reference
gate_metal_qwen35moe_e2e (qwen38d): PASS

Plain-decode compare against main (tools/metal-compare.sh 0d6f89cb8 --here <this tree>, second Mac):

added    main mtp_join
changed  main headnorm_rows: ... the device instructions are the same (text_sha256 equal)
summary: added=1 removed=0 changed=1 same=380 relinked=0
trace d0:   identical up to buffer numbering: 71 of 109 buffers renumbered ; generated ids: identical (134895,160927,412,163951)
trace d600: identical up to buffer numbering: 83 of 125 buffers renumbered ; generated ids: identical (211607,215232,220750,132615)

• headnorm_rows has its own source unchanged. mtp_join was appended to the same norm.metal, and only the library's metadata sections moved. The device code is equal by text_sha256.
• The renumbered buffers are all weight buffers. At depth 0, every renumbered id is ≤ 87. Each side's renumbered ids are the same set, permuted. Ids ≥ 88 (the arena and the activations) are unchanged.
• The weights load on several threads (arena.rs, in_order), so their allocation order, and therefore the trace's #N, can differ between any two runs. A follow-up makes the trace name a weight buffer independently of the load order, so this compare can return identical (rc 0).

After the first review (b52222d)

The review's four must-fix findings are fixed.

1. ids past u32. The dump reader refuses an id, position, target, pre or post value past u32, by name. Before, 4294967301 read as 5 under a valid checksum. The reader test covers 4294967301 and pre 4294967296. FAIL-first: with the old cast, the id case reads Ok.
2. Sidecar::part uses the checkpoint owner. It now calls hf::modelopt::normalize, as the session's own folder does.
   • Before deleting the copy, a temporary test compared the old and new paths. They were byte-equal on a 4-bit weight with scales and biases (216 B), a BF16 vector (32 B) and a BF16 vector at an odd offset (32 B).
   • A non-finite scale, bias or vector value is now refused by name. The copy accepted it silently.
   • Two tests remain: one pins the A4 layout and the dims, and one checks the three refusals. FAIL-first: the refusal test fails against the old part.
3. The copies of shared helpers are gone. bytes_of, all and fnv1a each had a second copy; each copy was checked equal to its owner first. The e2e gate imports mtp_check::fnv1a.
4. Unused API is removed. has_mtp, mtp_bytes, the two Mtp counters and Dump.target/Dump.picks had no readers.

Device re-run on b52222d (second Mac). Every line equals the run above, including the worst errors, so the sidecar's bytes did not move:

test result: ok. 167 passed; 0 failed
clause mtp_join / mtp_join_nan / mtp_join_refused: PASS ; mutants skips_output_norm / swaps_halves / swaps_gains: FAIL (same values)
PASS mtp draft logits, hidden pre:  worst 0.0035 ; pick 29 of 29
PASS mtp draft logits, hidden post: worst 0.0048 ; pick 28 of 28
INFO tapped 22/32 (pre), 19/32 (post), equal to the reference
gate_metal_qwen35moe_e2e (qwen38d): PASS

The compare as a proof, after #1146

#1146 made the trace's buffer numbering independent of the parallel load. This branch was rebased locally on main 5591bcd, which includes #1146, and compared on the second Mac:

added    main mtp_join
changed  main headnorm_rows: ... (text_sha256 equal, as above)
summary: added=1 removed=0 changed=1 same=380 relinked=0
identical    generated ids: identical (134895,160927,412,163951)
identical    generated ids: identical (211607,215232,220750,132615)
metal-compare: rc 3 (pipelines rc 3, added 1; trace depth ~0 rc 0, depth 600 rc 0; ids differ 0)

Both traces are identical. The plain decode's dispatch is unchanged by this PR; it adds one pipeline.

AI use: Claude Code wrote this change (a Sonnet round, reviewed and run on the device by the lead session).

🤖 Generated with Claude Code

https://claude.ai/code/session_01CjHkkntc2L1xRN21xB39mj

[github-actions]

Verdict: MUST_FIX: the dump reader truncates token IDs; the change also duplicates existing helpers and adds unused API and fields.

Must fix

1. Refuse token IDs that do not fit u32

crates/metal/src/arch/qwen3moe/mtp_check.rs:111 casts parsed u64 values to u32. In --mtp-check, changing a valid dump's ID 5 to 4294967301 leaves its float checksum valid; read returns 5, and check feeds it to draft_step (line 244) and step (line 327). Verification can pass for a different sequence than the header names.
Reject overflow before verification. Add an overflow-ID refusal test. Replace list:

fn list(v: &str) -> Result<Vec<u64>, String> {
    if v.is_empty() { return Ok(Vec::new()); }
    v.split(',')
        .map(|x| x.parse::<u32>().map(u64::from).map_err(|e| format!("{x:?}: {e}")))
        .collect()
}

2. Reuse the affine tensor normalizer (R14)

crates/metal/src/arch/qwen3moe/session/mtp.rs:443 repeats crates/hf/src/modelopt.rs:1804's affine packing. Its public normalize entry also widens BF16 vectors and rejects non-finite scales and biases.
Reuse that owner before the copies diverge. Replace Sidecar::part; delete a4_blocks, vector, drawn and the two repacking tests:

fn part(&self, name: &str) -> Result<Part<'static>, SessionError> {
    use hf::modelopt::{normalize, Plain, Use};
    use hf::quant_config::{Format, Placed};
    let quant = self.st.entry(&format!("{name}.scales")).is_some();
    let weight = if quant { format!("{name}.weight") } else { name.to_owned() };
    let placed = Placed {
        format: if quant { Format::MlxAffine { bits: 4, group: GROUP } } else { Format::Bf16 },
        companions: if quant {
            vec![format!("{name}.scales"), format!("{name}.biases")]
        } else { vec![] },
    };
    let normal = normalize(
        |s| self.st.raw(s), &[(&weight, &placed)], Use::Matrix(Plain::F32),
    ).map_err(|e| mtp_refuse(e.to_string()))?;
    let info = TensorInfo {
        name: name.to_owned(), dims: normal.dims,
        ty: if quant { GgmlType::MLX_A4 } else { GgmlType::F32 },
        offset: 0, nbytes: normal.bytes.len() as u64,
    };
    Ok(Part::made(info, normal.bytes, None))
}

3. Reuse the existing buffer and checksum helpers (R14)

crates/metal/src/arch/qwen3moe/session/mtp.rs:873 and its nested all repeat crates/metal/src/arch/qwen3moe/arena.rs:781 and crates/metal/src/arch/qwen3moe/ops.rs:87. The checksum at crates/metal/src/arch/qwen3moe/mtp_check.rs:51 repeats crates/metal/src/bin/gate_metal_qwen35moe_e2e/mlx.rs:61.
Share these owners before the copies diverge. Expose the existing buffer helpers:

pub(super) fn bytes_of<T: Elem>(b: &MetalBuffer<T>) -> u64 {
    bytes_of_len::<T>(b.pages().len().get())
}

pub(super) fn all<T: Elem>(b: &MetalBuffer<T>) -> Result<Window<'_, T, MetalDevice>, DeviceError> {
    b.window(Elems(0), b.len())
}

Delete mtp.rs's bytes_of and nested all; add:

use super::super::arena::bytes_of;
use super::super::ops::all;

Keep the checksum in mtp_check.rs as:

pub fn fnv1a(bytes: &[u8]) -> u64 {
    let mut h: u64 = 0xcbf2_9ce4_8422_2325;
    for &b in bytes {
        h ^= u64::from(b);
        h = h.wrapping_mul(0x0000_0100_0000_01b3);
    }
    h
}

Delete the gate's checksum definition and import it:

use bloomery_metal::arch::qwen3moe::mtp_check::fnv1a;

4. Remove unused queries and stored fields

crates/metal/src/arch/qwen3moe/session/mtp.rs:685's has_mtp and line 847's mtp_bytes have no callers. Only mtp_bytes reads Mtp.weight_bytes and Mtp.kv_bytes. crates/metal/src/arch/qwen3moe/mtp_check.rs:68's Dump.target and Dump.picks are never read.
Remove these before callers depend on them: delete both query methods, both stored counters, both Dump fields and their initializer entries. Keep open_mtp's returned totals and the reader's pick-list validation.

Checked

Duplicates: read both sides of the affine, byte-count, window and checksum copies; existing owners can replace them without new helpers.
Complexity: inspected reader validation and sidecar loading/dispatch; no finding rests on complexity alone.
Silent failure: followed truncated IDs through the reader to both step consumers.
Tests: reader assertions cover checksum and length refusals; no overflow-ID assertion exists. Inline replacement is valid Rust on an added line.
Gates and exclusions: no escape-catalog pattern or ledger lesson applies; this is a full review.

Notes

Metal gates were not rerun during validation.

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

[github-actions]

Reject integers above u32::MAX before the reader can truncate token IDs.

Suggested change

	.map(|x| x.parse::<u64>().map_err(|e| format!("{x:?}: {e}")))
	.map(|x| x.parse::<u32>().map(u64::from).map_err(|e| format!("{x:?}: {e}")))

[github-actions]

Since the last review:

• Token ID overflow: fixed. Parsing now rejects values outside u32; the reader test covers overflowing IDs and picks.
• Affine packing duplication: fixed. The sidecar delegates to hf::modelopt::normalize.
• Buffer and checksum duplication: fixed. Both callers reuse the existing owners.
• Unused queries and fields: fixed. The queries, stored counters and unused dump fields are removed.

Verdict: CLEAN: all four earlier findings are fixed, and the new commits introduce no must-fix defect.

Checked

Duplicates: 3 new test/helper definitions checked, none duplicated. All 3 dryer pairs read: two context-specific error constructors and one clause-table shape; none warrants another owner.
Complexity: crapper's read, Sidecar::open, open_mtp and encode_draft gain no branches in this delta. Its other 9 flagged functions are unchanged and outside this re-review's scope.
Silent failure: 3 production paths followed. Overflowing IDs return a field-named error before either step consumes them. Invalid affine planes and non-finite vectors return tensor-named errors through open_mtp. Two ignored test-directory cleanup results do not affect verification.
Tests: static mutation check of all 3 changed tests. Restoring the truncating parse at mtp_check.rs:82 makes the overflow assertions fail. Swapping companions at session/mtp.rs:431 fails the block-byte assertions. Restoring unchecked byte copying in Sidecar::part fails the non-finite refusal assertions.
Gates: no band or production clause moved. The replacement tests exercise the Part adapter; the shared normalizer retains its packing and malformed-plane coverage.
Escape catalog: no C1–C4 pattern matches.
Adversarial checks: overflow IDs, non-finite planes, unaligned BF16 vectors and missing companions are blocked. The shared window helper preserves the draft's buffer ranges; no launches or copies are added.

Notes

Metal tests and gates were not rerun in this environment. The PR supplies Mac gate output.

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