gem/docs/research/engine-research-record.md

Research record

Verified model facts

On 2026-08-23, the official Hugging Face configs and ggml-org GGUF headers were checked directly:

• 12B: 48 layers, width 3840, 16 query heads, 8 SWA KV heads, 1 full KV head, FFN 15360.
• 31B: 60 layers, width 5376, 32 query heads, 16 SWA KV heads, 4 full KV heads, FFN 21504.
• Both use a repeating five sliding-attention/one full-attention pattern, 1024 SWA, 262144 maximum context, 256-dimensional SWA heads, and 512-dimensional full-attention heads.
• Dense 12B/31B have no cross-layer shared KV. Full-attention layers reuse the K projection as the V input and omit attn_v.weight at every sixth layer, but the cache rows diverge after learned K normalization versus unweighted V RMS normalization and cannot be aliased directly.
• The normalization dataflow nevertheless exposes a future lossless compact-cache candidate: store the unweighted normalized V row once and reconstruct historical K as V * key_norm inside full-attention kernels. This must first pass an operator speed/differential gate and then cover short prefill, long prefill, decode, verifier, and rollback paths before changing cache layout.
• Sparse files constructed from the first 64 MiB of the published Q4_0 artifacts validated the authentic directory without downloading weight data. The loader accepted 667 tensors for 12B and 833 for 31B. Published Q4_0 layout consists of Q4_0 matrices, Q8_0 token embeddings, and F32 norms/scales.
• The native raw-UTF-8 BPE tokenizer matched the pinned llama-tokenize oracle on the authentic 12B vocabulary for ASCII, Unicode, whitespace/newline runs, byte fallback, and canonical Gemma 4 chat/control-token prompts.
• The authentic 31B Q4_0 tensor directory bound directly from mmap into 60 fixed native layer descriptions. The first custom Metal Q4_0 SIMD-group matvec matched the scalar C result locally; this is correctness evidence, not yet a throughput result.
• The native model-level path now owns Q8_0 embedding/logit projection, all 60 Metal layers, F32 KV, sampling, detokenization, canonical Gemma 4 chat/tool formatting, and OpenAI HTTP/SSE transport. A real ../hax/build/hax process passes discovery, /props, request, streaming, usage, and final output against a fixture backend using this exact server implementation. Authentic-weight token generation and tool-call quality remain open until the full 31B artifact is available locally.
• The official 2026-07-09 Gemma 4 canonical template uses <|turn>/<turn|> roles, dedicated thinking-channel tokens, <|tool> declarations, <|tool_call>/<|tool_response> handshakes, and <|\"|> string delimiters. The native formatter implements that structured subset directly; it does not embed a general Jinja runtime.

Native Metal diagnostic: buffer ownership

On the 40-core M4 Max, the native Q4_0 SIMD kernel was measured at the actual 31B FFN projection shape (5376 input, 21504 output, 65,028,096 matrix bytes). The only changed axis was buffer mode. Each raw sample contains 200 synchronized dispatches:

• copy-per-dispatch: 13.35 GB/s mean, 1.36% coefficient of variation;
• persistent/no-copy: 107.43 GB/s mean, 4.19% coefficient of variation;
• delta: +704.97%, above the 4.19% observed noise floor.

The evidence is stored in artifacts/2026-08-23-q4-matvec. This validates persistent buffers as an architectural requirement, but promotion.eligible remains false because the deterministic real-model generation gate and model-level gem print-config do not exist yet. The persistent result is also far enough below expected platform memory bandwidth to justify profiling SIMD row geometry and command submission before drawing a decode-throughput conclusion.

A second one-axis diagnostic held persistent buffers and the 31B FFN shape fixed while changing command-buffer span from one dispatch to seven. Three 1,000-dispatch repetitions measured 99.65 GB/s versus 138.61 GB/s, a +39.10% delta above the 2.75% noise floor. Evidence is under artifacts/2026-08-23-command-span; promotion remains withheld for the same model-level gates.

The native Metal correctness path now includes weighted RMSNorm over one or many heads and NeoX/proportional RoPE with optional frequency factors. A dependent RMSNorm-to-Q4 projection and a two-matvec chain execute in one command buffer and match the scalar C oracle.

The native scheduler's token-span path passed the authentic 31B Q4_0 gate. Against layer span it improved pp from 7.100 to 7.561 t/s (+6.48%) and tg from 6.946 to 7.427 t/s (+6.93%) over three repetitions, above 1.35%/1.23% noise floors, with bit-identical native logits. The analyzer reported the single command_span axis comparable and promotion eligible, so token span is the default. Raw evidence and configuration records are in artifacts/2026-08-23-native-command-span-31b.

For prompt Hello (tokens [2, 9259]), the native layer scheduler and pinned llama.cpp oracle both selected token 100 and had the same top-10 ordering except for one adjacent swap. Across all 262144 logits, correlation was 0.999852, NMSE 3.03e-4, mean absolute error 0.0391, and maximum absolute error 0.3312. The comparison used the exact same 31B Q4_0 file (SHA-256 031dc1c5fa9c5a0abbf3c39c5173fb2af65f5ac2dc2a090268561d3c72dcd834).

The first real-model roofline showed the one-row Q4_0 kernel at roughly the same 139 GB/s as its microbenchmark, accounting for native tg 7.43 t/s versus llama.cpp 21.56 t/s. After the pinned reference MTL0 MUL_MAT backend suite passed all 1154 cases, the native kernel independently adapted its four-output-row SIMD geometry. The real-shape microbenchmark reached 331.93 GB/s (+139.47%); model pp reached 19.80 t/s (+161.82%) and tg reached 18.99 t/s (+155.65%). The native logits kept the same argmax and top-10 with NMSE 2.60e-11 versus the prior kernel. A follow-up two-SIMD-group threadgroup changed only that scheduling axis and failed promotion (+1.04% tg inside 3.52% noise, -1.67% pp), so the one-SIMD-group/four-row geometry remains selected. Evidence is under artifacts/2026-08-23-q4-multirow.

The server sampler previously sorted all 262144 vocabulary entries for every non-greedy token, even with the default top_k=40. A fixed-capacity worst-first heap now selects only the requested top-k set in O(V log k), followed by the same descending sort and top-p sampling over those k entries. The exhaustive top_k=0 path is unchanged. Three 20-iteration samples fell from 26.525 ms/token to 0.959 ms/token (27.66x, -96.38% latency; 0.11% CV). A 128-sample seeded test matches the full-sort token sequence, and equal-logit selection is deterministic by token ID. Evidence is under artifacts/2026-08-23-sampler-topk.

The first batched-Q4 experiment accumulated four input rows with 16 scalar accumulators per SIMD lane. It regressed the 31B FFN-shape diagnostic from 5,694 to 2,291 rows/s (-59.76%) and was rejected. Its replacement independently specializes the pinned oracle's legacy Apple simdgroup_matrix strategy: each 128-thread group dequantizes one 64x32 Q4 weight tile to half in threadgroup memory and reuses it across as many as 32 prompt rows, with float accumulation. The native five-row CPU differential passes. At 32 rows the diagnostic measured 26,796 rows/s versus 4,385 at one row (6.11x, +511.11%); the reported 1.74 TB/s is equivalent weight bandwidth from reuse, not physical DRAM bandwidth. The crossover is eight rows, so smaller batches retain the four-output-row matvec kernel. The fixed graph now uses this operator for prompt tiles up to 32 rows. On a 19-token authentic 31B prompt, tiled versus scalar-in-time logits retained the same argmax and top-10 set, correlation 0.9999989, and NMSE 3.17e-7. A one-axis token-span comparison (n_ubatch 1 versus 32) improved pp from 19.65 to 23.60 t/s (+20.13%) over three repetitions; tg changed +1.49%, within the 2.05% candidate noise floor. The analyzer marked promotion eligible, so batch 32 is the production default. Raw model evidence is in artifacts/2026-08-23-native-prefill-31b; operator records and configuration are in artifacts/2026-08-23-prefill-matmul4.

The original decode-attention kernel assigned one 32-lane SIMD group per query head, scanned every token serially, and redundantly evaluated softmax work in every lane. At the 31B full-attention shape (32 Q heads, 4 KV heads, width 512), throughput collapsed from 100.11 dispatches/s at 1K tokens to 2.784 at 32K, only about 12 GB/s of logical KV traffic. After the pinned reference's complete MTL0 FLASH_ATTN_EXT suite passed 4792/4792 cases, the native kernel independently adopted token-parallel blocked scheduling: 256 threads distribute score dot-products, reduce max/sum in threadgroup memory, then distribute value dimensions. It switches at eight live tokens, the measured crossover. At 32K it reaches 35.767 dispatches/s (12.85x, +1184.74%) and 153.6 logical GB/s. Authentic 16-token 31B logits retain the same argmax and top ten, with correlation 0.9999992 and NMSE 1.41e-7. At 256 prompt tokens, model pp rises from 25.91 to 58.24 t/s (+124.78%); tg changes +2.24% inside noise. The one-axis analyzer marks it promotion eligible, so token-parallel attention is the production default.

An F16 KV candidate halves cache memory and raises the 32K attention diagnostic another 8.25%. Authentic logits retain argmax/top-ten with NMSE 2.11e-7, but a clean 256-token model comparison changed pp only +0.77% and tg +0.12%, both inside variance. It therefore remains unpromoted and the default cache stays F32 until a longer model-level gate justifies the representation change. Raw evidence is in artifacts/2026-08-24-attention-kv.

A follow-up GQA-reuse experiment split attention into 256-token online-softmax tiles and reused each K/V tile across the query heads sharing a KV head. The operator result was attractive: versus the promoted parallel kernel, the 31B full-attention shape improved 1.81x at 2K, 2.81x at 8K, and 3.57x at 32K; the 1K SWA shape improved 1.99x. The authentic gate rejected it. At 1,024 prompt tokens the 31B logits still matched argmax and top-10 order, but NMSE had already accumulated to 4.20e-3 and mean absolute error to 0.326. The first full-attention GQA8 dispatch at 1,792 tokens produced all NaN logits, and a SWA-only GQA2 restriction also produced all NaNs by 2,048 tokens. Normalizing tile partials worsened the 1,024-token NMSE to 1.18e-2. A second design retained one global softmax by separating GQA score generation from value accumulation, but extra dispatch/scratch traffic erased the isolated speed advantage. All candidate kernels and runtime switches were removed; the production path remains the previously promoted per-query-head parallel kernel. Raw measurements, failed logits, and the rejection record are in artifacts/2026-08-24-blocked-attention.

This failure tightens the next attention gate: a candidate must compare repeated wrapped-ring operator outputs at the exact 31B SWA/full shapes, then pass at least a 2K authentic model-logit differential before performance timing. A small synthetic operator comparison is insufficient for an online-softmax decomposition because layer-to-layer numerical error can compound catastrophically.

The verified token-parallel kernel was then profiled directly with Metal command-buffer GPU timestamps because this Command Line Tools installation has no xctrace, and the M4 Max exposes only the GPUTimestamp counter set with stage-boundary sampling. At 1K full attention, changing command span from one to twenty raised throughput from 1,125.78 to 1,393.11 dispatches/s and GPU wall fraction from 79.36% to 98.52%; short-context submission overhead is therefore material. At 32K, span one was already 99.25% GPU time and span five reached 99.84%, while throughput changed only 0.91%. Long-context optimization must reduce kernel work or memory traffic rather than launch count. F16 KV reached 38.23 dispatches/s versus 35.67 for F32 in this operator profile, consistent with the earlier capacity result but still not a model-level promotion.

An exact-softmax follow-up staged normalized score weights once per 256-token value tile in threadgroup memory. It passed the native CPU differential, and the reference MTL0 FLASH_ATTN_EXT suite passed 4,792/4,792 before the local gate. A longer fresh 32K comparison measured only +1.00% wall and +0.98% GPU throughput, close to the 0.87% baseline coefficient of variation and below the 3% promotion floor. The candidate was removed; raw profiles and the rejection record are in artifacts/2026-08-24-attention-profile.

The successful follow-up preserved the same global score, maximum, denominator, and normalized weights, but changed the larger traffic term: value accumulation dispatches one workgroup per KV-head, eight-query-head block, and 64 output dimensions. Each workgroup loads a 32-token x 64-dimension V tile once and shares it with the applicable GQA heads. This avoids online-softmax state merging entirely. A post-promotion stress audit found that both this score pass and the old parallel kernel lacked a phase barrier: a fast thread could overwrite the reduced maximum in shared[0] with a denominator partial before every thread consumed it. Adding the barrier to both kernels made 100 independent test processes deterministic, with two identical invocations of each kernel at every exact 12B/31B full/SWA wrapped-ring shape.

All gates were rerun after the fix. Against the corrected per-query-head kernel, full 31B attention improves 2.38x at 1K and 2.67x at 32K; the 1K SWA shape improves 1.74x. At 2,048 authentic 31B prompt tokens, all 262,144 logits remain finite, argmax and the top-10 set match, correlation is 0.999793, and NMSE is 4.11e-4. The three-repetition model gate raises pp from 23.97 to 31.95 t/s (+33.29%, above the 12.18% noise floor). The measured tg delta is -10.60% inside a 20.31% noise floor and therefore is not a detected regression. Normal and ASan/UBSan suites pass 22/22, the production binary remains free of llama/ggml linkage, and gem agent returned exactly OK through real hax at a 4K context with port cleanup verified. Corrected raw evidence and the promotion record are in artifacts/2026-08-24-split-gqa-race-fix; the earlier record is marked superseded.

The local llama.cpp checkout remains useful for its Gemma 4 graph, converter, tokenizer behavior, Metal kernels, and DSV4 optimization techniques. Its DSV4 operators are not emitted by the Gemma 4 graph and are reference material, not a linked dependency or automatic speedup.

After split-GQA promotion, prefill projection became the next target. A Q4 64x64 prototype used eight SIMD groups so each dequantized 64-output by 32-K weight tile served 64 prompt rows. It passed a 65-row CPU differential including the tail and the reference MTL0 MUL_MAT suite passed 1,154/1,154. The larger tile nevertheless regressed the 31B FFN shape by 7.23% at 64 rows and 30.03% at the production 32-row batch. Extra SIMD/register pressure outweighed weight reuse, so the candidate and runtime API were removed. Raw evidence is in artifacts/2026-08-24-q4-prefill-64x64.

The next attention pass targeted the remaining K traffic. The promoted split-GQA score kernel had still loaded each K vector once per query head. A cooperative score pass now stages K once per KV-head, eight-query-head block, and 32-token chunk; eight SIMD groups compute the associated query heads, followed by a separate exact global normalization pass and the existing shared-V pass. This keeps the same global softmax and supports the exact 12B/31B full/SWA ratios.

The reference MTL0 FLASH_ATTN_EXT suite passed 4,792/4,792, and 100 independent wrapped-ring stress processes were deterministic. At 1K/32K full attention the operator improves 27.41%/49.76%; 1K SWA improves 3.09%. Authentic 2K logits are all finite with matching argmax/top-10 set, correlation 0.999824, and NMSE 4.64e-4. The 2K model result (+4.48% pp) was correctly withheld because it sat inside a 12.56% noise floor. A stable 4K gate then improved pp from 31.74 to 34.60 t/s (+9.00%, 2.56% noise) while tg improved 0.92%. Normal and sanitizer suites pass 22/22, linkage is native-only, and real hax returned exactly OK with cleanup verified. The cooperative-K path is now the default; evidence is in artifacts/2026-08-24-gqa-score-reuse.

The next prefill pass isolated threadgroup layout while preserving Q4_0 storage, the 64-output by 32-prompt tile, four SIMD groups, F16 staging, and F32 accumulation. A bank-aware 8-wide swizzle, independently adapted from the pinned MIT llama.cpp legacy MUL_MM layout, improves the 31B FFN projection 44.78% and the down projection 95.02%; every tested authentic projection/batch shape cleared its variance floor. The reference MTL0 MUL_MAT suite passed 1,154/1,154. A 33-row CPU differential covers the complete MMA tile and tail, and authentic logits are bit-identical across all 262,144 entries. The promotion gate improves 32-token pp 35.84% with 64-token generation unchanged; a 512-token confirmation improves pp 26.45%. Normal/sanitizer suites pass 22/22, linkage remains native-only, and real hax returned exactly OK with port cleanup verified. The linear kernel was removed; evidence is in artifacts/2026-08-24-q4-prefill-swizzle.

With swizzled projections promoted, a native profiler was added because xctrace is unavailable in the installed command-line developer tools. Metal command-buffer timestamps show 32/512-token prefill spending 99.07%/99.33% of wall time on the GPU while issuing 11,732/200,312 dispatches. Batching RoPE positions into one grid per Q/K tensor removes 3,720 dispatches at 32 tokens and 59,520 at 512 tokens, improving the profiling workloads 5.37%/6.71%. The one-axis authentic gate improves pp 12.10%, logits are bit-identical, the reference MTL0 ROPE suite passes 448/448, and normal/sanitizer/linkage/hax gates pass. Evidence is in artifacts/2026-08-24-prefill-dispatch-profile.

A follow-up non-wrapping batched-KV prototype removed another 1,860 dispatches at 32 tokens and 29,760 at 512 tokens, while preserving exact logits and wrapped-store operator correctness. Its first sequential runs appeared 5.68%/6.45% slower, but the post-removal row-wise control drifted below the candidate, exposing thermal/run-order confounding. The candidate was removed and is classified as inconclusive rather than a negative result. A retry must interleave both layouts in one process; row-wise KV storage remains production behavior.

FlashAttention and online-softmax primary sources were revisited through Paperbridge before the next prompt-attention design. The retained principle is IO-aware fusion without materializing scores in device memory; the rejected tile-local recurrence was not reused. The new initial-prompt kernel is bounded to at most 32 non-wrapping causal rows, holds the entire per-row score vector in threadgroup memory, and performs the same global maximum and denominator reductions as the proven parallel path. It passes a CPU differential, the pinned MTL0 FLASH_ATTN_EXT oracle (4,792/4,792), and a 2K authentic logit gate with identical argmax/top-10 order (correlation 0.9997926, NMSE 7.51e-4). It reduces the 32-token model profile from 8,012 to 3,152 dispatches and improves interleaved pp 6.85% without a tg regression outside variance. The unfused split-GQA route remains selectable as an explicit differential and benchmark oracle.

Extending the one-dispatch causal kernel beyond the initial tile was rejected after authentic-logit testing. Although 64/128/512-token profiles initially looked faster, both the first prototype's incomplete score initialization above 256 tokens and the corrected implementation changed repeated layer numerics beyond the quality gate. Production therefore retains that kernel only for the first 32 tokens. The accepted later-tile design instead batches the proven rowwise arithmetic: cooperative GQA score generation, exact global normalization, and shared-V accumulation become three grids over all rows. The operator result is bit-identical to the rowwise oracle at 308 live tokens. Authentic 2K logits preserve argmax and top-10 order (correlation 0.999835, NMSE 2.54e-4). At 512 tokens, dispatches fall from 135,932 to 52,232 and the formal interleaved model gate improves pp 16.57% when comparing the combined prompt path with the fully rowwise oracle. A refined one-axis control holds the promoted initial tile constant and measures only later-tile batching: pp improves from 71.01 to 86.80 t/s (+22.24%, 6.91% noise floor), while tg changes -1.11% inside variance. The MTL0 FLASH_ATTN_EXT suite passes 4,792/4,792; evidence is in artifacts/2026-08-24-causal-prefill-tiles.

The first later-tile promotion deliberately capped batched scores at 1,024 live tokens. Real hax project prompts were approximately 2.5K tokens, so the ten 31B full-attention layers still returned to 32 rowwise score/normalize/value sequences per tile while SWA layers stayed batched. Raising only the non-wrapping ceiling to 4,096 requires a bounded 16 MiB score arena for the 31B head shape; a naive 262K extension would require roughly 1 GiB and was rejected. No MSL arithmetic changed. A new 1,032-token operator differential is bit-identical to rowwise attention. Sequential 2K/4K profiles improve 9.30%/9.17%, reducing dispatches by 29,760/89,280. The exact one-axis, five-pair 2K gate holds the 1K path constant and improves pp from 50.88 to 54.18 t/s (+6.50%, 2.93% noise); tg changes +0.28% inside variance. Authentic 4K logits are finite with the same argmax and top-10 set (correlation 0.999379, NMSE 5.34e-4), although two close top-10 alternatives swap order. The MTL0 suite passes 4,792/4,792, normal and ASan suites pass 23/23, and real hax returns exactly OK with launcher cleanup verified. Evidence is in artifacts/2026-08-24-prefill-split-4096.

Decode is now close to a weight-bandwidth problem: the 17.99 GB 31B Q4_0 model at roughly 18.8 generated tokens/s implies about 338 GB/s of model-weight traffic, consistent with the native four-row Q4_0 microbenchmark's 336--344 GB/s controls. A paired half-block SIMD-lane geometry, independently adapted from the adjacent DeepSeek-V4 llama.cpp reference after its MTL0 MUL_MAT suite passed 1,154/1,154, passed the native CPU differential but reached only 311.43 GB/s. That is 7.29% below the preceding production control and 9.58% below the following control, so the candidate and its API were removed. Paperbridge review of BaseRT also ruled out directly adopting its native base_q4: it is an asymmetric group-64 format, and the published converter rejects quant-to-quant conversion. Requantizing GGUF Q4_0 would change model quality and is out of scope; only a lossless derived layout could be considered. The next decode axis is per-stage GPU timing of one authentic 31B token before choosing between lossless repacking and projection/epilogue fusion. Evidence is in artifacts/2026-08-24-q4-decode-research.

The corresponding 12B baseline reaches 108.53 pp and 43.21 tg with native F32 KV. The adjacent llama.cpp default reaches 82.48 pp and 52.68 tg with F16 KV and automatic flash attention; this is competitor context rather than a one-axis comparison. An authentic position-32 native profile attributes 32.80% of sampled GPU time to FFN gate/up, 18.42% to FFN down/output, 11.93% to QKV, 8.04% to attention output, and 12.04% to tied logits. A 12B-only revisit of the paired half-block lane layout was nearly logit-identical to row4, but the formal five-run A/B/A model gate changed generation from a 41.364 t/s surrounding baseline to 41.144 t/s (-0.53%) with higher variance. It was removed. Evidence is in artifacts/2026-08-24-native-12b-baseline and artifacts/2026-08-24-q4-halfblock-12b-model-gate.

The batch-one Q4 harness now rotates explicit independently backed matrices and reports the total working set. Sustained 12B gate/up traffic falls from 335.64 GB/s for one resident 31.64 MiB matrix to 313.75 GB/s for a 253.13 MiB eight-matrix stream (-6.52%). A single 31B gate/up matrix is already 62.02 MiB: expanding to a 248.06 MiB four-matrix stream changes 372.48 to 376.85 GB/s (+1.17%, inside variance). At the 12B streaming shape, increasing only command span from seven to 336 raises 314.33 to 421.62 GB/s (+34.14%). Production already records an entire token in one command buffer; future kernel screens must therefore use both a nonresident working set and graph-like span. Evidence is in artifacts/2026-08-24-q4-decode-streaming-harness.

That corrected evidence shifted the next bounded axis to tied Q8 logits, which account for 12.04% of sampled 12B and 7.34% of sampled 31B decode GPU time. A four-output-row Q8 matvec passed the native CPU differential and the immediately preceding adjacent MTL0 MUL_MAT oracle passed 1,154/1,154. It nevertheless regressed the clean 12B exact-shape operator from a 329.18 rows/s surrounding baseline to 295.50 rows/s (-10.23%). The 31B long run entered thermal collapse; even its optimistic early samples implied only a single-digit stage gain, whose end-to-end ceiling is below 0.5%. The diagnostic kernel, API, test, and switch were removed before model integration. Evidence is in artifacts/2026-08-24-q8-row4-decode.

M4 Max exposes the timestamp counter set at compute-stage boundaries but not at arbitrary dispatch boundaries. A diagnostic-only profiler therefore keeps one command buffer and places 542 sampled compute passes around embedding, nine fixed groups in each of 60 layers, and logits. At decode position 32, three runs cover 97.64% of the command buffer's GPU interval. Projection-dominated groups account for 81.29% of sampled time, KV store/attention 6.04%, and tied-output logits 7.34%; explicit norm, transform, and activation groups together are only 5.30%. The five-run production control reaches 20.27 tg, or 49.34 ms/token. Sample boundaries add diagnostic overhead, but the profiled route's 262,144 logits are bit-identical to the same ordinary one-token schedule. This rules out norm/activation fusion as the next end-to-end 3% lever. Evidence is in artifacts/2026-08-24-decode-stage-profile.

The trace justified one lossless derived-layout experiment. Within every Q4_0 row, all binary16 scales were grouped before the nibble payload while preserving the exact 18 bytes per block. The candidate passed its CPU differential but reached 324.21 GB/s at the 31B FFN shape, versus 371.49/382.49 GB/s surrounding direct-GGUF controls. The sidecar layout and kernel were removed; direct Q4_0 remains production. Evidence is in artifacts/2026-08-24-q4-split-layout.

Paperbridge retrieval of speculative sampling and Draft & Verify identifies the next architectural lever: draft a short continuation, score it in parallel with the target, and use rejection sampling to preserve the target distribution. The necessary premise is that short parallel scoring approaches the latency of one target token. Gem's current 31B FFN projection takes 0.188 ms for one row, 0.627 ms for four rowwise matvecs, and 0.798/0.808 ms for the existing 64x32 MMA route at 8/16 rows. The fixed 32-row tile therefore costs 4.25x a decode projection at eight rows and cannot yet serve as an efficient verifier. The next bounded kernel axis is a 64-output x 8-token specialization; only after it clears correctness and latency gates should the engine add draft/accept/rollback state. The aligned 12B draft GGUF is not currently local, so acceptance-rate measurement remains a later gate. Evidence is in artifacts/2026-08-24-speculative-feasibility.

The fixed 64-output x 8-token Q4_0 verifier specialization passes the adjacent reference backend's 1,154/1,154 MUL_MAT cases and Gem's native CPU differential. It uses two SIMD groups and 4.5 KiB of threadgroup memory instead of four groups and 8 KiB, while preserving the direct GGUF layout, binary16 staging, and float accumulation. At the representative 31B FFN projection it lowers an eight-row call from 0.799/0.800 ms in surrounding 64x32 controls to 0.500 ms (+59.76% row throughput), and is 2.45x faster than eight rowwise matvecs. Authentic 31B logits after an eight-token prefill are bit-identical across all 262,144 vocabulary entries. The five-run full-model gate improves pp from 29.50 to 31.17 t/s (+5.65%, 0.073% noise floor); tg changes +0.21% inside variance, and a following old-kernel control returns to 29.48 pp. Production therefore selects 64x8 only at exactly eight rows, retaining matvec below eight and 64x32 above it. Normal and ASan suites pass 24/24. This clears the dense-projection premise for short verification, but not speculative decoding as a whole: batched tied-output Q8_0 logits, an authentic verifier profile, draft acceptance, and KV rollback remain gated work. Evidence is in artifacts/2026-08-24-q4-verify8.

The tied Q8_0 output projection now has a matching fixed-eight-row SIMD-matrix path. At the real 5,376 x 262,144 embedding shape it computes all eight vocabulary distributions in 4.96 ms versus 30.43 ms for eight rowwise matvecs. In the authentic 31B graph, changing only the final Q8 schedule reduces the complete eight-token verifier pass from 291.85 to 265.39 ms (-9.06%, five runs) and removes seven dispatches. All 2,097,152 logits are finite; relative to the rowwise target they retain the same global top-10 order with correlation 0.999999936, NMSE 1.92e-8, MAE 0.00112, and maximum absolute error 0.00952. The small difference is expected from binary16 SIMD operand staging versus the float matvec path. This kernel remains scoped to the explicit verifier API; ordinary one-token logits retain the Q8 matvec. Normal and ASan suites pass 24/24. Evidence is in artifacts/2026-08-24-q8-verify8.

Doubling the fixed Q4_0 64x8 tile from two to four SIMD groups splits the same 64 outputs across four 16-column slices. The direct GGUF layout, binary16 operand staging, float accumulation, and 4.5 KiB threadgroup allocation do not change. At the 31B FFN shape, four groups lower the call from 0.514 to 0.426 ms (+20.75% row throughput). Across the authentic target graph, with batched Q8 held constant, the effect compounds: eight-token verification falls from 265.63 to 166.22 ms (-37.42%, five paired runs), and all 2,097,152 logits are bit-identical. The normal five-run model gate improves eight-token pp from 31.24 to 51.42 t/s (+64.61%, 0.324% noise floor), while tg changes -0.21% inside variance; a following two-group baseline returns to 31.20 pp. Production now selects four groups at exactly eight rows only. Normal and ASan suites pass 24/24. This makes target verification about 3.4x the current one-token decode latency before draft cost, so speculative decoding has a plausible but acceptance-sensitive ceiling; the next empirical gate is a local aligned 12B draft. Evidence is in artifacts/2026-08-24-q4-verify8-4sg.

The official 12B Q4_0 GGUF is now locally available and hash-verified (7,219,673,216 bytes, SHA-256 3712b9bd32cae83a22f67ee7a4466d8d7a4f21646ac8a07d19bf9418e8767a70); it is accepted by the strict native 48-layer profile and is excluded from YAMS indexing. Native 12B decode reaches 44.92 t/s, or 22.26 ms/token. Eight draft tokens plus the promoted 31B verifier cost about 344.29 ms, so break-even against 49.08 ms direct target decode requires 7.02 emitted tokens per round, equivalently just over six accepted drafts on average. A new diagnostic replays shared prefixes and measures greedy 12B/31B agreement over four agent-style prompts and 64 K=8 rounds. Mean accepted prefixes are 0.375--1.188 for general prompts and 3.625 for a repetitive JSON prompt; none reaches the 6.015 break-even requirement. Even the best case estimates 74.44 ms/token before scheduler overhead, 51.7% slower than direct target decode. A greedy dual-model scheduler is therefore rejected. This does not yet reject stochastic speculative sampling: the next bounded gate is exact temperature-0.8/top-k-40/top-p-0.95 distribution overlap. If that fails, Gemma 4's small MTP artifact or self-speculation is structurally more promising than paying for eight full 12B draft passes. Evidence is in artifacts/2026-08-24-speculative-12b-gate.

Exact stochastic filtering confirms the full 12B draft is not viable: across three seeds and four agent-style prompts, its mean emitted tokens per K=8 round remains below the 7.015 break-even point. The official Gemma 4 assistant sidecars are structurally different. Their four 1,024-wide layers reuse the target's final SWA/full KV views, and the matching 31B Q4_0 artifact is only about 267 MiB (SHA-256 f8f1b8c89b72da2b7391d6fe29d26dd373efe9624a3d5b3f0e99e62579b12c08). Gem now validates both the 3,840-wide 12B and 5,376-wide 31B sidecar contracts and executes the assistant natively: target-sampled token plus post-final-norm hidden, 10,752-to-1,024 pre-projection, three shared-SWA query-only blocks, one shared-full query-only block, tied Q4_0 output, and recurrent 1,024-to-5,376 post-projection. Seven warm 31B assistant steps cost 32.75--34.37 ms total.

The native server now performs exact stochastic MTP rounds: one guaranteed target sample, seven assistant proposals, one fixed-eight target verifier, acceptance probability min(1, q/p), residual sampling from normalized (q-p)+, and a target bonus sample when all proposals pass. A deterministic streamed 64-token request transfers the same 3,255 bytes with and without MTP. After one discarded warm-up, three direct requests average 1.555971 s and three MTP requests average 1.478071 s, a 5.01% latency reduction / 1.053x speedup. The improvement exceeds both 3% and the larger 1.19% coefficient of variation. MTP is therefore promoted for gem serve/agent --mtp. Before each verifier the engine snapshots only SWA ring slots that would overwrite live entries; rejection restores the unaccepted tail before the correction token, so MTP remains correct past position 1,024 without copying the full cache. Evidence and raw JSON are in artifacts/2026-08-24-native-mtp-server.

An M4 Max stage trace then isolated the MTP drafter itself. Across three seven-token runs, the four assistant layers, pre-projection, tied output, and post-projection consumed 6.87--6.92 ms of sampled GPU time, while wall time remained 32.92--33.97 ms. A detailed run spent 30.69 ms in command submission/wait versus 0.52 ms encoding and 0.11 ms logits readback, so assistant matmul tuning was not the causal next axis. PaperBridge retrieved Wagner et al., Optimized Speculative Sampling for GPU Hardware Accelerators (arXiv:2406.11016), whose exact method likewise keeps probability work parallel and writes back only acceptance-relevant results. Gem does not adopt the paper's inexact sigmoid alternative.

The promoted response keeps seven recurrent drafts in one Metal command buffer. Its fixed top-k40 block sort/merge is independently adapted from the pinned MIT llama.cpp Metal TOP_K implementation; the M4 Max backend-op oracle passed 445/445 cases before adaptation. A local 262,144-token CPU differential and indirect Q8 embedding differential pass. Device nucleus draws export the same sparse probabilities used to generate each proposal, preserving exact speculative acceptance; unsupported top-k or repetition-penalty settings retain the CPU path. With one discarded warm-up and three interleaved repetitions, identical greedy proposals improve from 34.262 to 31.764 ms for seven drafts (-7.29%, 1.079x). The accompanying deterministic HTTP/SSE gate averages 1.793911 s direct and 1.701882 s device-MTP (-5.13%, 1.054x), with identical HTTP status and byte counts. Evidence is in artifacts/2026-08-24-mtp-device-chain.

The production-path stochastic gate now reports phase totals from the same device-resident draft chain used by the server. Across three identical five-round runs, seven-token drafting costs 12.46--12.69 ms per round and exact CPU target-distribution/acceptance work costs only 1.87--1.88 ms per round. Fixed-eight target verification instead costs 159.21--159.61 ms per round, of which 99.5% is measured GPU time; rejection/bonus reconciliation costs another 40.49--40.92 ms per round. A new 542-stage counter trace attributes 94.40% of sampled verifier GPU time to Q4 projections: 38.38 ms QKV, 21.10 ms attention output, 50.59 ms FFN gate/up, and 48.81 ms FFN down. Attention itself is 1.34 ms and tied logits are 4.90 ms. The profiled and ordinary verifier paths produce bit-identical 2,097,152-float output with SHA-256 72c9300b97c46705f24d4880cd1de7fde150dbe6183bc463aeac9c99d74d42d1.

The adjacent M4 Max Metal MUL_MAT oracle passed 1,154/1,154 cases before a bounded fixed-eight 4-to-8-SIMD-group experiment. The candidate reaches 17,028.65 rows/s against 18,659.37 rows/s across stable surrounding four-group controls, an 8.74% regression, so it was removed. More thread-level partitioning is therefore not the next lever; the four-group kernel remains production. Evidence is in artifacts/2026-08-24-verifier-stage-profile.

Two follow-up fixed-eight Q4 schedules were also rejected. A 128-output x 8-token tile retained four SIMD groups and halved threadgroup count, but doubled scratch from 4.5 to 8.5 KiB and doubled accumulators per group. It passes the native CPU differential and the immediately preceding 1,154/1,154 reference MUL_MAT gate, but regresses stable row throughput 9.84%, so it was removed. A second candidate made one thread load each 18-byte Q4 block and expand both nibbles instead of having two nibble threads reread the block. Its repeated-65-MiB operator benchmark improves about 11.3%, but the authentic 31B verifier improves only 0.30% across five candidate and ten surrounding baseline passes, inside variance and far below the 3% gate. It too was removed. This discrepancy shows that repeated-matrix microbenchmarks can overstate cache-sensitive load improvements for the full model's unique layer-weight stream. Evidence is in artifacts/2026-08-24-q4-verify8-128x8 and artifacts/2026-08-24-q4-verify8-singleload.

The prefill operator harness now accepts --matrix-slots and rotates each dispatch across distinct GGUF-layout matrices. It reports per-matrix and total working-set bytes, and its bandwidth metric now counts one matrix read per fixed-eight dispatch rather than multiplying traffic by eight output rows. A 31B FFN-shape characterization compared a repeated 65 MiB matrix against a four-slot 260 MiB stream. The repeated controls settled near 18.85k rows/s (about 153 GB/s); the rotating run initially matched that rate, then fell to 13.45k/12.91k rows/s, yielding excessive variance. This confirms that the mode exposes resident-working-set effects, but it does not support the narrower claim that cache reuse caused the single-load candidate's failure: stable portions of both modes were effectively equal. The remaining fidelity gap is the verifier's mixed projection dimensions and graph schedule. Raw characterization is in artifacts/2026-08-24-q4-streaming-harness.

A four-group K-tile experiment then staged four contiguous Q4_0 blocks (K=128) per barrier interval instead of one. Each thread expanded two complete blocks, eliminating duplicate nibble loads and reducing synchronization by four, while preserving the 64-output/eight-token geometry and GGUF bytes. The native CPU differential covered a 67-output tail, and the adjacent Metal MUL_MAT oracle passed 1,154/1,154 cases. The candidate nevertheless fell to 14,296 rows/s versus 18,983 rows/s across surrounding controls (-24.69%). Its 18 KiB threadgroup allocation, four times the production kernel's 4.5 KiB, costs more occupancy than the saved loads and barriers repay on M4 Max. It was removed before the authentic verifier gate; evidence is in artifacts/2026-08-24-q4-verify8-k128.

PaperBridge research after this fourth rejected schedule found two useful boundaries. Apple's 2026 Metal Performance Primitives guide recommends K=128 only as an M5 starting point and explicitly warns that barrier frequency trades instruction availability against cache working set; it is not evidence for M4. RIGEL's M4 Max characterization reports that Metal 4.1 tensor operations remain on the shader/simdgroup-matrix path and that FP8 is emulated, so a low-precision tensor-format detour is not a viable M4 optimization. BaseRT independently reports that its native runtime relies on dedicated GEMV/GEMM kernels and fusion, but its engine and kernels are proprietary and its .base weight format is outside Gem's GGUF ownership boundary. The actionable open-source direction is therefore a bounded register-resident Q4 fixed-eight kernel, independently adapting llama.cpp's MIT-licensed register dequantization pattern while keeping Gem's graph and storage unchanged.

That register-resident screen passed an eight-row CPU differential including a 67-output tail and the immediately preceding 1,154/1,154 oracle. One 64-thread group read each output row's Q4 block once and kept eight token sums in registers, with no threadgroup scratch or barriers. It reached only 10,886.60 rows/s against 18,828.53 rows/s surrounding MMA controls (-42.18%) and was removed. The simdgroup matrix path is therefore essential at verifier batch eight: relying on cache for the 64-fold repeated activation loads and replacing matrix instructions with scalar FMAs is much worse. Evidence is in artifacts/2026-08-24-q4-verify8-register.

The production-width specialization was also screened and removed. All accepted Gemma 4 projection outputs are multiples of 64, so the candidate deleted the per-block output-tail branch and the scratch-mediated tail store from the four-group MMA kernel while rejecting unaligned widths at the C boundary. The immediately preceding reference MUL_MAT oracle passed 1,154/1,154 cases and the aligned native CPU differential passed. The candidate reached 19,068.84 rows/s versus 18,805.56 rows/s across surrounding controls, only +1.40%; its own five-run standard deviation was 364.43 rows/s. This is below the 3% gate and inside variance, so the generic tail-safe production kernel was restored. Further schedule changes should be driven by compiler instruction and occupancy evidence. Raw evidence is in artifacts/2026-08-24-q4-verify8-aligned.

The next bounded loader change succeeded. The production four-group verifier now reads the sixteen Q4 payload bytes as eight aligned ushort words per nibble thread, extracts two quantized values from each word, and retains the exact half(quant - 8) * scale operation. The packed-word unpack is independently adapted from llama.cpp's MIT-licensed dequantize_q4_0 helper; Gem's fixed-eight schedule, scratch layout, barriers, and matrix accumulation are unchanged. The immediately preceding M4 Max oracle passed 1,154/1,154 MUL_MAT cases and the varied-nibble CPU differential passes.

The rotating 260 MiB operator screen improves from 18,882.25 rows/s across surrounding controls to 20,413.84 rows/s (+8.11%, 165.93 GB/s). More importantly, authentic 31B verifier GPU time falls from 166.45 ms across surrounding controls to 158.92 ms (-4.52%, 1.047x) over five candidate and ten control repetitions. Candidate CV is 0.93% and the larger control CV is 1.48%, so the result clears both the 3% threshold and variance floor. All 2,097,152 verifier logits are bit-identical.

An earlier packed-word formulation used algebraically equivalent d1/d2/midpoint half arithmetic and reached 22,447.49 rows/s, but changed half rounding enough to amplify across 60 layers (NMSE 0.1138 and some per-row argmax changes). It was rejected. This is a useful numerical boundary: preserve the original half operation sequence even when refactoring the unpack. Evidence for both variants is in artifacts/2026-08-24-q4-verify8-u16.

The same exact packed-word unpack does not transfer to batch-one decode. A separate four-row SIMD candidate preserved the production float accumulation order but loaded each Q4 payload as eight aligned ushort words instead of sixteen bytes. The varied-nibble CPU differential and immediately preceding 1,154/1,154 reference oracle passed. At the 31B FFN shape with persistent buffers and command span seven, it reached 369.32 GB/s versus 367.92 GB/s across surrounding controls (+0.38%) and was 1.73% slower than the stable post-control. It was removed before the authentic model gate. Unlike fixed-eight staging, batch-one already distributes contiguous blocks across SIMD lanes, so paired loads only exchange load instructions for extra shifts. Evidence is in artifacts/2026-08-24-q4-decode-u16.

The transfer to 32-row prefill does succeed. The production 64x32 MMA kernel now stages each Q4 payload through eight aligned ushort loads while preserving its scratch layout, barriers, simdgroup-matrix operations, and exact half(quant - 8) * scale sequence. The varied-nibble CPU differential includes a 33-row tail, and the immediately preceding reference oracle passed all 1,154 MUL_MAT cases.

On the rotating 260 MiB operator screen, the candidate reaches 40,261.50 rows/s versus 38,165.60 rows/s across surrounding controls (+5.49%). Authentic 31B prompt processing rises from 110.89 to 114.40 tokens/s (+3.16%) over three repetitions, while generation changes from 20.95 to 21.01 tokens/s (+0.29%). All 262,144 baseline/candidate logits are bit-identical. The separate diagnostic kernel and API were removed after promotion, leaving one production implementation. Evidence is in artifacts/2026-08-24-q4-prefill-u16.

The refreshed decode profile motivated a lossless FFN gate/up launch-fusion screen. A diagnostic kernel read the two original mapped Q4 matrices in one launch, kept independent accumulators, and preserved every output's float operation order. The native paired CPU differential passed, the immediately preceding MTL0 MUL_MAT oracle passed 1,154/1,154 cases, and authentic batch-one logits were bit-identical.

The isolated two-matrix screen improved from 384.83 to 408.96 GB/s (+6.27%), but authentic 31B generation regressed from 19.61 to 18.56 tokens/s across surrounding controls (-5.33%). Combining the work doubles per-thread accumulators and halves independent threadgroups, so launch savings do not repay the graph-level occupancy/scheduling loss. The candidate kernel, API, test route, and benchmark selector were removed. This is another case where an isolated bandwidth result cannot override the authentic model gate. Evidence is in artifacts/2026-08-24-q4-ffn-pair.

The resulting research block shifts the next architectural axis to tree-shaped MTP verification. Apple's family-9 profiling guidance directly supports the paired-kernel diagnosis: register and thread-local storage pressure can reduce occupancy. BaseRT describes effective decode fusion as removing both a launch and a global-memory round trip, which the rejected gate/up pairing did not do. Staged Speculative Decoding instead reduces target-model bandwidth by arranging alternative drafts as a tree within one verified batch.

Gem's MTP gate now records the official assistant's top-40 rank for the target token at each first greedy mismatch. Across twelve hax-like rounds each, linear K=8 accepts 1.92 proposals for C code, 5.50 for terminal-agent output, and 2.08 for prose. For code, the target correction is in the assistant top four on 9/11 mismatch rounds; for prose it is top four on 7/12. Terminal-agent output already has nine fully accepted rounds. This supplies direct model evidence for bounded branching: alternatives can help the low-acceptance cases without replacing the excellent linear fast path. Evidence and cached paper text are in artifacts/2026-08-24-mtp-acceptance-profile and artifacts/2026-08-24-current-stage-profile.

The diagnostic tree verifier now crosses the complete 60-layer 31B graph while remaining unreachable from production. A dedicated RoPE operator assigns each node base_position + depth; the fixed-eight F32 attention kernel then exposes only the windowed committed suffix plus that node's root-to-self ancestry. The explicit-position RoPE operator matches its CPU oracle with non-monotonic offsets, and a chain-shaped tree is bit-identical to the existing batched verifier across all 2,097,152 logits. For the branched topology, all eight full-vocabulary distributions agree with independent sequential ancestry runs: every top-1 token matches, cosine similarity is at least 0.999999882, and maximum absolute error is at most 0.01973. The small nonzero branched differences are the established batch-eight MMA versus batch-one matvec numerical path, not tree masking or position drift. Adjacent Metal RoPE and FLASH_ATTN_EXT backend-op correctness pass, as do normal and ASan suites at 26/26. The next gate is an adaptive server policy and one-axis target-throughput benchmark; production remains linear until that gate passes. Evidence is in artifacts/2026-08-24-mtp-tree-full-logits.

The diagnostic commit operation also relocates a selected non-contiguous branch ancestry into logical KV positions. After committing node 7 and evaluating a follow-up token, the complete 31B distribution matches a fresh sequential path with the same top-1, cosine 0.99999999646, and maximum absolute error 0.00343. This proves cache relocation independently of scheduler quality.

An acceptance-window MTP-to-direct scheduler was tested and rejected in two forms. A four-round warmup improved the low-acceptance prose replay by 16.45% but regressed code by 14.29% and agent output by 6.50%. Extending warmup to eight rounds preserved the agent's MTP route but still regressed code by 15.22% and erased the prose gain. Accepted-prefix history is nonstationary: direct rounds advance only one token and can miss a later high-acceptance region. Both policy implementations were removed. Future scheduling must optimize effective tokens advanced per target weight pass; tree verification is preferable to direct fallback because even a mistaken branch decision retains a multi-token main path. Raw interleaved three-repetition evidence is in artifacts/2026-08-24-mtp-adaptive-direct.

A fixed-eight branched scheduler was then exercised end to end using the now-verified branch-KV commit operation. The candidate retained a five-node main path and spread three rank-2 siblings over its first three edges, selecting it after a low-confidence assistant round. It rescued branches on prose, but the three-repetition 64-token gate regressed code by 6.41% and agent output by 7.54%; a 3.52% prose gain cannot offset those failures. More importantly, hybrid greedy generation was not deterministic: the first observed divergence at generated token 21 selected token 1437 in one run and the linear oracle's token 141562 in two runs. Switching batched numerical graph topology can flip a near-tied argmax even though the isolated node-logit oracle preserves top-1 on its test vector. The scheduler and harness were removed; tree attention, explicit RoPE, and branch commit remain diagnostic-only. Evidence is in artifacts/2026-08-24-mtp-tree-gate.

The next fixed-graph verifier experiment tested whether an aligned native half load could replace the explicit two-byte Q4_0 scale reconstruction without changing arithmetic. It was bit-identical across all 2,097,152 verifier logits, but the isolated operator improved only 2.85%, below the 3% promotion floor. Separate-process model timings showed severe thermal drift, so a temporary same-process A/B/A harness alternated baseline, candidate, and baseline around every sample. Across seven 31B repetitions, the direct load regressed verifier GPU latency by 4.10% against the mean of its surrounding controls. The candidate pipeline and harness were removed. Evidence is in artifacts/2026-08-24-q4-verify8-halfscale.

A new 56-stage timestamp trace now profiles the actual promoted one-command, device-sampled seven-token MTP chain rather than the older synchronized single-token loop. Across three runs, sampled GPU work averages 7.686 ms: the four assistant layers consume 52.34%, tied output plus post-projection 29.01%, top-k40 15.19%, and embedding plus input projection 3.47%. Encoder-boundary profiling intentionally perturbs total command time, so only the summed stage samples are used for attribution. A single-axis 512-thread top-k block preserved all seven greedy proposals and passed the 445-case reference oracle, but increased top-k time from 1.167 to 1.336 ms (+14.43%) by doubling the initial run count and adding a merge level. The 1024-thread production block was restored; the stage profiler remains available through gem-mtp-bench --device-chain. Evidence is in artifacts/2026-08-24-mtp-chain-stage-profile.

The trace now has 70 stages and splits each draft's output group into RMS normalization, tied vocabulary projection, and recurrent post-projection. In the clean three-run baseline the tied 1024-by-262144 Q4_0 projection averages 2.203 ms for seven drafts. It accounts for more than 91% of the output group; recurrent post-projection is only about 0.16 ms and normalization about 0.04 ms. An exact-shape candidate doubled the matvec row fanout from four to eight. It preserved the native CPU differential, all seven proposals, and all 1,154 adjacent Metal MUL_MAT oracle cases, but regressed stable isolated throughput 2.00%, the authentic tied-output stage 5.97%, and total sampled MTP work 3.69%. The candidate was removed. Higher row fanout likely spends too many registers or reduces occupancy, so the next tied-output experiment should retain four-row scheduling and specialize the fixed K=1024 control flow instead. Evidence is in artifacts/2026-08-24-mtp-output-row8.

That fixed-K experiment is also complete and rejected. Retaining four-row scheduling while replacing the generic block loop with one block per SIMD lane and removing output-tail tests reaches 446.51 GB/s across the four stable operator samples versus 448.60 GB/s for the control (-0.47%). The two thermally stable authentic runs reduce tied-output time only 1.41%, below the 3% promotion floor; all seven greedy proposals remain exact. The compiler already eliminates most fixed-geometry control cost. The specialization was removed. Further work on this projection must change its memory access or representation rather than merely constant-folding its control flow. Evidence is in artifacts/2026-08-24-mtp-output-k1024.

A same-size SoA layout then transposed each four-row group's Q4 scales and sixteen payload bytes so every load instruction addressed contiguous lanes. Separate-process measurements disagreed by run order, so a temporary same-process A/B/A harness alternated SoA between two row-major controls. After the first cold control, seven-repetition evidence gives 432.63 GB/s for SoA versus 456.77 GB/s for the surrounding controls (-5.28%). Packing also costs 27.36 ms at startup and duplicates 150,994,944 bytes. The ordinary 18-byte stream already exploits cache-line reuse across its sequential byte loads; transposition destroys that locality. The kernel, layout, API, and harness were removed. Evidence is in artifacts/2026-08-24-mtp-output-soa4.

Keeping the exact row-major bytes but copying the 144 MiB matrix into a GPU-private Metal resource was also screened. The private copy is bit-identical and takes 8.93 ms. Initial wall-only A/B/A samples showed scheduling phases large enough to obscure the result, so the harness added Metal command-buffer GPU timestamps. Across seven paired repetitions, private storage improves device throughput only 0.36% and wall throughput 0.69%. This misses the promotion floor while duplicating 150,994,944 bytes. The private-copy API and harness were removed; shared mmap/no-copy remains production. Evidence is in artifacts/2026-08-24-mtp-output-private.

Projection/top-k fusion was then tested as a stronger mechanism. The diagnostic fixed-shape kernel launched 256 threadgroups of 1,024 threads; its 32 SIMD groups retained the proven four-row Q4 arithmetic, produced 1,024 logits in 8 KiB of threadgroup storage, and emitted the same 40 (value, token) pairs as the existing first sort. All 10,240 finite and 10,240 NaN/tie comparison pairs matched exactly. However, stable same-process GPU timestamps measured 0.431587 ms fused versus 0.406014 ms for the surrounding projection-plus-sort controls, a 6.30% latency regression. The 1,024-thread grouping sacrifices more projection occupancy than the removed logits traffic and dispatch recover, so the fused kernel and diagnostic API were removed. Evidence is in artifacts/2026-08-24-mtp-fused-projection-topk.

That bounded top-k axis is promoted. The initial sort now writes 8-byte (value, token) pairs into the 256-by-40 frontier, and all eight merge levels plus final sampling consume the carried values instead of randomly gathering full logits. The diagnostic matched 40,960 ranked ID/value checks and 162 sampled-output checks across finite and NaN/tie-heavy inputs. Seven-repetition same-process A/B/A GPU time falls from 17.442905 to 14.208964 ms per 100 selections (-18.54%). On the authentic 31B seven-draft chain, three repetitions reduce top-k from 1.171972 to 0.991388 ms (-15.41%) with the same proposals; wall time changes from 31.588 to 31.752 ms (+0.52%), inside the 2% guardrail. The two scratch buffers grow by only 80 KiB total. Evidence is in artifacts/2026-08-24-mtp-topk-pairs.

Exact frontier truncation is promoted. Each pairwise node now emits only its best 40 entries, so the eight levels write 10,200 pairs instead of 81,920 and the second scratch buffer shrinks from 80 to 40 KiB. Final frontiers and sampled outputs match the full merge for finite and NaN/tie-heavy inputs at every power-of-two tree depth from 2 through 256 runs. Stable isolated A/B/A GPU latency improves 14.21%. A separate-process authentic screen had contradictory wall scheduling, so a stronger same-model seven-repetition A/B/A gate alternated full/compact/full chains inside one loaded 31B engine: top-k improves 13.79%, total GPU time 2.34%, and wall time 2.32%, with byte-identical sparse distributions and proposals. Evidence is in artifacts/2026-08-24-mtp-topk-compact.

The exact four-way merge is promoted. Each 80-thread group merges child runs 0/1 and 2/3 into two 40-pair threadgroup lists, synchronizes once, then emits their exact best 40. Twenty-one complete and missing-child shapes produce identical finite/NaN/tie frontiers and sampled outputs. Four levels now replace eight, global merge writes fall from 10,200 to 3,400 pairs, and scratch B shrinks from 40 to 20 KiB. Stable isolated A/B/A GPU latency improves 3.50%. In a seven-repetition same-model 31B gate, top-k improves 3.27%, total GPU time 0.54%, and wall time 0.48%, with identical sparse distributions and proposals. Evidence is in artifacts/2026-08-24-mtp-topk-fourway.

Terminal merge/sampling fusion is complete and rejected. It retained the final 40 pairs in threadgroup memory, removed one dispatch and the last global frontier write, and matched all 3,402 finite/NaN/tie sampled-output checks. Stable isolated A/B/A latency improved only 1.72% on-device and 1.82% by wall clock, below the 3% promotion floor, so no authentic-model gate was warranted. The fused kernel and diagnostic API were removed; evidence is in artifacts/2026-08-24-mtp-topk-terminal-fusion.

Top-k is now a bounded ~0.12 ms per-draft stage with no remaining dispatch-only proposal above the promotion floor. Return to the updated 31B MTP stage trace and select the largest fixed-graph cost that offers a stronger fusion or data-reuse mechanism.

The refreshed stable 31B trace assigns about 4.04 ms to the four assistant layers, 3.25 ms to output norm/tied output/post-projection, and 0.83 ms to top-k across seven drafts; one of three independent profiles suffered broad GPU interference and is retained as an outlier. The first layer-boundary experiment fused post-attention RMSNorm, residual add, and FFN RMSNorm. A 4 KiB threadgroup staging slice preserved the exact materialization boundary: 131,072 residual/normalized floats matched bit-for-bit, and the isolated three-to-one dispatch benchmark improved 39.1%. Yet the authoritative same-model 31B A/B/A gate improved assistant-layer time only 1.30%, total GPU time 0.69%, and wall time 0.69%. The tiny elementwise dispatches are hidden by surrounding graph scheduling. The kernel, API, and harness were removed; evidence is in artifacts/2026-08-24-mtp-post-attention-fusion and the refreshed profile summary is in artifacts/2026-08-24-mtp-stage-refresh-after-topk.

The assistant-specific gate/up-plus-GELU experiment is complete and rejected. One exact-shape kernel streamed the two canonical 1024-by-8192 Q4 matrices together, retained the production accumulation order, and wrote only GELU(gate) * up. It removed 56 dispatches per seven-draft chain and roughly 96 KiB of F32 intermediate traffic per layer. The native operator differential was bit-identical, all authentic proposals/candidate rows/probabilities matched, and the adjacent Metal MUL_MAT oracle passed 1,154/1,154 immediately before the gate. Nevertheless, the three stable same-model A/B/A triplets improved assistant-layer latency only 1.25%, sampled GPU and wall latency 0.95%, and command-buffer GPU latency 0.77%. The 9 MiB of Q4 weight traffic per layer dominates the eliminated activation traffic, while the combined kernel raises accumulator pressure. It was removed; evidence is in artifacts/2026-08-24-mtp-gate-up-gelu-fusion.

Dispatch-only and activation-only fusions are now exhausted for the assistant graph. Before another kernel rewrite, add fine-grained authentic assistant-layer timestamps separating Q projection, attention, output projection, gate/up, down projection, and epilogues. Select the next axis only if that trace identifies reusable weight work or a stage large enough to clear 3% end to end.

The first fine-grained trace was subsequently found to have used a hardcoded F16 KV cache in the benchmark harness despite its production-facing configuration record. Its five-token attribution is retained only as an F16 short-context diagnostic and is explicitly superseded. The harness now uses and reports the production F32 default. At a realistic 580-token context, the two stable 31B runs place 4.409 ms of 8.048 ms assistant-layer time in attention (54.79%); a third broad GPU interference run is retained as an outlier. All runs produced the same seven proposals. F16 attention mirrors were also slower at this context, by about 30% for SWA and 47% for the full layer, so cache conversion is not a shortcut. Corrected evidence is in artifacts/2026-08-24-mtp-layer-fine-profile-f32-580.

The first exact F32 attention screen changed only the 31B full-layer score geometry from eight heads over one token to four heads over two tokens per workgroup. Native attention and Metal suites passed, but median dispatch throughput improved only 0.33% (5713.68 to 5732.52/s), below the 3% isolated floor. The candidate was removed before server integration or an oracle/model gate. Evidence is in artifacts/2026-08-24-mtp-score-head-block-screen.

A second profiling pass split exact attention into cooperative-K score generation, global normalization, and shared-V accumulation. Across three stable 580-token F32 runs, those stages average 1.557, 0.136, and 2.730 ms respectively; shared-V alone is 34.41% of assistant-layer time and all proposal sequences match. This follows FlashAttention-2's work-partitioning lesson but uses target-hardware measurements to choose the axis. A GQA2 experiment then packed four KV heads into one workgroup so all eight SIMD groups accumulated outputs and workgroup count halved. It was bit-identical on wrapped 12B/31B SWA cases, but stable throughput regressed about 26% because added indexing and synchronization dominated. The candidate was removed before an oracle/model gate. Evidence is in artifacts/2026-08-24-mtp-attention-substage-profile and artifacts/2026-08-24-mtp-value-gqa2x4-screen.

A simpler GQA2 follow-up kept one KV head per workgroup and doubled only the value dimension tile from 64 to 128. It also remained bit-identical and halved workgroups, but regressed stable throughput 21.33%. The larger 17 KiB threadgroup tile and four accumulators per active lane reduced occupancy or issue efficiency more than launch reduction helped. It too was removed. Together the two screens show that the next shared-V candidate must reduce bytes or synchronization, not merely workgroup count. Evidence is in artifacts/2026-08-24-mtp-value-gqa2-d128-screen.

The synchronization-only control doubled the shared-V token tile from 32 to 64 while preserving one KV head, 64 dimensions, V traffic, and token accumulation order. It was bit-identical after the 64 weights per head were cooperatively staged. Stable wall/GPU throughput improved 0.97%/1.27%, showing that fewer barriers help, but the 18 KiB threadgroup footprint caps the gain below the 3% promotion floor. The candidate was removed. Evidence is in artifacts/2026-08-24-mtp-value-token64-screen.

Token tile 48 then tested the occupancy-balanced midpoint: 13.5 KiB of threadgroup memory and about one-third fewer barrier rounds than token 32. It was bit-identical and improved stable wall/GPU throughput 1.99%/1.78%, better than token 64 but still below promotion. It was removed. The measured 32/48/64 curve closes token-tile tuning for this value kernel; a future candidate needs a different dataflow rather than another nearby tile. Evidence is in artifacts/2026-08-24-mtp-value-token48-screen.

A structurally different GQA2 value pass mapped two normalized query-head rows by 128 output dimensions onto Metal simdgroup 8-by-8 matrix multiplies, using half inputs and float accumulation. It passed CPU-relative wrapped-ring 12B/31B checks, but stable wall/GPU throughput regressed 40.92%/45.57%. Padding two useful rows to eight and converting/staging F32 cache values to half cost far more than the matrix primitive saved. The candidate was removed before an oracle/model gate. Do not retry a padded matrix path at decode's two-row GQA shape without native low-bit KV storage or substantially more useful rows. Evidence is in artifacts/2026-08-24-mtp-value-mma-screen.

An implementation survey found an MIT-licensed Gemma-first Metal server with a custom Q3_0 format (14 bytes per 32 weights versus Q4_0's 18) and native matvec kernels. A local diagnostic independently adapted its split low/high bitplanes to Gem's four-row schedule. The native CPU/Metal differential passed, and the adjacent Metal MUL_MAT oracle passed 1,154/1,154 immediately before the confirmation screen. Nevertheless, exact assistant shapes regressed badly: 1024-by-8192 latency increased 71.76% and 8192-by-1024 increased 56.61%. Bitplane extraction and reconstruction overwhelm the 22.22% byte reduction on M4 Max. The entire diagnostic was removed before assistant conversion or acceptance testing. Do not retry this Q3_0 layout; any future reduced-weight format must first prove efficient native unpacking. Evidence is in artifacts/2026-08-24-mtp-q3-weight-screen.

The F16 KV path was revisited after direct Metal profiling showed that its remaining opportunity was prefill scheduling, not merely cache width. A direct F16 split-GQA decode diagnostic held half K/V tiles in threadgroup memory and improved the isolated 32K operator about 3.2%, but it reduced authentic 31B generation throughput 8.71%; that decode route was removed. Keeping the existing one-pass F16 decode and changing only batched prefill attention from one dispatch per row to an exact split-GQA score/normalize/value sequence improved an interleaved 1,024-token gate from 49.14 to 82.22 pp t/s (+67.31%). Generation changed from 18.99 to 18.68 tg t/s (-1.62%), below the 2% regression limit and inside the 4.04% baseline variation. The native batched CPU differential and wrapped-ring 12B/31B full/SWA tests pass, and the adjacent FLASH_ATTN_EXT oracle passed 4,792/4,792 immediately before the gate. This is promoted as an explicit --kv-type f16 server, agent, and generation option; F32 remains the default because the authentic 2K comparison preserves argmax but measures 0.998918 correlation and 1.91e-3 NMSE. Evidence is in artifacts/2026-08-24-f16-split-gqa.

Low-bit KV research does not currently justify an int4 production path. Open-TQ-Metal reports a fused Apple-Metal int4 attention design, but its own Gemma 4 31B quality study places the usable int4 KV ceiling around 950 tokens and reports compounding error beyond it. The broader SRFT/int4 Apple study evaluates much smaller models rather than Gemma 4 12B/31B long-context serving. No int4 kernel was implemented: hax project prompts routinely exceed the published quality boundary, so a future attempt first needs independent long-context Gemma 4 evidence rather than a speed-only operator result.

With batch-one decode at the Q4 memory-streaming ceiling, the next production lever moved to multi-turn hax scheduling. The server previously reset the target on every request, even though a tool result or follow-up turn resends nearly the entire prior conversation. A naive live-sequence cache correctly missed because raw generation contains channel-control tokens that OpenAI messages canonically re-render. The promoted design instead checkpoints the exact tokenized request prompt, streams the response, restores every K/V slot the bounded response could touch, and returns to the request boundary. A later request reuses state only when its token sequence strictly extends that exact prefix. Mismatches recompute from zero.

At approximately 3.1K prompt tokens, three real HTTP/SSE reset-based second turns took 70.216/81.469/70.722 seconds. Three growing cached turns took 1.024/1.050/1.096 seconds while reusing 3,101/3,123/3,145 tokens and evaluating 22 new tokens each: mean latency falls 98.57%, a 70.14x speedup. A 1,025-token wrapped-SWA differential discards eight generated positions, restores the prompt, and evaluates an alternate eight-token suffix. It preserves the same argmax with 3.29e-4 NMSE, slightly below the 3.56e-4 independent repeat-control NMSE; full-model Metal execution was not bit-repeatable under the measured system load. Checkpoint plus restore costs 2.57 ms for that tail. A real gem agent/hax tool loop then normalized Gemma's un-delimited scalar call:bash{command:pwd} to valid JSON, reused 2,533 tokens on the tool response request, returned the exact workspace path, and cleaned up the child server. Evidence is in artifacts/2026-08-24-exact-prefix-cache.

The fixed 128-token response limit is now removed. A checkpoint performs no K/V copy; every target evaluation lazily preserves the token-major slots it is about to write, and geometric storage growth continues through the remaining model context. This covers direct decode, batch-eight verification, correction, bonus, tree, and profiled evaluation paths at the inference boundary. If growth fails, the response still completes and only that next-turn cache opportunity is discarded. A wrapped 1,025-prefix/160-discard authentic check retained argmax in three normal runs. The final ASan/UBSan run measured 7.47e-5 restored NMSE versus 1.85e-4 independent-repeat NMSE. In a three-repetition HTTP/SSE gate, every preparation response reached 160 tokens and every second turn returned DONE; byte-identical 404-token full-prefill replays took 4.383/4.392/4.516 seconds, while exact hits took 2.305/2.338/2.529 seconds, reducing mean latency 46.04% (1.85x). Evidence is in artifacts/2026-08-24-lazy-prefix-rollback.

A longer linear MTP lookahead is rejected after a complete K=16 diagnostic. The generic 16-row Q4 path merely matched two fixed-eight calls because it computes a padded 32-row tile. A dedicated 64-output by 16-row Metal kernel then passed the CPU differential and the adjacent 1,154/1,154 MUL_MAT oracle, and improved equivalent projection throughput 35--38%. That operator win did not survive the authentic 31B gate. Across three deterministic 12-round runs, K=8 produced 28.44 t/s with a 4.5-token mean accepted prefix, while K=16 produced 22.44 t/s with a 5.67-token prefix: a 21.11% throughput regression. The larger verifier cost outweighed its extra accepted work. The fixed-16 kernel and all diagnostic inference APIs were removed; production remains K=8. Evidence is in artifacts/2026-08-24-q4-verify16.

A refreshed canonical 31B scorecard used context 64, eight prompt tokens, eight generated tokens, three repetitions, and native A/llama/A bracketing. Native F32 KV averaged 53.66 pp and 20.27 tg t/s with less than 0.5% bracket drift. The pinned llama.cpp oracle (commit d8df12e) averaged 31.16 pp and 23.52 tg t/s with its required F16 KV: native prefill is 72.21% faster, while llama.cpp direct decode is 16.05% faster. The already-promoted native K=8 MTP result of 28.44 t/s remains 20.92% above llama.cpp direct decode. llama-bench cannot select F32 KV, so this is competitor context, not a one-axis KV comparison. Longer 64-token diagnostics crossed into a slower sustained machine state and are excluded from the scorecard. Evidence is in artifacts/2026-08-24-current-31b-scorecard.

The direct-decode gap motivated a bounded encoder synchronization screen. Replacing only the global buffer-scope barrier before each one-row Q4 projection with a resource-specific activation barrier passed all 29 normal tests, but the machine slowed 8.86% between the preceding and following controls. Candidate decode was only 0.24% above the equally slow post-removal control, so no causal gain could be established and the change was removed. Do not promote synchronization changes without an interleaved selector that survives the 3% gate. Evidence is in artifacts/2026-08-24-q4-resource-barrier.

The fastest promoted engine path is now connected to the preferred one-command offline workflow. For gem agent with F32 KV and no explicit MTP option, a regular sibling named mtp-<target filename> is discovered and passed through the existing strict assistant manifest validation. Explicit --mtp wins, --no-mtp forces direct decode, F16 skips discovery, and gem serve remains explicit. A real 31B hax smoke returned exactly OK, advertised native MTP K=8, and left no listener behind. In the formal three-pair greedy gate, all six complete outputs were byte-identical; direct process wall time was 4.19/4.04/4.06 seconds versus 3.29/3.35/3.31 seconds with automatic MTP. Mean end-to-end latency fell 19.04% (1.235x), including model readiness, hax, generation, and cleanup. Evidence is in artifacts/2026-08-24-agent-mtp-autodiscovery.

A completion audit found that real-model cancellation was required by the correctness ladder but had no direct artifact. The server previously learned about disconnects only while writing visible model events; Gemma control/reasoning tokens could therefore leave the serial GPU stream occupied. The promoted fix sends content-free SSE keepalive comments before and after prompt evaluation and at each generation-round boundary, and hardens every socket write with MSG_NOSIGNAL. A failed heartbeat follows the existing prefix rollback/reset cleanup. In the authentic 31B MTP gate, curl cancelled after one second with status 28; the same process then returned /health 200 in 1.003 seconds. A same-binary three-pair real-hax gate measured 4.453 seconds without heartbeats and 4.347 seconds with them (-2.40%, within noise), with six byte-identical outputs. A follow-up poll check changed sequential recovery only from 1.003 to 0.956 seconds and was removed. Normal and sanitizer tests include a forced mid-stream disconnect. Evidence is in artifacts/2026-08-24-real-cancellation.

The deployed context-capacity gate now passes on the physical 128 GiB M4 Max for authentic 31B Q4_0 plus the production K=8 assistant at 32K, 64K, and 128K F32 contexts. The configured KV allocations are 7.05, 12.42, and 23.15 GB. Each server passed health and three requests using an eight-token tokenizer-confirmed prompt; all nine completions stopped normally with identical text. The two exact-prefix reuse requests at each capacity averaged 1.346, 1.350, and 1.351 seconds. First-request measurements are retained but are order-dependent Metal/residency warmup, not a context comparison. An ordered direct sweep reported 53.70/20.28, 53.67/20.02, and 53.45/18.39 pp/tg tokens per second, but bracketing runs alternated between roughly 18.3 and 20.1 tg at both 32K and 128K. That decode difference is therefore recorded as nonstationary drift, not a capacity regression. The gate also exposed that the required server-probe.sh request used stream=false while native serving rejected it. Native non-streaming completion aggregation is now implemented for text, reasoning, tool calls, finish reason, and usage without changing the hax streaming path. Normal and sanitizer suites pass 30/30 and the binary remains free of llama/ggml linkage. Evidence is in artifacts/2026-08-24-context-capacity-31b.

The remaining later-tile attention hotspot was score generation: at a 2,496-token position it was about 63% of both SWA and full-attention operator time. A fixed-31B q8k8 score candidate uses Metal 8x8 SIMD-group matrices to reuse each staged K tile across eight query rows. Full attention stores its regular 8x8 result directly; wrapped SWA uses a 16-token tile and a 1 KiB threadgroup scatter area to translate absolute ring positions back to each row's local score indices. The existing exact global normalization and shared-V passes are unchanged, and all other shapes fall back to the prior scalar score kernel.

After adding the required post-store SIMD-group visibility barrier, the exact 2,496/32 score stage falls 79.27% for full attention and 82.62% for SWA; complete three-pass operator time falls 50.28%/51.18%. The adjacent MTL0 FlashAttention suite passes 4,793/4,793 and 200/200 independent native full/wrapped-SWA differential processes pass. Authentic 2K logits retain finite values, identical argmax/top-10 order, correlation 0.999999687, and NMSE 4.50e-7. A pre-fix five-pair model gate had looked eligible but is invalid because the stress race was not yet fixed. Post-fix 256/512 gates improve mean pp 4.44%/4.27%, both inside 6.54%/20.14% noise floors while unrelated host load and thermal drift were high. Promotion is therefore withheld; production remains the scalar score path and q8k8 remains an explicit diagnostic selector pending a stable post-fix model gate. Evidence is in artifacts/2026-08-24-full-score-q8k8-31b.

A true streaming probe now separates curl's immediate header/assistant-role TTFB (roughly 0.6 ms) from the first SSE delta containing content, reasoning, or a tool call. Before scheduling changes, authentic warmed 31B K=8 visible TTFT was 523.4/519.6/502.5 ms at 32K/64K/128K, confirming that allocated capacity is not the short-prompt cause. Token tracing showed four suppressed channel pieces before the first visible token, so an early speculative-root experiment had no benefit and was removed. The promoted scheduler instead emits every already-verified token before evaluating the correction, bonus, or direct token that prepares future logits. A same-process interleaved 32K gate reduced MTP visible TTFT from 519.8 to 460.4 ms (-11.42%) and completion time 3.78%; direct TTFT fell 8.18%. All six normalized SSE pairs were identical. The diagnostic-free production gate measured 467.1 ms mean TTFT, 10.14% below baseline, with 30/30 normal and sanitizer tests and a real hax OK smoke. Evidence is in artifacts/2026-08-24-streaming-ttft-31b.

Long-prompt full attention first promoted a scalar long8 path: at positions from 4,096 onward, only the exact 31B F32 full-attention shape grouped eight query rows while 12B, SWA, F16, partial tiles, and every other shape remained rowwise. It was bit-identical to the oracle and improved a same-process 8K full-model gate from 32.90 to 37.08 pp tokens/s (12.70%).

The subsequent fused path failure was localized to selector routing, not attention math: its pre-4K fallback accidentally used serial decode attention. After routing those tiles through the rowwise oracle, the 4,128-token authentic 31B gate retains identical argmax and ordered top-10, correlation 0.999999929, NMSE 5.61e-8, and byte-identical fresh-process repeats. The original fused kernel also proved to be three complete K/V passes despite its online name and was slower than long8. The promoted replacement implements the true one-pass running-max recurrence, rescaling its denominator and 512-wide register accumulator as the maximum changes. Against long8 it improves 8K/32K/64K/128K full-attention tiles by 1.63x/1.61x/1.60x/1.59x and uses no score buffer. A same-process interleaved 8K model gate improves pp from 37.03 to 41.25 tokens/s (11.39%) against a 2.29% noise floor; generation improves 3.35%. Normal and sanitizer suites pass 32/32, adjacent MTL0 suites pass 4,793/4,793 FlashAttention and 1,154/1,154 matrix-multiply cases, native linkage is clean, and a direct-server hax smoke returns exactly OK. Full-model pp plus streaming visible TTFT at 64K/128K remain required before claiming end-to-end long-prompt completion. The promoted default completes the three-repetition authentic 32K prompt gate at 31.700 pp tokens/s with samples 31.701/31.750/31.650 (0.13% coefficient of variation) and 20.40 one-token tg tokens/s, using a 32,800-token context and 7.05 GB KV allocation. The authentic 64K gate completes at 23.884 pp tokens/s with samples 25.262/23.285/23.104 (4.09% coefficient of variation) and 20.04 one-token tg tokens/s, using a 65,568-token context and 12.42 GB KV allocation. The first-to-later drop is retained as sustained thermal/frequency drift rather than filtered as an outlier. An online 128K gate then completes at 14.705 pp tokens/s with samples 15.876/15.249/12.990 (8.43% coefficient of variation) and 16.25 one-token tg tokens/s, using a 131,104-token context. The full seven-hour ordered run therefore proves capacity and execution but also establishes a material endurance/thermal warning; the third sample is retained rather than normalized away. An online token-tile-8 experiment is rejected: its 32 KiB threadgroup allocation loses 2.6% at 8K and roughly 1.5-2.0% at 32K/64K versus the promoted tile-4 kernel, indicating occupancy costs exceed the saved barriers. Its post-endurance 128K operator samples are thermally unstable and are not evidence. Evidence is in artifacts/2026-08-24-long-context-online-attention.

The next long-context pass specialized the online recurrence by exact model geometry. For 31B/F32, one workgroup now evaluates two query rows while retaining the same causal masks, running maximum, denominator, and F32 accumulator. Seven interleaved samples show 1.58x attention speedup at 8K, 32K, and 128K (1.56x at 64K), still with zero score scratch. Repeated authentic 2K, 4K, and 4,128-token logits are bit-identical; at the routing boundary the candidate is byte-identical to the established online path across all 262,144 logits. Under a heavily throttled sustained run, the interleaved 8K model gate improves pp from 30.62 to 33.29 tokens/s (+8.72%), above its 5.80% noise floor. A final candidate decode sample fell to 9.74 tokens/s even though this selector does not execute during decode; it is retained as thermal drift and makes the secondary result non-significant. The analyzer still marks prompt-primary promotion eligible.

The 12B full-attention geometry received its own F32/F16 online kernels rather than inheriting the 31B shape. At 8K/32K/64K/128K, F32 attention is 6.05x/5.99x/5.86x/5.90x faster than its rowwise oracle; F16 is 27.00x/27.74x/26.66x/26.99x faster than its rowwise F16 oracle. The 12B/F32 candidate is deterministic at 2K, 4K, and 4,128 tokens. At the boundary it retains identical argmax and ordered top ten with correlation 0.9999997 and NMSE 6.31e-7. Its interleaved 8K model gate improves pp from 51.19 to 62.39 tokens/s (+21.88%) against a 2.45% noise floor; generation changes -1.90% inside 5.26% observed variation. F32 12B and F32 31B specialized kernels are now the production defaults. F16 and f16-full remain separately gated deployments: their operator results establish scalability and memory viability, but not yet model-level promotion. Final real gem agent/hax smokes supplied an 8K-plus raw prompt so both post-4K selectors actually executed; 31B and 12B each returned exactly OK, and the launcher released both server ports. A three-process 31B/F32 cold-server gate then confirmed exact 8,192-token usage in every response. First-visible SSE content arrived at 227.592/278.878/238.427 seconds (248.299-second mean, 8.89% CV), while completion averaged 249.788 seconds. The 22.5% first-to-second slowdown is retained as sustained thermal/frequency drift. The immediate ~1 ms HTTP header event is again excluded from model-visible TTFT. Fresh-process and prefix-extension gates at 32K/64K/128K remain separate endurance measurements; they are not inferred from this 8K result or from operator throughput.

The mixed f16-full capacity mode then received representation-specific promotion gates. It keeps SWA in F32 and stores only the context-growing full layers in F16, reducing 31B/12B 128K KV to 12.42/2.82 GB. For both sizes, candidate outputs are bit-identical across repeated 2K, 4K, and 4,128-token evaluations. At the first post-4K tile, online versus rowwise preserves argmax and ordered top ten; 31B correlation/NMSE is 0.999999922/7.67e-8 and 12B is 0.999999853/3.25e-7. Same-process 8K model gates show why shape-specific F16 online attention is mandatory: 31B pp rises from 17.25 to 29.19 tokens/s (+69.25%, 5.76% noise), and 12B rises from 29.14 to 57.28 (+96.59%, 4.57% noise). Decode means fall 7.41%/6.71%, but both changes remain inside 9.09%/11.20% observed variation and the selector does not execute during decode. The existing 31B mixed route is validated and the specialized 12B route is promoted. f16-full remains opt-in on the 128 GiB target: it is the capacity-efficient path, while F32 row2 remains the fastest absolute 31B attention kernel. Evidence is in artifacts/2026-08-25-long-context-f16-full. Final 8K-plus real hax smokes returned exactly OK for both mixed-KV model sizes and the launcher released both ports, confirming the selected kernels through the production HTTP/SSE path.

The final 31B/F32 row2 server then passed a three-process 32K cold plus exact-prefix gate. Cold first-visible TTFT was 1,163.851/1,192.594/1,185.567 seconds (1,180.671-second mean, 1.04% CV), with 1,183.040-second mean completion and exact 32,768-token usage. In the same process, a strict assistant/user extension reused all 32,768 prompt tokens and evaluated only 26 appended tokens. Its visible TTFT was 1.817/2.377/1.997 seconds (2.064-second mean), a 572x reduction from cold ingestion, and completion averaged 2.239 seconds. This validates the intended offline hax usage model: initial repository ingestion is expensive, but subsequent turns are suffix-only when the canonical conversation strictly extends. Evidence is in artifacts/2026-08-25-server-prefix-31b-row2-32768.

The equivalent final 31B/F32 row2 64K gate also passes. Three fresh cold requests report exact 65,536-token usage and first-visible TTFT of 2,809.768/2,837.049/2,830.016 seconds (2,825.611-second or 47.09-minute mean, 0.41% CV); completion averages 2,828.940 seconds. Each strict second turn reuses all 65,536 tokens and evaluates only 26 appended tokens, reaching visible output in 3.267/2.512/3.008 seconds (2.929-second mean) and completing in 3.191 seconds. Prefix reuse therefore reduces visible latency by 965x at this rung. Evidence is in artifacts/2026-08-25-server-prefix-31b-row2-65536.

Decisions

• GGUF is the canonical container, not the execution engine.
• Text-only Q4_0 is the first end-to-end target for both sizes because official artifacts exist and the block layout is small enough for a dedicated Metal path.
• 31B is the primary performance target; 12B shortens correctness/profiling iteration and remains a fully supported deployment profile.
• Multimodal artifacts remain a later phase. The optional matching MTP sidecar is now a native, promoted server acceleration path; direct decoding remains the fallback and correctness oracle.
• llama.cpp output provides differential evidence until the native engine has independent golden vectors. Matching its throughput is a benchmark goal, not an architectural dependency.

Open questions

1. Is M4 Max decode limited by Q4_0 weight bandwidth or command/dispatch overhead at batch one?
2. What command-buffer span minimizes launch overhead without harming cancellation latency?
3. For hax prompt shapes, when does a dedicated batched Q4_0 kernel overtake matrix-vector tiling?
4. Can the alternating 256/512 head dimensions share one Metal pipeline through function constants without losing occupancy?
5. Does a derived packed layout earn enough throughput to justify conversion and duplicate storage?

Primary sources

• Google Gemma 4 12B configuration
• Google Gemma 4 31B configuration
• ggml-org Gemma 4 12B GGUF
• ggml-org Gemma 4 31B GGUF
• Gemma 4 technical report
• Google Gemma 4 canonical chat template
• Google Gemma 4 prompt and tool-control format
• llama.cpp Gemma 4 graph
• FlashAttention: exact IO-aware tiled attention
• FlashAttention-2: improved work partitioning
• Online normalizer calculation for softmax
• Accelerating Large Language Model Decoding with Speculative Sampling
• Draft & Verify: Lossless Large Language Model Acceleration via Self-Speculative Decoding
• BaseRT: native Metal LLM inference
• Profiling LLM inference on Apple Silicon across quantization formats
• metal-llm-server Q3_0 reference implementation
• RIGEL: M4 Max Metal tensor-path characterization
• Apple Metal Performance Primitives Programming Guide
• Apple: Optimize Metal apps and games with GPU counters
• Apple: Discover Metal debugging, profiling, and asset creation tools
• llama.cpp Metal Q4_0 register dequantization
• Open-TQ-Metal: fused int4 attention on Apple Silicon
• When Quantization Is Free: Apple Metal int4 and SRFT study

2026-08-25: four-row global-attention sharing is an operator win, not a production win

A diagnostic 31B F32 global-attention kernel assigns one SIMD group to every (query row, query head) pair in a 1,024-thread group, so four causal rows share each staged K/V tile without quadrupling per-lane query and accumulator registers. The native differential reports the same 1.54594846192e-12 NMSE as the established row1/row2 online kernels and fresh repeats are bit-identical. The adjacent MTL0 FLASH_ATTN_EXT oracle passed 4,793/4,793 cases. Against row2, three-repetition operator speedups are 1.1574x at 8K, 1.1691x at 32K, 1.1622x at 64K, and 1.1549x at 128K.

The required full-model 8K interleaved gate did not reproduce that stability. Row2 averaged 30.0493 prompt tok/s and row4 averaged 31.4023 tok/s (+4.50%), but paired row4 outcomes were +16.28%, -0.43%, and -1.87%, with 5.89% candidate CV. Row4 therefore remains diagnostic and row2 remains the production default.

Primary-source review suggests the next axis should be sequence-parallel online attention, not a larger row tile. FlashAttention-2 parallelizes a single head over multiple thread blocks to improve occupancy. Gem's row4 geometry launches only 32 long-lived groups for a 32-row block; a two-segment split-K form would launch 64, retain total K/V reads, write about 4 MiB of partial F32 online-softmax state, then merge (maximum, denominator, accumulator) in a second kernel. Apple Family-9 guidance also makes two smaller independent screens worthwhile: direct device-buffer reads instead of threadgroup cache copies, and an explicit max_total_threads_per_threadgroup compiler constraint. Evidence is in artifacts/2026-08-25-long-context-row4.

2026-08-25: row4 follow-up rejects sequence splitting and retains row2

The proposed occupancy screens were implemented as isolated one-axis variants. A two-segment sequence-parallel row4 kernel increased the number of workgroups from 32 to 64 and merged roughly 4 MiB of (maximum, denominator, accumulator) state. It was slower than unsplit row4 by 0.69%/0.99%/0.13%/0.34% at 8K/32K/64K/128K. Reading K/V directly from device memory instead of cooperative threadgroup staging regressed 8K by 390.71%, and a 512-thread row4 geometry regressed 8K by 9.32%. An explicit 1,024-thread pipeline constraint produced only an approximately 1% operator improvement. These results reject the tested FlashAttention-2-style sequence split and Apple Family-9 occupancy/cache variants for this exact 31B geometry; the temporary kernels and APIs were removed.

The earlier row4 full-model result was resolved with a purpose-built long-prefix harness. It builds independent row2 and row4 31B inference states over the same mapped model, performs an unmeasured post-prefix warmup for each, restores exact prefix checkpoints, alternates three 32-token extensions, includes the final logits projection, and compares output logits. At a 32,768-token prefix, row2 took 2.370925/2.288364/2.186804 seconds (2.282031 mean), while row4 took 2.478137/2.320768/2.426694 seconds (2.408533 mean). Every pair favored row2; row4 regressed wall latency 5.54%. Logits were identical (correlation 1, NMSE 0, argmax 26745). A preliminary run that appeared to favor row4 is retained but marked invalid because it omitted warmup and final logits; its first row2 sample alone paid a 3.805-second cold-path cost.

Row2 therefore remains the production 31B F32 global-attention kernel. Base row4 remains only as a diagnostic comparator for the long-prefix harness. Attention row-sharing, direct-read, threadgroup size, and two-way sequence-splitting axes are closed for this dataflow; the next long-prompt optimization must begin with authentic stage-level attribution rather than another attention-only microbenchmark. Evidence is in artifacts/2026-08-25-long-context-row4-followup.

2026-08-25: authentic stage attribution finds SWA dominance, but exactness blocks reuse

A new long-prefill profiler uses Metal stage-boundary timestamps over one production 32-token tile, while preserving the authentic prefix position and separating the 50 sliding layers from the ten full-attention layers. At a 4,096-token 31B/F32 prefix, three production row2 samples average 0.907931 GPU seconds and account for 99.65% of GPU time through stage samples. KV attention is 0.541941 seconds (59.90% of sampled time). Surprisingly, bounded 1,024-token SWA contributes 0.482717 seconds (53.35%), while growing global attention contributes only 0.059224 seconds (6.55%). QKV projections, FFN gate/up, and FFN down/output contribute 8.38%, 13.67%, and 11.63%.

The cause is explicit routing: the long-context selector keeps the production global row2 kernel but evaluates each SWA row independently. Restoring the established 32-row batched split-GQA path for all SWA layers improved a corrected 4K prefix-extension gate 19.14%, but failed model-level correctness (correlation 0.998577, NMSE 3.65e-3). A new 64-thread online SWA row-pair kernel then passed its wrapped-ring operator differential at 7.09e-13 NMSE and the adjacent MTL0 FLASH_ATTN_EXT suite at 4,793/4,793. It improved the extension gate 37.62%, but recurrent use across 50 layers amplified its tiny local error to correlation 0.984060 and NMSE 3.29e-2.

Partial batched routing did not provide a safe compromise. Two-of-five SWA layers improved 9.28% but produced NMSE 2.30e-3; one-of-five improved 11.87% but produced NMSE 4.15e-3. The non-monotonic accuracy curve makes further subset selection prompt-specific overfitting. All temporary selectors, the online SWA kernel, its pipeline/API, and test extensions were removed. Production remains rowwise SWA plus row2 global attention. Future SWA work must preserve the established rowwise/global softmax arithmetic order while reducing dispatch or K/V traffic; isolated operator NMSE is not a sufficient proxy for repeated 50-layer stability. Evidence is in artifacts/2026-08-25-long-prefill-stage-attribution.

2026-08-25: barrier repair does not resolve batched-SWA model drift

Primary-source review reinforced two constraints for the next SWA design. Apple documents that threadgroup barriers are required for shared-memory visibility even when a threadgroup is only one SIMD group, and that a fixed launch geometry preserves SIMD grouping across the dispatch. The FlashAttention and FlashAttention-2 papers support exact IO-aware tiling and adding query/sequence parallelism, but they do not promise bitwise equivalence after reassociating softmax arithmetic.

Inspection found that attention_prefill_normalize_f32 lacked the barrier present in the rowwise normalizer between reading the reduced maximum and reusing shared memory for partial sums. The barrier was restored. The native authentic 31B SWA operator then reported identical score hashes and zero output NMSE for scalar-batched versus rowwise/chunked paths, and the adjacent MTL0 FLASH_ATTN_EXT suite passed 4,793/4,793 cases.

The real-weight 4,096-prefix, 32-token interleaved extension gate nevertheless reproduced the previous rejection exactly. Batched SWA averaged 1.077240 seconds versus 1.274846 seconds rowwise (15.50% lower wall time), but final logits had correlation 0.998576705593 and NMSE 0.00364940360302. The temporary selector was removed and production routing remains unchanged. The barrier remains because it is required synchronization, but the evidence shows that an isolated operator comparison does not reproduce the authentic multi-layer command-buffer failure. The next diagnostic must locate the first differing layer/tile and compare score, normalized probability, attention output, and post-commit KV state at that boundary. Evidence is in artifacts/2026-08-25-batched-swa-barrier.

Sources:

• Apple: Creating threads and threadgroups
• Apple: Porting Metal code to Apple silicon

2026-08-25: exact-order two-row K reuse further accelerates wrapped SWA

The promoted wrapped-SWA score kernel shares each staged K pair across two adjacent causal query rows while retaining the production lane accumulation loop and simd_sum reduction for every dot product. Normalization and P*V remain unchanged. It applies only after the 1,024-token ring wraps; tile zero stays rowwise and pre-wrap nonempty tiles retain scalar batching.

The native wrapped 31B scores and outputs are byte-identical to the scalar route. A 7x30 interleaved operator gate reduces score time from 2.492 to 1.522 ms (-38.91%) and complete SWA attention from 3.969 to 3.004 ms (-24.30%); the adjacent Metal suite passes 4,793/4,793. Authentic model gates improve 32-token extension latency 7.02% at 4K and 7.14% at 32K. All three 32K pairs favor the candidate, and repeated plus cross-route logits are bit-identical (correlation 1, NMSE 0, argmax 26745).

A three-repetition 4K full-graph profile attributes the gain to the intended stage. Sliding-layer KV-attention time falls from 190.399 to 154.146 ms (-19.04%), total KV-attention time falls 14.39%, and sampled GPU time falls 1.80%. Only the wrapped portion of a 4K prompt uses this kernel, so the 32K model A/B remains the promotion timing rather than this attribution profile.

The production route is now rowwise-gqa4096+31b-full-online-row2+rowwise-first-tile+batched-swa-row2-score. It passes 42/42 normal and 42/42 ASan/UBSan tests. Reprofile before choosing the next kernel because score and P*V are now approximately balanced. Evidence is in artifacts/2026-08-25-swa-score-row2.

The fresh-process 64K server gate confirms that the operator and shared-prefix gains survive cold deployment. Exact-order row2 scoring lowers visible TTFT from 2,809.768/2,837.049/2,830.016 seconds to 2,278.434/2,260.390/2,255.804 seconds. The means are 2,825.611 versus 2,264.876 seconds: -19.84%, or 1.248x, with every candidate sample faster than every baseline sample and only 0.43% candidate CV. All requests report exactly 65,536 prompt tokens. Each strict second turn reuses all 65,536 tokens and evaluates only 26 appended tokens. Its mean visible latency changes from 2.929 to 2.731 seconds (-6.75%), but 11.60% candidate CV makes that small prefix-performance delta inconclusive; exact suffix-only reuse is the acceptance result. Evidence is in artifacts/2026-08-25-server-prefix-31b-row2score-65536.

• FlashAttention: exact attention with IO-aware tiling
• FlashAttention-2: parallelism and work partitioning

2026-08-25: first-tile isolation makes batched 31B SWA exact and production-safe

The first-divergence harness resolved the repeated-layer failure. With a production-built prefix, rowwise and scalar-batched SWA produce bit-identical activations, logits, and committed K/V at 32-, 1,024-, 2,048-, and 4,096-token boundaries. A two-engine lockstep scan instead found the first persistent difference at position zero: layer-zero residual NMSE was only 5.4534158079e-13, but layer one then committed different K and V, and the perturbation amplified through the 60-layer stack. The earlier candidate had batched the empty-cache tile only when the configured context made the SWA cache capacity 1,024, explaining why small-context operator and model checks missed it.

The production route now evaluates the empty-cache 32-token tile rowwise and uses the existing three-dispatch scalar batched split-GQA path for every later 31B/F32 SWA tile. No encoder rotation or new approximation is required. A full 4,096-token lockstep scan across 128 tiles reports bit-identical K/V in all 60 layers. The shared-prefix 4K gate improves the 32-token extension from 0.834935 to 0.632096 seconds (-24.29%, 1.3209x); the shared-prefix 32K gate improves it from 1.384851 to 1.183060 seconds (-14.57%, 1.1706x). Both routes are independently bit-deterministic and final logits are identical.

Promoted stage profiling attributes the gain correctly. At 4K, the latest three-sample mean SWA attention falls from the recorded 0.482717 baseline to 0.190399 seconds (-60.56%), while full attention remains approximately 0.051 seconds. Profiled wall time is 0.554989 seconds versus the recorded 0.914530 baseline; these non-interleaved stage values are attribution rather than promotion evidence. The production 31B/F32 configuration is now rowwise-gqa4096+31b-full-online-row2+rowwise-first-tile+batched-swa. Evidence is in artifacts/2026-08-25-swa-first-divergence.

The promoted route also passes the deployed fresh-process 32K gate. Three cold visible-TTFT samples are 900.264/914.477/949.774 seconds (921.505-second mean, 2.26% CV), versus the final rowwise-SWA baseline mean of 1,180.671 seconds: -21.95%, or 1.281x. Exact-prefix extension samples are 1.664/2.013/1.732 seconds (1.803-second mean), -12.65% versus the 2.064-second baseline mean. All requests report the exact 32,768-token prompt, and every extension log records 32,768 reused plus 26 appended tokens. Evidence is in artifacts/2026-08-25-server-prefix-31b-batched-swa-32768.

2026-08-25: Metal 4 TensorOps are an M4 experiment, not an assumed replacement

Apple's current Metal 4 guidance adds one useful future axis but does not justify replacing the existing M4 kernels without measurement. Inline tensor operations can retain a matrix result in a cooperative tensor distributed across participating threads' register storage, avoiding an intermediate trip through threadgroup or device memory. Apple also recommends static slice extents when dimensions are known because they avoid dynamic per-element bounds handling. Those properties could eventually help fuse dequantization, projection, and activation work in the fixed Gemma 4 graph.

The hardware qualification matters: Apple documents dedicated neural acceleration for TensorOps on Apple10-family GPUs beginning with M5/A19, while older Apple GPUs run the same API through an optimized shader fallback. Gem's target machine is M4 Max, and the current kernels already use SIMD-group matrix operations. TensorOps therefore remains a separately gated compatibility and performance experiment, not a presumed speedup. It also must not be mixed into the active batched-SWA server comparison.

For SWA specifically, prior substage samples put normalization near 0.6% of attention time; fusing softmax alone has almost no useful ceiling. After the fresh-server gate, reprofile the promoted q8k8 score path and its value pass. Choose exactly one next axis from that evidence: improve Q/K score tiling if scoring still dominates, or reduce normalized-score/value traffic if P*V now dominates. Any fusion must preserve the now-established empty-cache first-tile arithmetic and pass the lockstep K/V and backend-op gates.

Apple's profiling guidance narrows that choice further. Occupancy alone is not a verdict: correlate it with limiter, utilization, and bandwidth counters. Low occupancy can mean resource exhaustion or simply too little work; very high occupancy can still accompany cache thrashing. For the current wrapped-SWA profile, score and P*V are now approximately balanced, so the next local experiment must wait for an Instruments replay of the retained trace. If register or threadgroup pressure caps occupancy, do not increase the row tile. If memory-read stalls dominate with headroom, screen exact- order four-row K reuse. If normalized-score/V traffic dominates, keep the two-row score kernel and screen a value-side reuse change instead. No local GPU experiment should overlap the 64K server gate because it would invalidate the timing.

MLX provides an independent implementation-level signal for the same locality axis. Its January 2026 two-pass vector-SDPA change was tuned on M3 Ultra, M4 Max, and M5, and changed work ownership so one threadgroup handles all queries associated with a given key. The stated purpose is to maximize locality for very long GQA sequences. Gem's two-row score kernel already applies a bounded form of that mapping by sharing each staged K pair across adjacent query rows. A four-row screen would extend the same ownership principle while retaining Gem's per-dot-product accumulation order; MLX's fused online-softmax arithmetic itself is not a correctness oracle because that arithmetic previously failed Gem's authentic repeated-layer quality gate.

Sources:

• Apple: Running inline ML operations in a shader with Metal 4
• Apple: Metal sample code
• Apple: Porting Metal code to Apple silicon
• Apple: Finding your Metal app's GPU occupancy
• Apple: Measuring the GPU's use of memory bandwidth
• Apple: Discover new Metal profiling tools for M3 and A17 Pro
• MLX PR 3023: Faster two-pass SDPA
• MLX Steel attention kernel

2026-08-25: Open-TQ-Metal audit keeps Gemma int4 KV closed but exposes a decode axis

A Paperbridge search surfaced Open-TQ-Metal, a 2026 Apple-Silicon study and MIT-licensed reference implementation. Its headline 128K result must be scoped carefully: the 128K quality/capacity claim is for Llama 3.1 70B, while the paper's own ablation says Gemma 4 31B int4 and int8 KV fail beyond roughly 950--1,000 tokens because its unit attention scale amplifies error through 60 layers. The paper also states that Gemma long-context perplexity and retrieval quality were not evaluated. This independently supports Gem's existing decision to keep F32 KV as the 128 GiB production default and reject int4 KV for 64K/128K hax use. Top-1 parity at a moderate prompt is not sufficient evidence for long-context quality.

The transferable part is sequence-parallel decode. Open-TQ-Metal partitions a single long KV scan into 512-token chunks, emits partial (maximum, denominator, accumulator) state, and merges it in a second ordered phase. That is distinct from Gem's rejected two-segment prefill row4 experiment: at one-token 128K decode, the current full-attention route exposes only one long-lived workgroup per query-row/head block and may have much more sequence-parallel headroom. Gem already preallocates KV storage, so it also avoids the reference engine's reported per-token cache-concatenation penalty. Before implementing anything, capture an authentic 31B decode-stage profile at a populated 64K or 128K position. Screen an F32 split-K full-attention decode kernel only if those ten global layers exceed the 3% end-to-end promotion ceiling. It must retain Gem's F32 KV, exact causal coverage, and the established logit/determinism gates; compressed-domain arithmetic is not part of that axis.

Sources:

• Open-TQ-Metal paper
• Open-TQ-Metal HTML, method and limitations
• MIT-licensed Gemma reference implementation

2026-08-26: populated decode profiling selects full-attention P*V sequence parallelism

gem-decode-profile now builds a requested prefix once and uses exact KV checkpoint/restore for its unmeasured warmup and every repetition. A 31B smoke reports one prefix build and four rollbacks for three samples, removing the previous four-ingestion cost at 64K/128K. The new gem-decode-attention-profile isolates the production F32 split-GQA score, normalization, and value passes at exact 12B/31B full/SWA shapes without claiming synthetic data as a model-level result. Both diagnostic binaries remain free of llama/ggml dependencies; normal and ASan/UBSan suites pass 43/43.

At the 31B full-attention 64K shape, seven 20-iteration samples put score at 3.154 ms (22.19%), normalization at 0.115 ms (0.81%), and value accumulation at 10.947 ms (77.01%). Effective unique K bandwidth is about 170.2 GB/s, while the value pass reaches only 49.0 GB/s. The 128K result scales linearly and preserves the same split: 6.314/0.356/22.379 ms and 21.74%/1.23%/77.04%, with value at 48.0 GB/s. The bounded 31B SWA operator remains roughly 0.255 ms at a 1,024-token live window, so the selected experiment must be gated to the ten context-growing full-attention layers.

This evidence refines the online-source transfer. llama.cpp's Metal vector decode divides KV over 32 workgroups and merges partial softmax state in a second kernel. MLX PR 3023 similarly increases two-pass parallelism for very long sequences. MLX PR 4077, merged August 18, 2026, improves GQA-8 decode by reading each K/V byte once and reports up to 14% operator and 7% full-model gains at 32K on 40-core Max hardware. Its register-pressure gate stops at head dimensions 64/128; Gemma 4 31B uses 512. Gem's current separated score/value passes already share each K and V tile across all eight query heads without holding eight 512-wide online accumulators in registers. Therefore the next axis is the value pass, not a verbatim fused MLX kernel and not another score change. Its exact 31B full geometry exposes a cheaper first screen than split-K: the current 64-dimension tile launches only 4 KV heads * 8 dimension blocks = 32 threadgroups on the 40-core target. A specialized 32-dimension tile would launch 64 groups, halve per-group V staging and accumulators, read every V element once, and preserve each output dimension's token accumulation order. Screen that exact-order geometry first. Sequence-parallel value accumulation plus partial-output reduction remains the fallback if 64 dimension groups still leave bandwidth unused, because it changes reduction order and requires extra scratch/dispatch. A populated full-model 64K profile must still establish the whole-token ceiling before either implementation is promoted.

The populated 64K profile now establishes that ceiling directly. One authentic 31B token at position 65,536 averages 243.076 ms of GPU time with 98.91% stage-counter coverage. KV attention uses 183.675 ms, or 76.40% of sampled time. Combining that model result with the exact-shape global operator profile puts the existing ten-layer full-attention value pass near 109.5 ms, approximately 45.5% of the sampled token. A value-only improvement of about 6.6% can therefore clear the 3% model promotion floor.

The exact-order 32-dimension candidate clears its operator gate decisively. It reduces value workgroup storage from 9,216 to 5,120 bytes and increases the exact global dispatch from 32 to 64 threadgroups without changing V bytes or per-output token accumulation order. An unwrapped 8,193-token CPU/current/candidate differential is bit-identical. In a 64K ABBA bracket with seven 20-iteration samples per arm, the existing value pass averages 11.621/11.147 ms and the candidate 8.009/8.001 ms, a bracketed improvement of roughly 29.7%; total three-stage attention improves about 23.2%. Normal and ASan/UBSan suites pass 43/43, and the adjacent MTL0 FLASH_ATTN_EXT oracle passes 4,793/4,793. Diagnostic selector 18 retains selector 17 prefill/SWA behavior and uses the candidate only for 31B F32 global decode at 8K+. Its populated 8K smoke passes; the decisive 64K model profile reduces token GPU time from 243.076 to 229.678 ms (-5.51%), exceeding both the 3% floor and the larger 3.11% run CV. KV attention falls 13.62%. At 128K, the stable closing operator bracket reduces P*V from 22.681 to 16.051 ms (-29.23%) and complete attention from 29.402 to 23.332 ms (-20.64%); all four raw arms are retained because the opening baseline frequency- ramped. Authentic final logits after 8,193 tokens are byte-identical across selectors 17/18 with matching SHA-256 b9ea77055de12bd4a327158a7feb3656c1997f228728d8299f47d67142c4c87e. Selector 18 is therefore the promoted 31B/F32 default. Normal and ASan/UBSan suites pass 43/43 after promotion, while selector 17 remains the diagnostic rollback.

Long-position iteration now also has an explicitly non-promotional shortcut. The diagnostic seeded-prefix API fills the already allocated KV buffer with finite bytes, advances the logical token position, and leaves the production 60-layer graph and dispatch routing unchanged. A 64-position smoke reports prefix_mode: seeded-kv, zero prefix builds, four checkpoint rollbacks, and the normal 1,383-dispatch decode graph. Normal and ASan/UBSan suites remain 43/43, including hax integration and all native dependency guards. Use this mode only to screen stage timing and dispatch geometry while the expensive evaluated-prompt 64K/128K gates remain authoritative.

PaperBridge retrieval of FlashInfer, FlashDecoding++, and Open-TQ-Metal refines the deferred sequence-parallel design. Attention states over disjoint token ranges are composable, and the Apple-specific Open-TQ implementation confirms that a separate partial/reduce dispatch avoids the single-threadgroup 128K serial scan. Gem should not copy its compressed online-softmax kernel: the production graph already computes globally normalized scores and shares each F32 V tile across all eight query heads. The isolated next axis is therefore a split normalized P*V pass. At the exact 31B shape, each token partition writes 32 * 512 F32 partial outputs and a deterministic second kernel sums partitions in ascending order. Split counts 2/4/8/16 require only 128 KiB/256 KiB/512 KiB/1 MiB of reusable scratch, raise the current 64 value threadgroups to 128/256/512/1,024, and keep aggregate V reads unchanged. Test each count independently; do not combine it with a dimension tile, token tile, score path, or KV representation change. Because partial summation reassociates F32 additions, this route needs a nonmultiple-tail differential, repeated model-logit determinism, and the ordinary quality thresholds rather than a claim of bit-identical arithmetic.

The 128K cold-server runner is now safely resumable because one repetition takes hours. Every new gate writes gate-config.json binding model and prompt paths/digests, context, exact prompt count, KV type, output limit, prefix-extension text, and repetition count. GEM_GATE_RESUME=1 additionally requires byte-equivalent canonical print-config, validates every completed measurement/SSE/log tuple, and reruns only incomplete repetitions in fresh processes. A CPU-only copied-artifact test accepts three valid completed samples while independently rejecting an occupied port, a changed native route, and a changed prompt digest. That check is now the permanent gem.long_context_server_gate_resume CTest rather than a one-off manual fixture. Shell syntax and ShellCheck are clean, and the normal plus ASan/UBSan suites pass 44/44.

Long decode profiling now exposes prefix progress without changing its execution schedule. The new progress API is a nullable callback inside the existing production evaluation loop and fires only after each already-completed batch; the ordinary API calls the same implementation with no callback. gem-decode-profile reports every 4,096 tokens to stderr by default and adds prefix wall time/rate to stdout JSON. A 64-token smoke emits exactly the requested 32/64-token reports and retains one prefix build, four checkpoint rollbacks, 60 layers, and 1,383 decode dispatches. Normal and ASan/UBSan suites pass 44/44 after adding the server-gate resume regression.

The current web audit adds a useful warning to the deferred split-PV plan. MLX's merged GQA-8 two-pass kernel reads each K/V subchunk for several query heads and reports a 7.0% full-model gain at 32K on a 40-core Max, but it gates only head dimensions 64/128 because register pressure changes the winning geometry; it does not validate Gemma 4's dimension 512. More importantly, a separate MLX issue localized deterministic generation collapse to the exact boundary where another GQA two-pass path became active. This is not evidence that Gem's separated normalized PV split is incorrect, but it makes threshold-crossing model tests mandatory. The screen must include positions immediately below/at/above the route threshold, a nonmultiple partition tail, every candidate split count, and repeated authentic logits. A block-count sweep alone is not evidence: a reported 100K GQA shape remained at roughly 25--28% of practical reduction bandwidth across its sweep. Apple also warns that neither high nor low occupancy diagnoses a bottleneck by itself; select the split count from P*V wall time plus limiter/bandwidth counters, not workgroup count alone.

Residual selection is now machine-checkable rather than an informal multiplication. The new analyze-decode-residual.sh binds an authentic populated-position profile to its exact-shape operator receipt and accounts for the ten 31B full-attention layers. On the promoted 64K route, current PV supplies a 34.83% theoretical whole-token ceiling, so an 8.61% value-pass improvement would be sufficient to reach the 3% whole-token floor. The operator accounts for 70.98% of measured KV-attention time, remaining physically consistent. This admits screening but explicitly does not promote a kernel. The same analyzer will consume the running authentic 128K receipt before any split-PV implementation begins.

The production server now reuses that batch-boundary progress API for long-prompt operability and cancellation. Streaming requests receive a content-free SSE keepalive after every already-completed prefill batch, so a disconnected hax client leaves at most one submitted batch before evaluation stops; non-streaming aggregation keeps the same callback with a no-op heartbeat writer. Prompts of at least 4,096 tokens log logical progress, evaluated-token throughput, elapsed time, and ETA every 4,096 appended tokens, including exact-prefix reuse in the logical numerator. The graph, batch size, command buffers, and first-visible-event definition are unchanged. A failed callback now exits the shared evaluation loop with an explicit cancellation error after preserving all already-committed batches; the next serial request takes the ordinary cache-miss reset path. The multi-hour server gate now treats the exact 100% prefill log record as part of each sample's atomic evidence tuple, both immediately after a fresh request and during resume validation. An exact-prefix extension must likewise log its extended logical prompt count at 100%, preventing a truncated or stale server log from being accepted alongside otherwise valid SSE/measurement files. The permanent CPU fixture now uses a 4,096-token sample and proves that deleting only this final log record forces a fresh server attempt instead of incrementing resumed_samples. An authentic regression runner is also prepared for the post-profile GPU window. It starts the native server, submits a canonical long prompt with curl's bounded timeout, requires status 28 and an observed SSE keepalive, waits a bounded interval for /health, and then requires a visible 16-token completion from a second request to the same process. It rejects a cancellation prompt that reaches the exact 100% prefill record. Its retained config, partial stream, logs, recovery measurement, and summary distinguish callback cancellation from merely killing and restarting the server. The complete runner is permanently covered by a serial CPU HTTP/SSE fixture: the first request emits a keepalive and remains busy past curl's timeout, health cannot respond until that handler returns, and the second request supplies visible content plus exact usage. Normal and sanitizer CTest configurations pass this orchestration gate; authentic GGUF execution remains intentionally separate. The queued 31B/F32 receipt uses the exact 131,072-token server-gate prompt at context 131,328, port 18085, a five-second disconnect, and a 60-second recovery ceiling; its command and runner hash are frozen in prefill-cancel-preflight-131072.json before execution.

A non-invasive process sample at 51:33 establishes liveness for the pre-heartbeat 128K profiler binary. Its 21.8 GiB physical footprint is stable below a 21.9 GiB peak, VM reports zero throttled pages, the main thread is waiting in MTLCommandBuffer waitUntilCompleted from the production spanned evaluation loop, and the Metal queue thread is inside IOGPU command-buffer submission. This is active GPU prefix evaluation rather than a silent host deadlock. The compact receipt is retained as live-profile-liveness-51m.json; the transient full process sample was not indexed.

The authentic populated 128K receipt completed after one prefix build and four exact checkpoint rollbacks. Production dim32 decode averages 416.378 ms/token across 414.381/409.375/425.378 ms samples (2.16% maximum mean deviation), with 99.43% timestamp coverage. KV attention averages 333.533 ms, 80.56% of sampled time. The matching exact-shape dim32 receipt attributes 233.321 ms to the ten global-attention layers and 160.513 ms to their normalized PV passes. That is a physically consistent 69.95% of authentic KV-attention time and a 38.55% theoretical whole-token value ceiling; only a 7.78% value-pass improvement is required to clear the 3% model floor. Sequence-partitioned PV screening is therefore admitted for splits 2/4/8/16 as separate diagnostic axes. No split is promoted by this attribution result.

The frozen authentic cancellation gate then passed at the exact 131,072-token prompt. Curl timed out with status 28 after observing four completed-batch keepalives; the still-running native process returned /health 200 one second later and served a second request successfully. That recovery request produced its first visible content event in 484.832 ms and completed in 567.813 ms with 14 prompt, 6 completion, and 20 total tokens. The cancellation request never logged 100% prefill. This closes the real-GGUF cancellation invariant, not the separate three-process 128K latency gate. Afterward all 46 normal CTests passed, and all 46 ASan/UBSan CTests passed.

Seeded-KV production-route cross-checks are also retained at 64K and 128K. They report 184.781 ms and 294.909 ms mean GPU time respectively, with the expected 60 layers, 1,383 dispatches, zero prefix builds, and four rollback checkpoints. These receipts prove that the production geometry is operational at both capacities without another multi-hour prefix build. Artificial finite KV bytes do not reproduce evaluated-model data, so these measurements are explicitly non-promotional and must not be compared against the authentic populated-position timings.

Evidence is in artifacts/2026-08-26-long-decode-profile. The cancellation receipt is in artifacts/2026-08-26-prefill-cancel-31b-131072.

The admitted split-P*V experiment is now implemented as diagnostic selector 19 without changing the production selector 18. Its first dispatch partitions the already-normalized probability/value dot product into contiguous token ranges while retaining the dim32 output geometry; a second dispatch sums partials in ascending partition order. Splits 2/4/8/16 use 128 KiB/256 KiB/512 KiB/1 MiB of reusable scratch. At 64K, retained three-repetition receipts reduce complete attention from 11.035 ms to 5.883 ms and value-plus-merge from 7.810 ms to 2.654 ms with split16. At 128K, the stable baseline is 22.344 ms total and 15.819 ms value; split16 reaches 11.852 ms and 5.324 ms in its stable closing arm, while split8 reaches 12.934 ms and 6.412 ms. Earlier arms show real GPU frequency nonstationarity, so the populated-model gate remains authoritative.

Correctness is green before that model timing. The native operator differential covers the exact 8,192 threshold, 8,193 nonmultiple tail, every split count, and deterministic repeats. Selector 19 is byte-identical to selector 18 below threshold at 8,191 tokens. At 8,192 it retains identical argmax and ordered top ten with correlation 0.999999999670 and NMSE 4.64e-10; at 8,193 the same checks pass with correlation 0.999999999754 and NMSE 8.28e-10. Repeated selector-19 8,193-token logits are byte-identical. The adjacent MTL0 FlashAttention oracle passes 4,793/4,793, and all 46 normal plus all 46 freshly rebuilt ASan/UBSan tests pass.

The current upstream recheck supports the observed geometry but does not supersede local evidence. MLX now routes long GQA decode to a two-pass kernel and selects 1,024 first-pass blocks on its large Apple device class at 65,536+ tokens; it also documents that its reduction consumes partials in SIMD-width chunks. Gem's split16 produces exactly 1,024 value workgroups at the 31B dim32 shape, which helps explain why it saturates where smaller splits do not. MLX's kernel still supports only head dimensions 64/128 on its specialized GQA path, whereas Gemma 4 uses 512, so copying its block policy or fused online-softmax state would be unjustified. Apple's current synchronization guidance also confirms that a write/read dependency between partition and merge dispatches requires the explicit buffer barrier already present here. With P*V reduced, the unchanged score pass becomes roughly half the 128K operator; score/K work is the next kernel axis only after selector 19 clears the populated-model and deployed gates.

The authentic populated 64K model gate now clears that promotion floor decisively. Selector 19 averages 177.717 ms/token across 157.732/204.886/170.532 ms samples, versus selector 18's retained 229.678 ms mean and 228.322/221.690/239.023 ms samples. This is a 22.62% whole-token reduction; every candidate sample is faster than every baseline sample despite the candidate's 15.29% maximum mean deviation. KV attention falls from 158.659 to 104.745 ms (-33.98%), timestamp coverage is 98.43%, and the receipt retains one authentic prefix build, four checkpoint rollbacks, 60 layers, one command buffer, and the expected 1,393 dispatches. This admits the populated 128K confirmation but still does not change the production default.

The matching authentic populated 128K receipt completes promotion. Selector 19 averages 304.832 ms/token across 296.518/300.203/317.775 ms samples, versus selector 18's retained 416.378 ms and 414.381/409.375/425.378 ms samples. Whole-token GPU time falls 26.79%, and every candidate sample is faster than every baseline sample. KV attention falls from 333.533 to 226.145 ms (-32.20%); stage-counter coverage is 98.79%. The receipt retains exactly one evaluated prefix build, four checkpoint rollbacks, 60 layers, one command buffer, and 1,393 dispatches. Together with the threshold/tail/repeat logits, 4,793/4,793 adjacent oracle, and complete normal and sanitizer suites, this promotes selector 19 as the 31B/F32 default. Selector 18 remains the unsplit dim32 rollback. The production route now reports its additional 1 MiB reusable value-partial scratch explicitly; no 12B, F16, prefill, SWA, tokenizer, GGUF, server, or hax protocol behavior changes in this promotion.

The post-promotion real hax acceptance smoke supplies an 8,203-token raw prompt through gem agent, forces direct target decoding, and returns exactly OK. The server reports complete prompt progress, hax exits successfully, and the launcher leaves no listener on its port. A preceding 8,448-context attempt also completed prefill but correctly returned a context-full length finish because hax's requested completion allowance exceeded the remaining envelope; the passing receipt therefore uses context 9,216 and retains that earlier outcome only as an admission check. Evidence is in artifacts/2026-08-26-long-decode-profile/hax-smoke-production-split16-8k.txt.

Sources:

• MLX PR 3023: Faster two-pass SDPA
• MLX PR 4077: Read each K/V byte once in GQA-8 decode
• MLX issue 4022: two-pass GQA routing-boundary correctness failure
• MLX issue 3837: long-context GQA bandwidth and block-count sweep
• Apple: Finding your Metal app's GPU occupancy
• FlashInfer recursive attention and KV sequence parallelism
• FlashInfer paper
• FlashDecoding++ paper
• Open-TQ-Metal paper
• llama.cpp Metal vector attention

2026-08-26: upstream recheck keeps long-prefill attribution ahead of kernel selection

A web recheck during the final 128K server gate found no upstream result that should bypass Gem's pending authentic stage profile. Current MLX uses a specialized two-pass GQA decode only for head dimensions 64/128 and selects 1,024 first-pass blocks at 65K-plus keys on its large Apple GPU class. That independently supports Gem's promoted split-16 decode parallelism, but it does not transfer directly to Gemma 4's 512-dimensional Q/K and 32-dimensional V geometry. MLX's fused full attention now covers dimensions through 256, while its own Gemma 4 head-dimension-512 report documents an unfused fallback. A separate MLX long-prefill report also identifies the need to chunk work across command-buffer boundaries to avoid multi-second GPU watchdog exposure.

BaseRT is useful architectural corroboration: its published native runtime uses sequence-regime specific SIMD-matrix GEMMs, bounded prefill chunks, and tiled online-softmax attention. Its public evaluation, however, stops at 2,048 prompt tokens and its production Metal engine is proprietary, so it supplies neither a 64K/128K performance claim nor adaptable kernel source. Open-TQ-Metal's fused int4 KV path remains a separate representation/quality axis; its own Gemma 4 analysis warns that the model's attention scale amplifies angular quantization error. It must not be combined with the current F32 kernel experiment.

The resulting order stays evidence-led: finish the frozen 128K cold/prefix gate; measure a seeded 128K production tile by projections, SWA attention, and full attention; then optimize the largest stage. If full attention dominates, screen a bounded BQ/BK tiled online kernel specialized for head dimension 512. If Q4 projections dominate, tune the existing 32-row Q4 GEMM instead. The decode-only score-token-tile experiment remains isolated and cannot establish prompt-processing speed.

Sources:

• MLX Metal SDPA dispatch and long-GQA block policy

• MLX head-dimension-512 fused-SDPA gap

• MLX long-context full-attention watchdog analysis

• BaseRT native Metal runtime paper

• BaseRT public repository and engine licensing boundary

• Open-TQ-Metal compressed-domain attention

• Apple Metal feature-set tables

2026-08-26: authentic 64K gate rejects direct, staged q8k8, and row4 prefill

Three bounded global-attention candidates were screened against the selector-19 row2 production path without changing GGUF, graph, KV representation, or online-softmax semantics. Direct device K/V reads were bit-identical but slower by 31.38% at 64K and 34.94% at 128K. The staged 8x8 score/value candidate was also bit-identical but slower by 67.86% and 67.69%. These results reject both Family-9 cache hypotheses for the fixed head-dimension-512 geometry.

Four-row ownership looked promising only in synthetic evidence: it reduced the isolated global operator by 14.92% at 64K and 14.35% at 128K, and a seeded 128K production-graph screen improved 11.05%. The authentic gate built one real 65,536-token selector-19 prefix, then interleaved three checkpointed 32-token extensions. Row2 took 1.716713/1.700924/1.608773 seconds (1.675470 mean), while row4 took 1.711669/1.698021/1.651547 seconds (1.687079 mean): row4 regressed 0.693% and failed the 3% promotion floor. Outputs were bit-deterministic and the cross-route logits were byte-identical, so correctness does not rescue the absent whole-model gain. Selector 21 and its CLI/harness routing were removed; selector 19 remains production.

The installed Xcode beta now provides a usable profiling path even though the active developer directory is Command Line Tools. scripts/profile-metal.sh discovers that Xcode without changing global xcode-select, accepts an exact diagnostic command, and can select an alternate template. A 128K seeded selector-19 Metal System Trace was captured. Adding Xcode's separate Metal GPU Counters instrument activates counter profile 3 and shader-PC sampling, although the headless XML does not expose occupancy or bandwidth aggregates. A new diagnostic MTLCaptureManager hook then captured exactly one untimed row2 dispatch to a .gputrace; the profiler's timing samples are completed before capture begins and its JSON marks capture:true. Use the 4K capture for resource, spill, occupancy, and heat-map inspection because occupancy is geometry-dependent and GPU capture serializes bound buffers. Retain the 128K system trace for long-context execution timing and memory behavior instead of serializing a multi-gigabyte KV buffer into the GPU document.

Evidence is in artifacts/2026-08-26-prefill-row2-direct-cache, artifacts/2026-08-26-prefill-row4-production-128k, and artifacts/2026-08-26-selector19-metal-trace-131040.

Sources:

• Apple: Capturing a Metal workload programmatically
• Apple: Metal capture descriptor
• Apple: Metal performance heat maps

2026-08-26: Family-9 constraints refine the bounded F32 prefill candidate

Google's released Gemma 4 model card establishes that both dense targets have a 1,024-token sliding window and support 256K context, not only the immediate 64K/128K acceptance rungs. Any production prefill candidate must therefore keep scratch bounded independently of context length. The existing q8k8 diagnostic's materialized score tensor remains useful attribution evidence but cannot become the long-context route.

Apple's Family-9 guidance says threadgroup, device, and constant storage use the same flexible on-chip cache hierarchy; explicit threadgroup staging is not inherently faster when the working set already benefits from that cache. Apple separately identifies threadgroup memory and registers as occupancy limiters. Consequently, the selector-20 16 KiB K tile remains a measured diagnostic, not an assumed improvement. A new standalone gem-metal-counter-info tool records the target device's public counter sets and supported sampling boundaries before counter instrumentation is chosen. Xcode-only occupancy or limiter counters must not be treated as programmatically available unless this receipt names them.

The existing all-F32 q8k8 score kernel and row2 online kernel expose a more specific fused design if the seeded 128K profile attributes the largest stage to global attention. Preserve the row2 GQA ownership: one 256-thread workgroup handles two causal rows and one KV head, with eight SIMD groups owning the eight query heads. For each eight-token KV tile, one SIMD-matrix operation can form the 8-query-head by 8-token score matrix for each row. The eight SIMD groups then retain the same two 512-wide F32 output accumulators they own today and share the staged V tile. This keeps K/V reuse across all eight GQA heads and both rows, halves tile barriers relative to the current four-token loop, and maintains only a constant approximately 16.5 KiB threadgroup footprint. Online maxima, denominators, and accumulator rescaling preserve exact dense attention without an O(context^2) score buffer. This geometry is preferable to a workgroup per query head, which would reread each shared K/V stream eight times. It is still only a design candidate until the 128K stage profile, pipeline resource limits, oracle correctness, and one-axis ABBA measurements admit it.

Sources:

• Google Gemma 4 model card
• Apple Family-9 Metal shader performance guidance
• Apple GPU occupancy guidance
• Apple SIMD-group matrix operations
• FlashInfer attention-state composability

2026-08-26: audited 128K service gate and stage profile select direct-cache full attention

The frozen selector-19 production binary passed the three-repetition 131,072-token HTTP/SSE gate. Cold first-visible latency was 6,418.141, 6,322.454, and 6,388.358 seconds, averaging 6,376.317 seconds with 0.768% CV. The immediately following extension requests reused exactly 131,072 tokens, appended 26, and reached their first visible event in 3.480, 3.394, and 3.342 seconds, averaging 3.405 seconds. The independent audit verifies the actual model, prompt, and frozen-binary digests, all raw/summary measurements, SSE usage and finish reasons, [DONE], and exact server log markers. This is a 1,872.4x cold-to-prefix first-visible speedup and establishes that prefix reuse is the critical operational path for repeated 128K hax work. The gate does not make cold ingestion fast: that remains the next performance problem.

The subsequent seeded-KV production-graph profile measures one 32-token tile at position 131,040 without claiming synthetic cache contents as correctness or throughput evidence. Across three repetitions, mean GPU time is 1.926 seconds with 0.395% CV and 99.865% timestamp coverage. KV attention occupies 1.605 seconds, or 83.45% of sampled GPU time. The ten global layers account for 1.455 seconds; all 50 sliding layers together account for only 0.151 seconds. Projection-family work is 0.303 seconds, or 15.77%. This selects only the global head-dimension-512 attention path for the next one-axis experiment.

The target M4 Max reports Apple Family 9, 32 KiB maximum threadgroup memory, stage-boundary sampling, and only GPUTimestamp in its public Metal counter sets. Occupancy, bandwidth, and limiter evidence therefore requires Xcode GPU Trace. Apple's Family-9 guidance also says a threadgroup-memory software cache can lose to direct device reads because both use the same flexible on-chip hierarchy; Apple's occupancy guidance identifies both threadgroup allocation and register pressure as possible launch limits. This corrects the earlier assumption that the approximately 16.5 KiB staged q8k8 layout should be the first implementation.

The first bounded candidate should retain row2/GQA ownership and exact online softmax, but load coalesced K/V directly from device storage so the Family-9 cache can reuse them across the eight SIMD groups. It must specialize the fixed 31B geometry with function constants and preserve context-independent scratch. Only after its resource report, correctness oracle, and ABBA screen should a second one-axis variant add explicit K/V staging. Conditional online-softmax rescaling from FlashAttention-4 is worth screening separately, but its Blackwell asynchronous pipeline and software exponential are hardware-specific and should not be transplanted. FlashInfer's bounded recursive state remains the right compositional model. Existing exact-shape data at 128K further prioritizes V reuse: value accumulation is 77.04% of global attention, scores 21.74%, and normalization 1.23%.

Evidence is in artifacts/2026-08-26-server-gate-31b-split16-131072 and artifacts/2026-08-26-prefill-stage-attribution-31b-131040.

Sources:

• Apple Family-9 Metal shader performance guidance
• Apple Family-9 GPU profiling workflow
• Apple Metal Performance Primitives programming guide
• FlashAttention-4
• FlashInfer
• MLX head-dimension-512 fused-SDPA gap

2026-08-26: focused trace inspection and first-party recheck keep limiter replay mandatory

A narrow inspection of the retained row2 GPU capture found compiler fields for uniform and temporary register counts, threadgroup memory, spilled bytes, and spill/reload passes. The capture also retains the complete runtime Metal library and all compiled attention functions. Its headless archive, however, does not expose a trustworthy mapping from the selected attention_prefill_online_gqa_31b_full_row2_f32 pipeline state to one compiler blob. Several blobs contain spill diagnostics, including records with five compiler spills, but attributing any of those records to row2 would be speculation. Replay the focused .gputrace in Xcode, group commands by pipeline state, and read the selected shader's Performance timeline/heat map. Use the retained 128K System Trace for the authentic long-context limiter and bandwidth view.

The capture path now removes the remaining command-identification ambiguity. While capture is active only, the production row2 dispatch pushes a debug group named exactly attention_prefill_online_gqa_31b_full_row2_f32; the new labeled GPU document retains that string in both its sorted and unsorted command streams. The profiler also queries the compiled pipeline state and records execution width 32, maximum 1,024 threads per threadgroup, zero static threadgroup memory, 16 KiB dynamic threadgroup memory, and the actual 256-thread launch. This rules out a compiler maximum-thread limit below the launch geometry. It does not rule out runtime occupancy-manager throttling from registers or the shared Family-9 cache, so counter replay remains the selection gate. The query and its fixed dispatch invariants are covered by normal and ASan/UBSan Metal tests.

The two authentically rejected row2 experiments are now fully retired rather than merely hidden from public routing. Their Objective-C pipeline fields, MSL source strings, library concatenation, pipeline creation, dispatch implementations, and cleanup references are removed. The resulting clean labeled capture contains the production row2 debug group and neither rejected shader name. This reduces the always-compiled Metal surface and prevents future compiler records from being mistaken for active candidates. The post-cleanup gate passes all 56 normal tests, all 56 freshly rebuilt ASan/UBSan tests, and the adjacent 4,793/4,793 FlashAttention oracle. No selector or production dispatch geometry changed.

An installed metalperftrace follow-up cannot replace Xcode replay for this command-line compute workload: live collection emitted no session and its one-minute always-on trace returned No session found. Within the existing Xcode trace, raw counter rows belong to the default RT Unit Active counter; the requested Performance Limiters profile has no decoded counter-info rows even though shader-PC samples are present. Do not derive occupancy or bandwidth from those zero-valued default rows.

Apple's current Family-9 workflow makes that distinction material. Low occupancy is actionable only after correlating Compute Shader Launch limiter, Occupancy Manager Target, ALU limiters, L1 eviction/residency, last-level-cache stalls, MMU stalls, and memory bandwidth. Registers, threadgroup storage, stack, and buffer data share the flexible on-chip cache hierarchy on this GPU; therefore neither more threadgroups nor explicit staging is an optimization by itself. The public pipeline properties can confirm execution width, maximum threadgroup size, and static threadgroup memory, but they cannot replace those counter correlations.

Google's released 31B configuration independently confirms the exact long-context shape already implemented by Gem: 60 layers in a five-local/one-global pattern, a 1,024-token local window, 262,144 maximum positions, 32 query heads, 16 local KV heads, four global KV heads, and 512-dimensional global heads. Thus 50 local layers remain bounded while ten global layers carry the quadratic cold-prefill cost. The config sets attention_k_eq_v but no cross-layer KV sharing, so reusing K as V inside each full layer is valid while sharing global caches across layers is not.

This recheck also closes a distracting hardware axis. Metal 4 TensorOps is portable, but Apple's new neural-accelerator acceleration is specific to M5/A19-class GPUs; M4 executes its optimized shader fallback. MPSGraph now exposes causal SDPA and could serve as a diagnostic oracle, but introducing a graph framework into the production engine would violate the small fixed-graph ownership goal unless a measured comparison first demonstrates a transferable kernel advantage. For M4, the next implementation remains a custom bounded-memory global-attention candidate selected from the captured limiter—not an M5 TensorOps rewrite.

Sources:

• Google Gemma 4 31B configuration
• Google Gemma 4 model overview
• Apple: Finding your Metal app's GPU occupancy
• Apple: Discover new Metal profiling tools for M3 and A17 Pro
• Apple: Learn performance best practices for Metal shaders
• Apple: Calculating threadgroup and grid sizes
• Apple: M5/A19 machine-learning acceleration
• Apple: MPSGraph causal scaled-dot-product attention

2026-08-26: full-graph gates reject a 128-row long-prefill batch

The production row2 operator was extended diagnostically from 32 through 128 rows without changing its kernel. At position 131,040, 128 rows complete in 0.509994 seconds versus 0.144215 seconds for 32 rows, so isolated attention delivers 13.04% more rows per second. A seeded stage-profiled 60-layer graph initially reinforced that hypothesis: one 128-row tile took 6.275 seconds versus an estimated 7.404 seconds for four 32-row tiles. Neither result is promotion evidence because the first omits the graph and the second changes encoder boundaries for timestamp sampling.

The ordinary single-pass graph rejected the candidate. A real 4,096-token shared prefix followed by interleaved 128-token extensions measured 6.765 seconds for four production 32-row tiles and 7.449 seconds for one 128-row tile, a 10.12% regression. Both routes were internally bit-deterministic and their final logits were byte-identical. Seeded interleaved screens then measured batch-128 deltas of -7.30% at 8K, -4.33% at 32K, +6.81% at 64K, and +22.29% at 128K. The critical long-context rungs therefore reject the larger batch. All widened workspace, scheduler, SWA, benchmark, and harness routing was removed; production remains 32 rows. The operator profiler alone retains rows through 128 so the scaling evidence remains reproducible without changing server behavior.

The divergence between stage-profiled and ordinary execution does not establish that a long command buffer is intrinsically slow. Apple's compute documentation permits multiple dispatches in one pass, and its MPS tuning guidance favors keeping command buffers queued instead of synchronously waiting. Family-9 profiling guidance instead requires correlating occupancy and cache limiters. The next bounded axis is therefore an interleaved batch-32 comparison of ordinary execution against selected encoder/pass boundaries, with no timestamp counters. Only a measured win would justify attributing the batch result to pass structure.

Evidence is in artifacts/2026-08-26-prefill-batch128-screen.

Sources:

• Apple: Compute passes
• Apple: Metal Performance Shaders tuning hints
• Apple: Family-9 Metal profiling workflow

2026-08-26: unprofiled layer-pass segmentation is neutral and retired

The batch-128 screen left a possible confound: stage-attribution creates many compute encoders in one command buffer, whereas production uses one compute encoder for the complete 60-layer tile. A diagnostic candidate therefore kept selector 19, batch 32, every kernel, every dispatch, and one command buffer fixed, but ended and reopened an ordinary compute encoder after embedding and each model layer. It allocated no counter sample buffer.

At the critical seeded position 131,040, the same-process interleaved screen measured production at 1.908965/1.909300/1.934500 seconds (1.917588 mean) and layer passes at 1.929068/1.930033/1.908960 seconds (1.922687 mean). The candidate regressed 0.266%, while repeated outputs were bit-identical and cross-route logits had correlation 1 and NMSE 0. Encoder boundaries therefore do not explain the earlier batch-128 discrepancy and do not clear the 3% promotion floor. The diagnostic Metal/inference APIs and harness routing were removed. Production remains a single encoder and 32-row tile; the next optimization axis returns to the row2 kernel's per-thread state and 16 KiB staged K/V footprint.

Evidence is in artifacts/2026-08-26-prefill-layer-pass-screen.

2026-08-26: serial K/V tile reuse is rejected

The production row2 kernel stages one four-token K tile and one four-token V tile concurrently, using 16,384 bytes of dynamic threadgroup memory. A one-axis diagnostic reused one 8 KiB tile for K and then V, with 256 bytes reserved for the 64 per-SIMD/per-row scores. This reduced the dynamic threadgroup footprint to 8,448 bytes while keeping 256 threads, row2 ownership, the four-token tile, F32 storage, and online-softmax update order fixed. The price was two additional uniform threadgroup barriers per four cache tokens and a score store/reload.

Finite-data output was bit-identical to production row2, and the adjacent FlashAttention backend oracle passed 4,793/4,793 cases. Interleaved three-repetition GPU timings decisively rejected the candidate: it was 62.44% slower at 4,096 tokens, 61.32% slower at 65,536, and 60.52% slower at 131,040. The nearly context-invariant penalty identifies synchronization and score spilling as the dominant cost; any occupancy benefit from halving threadgroup memory was insufficient. All candidate shader, pipeline, API, test, and profiler routing was removed after recording the raw JSON. Future row2 work should preserve concurrent K/V staging unless it removes barriers through a different ownership or tiling scheme.

Evidence is in artifacts/2026-08-26-row2-shared-stage-screen.

2026-08-26: row2 token tile 8 wins the operator but misses the graph gate

The next one-axis candidate retained row2 ownership, 256 threads, eight SIMD groups, concurrent K/V staging, F32 cache data, causal masking, and token-ordered online softmax while increasing the cache token tile from four to eight. Dynamic threadgroup memory increased from 16 KiB to 32 KiB, but the number of uniform threadgroup barriers over a long prefix was halved. Finite-data output remained bit-identical and the adjacent FlashAttention oracle passed 4,793/4,793 cases.

The operator screen cleared its gate: tile8 improved GPU time by 3.23% at 4K and, over seven repetitions, 4.51% at 64K and 4.41% at 128K. A seeded same-process full-graph harness then compared 32-token extensions through selector 19 with only the full-prefill tile changed. Logits were bit-identical across repeats and routes, but wall time improved only 1.96% at 64K and 2.64% at 128K. Both mandatory long-context gains are below the 3% promotion floor, so the candidate was rejected before an hours-long evaluated-prefix gate. Its shader, pipeline, API, selector, tests, and profiler route were removed.

The generic GEM_LONG_PREFIX_SEED=1 mode remains in gem-long-prefix-ab: it initializes finite KV state once and uses checkpoint rollback for interleaved full-graph screening. Its receipts explicitly report prefix_mode: seeded-kv and prefix_builds: 0; it is attribution evidence only and cannot replace evaluated-prefix correctness or prompt-throughput gates.

Evidence is in artifacts/2026-08-26-prefill-row2-tile8-screen.

2026-08-26: current upstream converges on matrix-tiled query blocks for long prefill

The post-tile8 research block found a structural axis that is materially different from another row2 staging or tile-size adjustment. FlashAttention-2 attributes low utilization to weak query parallelism, non-matrix work, and cross-warp communication. The adjacent llama.cpp Metal oracle implements that shape with eight query rows, 64 cache positions, eight SIMD groups, and 8x8 simdgroup_matrix multiply-accumulates. At the fixed Gemma 4 31B F32 head geometry (DK=DV=512), its dispatch formula requests 28,672 bytes of threadgroup memory. This is a reference design and correctness oracle only; Gem production remains independent C/Metal and does not link ggml.

A current upstream MLX recheck corrects an older note in this document. MLX's newer NAX attention path is limited to head dimensions through 256, but head dimension 512 now remains on a fused Steel Metal path rather than necessarily becoming unfused. That path selects BQ=32, BK=16, four SIMD groups, and one output threadgroup per query tile and query head. Thus two independent maintained implementations now support the same conclusion—batch query rows and perform QK/PV with matrix fragments—while disagreeing on the best fixed tile. The disagreement is useful: tile geometry must still be measured on Gem's exact 512/512, GQA-8, F32-cache workload.

The first bounded candidate will therefore be a specialized BQ=8, BK=64 online-attention kernel, owning one query head per threadgroup. It deliberately trades production row2's cross-head GQA cache reuse for much lower per-lane register state and matrix QK/PV. If correctness and the isolated long-context gate are sound, BQ=32, BK=16 is the next one-axis geometry comparison. F32 KV representation, causal semantics, selector-19 graph structure, and composable online softmax state remain fixed. Open-TQ-Metal's int4-KV and split-K decode claims are excluded from this prefill decision because they change cache representation and workload regime.

Metal 4.1 cooperative tensors are not the immediate target. Rigel's M4 Max measurements report that they execute through the existing shader-core/simdgroup-matrix path rather than a separate matrix unit, while their deployment requirements exceed this engine's current host path. The portable 8x8 SIMD matrix primitive is therefore the appropriate first implementation target.

Evidence and the exact promotion order are in artifacts/2026-08-26-long-prefill-research-block/summary.json.

Sources:

• FlashAttention-2
• llama.cpp Metal flash attention
• MLX Metal SDPA dispatch
• Rigel M4 Max Metal 4.1 study
• Open-TQ-Metal

2026-08-26: BQ8/BK64 matrix-tiled global prefill is promoted

The bounded BQ=8, BK=64 candidate validated the structural hypothesis. One 256-thread threadgroup owns one query head and eight query rows, stages 64 K/V positions at a time, and uses 8x8 SIMD-matrix operations for both QK and PV while retaining exact causal online softmax. Its 18,496-byte dynamic threadgroup footprint changes neither GGUF/KV representation nor the 31B graph outside eligible F32 global-prefill tiles.

Finite-data attention NMSE was approximately 1.1e-12. Isolated attention improved 59.0-59.3% at 4K, 64K, and 128K. More importantly, the authentic evaluated graph improved 7.12% at 4K and 33.38% at 64K over three interleaved repetitions. The 64K comparison preserved argmax and the top-10 set, produced correlation 0.99999958 and NMSE 8.35e-7, and was deterministic within each route. Seeded attribution predicts the gain continues at 128K: 45.36% lower extension wall time.

Selector 19 now owns the BQ8/BK64 production route. Selector 22 preserves the former row2 global prefill as a diagnostic rollback while sharing selector 19's SWA and split16 decode behavior. The normal and sanitizer suites pass 56/56, the adjacent FlashAttention oracle passes 4,793/4,793, and all native dependency guards remain green. This result follows the web-research direction directly: query blocking and matrix fragments mattered much more than further scalar row2 staging tweaks.

Evidence is in artifacts/2026-08-26-prefill-bq8-bk64-promotion.

2026-08-26: promoted BQ8/BK64 passes 8K and 32K native-server gates

The promoted selector was then exercised through the real hax-compatible HTTP/SSE server with a fresh process for every repetition and an exact-prefix extension on each process. At 8,192 prompt tokens, mean first-visible TTFT was 206.982 seconds and the 26-token appended turn completed its first visible event in 1.858 seconds after reusing all 8,192 prompt tokens. At 32,768 prompt tokens, the three cold samples were 847.953, 827.775, and 819.124 seconds (831.617 mean, 1.779% CV); the same exact-prefix extension averaged 2.804 seconds after reusing all 32,768 tokens.

Independent audits recomputed the aggregates from raw SSE and measurement files and verified every 100% progress record, prefix-hit record, model digest, prompt digest, and frozen production-binary digest. Historical same-prompt production gates are 16.64% slower at 8K and 9.75% slower at 32K, but those comparisons include intervening engine changes and are contextual rather than BQ8-only promotion evidence. The evaluated shared-prefix graph remains the one-axis evidence.

The gate work also closed two provenance defects before the measurements: the audit now derives the exact 31B F32 KV allocation from context and the fixed 50-SWA/10-global topology instead of hard-coding the 128K allocation, and the runner now records and fail-closed verifies the production binary digest on resume.

Evidence is in artifacts/2026-08-26-server-gate-31b-bq8-8192 and artifacts/2026-08-26-server-gate-31b-bq8-32768.

The critical 65,536-token rung subsequently passed with cold first-visible samples of 1,940.672, 1,957.977, and 1,914.561 seconds (1,937.737 mean, 1.128% CV). Every process then reused all 65,536 prompt tokens and evaluated only a 26-token appended turn; cached-extension first-visible latency averaged 5.208 seconds, a 372.1x cold-to-warm reduction. The independent raw-artifact audit passed all model, prompt, binary, progress, SSE, usage, and prefix-hit checks. This is 14.44% faster than the closest historical row2-score production gate under the same model/prompt/context layout, subject to the same intervening-change caveat. Evidence is in artifacts/2026-08-26-server-gate-31b-bq8-65536.

2026-08-26: 128K BQ8 reliability passes, but host contention invalidates timing

The first promoted-BQ8 131,072-token server sample completed end to end, returned 16 tokens, then reused the exact 131,072-token prefix and evaluated a 26-token extension. Cold first-visible time was 6,459.234 seconds and the cached extension reached visible output in 4.331 seconds. This is useful long-context capacity, correctness, streaming, and prefix-cache evidence, but not clean performance evidence. Two individual 4,096-token intervals stretched by hundreds of seconds while unrelated macOS services consumed multiple CPU cores. A second repetition encountered the same condition and was stopped at 32,768 tokens rather than promoted as a benchmark.

The runner now writes a macOS host-state receipt before every 64K-or-larger sample, rejects a one-minute load average above one runnable task per logical CPU, monitors that limit throughout the request, and aborts after three consecutive violations. It also owns a caffeinate -i -w assertion for the native server lifetime. These controls first appeared in schema 2; schema 3 added the short-window CPU guard. Current schema 4 additionally requires AC power and no thermal or performance warning before launch, monitors power and thermal state during the request, and makes any recorded violation audit-fatal. The contaminated schema-1 directory is retained but cannot be mixed with new repetitions.

Current primary-source research reinforces the existing matrix-tiled direction while narrowing the next experiment. Apple's Metal 4 inline-ML guidance says Apple silicon does not need an explicit software pipeline to overlap memory and compute, so Hopper/Blackwell warp-specialization should not be copied mechanically. Apple's current Metal Performance Primitives guide instead emphasizes a roofline model and multilevel SIMD-group/threadgroup tiling. FlashAttention-3's transferable idea is overlapping matrix work with softmax; FlashAttention-4 further shows that shared-memory traffic and exponential/rescale work can become dominant once matrix throughput improves. For Gem's already promoted BQ8/BK64 kernel, the next one-axis screens should therefore measure conditional online rescaling/exp reduction and a larger tile with explicit pipeline-resource receipts. BQ16 would exceed the target's 32 KiB threadgroup-memory limit with the current F32 output layout, so the bounded implemented geometry instead tests BK128 while retaining BQ8. CUDA-specific TMA, WGMMA, tensor memory, and 2-CTA mechanisms are not applicable to the M4 implementation.

Sources:

• Apple Metal 4 inline ML operations
• Apple Metal Performance Primitives programming guide
• Apple power-assertion guidance
• FlashAttention-3
• FlashAttention-4

Partial evidence is retained in artifacts/2026-08-26-server-gate-31b-bq8-131072.

2026-08-26: exact conditional rescaling is rejected; BK128 remains diagnostic

Paperbridge located FlashAttention-4's conditional online-softmax rescaling and the separate AMLA rescaling work. Gem first tested the conservative exact subset: skip exp(0) and the 4,096-float threadgroup output multiply only when a 64-token tile leaves the running maximum literally unchanged. Finite-data output was bit-identical to BQ8/BK64 and the adjacent Metal oracle passed 4,793/4,793. Seeded median deltas were +0.61% at 4K, +0.04% at 64K, and -1.01% at 128K. This is well below the 3% floor, so the candidate pipeline was removed. The result indicates that the threadgroup rescale is not the dominant standalone cost on this M4 Max path.

The operator profiler now supports --seed-prefix, which writes a deterministic sparse finite signal into Q/K/V. This prevents data-dependent optimizations from being judged only on the all-zero case, where the running maximum never changes after the first tile.

A second one-axis diagnostic doubles only the cache tile from BK64 to BK128 while retaining BQ8, F32 KV, exact causal online softmax, one-head threadgroup ownership, and all graph routing. It uses 20,544 bytes of dynamic threadgroup memory and achieves finite-data NMSE 1.09e-12 with deterministic repeats; normal and sanitizer suites pass 56/56 and the adjacent oracle again passes 4,793/4,793. An initial contaminated screen reported a 3.38% mean improvement at 128K, but cannot be promoted: the same 64K process drifted from 46.8 ms to 355.5 ms as macOS background work changed. BK128 is therefore retained only as a diagnostic pending a quiet-host interleaved rerun. Production remains BQ8/BK64.

Evidence is in artifacts/2026-08-26-bq8-conditional-rescale-screen and artifacts/2026-08-26-bq8-bk128-screen.

The clean retry is now frozen in scripts/run-bq8-bk128-gate.sh. It uses ten warmups and ten measured interleaved passes at 4K, 64K, and 128K, so BK64 and BK128 each run first equally often. It records binary/source digests plus macOS load, power, thermal, assertion, and process receipts before and after every context and rejects either a busy preflight or postflight. Operator-gate schema 2 now fails closed unless every receipt is on AC power and reports no thermal or performance warning; the analyzer validates those controls independently and embeds the digest/linkage receipts in its decision. It caps per-arm coefficient of variation and half-run drift at 3%, requires both 64K and 128K improvements to exceed 3% and observed noise, requires the paired 95% confidence lower bound to remain positive, and guards against a 4K regression beyond 2%. A pass only advances BK128 to the authentic evaluated-graph/logit gate; it cannot directly change selector routing. The normal and supported macOS ASan/UBSan suites now pass 60/60 including the analyzer's positive, regression, paired-noise, failed-power, failed-thermal, and failed-postflight fixtures. A live runner smoke correctly refused the current host at load 45.71 across 16 logical CPUs before launching any Metal work.

A later observation showed why load alone was insufficient: the one-minute value fell to 9.90 while YAMS indexing and media analysis still consumed about four CPU cores according to the process snapshot. The first schema-3 implementation gated that sum, but the macOS ps manual defines %CPU as a decaying average over up to one minute—not an instantaneous value. The corrected operator and server gates retain the sum only as provenance and gate a one-second top system-busy sample instead, with a 10% default ceiling. Operator postflight applies the same guard. This closes load-average lag without allowing stale process decay to reject an otherwise quiet host.

The adjacent DeepSeek-V4-enabled llama.cpp checkout at commit d8df12ebc4592b55dcecf97a32953623d031efdc independently converges on the same Apple Metal geometry for the relevant shape. Its FlashAttention constants use eight queries per threadgroup and 64 cache values per SIMD group (OP_FLASH_ATTN_EXT_NQPSG=8, OP_FLASH_ATTN_EXT_NCPSG=64), and its dispatcher selects eight SIMD groups when the head dimension is at least 512. Gemma 4 31B has a 512-element Q head, so Gem's production BQ8/BK64/8-SIMD-group shape matches all three choices. The adjacent code computes threadgroup memory from these constants and retains BK64 rather than increasing the cache tile. This is design-oracle evidence, not a Gem benchmark and not a runtime dependency; BK128 still receives the frozen local gate. If BK128 fails, the reference strengthens the decision to remove it rather than keep speculative compiled surface. The relevant reference locations are ggml-metal-impl.h:112-113 and ggml-metal-ops.cpp:2992-3105 in the adjacent checkout.

The live gem-metal-counter-info receipt reports the Apple9 limit as 32,768 bytes, matching Apple's current feature-set table. BK64 therefore uses 56.45% of the maximum memory available to one compute threadgroup and BK128 uses 62.70%. Apple's API describes this as a per-kernel threadgroup maximum, while maxTotalThreadsPerThreadgroup is a pipeline-specific dispatch limit; neither value states how many workgroups reside concurrently on a GPU core. The formal analyzer records the utilization and explicitly leaves occupancy unknown rather than deriving a false resident-workgroup count. Xcode GPU capture remains the occupancy source.

Sources:

• Apple maxThreadgroupMemoryLength
• Apple maxTotalThreadsPerThreadgroup
• Apple Metal feature-set tables

An explicit “enable fast math” experiment is closed before implementation. Gem compiles the Metal library with default options (options:nil). Apple documents fastMathEnabled as true by default and says that state maps to fast math mode and fast FP32 functions; the newer mathMode property replaces the deprecated Boolean. Therefore creating an options object merely to request fast math would not change the compiled axis. A future scalar-exponential experiment must compare a genuinely different implementation and demonstrate logit safety, not relabel the existing compiler default.

Sources:

• Apple fastMathEnabled
• Apple mathMode

A bounded quiet-window watcher then sampled the host every 30 seconds. All eleven observations failed the initial decaying-process-CPU guard, ranging from 186% to 362%; none began the Metal profiler, and no clean artifact directory was created. These values are retained as evidence of background activity, not described as instantaneous. The corrected short-window watcher must still find a valid interval before the BK128 decision can proceed.

2026-08-26: Metal startup compilation is measured and deprioritized

Gem's Metal backend previously compiled a single 131,435-byte source string and created every compute pipeline synchronously without phase attribution. gem_metal_open() now records device, source assembly, runtime library compilation, pipeline creation, command-queue creation, and total wall time. gem-metal-startup-profile emits those values as JSON without loading a GGUF or changing backend selection.

Three fresh-process diagnostic runs on the M4 Max measured 51.7--58.0 ms total. Device creation was 43.6--49.1 ms, runtime source compilation 1.79--1.99 ms, and all pipeline creation 4.27--5.52 ms. Three opens in one process fell from 60.7 ms to 4.3--4.5 ms after the first open. The host was not quiet enough for a promotion benchmark (one-second system busy was 15.0% and load was 30.68), but the result is sufficient for prioritization: even eliminating all observed compilation and pipeline creation would save only about 6--8 ms from fresh server readiness. Offline .metallib packaging and MTLBinaryArchive remain valid deployment improvements, not the next long-prompt performance axis.

Apple documents newLibraryWithURL:error: as loading a precompiled .metallib, whereas MTLBinaryArchive is a container for pipeline descriptors and associated compiled shader code. They address separate phases and should not be implemented as one unlabeled optimization. If hax readiness later becomes material, benchmark precompiled library loading first, then add a binary archive as an independent axis.

Sources:

• Apple precompiled Metal library loading
• Apple MTLBinaryArchive
• Apple binary-archive serialization

2026-08-26: BK128 now has an authentic-graph diagnostic route

The BK128 operator previously had no way to satisfy its required evaluated-graph/logit gate: it was callable only from gem-long-attention-profile. Diagnostic selector 23 now preserves selector 19's SWA, decode, batching, F32 KV, and rollback behavior while changing only the 31B global-prefill dispatch from BQ8/BK64 to BQ8/BK128. GEM_LONG_PREFIX_CANDIDATE=bk128 exposes a shared-prefix, interleaved selector-19/selector-23 comparison. Production still initializes selector 19 and no server option selects 23.

A deliberately non-promotable seeded 4K smoke exercised the complete graph. Both routes were bit-deterministic, retained argmax 100 and the same top-10 set, and measured correlation 0.999996027 with NMSE 2.31e-5. Timing varied too widely to interpret (individual route samples ranged from 3.12 to 5.36 seconds under heavy background work), so no performance artifact or claim was retained. This proves the next gate is executable; it does not bypass the quiet-host operator decision or the evaluated 64K prefix requirement.

The complete post-route suites pass 60/60 in both release and the supported macOS ASan/UBSan configuration. The added candidate-contract test ensures bk128 remains an accepted diagnostic mode while unknown modes fail before model loading.

The authentic graph step is now a frozen artifact gate rather than a manual harness command. run-bq8-bk128-graph-gate.sh independently re-analyzes the raw operator directory before loading the model, captures exact binary/source/model/configuration/linkage provenance, and samples the load, short-window CPU, AC-power, and thermal controls before, continuously during, and after the evaluated 64K prefix. Its analyzer rejects changed selector identities or model controls and requires correctness plus a gain above 3%, both arm CVs, and the paired df=2 95% lower confidence bound. Passing still only permits an internal selector-19 implementation change; it does not itself alter production.

A bounded post-verification watcher sampled the host twenty times from 05:38:51Z through 05:49:29Z on 2026-08-27. Load per logical CPU ranged from 0.274 to 1.792 and eventually remained well below the 1.0 ceiling, but the independent one-second CPU-busy sample never passed: it ranged from 13.01% to 29.65% against the 10% ceiling. The watcher therefore launched no Metal workload and created no 2026-08-26-bq8-bk128-clean directory. This is retained as host-state evidence, not a BK64/BK128 performance result.

2026-08-27: adjacent vector attention narrows the next decode design

The adjacent DeepSeek-V4 llama.cpp oracle remains at d8df12ebc4592b55dcecf97a32953623d031efdc. Its small-query Metal dispatcher selects the vector path below 20 query rows whenever the head size is divisible by 32. That path stripes the KV sequence across 32 workgroups, increases from one to four SIMD groups per workgroup as context grows, carries exact online-softmax (M,S,O) partials, and combines them in a second kernel.

It is not a drop-in Gemma 4 global-decode kernel. Although the adjacent source instantiates the correct symmetric DK=DV=512 F32 shape, it assigns only one query head per threadgroup. For global GQA-8 that rereads each shared K/V row for every query head. Gem's current score kernel instead stages each 512-float K row once and shares it across all eight query heads; its production "dim32" value name denotes a 32-element workgroup dimension tile, not DV=32. A future fused split-KV diagnostic should therefore combine the adjacent cross-workgroup online-softmax reduction with Gem's existing eight-head cache reuse, not copy the oracle kernel.

Before that larger axis, selector 20's already-implemented score-tile8 diagnostic must be resolved. Its runner previously lacked host controls and ran the heavyweight adjacent oracle before timing. Gate schema 2 now freezes 64K/128K ABBA controls, records pre/post load, one-second CPU, AC, and thermal receipts, aborts on changed postflight state, and moves native parity plus the adjacent oracle after all timing arms. The analyzer independently rejects invalid host receipts. This is gate hardening only; no selector or Metal kernel behavior changed and selector 19 remains production.

2026-08-26: primary-source review isolates the next 128K decode axis

Apple's 2026 Metal Performance Primitives guidance says its most optimized GEMMs generally read device memory directly and rely on sufficient occupancy for memory/compute overlap. That is useful context, but it does not reopen Gem's direct-read attention axis: the locally measured Family-9 direct-cache variants already regress 31--35%, and an online-softmax attention kernel is not a pure GEMM. Apple's advice to tune tile size against occupancy, edge quantization, and the number of GPU cores instead reinforces the frozen BK64/BK128 one-axis gate.

FlashAttention-3 and FlashAttention-4 do not supply a portable M4 implementation shortcut. Their main pipelines depend on NVIDIA Hopper/Blackwell asynchronous MMA, TMA, tensor memory, and warp-specialized scheduling. Gem already screened the one algorithmic sub-idea that transfers without that hardware--conditional online-softmax rescaling--and rejected it below the promotion floor. No FA3/FA4 mechanism is admitted without an Apple-specific trace showing the corresponding limiter.

The strongest applicable reference is Apple's current MLX two-pass vector GQA decode. Its first pass divides a long KV sequence among independent threadgroups, carries exact online-softmax partials, and has a duplication-free GQA variant that reuses each loaded K/V row across multiple query heads. A second kernel combines per-block maxima, sums, and outputs. This matches Gem's remaining post-split16 attribution: the P*V value side is already reduced, while the score/K pass is the next long-decode ceiling. MLX's duplication-free specializations stop at symmetric D64 and D128, while Gemma's global shape is symmetric DK=DV=512. Its algorithm and partial-state algebra remain useful, but its register/threadgroup-memory geometry is not a drop-in D512 kernel.

After the BK128 prefill decision and final production acceptance gate, the next admitted diagnostic is therefore a single fused split-KV GQA decode axis for only the ten 31B global layers. It should hold F32 KV, exact online softmax, selector 19's model graph, and all server behavior fixed; sweep only the number of sequence blocks at 64K and 128K. Required evidence is the existing threshold, tail, deterministic-logit, authentic whole-token, memory, and adjacent backend-op gates. KV quantization, sparse attention, paged layout, and M5/A19 TensorOps remain separate representation or hardware axes.

Sources:

• Apple Metal Performance Primitives Programming Guide
• Apple GPU performance and Metal System Trace workflow
• MLX Metal SDPA dispatcher
• MLX vector and two-pass GQA attention kernels
• FlashAttention-3 paper
• FlashAttention-4 paper

2026-08-27: current MLX GQA source confirms split-KV shape, not direct reuse

A fresh primary-source pass found that MLX 0.30.4 explicitly highlighted a much faster Metal vector GQA path for long context. Source inspected at d73eb752ef2e6288fd95b032c0bff0a15a4a9e93 makes the mechanism precise. For single-query GQA-8 with at least 8K keys, the dispatcher launches one threadgroup per KV head and sequence block, with one SIMD group per query head. The specialized first pass divides each block into contiguous token subchunks, lets one SIMD group compute multiple query heads per loaded K/V subchunk, merges their online-softmax state inside the threadgroup, and emits (max,sum,output) partials. A second kernel reduces the sequence-block partials. The Apple-family heuristic increases the block count with context, reaching 512 blocks at 64K and 1,024 beyond 64K on its current high-SIMD route.

This is stronger evidence for Gem's proposed fused split-KV design than the adjacent llama vector path because it preserves GQA reuse. It still cannot be copied or treated as a performance oracle: MLX admits the duplication-free path only at symmetric D64 and D128, whereas Gemma 4 global attention is symmetric DK=DV=512. MLX's instantiations process eight heads per loaded row for D64 and four heads at a time for D128; extending its G*HPT*V threadgroup partial array to D512 with multi-head reuse would exceed the M4 Max's 32 KiB threadgroup-memory limit. Gem should instead keep one D512 accumulator per SIMD group, stage each K/V tile once for all eight groups, and sweep only the sequence-split count after selector 20 is resolved. The source also validates a fail-closed requirement already in Gem's plan: the partial count must be a SIMD-width-compatible value because the reduction consumes 32 partials per chunk.

Apple's current profiler guidance also confirms that static threadgroup-memory limits are not an occupancy measurement. An Xcode Metal trace must correlate occupancy with limiter and bandwidth counters before choosing between larger K staging and more sequence splits. Apple Metal 4 cooperative tensors are useful for compute-bound matrix tiles, but the new neural-accelerator benefit is specific to M5/A19-class GPUs; it does not justify changing the M4 Max decode path.

Operationally, run-when-host-clean.sh now waits for three consecutive load, one-second CPU, AC, and thermal passes, logs every observation, and then execs the unchanged gate. It never kills or reprioritizes background processes. Deterministic fixtures prove streak reset, exact argument preservation, timeout-without-launch, and receipt non-overwrite behavior. All 63 release and 63 ASan/UBSan tests pass after the subsequent traffic- and ceiling-analyzer integration.

The parked fused design's traffic premise is now mechanically derived rather than hand-recorded. analyze-decode-attention-traffic.sh validates an exact-shape D512 profile and accounts for the raw score write, three normalization passes, and one complete normalized-score read by each of the sixteen dim32 output groups. Retained split16 profiles report 184,549,376 bytes at 64K and 369,098,752 bytes at 128K: 17.1875% additional logical traffic over the compulsory K+V reads. The analyzer labels this minimum logical traffic, not device traffic, because Apple GPU caches may satisfy some reads. It is non-promotional and explicitly requires Metal counters plus authentic whole-token A/B evidence.

Sources:

• MLX 0.30.4 release
• MLX two-pass vector/GQA kernel, inspected commit
• MLX SDPA dispatch and block heuristic, inspected commit
• Apple occupancy guidance
• Apple 2026 custom Metal ML operations

2026-08-27: current-main recheck adds a routing-boundary correctness gate

MLX current main still selects its duplication-free GQA first pass only for single-query GQA-8, no mask or sinks, symmetric D64/D128, and at least 8K keys. It also exposes a sequence-block override rounded to a multiple of 32 because the second pass reduces partials in SIMD-width chunks. These details reinforce Gem's proposed ownership and split-count constraints, but current MLX still has no fused D512 route and therefore does not answer the register-pressure question for Gemma 4 global attention.

A separate MLX report localizes deterministic bad generation to the exact point where a general two-pass GQA route activates. The report is not evidence that Gem's proposed kernel shares that bug, but it demonstrates that operator-level aggregate error alone is insufficient for a dispatch boundary. The parked Gem gate now requires exact 8,191/8,192/8,193-token CPU and forced-selector-19 differentials plus matching greedy multi-token continuations. This is a validation refinement only; no selector or Metal pipeline has been added.

Sources:

• MLX current Metal SDPA dispatch
• MLX issue 4022: two-pass GQA routing-boundary correctness failure
• MLX issue 3885: no fused D512 SDPA route

2026-08-27: measured whole-token ceiling admits one D512 fusion diagnostic

The retained selector-19 authentic stage profiles and exact-shape split16 operator profiles now feed a fail-closed ceiling analyzer rather than a hand calculation. Across the ten global layers, the current D512 attention route accounts for 33.11% of whole-token GPU time at 64K and 43.73% at 128K. Its score stage alone accounts for 17.54% and 22.02%. A candidate must therefore reduce the complete global-attention route by at least 9.06% at 64K and 6.86% at 128K to project a 3% whole-token improvement.

This is enough residual to justify one future fused split-KV diagnostic after the earlier gates close. It is not a speedup forecast: the candidate cannot eliminate compulsory K/V reads, dot-products, exponentials, or the second-pass merge. analyze-fused-split-kv-ceiling.sh rejects changed selectors, shapes, contexts, incomplete stage totals, weak sample coverage, and mismatched model profiles, records all four input digests, and marks its result non-promotional.

2026-08-27: Paperbridge audit finds D512 feasibility evidence, not a long-context oracle

A Paperbridge search surfaced Open-TQ-Metal and BaseRT as recent Apple-native inference references. BaseRT independently supports Gem's architecture choice--direct Metal, a zero-allocation decode loop, preallocated KV, specialized GEMV/GEMM, and operator fusion--but its published evaluation uses 128--2,048-token prompts on M4 Pro and different Gemma variants. It supplies no dense Gemma 4 31B 64K/128K comparison and is therefore architecture corroboration rather than a benchmark oracle.

Open-TQ-Metal is closer to the proposed decode kernel. Source inspected at 0f09466b7fde772a4876bf7bee3ccdeb34313304 instantiates D512 with 16 SIMD groups and 16 query plus 16 output floats per lane, showing that this accumulator shape compiles on Apple Silicon. However, it launches one threadgroup per query head and maps that head back to its KV head, so global GQA-8 rereads shared K/V eight times. It then combines SIMD-group outputs through a serialized barrier loop. Gem's planned one-threadgroup-per-KV-head ownership remains materially different.

The paper's Gemma performance table stops at 950 tokens, and the repository explicitly reports both int4 and int8 KV failure beyond roughly 950/1,000 Gemma tokens. Its appendix also lists GQA factors 16 and 32, conflicting with Gem's authenticated GGUF topology of 2 for sliding layers and 8 for global layers. Accordingly, neither its claimed Gemma speedups nor its KV quality generalize to this 64K/128K goal. Quantized KV remains a separate rejected axis; only the D512 register shape and online-softmax algebra inform the F32 candidate.

Sources:

• Open-TQ-Metal paper
• Open-TQ-Metal Gemma source
• Open-TQ-Metal Gemma limitations
• BaseRT paper
• BaseRT repository

2026-08-27: Metal execution guidance narrows the host-overhead question

Apple's current guidance still treats command-buffer frequency as a CPU/GPU synchronization tradeoff and recommends doing more work with fewer command buffers and encoders. Gem's authentic decode path already submits one command buffer and one compute encoder per token, but records 1,393 dispatches inside it. The remaining wall-minus-GPU gap therefore cannot be attributed to command buffer proliferation. It must be measured before changing ownership: likely contributors include encoding 1,393 dispatches, driver scheduling and synchronization, readback, and the single small gem_metal_batch control allocation. The allocation does contradict the roadmap's literal "generation performs no heap allocation" wording, but removing it is not yet a justified speed claim.

Apple's occupancy guidance likewise says that resource limits alone do not select a winning tile: register, threadgroup-memory, and cache pressure can each cap useful occupancy, while excessive occupancy can itself cause cache thrashing. The parked D512 fused split-KV diagnostic should therefore collect an M4 Max Metal trace/heat map for its candidate and selector-19 baseline before promoting a partition count. The 2026 Metal Performance Primitives tile advice is M5-specific and does not override measured M4 behavior.

Sources:

• Apple Metal command-buffer guidance
• Apple Metal Performance Shaders tuning hints
• Apple GPU occupancy guidance
• Apple GPU performance heat maps

2026-08-27: production-path timing rejects batch allocation as an optimization axis

gem-decode-profile now brackets the ordinary one-command-buffer/one-encoder execution separately from its deliberately fragmented stage-counter pass. Profiling is opt-in, and the recorded prepare, encode, submit/wait, cleanup, and logits-readback phases exactly partition the production wall time; GPU execution remains nested inside submit/wait.

An initial 31B/F32 seeded-KV sample at position 8,192 was intentionally retained as dirty-host, non-promotional attribution. Despite severe background interference making absolute GPU time nearly an order of magnitude slower than earlier clean samples, GPU time was 99.8977% of wall time. The complete wall-minus-GPU upper bound was 1.012 ms, while cleanup containing the batch free averaged 0.0053 ms. Removing the small control allocation therefore cannot plausibly clear the project's 3% whole-token floor. Kernel GPU time remains the correct optimization target. A clean-host 64K/128K sample may confirm the ratio, but is not required before rejecting allocation removal as a speed axis.

Evidence: artifacts/2026-08-27-decode-host-overhead/seeded-production-8192-dirty-host.json.

2026-08-27: fused split-KV CPU merge oracle landed without a Metal route

The parked D512 diagnostic now has an engine-owned CPU reference for online-softmax partial construction, exact pairwise merge, final normalization, and bounded scratch sizing. Tests cover non-power-of-two token and value dimensions, uneven tails, more partitions than tokens, empty partials, non-finite rejection, invalid empty-state encoding, and size overflow. The fixed P=32, QH=32, DV=512 layout is 2,105,344 bytes including maxima and denominators. No selector, Metal pipeline, workspace reservation, or production behavior was added; the earlier experiment ordering remains intact.

The independent gem-online-merge-check harness subsequently exercised the fixed D512/P32 algebra at the exact future routing boundary (8,191, 8,192, and 8,193 tokens) and the 65,536/131,072-token tails. All five cases were finite and bitwise repeat-deterministic with correlation at least 0.999999999999, NMSE no greater than 1.42e-11, and maximum absolute error below 1.64e-7 against a separate double-precision direct softmax. The digested audit is artifacts/2026-08-27-fused-split-kv-cpu-boundary/audit.json. This does not cover a Metal kernel, full-model logits, or greedy continuation.

The D512 compiled-resource schema is now fail-closed before a kernel exists. Its analyzer rejects changed tile/P/D/KV geometry, incorrect partial scratch, non-M4 device limits, an unlaunchable pipeline, changed dynamic threadgroup memory, and missing source/binary identities. The fixed first pass is 256 threads/eight SIMD groups and 16,384 dynamic bytes; the reducer is 256 threads and zero dynamic bytes. No fabricated production receipt was recorded—the checked fixture establishes only the future artifact contract, and an authentic receipt still requires compiled Metal pipelines.

The next correctness receipt is now equally constrained. Its analyzer consumes the existing CPU gate, future compiled-resource analysis, five raw D512 operator differentials, and three raw evaluated-prompt 31B model differentials. It rejects boundary/long-tail changes, route or selector drift, non-finite or nondeterministic output, weak correlation/NMSE, changed model identities, argmax/top-10 mismatch, fewer than eight continuation tokens, or any greedy-token divergence. Test fixtures prove those failure modes without fabricating an authentic candidate result. No Metal kernel, correctness claim, or promotion receipt exists yet.

2026-08-27: MLX 0.32.2 independently validates shared-K/V GQA ownership

MLX 0.32.2, released on August 25, includes PR 4077's specialized GQA-8 two-pass decode kernel. The first pass assigns contiguous token subchunks to SIMD groups, keeps multiple query heads in registers for each shared K/V stream, and merges their online-softmax partials through threadgroup memory. On the submitted M5 Pro measurements, this improved the isolated 32-query/4-KV-head D128 kernel by 1.07x, 1.09x, and 1.14x at 8K, 16K, and 32K. Independent M5 Max verification reported 1.16x, 1.08x, and 1.12x, with a 7.0% full-model decode gain at approximately 32K. These are not M4 Max or D512 performance claims, but they independently validate that avoiding redundant GQA K/V loads can survive wider Apple silicon and whole-model execution.

The implementation also sharpens Gem's resource constraint. MLX selects eight heads per SIMD group at D64 but only four at D128 because the larger accumulator spills at eight. It deliberately stops at D128 and leaves D512 unsupported. Gem therefore must not transplant that per-SIMD multi-head geometry. The parked candidate's one D512 query/output accumulator per SIMD group, eight SIMD groups per KV head, tile4 16 KiB shared K/V staging, and mandatory compiled-resource/occupancy receipt remain the defensible first experiment. The new release adds evidence for the ownership choice; it does not reorder the BK128, production-128K, selector-20, and fused-D512 gates.

Sources:

• MLX 0.32.2 release
• MLX PR 4077: read each K/V byte once in GQA-8 decode
• MLX current Metal SDPA dispatcher
• Apple threadgroup sizing guidance