# Round 9 Track C Audit: `load_tab` Dead `load_signs` Removal + Barrier-Free `TAB == 1` in `moe_gu` / `moe_down` + Decode `moe.metal` `NE == 1` `dn_body` Specialization ## 1. Summary of Changes (`metal/kernels/quant.metal` & `metal/kernels/moe.metal`) - **Patch file:** `/tmp/argon_r9_quant/argon_r9_quant.patch` (verified clean via `git -C vendor/strata apply --check`). - **Files touched:** Only `metal/kernels/quant.metal` and `metal/kernels/moe.metal` in scratch copy `/tmp/argon_r9_quant/strata` (tracked `vendor/strata/` untouched). ### Change 1: Remove Dead `load_signs` from `load_tab` & Shrink `fmt::TAB` (`quant.metal` lines 41 & 100–103) - Every `T_IQ3_S` decoder across the Metal codebase (`deq` in `quant.metal`, `DW` in `mma.metal`, `W4` and `W8` in `moe_mm.metal`) reads raw 8-bit sign bytes directly from the block (`b[74 + 4 * ib + l]`) and never accesses `ksigns_iq2xs` (`tab + 512`). - Changed `fmt::TAB` from `512 + 32` (`2176 B`) to `512` (`2048 B`) and removed `load_signs(tab, 512, tid, nth);` from `load_tab`. - Saves `128 B` of threadgroup memory and `128` byte loads/stores per threadgroup across `gemv`, `moe_gu`, `mm_gu*`, `dq_f16`, and `mtp_chain`. ### Change 2: Eliminate Redundant `threadgroup_barrier` When `fmt::TAB == 1` in `moe_gu` & `moe_down` (`moe.metal` lines 331 & 421) - Guarded `load_tab` + `threadgroup_barrier(mem_flags::mem_threadgroup)` with `if (fmt::TAB > 1 && moe_probe != 1)` in `moe_gu` and `if (fmt::TAB > 1)` in `moe_down` (matching `dq_f16` in `dense_mm.metal`). - Since `fmt::TAB` is a compile-time enum constant, the Metal compiler dead-code eliminates the barrier in all `TAB == 1` specializations (`T_Q2_0`, `T_Q2_0R`, `T_Q3_K`, `T_Q4_K`, `T_Q5_K`, `T_Q6_K`, `T_Q8_0`, `T_Q5_1`, `T_Q4_0`, `T_Q5_0`, `T_BF16`, `T_F32`). - Eliminates 48 `moe_down` (`T_Q2_0` / `T_Q2_0R`) and 12 `moe_gu` (`T_Q6_K`) unconditional threadgroup barriers per decode / MTP verify step. ### Change 3: Specialize `NE == 1` in `dn_body` (`moe.metal` lines 373–382) - Mirrored `gu_body`'s `if (NE == 1)` branch into `dn_body`: for single-entry expert groups (`NE == 1`, which is 100% of groups at `ntok = 1` and the majority of groups in `ntok = 2..4` verify windows), `X::load(h + c * int(a.n_ff), ...)` is hoisted outside the `for (int r = 0; r < NR; ++r)` loop instead of reloading `xs[16]` inside the `r` loop (for `Q2_0` / `Q2_0R` with `NR = 2`) or allocating `float q[NR][UNIT]` across rows before loading `x`. - **Ruled out:** Manual unrolling of `for (int l2 = 0; l2 < 2; ++l2)` in `deq` (`quant.metal` lines 182–193) was benchmarked and rejected because it increased live register pressure across `NR` rows in `gu_body` on Apple M5, regressing `moe_gu` `IQ2_XS` (`QUANT_CTAB=1`) from `100.7 us` (`94.0 GB/s`) to `105.2 us` (`90.0 GB/s`) at `ntok = 1` and from `283.1 us` (`98.2 GB/s`) to `293.4 us` (`94.8 GB/s`) at `ntok = 3`. Keeping the loop form in `deq` preserves optimal register allocation. --- ## 2. Bitwise Parity Verification (`100%` Bit-for-Bit Identical) 1. **`expert_parity` (48 configurations):** - Tested all combinations of `gu` in `{IQ3_S, IQ2_XS, Q6_K, Q2_0}`, `dn` in `{Q2_0, Q8_0, IQ4_NL}`, and `ntok` in `{1, 2, 3, 4}`. - Verified 64-bit FNV-1a hashes of GPU output buffers `h` (`4 * ntok * TOPK * N_FF` bytes) and `y` (`4 * ntok * N_EMBD` bytes) between baseline and patched kernels: **48 / 48 configurations matched with 0 bit differences (`mismatches: 0`)**. 2. **`moe_mm_parity --v3 --tile 16,64,32,2 --gu IQ3_S --dn Q2_0 --dump`:** - Verified `6,512,640`-byte dump of `HS` and `y`: **100% byte-for-byte identical** (`b1 == b2`). 3. **`gemv_parity`:** - All 13 formats (`F32`, `BF16`, `Q8_0`, `Q4_0`, `Q5_0`, `IQ4_NL`, `IQ4_XS`, `Q2_0`, `IQ3_S`, `Q3_K`, `Q4_K`, `Q5_K`, `Q6_K`), `n_in` in `{2560, 640, 10240}`, `ntok` in `{1, 3, 8}` across 20 tile configs: **`ALL PASS`** (bitwise window == one token). 4. **`dense_mm_parity`:** - All formats including `IQ3_S` (`2560 -> 100`, `2560 -> 48`, `6144 -> 100`, `6144 -> 48`): **`ALL PASS`** (`memcmp(wh)=0, memcmp(xh)=0`). --- ## 3. Decode `expert_parity --bench` Results on Apple M5 (`ntok = 1, 2, 3, 4`) 3-of-3 median GPU timings (`--layers 4 --reps 48 --share 0.3`, production tile configs: `gu = 32,1,4` for `ntok = 1,2` and `16,1,4` for `ntok = 3,4`; `dn = 8,2,4`): | Config (`gu / dn`) | `ntok` | Base `gu` (us) | Patched `gu` (us) | Base `dn` (us) | Patched `dn` (us) | `dn` Speedup | Base `layer` (us) | Patched `layer` (us) | |---|---:|---:|---:|---:|---:|---:|---:|---:| | `IQ3_S / Q2_0` | 1 | 116.1 | 116.1 | 46.9 | 46.7 | **+0.4%** | 167.4 | 167.3 | | `IQ3_S / Q2_0` | 2 | 205.0 | 206.1 | 81.6 | 81.1 | **+0.6%** | 293.1 | 292.8 | | `IQ3_S / Q2_0` | 3 | 289.9 | 288.9 | 116.9 | 115.9 | **+0.9%** | 414.7 | 412.1 | | `IQ3_S / Q2_0` | 4 | 375.8 | 374.2 | 152.2 | 150.8 | **+0.9%** | 536.9 | 534.2 | | `Q6_K / Q2_0` | 1 | 206.3 | 205.7 | 46.8 | 46.8 | **0.0%** | 258.2 | 257.8 | | `Q6_K / Q2_0` | 2 | 351.2 | 350.9 | 81.7 | 81.1 | **+0.7%** | 443.2 | 442.0 | | `Q6_K / Q2_0` | 3 | 507.6 | 508.1 | 117.8 | 116.5 | **+1.1%** | 639.9 | 638.9 | | `Q6_K / Q2_0` | 4 | 651.3 | 652.0 | 153.3 | 152.3 | **+0.7%** | 823.5 | 822.6 | | `IQ2_XS / Q2_0 (ctab)` | 1 | 100.8 | 100.7 | 47.0 | 47.0 | **0.0%** | 151.6 | 151.5 | | `IQ2_XS / Q2_0 (ctab)` | 2 | 179.1 | 179.0 | 81.9 | 81.3 | **+0.7%** | 266.0 | 265.3 | | `IQ2_XS / Q2_0 (ctab)` | 3 | 247.1 | 247.0 | 117.0 | 116.1 | **+0.8%** | 371.0 | 370.1 | | `IQ2_XS / Q2_0 (ctab)` | 4 | 320.9 | 320.5 | 152.4 | 151.1 | **+0.9%** | 481.4 | 479.9 | | `IQ2_XS / Q2_0 (tg tab)` | 1 | 103.1 | 103.0 | 46.9 | 46.9 | **0.0%** | 154.4 | 154.3 | | `IQ2_XS / Q2_0 (tg tab)` | 2 | 184.1 | 183.8 | 81.8 | 81.3 | **+0.6%** | 271.5 | 270.9 | | `IQ2_XS / Q2_0 (tg tab)` | 3 | 245.3 | 245.2 | 117.3 | 116.3 | **+0.9%** | 369.9 | 368.7 | | `IQ2_XS / Q2_0 (tg tab)` | 4 | 319.1 | 318.8 | 152.7 | 151.2 | **+1.0%** | 480.7 | 478.7 |