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 Optional, 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

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.