Skip to content

[CK] Fix FMHA sink dispatch when init_sink_value is set #7530

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
poyenc merged 2 commits into from
May 18, 2026
+19 −4
Merged

[CK] Fix FMHA sink dispatch when init_sink_value is set #7530

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c80318e
fix(fmha): include init_sink_value in has_sink trait dispatch
poyenc May 18, 2026
62f47e4
fix(fmha): gate sink smoke tests behind opt-in flags
poyenc May 18, 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/fmha_fwd_runner.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1138,7 +1138,7 @@ fwd_result fmha_fwd_run(mode_enum mode,
traits.has_logits_soft_cap = 0.f < logits_soft_cap;
traits.mask_type = mask.type;
traits.bias_type = bias.type;
traits.has_sink = mask.sink > 0 ? true : false;
traits.has_sink = (mask.sink > 0 || init_sink_value != 0) ? true : false;
traits.has_lse = lse;

if constexpr(std::is_same_v<fmha_fwd_traits, std::decay_t<decltype(traits)>>)
Expand Down
21 changes: 18 additions & 3 deletions projects/composablekernel/example/ck_tile/01_fmha/script/smoke_test_fwd.sh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,17 +28,27 @@ COMMON_ARGS='-v=1 -warmup=0 -repeat=1'

TEST_SPLITKV=0
TEST_APPENDKV=0
TEST_STREAM_SINK=0
TEST_GPTOSS_SINK=0
# options:
# -s: run splitkv tests
# -a: run appendkv tests
while getopts ":sa" opt; do
# -m: run StreamLLM sink mask tests (requires sink=true kernels)
# -g: run GPT-OSS sink init tests (requires sink=true kernels)
while getopts ":samg" opt; do
case "${opt}" in
s)
TEST_SPLITKV=1
;;
a)
TEST_APPENDKV=1
;;
m)
TEST_STREAM_SINK=1
;;
g)
TEST_GPTOSS_SINK=1
;;
*)
;;
esac
Expand Down Expand Up @@ -300,8 +310,13 @@ run_padding_smoke_tests
run_padding_basic_boundary_tests
run_fp8bf16_tests
run_fp8fp32_tests
run_sink_mask_tests
run_sink_init_tests
if [ $TEST_STREAM_SINK -eq 1 ] ; then
run_sink_mask_tests
fi

if [ $TEST_GPTOSS_SINK -eq 1 ] ; then
run_sink_init_tests
fi

if [ $TEST_APPENDKV -eq 1 ] ; then
run_fp16_appendkv_tests
Expand Down
Loading