Skip to content

[CK_TILE] Add LLC-aware FMHA head grouping and head-major scheduling on RDNA #5018

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
illsilin merged 20 commits into from
Mar 16, 2026
Merged

[CK_TILE] Add LLC-aware FMHA head grouping and head-major scheduling on RDNA #5018

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
fbf59e4
[CK_TILE] Add LLC-aware FMHA head grouping and head-major scheduling
hyoon1 Mar 1, 2026
9ab6720
macro placement and env API updates
hyoon1 Mar 3, 2026
db45ef0
Merge branch 'develop' into users/hyoon1/ck/fmha-group-head
hyoon1 Mar 3, 2026
c81d0b3
Merge branch 'develop' into users/hyoon1/ck/fmha-group-head
hyoon1 Mar 4, 2026
49322ec
formatting
hyoon1 Mar 4, 2026
8052c8a
Merge branch 'develop' into users/hyoon1/ck/fmha-group-head
hyoon1 Mar 5, 2026
ad1a5a2
Merge branch 'develop' into users/hyoon1/ck/fmha-group-head
hyoon1 Mar 6, 2026
b5ba2b0
switch bshd head grouping to single-launch remap
hyoon1 Mar 6, 2026
cb1de32
formatting
hyoon1 Mar 6, 2026
5495f97
remove unnecessary changes
hyoon1 Mar 6, 2026
0c97da4
Remove head-group count caps
hyoon1 Mar 6, 2026
aee7588
Revert single-launch head-grouping series
hyoon1 Mar 7, 2026
aecc364
Merge branch 'develop' into users/hyoon1/ck/fmha-group-head
hyoon1 Mar 9, 2026
73ce5d9
fix missing return in check_hdim compatibility rule
hyoon1 Mar 11, 2026
90eed51
gate head-major elemwise-bias path for gfx12 on ROCm 7.1
hyoon1 Mar 12, 2026
1cba80d
Merge remote-tracking branch 'upstream/develop' into users/hyoon1/ck/…
hyoon1 Mar 13, 2026
ebd0616
Restore executable bit on CMakeLists files
hyoon1 Mar 13, 2026
eceb7c9
Drop unintended untracked header include
hyoon1 Mar 13, 2026
064a944
Merge branch 'develop' into users/hyoon1/ck/fmha-group-head
eidenyoshida Mar 13, 2026
3981c1d
Merge branch 'develop' into users/hyoon1/ck/fmha-group-head
hyoon1 Mar 16, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion projects/composablekernel/example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -808,7 +808,7 @@ def check_hdim(problem_ctx: ProblemContext, kernel_ctx: KernelContext) -> bool:
kernel_ctx.pipeline.F_bias != "no"
or kernel_ctx.pipeline.F_dropout == "t"
):
False
return False
return True

def check_feature(
Expand Down
10 changes: 8 additions & 2 deletions projects/composablekernel/example/ck_tile/01_fmha/fmha_fwd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -299,6 +299,8 @@ struct fmha_fwd_args
ck_tile::index_t hdim_v;
ck_tile::index_t nhead_q;
ck_tile::index_t nhead_k;
ck_tile::index_t num_head_q_total = 0;
ck_tile::index_t head_start = 0;

float scale_s;
float logits_soft_cap;
Expand Down Expand Up @@ -733,7 +735,9 @@ auto fmha_fwd_create_kargs_and_grids(fmha_fwd_args args)
args.block_scale_size_kv,
args.cu_seqlen_q_ptr,
args.cu_seqlen_k_ptr,
args.sink_ptr);
args.sink_ptr,
args.num_head_q_total,
args.head_start);
}
else
{ // create batch mode kernel arguments
Expand Down Expand Up @@ -795,7 +799,9 @@ auto fmha_fwd_create_kargs_and_grids(fmha_fwd_args args)
args.block_scale_size_kv,
args.cu_seqlen_q_ptr,
args.cu_seqlen_k_ptr,
args.sink_ptr);
args.sink_ptr,
args.num_head_q_total,
args.head_start);
}
}();

Expand Down
Loading
Loading