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

Revert "early exit"

Browse Source 
This reverts commit e9b58d3307.
This commit is contained in:
zanzhang
2025-05-21 12:47:51 +08:00
parent e9b58d3307
commit be2d38fab3
1 changed files with 7 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -1067,6 +1067,13 @@ struct FmhaFwdKernel
const auto adjusted_seqstart_q_ptr = kargs.seqstart_q_ptr + i_batch;
kargs.seqlen_q = adjusted_seqstart_q_ptr[1] - adjusted_seqstart_q_ptr[0];
// # of required blocks is different in each groups, terminate unnecessary blocks
// earlier
if(kargs.seqlen_q <= i_m0)
{
return;
}
if constexpr(kIsChunkedPrefill)
{
if(kargs.seqlen_q <= kargs.min_seqlen_q)
@@ -1075,13 +1082,6 @@ struct FmhaFwdKernel
}
}
// # of required blocks is different in each groups, terminate unnecessary blocks
// earlier
if(kargs.seqlen_q <= i_m0)
{
return;
}
if(kargs.seqlen_k_ptr != nullptr)
{
kargs.seqlen_k = kargs.seqlen_k_ptr[i_batch];