metal: grouped-int4 (fmt=4) GEMV support + two latent fmt=4 fixes

[monotophic]

metal: grouped-int4 (fmt=4) GEMV support — the Metal twin of #298/#451

Cross-references: #298 and #451 (CUDA fmt=4 support, read READ-ONLY for
layout/naming conventions — gs/ng field names, "same packed layout as fmt=2, scale
folded into the accumulation" semantics, and the "new ABI entry rather than break the
existing symbol" caution both echoed below). #168 (Team A's int3-g64 reconciliation
— Stage 2 of this same effort, not started here; see §7). Campaign context, same framing
as #447: no maintainer of this repository has Metal hardware — this lab is the test
surface for the Apple GPU tier of every quantization format the CPU/CUDA tiers gain.

Branch: metal/grouped-gemv, based on origin/dev @ e9b3614. Touches
c/backend_metal.h, c/backend_metal.mm, c/colibri.c, c/tests/test_backend_metal.mm.

---

1. The problem: two latent defects that mutually masked each other at stock

Durable claim: dev's grouped-int4 format (fmt=4 — packed int4 nibbles, same layout
as fmt=2, but one f32 scale per gs-element group along I instead of one scale per row)
has a CPU reference (matmul_i4_grouped, c/quant.h:150) and an active CUDA tier
(#298/#451). The Metal GEMV shader (mm_gemv in c/backend_metal.mm) only branched on
fmt == 1|2|3; anything else — including fmt=4 — fell through to the unconditional
else arm and was decoded as raw f32. Taken in isolation, that is not a "CPU
fallback, no GPU acceleration" gap — it is a silent-wrong-answer defect: a code path that
reaches the Metal GEMV with an fmt=4 tensor gets a real, non-error, numerically-wrong
result. That structural claim is unchanged from the first revision of this document, and
the code evidence for it (below) is unchanged too. What changed, on review (credit:
review round 1, auditor — see §10), is the REACHABILITY claim that followed it.

Revised: this defect was NOT live at stock e9b3614. A second, independent defect
masked it, and §2(d) already contained the evidence — this section just failed to draw
the conclusion from it. §2(d) traces the actual allocation site of the fmt=4 scale
buffer: qt_from_disk's fmt=4 branch allocated it with falloc — a bare malloc, never
page-aligned, never coli_metal_register'd — while resolve() in backend_metal.mm
only finds pointers inside registered slabs. That means at stock, resolve(t->s, ...)
on any fmt=4 scale buffer always fails, so resolve_attn/bind_gemv always return
false and the caller always falls back to CPU — for every fmt=4 tensor,
unconditionally, before the shader's fmt=4 fallthrough ever gets a chance to execute.
The two defects mutually cancel at stock: the allocator bug is what kept the shader bug
from ever firing. Stock was accidentally safe, not correctly handled — an important
distinction, covered next — but that is not the same claim as "this was a live landmine
in stock e9b3614," which the first revision of this section overclaimed by tracing only
the shader/dispatch side (bind_gemv's missing fmt gate, quoted below) without also
propagating §2(d)'s own resolve()-fails finding back into this section's reachability
claim.

Where the shader defect would become live, and how this branch's commit ORDER avoids
it. Fix either defect alone, and depending on which one lands first, the other may stop
providing cover before the real fix arrives. If the allocator fix landed before the
shader fix, the intermediate commit would be ARMED: the scale buffer resolves
successfully, but the shader still doesn't know what fmt=4 means, so bind_gemv would
dispatch straight into the f32 fallthrough and silently compute garbage — a state
strictly worse than stock (stock's defect was masked; this one wouldn't be), reachable by
anyone who bisects or cherry-picks that single commit. An earlier revision of this branch
carried exactly that ordering — corrected on review (credit: review round 1, coordinator
— see §10) before this PR was finalized, not left as a residual risk.

This branch's actual, final commit order avoids the armed state by construction rather
than by discipline: f0fa5a9 (the mm_gemv shader fix + all host plumbing, §3/§4)
lands FIRST, and 2ba12d8 (the qt_from_disk allocator fix, §2(d)) lands SECOND. At
f0fa5a9 alone, the new shader branch is already correct, but the allocator bug still
keeps every real fmt=4 scale buffer from resolving — the capability lands inert, not
armed: verified (this branch, this session) that make glm METAL=0/METAL=1 build
clean at that commit and make metal-test still passes 27/27, because metal-test's fmt=4
cases construct their own tensors directly via posix_memalign+coli_metal_register and
never go through qt_from_disk — the test suite exercises and proves the new shader
correct at this commit, it just isn't reachable yet from a real loaded model (§6).
2ba12d8 then activates that already-correct capability by fixing the one thing that was
keeping it unreachable — verified clean/27-27 again at that commit. No commit in this
three-commit sequence is ever unsafe to check out alone: capability lands inert first,
the activation commit completes it. The third commit (fc0f5a5, the attn-test
blind-spot fix, §10) is a test-only strengthening layered on top afterward, with no
bearing on either commit's reachability argument. A reviewer or maintainer bisecting this
branch, or merging a prefix of it, never lands on the armed ordering described above — it
does not exist anywhere in this branch's actual history.

What doesn't change from the prior framing: both defects are real (validator
independently confirmed both with quotes — see §10), both are fixed by this PR, and the
trace below (including the QUOTED gate at colibri.c:1958-1960 showing attention_rows
never checked q_a/q_b/kv_a/o's fmt) is accurate as a description of the CODE'S
STRUCTURE — bind_gemv genuinely has no fmt validity gate, and the shader genuinely
mishandles fmt=4 if reached:

if(g_metal_enabled && !kvs && S<=4 && (g_absorb==1||(g_absorb<0&&S<=4)) && m->kv_start[layer]==0
   && D==6144 && H==64 && c->q_lora==2048 && c->kv_lora==512 && c->qk_nope==192
   && c->qk_rope==64 && vh==256 && l->kv_b.fmt==2){

(the pre-existing gate, git show e9b3614:c/colibri.c lines 1958-1960, unmodified in
shape by this PR — only the argument list changed, see §4). No known field model
currently ships an fmt=4 dense/attention weight next to an fmt=2 kv_b (fmt=4
checkpoints are new), so even the structural gap was latent, not observed-in-the-wild.
What was wrong in the prior revision was conflating "structurally reachable in general"
with "reachable at stock e9b3614 specifically" — the allocator bug traced independently
in §2(d) closes that gap at stock, even though the shader-side structural gap is real and
this PR still closes it (§3) for every path that reaches mm_gemv, not just the one this
PR's test matrix targets, and regardless of whether any known checkpoint currently
exercises it.

2. The trace: how a per-row scale reaches the shader today (quoted before any edit)

Three, and only three, host functions dispatch the mm_gemv kernel; all three pass
scale as buffer(1) and (before this PR) never told the shader anything about grouping:

(a) coli_metal_matmul (backend_metal.mm, the persistent-tensor entry point
metal-test's run()/run_grouped() harnesses call) — wraps weights and scales into
MTLBuffers on first use:

extern "C" int coli_metal_matmul(ColiMetalTensor **tp, float *y, const float *x,
                                 const void *weights, const float *scales,
                                 int fmt, int S, int I, int O) {
  if (!g_dev || fmt < 0 || fmt > 3) return 0;
  ...
      t->w = wrap(weights, t->wbytes);
      t->s = wrap(scales, (size_t)O * sizeof(float));     // <- hardcoded O floats

fmt > 3 was rejected outright, and the scale buffer was always wrapped at exactly O
floats — for fmt=4 the real scale array is O*ceil(I/gs) floats, so even lifting the
fmt>3 guard alone would have UNDER-wrapped the buffer (wrap() sizes the MTLBuffer
from this byte count) and the shader would read past it.

(b) coli_metal_gemm (the large-batch prefill GEMM path, reached from
matmul_qt_ex in colibri.c for S >= g_metal_gemm_min) — resolves already-registered
slabs rather than wrapping, so no buffer-size bug here, but the fmt gate excluded 4
outright and no gs was threaded through:

extern "C" int coli_metal_gemm(float *y, const float *x, const void *wp, const float *sp,
                               int fmt, int S, int I, int O) {
  if (!g_dev || (fmt!=1 && fmt!=2)) return 0;

reached from (colibri.c:421-424, matmul_qt_ex — the general-purpose matvec dispatch
every dense/attention/shared-expert projection goes through):

if(g_metal_enabled && S>=g_metal_gemm_min && !spec_pinned() && (w->fmt==1||w->fmt==2) && !omp_in_parallel()){
    const void *wp = w->fmt==1 ? (const void*)w->q8 : (const void*)w->q4;
    if(coli_metal_gemm(y,x,wp,w->s,w->fmt,S,w->I,w->O)) return;
}

(c) bind_gemv (the internal helper used by the fused single-command-buffer decode
kernels, encode_attention and coli_metal_layer_decode's shared-expert calls) — no fmt
gate at all (§1), and no gs parameter to pass even if it wanted to honor one:

static bool bind_gemv(id<MTLComputeCommandEncoder> e, const void* w, const float* s, int fmt, int I, int O,
                      id<MTLBuffer> xin, id<MTLBuffer> yout, int S){
  ...
  [e setBytes:&S ...]; [e setBytes:&I ...]; [e setBytes:&O ...]; [e setBytes:&fmt ...];
  [e setBytes:&NT ...];

reached, per weight, from colibri.c's attention_rows/layer_forward_rows via a
WP_(q) macro that already correctly resolves fmt=4 tensors to their .q4 pointer
(dev's own CPU-side dispatch does the same — qt_from_disk stores grouped weights in
t->q4, confirmed at colibri.c:824):

#define WP_(q) ((q).fmt==1?(const void*)(q).q8:(const void*)(q).q4)
int ok = coli_metal_attn_decode(x,
    WP_(l->q_a), l->q_a.s, l->q_a.fmt, l->q_a_ln, ...

(d) One more link in the chain, found by tracing to the actual allocation site, not
just the dispatch call sites: qt_from_disk (colibri.c:823-825, unmodified in shape
by this PR except the one-line fix in §4) allocates the fmt=4 scale array with plain
falloc — a bare malloc, not page-aligned and never coli_metal_register'd — while
every other fmt uses qsalloc (qalloc under the hood), which page-aligns and registers
under COLI_METAL:

else if(fmt==4){ int ng=(I+gs-1)/gs;
    if(t->fmt!=4||!t->q4){ t->fmt=4; t->O=O; t->I=I; t->gs=gs; t->q4=qalloc(nb); t->s=falloc((int64_t)O*ng); }

resolve() in backend_metal.mm only finds pointers inside coli_metal_register'd
slabs. A scale buffer allocated with falloc can never be found by resolve(), so
bind_gemv/coli_metal_gemm would return false/0 for it — safe (CPU fallback,
no wrong answer), but it means fixing (a)-(c) alone would not have mattered in practice:
every real fmt=4 dense/attention tensor's scale buffer would have permanently missed the
GPU path. Found while tracing the per-row-scale path this stage extends to per-group;
fixed alongside it (§4) since it is the literal allocation site of the buffer this whole
PR is about.

3. The shader change: per-group scale folded into the accumulation, not applied once

Unlike fmt 1-3 (y[row] = acc * scale[o], one multiply after the full dot product),
fmt=4's group scale must be applied per group, before summing across groups — this is
matmul_i4_grouped's actual math (quant.h:150-184), not a simplification of it. The new
fmt==4 branch in mm_gemv:

} else if (fmt == 4) {            // int4 GROUPED: same nibble packing as fmt=2,
                                   // one f32 scale per gsz-element group along I.
  int rb = (I+1)/2; int ng = (I+gsz-1)/gsz;
  device const uchar* wr = w + (long)o * rb;
  device const float* scl = scale + (long)o * ng;
  for (int i = slane*2; i < I; i += 64) {
    uchar b = wr[i>>1];
    int g0 = i / gsz; float sc0 = scl[g0];
    acc += float(int(b&0xF)-8) * xr[i] * sc0;
    if (i+1 < I) {
      int g1 = (i+1) / gsz; float sc1 = (g1==g0) ? sc0 : scl[g1];
      acc += float(int(b>>4)-8) * xr[i+1] * sc1;
    }
  }
} else { ... }
...
acc = simd_sum(acc);
if (slane == 0) y[row] = (fmt == 4) ? acc : acc * scale[o];   // fmt=4: don't scale again

Each of the 32 SIMD lanes owns one packed byte (2 elements) per 64-wide stride — memory
access stays coalesced (wr[i>>1] is contiguous across lanes), and a group never splits
a byte because gs is always even (dev's own loader only accepts gs from the candidate
list {16,32,48,64,96,128,192,256}, colibri.c:777). This is intentionally the
straightforward per-element form, not vectorized like fmt=2's uchar4/float4 path —
see §7 (Not done) for why that tradeoff was made deliberately, not by oversight.

A new constant int& gsz [[buffer(9)]] parameter carries the group size; it is ignored
for every other fmt (callers pass the tensor's QT.gs, which is 0 for non-grouped
tensors — harmless, since the shader only reads it when fmt==4).

4. The host-side plumbing: what changed at each traced site

• fmt_bytes gained an fmt==4 case (same packed-nibble byte count as fmt=2 — the
  weight layout is identical, only the scale layout differs, matching cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298's "row_bytes:
  fmt=4 = same packed int4 layout as fmt=2" note). New fmt_scale_bytes(fmt,I,O,gs)
  computes O*ceil(I/gs)*sizeof(float) for fmt=4, O*sizeof(float) otherwise — this is
  what coli_metal_matmul now uses to size the wrap() call in §2(a), fixing the
  under-wrap.
• coli_metal_matmul, coli_metal_gemm, bind_gemv all gained a gs
  parameter (appended at the call site, inserted right after the paired fmt in the
  fused-decode signatures — mirrors cuda: grouped-int4 (fmt=4) support in the expert-group kernels — opens the GPU tier to g64 and E8 containers (#334) #451's convention of adding gs fields alongside
  fmt), and their fmt gates now accept 4. coli_metal_matmul's doc comment in
  backend_metal.h states the new contract explicitly.
• AttnW (the per-layer weight-pointer bundle used by encode_attention) gained
  qa_gs/qb_gs/kva_gs/o_gs fields. coli_metal_attn_decode and
  coli_metal_layer_decode gained the matching parameters, threaded through to their
  bind_gemv calls (attention projections q_a/q_b/kv_a/o, plus shared-expert
  gate/up/down in the layer-decode variant).
• colibri.c call sites: matmul_qt_ex's Metal-GEMM gate now includes w->fmt==4
  and passes w->gs; attention_rows/layer_forward_rows pass .gs alongside each
  .fmt already being passed (the field defaults to 0 for non-grouped tensors, so this
  is a no-op for every model that isn't using fmt=4 anywhere).
• qt_from_disk: the one-line allocator fix from §2(d) — t->s for fmt=4 now comes
  from qalloc (page-aligned + coli_metal_register'd) instead of falloc.

kv_b is deliberately untouched. It never flows through bind_gemv/mm_gemv at
all — the MLA absorption kernels (a_qabs, a_ctx) dequantize it inline with a
hand-written, per-row-only helper (a_deqrow, backend_metal.mm), a structurally
separate kernel family from the GEMV this PR extends. The existing l->kv_b.fmt==2 gate
(§1) is left exactly as-is — not relaxed — specifically so an fmt=4 kv_b continues to
CPU-fallback safely rather than hitting a_deqrow's per-row-only math. See §7.

coli_metal_moe_block/moe_gemv (the batched routed-expert path, a third,
independent kernel family from mm_gemv) is also untouched. It was already safely gated
to fmt!=1 && fmt!=2 -> return 0/nil before this PR — fmt=4 experts already CPU-fall-back
correctly today, with no wrong-answer risk, so this was a coverage gap, not a
correctness bug, and out of scope for this stage. See §7.

5. What was verified, and how (RAN, this branch)

metal-test follows the existing run()/inline-gemm-test pattern the spec named,
extended with an fmt=4-specific CPU reference (cpu_ref_grouped) and harness
(run_grouped) — cpu_ref's [O] per-row scale layout cannot express fmt=4's
[O,ceil(I/gs)] layout, so it is a sibling reference, not a branch of the existing one.
cpu_ref_grouped mirrors quant.h's matmul_i4_grouped nibble-decode and group-index
math exactly, accumulating in double — the same construction
tests/test_i4_grouped.c (dev's own CPU-side fmt=4 oracle, unmodified by this PR)
already established as the "known exact" reference for this format.

Tolerance, stated and justified against the existing convention: the pre-existing
run() cases compare max|Δ| / max|y_ref| (ymax-relative) at 1e-4. The new
run_grouped/gemm-grouped/attn-grouped cases instead compare, per output element,
|Δ| / Σ|term_i| (magnitude-relative, summing |xs[i]*(nib-8)*scale| over every term of
that dot product) at the same 1e-4. Reason: a grouped dot product over many
groups/signed terms can land near zero from cancellation — an ordinary fixed
GPU-vs-double rounding difference then reads as an enormous relative error against a
near-zero result, which is the accumulator's precision, not a kernel defect.
Comparing against the summed magnitude of the terms instead means a wrong scale index or
wrong group boundary is still caught as an O(1) relative error (it shifts the result by a
fraction of the terms actually summed), while cancellation noise is not misreported as a
bug. This is exactly the convention tests/test_i4_grouped.c uses for the same kernel's
CPU-side oracle (see its ref_grouped/mag comment, unmodified by this PR) — reused
here for a GPU oracle of the same format, not invented fresh. 1e-4 (vs.
test_i4_grouped.c's 1e-6) carries the same slack the existing run() cases already
accept for GPU-vs-CPU float comparisons in this file, because the GPU sums in a different
order (byte-pair-strided across 32 SIMD lanes, then simd_sum) than the scalar double
reference.

Shape matrix (spec §"Validation"), all via coli_metal_matmul unless noted:

Case	Shapes	Evidence
I multiple of gs=64	gate/up I=6144, down I=2048 (real g64-checkpoint dims), S=1 and S=4	RAN, ok
I NOT a multiple of gs=64	I=200 (even), I=201 (odd, exercises the nibble tail)	RAN, ok
I≤64 degenerate (single group)	I=64 (ng=1, exact), I=40 (<gs, single partial group)	RAN, ok
S=1 and S>1	covered throughout (S∈{1,2,3,4})	RAN, ok
Random + outlier-heavy rows	one activation forced to 50.0 in group 0, I=512 and I=201, verifying grouped reconstruction end-to-end through the GPU path (not just that the kernel runs)	RAN, ok
Large-batch GEMM path (coli_metal_gemm)	O=2048,I=6144,S=64,gs=64 — the second of the two directly-named GEMV entry points	RAN, ok
Fused-decode path (bind_gemv via coli_metal_attn_decode)	q_a set to fmt=4/gs=64 inside the real fused attention command buffer (S=1 pos=37, S=4 pos=12/MTP — not pos=0, see §10: at T=1 softmax is provably independent of q_a). Each case additionally compares q_a's raw pre-RMSNorm GEMV output against the CPU oracle directly (qraw metric), closing RMSNorm's blindness to a uniform scale error — proves the plumbing in §4 end-to-end, not just the standalone kernel, and not just modulo RMSNorm's invariances	RAN, ok
Negative control: fmt 1/2/3 unchanged	all 8 original run() cases + 2 moe_block + 1 gemm(int4) + 4 attn cases, unmodified assertions	RAN, ok (identical nerr values to the pre-PR baseline run, see §6)

fmt=5 (stage 2): not attempted — blocked on Team A's int3/g64-rebased semantics, see §7.

6. Gates run on this branch (all RAN, 2026-07-19, macOS 26.5.2, Apple M5 Max 128 GB)

• make glm METAL=0 and make glm METAL=1: clean, 0 warnings (-Wall -Wextra), both before this PR's changes (stock e9b3614) and after — verified by
  capturing both build logs and grepping for warning (0 hits in all four: stock×2,
  branch×2).
• make metal-test: before this PR (stock e9b3614): 15/15 cases pass.
  After this PR: 27/27 cases pass — the original 15 unchanged (byte-identical nerr
  values, the negative control from §5), plus 12 new fmt=4 cases: 9
  run_grouped() shape-matrix cases, 1 grouped coli_metal_gemm case, 2
  run_attn_grouped() fused-decode-path cases. Case COUNT unchanged at 27/27 after
  review round 1 (§10) — the two run_attn_grouped() cases were strengthened in place
  (pos_base moved off the degenerate T=1 point; each case gained an internal raw
  pre-RMSNorm comparison against the CPU oracle, printed as qraw=...), not added
  alongside as new cases. See §10 for the mutation-detection matrix proving the
  strengthened cases actually catch what the original two could not.
• make test-c (all suites in TEST_BINS, including tests/test_i4_grouped — dev's
  own CPU-side fmt=4 oracle, unmodified by this PR): exit 0, 0 failures reported by any
  suite (e.g. test_topp: 123 cases run, 0 failure(s), test_dsa_select: 129 cases run, 0 failure(s), test_i4_grouped: ok).
• make test-python: 83/83 tests, OK, exit 0.
• git diff e9b3614 --numstat (this branch, after review round 1): 4 files —
  c/backend_metal.h (+25/-18), c/backend_metal.mm (+70/-36), c/colibri.c (+21/-14),
  c/tests/test_backend_metal.mm (+245/-5). No file outside c/backend_metal.{h,mm},
  c/colibri.c, and its own test file was touched.

7. Scope statement: what this PR does NOT do (read before assuming full coverage)

• coli_metal_moe_block/moe_gemv (batched routed-expert MoE path) does not gain
  fmt=4 support. It was already safely gated to CPU-fallback for fmt=4 before this PR
  (§4) and remains so — extending it means a second, independent per-group-scale kernel
  branch (moe_gemv is bindless/pointer-array-based, structurally different from
  mm_gemv) plus new per-expert gs array parameters on coli_metal_moe_block/
  coli_metal_moe_block_begin. Given fmt=4 is primarily an expert format in practice
  (per cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298's framing), this is real follow-up work, not a footnote — flagged here rather
  than attempted under this task's time/scope, with the header doc comment in
  backend_metal.h updated to say so explicitly.
• kv_b (MLA absorption core: a_qabs/a_ctx) does not gain fmt=4 support. These
  hand-written kernels dequantize inline with a per-row-only helper (a_deqrow) that
  never touches mm_gemv/bind_gemv — extending it is a different, larger shader change
  (new a_deqrow_grouped plus threading gs through a_qabs/a_ctx's own buffer
  lists), and the existing l->kv_b.fmt==2 gate (§1, §4) already keeps this safe by
  construction: an fmt=4 kv_b simply never reaches this GPU path. Not relaxed, not
  attempted.
• Stage 2 (fmt=5, int3-g64) was not started. Per the governing spec: build stage 1
  fully first if Team A's int3/g64-rebased reconciled semantics aren't yet available.
  This worker did not check out or read that branch — a finished stage 1 beats a
  speculative stage 2, and stage 2's addendum is explicitly gated on int3-g64 (fmt=5): per-group-scale 3-bit weights, the size/quality point your own ablation (#132) says the next container should sit at #168's merge, so
  starting it now would mean guessing at semantics likely to change. See UNCERTAINTIES.
• The shader arithmetic is not vectorized the way fmt=2's uchar4/float4 path is.
  fmt=4's per-group scale lookup means a byte's two nibbles can belong to different
  groups (g0/g1 in §3), which doesn't compose cleanly with the 8-wide vector loads
  the other formats use without either restricting gs to specific alignments beyond
  what the loader already guarantees, or adding a second code path for the aligned case.
  Given the spec's own framing ("the shader arithmetic is small" — correctness is the
  ask, not throughput) and no model-run performance measurement being permitted under
  this task's constraints anyway, the straightforward per-element form was chosen
  deliberately. A vectorized fast path for gs-aligned regions is natural follow-up work
  once real g64-checkpoint throughput numbers exist to justify it.
• No CUDA files were touched. cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298/cuda: grouped-int4 (fmt=4) support in the expert-group kernels — opens the GPU tier to g64 and E8 containers (#334) #451 were read read-only for convention only.

8. DEVIATIONS from the spec

• Spec's validation matrix names shapes generically ("I multiple of 64... I≤64
  degenerate"); this PR's run_grouped calls hardcode gs=64 throughout (the real
  g64-checkpoint value, per cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298's title and description) rather than also sweeping
  gs∈{16,32,...,256} the loader's detect_group_size candidate list supports. The
  shader itself is fully generic over gs (nothing in §3 assumes 64 specifically, only
  evenness) — this is a test-coverage choice, not an implementation limitation. Flagged
  as a deviation because the spec's shape matrix could be read as wanting the full
  candidate-list swept; it was not.
• Added a bind_gemv/fused-decode-path integration test (run_attn_grouped, §5) that
  the spec did not explicitly request (its validation section names coli_metal_matmul-
  style GEMV comparison and the gemm-test pattern). Added because §2/§4's plumbing work
  (AttnW, coli_metal_attn_decode/coli_metal_layer_decode signature changes) would
  otherwise be exercised by nothing — an untested code path in a PR whose entire premise
  is "a wrong PASS is the worst possible output" seemed like the wrong tradeoff to leave
  unaddressed even though not explicitly named.
• Fixed the qt_from_disk allocator bug (§2(d), §4) even though it lives outside
  c/backend_metal.* and is arguably "dev's pre-existing code, not this PR's shader
  work." Kept in-scope because it is the literal allocation site of the buffer this
  entire PR's plumbing traces and extends — leaving it unfixed would make §3/§4's shader
  and host-side work unreachable by any real fmt=4 dense tensor.

9. UNCERTAINTIES

• Whether any current field checkpoint actually has fmt=4 on q_a/q_b/kv_a/o/shared-
  expert weights (as opposed to only MoE routed experts) is unknown to this worker.
  §1's landmine and §4's plumbing are real given the loader's generic qt_resolve_fmt
  (any tensor loaded via qt_from_disk can be detected as fmt=4), but no specific
  checkpoint was inspected — this task ran no model and downloaded nothing. If in
  practice fmt=4 is expert-only today, §7's moe_gemv gap matters more in the field than
  this PR's bind_gemv fix does; both were done regardless, since the spec asked to trace
  "wherever the current per-row scale buffer flows into the shader," not to guess which
  path is hottest in practice.
• Real-model throughput/quality impact is entirely unmeasured. No model runs were
  permitted under this task's constraints; every number in §5/§6 is a synthetic
  kernel-vs-reference comparison. Whether fmt=4 dense tensors reaching the Metal GEMV
  produces a measurable decode-speed change (vs. today's CPU fallback) is UNVERIFIED.
• gs values outside {16,32,48,64,96,128,192,256} (the loader's candidate list) were
  not tested, though the shader and host plumbing place no restriction on gs beyond
  evenness. Untested does not mean unsupported, but it is UNVERIFIED beyond the tested
  set (which does include 64, the only value cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298/cuda: grouped-int4 (fmt=4) support in the expert-group kernels — opens the GPU tier to g64 and E8 containers (#334) #451's own checkpoints use).
• Stage 2 (fmt=5) semantics, timeline, and int3-g64 (fmt=5): per-group-scale 3-bit weights, the size/quality point your own ablation (#132) says the next container should sit at #168's state were not checked beyond
  confirming the branch name int3/g64-rebased exists in the spec's own text — this
  worker did not fetch or read it, per §7.

10. Review round 1 (credit — house style: the review record is part of the PR's evidence)

Three findings from independent review, all FIX-FIRST, all addressed on this branch. The
final commit sequence — f0fa5a9 (shader + host plumbing + original tests),
2ba12d8 (allocator activation), fc0f5a5 (attn-test blind-spot fix) — already
reflects all three; there was an earlier, different commit sequence during review that
the third finding below corrected, but it was never pushed and left no trace other than
the two SHAs (142ddd2/9315efa/5e3acd8) some of the discussion below still names for
context. git diff 5e3acd8 HEAD is empty and the two states' tree hashes are identical
(95ff3e6f...) — the reorder changed commit BOUNDARIES and ORDER only, not one byte of
final content.

Auditor: traced that §1 (first revision) overclaimed reachability at stock — the
falloc/resolve() evidence already present in §2(d) shows the allocator bug masks the
shader bug at stock e9b3614, so the two defects mutually cancel there rather than the
shader bug being independently live. §1 above is rewritten to reflect this; nothing about
either defect's reality was retracted, only the "live landmine in stock" framing.

Validator: found two proven blind spots in the run_attn_grouped() fused-decode-path
cases (§5): (a) the original S=1 case used pos_base=0, giving T=pos_base+S=1 — softmax
over a single key is identically 1.0 regardless of the score's value, so that case's
final output was provably independent of q_a's kernel entirely, not merely insensitive to
some subset of defects; (b) both cases fed q_a's output through RMSNorm before comparing
anything downstream, and RMSNorm(c·v) == RMSNorm(v) for any positive scalar c — a
whole-tensor uniform-scale-calibration bug in q_a's kernel would be completely invisible
to the fused-path comparison (nerr/cache) no matter how pos_base was chosen, since
(b) is independent of (a).

Fixes applied (both in tests/test_backend_metal.mm, third commit): the S=1 case now
uses pos_base=37 (mirroring run_attn's own non-grouped S=1 pos=37 case, which exists
for the identical reason); both cases now additionally capture q_a's raw GEMV output
before t_rms overwrites it in place, and compare that raw output directly against the
CPU oracle (cpu_ref_grouped, the same magnitude-relative construction run_grouped()
uses) via a standalone coli_metal_matmul call on the exact weight/scale/x data the
attention test generated — reported as qraw=... in the test's printed line, folded into
that case's overall pass/fail (not a separately counted case; §6).

Coordinator (final micro-round): read §1's own rewritten atomicity paragraph (the one
correcting the auditor's finding above) and noticed it proves too much against itself —
§1 stated the armed state ("allocator resolves, shader doesn't handle fmt=4") is reachable
by bisecting or cherry-picking a single commit, and this branch's THEN-current order
(allocator fix landing before the shader fix — the original 142ddd2→9315efa
sequence) put exactly that armed state at the first of the two commits: green build,
15/15 metal-test, and armed, checked out alone. Directed a reorder to shader-first so
the hazard §1 itself describes cannot occur anywhere in this branch's history, not just
in the branch's final (squashed-view) state. This is the reorder reflected in the commit
list at the top of this section and in §1's current text.

Verification of the reorder (RAN, this branch, this session): each of the three
resulting commits was built and tested standing alone (git reset/re-edit to that
commit's exact tree, not merely inspected) before being committed:

Commit (final SHA)	Content	METAL=0/METAL=1 build	make metal-test
f0fa5a9 (1st)	shader fmt=4 branch + host plumbing + original (blind-spotted) tests	clean, 0 warnings both	27/27 — fmt=4 cases construct tensors directly (posix_memalign+coli_metal_register), never through qt_from_disk, so none depend on the allocator fix landing first; the new shader is exercised and proven correct here, just not yet reachable from a real loaded model
2ba12d8 (2nd)	+ qt_from_disk allocator activation	clean, 0 warnings both	27/27 (unchanged from 1st — metal-test's cases never touched qt_from_disk, so this commit doesn't change metal-test's outcome, only real-model reachability)
fc0f5a5 (3rd)	+ attn-test blind-spot fix	clean, 0 warnings both	27/27 (qraw metric now present; see the mutation matrix below)

No test in the 12 new fmt=4 cases needed to become "self-allocating" or move commits —
they already self-allocate (they always have; qt_from_disk is a colibri.c model-load
function metal-test's binary doesn't even link against — make metal-test's build recipe
compiles only tests/test_backend_metal.mm + backend_metal.mm, never colibri.c).
Tree identity: git diff 5e3acd8 HEAD — empty; git rev-parse 5e3acd8^{tree} and
git rev-parse HEAD^{tree} — identical (95ff3e6f...). The reorder is a pure permutation
of the same three diffs across commit boundaries, verified byte-for-byte, not a
re-derivation that could have silently changed something.

Verification (RAN, this branch, this session): re-derived the validator's three
mutation classes by temporarily editing mm_gemv's fmt==4 branch in backend_metal.mm
(reverted after each run; final backend_metal.mm is byte-identical to the pre-mutation
state — confirmed with diff), and ran four configurations of a temporary diagnostic
harness (run_attn_grouped_diag, not committed) isolating fix (a) from fix (b):

Mutation	OLD: S=1 pos=0, no qraw check	fix (a) only: S=1 pos=37, no qraw	fix (b) only: S=1 pos=0, qraw	fix (a+b) — shipped: S=1 pos=37, qraw
Off-by-one group index (g = i/gs, biased down by 1 when >0)	not caught	caught (nerr 4.6e-6→1.34)	caught (qraw 1.8e-8→9.4e-2)	caught
Swapped nibbles (b&0xF/b>>4 swapped)	not caught	caught (nerr→2.38)	caught (qraw→9.6e-2)	caught
Uniform 1% scale error (scale *= 1.01)	not caught	not caught (nerr 5.78e-6, still passes)	caught (qraw 1.8e-8→6.4e-4)	caught

Reading the table: the OLD column (validator's claim) caught nothing, confirming the
blind spot was real, not theoretical. Fix (a) alone closes the off-by-one and
swapped-nibble classes but — exactly as the RMSNorm-invariance argument predicts —
still misses the uniform-scale class even at pos_base=37, empirically confirming that
(a) and (b) are independent fixes and neither subsumes the other. Fix (b) alone (the
qraw check) catches all three classes regardless of pos_base, because it inspects
q_a's kernel output directly rather than through the RMSNorm/attention pipeline; fix (a)
remains independently worth keeping because it makes the FUSED-PATH comparison itself
(not just the auxiliary qraw probe) sensitive to q_a, which matters for defects in the
integration (e.g. how qa_gs is threaded through bind_gemv's buffer indexing inside
encode_attention) rather than in the standalone GEMV kernel qraw isolates. The shipped
configuration (rightmost column) is what run_attn_grouped() now runs unconditionally —
3/3 mutation classes caught, make metal-test still exits 0 (27/27) with none of the
three mutations present. Also noted in passing: run_grouped()'s own shape-matrix cases
(§5), un-related to this attn-level fix, independently caught the uniform-scale mutation
too (worst_rel 7.1e-4 to 3.5e-3 across shapes, above the 1e-4 threshold) — the attn-level
blind spot was specifically about the FUSED attention output hiding the bug via RMSNorm,
not about fmt=4 GEMV correctness being unchecked elsewhere.

11. CUDA platform parity (#298 / #451) — added 2026-07-20

A structured comparison against the in-flight CUDA fmt=4 work (PR #298's design +
PR #451's expert-group kernels, both read at head on 2026-07-20) so the two platform
tiers can be reviewed as companions:

• Format contract: fully aligned. Same packed-nibble layout ((I+1)/2 row
  bytes), same [O, ceil(I/gs)] f32 scale shape with scale[o*ng+g] indexing,
  same g = i/gs group math. The one apparent decode divergence (CUDA's
  upload-time XOR-0x88 to two's-complement vs our direct offset-binary -8) is
  mathematically equivalent — and deliberately not interchangeable: CUDA mutates a
  device-private copy, while Metal buffers can alias host memory the CPU fallback
  still reads, so an in-place transform is not an option on this backend.
• Test oracle: our cases use the magnitude-relative tolerance
  (|Δ| / Σ|terms|), matching dev's own test_i4_grouped.c and the review
  guidance given on cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298 regarding cancellation in grouped dot products.
• Scope split: this PR covers the mm_gemv family (dense / attention /
  fused-decode). The batched routed-expert path (moe_submit/moe_gemv) remains
  gated to fmt {1,2} with CPU fallback — the Metal counterpart of cuda: grouped-int4 (fmt=4) support in the expert-group kernels — opens the GPU tier to g64 and E8 containers (#334) #451's
  expert-group work is planned as a follow-up ("stage 3"), for which cuda: grouped-int4 (fmt=4) support in the expert-group kernels — opens the GPU tier to g64 and E8 containers (#334) #451's
  per-tensor GroupDesc (per-expert fmt/gs) is the reference shape; Metal's
  current batch path assumes a single scalar fmt per block, which stage 3 must
  make explicit or generalize.
• End-to-end gate: all results in this PR are synthetic kernel-vs-oracle
  comparisons (§5–6). A CPU-vs-Metal token-stream identity check on a real
  grouped-int4 container (the determinism gate requested on cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298) is planned as a
  follow-up comment once a g64 container is minted from the FP8 source.

---

Coordinate stamp: 2026-07-19/20 (session spanned midnight), macOS 26.5.2, Apple M5 Max
128 GB (Mac17,7) — the same machine/lab referenced in #447. All RAN evidence in this
document (§5, §6, §10) was captured on this branch, this session, this machine. No model
runs; no GitHub writes; no pushes.

[JustVugg]

Rebase needed — and one substantive note: #298 just merged into dev, fixing the CUDA dense/attention kernels to apply fmt=4 scales per-group (absorb_scale, gs/ng params). Worth checking your two latent fmt=4 fixes against it: one of them may be the Metal twin of the same bug (per-row indexing of grouped scales), in which case your GEMV can mirror #298's approach 1:1. The conflict is semantic in the good way — your call, but you clearly know this area.

[monotophic]

Authored by Fable 5 in Claude Code, analysis in partnership with @monotophic.

Rebase: done — Current head 2277004
sits on dev @ 68ac9ff (post-#298/#192/#421), mergeable, all three commits re-verified on
the new base (43/43 metal-test per commit, zero warning delta, incl. a COLI_METAL_RESSET=1
run now that #426 is in the base).

On the substantive question — we checked against absorb_scale at source, and the picture is:

The GEMV scale math already mirrors #298 exactly. absorb_scale:
wscale[row*ng + k/gs], ng = ceil(K/gs). Our mm_gemv fmt=4 branch:
scale[o*ng + i/gsz], same ng. Only difference is your defensive tail clamp
(if (g >= ng) g = ng-1) vs our bound-by-construction (i < I ⟹ i/gsz ≤ ng-1, ceiling-division
identity — the review round proved it algebraically and the partial-tail test case exercises it).

Our two latent fixes are a different defect class than the one #298 fixed, though. They're
(1) the fmt=4 scale buffer being allocated outside the GPU-visible registration path
(falloc→qalloc — Metal zero-copy machinery, no CUDA analog), and (2) the dispatch falling to the
f32 arm. Neither is scale indexing.

The true Metal twin of your per-row-indexing bug is the attention-absorption path
(a_deqrow/a_qabs/a_ctx) — and this PR deliberately doesn't extend it: the pre-existing
l->kv_b.fmt==2 gate keeps grouped kv_b off the Metal absorb fast path entirely, so the bug is
unreachable rather than fixed on our side (disclosed in the PR body's scope notes). When Metal
absorption gets extended to grouped kv_b, absorb_scale is exactly the 1:1 blueprint we'd
mirror — happy to take that as a follow-up once this lands, since #298 has now settled the
reference semantics.