open-standards/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/BENCHMARKS.md

# jouleclaw-loader-gguf — measured numbers

Real-model numbers captured by running `cargo run --release
--example load_model` against models on an external SSD, plus
`cargo bench`. Hardware: Apple Silicon dev box (M-series).
Storage: USB-C external SSD (~3-4 GiB/s sequential read).

These numbers are reproducible — the bench harness is in
`benches/`, the loader example is in `examples/load_model.rs`,
the models are in any standard GGUF distribution. Numbers
shift with: storage class (NVMe vs HDD), page-cache state
(cold vs warm), and host kernel. Captured here as a checkpoint
the substrate's behaviour can be compared against.

## Load + advise wall-clock vs `PrefetchPolicy`

Same 9.5 GiB ZAYA1-8B Q8_0 GGUF, three policies:

| Policy | Load time | What ran |
|---|---:|---|
| `PrefetchPolicy::None` | **28 ms** | mmap + GGUF header parse only |
| `PrefetchPolicy::TopN(32)` | **139 ms** | header + WillNeed on top-32 by byte span + per-tensor hints on all 1283 |
| `PrefetchPolicy::All` (default) | **2,268 ms** | header + WillNeed on full 8.82 GiB region + per-tensor hints on all 1283 |

The `All` policy at 2.27 s corresponds to a sustained ~3.9 GiB/s
read rate from external SSD — close to the storage's
sequential ceiling. The substrate is correctly pushing the OS
to read pages eagerly.

The `TopN(32)` policy delivers most of the cold-load TTFT win
at **6% of the All policy's cost**. For latency-sensitive
paths (interactive inference, MoE expert swap), `TopN` is the
right choice when you know which tensors are first-touch hot.

Per-tensor advise overhead = (139 - 28) ms / 1283 tensors = **~87 µs per tensor**,
including syscall + byte-span computation.

## Load wall-clock by model size + format

All with `PrefetchPolicy::All` (default; full advise pass):

| Model | Format | Size on disk | Load + advise |
|---|---|---:|---:|
| LFM2.5 350M | Q4_K_M | 219 MiB | 53 ms |
| Bonsai 1.7B | Q1_0 (custom 1-bit) | 237 MiB | 67 ms |
| LFM2 350M | Q8_0 | 362 MiB | 93 ms |
| Ternary Bonsai 1.7B | Q2_0 (custom ternary) | 416 MiB | 113 ms |
| TinyLlama 1.1B chat | Q8_0 | 1.1 GiB | 254 ms |
| BitNet b1.58 2B | I2_S | 1.10 GiB | 285 ms |
| Gemma 4 E2B text | Q8_0 | 4.60 GiB | 1,081 ms |
| ZAYA1 8B | Q8_0 | 8.82 GiB | 2,268 ms |

Linear scaling with byte count — the per-tensor hint pass
dominates large-model load.

## Access-hint classifier — real-model behavior

The wave-3 classifier picks `Random` for embeddings,
`Sequential` for matmul weights, `WillNeed` for norms. Counts
from real models:

| Model | Random | Sequential | WillNeed | Total |
|---|---:|---:|---:|---:|
| LFM2.5 350M Q4_K_M | 2 | 72 | 74 | 148 |
| LFM2 350M Q8_0 | 2 | 72 | 74 | 148 |
| Bonsai 1.7B Q1_0 | 1 | 196 | 113 | 310 |
| TinyLlama 1.1B Q8_0 | 2 | 154 | 45 | 201 |
| BitNet b1.58 I2_S | 1 | 210 | 121 | 332 |
| Gemma 4 E2B Q8_0 | 2 | 245 | 354 | 601 |
| ZAYA1 8B Q8_0 | 1 | 360 | 922 | 1,283 |

`Random` count (1 or 2) tracks the number of distinct lookup
tables (token_embd + optional output projection). `Sequential`
tracks the matmul weights per block × block count.
`WillNeed` is everything else — small + always-touched.

## Matmul throughput — Q8_0 scalar vs NEON

`cargo bench --bench matmul_q8_0`. Apple Silicon dev box,
release build.

**`matmul_q8_0_scalar`** — scalar reference + wave-6 prefetch
hint only. Takes pre-quantized activation; no fp32 → i8
overhead.

**`matmul_q8_0`** — dispatching entry point. On aarch64 calls
the NEON kernel (`matmul_q8_0_neon`); also pays for
`quantize_a_per_row` (fp32 activation → i8 + per-row scale)
that the scalar bench skips. Closer to production cost.

| Shape (m × n × k) | Scalar | NEON dispatch | Speedup | NEON GFLOPS |
|---|---:|---:|---:|---:|
| 1 × 256 × 256       | 11.1 µs | **3.6 µs** | 3.1× | 36.4 |
| 1 × 1024 × 1024     | 181 µs  | **54 µs**  | 3.4× | 38.6 |
| 1 × 2048 × 2048     | 759 µs  | **218 µs** | 3.5× | 38.5 |
| 1 × 4096 × 4096     | 3.00 ms | **868 µs** | 3.5× | 38.6 |
| 8 × 1024 × 1024     | 1.42 ms | **432 µs** | 3.3× | 38.8 |

NEON path sustains **~38 GFLOPS** regardless of size — the
size-independent scaling holds for both kernels (wave-6
prefetch + the NEON kernel's own vld1q-loop prefetching keep
the L1 fed past the L2 ceiling).

Context: Apple Silicon's fp32 peak per core is ~200 GFLOPS for
pure fp32 MAD. The NEON Q8_0 kernel's effective rate is
bounded by the int8 widen-multiply + i32 accumulate sequence
(`vmull_s8` → `vpadalq_s16`), not pure fp32 MAD. 38 GFLOPS
against that effective ceiling is **~20%** of the pure-FMA
peak, which is in the right zone for an int8-dominated
quantized matmul.

The NEON dispatch ALSO pays per-call for `quantize_a_per_row`
(fp32 → i8 + scale, ~m·k operations); the matmul-only kernel
is faster than the 3.5× speedup above suggests. For batched
inference where activations stay i8 between layers, the
quantize cost amortizes.

## F32 matmul — scalar vs Apple Accelerate (AMX)

`cargo bench --features accelerate --bench matmul_q8_0
matmul_f32`. The `accelerate` feature routes f32 matmul
through `cblas_sgemm` which uses AMX (the M-series matrix
accelerator) on Apple Silicon. Direct apples-to-apples — no
quantization confounding the measurement.

| Shape (m × n × k) | Scalar | Accelerate (AMX) | Speedup | AMX GFLOPS |
|---|---:|---:|---:|---:|
| 1 × 256 × 256       |   27.2 µs |   **1.4 µs**  |   **19.8×** |   95 |
| 1 × 1024 × 1024     |  632   µs |  **13.5 µs**  |   **47.0×** |  156 |
| 1 × 2048 × 2048     |    2.73 ms |  **67.6 µs**  |   **40.4×** |  248 |
| 1 × 4096 × 4096     |   11.45 ms | **710   µs**  |   **16.1×** |   47 |
| 8 × 1024 × 1024     |    5.07 ms |  **83.8 µs**  |   **60.5×** |  200 |
| 32 × 1024 × 1024    |   20.30 ms |  **90.4 µs**  |  **224.6×** |  **742** |

AMX scales dramatically with batch dimension `m`:

- **m = 1** (single-token decode): peaks at ~250 GFLOPS at
  n=k=2048. The 4096² drop to 47 GFLOPS suggests the matrix
  exceeds AMX's working set — needs tiling.
- **m = 8** (small batch): 200 GFLOPS sustained.
- **m = 32** (medium batch / prefill): **742 GFLOPS** — the
  AMX-favoured regime where weight rows reuse across multiple
  activation rows. Approaches the per-core ceiling.

The 224× speedup at m=32 is the headline. Compute-bound +
weight-reuse-dominated workloads (think KV-cache prefill,
batched inference) get the full AMX win.

The single-token-gen path (m=1) caps at ~250 GFLOPS at
medium k — bandwidth-bound, not compute-bound. AMX can't
reuse weights when there's only one activation row; the
fetch from L1 dominates. Still ~50× faster than scalar
which is a real win.

## Q8_0 matmul — NEON direct vs AMX-tiled (wave 24)

`matmul_q8_0_amx_tiled` dequantizes a chunk of weight rows
into a small scratch buffer (128 rows × k floats) using the
new wave-24 NEON-tuned `dequantize_q8_0_neon`, then routes
the chunk through `cblas_sgemm`. Each chunk fits in L2 so
the f32 scratch never hits DRAM — only the Q8_0 weight
(4× smaller) crosses the DRAM boundary.

Side-by-side at the same shapes (m × n × k):

| Shape | NEON Q8_0 (dispatch) | AMX-Tiled Q8_0 | Winner |
|---|---:|---:|---|
| 1 × 256 × 256 | **3.6 µs** | 10.4 µs | NEON 2.9× |
| 1 × 1024 × 1024 | **54 µs** | 118 µs | NEON 2.2× |
| 1 × 2048 × 2048 | **218 µs** | 437 µs | NEON 2.0× |
| 1 × 4096 × 4096 | **868 µs** | 1.77 ms | NEON 2.0× |
| 8 × 1024 × 1024 | 432 µs | **202 µs** | **AMX-Tiled 2.1×** |
| 32 × 1024 × 1024 | ~1.4 ms (extrapolated) | **212 µs** | **AMX-Tiled ~6.6×** |

**The crossover is m ≈ 8.** Below that, the dequant cost
exceeds the AMX speedup — single-token decode (m=1) stays
on the NEON direct path. At m=8, AMX's weight reuse
amortizes the dequant. At m=32 (typical prefill /
batched), AMX-tiled wins decisively.

**Production dispatch decision:** route Q8_0 matmul to
NEON when m ≤ 4; to AMX-tiled when m ≥ 8; either is fine
in between. A consumer-level wrapper can make this
automatic; the substrate ships both kernels + lets the
caller pick.

## Multi-thread fan-out via rayon (wave 25)

Two new variants behind the `rayon` cargo feature:

  - `matmul_q8_0_threaded` — wraps the NEON Q8_0 kernel with
    rayon. Parallelizes over `n` for `m == 1`; over `m`
    otherwise.
  - `matmul_q8_0_amx_tiled_threaded` — wraps the AMX-tiled
    Q8_0 kernel; parallelizes over weight-row chunks.

**Full dispatch matrix** (Apple Silicon dev box, release build,
`--features "rayon accelerate"`):

| Shape (m × n × k) | NEON Single | NEON Threaded | AMX-Tiled Single | AMX-Tiled Threaded | **Best** |
|---|---:|---:|---:|---:|---|
| 1 × 256 × 256 | **3.6 µs** | 290 µs | 10.4 µs | 69 µs | NEON single |
| 1 × 1024 × 1024 | **54 µs** | 251 µs | 118 µs | 370 µs | NEON single |
| 1 × 2048 × 2048 | **218 µs** | 269 µs | 437 µs | 514 µs | NEON single |
| 1 × 4096 × 4096 | 868 µs | **308 µs** | 1.77 ms | 927 µs | **NEON threaded 2.8×** |
| 8 × 1024 × 1024 | 432 µs | 456 µs | **202 µs** | 337 µs | AMX-tiled single 2.1× |
| 32 × 1024 × 1024 | ~1.4 ms | 374 µs | **212 µs** | 311 µs | **AMX-tiled single 6.6× over NEON** |

**Two empirical results from the bench:**

1. **Rayon hurts small kernels.** Task-spawn overhead is
   ~200-300 µs on this M-series box. Below ~1 MFLOP of work,
   single-thread wins. Above 16 MFLOPs, rayon amortizes and
   delivers a real 2.8× speedup (NEON case).

2. **Rayon hurts AMX-tiled at every size measured.** Accelerate
   is ALREADY internally multi-threaded for sgemm calls of the
   chunk sizes we use. Wrapping with rayon oversubscribes the
   cores → contention. The AMX-tiled single-thread path stays
   the winner from m=8 onward.

**Production dispatch rules (revised, wave 25):**

| Workload | Choose | Why |
|---|---|---|
| Single-token decode, small/medium k (n,k ≤ 2048) | `matmul_q8_0` (NEON single) | Rayon overhead > parallelism win |
| Single-token decode, large k (n,k = 4096+) | `matmul_q8_0_threaded` (NEON threaded) | 2.8× speedup; kernel large enough to amortize |
| Batched / prefill (m ≥ 8) | `matmul_q8_0_amx_tiled` (AMX single) | Accelerate's internal threading saturates cores; rayon hurts |
| Multiple matmuls fanned out concurrently | Wrap multiple `matmul_q8_0` calls in `par_iter` instead | Coarser-grained parallelism > kernel-internal threading for many small matmuls |

This dispatch is a consumer-level wrapper. The substrate ships
all four kernels + the bench harness; the caller picks based
on shape + workload.

## Q4_K matmul — scalar fused (wave 19) vs AMX-tiled (wave 26)

Q4_K is the dominant format in distributed Llama-family models
(`Q4_K_M`). Wave 19 ships a scalar fused kernel (no f32
intermediate); wave 26 adds the AMX-tiled mirror of the
wave-24 Q8_0 approach.

`cargo bench --features accelerate --bench matmul_q4_k`:

| Shape (m × n × k) | Scalar Fused (W19) | AMX-Tiled (W26) | Speedup |
|---|---:|---:|---:|
| 1 × 256 × 256 | 40.8 µs | **33.0 µs** | 1.24× |
| 1 × 1024 × 1024 | 729 µs | **395 µs** | 1.84× |
| 1 × 2048 × 2048 | 2.94 ms | **1.46 ms** | 2.01× |
| 1 × 4096 × 4096 | 11.62 ms | **5.39 ms** | 2.16× |
| 8 × 1024 × 1024 | 5.78 ms | **473 µs** | **12.2×** |
| 32 × 1024 × 1024 | 22.61 ms | **524 µs** | **43.1×** |

**The Q4_K story differs from Q8_0.** For Q8_0, NEON direct
wins at m=1 (the dequant cost in the tiled path exceeds the
AMX speedup). For Q4_K there's no NEON kernel — only the
scalar fused — so the baseline is much slower (729 µs vs
54 µs at n=k=1024). Even with the dequant overhead, AMX-tiled
wins at every size measured.

### Wave 27 added: NEON Q4_K matmul

`matmul_q4_k_neon` — same fused walk as the scalar version
but the 32-element sub-block inner loop runs through NEON:
extract nibbles (`vandq_u8` / `vshrq_n_u8`), widen u8 → u16
→ u32 → f32, accumulate via `vfmaq_f32`. Uses the algebraic
split `sum(a * (d_sub * q - m_sub)) = d_sub * sum(a*q) -
m_sub * sum(a)` so each sub-block only needs two NEON
horizontal-sums (`vaddvq_f32`), not per-element scaling.

Full Q4_K dispatch table:

| Shape (m × n × k) | Scalar Fused (W19) | **NEON (W27)** | AMX-Tiled (W26) | Winner |
|---|---:|---:|---:|---|
| 1 × 256 × 256 | 41 µs | **7.3 µs** | 33 µs | **NEON 5.6×** |
| 1 × 1024 × 1024 | 729 µs | **115 µs** | 395 µs | **NEON 6.3×** |
| 1 × 2048 × 2048 | 2.94 ms | **458 µs** | 1.46 ms | **NEON 6.4×** |
| 1 × 4096 × 4096 | 11.62 ms | **1.84 ms** | 5.39 ms | **NEON 6.3×** |
| 8 × 1024 × 1024 | 5.78 ms | 911 µs | **473 µs** | AMX 1.9× |
| 32 × 1024 × 1024 | 22.6 ms | 3.66 ms | **524 µs** | **AMX 7×** |

**Same dispatch shape as Q8_0:**

  - m = 1 (decode): NEON wins
  - m ≥ 8 (prefill/batched): AMX-tiled wins
  - m = 4-8: either; bench-driven

Wave 27 closes the Q4_K m=1 gap that wave 26 left open. Q4_K
now has the full kernel family: scalar fused (W19) + NEON
(W27) + AMX-tiled (W26), parallel to Q8_0's scalar (W12) +
NEON (existing) + AMX-tiled (W24).

## Q5_K + Q6_K matmul — scalar fused vs AMX-tiled (wave 28)

Mirror of wave-26 Q4_K for the two larger K-quants. Same shape:
scalar fused + AMX-tiled. No NEON variant yet (mirror of the
gap wave 27 closed for Q4_K).

**Q5_K** (`cargo bench --features accelerate --bench matmul_q56_k`):

| Shape (m × n × k) | Scalar Fused | AMX-Tiled | Speedup |
|---|---:|---:|---:|
| 1 × 256 × 256 | **35.8 µs** | 43.9 µs | -1.2× (scalar wins) |
| 1 × 1024 × 1024 | 738 µs | **569 µs** | 1.30× |
| 1 × 2048 × 2048 | 2.95 ms | **2.16 ms** | 1.37× |
| 1 × 4096 × 4096 | 12.10 ms | **7.64 ms** | 1.58× |
| 8 × 1024 × 1024 | 5.83 ms | **660 µs** | **8.8×** |
| 32 × 1024 × 1024 | 23.15 ms | **669 µs** | **34.6×** |

**Q6_K**:

| Shape (m × n × k) | Scalar Fused | AMX-Tiled | Speedup |
|---|---:|---:|---:|
| 1 × 256 × 256 | 45.7 µs | **39.1 µs** | 1.17× |
| 1 × 1024 × 1024 | 754 µs | **462 µs** | 1.63× |
| 1 × 2048 × 2048 | 3.03 ms | **1.85 ms** | 1.64× |
| 1 × 4096 × 4096 | 12.25 ms | **6.70 ms** | 1.83× |
| 8 × 1024 × 1024 | 6.08 ms | **582 µs** | **10.4×** |
| 32 × 1024 × 1024 | 24.33 ms | **595 µs** | **40.9×** |

**Same shape as Q4_K dispatch:**
  - m = 1: marginal AMX wins (1.2-1.8×) for sizes ≥ 1024; for
    tiny sizes (n,k=256) the scalar fused path can match or
    beat the AMX overhead.
  - m ≥ 8: AMX-tiled wins decisively (8.8× and 10.4× at m=8;
    34.6× and 40.9× at m=32).

### Wave 30 added: NEON Q5_K + Q6_K matmul

Mirror of wave-27 (NEON Q4_K). Same algebraic split per
sub-block; Q5_K extracts 1 high bit from `qh` per element,
Q6_K extracts 2 high bits + applies a signed -32 offset.

**Full Q5_K dispatch table:**

| Shape (m × n × k) | Scalar | **NEON (W30)** | AMX-Tiled | Winner |
|---|---:|---:|---:|---|
| 1 × 256 × 256 | 35.8 µs | **10.6 µs** | 43.9 µs | NEON 3.4× |
| 1 × 1024 × 1024 | 738 µs | **168 µs** | 569 µs | NEON 4.4× |
| 1 × 2048 × 2048 | 2.95 ms | **671 µs** | 2.16 ms | NEON 4.4× |
| 1 × 4096 × 4096 | 12.10 ms | **2.71 ms** | 7.64 ms | NEON 4.5× |
| 8 × 1024 × 1024 | 5.83 ms | 1.35 ms | **660 µs** | AMX 2.0× over NEON |
| 32 × 1024 × 1024 | 23.15 ms | 5.35 ms | **669 µs** | AMX 8× over NEON |

**Full Q6_K dispatch table:**

| Shape (m × n × k) | Scalar | **NEON (W30)** | AMX-Tiled | Winner |
|---|---:|---:|---:|---|
| 1 × 256 × 256 | 45.7 µs | **7.9 µs** | 39.1 µs | NEON 5.8× |
| 1 × 1024 × 1024 | 754 µs | **123 µs** | 462 µs | NEON 6.1× |
| 1 × 2048 × 2048 | 3.03 ms | **492 µs** | 1.85 ms | NEON 6.2× |
| 1 × 4096 × 4096 | 12.25 ms | **1.97 ms** | 6.70 ms | NEON 6.2× |
| 8 × 1024 × 1024 | 6.08 ms | 988 µs | **582 µs** | AMX 1.7× over NEON |
| 32 × 1024 × 1024 | 24.33 ms | 3.95 ms | **595 µs** | AMX 6.6× over NEON |

Q6_K NEON edges Q5_K NEON at m=1 (~6× vs ~4.4× over scalar) —
no per-element min subtraction, just the signed -32 offset
that NEON's `vsubq_s8` handles in one op.

**Same dispatch pattern across Q4_K / Q5_K / Q6_K:**
  - m = 1 (decode): NEON wins
  - m ≥ 8 (prefill / batched): AMX-tiled wins

All three K-quant families now have the full three-kernel
suite (scalar / NEON / AMX-tiled) parallel to Q8_0.

## Attention — unchunked (wave 21) vs flash-chunked (wave 29)

Flash attention's online-softmax walks `seq_kv` in chunks of 256,
maintaining a running `(max, sum, output[d_head])` instead of
the full `[seq_kv]` scores buffer the unchunked kernel keeps.

`cargo bench --bench attention`, d_head=128:

| Shape (seq_q × seq_kv) | Unchunked | Flash | Verdict |
|---|---:|---:|---|
| 1 × 256 | 13.8 µs | 14.8 µs | tied (flash +7%) |
| 1 × 1024 | 55 µs | 60 µs | tied |
| 1 × 4096 | 218 µs | 232 µs | tied |
| 1 × 16384 | 883 µs | 927 µs | tied (flash +5%) |
| **1 × 65536** | 3.78 ms | **3.65 ms** | **flash wins 3%** |
| 32 × 4096 | 7.30 ms | 7.33 ms | tied |

**Compute is within 5% across all sizes.** The headline win
isn't speed — it's structural memory:

| Variant | Scratch per query | Scratch at seq_kv=1M |
|---|---|---|
| Unchunked | O(seq_kv) | 4 MB |
| Flash | O(d_head + 256) | 1.5 KB |

For batched long-context inference (seq_q=32 at seq_kv=1M),
unchunked allocates 128 MB of scratch total; flash allocates
50 KB. The flash kernel is the right choice for any context
length beyond ~16K — same compute, dramatically lower memory
footprint, no working-set pressure on L2/SLC.

The 3% speedup at seq_kv=65536 is the inflection point: the
unchunked scores buffer (256 KB) no longer fits cleanly in
the L2 cache. Below that the buffer is L2-resident in both
variants; above it, only flash's per-chunk buffer fits.

## BF16 kernel ladder on real Nemotron-3 weights (waves 42 + 46)

Verification harness `examples/bf16_bringup.rs` run against
**nvidia/NVIDIA-Nemotron-3-Nano-4B-BF16** — a real 7.4 GiB
single-file safetensors checkpoint (263 tensors, 94 2-D BF16
matrices) pulled from HuggingFace. `m=1` (single-token
decode shape), full checkpoint load via
`dispatch::read_any_model` in 1.69 s.

| tensor | shape | scalar | NEON | NEON speedup | dequant→f32 |
|---|---|---|---|---|---|
| `backbone.embeddings.weight` | 131072×3136 | 310.4 ms | 49.2 ms | **6.3×** | 423.4 ms |
| `lm_head.weight` | 131072×3136 | 312.4 ms | 48.4 ms | **6.5×** | 333.9 ms |
| `backbone.layers.0.mixer.in_proj.weight` | 17504×3136 | 42.1 ms | 6.6 ms | **6.4×** | 56.7 ms |

Max relative divergence across all three routes on real
weight bytes: **3.2 × 10⁻⁶**. The NEON decode
(`vshll_n_u16` widening shift = the entire bf16→f32
conversion) sustains ~16 GB/s of weight traffic
single-threaded; the dequant→f32 route loses to fused even
at this size because the 2× widened slab (1.53 GiB for the
embedding matrix) blows every cache level — it only wins
once the AMX `cblas_sgemm` backend is in play with `m ≥ 2`.

## NemotronH end-to-end decode (waves 48-51)

`examples/nemotron_decode.rs` against the real
**nvidia/NVIDIA-Nemotron-3-Nano-4B-BF16** (42 hybrid layers:
Mamba2 / GQA attention / relu² MLP; 131 072-token lm_head),
prompt `"The capital of France is"` → `" Paris."`,
greedy decode, weights BF16 in the mmap end to end:

| build | per-token wall clock | throughput |
|---|---|---|
| single-thread NEON (wave 49) | ~435 ms | 2.3 tok/s |
| `--features rayon` (wave 51) | ~80 ms | **12.5 tok/s** |

The rayon build fans each projection's output rows across
cores (wave 25 dispatch shape) with each thread running the
wave 42 NEON inner loop; outputs are bit-identical to the
single-thread kernel (chunking never splits a dot product).
5.5× on the dev box — short of core count because the
135 MiB/token of weight traffic begins to saturate memory
bandwidth, which is the right ceiling for a decode loop.

## Expert-batch dispatch (wave 63)

Follow-up to the wave 62 floor: the floor stops the waste
but leaves MoE expert matmuls single-core. `matmul_bf16_batch`
runs a layer's whole routed-expert stage (8 experts × gate+up,
then 8 × down) as ONE rayon scope — one pool wakeup amortized
over ~66 MB of work instead of 24 wakeups. Per-job math is the
same single-thread NEON kernel and accumulation stays in pick
order, so results remain bit-identical (parity suites pin it).
Wired into Mellum2 and Cohere2-MoE decode.

Controlled A/B (back-to-back, same heavy external load):

| config | Mellum2 decode (median) |
|---|---|
| RAYON_NUM_THREADS=1 | ~1.15 s/token |
| expert-batch + floor 128 | **~0.78 s/token** (1.45×) |

Best observed at quieter moments: ~0.5 s/token. Absolute
numbers float with box load; the A/B ratio is the honest
signal.

Wave 64 extends the batch shape to Granite h-tiny's fused
expert walk (expert matrices stay zero-copy slices of the
fused [experts, ·, ·] tensors). Same-load A/B: ~174 ms/token
single-thread → ~125 ms/token batched (1.4×) — h-tiny is
now the fastest MoE decode in the set, beating its own
wave-53 ~270 ms baseline.

## Prefill win, measured across the Mac fabric: 47–54× energy, 10–14× speed (wave 91)

Three quiet/loaded Metal-4 boxes, the matmul2d prefill path vs
threaded-NEON, on the canonical 256/512×4096×4096 projections.

**Energy (`mpp_energy` example, IOReport, idle-subtracted, QUIET
boxes) — the number that matters for an energy-metered substrate:**

```
            idle    matmul2d µJ/GFLOP   NEON µJ/GFLOP   CPU/GPU ratio
  M5 Max   20.6 W        6 630            358 906          54×
  M4 Max    4.45 W       5 052            238 990          47×
```

So a prefill projection costs **~47× fewer joules per GFLOP** on
the Metal-4 tensor path than on the CPU — an order-of-magnitude
energy collapse for the compute-bound part of inference.

**Honesty caveat (added after a second run):** trust the M4
number, not the "54×" first reported for M5. The M4 idle was
4.45 W (genuinely quiet) and its matmul2d µJ/GFLOP repeats
(~4.6–5.1k across runs → the calibration wrote 4615 Measured).
The M5 idle was 20.6 W — NOT truly idle (something was running),
so idle-subtraction is noisy there: a second run gave 12 385
µJ/GFLOP vs the first 6 630, a 2× spread. The RATIO (~47–54×) is
roughly stable because GPU and CPU are measured under the same
baseline, but the M5 ABSOLUTE is unreliable. Per-box Measured
calibrations written: M4 = 4615, M5 = 12385 µJ/GFLOP (the placer
uses each box's own; the seeded Estimate is 6000). Lesson: a box
is only "quiet" if its idle draw says so — check it, don't assume.

**Throughput (512×4096×4096, mpp_bt vs threaded-NEON):**

```
  M4 Max (quiet)   3240 GFLOP/s   14.0×
  M5 Max (quiet)   4073 GFLOP/s   12.3×
  M3 Ultra (load)  3022 GFLOP/s   10.2×
```

Portable across M3/M4/M5; the speedup is largest on M4 (slower
CPU baseline), the absolute GFLOP/s highest on M5.

**Token-drift sweep — 9/9 IDENTICAL.** NemotronH-4B, Granite-4.0-
h-micro, Falcon-H1-0.5B × 3 diverse long prompts each, OFF vs ON:
every output token-for-token identical. The ~5% bf16-activation
rounding flips no argmax on anything tested.

So the prefill path is correct (full-output shape sweep), fast
(10–14×), ~50× more energy-efficient, and output-safe across 3
families × 3 prompts. Still OPT-IN (`JOULECLAW_MPP_PREFILL=1`):
it is a real precision change and the substrate default stays the
exact path, but the case for flipping it on is now strong.

NOTE: the ledger still prices matmul2d ops with the OLD
MPSMatrix-based metal-gpu calibration; the measured 5–7 mJ/GFLOP
above is its true cost and should overlay the placement table —
a calibration-integration follow-up.

## matmul2d prefill DEFAULT-ON + 5/5 family-safe + q/k/v windows (wave 93)

Resolving the four open items from the prefill arc.

**Default-on decision: ON when the fabric is on** (opt OUT with
`JOULECLAW_MPP_PREFILL=0`). Rationale: turning the fabric on IS
the documented "caller is choosing joules over bit-exactness"
opt-in (the AMX path already rides that contract); a ~47× energy
win that is token-identical across all five families fits it
exactly, and fabric-OFF (the real default) stays byte-exact. The
opt-out keeps the more-precise resident-F16 GPU path available.

**Family token-safety: 5/5 (13/13 prompt pairs IDENTICAL).**
NemotronH-4B, Granite-h-micro, Falcon-H1 (×3 prompts each) +
North-Mini-Code-30B-A3B and Mellum2-12B-A2.5B (×2 each), default
vs opt-out: every output byte-for-byte identical. The bf16-
activation rounding flips no argmax across topologies (dense +
hybrid-Mamba + MoE) and prompt styles (Q&A, prose, code).

**q/k/v windows wired (`matmul_bf16_auto_multi`).** The fused
prefill windows (q/k/v, expert stages) now route each job
through matmul2d too, not just the single-weight `proj_m` path —
so the whole prefill projection set gets the win. All-or-nothing
per window; a fallback recomputes via the resident path.

**M4 quiet standard-member calibration** (idle 5.65 W, clean) for
context — prefill-batch (m=16) µJ/GFLOP: cpu-neon 335 836, AMX
29 523, metal-gpu (resident MPSMatrix F16) 13 253. So matmul2d at
~4 615 is **~2.9× lower energy than even the resident-F16 MPS
path**, ~6× under AMX, ~73× under NEON. Written to the M4 box's
calibration; the M3-Ultra dev calibration stays loaded-box (can't
quiet the dev box).

Verified default-on on Nemotron prefill: metal-mpp carries 99% of
projection GFLOPs by default; `=0` falls back to the resident
path; output identical either way. Backend 19 + loader 285 green
(fabric default-on AND opt-out AND fabric-off).

## matmul2d prefill WIRED IN — token-identical, opt-in (wave 90)

The fast non-square GEMM is now actually used by the model
forwards, and it's the first real prefill acceleration the M5
arc has delivered.

`matmul_bf16_auto` (the central projection dispatch every family
calls via `proj_m`) now routes bias-free, M≥32 GEMMs through
`matmul_bf16_mpp_bt` when **`JOULECLAW_MPP_PREFILL=1`**. Off by
default — matmul2d rounds activations to bf16 (~5% per-projection
rel error), so the operator opts into faster-but-lower-precision
prefill. M≥32 keeps the pad-to-64 waste under the GPU win; decode
(M=1) is untouched.

**Real-model A/B, NVIDIA-Nemotron-3-Nano-4B, M3 Ultra:** a long
coherent few-shot prompt ("…capital of Germany? Answer:") gives
the SAME output OFF and ON — `" Berlin. Question: What is the
capital of Canada? Answer"` — with matmul2d carrying **100% of
the prefill projection GFLOPs** (242.8 GFLOP). The ~5% activation
rounding did not flip a single argmax (also identical on a second
prompt). So prefill runs on the ~4–11× Metal-4 tensor path with
no observed token drift.

M5 Max throughput (vs threaded-NEON, wave 89): 4–11× on real
projection shapes. Decode still uses the existing bandwidth-bound
paths (matmul2d needs M%64 / pad, pointless for M=1).

**Cross-family token check (M3 Ultra, same long few-shot
prompt), all OFF==ON:**
- NemotronH 4B → " Berlin. … Canada? Answer" (gpu 100% of GFLOP)
- Granite-4.0-h-micro → " Berlin. … Australia" (gpu 100%)
- Falcon-H1-0.5B → " Berlin. … Italy" (gpu 57.8%)

So across three topologies the bf16-activation prefill path is
output-identical to the f32-activation baseline. Kept OPT-IN
anyway (one prompt each; a borderline argmax could still flip on
some input, and the substrate's default is the higher-precision
path) — the operator turns it on for the 4–11× prefill speedup.

Two real, integrated M5/Metal-4 wins now: wave 82 (2× decode, f16
MPS, on by default) + wave 90 (4–11× prefill, Metal-4 tensor op,
opt-in).

## NON-SQUARE matmul2d CRACKED — 3.6–10× faster prefill GEMM (wave 89)

The one-more-attempt paid off: the matmul2d kernel now does a
correct, fast NON-SQUARE GEMM in ONE dispatch — no tiling.

The missing piece was in Apple's example HOST code
(`matmul.swift`, fetched from the repo), which computes the
dispatch grid *conditionally on M>N vs M≤N*. That clue, plus a
single-element probe, revealed the fix: **declare the output
tensor TRANSPOSED — `dextents(N, M)` instead of `(M, N)`.** The
op then writes `O[m,n]` to `out[m*N+n]` (the natural row-major
read) for arbitrary shapes. The non-transposed C is correct
*only* for square M==N — which is exactly why waves 84–88 saw
"square-only." Constraints are now just M%64, N%32, K%32.

Verified `WRONG=0` over the FULL output across a non-square
shape sweep (128×64, 128×256, 256×128, 192×96, 64×256,
320×128, …) on the M3 Ultra dev box.

`matmul_bf16_mpp_bt` rewritten to a SINGLE dispatch (transpose
weight `[N,K]→[K,N]` cached, pad M to %64, one matmul2d). The
per-block overhead that made wave 88 slow is gone. Measured vs
threaded-NEON on M3 Ultra, realistic projection shapes:

```
  M×N×K             mpp_bt GFLOP/s   NEON   speedup
  128×4096×4096          954         265     3.6×
  256×4096×4096         1672         294     5.7×
  256×11008×4096        2121         285     7.4×
  512×4096×4096         3022         297    10.2×
```

So prefill projections are **3.6–10× faster** on the Metal-4
tensor path than the CPU path, and the speedup grows with M
(more tokens) and N. Caveat: matmul2d rounds activations to
bf16 → ≈5% rel error vs the f32-activation NEON path — a
precision-for-speed trade, acceptable for prefill (same class
as the existing f16-resident GPU path).

NOT yet wired into the family forwards — doing so changes
prefill numerics (the bf16-activation rounding), so it needs a
real-model output check (token drift) before flipping it on.
That is the next step; the primitive + wrapper are correct,
fast, and tested.

## Square-tiling makes matmul2d correct for non-square — but 3× SLOWER than CPU (wave 88)

The square-tiling wrapper works and is fully validated, and it
still doesn't help. Honest end of this thread (for now).

`matmul_bf16_mpp_bt` now tiles a non-square `out[M,N]` into
S×S=128² SQUARE blocks, each a square matmul2d call over a
zero-padded `[S,K]` activation tile and a `[K,S]` weight-column
tile. **Correctness: validated against the NEON oracle over the
FULL output across a SHAPE SWEEP** — square, non-square, and
non-tile-aligned (130×300, 192×448, 300×130) all pass. So
non-square GEMM via this kernel is finally correct.

**But it is ~3× SLOWER than the existing threaded-NEON CPU
path** (M3 Ultra, realistic projection shapes):

```
  M×N×K            mpp_bt GFLOP/s   threaded-NEON   speedup
  128×4096×4096        94             275           0.34×
  256×4096×4096        95             288           0.33×
  256×11008×4096       76             279           0.27×
  512×4096×4096        95             292           0.33×
```

The kernel hits ~24 TFLOP/s at large SQUARE shapes, but tiling
into 128² blocks destroys it: per block there is a CPU-side
strided weight-column extraction and a SEPARATE command buffer,
so a single projection becomes dozens of small dispatches +
copies. It's overhead-bound (~95 GFLOP/s regardless of shape),
not compute-bound. On top of that, matmul2d rounds activations
to bf16 → ~5% relative error vs the f32-activation NEON path.

**Decision: NOT wired into the families.** The wrapper stays as
a correct, tested primitive, but it currently loses to the CPU
path, so using it would make prefill slower and less precise.
To ever win it needs all three: (1) GPU-RESIDENT bf16 weights
(no per-call re-upload / per-block extraction), (2) all blocks
in ONE command buffer, (3) larger square tiles (256/512) — i.e.
essentially the correct NON-SQUARE kernel (one dispatch, no
tiling), which still needs Apple's host grid that I haven't
derived. Net deliverable to the substrate from the whole M5
matmul2d arc: **a fast correct SQUARE GEMM primitive, and zero
prefill acceleration.**

## matmul2d is SQUARE-ONLY — not usable for prefill yet (wave 87)

A second, harder correction (the wave-86 fix was still wrong).

A full-output **shape sweep** showed the kernel is correct ONLY
for **square M==N** (verified `WRONG=0` at 128², 192², 256²,
384², …, 2048² — including non-powers-of-2) and wrong for EVERY
non-square shape (128×256, 256×128, 192×128, 128×160 … all
~95% wrong). Wave 86's "correct for M≥128 & N≥128" was itself
based on shapes that happened to be square; it was still an
overclaim. The benchmarked ~2–5× throughput numbers stand —
but only for square GEMMs.

**Consequence: this cannot accelerate prefill.** Real
projection GEMMs are non-square (M = prompt tokens ≠ N =
hidden), so `matmul_bf16_mpp_bt` now requires M==N and **falls
back for all real prefill** — it provides zero acceleration to
the substrate today. The cause is MY grid/slice mapping, not
the MPP op: the op supports non-square (Apple's example does
128×64), but the correct host grid for the non-square case is
not in the example's README and I have not derived it. Both
`matmul_bf16_mpp` and the wrapper are constrained to square
M==N (M%64≥128, K%32) and documented as such; the backend test
asserts square-correct over the full output AND that a
non-square call is rejected.

Net honest status of the M5/Metal-4 matmul2d arc: a **correct,
fast SQUARE bf16 GEMM** (~4× simdgroup, portable M3+M5,
full-output verified) that is **not yet a general GEMM and not
yet useful to the loader**. Next requires the non-square host
mapping (Apple sample code with the dispatch grid, or a square-
tiling wrapper that blocks a non-square problem into square
sub-GEMMs — feasible but adds per-block dispatch + strided
weight slicing).

## matmul2d: full-output validation + prefill wrapper (wave 86)

A methodology correction and the loader bridge.

**Validation was corner-only — fixed.** Waves 84–85 verified
the matmul2d output by checking only the top-left 8×8 corner.
That was insufficient: a prefill-wrapper test at M=64 then
showed **96% of elements wrong**. Re-checking the FULL output
(every element) across the benchmarked shapes shows the wave-84
conclusion still HOLDS — `WRONG=0 / max_abs=0.0` at 256²×512,
512²×2048, 1024²×4096, and **2048²×4096 (all 4.2 M elements)** —
but the kernel is **correct only for M≥128 and N≥128** (with
M%64, N%32, K%32). Smaller M or N mis-tile and give wrong
results; the corner check happened to sit in the one region
that stays right. Bench, backend test, and wrapper test now all
assert over the full output. Lesson logged: GEMM correctness
must be checked across all tiles, never a corner.

**Prefill wrapper** `fabric_dispatch::matmul_bf16_mpp_bt`
(`out[M,N] = a[M,K]·w[N,K]ᵀ`): transposes the loader's `[N,K]`
weight to the `[K,N]` the kernel wants (once, cached by ptr),
pads M up to a multiple of 64, dispatches `matmul_bf16_mpp`,
copies the live rows back, and records the GFLOPs in the
attribution ledger. Returns false → caller falls back when the
fabric is off, the GPU member is uncertified, the shape is
outside M≥128/N≥128, or the device is pre-Metal-4. Tested
against the NEON oracle over the full output (looser tolerance:
matmul2d rounds activations to bf16, a deliberate
precision-for-speed trade). Not yet called by the families —
wiring it into a real prefill path + measuring is next.

## M5 Neural-Accelerator matmul: CRACKED — correct + ~5× (wave 84)

The matmul2d GEMM now works, bit-exact, on the M5 GPU Neural
Accelerator (`mpp_gemm_bench.rs`). The unblock came from Apple's
own example (github liuliu/example_matmul_metal4, found by web
search) and was a single wrong token: the descriptor's K must be
**`dynamic_length_v<int>`**, NOT `dynamic_extent`/`0`. The wrong
value silently caps the K-sum at one ~32-wide tile (the bug that
defeated every header-derived attempt in waves 83–83d). The rest
of the recipe is plain: standard `C[M,N]=A[M,K]·B[K,N]` (no
transpose, `tensor_inline` over the raw buffers), `slice()` for
the tile origin, ONE `run()` that loops K internally,
`execution_simdgroups<4>`, descriptor tile `(64, 32)`.

The layout was pinned with single-element probes (set one input
element to 1, see where the output lights up): A is effectively
`[M,K]`, B `[K,N]`, C `[M,N]`. Constraints: M%64, N%32 (N must be
≥128 — N=64 mis-tiles), K%32; M=1 decode is NOT a valid shape, so
this is a **prefill / batched** accelerator, not a decode-GEMV one
(consistent with the wave-67 finding that decode is bandwidth-
bound).

**Measured on M5 Max, every shape `max_abs = 0.0` (bit-exact vs
the f32 CPU reference over bf16 inputs):**

```
              matmul2d (Neural Accel)   simdgroup_matrix (Metal 3)
  512²×2048     ~2.3 TFLOP/s
  1024²×4096    ~9.3 TFLOP/s
  2048²×4096    ~16–25 TFLOP/s            ~4.8 TFLOP/s (peak)
```

So **beyond-2× is real: ~2–5× the simdgroup path depending on
shape, and bit-correct** (throughput varies with box load). This
SUPERSEDES the withdrawn wave-83 figures, which came from
partial-K (incorrect) kernels — these numbers are from a kernel
verified `max_abs = 0.0`.

**Correction to the "Neural-Accelerator" framing: this is a
Metal-4 win, not an M5 exclusive.** The same kernel runs
bit-correct on the **M3 Ultra** dev box at ~13–20 TFLOP/s — the
same ~4× over its own simdgroup ~4.8. So the headline win is the
**MPP `matmul2d` tensor-op path itself, portable across M3 Ultra
and M5 Max**, not specifically the M5 GPU Neural Accelerators.
Whether the M5 NA adds anything *beyond* M3's matrix units is NOT
shown by these loaded-box numbers (M3 ~13–20, M5 ~16–25 TFLOP/s
overlap within run-to-run noise) and would need a quiet-box A/B
to isolate. The portability is good news: the win applies to the
whole Apple-silicon Mac fabric with Metal 4, including the dev
box where tests run — not stranded on the one M5 host.

Not yet wired into the loader. Next: a bf16-native resident
prefill path that dispatches the projection GEMMs through this
kernel (and a real µJ/GFLOP calibration of it), then attribution.

## M5 Neural-Accelerator matmul: reachable + fast, not yet correct (wave 83)

Chasing "beyond 2×" on the M5 (wave 82 got 2× for free on the
existing f16 MPS path). Findings, in `backend-apple/examples/`:

- **bf16 `simdgroup_matrix` GEMM** (`bf16_gemm_bench.rs`):
  bit-correct (max_rel 0.0, f32 accumulation), but **ties M3**
  — ~4.78 TFLOP/s on M5 vs ~4.80 on M3 Ultra. So
  `simdgroup_matrix` is the Metal-3 matrix-unit path and does
  NOT engage M5's GPU Neural Accelerators. A real, correct,
  bf16-native kernel — just not the accelerator.

- **The accelerator IS the Metal-4 Metal-Performance-
  Primitives path** `mpp::tensor_ops::matmul2d`
  (`mpp_gemm_bench.rs`). Confirmed reachable from our pure-Rust
  + runtime-MSL stack; the full API was mapped by reading the
  on-device framework headers
  (`matmul2d_descriptor(M,N,K,transL,transR,relaxed,mode)` with
  M or N a multiple of 16, `execution_simdgroups<N>`, tensor
  `slice<>()`, `run`, cooperative tensors with explicit
  zero-init).

- **Contraction DECODED + verified correct at a square tile**
  (`mpp_diag.rs`, one-hot + full-K probes): the op computes
  `C[a,b] = Σ_c L[a,c]·R[c,b]` (L's dim1 contracted against R's
  dim0). For `C[m,n]=Σ_k A[m,k]·B[n,k]`: L = A`[M,K]`, R = Bᵀ
  via `transpose_right` on the `[N,K]` weight, C `[M,N]`. At
  32×32×32 the output is exact (within bf16 rounding).

- **Non-square (K≠M): layout SOLVED, full-K reduction NOT.**
  Two corrections to earlier drafts, both from reading the
  actual on-device headers:
  - `static_slice` is **fiction** — it appears only in Apple's
    MPP example *comment*, never in `metal_tensor`. The real
    primitive is templated **`slice<Extents...>(indices...)`**:
    the template params ARE the static tile extents. So the
    earlier "needs host `MTLTensor`/`tensor_handle`" framing
    (waves 83b/83c) was WRONG; and there is no `setTensor` in
    the Metal API at all, so the "binding-layer blocker" was a
    misdiagnosis.
  - With templated `slice<TM,K>` on full in-kernel tensors and
    the correct layout (A`[M,K]` + `transpose_right` on
    B`[N,K]` → C`[M,N]`), a SINGLE `run()` produces the right
    dominant output structure at non-square (a one-hot A row
    lands in the right C row) — but the K-sum **caps at one
    ~32-wide hardware tile** (full-K-512 test gave 32, not 512)
    and bleeds into neighbouring rows. Multi-`run()` K
    accumulation — both device-memory output and the
    cooperative-tensor idiom (zero-init → loop → store) —
    yields all-zero.
  - Net: the contraction/layout is understood and correct for
    K ≤ one tile; the multi-tile K-reduction idiom is the
    unresolved core and could NOT be derived from headers
    alone. Next step needs Apple MPP sample code / docs — not
    more blind kernel iteration, and not binding work.

**CORRECTION to an earlier draft of this entry:** the
"~7.4–12 TFLOP/s, beyond 2×" figures were measured on kernels
that did NOT sum the full K (partial work), so they OVERSTATE
throughput and are withdrawn. A valid beyond-2× number needs
the correct tiled kernel (host `MTLTensor`) first. What is
honestly established: the accelerator path is reachable and
the matmul is correct at a square tile; the throughput win is
plausible but not yet validly measured.

Honest status: reachable + semantics solved + correct at a
square tile; correct non-square tiling and a valid throughput
number are the remaining work. Not integrated into the
substrate.

## Cross-silicon: the substrate on an M5 Max (wave 82)

The pure-Rust substrate is portable across Apple-silicon
generations with zero code changes. Cloned onto an M5 Max
(40-core GPU, 128 GB, macOS 26.5.1, Metal 4), it builds in
~30 s and decodes Falcon-H1-0.5B **token-identical** to the
M3 Ultra dev box — at **~42 ms/step vs ~84 ms**, a clean 2×
on identical code. The existing MPS f16 path already rides
M5's faster GPU; no `MTLTensor` / BF16-native / Neural-
Accelerator work was needed to get the speedup.

Measured M5 calibration (IOReport, idle-subtracted, **0.67 W
idle** — a genuinely quiet box):

```
                 prefill-batch (m=16)   decode-gemv (m=1)
  local/metal-gpu   21.8 mJ/GFLOP          118 mJ/GFLOP
  local/amx        107   mJ/GFLOP          941 mJ/GFLOP
  local/cpu-neon  1031   mJ/GFLOP         1150 mJ/GFLOP
```

Falcon-0.5B decode attribution on M5 (measured): GPU's
9.5 GFLOP costs **0.2 J**, CPU's 17.6 GFLOP costs 18.1 J —
the GPU member is ~50–90× more efficient per GFLOP, which is
the whole reason the placer routes the big GEMMs there.

**Honest caveat — NOT a fair energy A/B yet.** The standing
M3 Ultra calibration was captured on a LOADED box (100 W
idle); idle-subtraction over a 100 W baseline is a different
regime than over 0.67 W. The 2× LATENCY result is clean
(same binary, both boxes), but a clean energy comparison
needs a quiet M3 Ultra re-calibration. Recorded as a number
to earn, not a number to quote.

The MTLTensor/BF16-native lift is now "can we go further on
M5," not "is M5 usable" — the floor is already 2×.

## The ledger was under-crediting the GPU (wave 81b)

The attribution counted only the four GEMM weights in the
attention block — the on-device `attn_decode` compute (score
= q·kᵀ over the cache, softmax, weighted sum over v) was real
GPU work credited to nobody. ≈ 4·n_q·hd·seq flops per layer
per token, now recorded as two GEMM-equivalents of
`(n_q, seq, hd)` so the same measured µJ/GFLOP prices it.
Magnitude is small at short context (North 48-tok: GPU
302.7 → 303.8 GFLOP, share 87.7 → 87.8 %) but it grows
linearly with sequence length, and silently dropping it
inflated the CPU's apparent share in every prior wave's
numbers. Honest accounting over flattering accounting.

## Attention block reaches Granite-H — all five families on-device (wave 81)

Granite-H is the last family in. Its attention is NoPE with
a muP `attention_multiplier` folded into a q pre-scale
(`q × attention_multiplier·√d` so the decode kernel's `1/√d`
cancels to `(q·k)·multiplier`). That is neither a rope nor a
norm — it's a fourth qk stage, `ScaleQk { q_scale, k_scale }`,
a one-line MSL kernel that scales q and k per kind with no
rotation. Granite's 4 attention layers (the rest are Mamba)
run as one command buffer each; the device cache coexists
with the MoE expert windows that already owned most of this
model's GPU GFLOPs.

Verified token-identical on BOTH micro (dense) and h-tiny
(MoE), baseline vs fabric. h-tiny attribution:

```
local/cpu-neon:     208 ops,  7.4 GFLOP (16.9%), 28.1 J projected
local/metal-gpu:  15376 ops, 36.7 GFLOP (83.1%),  5.9 J projected
```

**The qk-stage taxonomy is now closed over the set** — every
family's attention front-half is one of:

| stage           | rope            | norm | scale            | family       |
|-----------------|-----------------|------|------------------|--------------|
| `Skip`          | none (NoPE)     | no   | none             | NemotronH; Cohere2 full layers |
| `NormRotateHalf`| rotate-half     | yes  | none             | Mellum       |
| `InterleavedRope`| interleaved    | no   | none             | Cohere2 sliding |
| `RotateHalf`    | rotate-half     | no   | q/k post-scale   | Falcon-H1    |
| `ScaleQk`       | none            | no   | q/k scale only   | Granite-H    |

Five MSL kernels back the whole-block path: {silu·mul,
rmsnorm, qknorm_rope, rope_interleaved, rope_rotate_half,
scale_qk, kv_append, attn_decode}. Each rope/scale variant
has a ring-wrap CPU-reference test (backend suite now 18).

## Attention block reaches Cohere2/North — interleaved rope (wave 79)

The qk stage of the fused block is now a three-way enum:
`Skip` (NoPE), `NormRotateHalf` (Mellum), and a new
`InterleavedRope` — a second small MSL kernel implementing
the Cohere convention (pair `(x[2i], x[2i+1])` rotates by
`position · inv_freq[i]`, no per-head norm). Cohere2/North
adopts the block on ALL its layers: sliding layers ride the
ring cache with interleaved rope; full-attention layers are
NoPE and use `Skip`; `force_rope` dense-prefix layers follow
`uses_rope` per layer. RoPE is absolute and applied before
caching, so the device ring IS the sliding mask — the same
argument that justified the CPU shadow's front-trim.

Correctness pinned two ways: a ring-wrap CPU-reference test
of the interleaved path (backend suite), and
live North Mini Code 30B-A3B — token-identical to the CPU
baseline, coherent 48-token completion. Attribution:

```
local/cpu-neon:   8544 ops,  42.3 GFLOP (12.3%), 159.7 J projected
local/metal-gpu: 61452 ops, 302.7 GFLOP (87.7%),  48.9 J projected
```

Wall-clock: early steps pay the one-time f16 residency
conversion of each layer's q/k/v/o (first ~20 steps climbed
to ~1.6 s); steady state settles to ~356–488 ms vs ~393 ms
CPU-only on the same loaded box — wall parity, at a third of
the projected joules. The MoE expert windows already owned
most of this model's GPU GFLOPs; wave 79's win here is
energy, not latency.

Falcon-H1 remains the one family outside the block: its μP
multiplier suite scales q/k/v outputs per weight, which the
window has no hook for yet (cleanest path: per-job alpha in
the cached MPS kernels, keyed by (m, n, k, alpha_bits)).

## Attention block reaches NemotronH + the encode race (wave 78)

The wave-77b fused block generalizes: the qk-norm+rope stage
is now `Option`al, and NemotronH — NoPE, no QK-norm — is the
first adopter (`qk = None`: the rope kernel never encodes,
q/k flow raw into the cache). Its 4 attention layers each run
as one command buffer against a device-resident f16 KV cache;
non-rolling, so at device capacity the layer retires to CPU
and the always-maintained shadow carries on with full history.

Live (Nemotron-3-Nano-4B, loaded box): token-identical
output, ~80-85 ms/step fabric-on vs ~91 ms CPU-only —

```
local/cpu-neon:    37 ops,  7.9 GFLOP (17.6%), 29.9 J projected
local/metal-gpu:  423 ops, 37.0 GFLOP (82.4%),  6.0 J projected
```

GPU ops = 4 attn layers × 4 weights per step + prefill GEMMs;
the placer correctly keeps the sub-GFLOP Mamba/MLP GEMVs on
CPU (below the measured ~20 GFLOP solo crossover).

**Env-knob footgun, recorded the hard way**: the first run
inherited `JOULECLAW_PAR_MIN_WEIGHT_MB=128` (right for Mellum
on a loaded box, where the GPU carries nearly everything) and
clocked ~380 ms/step — the knob had single-threaded all of
NemotronH's sub-128 MB Mamba/MLP weights, which here are most
of the per-token CPU work. Same binary without the override:
~83 ms. The knob tunes the CPU remainder, and the remainder's
share differs per family.

**Parallel-encode race fixed**: backend tests SIGSEGV'd when
the harness ran them on multiple threads — bisected to
concurrent ENCODE of cached MPS kernel objects (the mutexes
guarded the cache MAPS, not the kernels; pre-existing in
77b, reproduced on the unmodified tree). One `encode_lock`
in the context now serializes encode+commit at every
dispatch entry. Free in production (single dispatcher);
5× parallel suite runs green.

## The whole attention block on-device (wave 77b)

Mellum's decode attention is now ONE command buffer per
layer: rmsnorm → q/k/v GEMMs → qk-norm+rope → KV append →
attention → o_proj, against DEVICE-RESIDENT f16 KV ring
caches (rolling on the 21 sliding layers, capacity-bounded
on the full layers). Two new kernels: kv_append and a
two-pass-softmax decode attention (one threadgroup per
q-head, GQA mapping, score vector in threadgroup memory,
ring indexing). The kernel set: {silu·mul, rmsnorm,
qknorm_rope, kv_append, attn_decode}.

Discipline that keeps every tier correct mid-stream: the
CPU shadow cache is appended from the readback every token
(k/v rows ≈ 4 KB); after a CPU prefill the device cache
BACKFILLS from the shadow (direct shared-memory writes);
any fused failure marks the layer dead for the session and
the CPU tiers continue seamlessly. Correctness pinned by a
multi-token CPU-reference test INCLUDING a ring wrap.

Live (Mellum, loaded box): identical text, steps down to
~155 ms between load spikes, attribution:

```
local/cpu-neon:   2879 ops, 24.0 GFLOP (26.4%),  90.4 J projected
local/metal-gpu: 11776 ops, 66.7 GFLOP (73.6%),  10.8 J projected
```

Per decode layer: one attention cb + one MoE cb. The CPU's
remaining per-token work: residual adds, router top-k,
embedding gather, sampling — and the receipt.

## Attention front half on-GPU (wave 77a)

Third compute kernel: fused per-head QK-norm + rotate-half
RoPE (one threadgroup per head; the inv_freq table is a
device buffer, so YaRN rides for free). Mellum's window is
now `[rmsnorm → q/k/v GEMMs → qk-norm → rope]` — ONE
command buffer, q and k returning ready-to-cache. Verified
live: text and attribution identical to wave 76; chain
correctness pinned vs a CPU reference (GEMMs + per-head
norm + rotation).

Kernel set: {silu·mul, rmsnorm, qknorm_rope}. The remaining
CPU steps per attention block are the KV append, the
attention kernel itself, and o_proj — the wave-77b targets
(device-resident KV cache is the prerequisite).

## RMSNorm joins the GPU — norm→q/k/v in one cb (wave 76)

Second compute kernel on the GPU member: a single-
threadgroup tree-reduction RMSNorm in MSL, writing the
normed activation as f16 in place — the layer input
uploads ONCE as f32 and the q/k/v GEMMs read the on-device
normed value directly. One command buffer for
norm → projections; the CPU-side normed buffer and all
three activation casts disappear from the fused path.
Correctness pinned vs a CPU reference of the chain;
three-tier fallback intact (norm window → CPU norm +
wave-71 window → CPU batch).

Live (Mellum, loaded box): correct text, decode dipping to
~190 ms/token between load spikes, GPU at 64.8% of GFLOPs
for 9.5 J projected vs CPU 120 J.

The kernel set is now {silu·mul, rmsnorm}; each addition
widens the single-cb span. Remaining for whole-layer-on-
device: QK-norm, RoPE, attention, and keeping h resident
across the residual adds.

## Whole-layer MoE window — one command buffer (wave 75)

`gemm_moe_window_f16`: gate/up GEMMs → an on-GPU silu·mul
kernel (12 lines of MSL, writing f16 so the down-GEMMs read
it in place) → down GEMMs, all in ONE command buffer. The
intermediate never returns to the CPU; the only readback is
the per-pick output rows. Mellum tries this first and falls
back to the two-window path, which itself falls back to the
CPU batch scope — three tiers, all bit-or-tolerance
verified.

Live: identical text and identical attribution to wave 72
(same op counts, half the command buffers per MoE layer);
the cb-merge gain is within load noise on this box today —
recorded as structure shipped, wall-clock claim deferred to
a quiet-box measurement.

## Fused windows complete the family set (wave 74)

NemotronH's attention layers get the q/k/v window; and
Falcon-H1's PARALLEL topology yields the widest window in
the set — in_proj + q/k/v, FOUR ops off the same layer
input (with their distinct μP pre-scales as per-job
activations), one command buffer. Falcon's tiny 0.5B
projections sit below the 2 MiB window floor on this
checkpoint, so its window lowers to the single CPU batch
scope — the structure is in place for the larger Falcon-H1
sizes where the floor passes. Both verified live with
their established outputs (Nemotron " Paris.\"] output",
Falcon " Paris. It is known for its historical").

All five families now share the fused-window decode shape
wherever their topology allows; dependent projections
(o_proj, out_proj, MLP down) remain sequential by data
flow — a GPU-side activation/attention op is the lever
that would widen windows further.

## Fused windows reach all three MoE families (wave 73)

Cohere2 (North) and Granite h-tiny's decode expert stages
lower onto the same `matmul_bf16_auto_multi` windows as
Mellum — per-tensor experts (North) and fused-tensor expert
slices (Granite) both ride the path; the fabric-off
fallback stays the wave-63 single CPU batch scope, bit-for-
bit. Verified live: h-tiny " Paris. Paris is the largest
city in France…" (~180–220 ms/token under fabric), North
" Paris, and" (decode improving as residency warms; its
1.5 MB down-projections sit below the 2 MiB window floor —
partial coverage, recorded rather than hidden).

## Expert FFNs join the fused windows (wave 72)

Mellum's decode MoE stages — 16 gate/up GEMMs, CPU silu,
8 down GEMMs — lower onto `matmul_bf16_auto_multi`: two
GPU command buffers per layer when the experts are F16-
resident, falling back to the wave-63 single CPU batch
scope otherwise (the fallback preserves that shape exactly
— never one dispatch per job).

Steady state (loaded box, fabric ON, 24 GiB residency,
24 tokens):

```
decode ~270–460 ms/token (load-dependent)
local/cpu-neon:   3747 ops, 36.2 GFLOP (28.7%), 136.5 J projected
local/metal-gpu: 17412 ops, 89.9 GFLOP (71.3%),  14.5 J projected
```

71% of the work on the GPU for ~10% of the projected
energy — the attribution line the fabric arc was built to
produce. Early tokens pay the one-time expert residency
conversion (~22 GiB f16 across 64×3×28 matrices, lazily on
first routing). Remaining knob: two cbs per MoE layer
(silu splits the stages); a GPU activation op would fuse
them to one.

## Fused dispatch windows — q/k/v in one command buffer (wave 71)

`gemm_bt_resident_f16_multi` encodes MANY GEMMs into ONE
command buffer; `matmul_bf16_auto_multi` is the fused
dispatch window (all-or-nothing F16 residency, 2 MiB floor
for window members vs 64 MiB for solo ops — the cb cost
amortizes across the window). Mellum's decode q/k/v — three
projections off the same normed input — ride one cb.

Live (loaded box, fabric ON, floor 128):

```
decode: ~161 ms/token            (arc start, wave 62: 2.8 s — 17×)
local/cpu-neon:   2851 ops, 27.6 GFLOP (77.0%), 104.1 J projected
local/metal-gpu:   596 ops,  8.2 GFLOP (23.0%),   1.3 J projected
```

THE BUG THE BRING-UP CAUGHT: first fused run produced
fluent garbage (" Estab Estab…"). Not the cb batching — a
same-shape-twice-in-one-cb test exonerated kernel-object
reuse — but the hand-rolled f32→f16 SUBNORMAL encoder
shifting 10 bits too far: every weight in [6e-8, 6.1e-5]
(plentiful in q/k/v, absent in lm_head, which is why wave
70 looked fine) collapsed to garbage. Now pinned by
known-answer bit patterns including the subnormal range.
Also: the F16 upload no longer REFUSES tensors for
sub-subnormal stragglers — they flush to zero (bounded
6e-8 absolute, standard f16 underflow); refusal is
reserved for overflow.

## Half-width GPU residents + the BF16 negative result (wave 70)

Context: WWDC26 / the M5 cycle put BF16 tensors front and
center (Metal 4 `MTLTensor`, `MTLTensorDataType.bfloat16`,
M5 GPU Neural Accelerators unlocked in macOS 26.2, ~4× M4
LLM peak via MLX). This box is an M3 Ultra on macOS 26.5 —
no Neural Accelerators, but the API surface applies.

PROBED FACT (cost a SIGABRT to learn): classic
`MPSMatrixMultiplication` hard-asserts on BF16 inputs —
"Input data type must be one of Float32, Float16, Int8,
Int16". BF16-native GEMM lives on the Metal-4
MTLTensor / MPSGraph surface (future bindings lift), and
the assert is an abort, NOT a catchable ObjC exception:
capability on this API is gated by knowledge, not
probe-by-attempt.

The assertion message hands over the win available today:
**F16 residents.** Weights convert BF16 → F16 once at
upload — EXACT in range (bf16's 7 mantissa bits widen into
f16's 10; the only approximation in the path is the per-call
f32→f16 activation cast) — for HALF the resident bytes and
half the per-call GPU streaming of the f32 path. Conversion
REFUSES any weight outside f16's faithful range
(|v| > 65504 or < 2⁻²⁴) and the loader falls back to the
f32 resident — honest fallback, never silent truncation.
Mellum2 lm_head resident: 0.9 GiB f32 → 0.45 GiB f16,
identical generated text, GPU certification tolerance
unchanged (the strict abs+rel check still passes).

When M5 hosts join the fleet, the Metal-4 tensor path is
the upgrade that reads checkpoint BF16 with no conversion
at all — and lights the Neural Accelerators. The fabric's
probe-per-host doctrine already handles the heterogeneity.

## Fabric attribution — the run carries its breakdown (wave 68)

Every dispatched matmul records (member, GFLOPs) in an
attribution ledger priced by the SAME calibration figures
the placer used — attribution and placement cannot drift
apart. `ledger_snapshot()` is serializable; the decode
examples print it per run.

First live attribution (Mellum2, fabric ON, 6 tokens):

```
local/cpu-neon:   3159 ops,  29.8 GFLOP (91.6%), 112.5 J projected
local/metal-gpu:     6 ops,   2.7 GFLOP ( 8.4%),   0.4 J projected
```

The receipt itself shows the next lever: the GPU carried
8.4% of the work for 0.3% of the projected energy — the
~3 000 small CPU ops (experts, q/o at m=1) are where the
joules live, and they need command-buffer batching before
the GPU can take them (one cb per 4-MFLOP expert op would
cost more than it saves; the ledger is how we know).

## Measured placement — the fabric runs on this machine's joules (wave 67)

`fabric_calibrate` (--features rayon,fabric-meter) runs the
same BF16 matmul workloads on every serving member under
the IOReport SoC counter (idle-subtracted — a MODEL on a
hardware reading; the tool prints the caveat and records
the method in the output file). `Fabric::discover()` loads
`~/.jouleclaw/fabric-calibration.json` and overlays the
measured µJ/GFLOP as `Measured` provenance — placement
runs on this machine's numbers, not a table written for
another one.

First calibration (loaded box — idle read 100 W with
parallel sessions running; RATIOS are the signal):

| op | member | ms/op | mJ/op |
|---|---|---:|---:|
| decode GEMV (m=1, lm_head class) | NEON | 43.0 | 924 |
| | AMX | 73.3 | 2 375 |
| | **GPU (resident)** | **2.8** | **242** |
| prefill batch (m=16) | NEON | 28.5 | 1 140 |
| | AMX | 6.2 | 262 |
| | **GPU (resident)** | **1.4** | **49** |

Wave 66's open question, answered by a meter: the GPU wins
decode GEMV on BOTH axes — 15× wall-clock AND 3.8× joules
vs NEON — and prefill by 20×/23×. AMX is honestly the worst
decode member (the per-call widen) and a solid middle for
batch. The energy-optimal member was the GPU all along;
now the placer knows it from measurement, not assertion.

## Metal GPU member certified + placed (wave 66)

The third local fabric member goes live. `backend-apple`
gains a RESIDENT-WEIGHT path: cached device+queue, weights
widened once into shared MTLBuffers (UMA — the GPU reads in
place, results are CPU-visible at command-buffer
completion), `transposeRight` GEMM so weights stay in their
natural `[n,k]` layout, and — the fix that mattered —
**MPS kernel objects cached per shape**: building
`MPSMatrixMultiplication` per call cost ~0.9 s on a loaded
box and single-handedly erased the GPU win in the first
cut. Encode is cheap; pipeline setup is not.

Certification: `gpu_member_matches_neon_oracle` (same
abs+rel contract as AMX). `jouleclaw-fabric` flips
`local/metal-gpu` to ToleranceCertified citing it, and the
placer's GPU dispatch floor is CALIBRATED to the measured
~20 GFLOP AMX↔MPS crossover (~100 mJ effective — the
placer must reproduce the measured crossover or it is
fiction). An uncertified-member synthetic test pins that
the gate stays structural.

Loaded-box A/B (fabric OFF → ON), Mellum2, identical text:

| | OFF | ON |
|---|---|---|
| prefill (6 tok) | 15.4 s | 6.7 s (2.3×) |
| steady decode | ~1.2 s/tok | ~1.2 s/tok (neutral) |

Decode wall-clock is NEUTRAL under load with the lm_head
resident on GPU — the decode question is now JOULES at
equal latency (GPU GEMV at ~20 W vs 10 CPU cores at ~5 W
each), which requires the IOReport meter wave to answer
honestly. No claim until measured.

## Fabric placement — the forwards become hardware-aware (wave 65)

`jouleclaw-fabric` gives the runtime its compute awareness:
detect local silicon (NEON / AMX / Metal — probed live, with
honest verified flags), enumerate remote reach from
credentials (CUDA fleets, LPU/WSE/RDU/Tenstorrent — presence
of keys, never claimed as verified), and place each op by
projected µJ. All five model families' projections now route
through `fabric_dispatch::matmul_bf16_auto`; with
`JOULECLAW_FABRIC=1` (+ `--features fabric`), batch-shaped
matmuls run on the AMX matrix unit via Accelerate
`sgemm_f32_bt`, with NEON row-major BF16→f32 tile widening.

Same-load A/B (fabric OFF → ON):

| model prefill | OFF | ON | gain |
|---|---|---|---|
| Granite h-micro, 15 tok | 5 992 ms | 2 949 ms | 2.0× |
| Mellum2, 6 tok | 12 567 ms | 4 396 ms | 2.9× |

Identical generated text. Numerics contract is per-member:
NEON stays the bit-exact oracle (fabric OFF default is
byte-identical to wave 64); the AMX member is certified
against the oracle to abs+rel tolerance — sgemm accumulation
order differs, and the docs say so instead of pretending.

Lesson the first attempt taught: a scalar TRANSPOSED widen
made fabric-ON 2.3× SLOWER than no AMX at all — strided
stores ate the GEMM win. NEON row-major widen +
transposed-B sgemm (`sgemm_f32_bt`) flipped it to 2-3×
faster. The widen is the product; the GEMM is table stakes.

## Decode dispatch floor (wave 62)

Mellum2 decode exposed a dispatch pathology: ~800 threaded
matmul calls per token (24 small expert projections × 28
layers + attention), each paying a rayon pool wakeup. On an
OVERSUBSCRIBED box (parallel build/agent sessions; wakeups
balloon to ~25 ms of sys time each under park/unpark
contention) that overhead dwarfed the math:

| config | Mellum2 decode |
|---|---|
| 10 threads, no floor | ~2.8 s/token (21.5 s sys vs 8.1 s user) |
| RAYON_NUM_THREADS=1 | ~650 ms/token |
| 16 MiB floor (new default) | small expert matmuls inline |
| `JOULECLAW_PAR_MIN_WEIGHT_MB=128` | **~0.5 s/token** (only lm_head fans out) |

`matmul_bf16_threaded` now runs single-thread NEON inline
below a weight-bytes floor (default 16 MiB ≈ 1 ms of
single-core streaming — an order of magnitude above a
healthy dispatch, so quiet-box behavior of the existing
families is unchanged; their big projections stay
threaded). On a loaded box raise the env knob. All numbers
above were measured under heavy external load — treat them
as the loaded-box regime, not the model's ceiling.

## Mellum2 / JetBrains 12B-A2.5B (wave 61)

Fifth topology consumer: JetBrains/Mellum2-12B-A2.5B-Instruct
(released 2026-05-26; Qwen3-MoE attention lineage + Laguna
block per the reference modular file). New ground vs the
North bring-up: per-head **QK-norm** before RoPE,
**per-layer-type RoPE** — YaRN (factor 16 over an 8192
window, explicit attention_factor on cos/sin) on the 7
full-attention layers, default RoPE on the 21 sliding
layers (window only 1024) — and the **Qwen3 gating
lineage**: softmax over all 64 router logits → top-8 →
renormalize. Every layer is sparse. YaRN inv_freq blend is
spot-checked against the reference formula in tests;
prefill parity is bit-identical across the window trim AND
the YaRN layer.

Smoke: "The capital of France is" → " Paris.  \nParis is
located in". ~2.7 s/token cold off the external USB drive
(22.6 GiB, 5 shards); warm decode untested.

Tokenizer: second pre-tokenizer gap in two waves — Mellum
uses `Digits` (isolate each digit) + `ByteLevel
{ use_regex: true }` (the GPT-2 split applied internally)
instead of explicit `Split` entries. Both now map onto the
chained-split machinery from wave 60.

## North Mini Code on the CPU: decode 3.3–5.2× and prefill 5.9–10.6× against base at its default thread count, bit-identical (2026-11)

The pure-Rust CPU path (`--features rayon`) of North Mini Code 1.0
(Cohere2-MoE, 30B total / 3B active, BF16 safetensors), made
faster without changing a bit of any logit. Measured with
`examples/north_bench.rs` (prints every step's token id and a hash
of all 262144 logits' bits), base = `main` before this change. The
ratios in the title are median against median in the table under
**Speed**, with base at rayon's default thread count (every logical
CPU) and the new code at its own default; that section also gives
base at 12 threads and says what was not compared.

**Where the time went** (`JOULECLAW_PROF=1`, base, M3 Ultra under
load, ~600-token prompt): prefill 45.3 s = attention 36% (every
head on one thread, a scalar add chain per score), experts 50%
(one expert after another), q/k/v + o 11%; decode 212 ms/token =
attention 26%, q/k/v 20%, o_proj 20%, experts 28%, lm_head 3%.
lm_head streams its 1 GB in ~7 ms (~140 GB/s) while each layer's
16–50 MB projections moved 12–40 GB/s: decode was bound by its ~7
dispatches per layer (three of them single-threaded GEMVs) and a
GEMV kernel held to ~16–20 GB/s per core by its two accumulator
chains, not by memory bandwidth. Base decode under load was also
fastest at 8 threads (94 ms) and slowest at the default 32
(174 ms).

**What changed** (each output keeps its exact reduction order):

- the BF16 kernel computes blocks of 4 activation rows × 2 weight
  rows, 2 × 4 or 1 × 4 (up to 16 accumulator chains instead of 2)
  with a 1 KiB prefetch;
- decode runs three dispatches per layer (q/k/v + router;
  attention heads + expert gate/up; o_proj + expert down) of
  ~256 KiB row chunks on a pool whose threads stay awake between
  dispatches (`spin_pool`);
- attention runs per head (decode) and per (head, 16 query rows)
  (prefill) in parallel, with a bit-exact NEON score kernel;
- prefill runs the experts in parallel and the projections as
  64 × 64 output tiles; RoPE angles are computed once per position;
- `forward_extend` batches tokens into a state that already holds
  a context, bit-identical to feeding them one decode step at a
  time;
- threads: one per performance core, at most 16
  (`JOULECLAW_THREADS` overrides).

**Identity.** Five prompts (a raw completion, Python, OpenQASM, a
596-token and a 2698-token Rust prompt), 72 greedy tokens each:
the same token ids and the same hash of every step's logits as
base, on an M3 Ultra and on an M4 Max, and the same on both
machines; every repeat of the timing runs below gave that one hash
sequence too. Unit tests pin the kernels and both forward passes
against the old code bit for bit.

**What is re-checked without the checkpoint.**
`src/cohere2_moe_reference.txt` holds logit hashes of three synthetic
Cohere2-MoE checkpoints, generated by base's `forward_prompt` and
`forward_token` (with its one-row BF16 kernel) on macOS arm64, Linux
arm64 and Linux x86_64. The three platforms give different bits, so
each has its own lines. One checkpoint is wide enough that decode
splits its projections into several row chunks and prefill runs the
2-D tiles. The test `stored_reference::logits_match_the_stored_reference`
recomputes the hashes with the reference passes, the CPU passes and
`forward_extend`, at the default thread count and at 1 and 3 threads.
CI runs it in the Linux arm64 workspace job. That job's container has
glibc 2.35, and the Linux arm64 lines were made with 2.39: with 2.35's
libm loaded in its place on the host that made them, the test gives
the stored lines too, and CI's own first run of it was pending when
this was written. The macOS lines are checked by running the tests on
a Mac. `examples/cohere2_reference.rs` generates the lines, at commit
30c4cf1ee, for a platform not yet in the file. **What still needs the checkpoint:** identity at North's
own sizes (hidden 2048, 128 experts with 8 routed, a 262144-row
lm_head, contexts past the 4096-token window) and on the five prompts
above, whose text and hash sequences are not stored, and every speed
figure here. The per-run logs are not in the repository.

**Speed.** Medians of 4 runs (the 72-token identity run and three
48-token timing runs; base at 12 threads had only the three), base
and new interleaved in the same conditions, each at its default
thread count unless the column says otherwise: base at rayon's
default, every logical CPU (32 on the M3 Ultra, 16 on the M4 Max),
the new code at one thread per performance core (16, its cap, on the
M3 Ultra; 12 on the M4 Max). M3 Ultra (24P+8E): shared, load average
30–130, with another North benchmark and CI running. M4 Max
(12P+4E): lightly loaded (load average 7–29).

| | M3 Ultra base, 32 thr | M3 Ultra base, 12 thr | M3 Ultra new, 16 thr | M4 Max base, 16 thr | M4 Max base, 12 thr | M4 Max new, 12 thr |
|---|---|---|---|---|---|---|
| decode at ~600 context | 267 ms | 168 ms | **54 ms** | 116 ms | 106 ms | **35 ms** |
| decode at ~2700 context | 405 ms | 364 ms | **80 ms** | 246 ms | 224 ms | **48 ms** |
| prefill, 596 tokens | 9.9 tok/s | 12.0 tok/s | **91 tok/s** | 17.8 tok/s | 19.8 tok/s | **105 tok/s** |
| prefill, 2698 tokens | 5.9 tok/s | 4.5 tok/s | **63 tok/s** | 9.0 tok/s | 9.2 tok/s | **79 tok/s** |

Paired run-for-run ratios against base at its default thread count
(base / new, all four pairs of each cell). M3 Ultra decode 4.3–4.8×
at ~600 and 5.0–6.6× at ~2700, apart from one pair in each column,
20.2× and 21.2×, where base's run took 3.6× and 5.8× its median
step (967 ms at load 71, 2358 ms at load 131); those two runs are
inside the table's medians.
Prefill 5.4–10.3× and 9.7–11.8×. M4 Max decode 3.3–3.5× at ~600
apart from the identity run's pair, 2.0× (the new code's median step
was 59.7 ms in that run, 34–37 ms in the others), and 5.1–6.3× at
~2700; prefill 5.2–6.8× and 7.0–9.0×.

Against base at 12 threads (the three timing pairs): M3 Ultra decode
2.9–3.8× and 3.3–5.8×, prefill 5.9–10.6× and 11.9–14.1×; M4 Max
decode 3.1× and 4.5–5.2×, prefill 5.8–6.1× and 8.4–9.5×. Base at 8
threads, its fastest in the thread sweep above (94 ms), was not run
in pairs with the new code: against this table's 54 ms that is about
1.7×, but measured under other load, so it is not a ratio this
benchmark establishes.

A 69-token follow-up turn fed into a ~630-token context
(`north_bench --then` and `--then-seq`): `forward_extend` 1.23 s
(M3 Ultra) / 0.79–0.80 s (M4 Max), against 3.3–4.1 s / 2.4 s one
`forward_token` at a time on the new code, and an estimated ~18 s at
base's decode speed (69 × its 267 ms median step; not measured);
same logits. A caller gets the first figure only by
feeding its turn through `forward_extend`; one that feeds
`forward_token` per token gets the second.

**Did not help, or not enough:** more threads (base and new were
both slower above 16 under load); for decode, running each step
inside a rayon pool rather than handing each of its loops to the
pool (no clear win once measured in-process; decode's loops now run
on the spin pool, while prefill, whose projections and experts are
rayon loops, does run inside its pool, `in_cpu_pool`); a rayon pool
of any size against the spin pool for decode (2.5–4.4× slower under
load: 164–465 ms against 65–106 on an M3 Ultra; on a quiet M4 Max
40–42 ms against 34 with the same 12 threads, 18–24% slower);
prefetch distances other than 1 KiB (256 B–4 KiB swept).

**Next levers:** the sliding layers trim their KV cache by moving
the whole window each token past 4096 positions (a ring buffer
removes ~0.6 GB of copies per token there); decode now moves ~6.5
GB of weights per token near what the CPU side reads from memory,
so the next step is fewer bytes (narrower weights, which would
change logits) or several sequences sharing each weight read; and
prefill experts are bound by the 4 × 2 NEON block's arithmetic.

## Cohere2-MoE / North Mini Code (wave 60)

Fourth topology, first attention-only one:
CohereLabs/North-Mini-Code-1.0 (30B total / 3B active)
decodes end-to-end. 49 layers in a 1-full-per-4 pattern
(sliding window 4096, ROLLING per-layer KV cache), Cohere
INTERLEAVED RoPE applied only on sliding (+ forced
dense-prefix) layers — full-attention layers carry no
position encoding. MoE is 128 tiny experts, top-8,
**sigmoid** gating (unnormalized — pinned by test). Block
is Cohere's parallel residual (one norm, attn + mlp summed
into one add). lm_head is the tied 262K×2048 embedding
table. Prefill reuses the wave 54/56 recipe (batched
projections + expert grouping), bit-identical parity
pinned across the rolling-window trim with a window-3 tiny
config.

Smoke: "The capital of France is" → " Paris, and the
capital of the United". Decode ~5 s/token cold — 57 GiB
over 49 shards streamed from the external USB drive; this
checkpoint is a correctness trophy, not a speed one.

Tokenizer: Cohere's tokenizer.json exposed a real bug —
the pre-tokenizer is a SEQUENCE of Split regexes (digit
grouping, then the GPT-style split), and the loader used
only the FIRST, dropping all non-digit text (the prompt
encoded to just BOS). Splits now chain (matches AND gaps
survive, Isolated behavior); the existing byte-equality
suites still pass.

## Falcon-H1 decode (wave 58)

Third hybrid topology: Falcon-H1-0.5B-Instruct runs
attention AND Mamba2 **in parallel inside every layer**
(NemotronH interleaves across layers; Granite stacks mixer
+ shared MLP). First RoPE hybrid here (`rope_theta = 1e11`,
rotate-half), no gated RMSNorm (plain `silu(z)` gate), and
a μP multiplier suite (embedding / ssm in+out / attention
in+out / key / five-section in_proj / mlp gate+down /
lm_head). Wave 47's conv/SSD primitives and the wave 21
attention kernel drop in unchanged; prefill follows the
wave 54 recipe with RoPE applied per position
(bit-identical parity pinned, including continued decode —
which would catch any RoPE position misalignment).

| model | decode | prefill |
|---|---|---|
| Falcon-H1-0.5B (36 layers, h=1024) | ~90 ms/token | ~49 ms/token |

Smoke: "The capital of France is" → " Paris. It is known
for its historical" — correct fluent text on the FIRST run
of the new architecture; the shared primitives carried it.

## NemotronH batched prefill (wave 54)

`forward_prompt` consumes the whole prompt in one pass:
projections run as one `m = T` matmul per weight (the
recurrent conv/SSD parts stay sequential per token — no
weight traffic there), attention prefills causal with
`seq_q = seq_kv = T`. Same per-row kernels → the last
position's logits and the resulting state are
BIT-IDENTICAL to sequential decode (pinned by test).

16-token prompt, rayon build, real 4B checkpoint:

| path | per-token | throughput |
|---|---|---|
| sequential decode | ~80 ms | 12.5 tok/s |
| batched prefill | ~42 ms | 23.8 tok/s |

The enabler was a threading-axis fix in
`matmul_bf16_threaded`: parallel-over-m made every thread
stream the WHOLE weight matrix (traffic = threads ×
weights); for `m < n / (2·threads)` it now parallelizes
over n (weights read once in total) and scatters the
per-chunk results. Prefill's remaining cost is the
sequential SSD recurrence — the chunked SSD form is the
next lever if long-prompt prefill ever dominates. (A
parallel-over-heads SSD path exists but is gated above
today's shapes: at 96×80×128 ≈ 1 MFLOP/layer, rayon
dispatch eats the gain — measured, not assumed.)

## Granite-4.0-H decode (waves 52-53)

Same harness (`examples/granite_decode.rs`), same prompt,
rayon builds, real IBM checkpoints:

| model | shape | per-token | continuation |
|---|---|---|---|
| granite-4.0-h-micro | 3B dense, 36 M + 4 attn | ~95 ms | " Paris. Paris is located in the northern" |
| granite-4.0-h-tiny | 7B-A1B MoE, 64 experts top-6 | ~200 ms | " Paris. Paris is the largest city in" |

h-micro pays more per parameter than Nemotron because every
layer carries a shared MLP (two residual sub-blocks). h-tiny
adds the routed-expert walk: router + 6 × (2 small expert
matmuls) per layer; expert matrices are sliced straight out
of the fused `[64, ·, ·]` mmap tensors, zero copies. The
small per-expert matmuls amortize rayon poorly — per-token
cost is router-dispatch latency, not bandwidth, so this is
the first decode shape where batching experts (not threads)
is the next lever.

Wave 55 adds `forward_prompt` (same recipe as NemotronH's
wave 54, bit-identical parity pinned for dense AND MoE):
h-micro prefill lands ~39 ms/token (~2.5× over its decode).
h-tiny initially got NO prefill speedup — its per-token
expert FFNs are inherently m=1 (each position routes to
its own experts), so batching the dense projections barely
moved the total.

Wave 56 closes that gap with expert batching: prompt
tokens are grouped by chosen expert, each routed expert's
weights stream ONCE for its whole token batch, and the
per-(token, pick) outputs are accumulated back in the
token's original pick order — so the f32 accumulation
sequence, and therefore the result, stays bit-identical
to the per-token walk (the MoE parity test pins it).

| h-tiny prefill, 15 tokens | total | per token |
|---|---|---|
| wave 55 (per-token experts) | ~3.0 s | ~200 ms |
| wave 56 (expert-batched) | ~1.4 s | ~94 ms (2.1×) |

Decode (single token) is untouched — there is no batch to
group at m=1.

## What this still leaves on the table
- **f16 accumulation via `cblas_sgemm` family + half-
  precision variants.** Apple Accelerate exposes
  `cblas_hgemm`-like paths via the `vDSP_*` family; not as
  obvious as sgemm. Future wave.
- **Multi-thread fan-out** — Accelerate is already
  multi-threaded internally for large enough sizes; rayon
  on top would help when many independent matmuls fire
  concurrently (e.g., multi-head attention with separate
  heads per worker).
- **Tile + transpose for the m=1 case** to recover the
  large-k single-token throughput that drops to 47 GFLOPS
  at 4096². Either better data layout or
  `cblas_sgemv` (single-vector specialization).

Each is its own bench-driven wave. The substrate primitives
(prefetch, tier store, mmap advice) are already in place;
the wins above prove they don't get in the way of any SIMD /
AMX / multi-thread acceleration.

## TierStore hot-path latency

`cargo bench --bench tier_store`. Same dev box:

| Operation | Time |
|---|---:|
| `pin_hot(first_tensor)` | **293 ns** |
| `bytes_borrow` (zero-copy &[u8]) | **10.6 ns** |
| `promote_to_warm(first_tensor)` idempotent | ~10 ns |

The `bytes_borrow` 10.6 ns is the hot path during inference —
every kernel call goes through it. Tiny.

## What these numbers DON'T capture

- **Cold-cache TTFT** for the `WillNeed` policy itself.
  Measured from a warm box; the page cache absorbed the
  earlier load. For real cold-load measurements: reboot or use
  Linux `drop_caches`. The 2.27 s `All` on ZAYA1 is the right
  order-of-magnitude for an actually-cold load.
- **SIMD kernel performance** (NEON / AVX-512 / AMX). All
  numbers above are scalar reference. SIMD is a follow-on.
- **Multi-thread scaling.** All single-threaded; rayon over
  the m-loop is straightforward to add.
- **GPU overlay performance.** Wave 5 (Metal UMA) and wave 9
  (CUDA UVA) shipped but require a Metal compute pipeline /
  CUDA kernel above them — measured by the consumer, not the
  substrate.

## Reproducing

```bash
cd jouleclaw/jouleclaw-rs

# Load any GGUF + see advise + classify behaviour:
cargo run --release -p jouleclaw-loader-gguf --example load_model -- \
    path/to/model.gguf

# Cold-load vs full-advise comparison:
for p in none top-n:32 all; do
  cargo run --release -p jouleclaw-loader-gguf --example load_model -- \
      path/to/model.gguf --policy=$p
done

# Kernel benches:
cargo bench -p jouleclaw-loader-gguf --bench matmul_q8_0
cargo bench -p jouleclaw-loader-gguf --bench mmap_strategy
cargo bench -p jouleclaw-loader-gguf --bench tier_store
```

All numbers above were captured against models in
`/Volumes/macos_4TB_external_1/neuronexus_ai/model_library/`
on the original dev box. Reproducing requires GGUF files in
the same formats; any standard llama.cpp / HuggingFace
distribution will do.