#3226 [Engine] perf(hip): optimize GDN verify kernels with LDS-only barriers and enable batched commit under STRATA_GDN_SPLIT=1

本文

## Summary

Optimizes the Gated Delta Net (GDN) speculative-verify forward and commit kernels in `src/kernels/cuda/verify_kernels.cu` while preserving **100% bit-for-bit output and state parity**:

1. **Enable 2-launch batched `gdn_commit_layers` when `STRATA_GDN_SPLIT=1` (`verify_kernels.cu:1073`):**
   - `gdn_commit_layers` previously checked `(sp && sp[0] == '1')` (`STRATA_GDN_SPLIT=1`) and returned `false`, forcing `verify.cpp:2419` to fall back to **72 per-layer kernel launches** (`gdn_conv_commit` + `gdn_state_replay_kernel` across 36 GDN layers) instead of the 2-launch 36-layer batched commit (`gdn_conv_commit_layers_kernel` + `gdn_step_commit_layers_kernel`).
   - Because `gdn_step_commit_layers_kernel` and `gdn_state_replay_kernel` share the exact same `32 × 4` column-split state layout and `red[0..3][col_local]` reduction order, their committed states are **100% bitwise identical** (`gdn_parity` Section 6: 8 `n_keep` cases × 5 layers, 0 mismatches).
   - Notably, `STRATA_GDN_SPLIT=1` is enabled by default on `gfx1151` (`Strix Halo`) in `src/core/arch_defaults.cpp`, so `gfx1151` was silently falling back to 72 per-layer commit launches on every verify window.
   - `STRATA_GDN_COMMIT_SPLIT=0` and `STRATA_DBG_GDN=1` remain available as explicit per-layer commit overrides, and `gdn_parity.cpp` now fails if `gdn_commit_layers` is unexpectedly skipped when neither override is set.

2. **LDS-only workgroup barriers (`vk_sync_lds()`) in GDN verify kernels (`verify_kernels.cu:37`):**
   - On RDNA4 (`gfx1200`/`gfx1201`) and RDNA3 (`gfx11xx`/`gfx1030`), `__syncthreads()` compiles to a full workgroup memory fence (`s_wait_loadcnt 0x0; s_wait_storecnt 0x0; s_wait_dscnt 0x0; s_barrier_signal -1; s_barrier_wait -1; global_inv scope:SCOPE_SE` on `gfx12`), draining all in-flight global memory loads and invalidating the shader engine's L1 vector cache at every LDS synchronization point (4× per token in `gdn_step_norm_multi_kernel`, 5× per token in `gdn_step_split_kernel`, 3× per token in `gdn_step_commit_body`).
   - Since inter-thread communication in `gdn_conv_l2_body`, `gdn_step_commit_body`, `gdn_step_norm_multi_kernel`, `gdn_step_split_kernel`, `gdn_state_replay_kernel`, and `gdn_out_norm_kernel` goes strictly through `__shared__` memory (LDS), `vk_sync_lds()` uses `s_wait_dscnt 0x0; s_barrier_signal -1; s_barrier_wait -1` on `gfx12` and `s_waitcnt lgkmcnt(0); s_barrier` on `gfx11`/`gfx1030` (falling back to `__syncthreads()` on CUDA and other targets).

3. **Prologue state-load ordering and unified next-token prefetch in `gdn_step_norm_multi_kernel` (`verify_kernels.cu:328`):**
   - Issuing the 32 per-thread state loads (`s[0..31]`) before the initial `sk[0]`/`sq[0]` LDS store allows `vk_sync_lds()` at `t = 0` to overlap all 32 global state loads with the first barrier.
   - Unifying the `rg == 0` (`sk`) and `rg == 1` (`sq`) next-token prefetch under a single predicate (`load_kq`) overlaps the global load of token `t + 1`'s `k`/`q` and token `t`'s `v`/`z` with the 32-FMA `kv` and `o` loops while reducing register pressure by **4–7 VGPRs** across all four `gdn_step_norm_multi_kernel` instantiations.

4. **Double-buffered 3-barrier `gdn_step_split_kernel` (`verify_kernels.cu:417`):**
   - `gdn_step_split_kernel<true>` previously used single-buffered `sk[S]`, `sq[S]`, and `red[RG][GS_COLS]`, requiring **5 barriers per token** and stalling on `ht` global loads inside the barrier pair.
   - Double-buffering `sk[2][S]` and `sq[2][S]` with separate `red_kv[RG][GS_COLS]` and `red_o[RG][GS_COLS]` (LDS `1536 B -> 3072 B`, preserving full 16-wave / 4-CTA-per-WGP occupancy) reduces the loop to **3 `vk_sync_lds()` barriers per token** and overlaps next-token `k`/`q` global loads with the 32-FMA `kv` dot product.

---

## Root Cause & RED / GREEN Verification

On `origin/main` (`477ac13c`), running `gdn_parity` with `STRATA_GDN_SPLIT=1` skipped Section 6 (`gdn_commit_layers == per-layer commit`) because `gdn_commit_layers` returned `false`. With the tightened assertion in `src/kernels/gdn_parity.cpp`:

- **RED (`origin/main` `verify_kernels.cu` + `STRATA_GDN_SPLIT=1 ./gdn_parity`):**
  ```text
  gdn_ab_multi T=1..8 == fused_gdn_ab        bitwise (8 cases)
  gdn_step_norm_multi T=1..8 == single-token bitwise (47 cases)
  gdn_commit_layers == per-layer commit      *** NO (unexpectedly skipped) ***

  gdn: 1 failures
  ```
- **GREEN (this PR + `STRATA_GDN_SPLIT=1 ./gdn_parity`):**
  ```text
  gdn_ab_multi T=1..8 == fused_gdn_ab        bitwise (8 cases)
  gdn_step_norm_multi T=1..8 == single-token bitwise (47 cases)
  gdn_commit_layers == per-layer commit      bitwise (8 cases, 5 layers)

  gdn: 0 failures
  ```
- `STRATA_GDN_SPLIT=0 ./gdn_parity`, `STRATA_GDN_COMMIT_SPLIT=0 ./gdn_parity`, `./verify_batch_parity --selftest`, `./spec_verify_parity`, and `./arch_defaults_test` all pass with 0 failures.

---

## Code Object Resource Usage (`gfx1201`, `llvm-readobj --notes`)

| Kernel (`gfx1201`) | `origin/main` (VGPR / SGPR / LDS / Spills) | This PR (VGPR / SGPR / LDS / Spills) |
|---|---|---|
| `gdn_step_norm_multi_kernel<true, false>` | `89 / 32 / 6160 B / 0` | **`83 / 32 / 6160 B / 0`** (`-6 VGPRs`) |
| `gdn_step_norm_multi_kernel<true, true>` | `92 / 38 / 6160 B / 0` | **`88 / 38 / 6160 B / 0`** (`-4 VGPRs`) |
| `gdn_step_norm_multi_kernel<false, false>` | `168 / 38 / 6160 B / 0` | **`162 / 34 / 6160 B / 0`** (`-6 VGPRs, -4 SGPRs`) |
| `gdn_step_norm_multi_kernel<false, true>` | `173 / 38 / 6160 B / 0` | **`166 / 38 / 6160 B / 0`** (`-7 VGPRs`) |
| `gdn_step_split_kernel<true>` | `140 / 30 / 1536 B / 0` | **`142 / 30 / 3072 B / 0`** (`5 -> 3` barriers/tok) |
| `gdn_step_commit_layers_kernel` | `141 / 28 / 1024 B / 0` | **`141 / 28 / 1024 B / 0`** (LDS-only barriers) |

Compiled cleanly with zero errors across HIP targets (`gfx1100, gfx1101, gfx1102, gfx1103, gfx1151, gfx1200, gfx1201`) and CUDA targets (`sm_75, sm_89, sm_100, sm_120`).

---

## Benchmark Evidence (AMD Radeon AI PRO R9700 `gfx1201`, ROCm 10.2)

### 1. Interleaved 36-Layer GDN Microbenchmark (5 interleaved pairs)
Geometry matches Qwen3.8-Flash-Next 397B (`S=128, h_k=16, h_v=48, C=10240, value_dim=6144, MT=8, L=36` layers; 117.7 MB total GDN state across 36 layers so every window streams cold from VRAM beyond the 64 MB Infinity Cache). Each trial averages 50 full 36-layer verify windows after 10 warmup windows.

#### `STRATA_GDN_SPLIT=0` (default)
| Operation (36 layers total, us) | `origin/main` median (range) | This PR median (range) | Median Delta |
|---|---|---|---|
| `forward_unfused` (`T=1`) | `452.32` (`451.92..453.15`) | **`445.76`** (`445.07..445.82`) | **`-6.56 us` (`-1.5%`)** |
| `forward_unfused` (`T=2`) | `505.54` (`504.86..506.28`) | **`491.39`** (`490.32..491.66`) | **`-14.15 us` (`-2.8%`)** |
| `forward_unfused` (`T=4`) | `613.94` (`613.03..614.64`) | **`586.45`** (`585.36..587.67`) | **`-27.49 us` (`-4.5%`)** |
| `forward_unfused` (`T=5`) | `667.81` (`667.07..668.67`) | **`634.13`** (`633.48..635.22`) | **`-33.68 us` (`-5.0%`)** |
| `forward_qfuse` (`T=1`) | `342.04` (`341.81..342.26`) | **`337.62`** (`337.37..337.64`) | **`-4.42 us` (`-1.3%`)** |
| `forward_qfuse` (`T=2`) | `401.96` (`401.82..402.47`) | **`389.40`** (`387.85..390.58`) | **`-12.56 us` (`-3.1%`)** |
| `forward_qfuse` (`T=4`) | `523.24` (`522.12..523.72`) | **`497.17`** (`495.51..498.18`) | **`-26.07 us` (`-5.0%`)** |
| `forward_qfuse` (`T=5`) | `580.32` (`578.97..580.56`) | **`557.36`** (`555.50..559.85`) | **`-22.96 us` (`-4.0%`)** |
| `commit_per_layer` (`nk=5`) | `768.05` (`768.05..768.84`) | **`759.76`** (`758.62..759.97`) | **`-8.29 us` (`-1.1%`)** |
| `commit_batched` (`nk=5`) | `422.14` (`421.90..422.33`) | **`420.01`** (`419.88..420.10`) | **`-2.13 us` (`-0.5%`)** |

#### `STRATA_GDN_SPLIT=1` (also `gfx1151` default in `arch_defaults.cpp`)
| Operation (36 layers total, us) | `origin/main` median (range) | This PR median (range) | Median Delta |
|---|---|---|---|
| `forward_unfused` (`T=1`) | `452.37` (`452.02..452.97`) | **`446.02`** (`444.74..447.13`) | **`-6.35 us` (`-1.4%`)** |
| `forward_unfused` (`T=2`) | `593.98` (`593.03..594.65`) | **`581.03`** (`580.32..581.24`) | **`-12.95 us` (`-2.2%`)** |
| `forward_unfused` (`T=4`) | `695.61` (`694.78..696.42`) | **`657.51`** (`656.41..658.53`) | **`-38.10 us` (`-5.5%`)** |
| `forward_unfused` (`T=5`) | `746.94` (`746.28..748.58`) | **`698.87`** (`697.13..700.95`) | **`-48.07 us` (`-6.4%`)** |
| `forward_qfuse` (`T=1`) | `341.94` (`341.64..342.27`) | **`337.60`** (`337.59..337.75`) | **`-4.34 us` (`-1.3%`)** |
| `forward_qfuse` (`T=2`) | `483.98` (`482.91..484.31`) | **`470.88`** (`470.36..471.35`) | **`-13.10 us` (`-2.7%`)** |
| `forward_qfuse` (`T=4`) | `584.12` (`583.46..585.01`) | **`548.27`** (`546.97..548.80`) | **`-35.85 us` (`-6.1%`)** |
| `forward_qfuse` (`T=5`) | `641.35` (`639.26..642.32`) | **`591.03`** (`589.98..591.54`) | **`-50.32 us` (`-7.8%`)** |
| `Effective 36L commit` (`nk=1`) | `677.49` (`674.59..678.83`, 72-launch fallback) | **`456.33`** (`456.09..456.87`, 2-launch batched) | **`-221.16 us` (`-32.6%`, `1.48x`)** |
| `Effective 36L commit` (`nk=2`) | `710.57` (`710.15..711.48`, 72-launch fallback) | **`451.44`** (`450.61..452.58`, 2-launch batched) | **`-259.13 us` (`-36.5%`, `1.57x`)** |
| `Effective 36L commit` (`nk=3`) | `725.54` (`724.37..730.42`, 72-launch fallback) | **`448.58`** (`443.12..452.33`, 2-launch batched) | **`-276.96 us` (`-38.2%`, `1.62x`)** |
| `Effective 36L commit` (`nk=4`) | `749.85` (`748.71..750.47`, 72-launch fallback) | **`415.29`** (`415.17..418.40`, 2-launch batched) | **`-334.56 us` (`-44.6%`, `1.81x`)** |
| `Effective 36L commit` (`nk=5`) | `771.88` (`770.62..778.58`, 72-launch fallback) | **`420.15`** (`419.96..420.41`, 2-launch batched) | **`-351.73 us` (`-45.6%`, `1.84x`)** |

### 2. End-to-End 397B `--serve` Decode A/B (`gfx1201`)
Workload: `Swift-Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S` (`--spec 4 --spec-min-p 0.5 --mtp .../rt --pcie-frac 0.30 --adapt-every 0 --kv int8 --prefill 4096 --prompt-cache 0`, 512 prompt tokens, 256 generated tokens, 1 warmup + 3 measured trials per arm, 15s cooldown between arms). All arms use `STRATA_HIPBLASLT_TUNING=tools/hip/gfx1201-hipblaslt-100500.txt STRATA_DENSE_MMQ=1 STRATA_PF_FUSED=1 STRATA_HC_Q8=1 STRATA_HIP_WMMA=1`.

| Arm | Median Decode (tok/s) | Median Decode (ms) | Per-Trial Decode (tok/s) | Exact Token Match (768/768) |
|---|---:|---:|---|---|
| `origin/main` (`STRATA_GDN_SPLIT=0, STRATA_QFUSE=0`) | `91.33` | `2803.1` | `[91.61, 91.33, 91.18]` | `True` (ref) |
| **This PR** (`STRATA_GDN_SPLIT=0, STRATA_QFUSE=0`) | **`91.97`** (`+0.7%`) | **`2783.6`** (`-19.5 ms`) | `[92.13, 91.97, 91.36]` | `True` |
| **This PR** (`STRATA_GDN_SPLIT=0, STRATA_QFUSE=1`) | **`92.12`** (`+0.9%`) | **`2779.0`** (`-24.1 ms`) | `[92.23, 92.12, 91.51]` | `True` |
| **This PR** (`STRATA_GDN_SPLIT=1, STRATA_QFUSE=0`) | **`91.66`** | **`2792.8`** | `[91.76, 91.66, 91.53]` | `True` |

*(Note: Because the 36-layer GDN recurrence step accounts for ~0.65 ms of a ~25.0 ms 397B verify window on R9700 when experts stream over PCIe/CPU, the -34 us to -352 us kernel-level savings translate to a modest ~0.4–0.9% end-to-end wall-time improvement on 397B with overlapping trial ranges on a 3-trial sample; the primary isolated evidence is the 5-pair interleaved 36-layer GDN microbenchmark above.)*

---

## What This Does NOT Fix / Scope
- Does not change default environment switches in `src/core/arch_defaults.cpp` (`STRATA_GDN_SPLIT` and `STRATA_QFUSE` remain opt-in on `gfx12` and default-on on `gfx1151`).
- Does not modify GDN prompt prefill kernels (`src/kernels/cuda/gdn.cu`) or MTP / S2 grouped expert kernels.
- On `gfx11xx` (`gfx1151` Strix Halo) and CUDA targets (`sm_75, sm_89, sm_100, sm_120`), compilation is verified; target-device runtime validation on `gfx1151` and CUDA hardware was not run on this `gfx1201` host.

## Reproduction Commands
```bash
cmake --build build-copilot-hip --target gdn_parity verify_batch_parity spec_verif