# Round 17 Track C: Decode `qsa_attn` Vectorized `char4` Loads / Hoisting & `gdn_step` `ib == 0` Readout Pruning + Loop Unrolling

## Summary of Changes

### 1. `metal/kernels/gdn.metal` (`gdn_step_t`)
- **`ib == 0` Readout Pruning:** Moved the readout reduction `for (int k = 0; k < IB; ++k) o += red2[k][j]`, scaling `o *= scale`, `simd_sum(o * o)`, and `part2[cb] = osq` write inside `if (ib == 0)` (`tid < 128`), keeping the surrounding `threadgroup_barrier(mem_flags::mem_threadgroup)` uniform across all threads. Threads `ib = 1, 2, 3` (`tid = 128..511`, 75% of the threadgroup) never read `o` or `part2` after `red2[ib][j] = o`, so pruning their redundant readout eliminates `1,536` shared-memory loads and `384` `simd_sum`s per token per value head (`73,728` shared-memory loads per GDN layer).
- **State Base Pointer Hoisting & `#pragma clang loop unroll(full)`:** Hoisted `device float* st_base = state + (i0 * hv + h) * S + j` with stride `st_stride = hv * S` across the initial state load and commit writeback, and added `#pragma clang loop unroll(full)` to all fixed-trip loops (`k < 3`, `k < 4`, `r < RPT`, `k < IB`, `k < 16; k += 2`).
- **Threadgroup Occupancy Preserved:** All `GDN_NTOK = 1..8` specializations compile with `maxTotalThreadsPerThreadgroup = 512` (`gdn_step`) on Apple M5.

### 2. `metal/kernels/qsa.metal` (`qsa_attn`, lines 327–398)
- **Vectorized `char4` Key Loads & Query Pointer Hoisting:** Replaced the 8 scalar `char` loads from `kq + rowi * QD + lane * 8` (which is 8-byte aligned since `QD = 256` and `lane * 8` is a multiple of 8) with two 32-bit `char4` loads (`k0`, `k1`), hoisted `threadgroup const float* q_lane = &qs[0][lane * 8]` outside the cell loop with stride `QD` per query head, and added `#pragma clang loop unroll(full)` on the `HPG = 12` and `i < 8` loops in both the score and value-accumulation phases.

---

## Verification & Bitwise Parity

1. **Standalone Test Suite (`/tmp/argon_r17_gdn/strata/build-metal`):**
   - `gdn_parity`: `ALL PASS` (all `GDN_NTOK = 1..8` pipelines report `max 512 threads per threadgroup`; `ntok = 1..4` verify/commit/partial-commit checks pass).
   - `gdn_prompt_parity`: `ALL PASS` (`T = 1..2048` single-pass vs TB=16 and `CB` scan checks pass).
   - `qsa_parity`: `ALL PASS` (`prep`, `score`, `top-k`, `attention + gate` pass).
   - `qsa_prompt_parity`: `ALL PASS` (`pos0 = 0..30000` `qsa_attn_mm` / `qsa_flash` / `qsa_prep` checks pass).
2. **Paired A/B Bitwise Parity (`/tmp/argon_r17_gdn/ab_bench`):**
   - `gdn_step`: `100% bitwise memcmp == 0` across `state`, `hist_out`, and `y` (including all 4 prepared `xprep` planes + per-16 sums when `plane = MAXT * 6144 = 49152`) for all `ntok in {1, 2, 3, 4}`, `commit in {0, 1}`, `plane in {0, 49152}`.
   - `qsa_attn`: `100% bitwise memcmp == 0` across `part` (`[n_tok][chunks][QH][QD + 2]`) for all `pos0 in {256, 2048, 8192}`, `ntok in {1, 2, 3, 4}`.

---

## Paired Interleaved A/B Benchmark Results (Apple M5)

### 1. `gdn_step` (36 layers rotated, 113 MB state, median of 9 interleaved runs)

| `ntok` | `commit` | `plane` | Baseline (us/layer) | Candidate (us/layer) | Speedup | Delta | Bitwise `memcmp` |
|---:|---:|---:|---:|---:|---:|---:|---|
| 1 | 0 | 0 | 28.78 | 27.73 | **1.04x** | -3.63% | `0 (PASS)` |
| 1 | 0 | 49152 | 28.94 | 28.62 | **1.01x** | -1.11% | `0 (PASS)` |
| 1 | 1 | 0 | 55.17 | 55.34 | **1.00x** | +0.32% | `0 (PASS)` |
| 1 | 1 | 49152 | 55.71 | 55.95 | **1.00x** | +0.44% | `0 (PASS)` |
| 2 | 0 | 0 | 37.05 | 33.69 | **1.10x** | -9.05% | `0 (PASS)` |
| 2 | 0 | 49152 | 37.66 | 34.58 | **1.09x** | -8.18% | `0 (PASS)` |
| 2 | 1 | 0 | 58.45 | 58.09 | **1.01x** | -0.62% | `0 (PASS)` |
| 2 | 1 | 49152 | 58.51 | 58.27 | **1.00x** | -0.42% | `0 (PASS)` |
| 3 | 0 | 0 | 42.89 | 38.62 | **1.11x** | -9.95% | `0 (PASS)` |
| 3 | 0 | 49152 | 42.62 | 39.20 | **1.09x** | -8.02% | `0 (PASS)` |
| 3 | 1 | 0 | 62.42 | 60.81 | **1.03x** | -2.58% | `0 (PASS)` |
| 3 | 1 | 49152 | 62.10 | 60.61 | **1.02x** | -2.39% | `0 (PASS)` |
| 4 | 0 | 0 | 48.47 | 44.66 | **1.09x** | -7.87% | `0 (PASS)` |
| 4 | 0 | 49152 | 49.05 | 45.28 | **1.08x** | -7.70% | `0 (PASS)` |
| 4 | 1 | 0 | 65.13 | 62.29 | **1.05x** | -4.36% | `0 (PASS)` |
| 4 | 1 | 49152 | 66.28 | 63.75 | **1.04x** | -3.81% | `0 (PASS)` |

### 2. `qsa_attn` (24 dispatches per timing batch, median of 11 interleaved runs)

| `pos0` | `ntok` | Baseline (us/layer) | Candidate (us/layer) | Speedup | Delta | Bitwise `memcmp` |
|---:|---:|---:|---:|---:|---:|---|
| 256 | 1 | 20.45 | 18.77 | **1.09x** | -8.21% | `0 (PASS)` |
| 256 | 2 | 23.16 | 21.55 | **1.07x** | -6.95% | `0 (PASS)` |
| 256 | 3 | 28.87 | 26.53 | **1.09x** | -8.13% | `0 (PASS)` |
| 256 | 4 | 40.03 | 35.46 | **1.13x** | -11.41% | `0 (PASS)` |
| 2048 | 1 | 64.69 | 59.60 | **1.09x** | -7.87% | `0 (PASS)` |
| 2048 | 2 | 115.09 | 106.63 | **1.08x** | -7.35% | `0 (PASS)` |
| 2048 | 3 | 170.40 | 158.91 | **1.07x** | -6.74% | `0 (PASS)` |
| 2048 | 4 | 222.10 | 205.59 | **1.08x** | -7.43% | `0 (PASS)` |
| 8192 | 1 | 64.82 | 59.42 | **1.09x** | -8.32% | `0 (PASS)` |
| 8192 | 2 | 114.44 | 106.70 | **1.07x** | -6.76% | `0 (PASS)` |
| 8192 | 3 | 171.15 | 159.12 | **1.08x** | -7.03% | `0 (PASS)` |
| 8192 | 4 | 223.25 | 206.02 | **1.08x** | -7.72% | `0 (PASS)` |

---

## Artifacts
- Unified patch: `/tmp/argon_r17_gdn/argon_r17_gdn.patch` (verified with `git -C vendor/strata apply --check /tmp/argon_r17_gdn/argon_r17_gdn.patch`)
- Paired A/B benchmark source & binary: `/tmp/argon_r17_gdn/ab_bench.cpp`, `/tmp/argon_r17_gdn/ab_bench`
- Scratch build & test tree: `/tmp/argon_r17_gdn/strata/build-metal`