Round 34 blitz: gemm_hc_down TN=96 tail (-5.0..-7.2%), grd_acc ROLL=true at KS=4 (-1.9..-4.6% on _r2), W8<T_IQ3_S>::block + half4 HS store + IQ2_XS sboff -> 432.78 tok/s 2K prefill, 42.27 tok/s decode, 25.04 tok/s @10.2K
1234567891011121314151617181920212223242526272829303132333435# Round 34 Track C Audit: `mm_gu3` `T_IQ3_S` & Super-Block / Loop-Invariant Hoisting on M5
## 1. Executive Summary- **Patch**: `/tmp/argon_r34_moemm/argon_r34_moemm.patch` (`git -C vendor/strata apply --check` passes with exit code `0`).- **Bitwise Parity**: 100% bitwise identical (`HS_memcmp=PASS YE_memcmp=PASS Y_memcmp=PASS`) across `T_IQ3_S`, `T_IQ2_S`, `T_IQ2_XXS`, and `T_IQ2_XS` at `T = 512, 1024, 2048`.- **Key Architectural Finding (M5 Cooperative Matrix VGPR Cliff in `mm_gu3`)**: - In `mm_gu3`, each SIMD lane simultaneously holds two FP32 destination cooperative tensors (`accg`, `accu`: `2 x 16 = 32` 32-bit registers) plus the FP16 right-input cooperative tensor (`rw`: `32` halves = `16` 32-bit registers) across `op.run(tA, rw, accg)` and `op.run(tA, rw, accu)`. - At `TM=16, TN=32, TK=32, SG=4` (`128` threads/threadgroup), `mm_gu3` sits at the exact M5 VGPR occupancy cliff: 1. Adding `F == T_IQ3_S` to `sb` (caching `float dg[4], du[4]` = 8 extra 32-bit VGPRs live across all 8 `ki` steps and 16 `op.run` calls) or even `half dg[4], du[4]` spills/drops occupancy, regressing `T_IQ3_S` by `+1.55%` to `+5.39%`. 2. Similarly, hoisting `l = kl >> 3`, `sh0`, `sh1`, `sh`, `th`, or per-lane `g0_l`/`u0_l` pointer increments across `op.run` extends their live ranges across the matrix accumulate instructions, regressing `T_IQ2_S` (`+13.8%` to `+67.3%`) and `T_IQ2_XXS` (`+25.2%` to `+28.8%`). Recomputing `l` and shifts inside `W8<F>::block` allows those temporaries to die before `op.run`.- **What Shipped (Proven Net Wins with Zero Occupancy Penalty)**: 1. **`W8<T_IQ3_S>::block` + `F == T_IQ3_S` Vectorized `half4` Epilogue Store**: - Added `W8<T_IQ3_S>::block(bl, k, d, tab, hh, a, b)` and dispatched it in `mm_gu3` with inline `rdh(rgw)` / `rdh(ruw)` and the already-live `hh` (zero extra registers live across `op.run`), eliminating redundant `row + (j_in >> 8) * 110`, `(k & 255)` masks, and `(k & 7) >> 2` recomputations on every inner-loop step. - Added `F == T_IQ3_S` to `mm_gu3`'s vectorized `*(device half4*)hs_p` epilogue store (line 1028), which executes after the K loop when `rw` and `accu` are dead. 2. **`F == T_IQ2_XS` Super-Block Offset (`sboff`) Reuse**: - Unlike `T_IQ2_S` and `T_IQ2_XXS`, `W8<T_IQ2_XS>::block` has a lower internal register footprint (no `l` or `qh` temporaries), leaving headroom to reuse `sboff = uint(kb >> 8) * fmt<F>::BB` across the `dg`/`du` preload and all 8 `ki` steps (`const uint boff = F == T_IQ2_XS ? sboff : ...`). 3. **`gdn_prompt.metal` Check**: - Verified `gdn_prompt_parity --ab` at `T = 2048` (`f32 = 0`); `gdn_pre` (`47.32 ms`), `gdn_scanr_4` (`187.46 ms`), and `gdn_post` (`34.93 ms`) are already at their register optimum (`268.76 ms` total across 36 layers), so `gdn_prompt.metal` was left untouched.
## 2. Quiet-GPU A/B Benchmark Results (`moe_mm_parity --v3 --dn Q2_0 --tile 16,32,32,4 --dtile 64,32,2 --gu-bench`)
| Format | `T` | Bitwise Parity (`HS / YE / Y`) | Baseline Med (Min) `us` | Candidate Med (Min) `us` | Med Diff | Min Diff ||---|---:|---|---:|---:|---:|---:|| `IQ3_S` | 512 | `PASS / PASS / PASS` | `9864.6` (`9814.7`) | `9844.9` (`9797.7`) | **-0.20%** | **-0.17%** || `IQ3_S` | 1024 | `PASS / PASS / PASS` | `16247.3` (`16210.3`) | `16144.0` (`16134.4`) | **-0.64%** | **-0.47%** || `IQ3_S` | 2048 | `PASS / PASS / PASS` | `27632.6` (`27613.1`) | `27535.3` (`27490.0`) | **-0.35%** | **-0.45%** || `IQ2_XS` | 512 | `PASS / PASS / PASS` | `7726.8` (`7696.6`) | `7728.9` (`7698.7`) | +0.03% | +0.03% || `IQ2_XS` | 1024 | `PASS / PASS / PASS` | `12758.6` (`12377.6`) | `12670.6` (`12271.4`) | **-0.69%** | **-0.86%** || `IQ2_XS` | 2048 | `PASS / PASS / PASS` | `22241.0` (`21177.6`) | `21962.3` (`20952.1`) | **-1.25%** | **-1.06%** || `IQ2_S` | 512 | `PASS / PASS / PASS` | `8846.8` (`8797.6`) | `8866.7` (`8796.8`) | +0.22% | **-0.01%** || `IQ2_S` | 1024 | `PASS / PASS / PASS` | `14642.3` (`14624.9`) | `14644.5` (`14629.9`) | +0.02% | +0.03% || `IQ2_S` | 2048 | `PASS / PASS / PASS` | `25110.4` (`25040.7`) | `25114.6` (`25019.1`) | +0.02% | **-0.09%** || `IQ2_XXS` | 512 | `PASS / PASS / PASS` | `7530.3` (`7521.9`) | `7543.7` (`7518.6`) | +0.18% | **-0.04%** || `IQ2_XXS` | 1024 | `PASS / PASS / PASS` | `12481.9` (`12457.5`) | `12493.6` (`12471.8`) | +0.09% | +0.11% || `IQ2_XXS` | 2048 | `PASS / PASS / PASS` | `21520.8` (`21184.5`) | `21586.4` (`21177.1`) | +0.30% | **-0.03%** |123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081# opus_r34_gr — Round 34 Track A: multi-row / coalesced `gr_down` (ntok 2..8) + `Window::gr_read` dispatch
Patch: `/tmp/opus_r34_gr/opus_r34_gr.patch` (`git -C vendor/strata apply --check` OK; touches `metal/kernels/gr.metal`,`metal/src/engine/window.cpp` (`Window::gr_read` dispatch line only; `gr_read_m` unchanged) and `metal/tests/fuse_parity.cpp`).Scratch tree `/tmp/opus_r34_gr/strata`, baseline kernels `/tmp/opus_r34_gr/base_kernels`, full logs`/tmp/opus_r34_gr/final_bench.txt` (+ `final_bench_run1.txt`), every experimental variant kept in`/tmp/opus_r34_gr/experiments_gr.metal` / `experiments_fuse_parity.cpp`. Device: Apple M5 10-core GPU, live server sharing the GPU(a few late-run samples were trashed by it; medians of 16 paired ratios used throughout).
## tl;dr
The big-idea candidates (j-outer token loop, NR = 4, larger RS, threadgroup-staged coalesced xn) are all bitwise PASS but**slower** at every ntok. The only win is small: keep the packed Q8_0 NR = 2 token loop *rolled* at every ntok for KS = 4, andlet `gr_read` switch to `_r2` from ntok 4 instead of 5.
## Shipped (bitwise PASS ntok 1..8 in `--gr-unroll`, default `fuse_parity` ALL PASS)
1. `gr.metal`: `grd_acc<F, NR, ROLL>`; `gr_down_f` sets `ROLL = F == Q8_0 && NR == 2 && KS == 4` (the 4_4_r2 / q8f_4_4_r2 down rows). Same fmas, same order, just no `unroll(full)` over tokens. KS = 8 (`8_8_r2`, the MTP head) measured +2..6% slower rolled, so it keeps the unrolled loop (now ±0.2%).2. `window.cpp` `Window::gr_read`: `r2 = ... ntok > (m5_gr ? 3 : 4)`. `4_4_r2` is bitwise `4_4` (same KS), so the switch is free for parity. `gr_read_m` (KS = 8) stays at `> 4`: 8_8 beats 8_8_r2 at ntok 4 (45.5 vs 46.7 µs).3. `fuse_parity --gr-unroll`: benches for the dispatch choice (`4_4` vs `4_4_r2`, `q8f_4_4` vs `q8f_4_4_r2`, `8_8` vs `8_8_r2`).
### Final paired timings (`final_bench.txt`, median of 16 pair ratios)
| kernel (base → new) | ntok | base µs | new µs | median Δ ||---|---|---|---|---|| 4_4_r2 → 4_4_r2 (rolled) | 4 | 41.34 | 39.55 | -4.4% || 4_4_r2 → 4_4_r2 | 5 | 49.50 | 48.18 | -2.7% || 4_4_r2 → 4_4_r2 | 6 | 58.43 | 57.44 | -1.7% || 4_4_r2 → 4_4_r2 | 7 / 8 | 68.14 / 79.77 | 68.14 / 79.79 | 0.0% (already rolled) || q8f_4_4_r2 → same | 4 / 5 / 6 | 41.24 / 49.51 / 58.96 | 39.96 / 48.48 / 57.69 | -3.0 / -2.1 / -1.9% || 8_8_r2, 4_4, q8f_4_4, 8_8 | 1..8 | | | ±0.2% (unchanged code) || **engine ntok 4: 4_4 → 4_4_r2 (rolled)** | 4 | 40.64 | 39.63 | **-2.5%** (run 1: -0.8%) || engine ntok 3: 4_4 → 4_4_r2 | 3 | 31.67 | 33.03 | +4.3% → stays 4_4 || engine ntok 2: 4_4 → 4_4_r2 | 2 | 23.84 | 23.82 | -0.3% (noise) → stays 4_4 || engine ntok 5..8: old r2 → rolled r2 | 5/6 | 49.50/58.43 | 48.18/57.44 | -2.7/-1.7% |
Net at the 3-draft verify (ntok 4, 97 calls): ~0.3-1.0 µs/call → ~0.03-0.1 ms per verify step. Small, but free and bitwise.
## Measured and rejected (all vs base `gr_down_q8_4_4` NR = 1 at the same ntok, median Δ; all bitwise PASS unless noted)
| variant | 2 | 3 | 4 | 5 | 6 | 7 | 8 ||---|---|---|---|---|---|---|---|| 4_4_r2 (old, unrolled ≤6) | -0.9 | +4.4 | +1.8 | -11.0 | -22.4 | -30.0 | -33.5 || 4_4_j (NR 1, j-outer) | +1.5 | +0.5 | +25.9 | +5.5 | +6.9 | +5.2 | +8.9 || 4_4_r2j (NR 2, j-outer, task idea #1) | +10.5 | +8.5 | +13.4 | +56.0 | +78.7 | +32.6 | +52.8 || 4_4_r4 (NR 4 packed, 81 TGs) | +32.6 | +21.2 | +17.5 | +28.6 | +28.7 | +9.9 | +0.0 || 4_4_r4j | +16.6 | +20.3 | +42.5 | +40.9 | +49.8 | +34.4 | +37.8 || 8_4 (RS 2) | +2.1 | +2.3 | +2.2 | -0.3 | -3.7 | -0.1 | -0.6 || 8_4_r2 | +2.4 | +8.6 | +5.7 | -7.9 | -21.7 | -30.1 | -34.2 || 8_4_r4j | +22.6 | +26.0 | +50.0 | +46.6 | +48.0 | +38.2 | +36.0 || 16_4 (RS 4) | +8.1 | +5.8 | +3.8 | -1.0 | -10.7 | -12.9 | -15.6 || 16_4_r2 | | +20.4 | +21.0 | 0.0 | -15.8 | | -30.3 || 32_4 (RS 8) | | +17.4 | +13.3 | +7.7 | -8.3 | | || 4_4_u2 (`unroll_count(2)` tokens) | | | +9.5 | +9.7 | +0.2 | | || 4_4_r2u2 | | | +16.2 | +27.2 | -14.1 | | || 4_4_st (per-simdgroup threadgroup-staged, coalesced xn) | | | +96.3 | | +49.9 | | +32.4 || 4_4_r2st (same, NR 2) | | | +30.5 | | -0.9 | | -9.0 || *probe* 4_4_pc (coalesced xn, WRONG math) | -1.9 | -15.1 | -0.1 | -11.2 | -22.0 | -29.9 | -22.9 || *probe* 4_4_r2pc (coalesced, WRONG math) | +4.3 | -2.5 | -7.5 | -14.4 | -23.9 | -34.3 | -36.2 |
Notes:- **j-outer (task idea #1) is bitwise identical but 5-80% slower.** Holding `s[NR][ntok][4]` live across the j loop costs more registers than the 2-float4-per-row it saves; the CSE-hoist theory doesn't pay off in practice (the compiler already widens the packed `char4` per use well enough — rolling the token loop is what helps).- **NR = 4 / more RS:** fewer, fatter simdgroups lose at ntok ≤ 6 (latency/occupancy: 81 TGs × 4 SGs isn't enough in flight). Bigger threadgroups (16_4, 32_4) only help where NR = 1 is already much worse than `_r2`.- **Coalesced xn:** the wrong-math probe bounds the gain: ≤ 0% at ntok 4 for NR 1 (confirms r33), -7.5% for NR 2 at ntok 4, -2..-3% vs `_r2` at 7-8. A real transpose through threadgroup memory (8 coalesced loads → padded 9-float4/unit tile → `simdgroup_barrier` → 8 tile reads per token) costs far more than that (+30..+96%), so staging is a dead end here. Cross-simdgroup cooperative staging needs RS > 1, which is itself slower (above), plus two threadgroup barriers per token.
## Where ntok 4 is stuck
`gr_down_q8_4_4` at ntok 4 is ~40 µs vs a ~23 µs DRAM floor. Coalescing doesn't move it (probe), halving xn traffic (NR 2) movesit ~2.5% only when rolled, and more/fatter simdgroups lose: it looks latency-bound per simdgroup (3 serial u-iterations ofDRAM weight load → L2 xn loads, with lanes 16-31 idle on the third). The bitwise constraint pins the unit→lane map (each lane's`acc` order and the `simd_sum` tree), so the remaining lever is a structurally different grid (e.g. split-K acrossthreadgroups with an ordered second-pass reduce) — out of this file/line scope.123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778# round 34 track b: hc-down tail narrowing
## decision
ship a **single 96-column tail tile**, not serial 64+16. patch: `/tmp/sol_r34_prompt/sol_r34_prompt.patch`. `git -C vendor/strata apply --check` exits 0. tracked vendor files were never edited.
production hc-down tiles improve roughly **5% at t=512, 6–7% at t=1024, 6–7% at t=1983/2048** in short paired a/b measurements on apple m5. all requested bitwise checks pass. hc-up and swiglu-up stay unchanged: their attempted guard/pointer optimization did not establish a win.
this computes nominal 128+128+96 = **352 rather than 384 columns (-8.33%)**, not the proposed 336 (-12.5%). minimizing nominal arithmetic alone lost badly. no full-model performance claim is made.
## design and scope
only original `metal/kernels/prompt.metal` lines 559–589 change. extracted the existing GEMM/register epilogue to one inline template; the existing entry point selects width 96 only when TN=128, hc=4, and the tile contains exactly 64 remaining down columns. all other shapes use the existing width. no ABI, host grid, buffer layout, allocations, arithmetic, or output guards change. no duplicate epilogue. tests already covered the required cases, so no test changes ship.
final visual review used `view_diff` and `view_code` over the full modified helper and entry point. searched for existing helper equivalents; none found.
## reproducibility
- vendor head at snapshot: `296b59d2f73c8f04a805df29c2fa97655adace78`; baseline is the actual working-tree file, not assumed identical to head.- isolated scratch tree: `/tmp/sol_r34_prompt/strata`, baseline committed in its local git repo.- exact baseline: `/tmp/sol_r34_prompt/baseline-prompt.metal`, sha256 `86adff9cfdbdf1af34ecfbdab0bef8b9e0668b3884a577469ac7bc58140f13a3`.- final source: `/tmp/sol_r34_prompt/candidate-final.metal`, sha256 `97a3391e86fc1fa0ba5089263e4d5bdb15e6ef4b4c2b5c375f3da5e859013921`.- evaluator: existing `vendor/strata/build-metal/comb_fold_parity`, sha256 `ba877816cf682d912ebdcd5fcd30e83dfcdcb9ff024b06ff58c8ad3d9547bf30`; C++ evaluator unchanged.- every run sets `STRATA_METAL_KERNELS=/tmp/sol_r34_prompt/strata/metal/kernels`; no math override.- no children spawned, no port 8080 access, no full-model CLI runs. every test invocation completed under 6 seconds, with timeouts 10 or 20 seconds.- baseline check: `comb_fold_parity 256 epi` passed before edits, log `baseline-256-epi.log`.
## correctness
all final runs exited 0 with bitwise PASS:
| command arguments | log basename | elapsed seconds ||---|---|---:|| `1024` | `final96-main1024.log` | 4.98 || `0 epi` | `final96-epi.log` | 0.94 || `256 hcup` | `final96-hcup.log` | 0.05 || `512 ab /tmp/sol_r34_prompt/baseline-prompt.metal` | `final96-full-ab-512.log` | 1.35 || `1024 ab /tmp/sol_r34_prompt/baseline-prompt.metal` | `final96-full-ab-1024.log` | 1.93 || `2048 ab /tmp/sol_r34_prompt/baseline-prompt.metal` | `final96-full-ab-2048.log` | 3.74 |
all logs reside under `/tmp/sol_r34_prompt/`. `0 epi` sweeps t=1,7,64,256,512,1024,1983,2048 and seven tile shapes against separate plain GEMM + epilogue. a/b additionally compares the old fused implementation directly, using a different random seed family. hcup tests mixed both bound and unbound.
## measured performance
existing a/b evaluator: 12 alternating samples, 8 repeated dispatches per sample, swaps output buffer placement; reports median paired ratios. lower is better. exploratory metric/evaluator was held fixed; confirmation repeats use fresh invocations, not new inputs.
final single-tail96 confirmation ratios (`final-confirm{0,1,2}-ab-T.log`):
| t | production tile | repeat 0 | repeat 1 | repeat 2 ||---:|---|---:|---:|---:|| 512 | 32_128_4 | 0.9489 | 0.9535 | 0.9518 || 1024 | 64_128_4 | 0.9290 | 0.9268 | 0.9097* || 1983 | 32_128_2 | 0.9367 | 0.9360 | 0.9281 || 2048 | 32_128_2 | 0.9325 | 0.9337 | 0.9408 |
representative first confirmation candidate/baseline ms: t512 0.2739/0.2890; t1024 0.6261/0.6739; t1983 1.1007/1.1754; t2048 1.1366/1.2152.
*timing contamination is visible, not suppressed: t1024 repeat 2 had 1.7585/1.8084 ms and ratio sample range [0.8752,1.1092], unlike normal ~0.63/0.67 ms. final full-a/b t2048 also gave a regression ratio 1.0411 [1.0041,1.1757] while unchanged hc-up jumped from ~1.13 to ~2.0 ms. this is evidence of changing shared GPU conditions, not proof of its precise cause. three immediate t2048 rechecks returned 0.9315, 0.9360, 0.9739; last recheck remained noisy [0.9429,1.0189]. retained in `final96-recheck{0,1,2}-ab-2048.log`. the win is supported by repeated stable short tests, but loaded-system/full-model confirmation remains the parent's job.
## rejected experiments
snapshots and logs are retained; nothing below ships.
1. `candidate-down-v1.metal`: fast interior helper plus serial 64+16 (or 16-only for TN64), including full-row guards. all `0 epi` checks passed, **but direct fused a/b at t512 failed bitwise** and latency ratios were 1.4482, 1.2021, 1.5143 at t512/1024/2048. do not infer decomposition bitwise safety from one seed family. the exact source of the bitwise difference was not isolated.2. `candidate-up-sw-v1.metal`: full-row guards and pointer/index hoisting. a/b parity passed. hc-up ratios 1.0042/1.0058/1.0041; swiglu-up 1.0541/1.0218/0.9906. rejected for no net win. raw `up-sw-v1-ab-T.log`.3. `candidate-down-v2.metal`: single width80 TN128 tail, width16 TN64 tail. parity sweep passed. production TN128 ratios 0.9472/0.9346/**1.1276**. SG2 width80 regressed.4. `candidate-down-v3.metal`: width80 only for SG4. confirmed SG4 improvements and SG2 neutrality; superseded by width96 to recover SG2 gains without differing tail widths.5. `candidate-down-v4.metal`: SG4 width80 / SG2 width96. t1983/2048 ratios 0.9268/0.9324. `candidate-down-v5.metal` switched both to width96, matching SG4 latency within noise, leading to the simpler final selection.
## application
```textgit -C vendor/strata apply /tmp/sol_r34_prompt/sol_r34_prompt.patch```
no unrelated vendor lines or files are in the patch. report and immutable scratch evidence are the deliverables; parent owns integration and any full-model measurement.
parent subsequently confirmed tracks a/c ran GPU parity concurrently during the late timing spike; parent plans an exclusive-GPU integration confirmation.