fix(attention): Track E PV repack — m16n8k16 A-operand layout (re-enables Track E)

[kekzl]

Context

PR #352 was meant to ship this fix but auto-merge fired on the disable-hotfix commit before the actual fix commit was pushed. Main is currently running with Track E unconditionally disabled (early `return false;` at launcher entry).

This PR re-enables Track E with the actual root-cause fix.

Root cause

The PV repack (S_frag → P_frag for the second mma) used the wrong A-operand interleaving for `mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32`.

Correct PTX m16n8k16 A-frag per-lane layout (row × k-half interleaved):

• `a[0] = (row_a, k<8)` → accumulates into `d[0,1]` (row_a output)
• `a[1] = (row_b, k<8)` → accumulates into `d[2,3]` (row_b output)
• `a[2] = (row_a, k≥8)` → accumulates into `d[0,1]`
• `a[3] = (row_b, k≥8)` → accumulates into `d[2,3]`

Buggy code grouped by row-then-k:

• `a[1] = S[row_a, k≥8]` ← WRONG (should be row_b k<8)
• `a[2] = S[row_b, k<8]` ← WRONG (should be row_a k≥8)

Row_a softmax weights ended up in d[2,3] (row_b output) and vice versa. Original test fill (magnitude 0.125) was too uniform to expose this — error stayed within 5e-3 tolerance. Real Qwen3 attention scores produced ~0.116 abs error → garbage output.

Fix

1. P_frag construction — corrected interleaving order in PV mma call.
2. O_frag rescale & epilogue store — restored standard mappings (`frag[0,1] → row_a = lane/4`, `frag[2,3] → row_b = lane/4 + 8`).
3. Test fill magnitude — bumped from `* 0.125f` to `* 1.0f` so future bugs of this class don't slip through.
4. Removes the unconditional `return false` from fix(attention): Track E PV repack — m16n8k16 A-operand layout #352.

Verification

```
Prompt: "What is 17 + 25?"
With #352 (Track E disabled, runs cuBLAS): coherent output
With this fix (Track E enabled): "Okay, I need to figure out how to solve this problem..." (coherent)
```

Test suite at stricter magnitude=1.0 fill:

• 6/6 TrackE_Correctness tests PASS (1 hd=512 SKIP, expected)
• 3/3 TrackE_Features tests PASS (SWA, softcap, combined)
• max_abs = 0.000244 (FP16 rounding floor only)

Test plan

• All 9 TrackE_* tests pass at magnitude=1.0
• Real Qwen3-8B Q8_0 inference produces coherent output
• verify-fast pre-push hook green
• Re-validate perf gains from perf(attention): Track E warp-spec 4+4 + perf validation report #351 are intact after the fix (no reason to expect regression, but worth a quick A/B post-merge)

🤖 Generated with Claude Code