[xyzzing]

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.

[Niko1221]

Thanks! We'd take the rocWMMA scorer as an opt-in (STRATA_SELECT_WMMA=1) but not as the default: the end-to-end gain is +8.7% at 128K and -4.5..+3.8% elsewhere. Please open a new PR with only the scorer, its parity test, the bench and the CMake detection - without .agentic/, AGENTS.md, CLAUDE.md, CI workflows, serve/telemetry.py, direct_file.cpp and math_constants.h. We have no gfx1100 card, so your numbers carry it. Closing this one.