Skip to content

Commit 9d482ad

Browse files
updates to has_work logic, tile scheduler selection, and varlen test suite
1 parent c786e49 commit 9d482ad

4 files changed

Lines changed: 132 additions & 76 deletions

File tree

benchmarks/benchmark_varlen_sched.py

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -243,7 +243,7 @@ def fn():
243243
# fmt: off
244244
MODES = [
245245
("dense", setup_dense),
246-
("static", make_varlen_setup(clc=False, prep="none")),
246+
("single-tile", make_varlen_setup(clc=False, prep="none")),
247247
("clc", make_varlen_setup(clc=True, prep="none")),
248248
("clc-prep", make_varlen_setup(clc=True, prep="precompute")),
249249
("dynamic-prep", make_varlen_setup(clc=False, prep="precompute")),
@@ -368,7 +368,7 @@ def main():
368368
("mode", 14),
369369
("mean_us", 10),
370370
("tok/us", 9),
371-
("rel_static", 11),
371+
("rel_st", 7),
372372
("rel_clc", 9),
373373
]
374374
widths = [w for _, w in cols]
@@ -424,14 +424,14 @@ def main():
424424
None if samples is None else sum(samples) / len(samples) * 1e3
425425
)
426426

427-
static_us = results.get("static")
427+
single_tile_us = results.get("single-tile")
428428
clc_us = results.get("clc")
429429
for cli, _ in selected_modes:
430430
us = results.get(cli)
431431
if us is None:
432432
continue
433433
tok_per_us = (total_q / us) if us > 0 else 0.0
434-
rel_st = f"{static_us / us:.3f}" if static_us else "-"
434+
rel_st = f"{single_tile_us / us:.3f}" if single_tile_us else "-"
435435
rel_cl = f"{clc_us / us:.3f}" if clc_us else "-"
436436
print(
437437
_format_row(

flash_attn/cute/flash_fwd_sm100.py

Lines changed: 39 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -220,8 +220,7 @@ def __init__(
220220
use_clc_scheduler
221221
and self.use_tma_KV
222222
)
223-
self.static_persistent = is_static_persistent
224-
self.is_persistent = self.dynamic_persistent or self.static_persistent
223+
self.is_persistent = self.dynamic_persistent or self.is_static_persistent
225224
self.sched_stages = 1
226225
if self.use_clc_scheduler:
227226
assert self.cluster_shape_mn[1] == 1, f"CLC requires cluster N == 1: {self.cluster_shape_mn}"
@@ -236,14 +235,19 @@ def __init__(
236235
else SchedulingMode.STATIC
237236
)
238237

238+
self.use_varlen_scheduler = False
239239
if is_varlen_q:
240240
if self.dynamic_persistent and not self.use_clc_scheduler:
241+
self.use_varlen_scheduler = True
241242
self.TileScheduler = DynamicPersistentVarlenScheduler
243+
elif self.is_static_persistent:
244+
self.TileScheduler = StaticPersistentTileScheduler
242245
else:
246+
self.use_varlen_scheduler = True
243247
self.TileScheduler = SingleTileVarlenScheduler
244248
elif self.is_causal or self.is_local or self.use_clc_scheduler:
245249
self.TileScheduler = SingleTileLPTScheduler
246-
elif self.static_persistent:
250+
elif self.is_static_persistent:
247251
self.TileScheduler = StaticPersistentTileScheduler
248252
else:
249253
self.TileScheduler = SingleTileScheduler
@@ -1544,11 +1548,15 @@ def load(
15441548

15451549
if const_expr(not self.use_block_sparsity):
15461550
n_block_min, n_block_max = block_info.get_n_block_min_max(
1547-
seqlen, m_block, split_idx=split_idx, batch_idx=batch_idx, num_splits=num_splits,
1551+
seqlen,
1552+
m_block,
1553+
split_idx=split_idx,
1554+
batch_idx=batch_idx,
1555+
num_splits=num_splits,
15481556
)
15491557
if const_expr(self.is_split_kv and block_info.num_splits_dynamic_ptr is not None):
15501558
split_idx = split_idx & 0xFFFF
1551-
if const_expr(not self.is_split_kv) or n_block_min < n_block_max:
1559+
if self.process_work_tile(seqlen, n_block_min, n_block_max):
15521560
n_block_first = n_block_max - 1 if n_block_max > 0 else 0
15531561
page_idx = (
15541562
mPageTable[batch_idx, n_block_first]
@@ -1750,15 +1758,14 @@ def mma(
17501758
process_tile = block_iter_count > Int32(0)
17511759
else:
17521760
n_block_min, n_block_max = block_info.get_n_block_min_max(
1753-
seqlen, m_block, split_idx=split_idx, batch_idx=batch_idx, num_splits=num_splits,
1754-
)
1755-
if const_expr(self.is_split_kv and block_info.num_splits_dynamic_ptr is not None):
1756-
split_idx = split_idx & 0xFFFF
1761+
seqlen,
1762+
m_block,
1763+
split_idx=split_idx,
1764+
batch_idx=batch_idx,
1765+
num_splits=num_splits,
1766+
)
17571767
block_iter_count = n_block_max - n_block_min
1758-
if const_expr(not self.is_split_kv):
1759-
process_tile = True
1760-
else:
1761-
process_tile = n_block_min < n_block_max
1768+
process_tile = self.process_work_tile(seqlen, n_block_min, n_block_max)
17621769

17631770
if process_tile and is_leader_cta:
17641771
for stage in cutlass.range_constexpr(self.q_stage):
@@ -2125,7 +2132,7 @@ def softmax_loop(
21252132
has_work = tile_block_count > Int32(0)
21262133
else:
21272134
tile_block_count = n_block_max - n_block_min
2128-
has_work = const_expr(not self.is_split_kv) or tile_block_count > Int32(0)
2135+
has_work = self.process_work_tile(seqlen, n_block_min, n_block_max)
21292136

21302137
softmax_step = partial(
21312138
self.softmax_step,
@@ -2204,7 +2211,7 @@ def softmax_loop(
22042211
sm_stats_barrier.arrive_w_index(index=stage * 4 + warp_idx)
22052212
# if tidx == 0: cute.printf("softmax row sum stage %d: %f\n", stage, softmax.row_sum[0])
22062213
else:
2207-
if const_expr(not self.is_split_kv) or tile_block_count > Int32(0):
2214+
if has_work:
22082215
mma_si_consumer_phase, sm_stats_producer_phase, s0_s1_sequence_phase = softmax_step(
22092216
mma_si_consumer_phase,
22102217
sm_stats_producer_phase,
@@ -2547,7 +2554,7 @@ def correction_loop(
25472554
has_work = total_block_count > Int32(0)
25482555
else:
25492556
total_block_count = n_block_max - n_block_min
2550-
has_work = const_expr(not self.is_split_kv) or total_block_count > Int32(0)
2557+
has_work = self.process_work_tile(seqlen, n_block_min, n_block_max)
25512558

25522559
if has_work:
25532560
# Ignore first signal from softmax as no correction is required
@@ -2977,7 +2984,7 @@ def epilogue_s2g(
29772984
if const_expr(self.is_split_kv and block_info.num_splits_dynamic_ptr is not None):
29782985
split_idx = split_idx & 0xFFFF
29792986

2980-
if const_expr(not self.is_split_kv) or n_block_min < n_block_max:
2987+
if self.process_work_tile(seqlen, n_block_min, n_block_max):
29812988
if const_expr(self.is_split_kv):
29822989
mO_cur = seqlen.offset_batch_Q(mO, batch_idx, dim=3)[None, None, head_idx, split_idx]
29832990
else:
@@ -3249,3 +3256,18 @@ def apply_score_mod(
32493256
constant_q_idx=q_idx_logical,
32503257
qhead_per_kvhead=self.qhead_per_kvhead if cutlass.const_expr(self.pack_gqa) else 1,
32513258
)
3259+
3260+
@cute.jit
3261+
def process_work_tile(
3262+
self,
3263+
seqlen_info: SeqlenInfoQK,
3264+
n_block_min: Int32,
3265+
n_block_max: Int32,
3266+
):
3267+
is_varlen_q = seqlen_info.has_cu_seqlens_q or seqlen_info.has_seqused_q
3268+
process_work_tile_k = const_expr(not self.is_split_kv) or n_block_min < n_block_max
3269+
if const_expr(is_varlen_q and not self.use_varlen_scheduler):
3270+
process_work_tile_q = seqlen_info.seqlen_q > 0
3271+
else:
3272+
process_work_tile_q = True
3273+
return process_work_tile_k and process_work_tile_q

flash_attn/cute/interface.py

Lines changed: 11 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -536,6 +536,7 @@ def _flash_attn_fwd(
536536
q_stage = 1
537537

538538
m_block_size_effective = q_stage * tile_m
539+
max_m_blocks_leq_one = seqlen_q_packgqa <= m_block_size_effective
539540
seqlen_k_loaded = max_seqlen_k if not local else max(0, min(max_seqlen_k, (window_size_right or max_seqlen_k) + (window_size_left or max_seqlen_k) + 1 + tile_m))
540541
num_m_blocks = (seqlen_q_packgqa + m_block_size_effective - 1) // m_block_size_effective
541542
total_mblocks = batch_size * num_head_kv * num_m_blocks
@@ -744,6 +745,14 @@ def _flash_attn_fwd(
744745
num_nheads_in_l2 = None
745746
tile_count_semaphore = None
746747

748+
is_static_persistent = (
749+
not causal
750+
and not local
751+
and cu_seqlens_q is None
752+
and seqused_q is None
753+
and not is_split_kv
754+
) or (max_m_blocks_leq_one and not is_split_kv)
755+
747756
compile_key = (
748757
dtype,
749758
head_dim,
@@ -787,6 +796,7 @@ def _flash_attn_fwd(
787796
virtual_batch_idx is not None,
788797
num_nheads_in_l2 is not None,
789798
tile_count_semaphore is not None,
799+
is_static_persistent,
790800
qv is not None,
791801
gather_kv_length,
792802
sparse_kv,
@@ -984,11 +994,7 @@ def _flash_attn_fwd(
984994
m_block_size=tile_m,
985995
n_block_size=tile_n,
986996
q_stage=q_stage,
987-
is_static_persistent=not causal
988-
and not local
989-
and cu_seqlens_q is None
990-
and seqused_q is None
991-
and not is_split_kv,
997+
is_static_persistent=is_static_persistent,
992998
score_mod=score_mod,
993999
mask_mod=mask_mod,
9941000
has_aux_tensors=aux_tensors is not None,

tests/cute/test_flash_attn.py

Lines changed: 78 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -769,10 +769,15 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
769769
# num_splits_vals = [1, 3]
770770
# SplitKV is not supported for hdim >= 192
771771
num_splits_vals = [1, 3] if d < 192 and not DISABLE_SPLIT and not TEST_BWD_ONLY else [1]
772-
for pack_gqa, num_splits in itertools.product(pack_gqa_vals, num_splits_vals):
772+
precompute_metadata_vals = [False, True]
773+
for pack_gqa, num_splits, precompute_metadata in itertools.product(
774+
pack_gqa_vals, num_splits_vals, precompute_metadata_vals
775+
):
773776
# SplitKV not supported on SM90 - skip this iteration
774777
if IS_SM90 and num_splits > 1:
775778
continue
779+
if precompute_metadata and is_fake_mode():
780+
continue
776781
# TODO(wangsiyu): SM100 head_dim=256 2CTA kernel does not support pack_gqa yet.
777782
# pack_gqa=None means auto-enable for GQA/MQA (qhead_per_kvhead > 1)
778783
# Remove this when support is added.
@@ -781,56 +786,79 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
781786
continue
782787
if pack_gqa is None and mha_type != "mha":
783788
continue
784-
out_unpad, lse = flash_attn_varlen_func(
785-
q_unpad if unpad_q else q,
786-
k_unpad if unpad_kv else k,
787-
v_unpad if unpad_kv else v,
788-
cu_seqlens_q=cu_seqlens_q if unpad_q else None,
789-
cu_seqlens_k=cu_seqlens_k if unpad_kv else None,
790-
max_seqlen_q=seqlen_q,
791-
max_seqlen_k=seqlen_k,
792-
seqused_q=seqused_q if not unpad_q else None,
793-
seqused_k=seqused_k if not unpad_kv else None,
794-
causal=causal,
795-
# qv=qv_unpad,
796-
# q_descale=q_descale,
797-
# k_descale=k_descale, v_descale=v_descale,
798-
window_size=window_size,
799-
# attention_chunk=attention_chunk,
800-
learnable_sink=learnable_sink,
801-
softcap=softcap,
802-
num_splits=num_splits,
803-
pack_gqa=pack_gqa,
804-
deterministic=deterministic,
805-
)
806-
out = output_pad_fn(out_unpad) if unpad_q else out_unpad
807-
if is_fake_mode():
808-
# no more flash_attn cutedsl calls for the rest of the loop
809-
# skip data-dependent postprocessing
810-
continue
811-
if query_unused_mask is not None:
812-
out.masked_fill_(q_zero_masking, 0.0)
813-
# When unpad_q=False with seqused_q, the kernel doesn't write positions
814-
# beyond seqused_q, so those contain uninitialized values. Mask them out
815-
# before comparing.
816-
out_cmp, out_ref_cmp, out_pt_cmp = out, out_ref, out_pt
817-
if not unpad_q and seqused_q is not None:
818-
seqused_mask = torch.arange(seqlen_q, device=device)[None, :] < seqused_q[:, None]
819-
seqused_mask = rearrange(seqused_mask, "b s -> b s 1 1")
820-
out_cmp = out.clone().masked_fill_(~seqused_mask, 0.0)
821-
out_ref_cmp = out_ref.clone().masked_fill_(~seqused_mask, 0.0)
822-
out_pt_cmp = out_pt.clone().masked_fill_(~seqused_mask, 0.0)
823-
print(f"Output max diff: {(out_cmp - out_ref_cmp).abs().max().item()}")
824-
print(f"Output mean diff: {(out_cmp - out_ref_cmp).abs().mean().item()}")
825-
# if not causal:
826-
# print(f"LSE max diff: {(lse - lse_ref).abs().max().item()}")
827-
# breakpoint()
789+
if precompute_metadata:
790+
scheduler_metadata = get_scheduler_metadata(
791+
num_batch=batch_size,
792+
max_seqlen_q=seqlen_q,
793+
max_seqlen_k=seqlen_k,
794+
nheads=nheads,
795+
nheads_kv=nheads_kv,
796+
headdim=d,
797+
headdim_v=dv,
798+
num_splits=num_splits,
799+
tile_m=128,
800+
tile_n=128,
801+
causal=causal,
802+
cu_seqlens_q=cu_seqlens_q if unpad_q else None,
803+
cu_seqlens_k=cu_seqlens_k if unpad_kv else None,
804+
seqused_q=seqused_q if not unpad_q else None,
805+
seqused_k=seqused_k if not unpad_kv else None,
806+
)
807+
else:
808+
scheduler_metadata = None
809+
# Repeat to exercise metadata reuse across calls.
810+
for _ in range(1 if not precompute_metadata else 2):
811+
out_unpad, lse = flash_attn_varlen_func(
812+
q_unpad if unpad_q else q,
813+
k_unpad if unpad_kv else k,
814+
v_unpad if unpad_kv else v,
815+
cu_seqlens_q=cu_seqlens_q if unpad_q else None,
816+
cu_seqlens_k=cu_seqlens_k if unpad_kv else None,
817+
max_seqlen_q=seqlen_q,
818+
max_seqlen_k=seqlen_k,
819+
seqused_q=seqused_q if not unpad_q else None,
820+
seqused_k=seqused_k if not unpad_kv else None,
821+
causal=causal,
822+
# qv=qv_unpad,
823+
# q_descale=q_descale,
824+
# k_descale=k_descale, v_descale=v_descale,
825+
window_size=window_size,
826+
# attention_chunk=attention_chunk,
827+
learnable_sink=learnable_sink,
828+
softcap=softcap,
829+
scheduler_metadata=scheduler_metadata,
830+
num_splits=num_splits,
831+
pack_gqa=pack_gqa,
832+
deterministic=deterministic,
833+
)
834+
out = output_pad_fn(out_unpad) if unpad_q else out_unpad
835+
if is_fake_mode():
836+
# no more flash_attn cutedsl calls for the rest of the loop
837+
# skip data-dependent postprocessing
838+
continue
839+
if query_unused_mask is not None:
840+
out.masked_fill_(q_zero_masking, 0.0)
841+
# When unpad_q=False with seqused_q, the kernel doesn't write positions
842+
# beyond seqused_q, so those contain uninitialized values. Mask them out
843+
# before comparing.
844+
out_cmp, out_ref_cmp, out_pt_cmp = out, out_ref, out_pt
845+
if not unpad_q and seqused_q is not None:
846+
seqused_mask = torch.arange(seqlen_q, device=device)[None, :] < seqused_q[:, None]
847+
seqused_mask = rearrange(seqused_mask, "b s -> b s 1 1")
848+
out_cmp = out.clone().masked_fill_(~seqused_mask, 0.0)
849+
out_ref_cmp = out_ref.clone().masked_fill_(~seqused_mask, 0.0)
850+
out_pt_cmp = out_pt.clone().masked_fill_(~seqused_mask, 0.0)
851+
print(f"Output max diff: {(out_cmp - out_ref_cmp).abs().max().item()}")
852+
print(f"Output mean diff: {(out_cmp - out_ref_cmp).abs().mean().item()}")
853+
# if not causal:
854+
# print(f"LSE max diff: {(lse - lse_ref).abs().max().item()}")
855+
# breakpoint()
828856

829-
# Check that FlashAttention's numerical error is at most 3x the numerical error
830-
# of a Pytorch implementation.
831-
assert (out_cmp - out_ref_cmp).abs().max().item() <= rtol * (
832-
out_pt_cmp - out_ref_cmp
833-
).abs().max().item() + fwd_atol
857+
# Check that FlashAttention's numerical error is at most 3x the numerical error
858+
# of a Pytorch implementation.
859+
assert (out_cmp - out_ref_cmp).abs().max().item() <= rtol * (
860+
out_pt_cmp - out_ref_cmp
861+
).abs().max().item() + fwd_atol
834862

835863
if (
836864
dtype != torch.float8_e4m3fn

0 commit comments

Comments
 (0)