[xyzzing]

Re-file of #888 — that PR was auto-closed on 2026-10-06 when this repository's main history was cleaned up (per the maintainer's note there, not a rejection). Same two commits, re-based onto the new main. Original discussion: #888.

What the re-base changed, explicitly:

• The old branch's first commit had swept in two local agent-onboarding files (.agentic/, CLAUDE.md); they are dropped here — they were never part of the PR.
• src/prefill/kernels.cu had grown the Aurora gfx11 HEAD/quad kernels in the same region the WY arm lands. The resolution keeps both: this branch adds the WY kernels + the STRATA_GDN_WY branch of the dispatch on top of the current Aurora code, with the composition untouched (the WY kernels write the same pre-norm output the cols kernels write; the following gdn_out_norm_kernel call is the existing one). Default path (no env, STRATA_GDN_HEAD default on this card) is byte-identical to main.
• Gates re-run on this base: full HIP build green; gdn_wy_check all checks PASS (0 mismatches, fp64 oracle 9/9 fixtures); gdn_wy_wmma_check PASS (worst bound utilisation 3.5%, negative control ok).

The measured verdict in the original description stands: on gfx1100 both arms lose to the serial cols_pipe kernel (fma ~144x slower, tensor-core folds ~6x slower per layer-chunk, engine prefill -39..-40%) — this ships opt-in default-off + parity-protected so the next attempt starts from evidence, not folklore.

---

What this is

The gated-delta-net recurrence (~7.4% of a 128K prefill, latency-bound: one serial token chain per head) rewritten in the chunk-parallel WY form (fla's recompute_w_u_fwd structure, reimplemented against this engine's exact contract): per 64-token sub-chunk one C×C unit-lower-triangular solve plus GEMM-shaped folds, one block per (value head, state column slice), state register-resident.

Two arms behind one opt-in env, STRATA_GDN_WY=1 (env absent = today's serial cols_pipe, dispatch byte-identical):

• fma arm — pure-fp32 folds.
• wmma arm (gfx1100, runtime-arch-guarded, falls back to fma with a printed arm line) — rocWMMA 16×16×16 fp16×fp16→f32 folds for the A-build (K̃ᵀK̃), the rhs/pq folds, P = QK̃ᵀ, O_intra, and the rank-BT state fold; the triangular solve stays fp32 scalar; ~61.4 KiB LDS via barrier-separated overlays.

The parity discipline (the substance)

Two HIP-visible CTest harnesses ship with the kernels:

• gdn_wy_check — all-heads fp64 oracle + a derived reassociation bound (T/WY+1)·2⁻¹⁴·(1+scale), NaN-aware metrics, negative controls, 9 fp64 host fixtures.
• gdn_wy_wmma_check — the fp16-class bound ⌈T/BT⌉·(2⁻⁶+2⁻¹⁴)·(1+scale) was derived and recorded before the wmma kernel was written (the precision-contract rule), with an explicit +2.5·B negative control. Green: worst oracle-bound utilisation 3.5% across T ∈ {1,37,64,65,200,4099} × decay modes.

Measured verdict on gfx1100 (7900 XTX, nightly ROCm): the WY form loses on this geometry

32k-token engine spot pairs, arms differing only in the env:

arm	gdn recurrence (32k)	result
serial cols_pipe	1,318 ms	baseline
WY fma folds	189,878 ms	144× slower — REJECT
WY wmma folds	~3,050 ms eq.	~6× slower per layer-chunk than cols_pipe (~127 ms vs ~13.7 ms); engine prefill −39..−40% reproducibly — REJECT

rocWMMA 16×16×16 at 64×64/64×32 shapes with 4-warp blocks does not amortize fragment traffic at S=128/HV=48. The kernels ship opt-in default-off and parity-protected so the next attempt (C=32 LDS-side restructure, GEMM fusion) starts from evidence instead of folklore.

gfx1100 / rocWMMA 7.1.1 notes baked into the code

• rocwmma headers must be included at global scope (their std:: names resolve against an enclosing namespace otherwise).
• store_matrix_sync static-asserts fragment/pointer type equality — no converting stores; fold results stay fp32 in LDS overlays and P16 is hand-converted through an f32 scratch region.
• An mma operand tile must never overlap the same GEMM's store target (the T=1 oracle-FAIL class this campaign hit and fixed).

Testing

gdn_wy_check and gdn_wy_wmma_check (unconditional CTest, HIP-visible scope); both green on this branch. Engine runs: STRATA_GDN_WY=1 prints gdn wy arm = wmma (gfx1100, opt-in) (or the fma line where rocwmma is absent); unset behaves byte-identically to main.

Related context from the same campaign: the model's rope scaling (--rope-scaling yarn) was validated to 512K prompt positions end-to-end on this hardware (needle recall at 256K/524K depths, 1.5-1.9K tok/s prefill via --kv-resident streaming) — decode at that depth is gated by MTP draft acceptance (0.80 → 0.06), a separate problem from this PR.

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

[Niko1221]

Thank you for the honest verdict in the description. As the WY arms lose to the serial kernel on gfx1100 (-39..-40% prefill by your own measurement) we are not merging them; #2299 now adds a chunked WMMA recurrence for gfx11 that wins on gfx1151. If you find a geometry where the WY form beats the serial kernel (your C=32 LDS idea), please re-open with the numbers.

[xyzzing]

Accepted. For the record on gfx1100: the WY arms lose to the serial kernel on this card (our own -39..-40% prefill measurement, as you quoted) and we never found a geometry where the WY form wins. The fork keeps the arm dormant and parity-protected; no supersession claim either way. #2299's chunked WMMA gfx11 path is in main, and we will measure STRATA_GDN_CHUNKED against the serial kernel on this card as our own cell, not as an argument for #1296. Closing.