Make KVarN usable on ROCm RDNA3 (toward #122: D256 race, portable prefill, f32 WMMA prototype)

[raufaser]

Overview

Toward #122 (KVarN on ROCm is unusable: wrong results plus prefill far behind). This branch makes KVarN work on RDNA3 step by step, validated on gfx1100 (RX 7900 XTX, ROCm 7.14) with Qwen3.6-27B-Q5_K_S and kvarn6:

• D256 k00-combine race fix in fattn-mma-f16.cuh (tile_Q reused across k00 iterations, barrier only fired for np > 1), plus inter-tile sync. Enables the RDNA D256 route.

• HIP prompt-prefill (nq > 16) routes through portable-native direct-record attention, since RDNA WMMA VKQ accumulates in fp16 (~3e-4/call) and compounds over depth into a KLD collapse. Decode stays on WMMA. Opt out with GGML_KVARN_AMD_PROMPT_PORTABLE=0.

• Portable prefill optimizations (warp-shuffle reduction, fp32 rescale, QB=8 queries per block, hoisted token resolve): pp4096 107 to 291 t/s at identical accuracy.

• Accuracy ladder harness (GGML_KVARN_TEST_NKV_LADDER_ONLY): GPU-native vs CPU-materialized error over D128/256/512 and n_kv 256 to 8192.

• f32 WMMA accumulator prototype for DV=128/256 prefill (16-wide A, single wmma_f32 per K step, persistent f32 acc, mirroring DV=80/112), plus the VKQ_C entry-count fix it required. Still gated behind PORTABLE=0, default serving path unchanged.

• ub>64 fix: the three whole-tile dst_final_meta stores in flash_attn_ext_f16_process_tile were gated !is_kvarn_kv, so KVarN whole-tile blocks never published their (max, rowsum). Split tiles got meta from the stream-k fixup, but that skips boundary-aligned tiles, so those rows kept zero meta and the tail merge silently discarded their correct body values. Symptom: KVarN WMMA prompt path clean at ub<=64, garbage at ub>=96. Removing the gate makes whole-tile blocks publish their meta like dense FA. Validated: ub512 KLD on 4B / 27B / 35B-MoE at portable parity (4B bit-exact); 32k KLD 2.118 -> 0.029.

Results: ladder RMSE ~1e-5; 32k KLD 0.022 vs portable 0.023 (was 2.14 on the fp16 path), same-top 97.0%; prefill pp4096 286 (portable) vs 497 (f32-WMMA); decode ~33.5 both; reference q8_0 at 804.

Base: v0.4.7

Retargeted to v0.4.7 per maintainer request. The branch merges v0.4.7 (D=64 KV-head support, etc.) with this work; the only merge conflict was fattn-kvarn-portable.cuh, resolved by porting v0.4.7's D=64 support (RECORD_DIM, the v_original_domain original->rotated V Hadamard, D=64 dispatch) onto the portable QB-batching + warp-shuffle kernel.

• D=64 is CUDA-only. v0.4.7 fail-closes the KVarN D=64 route on HIP/ROCm (fattn-kvarn-dispatch.cu: D64 stripped from capabilities on non-CUDA backends "until each backend has independent runtime qualification"). So on this gfx1100 build D64 is never dispatched; the D=64 kernel here is compile-validated and mirrors v0.4.7's, but runtime qualification on HIP is out of scope. D128/256/512 are the validated ROCm paths.

Additional information

• Progress toward Bug Report: kvarn6 KV-quantization causes massive performance regression on ROCm (RX 7900 XTX) #122 (update posted there during the work).

• The f32 approach follows https://github.com/stew675/llama-cpp-rdna-boosts (question thread; their answer pointed at the DV=80/112 fp32-PV tiles as reference).

• Validate with: ctest KVarN subset, ladder env above, llama-perplexity KLD flow and llama-bench -p 4096 from AGENTS.md.

• Known, kept out of scope: the f32-WMMA default flip (still behind PORTABLE=0); two flaky AVs under long KLD runs share one fault offset with dumps saved, zero occurrences on portable, PORTABLE=1 is the fallback. (The ub>=96 serving divergence previously listed here is now fixed by the ub>64 commit above.)

Requirements

• I have read and agree with the contributing guidelines

• AI usage disclosure: YES. An AI coding assistant (Muse Spark supervised by the author) wrote the tile-spec and count changes plus test-harness knobs, ran the builds, ladder/KLD/bench measurements on the author's RX 7900 XTX, and drafted this description. All numbers above are measured, not estimated; the author reviewed the diff before submitting.

Update 2026-09-09: F32-WMMA is now the default + CUDA build fix (2 new commits)

This supersedes the "f32-WMMA default flip (still behind PORTABLE=0)" note above.

• bad6f86c1 Ôö£├ÂÔö£├ºÔö£├é F32-WMMA is now the default for HIP KVarN prompt-prefill

  (was: portable-native by default, WMMA opt-in). The ub-geometry fix plus the

  DV=128/256 fp32 accumulators make WMMA both fast and exact, validated at the

  serving default (-ub 512, gfx1100, kvarn6, ngl99): 4B KLD 0.006822 (bit-exact

  vs fp16 base), 27B 0.004814, 35B-MoE 0.004093; 32k KLD 0.029 (was 2.118);

  ladder D128/256/512 ~1e-5. Portable-native remains for unsupported shapes and

  as an explicit opt-in (GGML_KVARN_AMD_PROMPT_PORTABLE=1). Decode (nq<=16)

  unchanged. HIP-gated; CUDA/Vulkan untouched.

• edc1d2f64 Ôö£├ÂÔö£├ºÔö£├é CUDA build fix (pre-existing, found by the thermis RTX 4090

  build): the GGML_KVARN_PORTABLE_ATTRS debug print used HIP-only

  func-attribute APIs, breaking nvcc. Now HIP API on HIP / CUDA API on CUDA

  (debug-only, no behavior change). **This PR as-merged did not build on CUDA;

  this fix is required.**

• CUDA gate (thermis RTX 4090, GGML_CUDA=ON): full build clean; KVarN/FA ctest

  (kvarn|fattn-route|cuda-fattn) 14/14 pass, incl. test-kvarn. The flip is

  compiled out on CUDA (HIP-gated) Ôö£├ÂÔö£├ºÔö£├é expected no-op, confirmed.

Community validation needed: gfx1151 (Strix Halo) / gfx1201 (RDNA4)

We only have gfx1100 (RX 7900 XTX) plus CUDA/4090. If you run **gfx1151 or

gfx1201**, please validate this PR on your hardware and report back:

• Build clean (GGML_HIP=ON, GPU_TARGETS set for your arch at initial configure).

• The ub512 KLD trio (see "Validate with" above) Ôö£├ÂÔö£├ºÔö£├é especially whether F32-WMMA

  as-default is accurate at your serving -ub (state the -ub value).

• The accuracy ladder (GGML_KVARN_TEST_NKV_LADDER_ONLY=1).

• ctest -R "kvarn|fattn-route".

• RDNA4 note: D64 KVarN is CUDA-only by design (fail-closed on HIP); D128/256/512

  are the expected ROCm paths. Thank you!

Update 2026-09-09 (pm): merged v0.4.7 "complete optimized D64", conflict resolved

• origin/v0.4.7 advanced (e1f6d6fe6, complete optimized D64) after this PR's

  earlier v0.4.7 merge, conflicting in portable.cuh. Resolved by merging

  (cc566c899) and taking upstream's complete D64 rewrite.

• Our interim D64 hand-merge + portable QB-batching/warp-shuffle are superseded

  (upstream covers D64 natively; portable is fallback-only post-flip, so peak

  fallback perf is acceptable; re-propose batching as a follow-up if it matters).

  Our CUDA func-attributes fix proved unnecessary (upstream removed the HIP-only

  debug block).

• Preserved through the merge: the F32-WMMA default flip and the ub-geometry fix

  (both auto-merged). One 3-line follow-up (37ce9e337) drops the now-dead

  batched-portable call (fallback uses standard portable ├ö├ç├ correct, unbatched).

• Validation on the merged tree (gfx1100, default WMMA): ub512 trio identical to

  pre-merge (4B 0.006822 bit-exact, 27B 0.004814, 35B 0.004093); key-target build

  clean. PR is MERGEABLE/CLEAN.

---

Follow-up: validated RDNA D256 tiles + portable shared OOB fix (570cb2c)

• RDNA table carries the validated D256 WMMA tile configs (256/320/512/576 keys, KLD-clean on gfx1100) and the AMD WMMA device guard is back at the largest selectable head (576), so KVarN D256 shapes use the fast WMMA path.
• Portable kernel shared scratch sized by head dim (reduction[D]/transform[D]): the D64/D256/D512 V-domain transform indexed full-head dims into RECORD_DIM-sized arrays (OOB for D256/D512). Latent on HIP (always rotated domain) but reachable on CUDA portable+MIXED shapes.
• Validation (gfx1100, kvarn6, ub512, defaults): 27B 0.0046/95.1% via prompt-generic-mma with zero kernel traps (was 10.8 garbage); 35B-MoE 0.0037/97.1%; 4B 0.0072/96.9% unchanged. CUDA 4090: 27B 0.0033, KVarN/FA ctest 16/16, backend-ops MUL_MAT + TOP_K clean.
• Upstream v0.4.7 tip merged (8b59e71).

Update 2026-09-11: review round 2 addressed (force-pushed as f7f3350)

Commits on top of the previous pushed head (bebc08f):

• 69c616d - validated RDNA (256,256,64) tile + portable shared OOB fix (message corrected: only the live row is listed, guard back at 256).
• 3c690bc - round 2 part 1: unreachable tile rows (320/512/576) reverted to base, fp32 tile scoping, 2 committed tests, tile boundaries, feature matrix.
• c5f2c02 + f7f3350 - round 2 part 2: domain_ok admits MIXED at all dims (required enabling fix - without it the portable path rejects every D256/D512+MIXED dispatch and is unreachable), fp32 tile scoped to KVarN via mma_tile_sizes_kvarn (dense keeps primary tiles on all arches), RDNA3/3.5 vs RDNA4 matrix rows split, RDNA4 fail-closed.

Validation on the final tree (clean rebuild from f7f3350, gfx1100 RX 7900 XTX, kvarn6, ngl99, -ub 512):

• llama-bench -p 4096, 27B Q5_K_S: pp4096 649.6 t/s (kvarn6) vs 937.8 t/s (dense-f16 reference); tg128 34.6 both. The (256,256,64) live row shows no spill collapse at long prompt.
• ub512 KLD trio: 4B 0.00624/97.8%, 27B 0.00571/94.5%, 35B-MoE 0.00386/97.5%; 4B dense-f16 0.00688/96.7%. (Note: the earlier 0.006822 figure for 4B on this branch predates the clean rebuild; 0.00624 is the clean-rebuild number, reproduced bit-exact on repeat.)
• New tests on HIP: test_kvarn_d256_prompt_tail_regression OK (route-pinned, prompt_prefill=1); test_native_flash_attention_portable_original_v D256 rmse 5.4e-4 OK, D512 rmse 5.1e-4 OK. After them the run still stops at the known pre-existing head-wide store-route abort (gfx1100 device-specific, identical with and without this branch).
• CUDA gate on the final tree (RTX 4090 sm_89, clean rebuild from f7f3350): full build clean; ctest kvarn|fattn-route|cuda-fattn 16/16 pass incl. test-kvarn; 27B Q5_K_S kvarn6 KLD 0.0042-0.0044/97.5-97.8% on the default prompt-generic-mma route (round-1 default was 0.0062/96.3% against its own base; portable-forced 0.0046/96.3%); dense f16 parity round-1 vs round-2 bit-identical (mean KLD -0.0, max 5.5e-5, 100% same-top); backend-ops MUL_MAT+TOP_K 2190 cases CUDA0 OK.

Two honest caveats:

• Tail-test fail-acceptance is unmet single-shot: an exhaustive fail-demo sweep on the gate-restored (OOB) tree (nq 32-256 x n_kv 256-32768, HIP + CUDA) never trips; the trigger needs incremental multi-ubatch cache-fill dynamics (serving KLD caught the original at 2.118). Serving KLD stays the acceptance unless you want an incremental-fill harness as follow-up.
• The D512+tail abort (GGML_KVARN_TEST_NKV_LADDER_ONLY=1 GGML_KVARN_TEST_LADDER_TAIL=128, D512/256+) is pre-existing fail-closed, identical pre/post - filing as a separate issue.

[Anbeeld]

Please change the target branch to v0.4.7, and ensure changes are compatible with it.

[Anbeeld]

Requesting changes, one blocking item on RDNA4 plus four smaller asks. Reviewed at 37ce9e337 against base 8b59e71ba.

1. RDNA4 gets the half2 tiles the fp32 work was meant to replace (blocking)

fattn-kvarn-route-policy.h:63-64 raises the head-dim limit for everything classified GGML_CUDA_FATTN_KVARN_AMD_RDNA_WMMA, and that class is amd_wmma_available(cc), so RDNA3, RDNA3.5 and RDNA4 (common.cuh:347-349, hip.h:231-237). The fp32 accumulator tiles that justify WMMA by default are compiled only inside #ifdef RDNA3 (fattn-mma-f16.cuh:1123-1191; new <128>/<256> at :1151-1166). RDNA4 falls through to the #else primary at :1167-1191 with T_C_VKQ = tile<16,16,half2, DATA_LAYOUT_I_MAJOR_SCRAMBLED> and the AMD_WMMA_AVAILABLE || AMD_MFMA_AVAILABLE VKQ_C sizing at :1312-1313. On gfx1200/gfx1201, D256 prompt prefill therefore now defaults to fp16 PV accumulation with an fp32 scale only, which is the configuration your 32k KLD measurement rejects (2.118 before the fp32 fix), while the previous behaviour was portable-native. The comment at fattn-kvarn-dispatch.cu:1238-1240 also states the fp32 property as family-wide.

Please gate the raised limit on the arches that actually compile the fp32 tiles (the D64 path already does this per-backend fail-closed clearing at fattn-kvarn-dispatch.cu:88-98), or port the fp32 tiles to RDNA4 before admitting 256 there, and correct that comment.

2. The D128 fp32 tile also changes dense HIP attention

mma_tile_sizes is shared, and dense HIP selects MMA for Q->ne[0] <= 128 (fattn.cu:693). Head_dim 128 dense on RDNA3/RDNA3.5 (Llama-3-8B, Mistral, Qwen class) moves from the 32-wide half2 tile to this float tile, and VKQ_C grows from 4 half2 entries to 8 float entries (fattn-mma-f16.cuh:1308-1311). Your ladder, KLD and llama-bench -p 4096 evidence is KVarN-only. Can you add dense-HIP gfx1100 evidence for D128 (accuracy plus prefill and decode throughput), or scope the fp32 tile to the KVarN path?

3. The regression guard for this geometry does not run on HIP

test_native_flash_attention_prefill_route_parity(), called by default at tests/test-kvarn.cpp:5621, covers D256 prefill with 1536/2560-candidate exact tails. It returns early at :4262-4265 when specialized_generic_mma && !original_v_domain, and original_v_domain is backend == CUDA && generic_mma (fattn-kvarn-route-policy.h:173), which is true on HIP with KVarN matrix capability. That is the ub>=96 tail geometry this PR fixes, untested on the backend whose default it flips. The new ladder is env-gated and passes exact_tail_tokens = ladder_tail, default 0 (:4334-4337), so it carries no tail either. Please cover D256 WMMA prefill with an attached tail on HIP (the ladder already has GGML_KVARN_TEST_LADDER_TAIL), or explain why the skip is required.

4. The portable opt-in still pays for the WMMA pass

ggml_cuda_flash_attn_ext_mma_kvarn at fattn-kvarn-dispatch.cu:1229 launches the WMMA kernel before the new HIP block at :1237-1256 can divert to portable, so with GGML_KVARN_AMD_PROMPT_PORTABLE=1 prompt prefill runs both and discards the first. D256 prompts never reached that launch before this PR. Moving the opt-in check ahead of the probe would remove the duplicate prompt pass.

5. Document the new knob

GGML_KVARN_AMD_PROMPT_PORTABLE appears only at its read site (fattn-kvarn-dispatch.cu:1247), while comparable knobs are documented (docs/beellama-args.md:47, docs/beellama-features.md:201). Its accepted value is strict (atoi(...) == 1, so =true silently does nothing), unlike the force_portable check in the same function. Please document the exact accepted values and align the boolean form.

What I verified and did not: this was a read-only review; no build or test ran, and there is no ROCm toolchain on this machine, so the AMD paths are neither compiled nor executed here. Your gfx1100 ladder, KLD and bench numbers are taken as reported. The merge with the current base tip is clean, but 53a68d3c3 and 8b59e71ba landed after e1f6d6fe6 and touch the same test files, so re-run validation on the merged tree.

Checked and found sound, in case it saves time: the dst_final_meta gate removal is disjoint from the fixup writers (the general fixup returns early for wrote_beginning_of_tile, fattn-common.cuh:1385-1388, and a tile served by one whole-tile block has no needs_fixup peer), and for parallel_blocks > 1 without stream-k the combine kernel overwrites the same buffer later on the same stream; the two new __syncthreads() cannot deadlock (block-uniform conditions, no thread-divergent early returns, and nbatch_combine <= DV/2 in every RDNA config entry); VKQ_C_COUNT matches all three index sites and fixes a real under-allocation for the float tiles; and the device-guard raise cannot change dense dispatch because fattn.cu:693 still caps AMD WMMA at ne[0] <= 128.

[raufaser]

Thanks for the careful read — all five addressed in bebc08fa3 (pushed), validated below on gfx1100 + RTX 4090.

1. RDNA4 fail-closed (blocking): fixed as suggested (gate, not tile port).
New GGML_CUDA_FATTN_KVARN_AMD_RDNA4_WMMA arch: amd_arch() returns it for IS_RDNA4, eligibility keeps the 128 cap there (D64-style fail-closed) while RDNA3/3.5 admit 256. I can't qualify fp32 tiles without RDNA4 hardware, so no tile port. Dispatch comment corrected to RDNA3 scope, route-policy test now asserts RDNA4-256 rejection. Device guard stays wide as backstop — unreachable (KVarN eligibility rejects RDNA4-D256; standard FA keeps upstream's D128 cap on this branch).

2. Dense D128: evidence (no scoping).
PR-branch build, gfx1100, Qwen3-4B: dense f16 KLD 0.0074 / same-top 96.86% (vs kvarn6 0.0068/96.28% on the same build — healthy), llama-bench pp512 6322 t/s, tg128 177 t/s. The fp32 tile is accuracy-positive by construction; happy to scope it to KVarN if you still prefer that.

3. HIP tail coverage: provided.
GGML_KVARN_TEST_LADDER_TAIL=128 on the ladder: D256 n_kv 256→8192 RMSE 6.8e-05→5.0e-06 on the WMMA route (prompt_prefill=1), i.e. the ub>=96 attached-tail geometry is now covered on HIP. Note: D512+tail aborts in fattn-tail.cuh:713 ("unsupported structured body") — verified pre-existing by running the identical ladder on the pre-change build (same RMSEs, same abort point), so out of scope here.

4. Opt-in double pass: fixed.
The HIP opt-in block now precedes the generic probe. Validated: GGML_KVARN_AMD_PROMPT_PORTABLE=1 27B run routes portable-native/hip-prompt-precision-optin, KLD 0.0049/95.5%, zero traps; default path bit-identical (0.0057 both before/after).

5. Knob docs: done.
docs/beellama-args.md documents exact values (any nonzero opts in, conventionally 1; unset/0/non-numeric keeps WMMA default) and notes the pre-probe ordering. Boolean aligned to != 0 like force_portable.

Re-validation on the merged tree (tip 8b59e71ba now in-branch): HIP trio 27B 0.0057/94.5%, 4B 0.006822 bit-exact, 35B 0.0039/97.5%; ladder ~1e-5; ctest green modulo the known gfx1100 test-kvarn store-route abort (device-specific, passes on 4090). CUDA 4090: 27B 0.0035, KVarN/FA ctest 16/16.

[Anbeeld]

Thanks for the turnaround on the five items. I checked the new head bebc08fa3 and the RDNA4 gate, the pre-probe opt-in ordering and the knob docs all hold up, including the unit assertion for the RDNA4-256 rejection. Rather than merging and following up, I would like the remainder closed in this branch. Six items, each with what I would accept.

1. Drop the config rows no route can reach, or name the route that reaches them. 570cb2cf0 retunes (320,256,32/64), (512,512,16/32) and (576,512,16/32) in ggml_cuda_fattn_mma_get_config_rdna (fattn-mma-f16.cuh:167-186) and widens the AMD WMMA device guard to DKQ > 576 (:1967-1972). The KVarN MMA switch only has 128, 256 and 512 (fattn-kvarn-dispatch.cu:993-1003), standard FA keeps the Q->ne[0] <= 128 cap on RDNA (fattn.cu:693), eligibility rejects above 256 on RDNA3 and above 128 on RDNA4 (fattn-kvarn-route-policy.h:65-66), and get_config_rdna is selected only for AMD WMMA devices (:251-253). So six of the seven edited rows are unreachable, while the guard comment above them still claims to be mirrored by the host eligibility check, which is no longer true for 256 < DKQ <= 576. I would accept either the base values back with the guard restored to 256, or a one-line statement per row naming the reachable route and the measurement that covers it.

2. Give the one live row a number. (256,256,64) (fattn-mma-f16.h:167) is the only edited row the new default actually uses: nthreads 128 to 256, nbatch_fa 64 to 32, nbatch_combine 64 to 32, so __launch_bounds__(256, 2). The RDNA3 fp32 accumulator is 16 tiles of 8 lanes (mma.cuh:184), 128 fp32 values per thread against 32 registers' worth for the previous half2 tile, so spilling is the first thing I would want ruled out. A gfx1100 llama-bench -p 4096 or KLD run at the geometry that selects ncols=64 would settle it.

3. Fix the 570cb2cf0 claim now, while it is still a revision. The message says "256/320/512/576 keys, measured KLD-clean on gfx1100". D320, D512 and D576 cannot be selected on gfx1100 by this code, so whatever was measured, it was not those routes. Amend the message on your branch, or add a correction in the PR description stating which dims and which route were measured. I would rather not have the merged history assert something the merged code cannot reproduce.

4. Make the HIP tail coverage a committed test, not a manual run. test_native_flash_attention_prefill_route_parity() still returns early on HIP at tests/test-kvarn.cpp:4262-4265, and the ladder stays env-gated with exact_tail_tokens at 0 (:4334-4337), so the ub>=96 fix rests on one GGML_KVARN_TEST_LADDER_TAIL=128 run. Please land an always-run case that covers D256 WMMA prompt prefill with an attached exact tail on HIP, comparing against the CPU reference the way the ladder does, for example n_kv 512 with a 128-candidate tail. Acceptance: it should fail with the !is_kvarn_kv gate restored in fattn-mma-f16.cuh and pass on your head. If it passes with the gate restored, it is not covering the fix.

5. Pin the portable original-V-domain path. I verified the reduction[D]/transform[D] fix is correct and I withdraw my earlier claim that it affects D256/D512 HIP routes: the writes sit under if (v_original_domain), so it is latent on HIP and reachable on CUDA portable with the original V domain. Since v0.4.6 has reduction[THREADS], this is a v0.4.7-dev regression with no earlier branch to backport to, so it belongs here. It just needs a committed test that exercises portable with original_value_domain = true at D256 and D512; the HIP ladder cannot reach that path.

6. Close or scope the dense D128 gap. The mma_tile_sizes<128> float tile also applies to dense HIP attention (fattn.cu:693). Your Qwen3-4B gfx1100 run shows it is not broken, but the pp512 figure has no pre-change counterpart, so a regression is not excluded. Either run the same dense f16 KLD and llama-bench on the pre-change build and put both numbers in the PR, or scope the fp32 tile selection to the KVarN path.

Update the feature matrix too. docs/beellama-features.md:193 still lists RDNA3, RDNA3.5 and RDNA4 together as "WMMA generic/prefill", and :500-503 still calls AMD routing source-policy verified pending hardware reports, while the docs/beellama-args.md paragraph you added says RDNA4 keeps the portable fallback and you now have gfx1100 results. The per-arch contract should read the same in both files.

Out of scope, please file separately: the D512 + attached-tail GGML_ABORT at fattn-tail.cuh:713. You reproduced it on the pre-change build, so it is not this PR's regression, but aborting the process instead of degrading to materialization is worth its own issue with the ladder command that triggers it.

After these I have no remaining blockers on the change itself. For the record: my review is read-only. There is no ROCm toolchain on this machine, so the HIP and AMD paths are neither compiled nor run here, and all of your gfx1100 and RTX 4090 numbers are taken as reported.

[raufaser]

All six items are closed on f7f3350 (force-pushed; 4 commits on top of bebc08f). Per-item, with the acceptance evidence:

1. Reachable rows only: 320/512/576 reverted to base values, AMD WMMA device guard back to 256. The only live row is (256,256,64).

2. Live-row number: llama-bench -p 4096, 27B Q5_K_S, gfx1100, ngl99, final tree: pp4096 649.6 t/s (kvarn6) vs 937.8 t/s (dense-f16 reference); tg128 34.6 both. No spill collapse at long prompt.

3. Message fixed via amend (69c616d now lists only the live row; the 320/512/576 claim is gone from history).

4. Committed test_kvarn_d256_prompt_tail_regression: route-pinned WMMA (prompt_prefill=1 in the route stats), body_meta den>0 over all rows, serving-geometry (n_kv 512 + 128 tail) parity vs CPU reference. Passes on gfx1100. One honest caveat: the fail-acceptance sweep on the gate-restored tree (nq 32-256 x n_kv 256-32768, HIP + CUDA) never trips single-shot, so the test as committed guards the route plus the metadata invariant rather than reproducing the original failure. The trigger needs incremental multi-ubatch cache-fill dynamics (serving KLD caught the original at 2.118). Serving KLD stays the acceptance unless you want an incremental-fill harness as follow-up.

5. Committed test_native_flash_attention_portable_original_v: D256 rmse 5.4e-4 OK, D512 rmse 5.1e-4 OK on gfx1100; fail-demo verified (gate-restored tree aborts with -inf). This required an enabling fix that is part of the branch: domain_ok now admits MIXED at all dims (was D64-only); without it portable_supported rejects every D256/D512+MIXED dispatch and the path the test pins is unreachable.

6. Dense gap closed both ways: the fp32 tile is scoped to KVarN via mma_tile_sizes_kvarn (dense keeps the primary tiles on all arches), and measured - HIP 4B dense-f16 KLD 0.00688/96.7%, plus CUDA dense parity round-1 vs round-2: bit-identical (mean KLD -0.0, max 5.5e-5, 100% same-top). Trio on the final tree (gfx1100, ub512, kvarn6): 4B 0.00624/97.8% (bit-exact on repeat; the earlier 0.006822 figure predates the clean rebuild), 27B 0.00571/94.5%, 35B-MoE 0.00386/97.5%.

Feature matrix: RDNA3/3.5 vs RDNA4 rows are split and read the same in both files.

CUDA gate on the final tree (RTX 4090, clean rebuild): 16/16 ctest incl. test-kvarn; 27B kvarn6 0.0042-0.0044/97.5-97.8% on the default route (round-1 default was 0.0062/96.3% against its own base); backend-ops MUL_MAT+TOP_K 2190 cases OK.

D512+tail filed as #158. Note for the record: the fattn-tail.cuh:713 abort reproduces on base and mid-branch trees but no longer reproduces on f7f3350 (tail ladders 128 and 256 fully green) - details in the issue.