[CK_TILE] FMHA forward: add head dim 512 support (qr pipeline)

[Reis-McMillan]

Suggestion Description

Problem

Every ck_tile FMHA forward pipeline static-asserts kSubQKHeaddim <= 256, and
fmha_fwd.py has no tile above 256. Single-head attention over a 512-wide head
therefore has no fused path and falls back to a materialised score matrix.

The concrete case is the VAE mid-block attention shared by the SD/FLUX family of
autoencoders (FLUX.1 Kontext included): one head of width 512 over every latent
pixel. At a 7 MP decode that is ~111k tokens, i.e. a 12.3-billion-element score
matrix which neither fits in memory nor launches as a single elementwise region.
PyTorch handles this shape through its memory-efficient SDPA kernel; ck_tile
currently cannot.

Proposal

Enable hdim 512 for fp16/bf16 (and fp32) on gfx9 through the qr pipeline only,
following the precedent of the hdim 256 enablement (#6846):

• relax the assert in block_fmha_pipeline_qr_ks_vs.hpp and encode the real
  constraint: Q and the f32 O accumulator are register resident, so
  kM0 / NumWarps <= 16;
• add (512, 512) tiles at 16 rows per warp (64 rows / 4 warps and
  128 rows / 8 warps) and route between them on block count;
• qr variants only: qr_async needs ~110 KB of LDS at this width;
• mirror the rules in the dispatcher codegen and its tests;
• extend test_fmha_fwd.cpp, the smoke and benchmark scripts, and the README.

Notes

The other pipelines (qr_async, qr_async_trload, qr_tdm, v3, splitkv,
pagedkv, batch_prefill) keep their <= 256 assert. Expected throughput on
gfx942 is below hdim 256: at 16 rows per warp the kernel is LDS-read bound
(every warp streams the full K and V tile), which is a property of the qr
design at this width and a follow-up optimisation, not a blocker.

Operating System

No response

GPU

No response

ROCm Component

composablekernel