#228 QSA select on tensor cores (rocWMMA, opt-in, gfx1100) + AMD RX 7900 XTX rates row + verify profiler columns fix

Description

Three independent parts; each is separately reviewable. Built on PR #187's multi-query
block scorer (cherry-picked as `ed70478` in this branch — if you'd rather merge #187
first, the branch rebases onto it cleanly; `git merge-tree` shows no other overlap).

### 1. `qsa_block_scores_wmma`: the prompt scorer arm on tensor cores

The prompt path's QSA block scores (`score(qi,b) = sum_h relu(sum_i k_b[i] · q_qi,h[i])`)
now have a rocWMMA 16×16×16 FP16→FP32 implementation beside the warp kernel. gfx11 has no
`mma.sync`, and this is the tensor-core path that actually runs there — verified at the
instruction level: the gfx1100 code object executes 8 `v_wmma_f32_16x16x16_f16` (the 8
K-slices of the mma loop) with `block_scores_tail_kernel` and `f32_to_f16_poison_kernel`
alongside.

**Opt-in, with an honest fallback.** `STRATA_SELECT_WMMA=1` tries the arm first; the host
entry refuses (wrong geometry, not gfx1100, no rocWMMA) and the warp kernel keeps the
work. The arm that actually ran prints once (`strata select: prompt scorer arm = ...`) so
a fallback cannot masquerade as the fast path. **No other device changes behavior**;
every non-gfx1100 device keeps the warp kernel by construction.

**The arithmetic contract was derived before the kernel** (PORTING.md §18 in our rollout
repo, excerpt below): FP16 operands and an unspecified accumulation order put every
score within

```
B(qi,b) = 2^-9.5 · S(qi,b) + 8·2^-24 · |score_ref|,   S = the absolute product sum
```

of the float64 reference — and the tail block is bitwise the warp kernel's (the shared
tail kernel runs unchanged). No tolerance was moved at any point.

**Parity suite** (`src/kernels/qsa_select_wmma_parity.cpp`, in the test suite): two
references (a float64 oracle and a warp-order host model pinned to the observed FMA
contraction), five cases — ragged query/block tiles, whole tiles, ×8 magnitude headroom
(near FP16 overflow), degenerate `n_kv=0`, and one-hot/coordinate-coded adversarial
fixtures that expose lane-duplication/transpose/store errors at zero tolerance. Plus: the
write set is exactly the warp kernel's (canary-asserted), selection IDs agree, and a
negative control (+2.5·B perturbation) fails the same checker. Measured bound
utilisation 0.0–11.0% worst case.

**Bench** (`qsa_select_bench`, nq=256, median of 9 after 3 warmups) — scores only:

| context | warp | wmma | speedup |
| --- | ---: | ---: | ---: |
| 1K | 0.041 ms | 0.028 ms | 1.5× |
| 4K | 0.072 ms | 0.023 ms | 3.2× |
| 32K | 0.655 ms | 0.066 ms | 9.9× |
| 64K | 1.398 ms | 0.107 ms | 13.0× |
| 128K | 2.627 ms | 0.252 ms | 10.5× |

Selection cells differing vs the warp arm: 0–0.009%.

To keep the numbers honest at both levels: the bench table is the **scores kernel**; the
whole `qsa select` phase (scores + selection/topk + the fp16 conversion + the batch loop)
speeds up less — at 128K it drops 13,592 ms → 5,945 ms (~2.3×) on the measured e2e runs.
The kernel speedup is the arm; the phase speedup is what the prompt actually sees; the
end-to-end row below is what the request sees.

**End-to-end, paired** (frozen contract: 5 interleaved pairs × 2 slots per length,
order alternating per pair, the tuned AMD engine, prompt GPU timeline):

| 1K | 4K | 32K | 64K | 128K |
| ---: | ---: | ---: | ---: | ---: |
| **+0.2%** (paired, n=10; `qsa select` 1 ms -> 1 ms - flat: the phase is a sliver of the timeline) | **+0.9%** (paired, n=10/10 pairs, zero exclusions (26 failed attempts absorbed by the retry-until-valid policy, all declared); phase 23 -> 18 ms) | **-4.5%** (paired, n=10; phase 652 -> 266 ms. Within this box's interference noise at ~22 s cells: 6 clean pairs flat at 0.97-1.06, 4 pairs render-node-interference-inflated and flagged in the record, never dropped) | **+3.8%** (gate 3 holdout, paired ratio 0.964, n=10/10 pairs; phase 2,414 -> 746 ms) | **+8.7%** (paired ratio 0.913, n=10/10 pairs; warp median 97,242 ms vs 85,712 ms GPU timeline; phase 13,624 -> 5,991 ms) |

Promotion gate: >=5% at 128K on the prompt GPU timeline, decode protected to 3%.
**All three gates PASS -> PROMOTE (default flip earns its gates)** on the complete contract
(10 pairs per length, exclusions declared in the record):
- **Gate 1** (>=5% at 128K): **PASS at +8.7%** (paired ratio 0.913, n=10).
- **Gate 2** (decode <=3% regression): **PASS at +2.2%** (paired ratio 1.022, n=10). Decode is
  shared code between the arms, so the expectation is zero regression; the clean pairs read
  wmma +2-3% faster, the interference-dipped pairs scatter both ways. The decision statistic is
  the paired per-cell ratio median - the campaign's pre-declared drift-robust estimator, the one
  gates 1 and 3 already use. For full transparency the record also carries the ratio-of-arm-medians
  (0.937): on this bimodal decode distribution (5 clean / 5 interference-dipped cells) the arm
  medians land in the distribution gap and flip with sampling order, so that variant is reported
  but is not the decision statistic. The <=3% threshold itself never moved.
- **Gate 3** (64K holdout >=2.5%): **PASS at +3.6%** (paired ratio 0.964, n=10).

### 2. The AMD RX 7900 XTX (gfx1100) IQ3_S row in the rates tables

Prompt (tokens/s): collecting (the 1K-128K sweep record; 262K "-")
Output (tokens/s): collecting (the same runs)
(262K: "-", per the IQ3_S row's convention.)

Methodology follows `bench/results/2026-09-29-speed-0126`: one-shot `strata generate`
runs, greedy, 256 generated tokens, the tuned AMD engine config. Every cell is a row in
`artifacts/rocm/results/amd-rates-20260930.json` with its raw engine log kept.

### 3. `verify`: the profiler's stage columns stop lying (PORTING.md §19)

Two integer-arithmetic defects made `STRATA_VERIFY_PROFILE`'s report print ~1e14
ms/window for the "(gap)"/"head"/hc0 columns:

1. unsigned stamp differences wrapped to ~1.8e19 ns whenever a slot was unstamped;
2. `kProfPer = 32` left the hc-read second half's up-stamp at slot 32 — the next layer's
   stage-0 base stamp — corrupting the next layer's first delta.

`kProfPer` is now 33 and a delta counts only when both stamps are nonzero and monotonic.
Metrics-only: no computed model value changes. The GPU-side attribution of the residual
(the no-doorbell verify-window floor for native packs) is recorded in our rollout repo.

### What this is not

- Not a decode-path change (the decode scorer on tensor cores was measured slower and is
  not promoted).
- Not a claim for other GPUs: the arm's runtime gate matches its evidence (gfx1100
  measured). Wider RDNA3/4 support is a per-SKU measurement, not a flag flip.

Evidence lives in the xyzzing/strata-amd rollout repo: PORTING.md 18/18.1/19, `record-qsa-select-wmma-parity-*.json`, `record-qsa-select-e2e-*.json`, `record-prepr-merge-audit-*.json`, `record-d8-stamp-fix-*.json`, `ab-qsa-select-20260930.json`, and the D8 capture (baseline vs no-doorbell, `STRATA_VERIFY_DEVICE_PLAN=1`) that attributes the ~26 ms/window residual.

## Engine attribution (after the v0.1.29 merge)

The branch now includes `upstream/main` at v0.1.29 (merge `0975f72`); upstream's reviewed
implementation of the multi-query scorer (#187) supersedes the cherry-pick this PR was
originally stacked on. The two merge conflicts were comment-only in `qsa_select.cu`
(upstream's comments kept) and one informative comment in `CMakeLists.txt` (ours kept);
the merged tree builds 236/236.

The **measurements above were produced by the pre-merge, engine-0.1.28-based build** (the
campaign was mid-run on it); each record names its binary and engine. GPU parity
re-validation of the merged build is scheduled at the campaign's end and will be posted
here before this draft is marked ready.