FA4 consumer Blackwell (sm_120) integration: forward + backward + dispatcher fixes - #1

Closed

thad0ctor wants to merge 224 commits into

mainfrom

sm120-integrate

Commits on May 24, 2026

[Cute,Fwd,Sm120] Disable use_tma_O in SM80 base class

authored and
thad0ctor
committed
ead67d3
[Cute,Fwd,Sm120] Use universal smem-store atom for SM80 MMA layout

authored and
thad0ctor
committed
bc67a9c
[Cute,Fwd,Sm120] Default pack_gqa=False on consumer Blackwell

authored and
thad0ctor
committed
c0250ca
Add SM120 TMA forward kernel with warp specialization

authored and
thad0ctor
committed
e595c95
Fix arch override and add SMEM fallback in SM120 TMA forward

authored and
thad0ctor
committed
a6b9896
Move `stream` to end of FlashAttentionForwardSm120Tma.__call__

blake-snc
authored and
thad0ctor
committed
bf10b6d
Add SM80/SM120 block-sparse forward attention support

authored and
thad0ctor
committed
b7be538
Use get_total_block_count utility for block-sparse early exit check

authored and
thad0ctor
committed
3cd8cb6

Commits on May 25, 2026

Commits on May 26, 2026

cute: enable pack_gqa in SM120 backward (interface + kernel port)
thad0ctor
committed
dd8ef9b
cute: keep SM120 bwd pack_gqa opt-in
thad0ctor
committed
8ec45d9
cute: store SM120 pack_gqa dkv accum directly
thad0ctor
committed
78870f9
cute: split SM120 pack_gqa backward m-loop
thad0ctor
committed
e63118b
cute: coalesce SM120 pack_gqa dkv accum zero
thad0ctor
committed
0e821f3
cute: fuse SM120 pack_gqa dkv postprocess
thad0ctor
committed
bf6861f
cute: skip aligned SM120 pack_gqa row guards
thad0ctor
committed
d93408a
cute: keep SM120 GQA forward on nonpacked path
thad0ctor
committed
c43a7b4
cute: restore SM120 packed GQA forward path
thad0ctor
committed
1b7db10
bench: mark SM120 forward current winner
thad0ctor
committed
d2d0ec2
bench: record SM120 current FA2 comparison
thad0ctor
committed
9b46c42
cute: retune SM120 Qwen forward tile
thad0ctor
committed
37f3b27
cute: optimize SM120 local forward dispatch
thad0ctor
committed
f609c64
cute: skip SM120 dense TMA seqlen masks
thad0ctor
committed
0af9a4c
cute: skip SM120 dense base seqlen mask
thad0ctor
committed
eb6e4b8
cute: retune SM120 qpkv5 and local forward tiles
thad0ctor
committed
e65b67a
cute: retune SM120 Qwen forward tiles
thad0ctor
committed
4c0b291
cute: retune SM120 long Qwen forward tiles
thad0ctor
committed
37840d3
cute: retune SM120 long Qwen forward tiles
thad0ctor
committed
14459b8
cute: skip unneeded SM120 causal mask work
thad0ctor
committed
5174f9e
cute: stage SM120 long qpkv5 Q in registers
thad0ctor
committed
7466e10
cute: retune SM120 long qpkv5 forward
thad0ctor
committed
f25e87e

Commits on May 27, 2026

cute: extend SM120 qpkv5 long causal retune
thad0ctor
committed
48c7d4d
cute: tune SM120 qpkv6 TMA stages
thad0ctor
committed
4e895b5
cute: tune SM120 qpkv5 noncausal TMA stages
thad0ctor
committed
ec6dd92
cute: narrow SM120 qpkv5 TMA retune
thad0ctor
committed
ca7f03e
cute: keep SM120 qpkv6 TMA retune noncausal
thad0ctor
committed
08c184f
cute: revert unstable SM120 TMA stage retune
thad0ctor
committed
6a8f9ab
cute: tune SM120 Gemma qpkv4 local forward
thad0ctor
committed
3dc3f22
cute: tune SM120 qpkv5 midseq causal forward
thad0ctor
committed
6b77ace

Commits on May 28, 2026

cute: tune SM120 qpkv6 D256 forward load overlap
thad0ctor
committed
6459f7d
cute: retune SM120 qpkv5 16k forward
thad0ctor
committed
532331f
docs: add SM120 win ledger
thad0ctor
committed
4774a0c
docs: tighten SM120 win ledger scope
thad0ctor
committed
c6051aa
cute: retune SM120 qpkv8 8k forward
thad0ctor
committed
60354fd
docs: record SM120 qpkv6 profiling note
thad0ctor
committed
a986d34
docs: record SM120 qpkv6 hook scheduling result
thad0ctor
committed
bbe6294
docs: record SM120 qpkv6 predicate probes
thad0ctor
committed
3c34c6a
cute: address SM120 review correctness
thad0ctor
committed
fb5d416
cute: close remaining SM120 review gaps
thad0ctor
committed
56310a8
cute: retune SM120 qpkv4 forward lookup
thad0ctor
committed
ced523f
docs: record SM120 qpkv4 lookup validation
thad0ctor
committed
fdd8c12
docs: record SM120 five-repeat broad sweep
thad0ctor
committed
09b885d
docs: record rejected SM120 qpkv2 dense probe
thad0ctor
committed
b9261e6
docs: record SM120 five-repeat rerun
thad0ctor
committed
f6cf982
docs: record SM120 noisy miss repeat
thad0ctor
committed
d6bdf7e
docs: record SM120 ten-repeat broad sweep
thad0ctor
committed
4275428

Commits on May 29, 2026

Tune SM120 D64 MHA forward lookup
thad0ctor
committed
e93f8ef
Tune SM120 D128 qpkv4 forward lookup
thad0ctor
committed
3dcf691
Add SM120 paged KV for D192 and D256
thad0ctor
committed
85d75e0
Fix SM120 local test worktree import
thad0ctor
committed
1b7c366
Add SM120 D256 backward baseline
thad0ctor
committed
7e29608
docs: record SM120 forward tail reruns
thad0ctor
committed
3b236a5
cute: auto-pack SM120 qpkv8 noncausal backward
thad0ctor
committed
610a666
docs: record SM120 D256 N32 backward rejection
thad0ctor
committed
ed94071
cute: auto-pack SM120 qpkv4 S8192 backward
thad0ctor
committed
bd5b29f
cute: expand SM120 qpkv4 S8192 backward auto-pack
thad0ctor
committed
74c0b4d
docs: record SM120 D256 backward follow-up probes
thad0ctor
committed
5bbee40
cute: expand SM120 D256 backward autopack rows
thad0ctor
committed
08ea5fc
docs: record qpkv2 backward autopack extension
thad0ctor
committed
f5b5ea1
cute: tune SM120 qpkv5 forward hooks
thad0ctor
committed
e01e891
docs: record D64 MHA override rejection
thad0ctor
committed
effed34
cute: tune SM120 qpkv8 S4096 forward
thad0ctor
committed
e20ee42
docs: record SM120 forward no-change probes
thad0ctor
committed
76746ef
cute: add SM120 fused dKV profiling gate
thad0ctor
committed
a05b301
cute: oversplit SM120 qpkv4 causal backward
thad0ctor
committed
3b1f43f
cute: fuse SM120 qpkv8 causal dkv postprocess
thad0ctor
committed
1432c92
cute: gate SM120 qpkv8 fused dkv to bf16
thad0ctor
committed
e943809
docs: record SM120 D256 backward retest results
thad0ctor
committed
7d12686
docs: record SM120 D256 source probe rejections
thad0ctor
committed
30db028
tests: harden SM120 review guards
thad0ctor
committed
8dc70d4
cute: split SM120 D256 qpkv6 qpkv8 causal bwd
thad0ctor
committed
229f87f
docs: record SM120 split2 B1 validation
thad0ctor
committed
2d7ce45
docs: record SM120 empty split rejection
thad0ctor
committed
89fe103
cute: tune SM120 qpkv8 h8 causal bwd split
thad0ctor
committed
3b276ae
docs: record SM120 qpkv8 h8 split3 sweep
thad0ctor
committed
17ce38a
docs: record SM120 combined zero rejection
thad0ctor
committed
986d732
docs: record SM120 qpkv4 split rejection
thad0ctor
committed
e59fa22
docs: record SM120 fused dKV rejections
thad0ctor
committed
cf39944
cute: tune SM120 qpkv4 causal pack split
thad0ctor
committed
4f1ea4d
test: update SM120 qpkv4 split policy
thad0ctor
committed
541cd4b
docs: record SM120 qpkv6 bwd NCU
thad0ctor
committed
32f1d31
docs: record SM120 split3 wide rejection
thad0ctor
committed
f5da596
cute: tune SM120 qpkv8 long causal
thad0ctor
committed
10dbd84
cute: extend SM120 qpkv8 long causal tile
thad0ctor
committed
d7f8a5d
docs: record SM120 qpkv5 tile rejection
thad0ctor
committed
eb91374
docs: record SM120 qpkv5 non-TMA rejection
thad0ctor
committed
557ce60
docs: record SM120 D256 schedule rejections
thad0ctor
committed
a11ddad
docs: record SM120 long forward sweep
thad0ctor
committed
a8632f8
docs: record SM120 long forward rejections
thad0ctor
committed
38fb046
docs: record SM120 focused long no-patch probes
thad0ctor
committed
0d709fb
docs: record SM120 post-qpkv8 broad sweep
thad0ctor
committed
d11c0a1
docs: record SM120 S4096 forward rejections
thad0ctor
committed
a0fb2e2
cute: route SM120 qpkv5 S4096 noncausal off TMA
thad0ctor
committed
8320cf4
docs: record SM120 qpkv5 non-TMA sweep
thad0ctor
committed
1ca5ca7
docs: record SM120 qpkv4 pack-off rejection
thad0ctor
committed
6507984
docs: record SM120 qpkv8 noncausal rejections
thad0ctor
committed
1f1667d
cute: skip full causal masks for SM120 D256 bwd rows
thad0ctor
committed
f4bf650
cute: extend SM120 D256 causal mask skip to qpkv2
thad0ctor
committed
e448d3e
cute: reject SM120 varlen block sparsity on SM80 path
thad0ctor
committed
9078c7a
docs: reject SM120 D256 m-load predicate elision
thad0ctor
committed
29ad06c
docs: update SM120 qpkv6 bwd profile baseline
thad0ctor
committed
65b4b5d
docs: reject SM120 PackGQA row-map probe
thad0ctor
committed
fb21d97
docs: record SM120 qpkv8 H16 bwd profile
thad0ctor
committed
a7c754a
docs: reject SM120 D256 non-alias schedule
thad0ctor
committed
682a020
docs: reject SM120 D256 cp.async cache-always
thad0ctor
committed
e09d6a2
docs: reject SM120 B1 qpkv6 mask skip
thad0ctor
committed
1a4a7ab
docs: reject SM120 TMA no-inf-check probe
thad0ctor
committed
9ed3b49
tune SM120 qpkv5 S16384 B1 hooks
thad0ctor
committed
9f29cce
tune SM120 qpkv5 S16384 qregs
thad0ctor
committed
8ed047f
docs: reject SM120 qpkv2 split probe
thad0ctor
committed
51a3c5b
tune SM120 qpkv5 S16384 B2 qregs
thad0ctor
committed
d6bf449
tune SM120 qpkv5 S16384 qregs all batches
thad0ctor
committed
d6dc1e9
tune SM120 qpkv8 S65536 causal
thad0ctor
committed
1278426
tune SM120 D256 qpkv8 qpkv16 noncausal
thad0ctor
committed
83d9968
tune SM120 qpkv8 long noncausal
thad0ctor
committed
246a798
tune SM120 qpkv6 long hooks
thad0ctor
committed
4f7e0f4
limit SM120 qpkv6 hooks to stable rows
thad0ctor
committed
2b988f5
docs: reject SM120 qpkv5 long nc probes
thad0ctor
committed
3db93aa
tune SM120 qpkv6 causal block bounds
thad0ctor
committed
fcb6352
docs: record SM120 qpkv8 causal probe results
thad0ctor
committed
2fedf02
tune SM120 qpkv8 D256 causal qregs
thad0ctor
committed
8f02987
docs: reject SM120 qpkv8 D256 hook probe
thad0ctor
committed
517282c
tune SM120 qpkv6 D256 long qregs
thad0ctor
committed
77a7cb4
tune SM120 qpkv8 D256 causal qregs
thad0ctor
committed
fcc906b
docs: update SM120 win ledger after qregs probes
thad0ctor
committed
8d7c388
tune SM120 qpkv16 D256 causal qregs
thad0ctor
committed
1d2d944
docs: record current SM120 long sweep
thad0ctor
committed
fd8ddd6
docs: reject SM120 qpkv5 S8192 nonTMA probe
thad0ctor
committed
e430ad2
docs: record SM120 qpkv8 D128 profile
thad0ctor
committed
93d200b
docs: record post-qpkv16 broad sweep
thad0ctor
committed
3960316
tune SM120 qpkv6 B2 qregs rows
thad0ctor
committed
43f1ea9
docs: record qpkv6 B2 broad sweep
thad0ctor
committed
641a895
docs: record rejected SM120 forward gap probes
thad0ctor
committed
cf0a7f0
docs: record SM120 pause-point forward sweep
thad0ctor
committed
5cbf8fc
docs: record rejected SM120 D128 forward probes
thad0ctor
committed
3f7b69c
optimize SM120 packed GQA valid rows
thad0ctor
committed
6125775
record SM120 pause benchmark
thad0ctor
committed
af5a57d
tune SM120 qpkv6 S8192 noncausal
thad0ctor
committed
d49655e
tune SM120 qpkv4 backward mask skip
thad0ctor
committed
b782823
docs: record SM120 rejected probes
thad0ctor
committed
5dc126d
docs: record SM120 S2048 backward pause point
thad0ctor
committed
cd1b3e3
tune SM120 qpkv4 S2048 backward
thad0ctor
committed
dfb7a24

Commits on Jun 2, 2026

Commits on Jun 3, 2026

Commits on Jun 4, 2026