Skip to content

[MHA V3] Remove uses of AITER_ASM_DIR and embed code objects - #1597

Merged
draganmladjenovic merged 1 commit into
mainfrom
draganm/embed_hsaco
Dec 12, 2025
Merged

draganmladjenovic merged 1 commit into
mainfrom
draganm/embed_hsaco

Conversation

@draganmladjenovic

@draganmladjenovic draganmladjenovic commented Dec 9, 2025

Copy link
Copy Markdown
Contributor

Motivation

Remove the need for users to set AITER_ASM_DIR

Technical Details

Adapted generate.py scripts to embed code object into source files

Test Plan

Build check and manual testing of aiter/op_tests/cpp/mha on gfx942 + TE tests/jax/test_fused_attn.py + op_tests/test_mha.py

Test Result

Passing op_tests/test_mha.py on gfx942 and gfx950.

Submission Checklist

@draganmladjenovic
draganmladjenovic requested review from minmengdie and slippedJim and removed request for slippedJim December 9, 2025 13:44
@draganmladjenovic
draganmladjenovic force-pushed the draganm/embed_hsaco branch 2 times, most recently from 9082c2c to 6c573e7 Compare December 9, 2025 14:14
@i-chaochen
i-chaochen requested a review from valarLip December 9, 2025 15:09
Comment thread 3rdparty/composable_kernel
@valarLip

Copy link
Copy Markdown
Collaborator

looks not works well ..
/workspace/aiter/jit/build/module_fmha_v3_fwd/blob/asm_fmha_fwd_v3_gfx950.cpp:46:186: error: use of undeclared identifier 'fwd_hd192_hd128_bf16_causal'
46 | template<> struct FmhaFwdV3Buf<fmha_fwd_kernel_selector<FmhaFwdBf16, 192, 1, false, false, 1, GPUArch::gfx950>> { static constexpr const void * fwd_v3_buf = fwd_hd192_hd128_bf16_causal; };

@draganmladjenovic
draganmladjenovic requested a review from a team December 10, 2025 15:39
@draganmladjenovic

Copy link
Copy Markdown
Contributor Author

looks not works well .. /workspace/aiter/jit/build/module_fmha_v3_fwd/blob/asm_fmha_fwd_v3_gfx950.cpp:46:186: error: use of undeclared identifier 'fwd_hd192_hd128_bf16_causal' 46 | template<> struct FmhaFwdV3Buf<fmha_fwd_kernel_selector<FmhaFwdBf16, 192, 1, false, false, 1, GPUArch::gfx950>> { static constexpr const void * fwd_v3_buf = fwd_hd192_hd128_bf16_causal; };

Where do you observe this? I see fwd_hd192_hd128_bf16_causal.co exists.

@draganmladjenovic
draganmladjenovic force-pushed the draganm/embed_hsaco branch 4 times, most recently from 8526a39 to f25e739 Compare December 11, 2025 14:05
@draganmladjenovic

Copy link
Copy Markdown
Contributor Author

@valarLip testing passed. Can we merge this. I have rebase this daily due to new change to main.

@draganmladjenovic

Copy link
Copy Markdown
Contributor Author

@ valarLip can we merge this before I get another merge conflict?

@draganmladjenovic
draganmladjenovic merged commit ba9047e into main Dec 12, 2025
22 checks passed
@draganmladjenovic
draganmladjenovic deleted the draganm/embed_hsaco branch December 12, 2025 22:21
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.

2 participants