fix(ck_tile/fmha): pass kHasSink to batch_prefill traits

[JinHe-102]

Fixes #12093

Summary

The dispatcher's batch prefill codegen emits an ill-formed traits list
on every architecture. kHasSink_ is never emitted, so every argument
after it shifts one slot: the page size lands in a bool parameter, and
two scoped enums bind to parameters of unrelated types. The batch prefill
path does not compile anywhere.

• emit kHasSink_ in both batch prefill traits lists in dispatcher/codegen/fmha/codegen.py — two lines
• add two unit tests that pin the trailing traits slots to their positions

Only the dispatcher is affected. The example generator
example/ck_tile/01_fmha/generate.py emits the list correctly.

What goes wrong

TileFmhaBatchPrefillTraits and fmha_fwd_batch_prefill_traits_ both
declare kHasSink_ between kSkipMinSeqlenQ_ and kPageBlockSize_.
Codegen skipped it, so a generated instance expanded to:

kBlockPerCu_     = -1
kSkipMinSeqlenQ_ = false
kHasSink_        = 16                      // page size
kPageBlockSize_  = LINEAR_LAYOUT           // memory layout enum
kKVMemoryLayout_ = VLLM_BLOCK_TABLE_2D     // lookup table enum
kKVLookupTable_  = <default>

int to bool is a narrowing conversion in a converted constant
expression, and the two enums bind to unrelated parameter types — three
separate errors from one missing argument.

The fix emits {_bool_cpp(sig["sink"])} in both lists. The argument that
was missing is the false now at index 12, immediately before the page
size.

Test Plan

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

Both tests share a _batch_prefill_traits_args helper that runs codegen
for a gfx950 batch prefill config and parses the generated
TileFmhaBatchPrefillTraits<...> argument list.

• test_batch_prefill_traits_pass_sink_before_page_size pins each trailing parameter to its slot
• test_batch_prefill_sink_slot_tracks_the_signature sweeps sink over both values and requires slot 12 to follow it

The second test is needed because slots 11 (kSkipMinSeqlenQ) and 12
(kHasSink) are adjacent bools: asserting only that each is
"true"/"false" still passes if the two are swapped, or if the sink
flag is emitted as a constant.

gfx950 rather than an RDNA target on purpose. The traits list is emitted
identically everywhere, so the defect is not arch-specific — but on RDNA
validate_config rejects the batch prefill pipeline before codegen runs,
leaving the test nothing to parse. A CDNA target keeps this PR
independently testable.

Test Result

before   FAILED - AssertionError: 15 != 16
         ['true','true','true','true','false',NO_BIAS,'false','false',
          'false',NO_SCALE,'1','false','16',LINEAR_LAYOUT,VLLM_BLOCK_TABLE_2D]
after    Ran 24 tests - OK

Each test was checked against an injected mutant rather than trusted on
its face:

injected mutation	caught by
revert the two-line codegen fix	test 1, 15 != 16
re-emit slot 12 as a literal false	test 2 only
swap slots 11 and 12	test 2 only

The last two rows are the blind spot the second test was added for.

Submission Checklist

• tests added and passing locally (24 tests, OK)
• inline documentation for the motivation
• ruff-format clean; no C++ changed; all 8 pre-commit hooks pass
• rebased onto current origin/develop 4606f83174
• REGRESSION_TESTS entry — not needed, the test runs well under 30s
• release notes — internal codegen fix, not user facing

[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.

[JinHe-102]

Friendly ping for a maintainer — this PR is stalled on one thing that needs write access, and it is not in the diff.

The workflow runs need approval. Every pull_request-triggered workflow here is action_required — pre-commit, clang-tidy, Component CI, TheRock CI, TheRock Multi-Arch CI and Multi-Arch ASAN. That is GitHub's first-time external contributor gate, so none of them have actually run. Reviews were auto-requested from the CODEOWNERS teams when this was marked ready, so that side is in place.

About the red therock-pr-bot check

It is not a policy failure. The table the bot posted reads All Policy Checks Passed (PR Description / Forbidden Files / Unit Test), and that comment's last update is timestamped three seconds into the job's Enforce policy step — the rules were evaluated and published immediately.

The step then ran for 907 s before failing. The workflow invokes policy_check.py --timeout-seconds 900 --poll-seconds 15, so that is the timeout, not a crash: the bot polls for the required checks above to conclude, and they cannot conclude while they sit in the approval gate. All three runs ended the same way (~19.5 min each), which is what a deterministic wait looks like rather than flakiness.

The side effect is that the Not ready to Review label is never removed — in policy_check.py the remove_label(...) call sits after the point where the timeout returns. So a PR whose policies all passed is left labelled as not reviewable. Approving the workflows should clear both on the next run.

Worth noting for the bot's own sake: any first-time external contributor will hit this, since the approval gate and the required-check wait are mutually exclusive by construction.

Nothing here involves the PR contents — libraries-pr-bot.yml checks out github.event.pull_request.base.sha by design and never sees the head.

The change itself

2 files, +69, no C++ touched. It adds the missing kHasSink_ argument to the dispatcher's batch prefill traits lists, which had shifted every later trait by one slot and made the generated instance ill-formed on every architecture. Two unit tests come with it, both checked against injected mutants rather than trusted on their face. Details and the before/after output are in the description.

Happy to rebase, split, or adjust the tests on request.