feat(ck_tile/fmha): add gfx1100 batch_prefill_gfx11 pipeline policy

[JinHe-102]

ISSUE ID : #12100

Summary

gfx1100 has no batch prefill pipeline it can run. Emission hands it the gfx9 rule set, so all 277,120 configs select qr_async — which gfx1100 cannot compile — and every one is rejected downstream. Today the architecture has zero working batch prefill kernels.

• add BlockFmhaBatchPrefillPipelineQRKSVSAsyncGfx11Policy, derived from the existing default policy
• select it for gfx1100 only, on a narrow predicate: block 256, tile 128x32x32, hdim 128, linear page KV, no dropout
• leave every other arch, tile, head dim and dtype on qr_async, untouched
• fix three tile-engine harness defects that made this path impossible to validate
• make a sweep that builds nothing exit non-zero instead of reporting success

Emission for gfx1100 goes from 277,120 configs that cannot compile to 192 that all build. CDNA is unchanged, config for config.

Partially addresses #12100: gfx1100 gets a pipeline it can run, but get_pipelines_for_config is not changed and fwd is not touched.

Stacked on #12092 — the diff shows two commits, the first being
that PR. Please review 66469c3f8f only; will rebase once #12092
merges.

Why the existing pipeline cannot be reused

Two independent blockers, both in device code.

K staging needs a CDNA3 instruction. qr_async moves K to LDS via async_load_tile_raw → async_global_load_lds_dwordxn, which is guarded #if !defined(__gfx94__) && !defined(__gfx950__) with "global_load_lds
requires CDNA3+" (amd_buffer_addressing_builtins.hpp:1503).

gemm1 cannot consume a gfx11 P tile. BlockGemmARegBSmemCRegV2 expects a thread buffer of 16; gfx11 WMMA produces 32.

So this is not a tuning gap: the pipeline is unbuildable on gfx1100 by construction, and supported_pipelines is right to exclude it.

What the new policy does

change	blocker it answers
kUseSyncKLoad = true — K staged with ordinary loads plus LDS stores	no CDNA3 instruction needed
vec4 K packing, row-major K LDS descriptor on N0=32 / K0=32 tiles	keeps that path's LDS traffic coalesced
Independent-V gemm1 BlockGemmARegBRegCRegV2 over a tiled V LDS layout	V arrives in registers, avoiding the unusable gemm1
P permuted from the WMMA C distribution to the A distribution	makes gfx11 WMMA output consumable by that gemm1

The permutation reuses the existing PermuteWarpGemmCToA helper that block_fmha_pipeline_qr_ks_vs.hpp and four other pipelines already use.

How CDNA is kept out

where	effect
instance_gen.py	emits nothing outside the predicate
validation.py	rejects an out-of-predicate config with a named error
pipeline header	device-side static_assert keyed on kUseSyncKLoad

The assert keys on kUseSyncKLoad, not on the architecture. No CDNA instance sets that flag, so none can reach the new path — a structural guarantee, not a measurement that happens to come out clean.

It sits inside #if defined(__gfx11__) for a separate reason: amdgcn_target_wave_size_id::HOST is 64, so in the host compilation pass Problem::kBlockSize evaluates to 512 for an 8-warp block and a class-scope assert would fire on every instance, CDNA included.

Emission

expand_sweep on batch_prefill.json:

arch	before	after	pipeline after
gfx90a	277,120	277,120	qr_async
gfx942 / gfx950	288,964	288,964	qr_async
gfx1100	277,120	192	batch_prefill_gfx11

Resource usage

From the code objects of all 144 built kernels, not a sample (llvm-objdump --offloading, then llvm-readelf --notes):

	value
LDS per workgroup	27,648 B, identical across all 144
VGPR / SGPR	205-230 / 51-83
spills, scratch	0

27,648 B is 42% of the 65,536 B per-workgroup LDS limit, so two workgroups fit concurrently.

Harness fixes

The tile engine runner could not exercise this path at all:

• fmha_ctypes_lib.cpp zeroed the paged K/V buffers and never copied the host pages — every verification compared against zeros
• the same file hardcoded window_size_right = -1, a full row, so top_left and bottom_right did not mask
• fmha_utils.py passed dense BHSD without the linear page pack or the [B,H,S,D] <-> [B,S,H,D] transpose the kernel expects

fmha_benchmark.py now exits non-zero when kernels fail to build or any verification reports FAIL; previously it printed the failure and exited 0. This is a global tile engine change, not gfx11 only.

Test Plan

cd projects/composablekernel/dispatcher/tests
python3 -m unittest test_fmha_rules test_fmha_codegen

cd ../..
python3 tile_engine/ops/fmha/fmha_benchmark.py \
  tile_engine/ops/fmha/configs/batch_prefill.json --arch gfx1100

Test Result

Host: 33 tests, all pass.

Device: gfx1100 / W7900, ROCm 10.0.0, HIP 7.15.26333:

gate      192 emitted, 0 off-predicate
compile   144/144 built, 0 failed
verify    24 configurations x 3 problems = 72 measurements, 72 PASS / 0 FAIL

Problems 1,8,128,128, 1,8,2048,128, 1,8,8,8,1024,128, all D128.

Extended shapes on the same kernels — 32 PASS / 0 FAIL: batch 4 and 8, 4096 context, seqlen 130 straddling the 128 tile, 1000 and 777, GQA hq=4 / hk=2 with q=17 against k=2048, and q=k=1. Across four seeds: 1728 measurements, 0 FAIL. Also reproduced in rocm/composable_kernel:ck_ub24.04_rocm10.0.

Scope and follow-up

• 144 of the 192 emitted instances are verified numerically. The other 48 carry logits soft cap and do not compile in the dispatcher JIT — that is [CK_TILE] dispatcher JIT sets CK_TILE_FMHA_FWD_FAST_EXP2=0 on RDNA targets, which the logits-soft-cap static_assert rejects #12102, pre-existing and independent of this change.
• The exit-code change is not covered by a test.
• Not executed on CDNA hardware; the argument above is structural.
• No performance comparison: nothing builds on gfx1100 today, so there is no baseline.

Submission Checklist

• tests added and passing locally (33 tests, OK)
• device evidence on the target architecture
• clang-format / ruff-format clean; all 8 pre-commit hooks pass
• rebased onto current origin/develop 4606f83174
• depends on fix(ck_tile/fmha): pass kHasSink to batch_prefill traits #12092

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.