From ec2db01e4acd672ef3ba3a01e422b8b8f046d567 Mon Sep 17 00:00:00 2001 From: root Date: Wed, 1 Apr 2026 15:54:17 +0000 Subject: [PATCH] Fix fmha_fwd early-exit bug: seqlen_q <= min_seqlen_q should be < 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 --- include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp b/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp index 16f5b00bb1..bc40d2565f 100644 --- a/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp +++ b/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp @@ -1509,7 +1509,7 @@ struct FmhaFwdKernel if constexpr(kSkipMinSeqlenQ) { - if(kargs.seqlen_q <= kargs.min_seqlen_q) + if(kargs.seqlen_q < kargs.min_seqlen_q) { return; }