This commit is contained in:
Juuso Korhonen
2025-10-16 09:23:39 +00:00
parent aa4908ac14
commit 072de3842f

View File

@@ -460,10 +460,10 @@ struct UnifiedAttentionKernel
FmhaMask mask = [&]() {
if constexpr(kHasMask)
return ck_tile::make_generic_attention_mask_from_lr_window<FmhaMask>(
cur_batch_query_len, // extend length
seq_len - cur_batch_query_len, // context length
cur_batch_query_len, // extend length
seq_len, // key length (context + extend)
cur_batch_query_len, // x (i.e. extend)
seq_len - cur_batch_query_len, // y (i.e. context)
cur_batch_query_len, // x_total
seq_len, // y_total (x + y)
num_queries_per_kv, // the same sequence index is repeated num_queries_per_kv times along x dim of the tile
kargs.mask_type == GenericAttentionMaskEnum::MASK_FROM_TOP_LEFT);
else