Skip to content

Add 2048-token pages to FMHA batch prefill - #3760

Draft
rlrs wants to merge 1 commit into
ROCm:developfrom
rlrs:feat/fmha-batch-prefill-page-2048
Draft

Add 2048-token pages to FMHA batch prefill#3760
rlrs wants to merge 1 commit into
ROCm:developfrom
rlrs:feat/fmha-batch-prefill-page-2048

Conversation

@rlrs

@rlrs rlrs commented Jul 26, 2026

Copy link
Copy Markdown

Proposed changes

Add 2048-token pages to the FMHA batch-prefill generator.

vLLM uses 2048-token cache pages for hybrid models whose recurrent layers have
large state-alignment requirements. The batch-prefill kernel already accepts a
compile-time page size, so this only adds 2048 to the generated set.

On an MI250X with a 128K context, the generated kernel was 1.70x faster than
the existing Triton paged-prefill path. The same comparison at 256K was 1.69x.

This does not change kernels generated for existing page sizes or for non-gfx9
targets.

Validation:

python example/ck_tile/01_fmha/generate.py \
  -d batch_prefill -r 200 -f '*ps2048*' -o <output>

1681 generated files, all containing the 2048 page-size specialization

I also checked the generated kernel against the PyTorch attention reference for
the vLLM BF16 MQA shape and ran an end-to-end 128K-context serving benchmark.

Checklist

  • I have added tests relevant to the introduced functionality, and the unit tests are passing locally
  • I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, IF the test takes more than 30 seconds to run.
  • I have added inline documentation which enables the maintainers with understanding the motivation
  • I have removed the stale documentation which is no longer relevant after this pull request
  • (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request
  • I have run clang-format on all changed files
  • Any dependent changes have been merged

Discussion

This change was prepared with AI assistance and reviewed and tested by the
submitter.

Co-authored-by: OpenAI Codex <noreply@openai.com>

Signed-off-by: Rasmus Larsen <rasmus.larsen@alexandra.dk>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant