ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 11:16:59 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix fmha_fwd early-exit bug: seqlen_q <= min_seqlen_q should be <

Browse Source 
The kSkipMinSeqlenQ optimization incorrectly used <= comparison, causing
the kernel to skip batches where seqlen_q equals min_seqlen_q. This
happens in the common case of no padding (all batches have the same
seqlen_q == min_seqlen_q), producing all-zero output silently.

Changed to strict < so batches with exactly min_seqlen_q tokens are
still processed.

Made-with: Cursor
This commit is contained in:
root
2026-04-01 15:54:17 +00:00
parent cb6fb2802d
commit ec2db01e4a
1 changed files with 1 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp
View File

@@ -1509,7 +1509,7 @@ struct FmhaFwdKernel
if constexpr(kSkipMinSeqlenQ) if constexpr(kSkipMinSeqlenQ)
{ {
if(kargs.seqlen_q <= kargs.min_seqlen_q) if(kargs.seqlen_q < kargs.min_seqlen_q)
{ {
return; return;
} }