Skip to content

Added custom FMHA codegen receipt for TransformerEngine - #6867

Merged
poyenc merged 5 commits into
developfrom
zain/te-receipt
May 14, 2026
Merged

Added custom FMHA codegen receipt for TransformerEngine#6867
poyenc merged 5 commits into
developfrom
zain/te-receipt

Conversation

@Micky774

Copy link
Copy Markdown
Contributor

Motivation

TE uses AITER to build static MHA libraries, which ultimately rely on CK kernels. We use the 600 receipt which generates more kernels than TE truly needs. This bespoke receipt allows us to minimize the kernel count, compile time, and memory footprint of our MHA library.

Technical Details

Extended the receipt mechanism to include a custom 700 receipt for TE's needs

Test Plan

Test by building TE using the same receipt profile

Test Result

Build validated in TE using a custom feature branches of AITER/CK to temporarily apply the patch

Submission Checklist

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Note

Copilot was unable to run its full agentic suite in this review.

Adds a dedicated “700” receipt to constrain FMHA codegen output for TransformerEngine, aiming to reduce kernel count, compile time, and memory footprint compared to the broader “600” receipt.

Changes:

  • Documented a new 700 receipt in the generator CLI help text.
  • Added receipt == 700 filtering rules for FMHA forward kernel selection.
  • Added receipt == 700 filtering rules for FMHA backward blob generation.

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated 3 comments.

File Description
projects/composablekernel/example/ck_tile/01_fmha/generate.py Adds CLI help text describing the new TE-specific receipt and its invariants.
projects/composablekernel/example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py Introduces TE-specific forward-kernel compatibility rule under receipt 700.
projects/composablekernel/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py Introduces TE-specific backward blob filtering under receipt 700.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment thread projects/composablekernel/example/ck_tile/01_fmha/generate.py

@poyenc poyenc left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Review of receipt 700 addition for TransformerEngine integration.

@poyenc poyenc left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM. All review comments addressed — zero-output guards added for unsupported APIs, and remaining questions answered satisfactorily.

@poyenc
poyenc requested a review from a team May 12, 2026 16:19
@poyenc
poyenc requested a review from a team May 13, 2026 03:52
@poyenc
poyenc enabled auto-merge (squash) May 14, 2026 02:02
@poyenc
poyenc requested a review from a team May 14, 2026 02:02
@poyenc
poyenc merged commit 3cb0219 into develop May 14, 2026
30 checks passed
@poyenc
poyenc deleted the zain/te-receipt branch May 14, 2026 14:32
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request May 14, 2026
Added custom FMHA codegen receipt for TransformerEngine
 (#6867)

## Motivation

TE uses AITER to build static MHA libraries, which ultimately rely on CK
kernels. We use the `600` receipt which generates more kernels than TE
truly needs. This bespoke receipt allows us to minimize the kernel
count, compile time, and memory footprint of our MHA library.

## Technical Details

Extended the receipt mechanism to include a custom `700` receipt for
TE's needs

## Test Plan

Test by building TE using the same receipt profile

## Test Result

Build validated in TE using a custom feature branches of AITER/CK to
temporarily apply the patch

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
aledudek pushed a commit that referenced this pull request May 20, 2026
## Motivation

TE uses AITER to build static MHA libraries, which ultimately rely on CK
kernels. We use the `600` receipt which generates more kernels than TE
truly needs. This bespoke receipt allows us to minimize the kernel
count, compile time, and memory footprint of our MHA library.

## Technical Details

Extended the receipt mechanism to include a custom `700` receipt for
TE's needs

## Test Plan

Test by building TE using the same receipt profile

## Test Result

Build validated in TE using a custom feature branches of AITER/CK to
temporarily apply the patch

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

---------

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
shumway pushed a commit to ROCm/composable_kernel that referenced this pull request May 27, 2026
Added custom FMHA codegen receipt for TransformerEngine (#6867)

## Motivation

TE uses AITER to build static MHA libraries, which ultimately rely on CK
kernels. We use the `600` receipt which generates more kernels than TE
truly needs. This bespoke receipt allows us to minimize the kernel
count, compile time, and memory footprint of our MHA library.

## Technical Details

Extended the receipt mechanism to include a custom `700` receipt for
TE's needs

## Test Plan

Test by building TE using the same receipt profile

## Test Result

Build validated in TE using a custom feature branches of AITER/CK to
temporarily apply the patch

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

---------

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants