Added custom FMHA codegen receipt for TransformerEngine - #6867
Conversation
There was a problem hiding this comment.
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
700receipt in the generator CLI help text. - Added
receipt == 700filtering rules for FMHA forward kernel selection. - Added
receipt == 700filtering 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.
poyenc
left a comment
There was a problem hiding this comment.
Review of receipt 700 addition for TransformerEngine integration.
poyenc
left a comment
There was a problem hiding this comment.
LGTM. All review comments addressed — zero-output guards added for unsupported APIs, and remaining questions answered satisfactorily.
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.
## 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>
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>
Motivation
TE uses AITER to build static MHA libraries, which ultimately rely on CK kernels. We use the
600receipt 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
700receipt for TE's needsTest 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