Files
composable_kernel/include/ck_tile/ops
juuso-oskari 1864287f95 CK-UA: issue WG1's next-tile prefetch from its MATRIX slot; drop PREFETCH_EARLY knob
WG0 already issues its cooperative K/V prefetch in its MATRIX slot; WG1 issued
its prefetch in its SOFTMAX slot (slot A). Move WG1's prefetch to its MATRIX slot
(slot B) so the async DRAM load issue overlaps the partner group's SOFTMAX phase.

Verified on the canonical fp8 prefill shape (b1 sq=sk=75600 hq=hk=5 d128 non-causal,
GPU2, interleaved A/B): correctness bit-exact (0% mismatch), ~+0.5% mean (up to
+1.5%), and -- the risk with issuing later -- the DRAM wait does NOT get exposed:
the vmcnt-wait stays at ~24 cyc/window (DRAM still fully hidden by the K/V prefetch).
The exposed wait floor remains lgkmcnt (LDS roundtrip), untouched by this change.

Also removes the UA_FA4_PREFETCH_EARLY compile knob (write-before-barrier variant)
from both branches -- it was off-by-default, unsafe for bf16/paged, and only added
noise. Both branches now use the single clean `barrier(); prefetch();` ordering.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-06-15 08:26:51 +00:00
..
2026-01-13 09:21:29 -08:00