Skip to content

hip: the gfx110x fp16 (rocWMMA) qsa-select scorer arm (STRATA_SELECT_WMMA) - #1298

Closed
xyzzing wants to merge 1 commit into
Niko1221:mainfrom
xyzzing:refile-select-fp16
Closed

xyzzing wants to merge 1 commit into
Niko1221:mainfrom
xyzzing:refile-select-fp16

Conversation

@xyzzing

@xyzzing xyzzing commented Oct 7, 2026

Copy link
Copy Markdown
Contributor

Per the offer in the #944 thread: the gfx110x fp16 (rocWMMA) qsa-select scorer arm, as the gfx110x branch of the existing STRATA_SELECT_WMMA opt-in, alongside the bf16 hand-rolled kernel that 0.1.40 ships.

What it is. One commit: the fp16 rocWMMA scorer kernel for gfx110x (rocwmma.hpp, 16x16x16, fp16 operands / fp32 accumulate) behind the same STRATA_SELECT_WMMA env, plus its parity harness src/kernels/qsa_select_wmma_parity.cpp (fp64 oracle + the bitwise warp host model, pre-declared bounds, negative controls). The bf16 kernel (gfx12 + the other gfx11 layouts) and the warp kernel stay the fallbacks everywhere else — the default path (env absent) is byte-identical to main; on cards without the rocWMMA headers the arm compiles out and the host entry refuses.

Why fp16 here. Paired A/B on gfx1100 (2026-10-06, same engine build, arms differing only in the select kernel): the fp16 arm's qsa-select phase ran 1.64x faster than the shipped bf16 arm at 128K (measured on the fork's 0.1.40 refresh line; the fork carries this exact kernel in production). The re-land on this base measured parity-neutral overall vs our previous production reads (x1.0009 on the canonical stream tail).

Gates on this base. Full HIP build green; qsa_select_wmma_parity all cases PASS (worst fp64-oracle bound utilisation 11.0%, one-hot/coordinate-coded adversarial case bitwise vs its host model, 0 differing selection ids, negative control fails as designed).

Prepared with an AI engineering agent (GLM-5.3 / Z.ai) under human direction; every number is from our own recorded runs.

…1.40 line

Measured 2026-10-06 on this card: the qsa-select phase runs 1934 ms (3.2% of
128K prefill) here vs 3176 ms (5.1%) on the bf16 arm upstream 0.1.40 ships
(92dde2f) - fp16 is 1.64x faster at the phase level, and that phase is the
entire -2.27% e2e delta of the 0.1.40 line vs production. This arm takes
gfx1100 first under STRATA_SELECT_WMMA=1 and prints the arm that ran; false
(no rocWMMA headers, another gfx110x geometry) falls through to the bf16
kernel, which keeps serving gfx115x and gfx12 unchanged. Parity harness
re-landed with it (fp64 oracle + the bitwise warp host model, PORTING.md 18).
Maddoxkkm added a commit to Maddoxkkm/Strata that referenced this pull request Oct 11, 2026
What the release brought for this box (Niko1221#1709 Windows HIP VRAM floor, the HIP stream-capture fix, the
Linux-only PCIe telemetry), the merge table with what was kept and why Niko1221#1439/Niko1221#1841/Niko1221#1656 went, the
measurements (0.1.42+local +5.6% at 16K; the four new PRs +4.4% on top), and the ranked shortlist for
the next cycle (Niko1221#1297 bisect, Niko1221#1909, Niko1221#1324, Niko1221#1167, Niko1221#1298/Niko1221#1151) with the gfx1100 reading list.
Maddoxkkm added a commit to Maddoxkkm/Strata that referenced this pull request Oct 11, 2026
…in (STRATA_SELECT_WMMA)

Conflicts were additive: upstream's FP32 tiled block for CUDA cards without tensor cores and the PR's AMD
tensor-core arm coexist, and the parity target is registered beside upstream's entries.
Maddoxkkm added a commit to Maddoxkkm/Strata that referenced this pull request Oct 11, 2026
…STRATA_SELECT_SIMT)

Additive conflicts: this arm, Niko1221#1298's rocWMMA arm and upstream's CUDA FP32 block all coexist behind their
own switches; the docs merge too. Applied for a standalone measure - inside batch 2 the SIMT arm recovered
most of that set's loss (836.8 vs 666.9 tok/s at 4K).
Maddoxkkm added a commit to Maddoxkkm/Strata that referenced this pull request Oct 11, 2026
…10.9%

Six PRs merged individually (Niko1221#1297 Niko1221#1909 Niko1221#1324 Niko1221#1167 Niko1221#1298 Niko1221#1151), none of which changes the default path
(the arms are opt-in and Niko1221#1297 needs a tuning table). Reading the issue tracker produced the real win:
issue Niko1221#1549 documents ring_default_slots() sizing the staging ring at 384 (1024 for a fused ring) when
g_pinned_share >= 0.9, and our STRATA_PREFILL_RING=96 override - left from tuning against a smaller budget
- was costing 16.5% at 4K and 10.9% at 16K. The override is now gone, so the engine's own rule applies.
Measured, with the LRU left off (it buys prefill but eats the resident budget: 11.9 -> 9.1 GiB).
Maddoxkkm added a commit to Maddoxkkm/Strata that referenced this pull request Oct 11, 2026
…nd the WMMA select scorer

Issue Niko1221#1712's trace showed the prompt chunks running the decode attention kernel on gfx1100 because
qsa_prompt refused. The kernel was compiled for this card all along (STRATA_PA_GFX11 layout); it needed
STRATA_HIP_WMMA=1. Measured: +11.5% at 4K, +22.7% at 16K; adding Niko1221#1298's select arm
(STRATA_SELECT_WMMA=1) takes it to +5.0% / +11.4% on top. Validated with the HIP parity test built for
the purpose: on gfx1100 the matrix-core kernel is at FP32 level (4.05e-06 vs FP64, the old kernel's
2.51e-06), agrees with the old kernel to 2.2e-06 of output scale, and the attention chunk is 4.1-6.2x
faster. Both switches are in the live config; the path needs --kv int8.
@Niko1221

Copy link
Copy Markdown
Owner

Thank you. The qsa-select phase is now covered by the exact thread scorer (#1943, bitwise equal scores, no rocWMMA dependency, 3.6-4x on the phase), and your own end-to-end number on this base is neutral, so we are not adding a rounded fp16 arm next to it. If the fp16 arm beats the exact scorer end to end on a gfx110x we would reconsider.

@xyzzing

xyzzing commented Oct 18, 2026

Copy link
Copy Markdown
Contributor Author

Measured on gfx1100 (build 0.1.45, 128K prompt, RX 7900 XTX): the fork's fp16 gfx110x arm runs the QSA select phase in 1,714 ms, the exact thread scorer in 4,320 ms, and the default radix in 9,859-10,440 ms. End to end prefill ties within noise (2,066 vs 2,055 tok/s) because the select is a small share of the 128K prompt; the 262K tier was not tested. Your condition to reconsider was an end-to-end win for the fp16 arm: the phase win is 2.5x but the total is a tie at 128K, so we accept the decision and close. The fork keeps the arm (carried since 700c026) and will re-test if the exact scorer's dispatch changes.

@xyzzing xyzzing closed this Oct 18, 2026
xyzzing added a commit to xyzzing/Strata that referenced this pull request Oct 18, 2026
…, 262K closed as config, HC fusion no-go

Adds "Follow-ups after the four windows (2026-10-18, later)" and the three adoptions to the
ledger's live set:
- STRATA_ROUTE_PRIOR=0.5: 1K +22.20% (5/5), 32K +25.56%, 128K +0.24% with -15.4% ms/round;
  CPU experts/layer 1.61 -> 0.44-0.73, PCIe 1.19 -> 0.20-0.89; battery 7/8=7/8=7/8.
- STRATA_GDN_CHUNKED=1: +5.85% prefill @32k (parity + 8K needle).
- STRATA_PF_GEMM=1 + STRATA_PF_PAD=1: +4.41% prefill @128k (same gate).
And the closed items: the 262K suspected regression is config (0.1.45 +3.96% vs 0.1.43.1 on the
same workload), the BELOW sweep is rejected (-2.48%/-3.96%), the gfx11 HC fusion is a documented
no-go, the exact scorer is 2.5x slower than the fp16 arm at the phase (closes Niko1221#1298's question),
and the four disposed PRs are closed with their numbers.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants