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](https://huggingface.co/google/gemma-4-12B-it/blob/main/config.json)
- [Google Gemma 4 31B configuration](https://huggingface.co/google/gemma-4-31B-it/blob/main/config.json)
- [ggml-org Gemma 4 12B GGUF](https://huggingface.co/ggml-org/gemma-4-12B-it-GGUF)
- [ggml-org Gemma 4 31B GGUF](https://huggingface.co/ggml-org/gemma-4-31B-it-GGUF)
- [Gemma 4 technical report](https://arxiv.org/abs/2607.02770)
- [Google Gemma 4 canonical chat template](https://huggingface.co/google/gemma-4-E2B-it/blob/main/chat_template.jinja)
- [Google Gemma 4 prompt and tool-control format](https://ai.google.dev/gemma/docs/core/prompt-formatting-gemma4)
- [llama.cpp Gemma 4 graph](https://github.com/ggml-org/llama.cpp/blob/master/src/models/gemma4.cpp)
- [FlashAttention: exact IO-aware tiled attention](https://arxiv.org/abs/2205.14135)
- [FlashAttention-2: improved work partitioning](https://arxiv.org/abs/2307.08691)
- [Online normalizer calculation for softmax](https://arxiv.org/abs/1805.02867)
- [Accelerating Large Language Model Decoding with Speculative Sampling](https://arxiv.org/abs/2302.01318)
- [Draft & Verify: Lossless Large Language Model Acceleration via Self-Speculative Decoding](https://arxiv.org/abs/2309.08168)
- [BaseRT: native Metal LLM inference](https://arxiv.org/abs/2607.00501)
- [Profiling LLM inference on Apple Silicon across quantization formats](https://arxiv.org/abs/2508.08531)
- [metal-llm-server Q3_0 reference implementation](https://github.com/githubpradeep/metal-llm-server)
- [RIGEL: M4 Max Metal tensor-path characterization](https://arxiv.org/abs/2606.12765)
- [Apple Metal Performance Primitives Programming Guide](https://developer.apple.com/download/files/Metal-Performance-Primitives-Programming-Guide.pdf)
- [Apple: Optimize Metal apps and games with GPU counters](https://developer.apple.com/videos/play/tech-talks/111373/)
- [Apple: Discover Metal debugging, profiling, and asset creation tools](https://developer.apple.com/videos/play/tech-talks/111374/)
- [llama.cpp Metal Q4_0 register dequantization](https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-metal/ggml-metal.metal)
- [Open-TQ-Metal: fused int4 attention on Apple Silicon](https://arxiv.org/abs/2604.16957)
- [When Quantization Is Free: Apple Metal int4 and SRFT study](https://arxiv.org/abs/2605.05699)

### 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](https://developer.apple.com/documentation/metal/creating-threads-and-threadgroups)
- [Apple: Porting Metal code to Apple silicon](https://developer.apple.com/documentation/apple-silicon/porting-your-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](https://arxiv.org/abs/2205.14135)
- [FlashAttention-2: parallelism and work partitioning](https://arxiv.org/abs/2307.08691)

### 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](https://developer.apple.com/documentation/metal/running-inline-ml-operations-in-a-shader-with-metal-4)
- [Apple: Metal sample code](https://developer.apple.com/metal/sample-code/)
- [Apple: Porting Metal code to Apple silicon](https://developer.apple.com/documentation/apple-silicon/porting-your-metal-code-to-apple-silicon)
- [Apple: Finding your Metal app's GPU occupancy](https://developer.apple.com/documentation/xcode/finding-your-metal-apps-gpu-occupancy)
- [Apple: Measuring the GPU's use of memory bandwidth](https://developer.apple.com/documentation/xcode/measuring-the-gpus-use-of-memory-bandwidth)
- [Apple: Discover new Metal profiling tools for M3 and A17 Pro](https://developer.apple.com/videos/play/tech-talks/111374/)
- [MLX PR 3023: Faster two-pass SDPA](https://github.com/ml-explore/mlx/pull/3023)
- [MLX Steel attention kernel](https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/kernels/steel/attn/kernels/steel_attention.h)

### 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](https://arxiv.org/abs/2604.16957)
- [Open-TQ-Metal HTML, method and limitations](https://arxiv.org/html/2604.16957v1)
- [MIT-licensed Gemma reference implementation](https://github.com/mutable-state-inc/gemma4metal)

### 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-P*V 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 P*V 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 P*V 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-P*V 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 P*V 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 P*V 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](https://github.com/ml-explore/mlx/pull/3023)
- [MLX PR 4077: Read each K/V byte once in GQA-8 decode](https://github.com/ml-explore/mlx/pull/4077)
- [MLX issue 4022: two-pass GQA routing-boundary correctness failure](https://github.com/ml-explore/mlx/issues/4022)
- [MLX issue 3837: long-context GQA bandwidth and block-count sweep](https://github.com/ml-explore/mlx/issues/3837)
- [Apple: Finding your Metal app's GPU occupancy](https://developer.apple.com/documentation/xcode/finding-your-metal-apps-gpu-occupancy)
- [FlashInfer recursive attention and KV sequence parallelism](https://docs.flashinfer.ai/tutorials/recursive_attention.html)
- [FlashInfer paper](https://arxiv.org/abs/2501.01005)
- [FlashDecoding++ paper](https://arxiv.org/abs/2311.01282)
- [Open-TQ-Metal paper](https://arxiv.org/abs/2604.16957)
- [llama.cpp Metal vector attention](https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-metal/ggml-metal.metal)

### 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](https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/scaled_dot_product_attention.cpp)
- [MLX head-dimension-512 fused-SDPA gap](https://github.com/ml-explore/mlx/issues/3885)

- [MLX long-context full-attention watchdog analysis](https://github.com/ml-explore/mlx/issues/3302)
- [BaseRT native Metal runtime paper](https://arxiv.org/abs/2607.00501)
- [BaseRT public repository and engine licensing boundary](https://github.com/basecompute/baseRT)
- [Open-TQ-Metal compressed-domain attention](https://arxiv.org/abs/2604.16957)
- [Apple Metal feature-set tables](https://developer.apple.com/metal/limits/)

### 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](https://developer.apple.com/documentation/xcode/capturing-a-metal-workload-programmatically)
- [Apple: Metal capture descriptor](https://developer.apple.com/documentation/metal/mtlcapturemanager/startcapture%28with%3A%29)
- [Apple: Metal performance heat maps](https://developer.apple.com/documentation/xcode/analyzing-apple-gpu-performance-using-performance-heatmaps-a17-m3/)

### 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](https://ai.google.dev/gemma/docs/core/model_card_4)
- [Apple Family-9 Metal shader performance guidance](https://developer.apple.com/videos/play/tech-talks/111373/)
- [Apple GPU occupancy guidance](https://developer.apple.com/documentation/xcode/finding-your-metal-apps-gpu-occupancy)
- [Apple SIMD-group matrix operations](https://developer.apple.com/videos/play/tech-talks/10858/)
- [FlashInfer attention-state composability](https://arxiv.org/abs/2501.01005)

### 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](https://developer.apple.com/videos/play/tech-talks/111373/)
- [Apple Family-9 GPU profiling workflow](https://developer.apple.com/videos/play/tech-talks/111374/)
- [Apple Metal Performance Primitives programming guide](https://developer.apple.com/download/files/Metal-Performance-Primitives-Programming-Guide.pdf)
- [FlashAttention-4](https://arxiv.org/abs/2603.05451)
- [FlashInfer](https://arxiv.org/abs/2501.01005)
- [MLX head-dimension-512 fused-SDPA gap](https://github.com/ml-explore/mlx/issues/3885)

### 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](https://huggingface.co/google/gemma-4-31B/blob/main/config.json)
- [Google Gemma 4 model overview](https://deepmind.google/models/gemma/gemma-4/)
- [Apple: Finding your Metal app's GPU occupancy](https://developer.apple.com/documentation/xcode/finding-your-metal-apps-gpu-occupancy)
- [Apple: Discover new Metal profiling tools for M3 and A17 Pro](https://developer.apple.com/videos/play/tech-talks/111374/)
- [Apple: Learn performance best practices for Metal shaders](https://developer.apple.com/videos/play/tech-talks/111373/)
- [Apple: Calculating threadgroup and grid sizes](https://developer.apple.com/documentation/metal/calculating-threadgroup-and-grid-sizes)
- [Apple: M5/A19 machine-learning acceleration](https://developer.apple.com/videos/play/tech-talks/111432/)
- [Apple: MPSGraph causal scaled-dot-product attention](https://developer.apple.com/documentation/metalperformanceshadersgraph/mpsgraphsdpadescriptor)

### 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](https://developer.apple.com/documentation/metal/compute-passes)
- [Apple: Metal Performance Shaders tuning hints](https://developer.apple.com/documentation/metalperformanceshaders/tuning-hints)
- [Apple: Family-9 Metal profiling workflow](https://developer.apple.com/videos/play/tech-talks/111374/)

### 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](https://arxiv.org/abs/2307.08691)
- [llama.cpp Metal flash attention](https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-metal/ggml-metal.metal)
- [MLX Metal SDPA dispatch](https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/scaled_dot_product_attention.cpp)
- [Rigel M4 Max Metal 4.1 study](https://arxiv.org/abs/2606.12765)
- [Open-TQ-Metal](https://arxiv.org/abs/2604.16957)

### 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](https://developer.apple.com/documentation/metal/running-inline-ml-operations-in-a-shader-with-metal-4)
- [Apple Metal Performance Primitives programming guide](https://developer.apple.com/download/files/Metal-Performance-Primitives-Programming-Guide.pdf)
- [Apple power-assertion guidance](https://developer.apple.com/library/archive/documentation/Performance/Conceptual/power_efficiency_guidelines_osx/PrioritizeWorkAtTheAppLevel.html)
- [FlashAttention-3](https://arxiv.org/abs/2407.08608)
- [FlashAttention-4](https://arxiv.org/abs/2603.05451)

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`](https://developer.apple.com/documentation/metal/mtldevice/maxthreadgroupmemorylength)
- [Apple `maxTotalThreadsPerThreadgroup`](https://developer.apple.com/documentation/metal/mtlcomputepipelinestate/maxtotalthreadsperthreadgroup)
- [Apple Metal feature-set tables](https://developer.apple.com/metal/feature-sets/)

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`](https://developer.apple.com/documentation/metal/mtlcompileoptions/fastmathenabled)
- [Apple `mathMode`](https://developer.apple.com/documentation/metal/mtlcompileoptions/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](https://developer.apple.com/documentation/metal/mtldevice/makelibrary%28url%3A%29?language=objc)
- [Apple `MTLBinaryArchive`](https://developer.apple.com/documentation/metal/mtlbinaryarchive)
- [Apple binary-archive serialization](https://developer.apple.com/documentation/metal/mtlbinaryarchive/3553928-serializetourl?language=objc)

### 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](https://developer.apple.com/download/files/Metal-Performance-Primitives-Programming-Guide.pdf)
- [Apple GPU performance and Metal System Trace workflow](https://developer.apple.com/videos/play/tech-talks/111432/)
- [MLX Metal SDPA dispatcher](https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/scaled_dot_product_attention.cpp)
- [MLX vector and two-pass GQA attention kernels](https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/kernels/sdpa_vector.h)
- [FlashAttention-3 paper](https://arxiv.org/abs/2407.08608)
- [FlashAttention-4 paper](https://arxiv.org/abs/2603.05451)

### 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 `exec`s 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](https://github.com/ml-explore/mlx/releases/tag/v0.30.4)
- [MLX two-pass vector/GQA kernel, inspected commit](https://github.com/ml-explore/mlx/blob/d73eb752ef2e6288fd95b032c0bff0a15a4a9e93/mlx/backend/metal/kernels/sdpa_vector.h)
- [MLX SDPA dispatch and block heuristic, inspected commit](https://github.com/ml-explore/mlx/blob/d73eb752ef2e6288fd95b032c0bff0a15a4a9e93/mlx/backend/metal/scaled_dot_product_attention.cpp)
- [Apple occupancy guidance](https://developer.apple.com/documentation/xcode/finding-your-metal-apps-gpu-occupancy)
- [Apple 2026 custom Metal ML operations](https://developer.apple.com/videos/play/wwdc2026/330/)

### 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](https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/scaled_dot_product_attention.cpp)
- [MLX issue 4022: two-pass GQA routing-boundary correctness failure](https://github.com/ml-explore/mlx/issues/4022)
- [MLX issue 3885: no fused D512 SDPA route](https://github.com/ml-explore/mlx/issues/3885)

### 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](https://arxiv.org/abs/2604.16957)
- [Open-TQ-Metal Gemma source](https://github.com/mutable-state-inc/gemma4metal/blob/0f09466b7fde772a4876bf7bee3ccdeb34313304/lib/turboquant.metal)
- [Open-TQ-Metal Gemma limitations](https://github.com/mutable-state-inc/gemma4metal)
- [BaseRT paper](https://arxiv.org/abs/2607.00501)
- [BaseRT repository](https://github.com/basecompute/baseRT)

### 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](https://developer.apple.com/library/archive/documentation/3DDrawing/Conceptual/MTLBestPracticesGuide/CommandBuffers.html)
- [Apple Metal Performance Shaders tuning hints](https://developer.apple.com/documentation/metalperformanceshaders/tuning-hints)
- [Apple GPU occupancy guidance](https://developer.apple.com/documentation/xcode/finding-your-metal-apps-gpu-occupancy)
- [Apple GPU performance heat maps](https://developer.apple.com/documentation/xcode/analyzing-apple-gpu-performance-using-performance-heatmaps-a17-m3/)

### 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](https://github.com/ml-explore/mlx/releases/tag/v0.32.2)
- [MLX PR 4077: read each K/V byte once in GQA-8 decode](https://github.com/ml-explore/mlx/pull/4077)
- [MLX current Metal SDPA dispatcher](https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/scaled_dot_product_attention.cpp)
- [Apple threadgroup sizing guidance](https://developer.apple.com/documentation/metal/calculating-threadgroup-and-grid-sizes)