# 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`**, 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`, 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(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` 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.