Skip to content

Bump nvidia-cutlass-dsl to >=4.5.2 and quack-kernels to >=0.5.0 - #2605

Merged
Johnsonms merged 1 commit into
Dao-AILab:mainfrom
Johnsonms:bump-cutlass-4.5.2-quack-0.5.0
May 30, 2026
Merged

Bump nvidia-cutlass-dsl to >=4.5.2 and quack-kernels to >=0.5.0#2605
Johnsonms merged 1 commit into
Dao-AILab:mainfrom
Johnsonms:bump-cutlass-4.5.2-quack-0.5.0

Conversation

@Johnsonms

@Johnsonms Johnsonms commented May 30, 2026 •
edited
Loading

Copy link
Copy Markdown
Collaborator

Bump nvidia-cutlass-dsl to 4.5.2 and quack-kernels to 0.5.0

Summary

Upgrades the FA4 (flash_attn/cute/) runtime dependencies:

  • nvidia-cutlass-dsl: >=4.4.2 → >=4.5.2 (dependencies + cu13 extra)
  • quack-kernels: >=0.4.0 → >=0.5.0

Per discussion: cutlass 4.5.2 is safe to update, and quack 0.5.0 has been published, so we bump the flash-attention requirement to match. CLAUDE.md is updated to keep the documented dependency versions in sync with the manifest.

This is the Python CuTeDSL dependency only; the csrc/cutlass C++ submodule (legacy FA2/FA3) is unchanged.

Changes

  • flash_attn/cute/pyproject.toml: bump nvidia-cutlass-dsl (dependencies + cu13 extra) and quack-kernels version floors.
  • CLAUDE.md: update the documented dependency versions.

Environment

Common methodology: clocks locked for reproducibility (GB300: nvidia-smi -i 0 -lgc 1965,1965 → effective sustained 1957 MHz, throttle 0x0; reset with -rgc after). --warmup 25 --rep 100. Python 3.12.

GB300 (after) B200
GPU NVIDIA GB300 (Blackwell, SM100) ×4 NVIDIA B200 (Blackwell, SM100) ×1
Driver / CUDA 580.126.20 / CUDA 13.2 (V13.2.78) 580.159.03 / CUDA 13.1 (V13.1.115)
torch 2.12.0a0+0291f960b6.nv26.4 2.11.0a0+eb65b36914.nv26.02
Clock lock 1965 MHz → eff. 1957 MHz 1965 MHz max → locked 1867 MHz (~95%)

Dependency versions are the same on both archs: before = nvidia-cutlass-dsl 4.4.2, after = nvidia-cutlass-dsl 4.5.2 (base + cu13) / quack-kernels 0.5.0 / apache-tvm-ffi 0.1.10.

Correctness

Verified on the environment above with nvidia-cutlass-dsl==4.5.2 / quack-kernels==0.5.0:

  • Deps resolve cleanly (no conflicts); from flash_attn.cute import flash_attn_func, flash_attn_varlen_func imports OK against cutlass 4.5.2.
  • Representative GPU sample of tests/cute/test_flash_attn.py::test_flash_attn_output (12 cases spanning head_dim 64/96/128/192, causal both, MHA/GQA/MQA, fwd + sm100 bwd): 6 passed, 6 skipped, 0 failed.

Curated key-test sample (~467 cases across test_flash_attn, _varlen, mask_mod, mask_mod_varlen, score_mod, block_sparsity), GB300:

after (4.5.2):  250 passed, 212 skipped, 5 failed

Same sample re-run on B200 (with byte-identical sampled test IDs) confirms the GB300 finding — both before and after cutlass versions report the same 5 failures in test_score_mod.py, no other deltas:

b200_before (4.4.2): 250 passed, 212 skipped, 5 failed  (0:42:10)
b200_after  (4.5.2): 250 passed, 212 skipped, 5 failed  (1:16:07)

All 5 failures are in test_score_mod.py (the tight score_mod-vs-flex-attention tolerance check: cute_error <= 2× torch_error + atol); no failures in test_flash_attn, varlen, mask_mod, or block_sparsity. None are caused by this bump — confirmed against the before stack (cutlass 4.4.2 + quack 0.4.0):

  • 4 of 5 also fail on 4.4.2 → pre-existing (test_cute_vs_flex_attention[score_mod_pair{2,3,6}], test_score_mod_with_paged_kvcache[score_mod_pair6...]).
  • 1 backward case (test_cute_vs_flex_attention_backward[False-score_mod_triple1-dtype1-64-256-128]) passes in isolation on both 4.4.2 and 4.5.2 with identical error margins (dQ 3.91e-3, dK 9.28e-3, dV 1.56e-2) → the single failure in the full sample was test-ordering flakiness, not a numerical change from the bump.

Net: across the curated sample on both Blackwell arch and both cutlass versions, the pass/skip/fail counts are identical and every failure is a pre-existing test_score_mod.py tolerance case unrelated to the bump. The full parametrized suite (~417k cases) runs in CI.

Benchmark — before/after the cutlass bump, per architecture

Purpose: confirm the nvidia-cutlass-dsl 4.4.2 → 4.5.2 bump does not regress FA4 performance. Before = cutlass 4.4.2, after = cutlass 4.5.2, measured on each Blackwell arch. FA4 forward+backward, dense + varlen, head dims {64,128,256} × seqlens {1k,2k,4k,8k,16k}. Cell = ms / TFLOPS / MFU% / TB/s / BW%. Heads scale with head_dim (hd64→32, hd128→16, hd256→8).

cuDNN backward is omitted because cuDNN doesn't support head_dim > 128 backward (cudnnGraphNotSupportedError at hd256), so the FA4 baseline is FA4-only fwd+bwd; cuDNN is kept as a forward-only reference.

All four quadrants — {GB300, B200} × {before, after} — are reported below, each as collapsible dense / varlen / cuDNN-reference tables.

GB300 (SM100)

After — cutlass 4.5.2

FA4 peaks: Dense FWD 1882 TFLOPS / 75.3% MFU, Dense BWD 1530 / 61.2%, Varlen FWD 1745 / 69.8%, Varlen BWD 1508 / 60.3% (all hd128, non-causal, 16k). FA4 beats the cuDNN forward reference clearly at hd256 (~1.7× at long seqlens, e.g. dense nc 16k: 1792 vs 1029 TFLOPS) and slightly at hd128 non-causal. Weak spot: short-seqlen varlen hd256 — fwd 177 TFLOPS (7.1% MFU), bwd 43 TFLOPS (1.7% MFU, ~16 ms) at seqlen 1024, recovering only at long seqlens (small-tile/scheduling inefficiency; baseline characteristic, follow-up).

GB300 after (4.5.2) — FA4 dense, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024   0.25/1087/43.5%/2.12/26.5%
       64  False    16   2048   0.47/1167/46.7%/1.14/14.2%
       64  False     8   4096    0.91/1213/48.5%/0.59/7.4%
       64  False     4   8192    1.79/1229/49.2%/0.30/3.7%
       64  False     2  16384    3.54/1243/49.7%/0.15/1.9%
       64   True    32   1024    0.25/542/21.7%/2.12/26.5%
       64   True    16   2048    0.35/785/31.4%/1.53/19.2%
       64   True     8   4096   0.54/1011/40.4%/0.99/12.3%
       64   True     4   8192    0.94/1169/46.8%/0.57/7.1%
       64   True     2  16384    1.76/1251/50.1%/0.31/3.8%
      128  False    32   1024   0.18/1540/61.6%/3.01/37.6%
      128  False    16   2048   0.32/1706/68.2%/1.67/20.8%
      128  False     8   4096   0.61/1796/71.8%/0.88/11.0%
      128  False     4   8192    1.20/1837/73.5%/0.45/5.6%
      128  False     2  16384    2.34/1882/75.3%/0.23/2.9%
      128   True    32   1024    0.17/829/33.2%/3.24/40.5%
      128   True    16   2048   0.23/1172/46.9%/2.29/28.6%
      128   True     8   4096   0.38/1431/57.2%/1.40/17.5%
      128   True     4   8192    0.68/1610/64.4%/0.79/9.8%
      128   True     2  16384    1.28/1718/68.7%/0.42/5.2%
      256  False    32   1024   0.27/1037/41.5%/2.03/25.3%
      256  False    16   2048   0.40/1369/54.8%/1.34/16.7%
      256  False     8   4096    0.70/1571/62.8%/0.77/9.6%
      256  False     4   8192    1.29/1704/68.2%/0.42/5.2%
      256  False     2  16384    2.48/1777/71.1%/0.22/2.7%
      256   True    32   1024    0.21/653/26.1%/2.55/31.9%
      256   True    16   2048    0.28/995/39.8%/1.94/24.3%
      256   True     8   4096   0.42/1323/52.9%/1.29/16.2%
      256   True     4   8192    0.72/1531/61.3%/0.75/9.3%
      256   True     2  16384    1.34/1644/65.7%/0.40/5.0%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    1.04/661/26.5%/0.90/11.3%
       64  False    16   2048     1.69/814/32.5%/0.56/7.0%
       64  False     8   4096     2.99/921/36.8%/0.31/3.9%
       64  False     4   8192     5.52/996/39.9%/0.17/2.1%
       64  False     2  16384   10.66/1031/41.3%/0.09/1.1%
       64   True    32   1024    0.77/447/17.9%/1.22/15.3%
       64   True    16   2048    1.15/598/23.9%/0.82/10.2%
       64   True     8   4096     1.91/718/28.7%/0.49/6.1%
       64   True     4   8192     3.44/799/31.9%/0.27/3.4%
       64   True     2  16384     6.48/849/33.9%/0.15/1.8%
      128  False    32   1024    0.73/942/37.7%/1.29/16.1%
      128  False    16   2048   1.14/1207/48.3%/0.82/10.3%
      128  False     8   4096    1.99/1384/55.4%/0.47/5.9%
      128  False     4   8192    3.69/1492/59.7%/0.25/3.2%
      128  False     2  16384    7.19/1530/61.2%/0.13/1.6%
      128   True    32   1024    0.56/610/24.4%/1.67/20.8%
      128   True    16   2048    0.76/906/36.2%/1.24/15.5%
      128   True     8   4096    1.18/1163/46.5%/0.79/9.9%
      128   True     4   8192    2.02/1359/54.3%/0.46/5.8%
      128   True     2  16384    3.79/1451/58.1%/0.25/3.1%
      256  False    32   1024    1.02/671/26.9%/0.92/11.5%
      256  False    16   2048     1.56/879/35.2%/0.60/7.5%
      256  False     8   4096    2.65/1038/41.5%/0.35/4.4%
      256  False     4   8192    4.85/1134/45.4%/0.19/2.4%
      256  False     2  16384    9.32/1180/47.2%/0.10/1.3%
      256   True    32   1024    0.80/430/17.2%/1.18/14.7%
      256   True    16   2048    1.08/639/25.6%/0.87/10.9%
      256   True     8   4096     1.65/831/33.2%/0.57/7.1%
      256   True     4   8192     2.80/983/39.3%/0.34/4.2%
      256   True     2  16384    5.24/1049/42.0%/0.18/2.2%
GB300 after (4.5.2) — FA4 varlen, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    0.35/776/31.1%/1.52/19.0%
       64  False    16   2048    0.56/982/39.3%/0.96/12.0%
       64  False     8   4096    0.99/1112/44.5%/0.54/6.8%
       64  False     4   8192    1.87/1179/47.1%/0.29/3.6%
       64  False     2  16384    3.61/1218/48.7%/0.15/1.9%
       64   True    32   1024    0.30/460/18.4%/1.80/22.5%
       64   True    16   2048    0.39/699/28.0%/1.36/17.1%
       64   True     8   4096    0.61/904/36.2%/0.88/11.0%
       64   True     4   8192    1.03/1065/42.6%/0.52/6.5%
       64   True     2  16384    1.90/1160/46.4%/0.28/3.5%
      128  False    32   1024   0.24/1145/45.8%/2.24/28.0%
      128  False    16   2048   0.40/1382/55.3%/1.35/16.9%
      128  False     8   4096    0.71/1559/62.3%/0.76/9.5%
      128  False     4   8192    1.34/1646/65.8%/0.40/5.0%
      128  False     2  16384    2.52/1745/69.8%/0.21/2.7%
      128   True    32   1024    0.18/785/31.4%/3.06/38.3%
      128   True    16   2048   0.24/1132/45.3%/2.21/27.6%
      128   True     8   4096   0.40/1381/55.2%/1.35/16.9%
      128   True     4   8192    0.69/1590/63.6%/0.78/9.7%
      128   True     2  16384    1.29/1707/68.3%/0.42/5.2%
      256  False    32   1024      1.56/177/7.1%/0.35/4.3%
      256  False    16   2048     1.25/441/17.6%/0.43/5.4%
      256  False     8   4096     1.28/857/34.3%/0.42/5.2%
      256  False     4   8192    1.70/1297/51.9%/0.32/4.0%
      256  False     2  16384    2.83/1555/62.2%/0.19/2.4%
      256   True    32   1024       1.52/91/3.6%/0.35/4.4%
      256   True    16   2048      1.15/238/9.5%/0.47/5.8%
      256   True     8   4096     1.07/515/20.6%/0.50/6.3%
      256   True     4   8192     1.20/916/36.6%/0.45/5.6%
      256   True     2  16384    1.68/1307/52.3%/0.32/4.0%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    1.18/585/23.4%/0.80/10.0%
       64  False    16   2048     1.88/730/29.2%/0.50/6.2%
       64  False     8   4096     3.34/824/33.0%/0.28/3.5%
       64  False     4   8192     6.28/876/35.0%/0.15/1.9%
       64  False     2  16384    12.15/905/36.2%/0.08/1.0%
       64   True    32   1024    0.88/392/15.7%/1.07/13.4%
       64   True    16   2048     1.25/548/21.9%/0.75/9.4%
       64   True     8   4096     2.02/679/27.2%/0.46/5.8%
       64   True     4   8192     3.56/771/30.9%/0.26/3.3%
       64   True     2  16384     6.62/830/33.2%/0.14/1.8%
      128  False    32   1024    0.84/816/32.7%/1.12/14.0%
      128  False    16   2048    1.22/1129/45.1%/0.77/9.6%
      128  False     8   4096    2.07/1328/53.1%/0.45/5.7%
      128  False     4   8192    3.81/1444/57.8%/0.25/3.1%
      128  False     2  16384    7.29/1508/60.3%/0.13/1.6%
      128   True    32   1024    0.67/515/20.6%/1.41/17.6%
      128   True    16   2048    0.84/822/32.9%/1.12/14.1%
      128   True     8   4096    1.23/1117/44.7%/0.76/9.5%
      128   True     4   8192    2.07/1329/53.1%/0.45/5.7%
      128   True     2  16384    3.74/1469/58.8%/0.25/3.1%
      256  False    32   1024      16.03/43/1.7%/0.06/0.7%
      256  False    16   2048     13.64/101/4.0%/0.07/0.9%
      256  False     8   4096     12.82/214/8.6%/0.07/0.9%
      256  False     4   8192    13.28/414/16.6%/0.07/0.9%
      256  False     2  16384    14.98/734/29.4%/0.06/0.8%
      256   True    32   1024       4.38/79/3.1%/0.21/2.7%
      256   True    16   2048      2.84/242/9.7%/0.33/4.1%
      256   True     8   4096     2.52/545/21.8%/0.37/4.7%
      256   True     4   8192     3.24/848/33.9%/0.29/3.6%
      256   True     2  16384    5.49/1002/40.1%/0.17/2.1%
GB300 after (4.5.2) — FA4 vs cuDNN dense, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.28/979/39.2%/1.91/23.9%   0.27/1025/41.0%/2.00/25.0%
       64  False    16   2048   0.51/1069/42.7%/1.04/13.0%   0.50/1095/43.8%/1.07/13.4%
       64  False     8   4096    0.99/1113/44.5%/0.54/6.8%    0.97/1128/45.1%/0.55/6.9%
       64  False     4   8192    1.92/1147/45.9%/0.28/3.5%    1.92/1144/45.8%/0.28/3.5%
       64  False     2  16384    3.80/1157/46.3%/0.14/1.8%    3.81/1153/46.1%/0.14/1.8%
       64   True    32   1024    0.22/627/25.1%/2.45/30.6%    0.25/544/21.8%/2.13/26.6%
       64   True    16   2048    0.33/827/33.1%/1.61/20.2%    0.35/785/31.4%/1.53/19.2%
       64   True     8   4096    0.56/977/39.1%/0.95/11.9%   0.54/1012/40.5%/0.99/12.4%
       64   True     4   8192    1.03/1072/42.9%/0.52/6.5%    0.94/1167/46.7%/0.57/7.1%
       64   True     2  16384    1.96/1120/44.8%/0.27/3.4%    1.77/1246/49.8%/0.30/3.8%
      128  False    32   1024   0.19/1428/57.1%/2.79/34.9%   0.18/1537/61.5%/3.00/37.5%
      128  False    16   2048   0.34/1616/64.6%/1.58/19.7%   0.32/1710/68.4%/1.67/20.9%
      128  False     8   4096   0.64/1712/68.5%/0.84/10.4%   0.61/1794/71.8%/0.88/10.9%
      128  False     4   8192    1.25/1764/70.6%/0.43/5.4%    1.19/1848/73.9%/0.45/5.6%
      128  False     2  16384    2.47/1780/71.2%/0.22/2.7%    2.34/1880/75.2%/0.23/2.9%
      128   True    32   1024    0.14/957/38.3%/3.74/46.7%    0.17/828/33.1%/3.23/40.4%
      128   True    16   2048   0.22/1257/50.3%/2.46/30.7%   0.24/1165/46.6%/2.28/28.4%
      128   True     8   4096   0.38/1463/58.5%/1.43/17.9%   0.39/1426/57.0%/1.39/17.4%
      128   True     4   8192   0.67/1641/65.6%/0.80/10.0%    0.68/1613/64.5%/0.79/9.8%
      128   True     2  16384    1.30/1696/67.8%/0.41/5.2%    1.28/1718/68.7%/0.42/5.2%
      256  False    32   1024    0.31/891/35.6%/1.74/21.7%   0.26/1041/41.6%/2.03/25.4%
      256  False    16   2048    0.57/961/38.4%/0.94/11.7%   0.40/1372/54.9%/1.34/16.8%
      256  False     8   4096    1.09/1007/40.3%/0.49/6.1%    0.70/1568/62.7%/0.77/9.6%
      256  False     4   8192    2.13/1033/41.3%/0.25/3.2%    1.28/1715/68.6%/0.42/5.2%
      256  False     2  16384    4.27/1029/41.2%/0.13/1.6%    2.45/1792/71.7%/0.22/2.7%
      256   True    32   1024    0.22/633/25.3%/2.47/30.9%    0.21/654/26.2%/2.56/32.0%
      256   True    16   2048    0.35/778/31.1%/1.52/19.0%    0.28/998/39.9%/1.95/24.4%
      256   True     8   4096    0.61/900/36.0%/0.88/11.0%   0.42/1314/52.6%/1.28/16.0%
      256   True     4   8192     1.13/971/38.8%/0.47/5.9%    0.71/1546/61.8%/0.75/9.4%
      256   True     2  16384    2.19/1003/40.1%/0.24/3.1%    1.33/1654/66.2%/0.40/5.0%
GB300 after (4.5.2) — FA4 vs cuDNN varlen, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.28/980/39.2%/1.91/23.9%    0.35/787/31.5%/1.54/19.2%
       64  False    16   2048   0.51/1073/42.9%/1.05/13.1%    0.55/994/39.8%/0.97/12.1%
       64  False     8   4096    0.98/1119/44.8%/0.55/6.8%    0.97/1130/45.2%/0.55/6.9%
       64  False     4   8192    1.91/1151/46.1%/0.28/3.5%    1.83/1203/48.1%/0.29/3.7%
       64  False     2  16384    3.79/1159/46.4%/0.14/1.8%    3.56/1236/49.5%/0.15/1.9%
       64   True    32   1024    0.22/627/25.1%/2.45/30.6%    0.29/475/19.0%/1.85/23.2%
       64   True    16   2048    0.33/826/33.1%/1.61/20.2%    0.37/733/29.3%/1.43/17.9%
       64   True     8   4096    0.56/977/39.1%/0.95/11.9%    0.57/963/38.5%/0.94/11.8%
       64   True     4   8192    1.02/1074/42.9%/0.52/6.6%    0.97/1131/45.3%/0.55/6.9%
       64   True     2  16384    1.97/1117/44.7%/0.27/3.4%    1.79/1229/49.2%/0.30/3.8%
      128  False    32   1024   0.19/1427/57.1%/2.79/34.8%   0.24/1151/46.0%/2.25/28.1%
      128  False    16   2048   0.34/1600/64.0%/1.56/19.5%   0.40/1386/55.4%/1.35/16.9%
      128  False     8   4096   0.65/1701/68.0%/0.83/10.4%    0.71/1554/62.1%/0.76/9.5%
      128  False     4   8192    1.24/1772/70.9%/0.43/5.4%    1.32/1662/66.5%/0.41/5.1%
      128  False     2  16384    2.45/1793/71.7%/0.22/2.7%    2.50/1762/70.5%/0.22/2.7%
      128   True    32   1024    0.14/958/38.3%/3.74/46.8%    0.17/788/31.5%/3.08/38.5%
      128   True    16   2048   0.22/1260/50.4%/2.46/30.8%   0.24/1136/45.4%/2.22/27.7%
      128   True     8   4096   0.38/1451/58.0%/1.42/17.7%   0.40/1382/55.3%/1.35/16.9%
      128   True     4   8192    0.68/1626/65.1%/0.79/9.9%    0.69/1591/63.6%/0.78/9.7%
      128   True     2  16384    1.32/1671/66.8%/0.41/5.1%    1.30/1693/67.7%/0.41/5.2%
      256  False    32   1024    0.31/887/35.5%/1.73/21.7%      1.56/177/7.1%/0.35/4.3%
      256  False    16   2048    0.57/961/38.5%/0.94/11.7%     1.25/441/17.6%/0.43/5.4%
      256  False     8   4096    1.09/1005/40.2%/0.49/6.1%     1.28/856/34.2%/0.42/5.2%
      256  False     4   8192    2.14/1025/41.0%/0.25/3.1%    1.69/1299/52.0%/0.32/4.0%
      256  False     2  16384    4.30/1024/41.0%/0.12/1.6%    2.84/1549/62.0%/0.19/2.4%
      256   True    32   1024    0.22/633/25.3%/2.47/30.9%       1.52/91/3.6%/0.35/4.4%
      256   True    16   2048    0.35/783/31.3%/1.53/19.1%      1.16/238/9.5%/0.46/5.8%
      256   True     8   4096    0.61/900/36.0%/0.88/11.0%     1.07/515/20.6%/0.50/6.3%
      256   True     4   8192     1.14/962/38.5%/0.47/5.9%     1.20/915/36.6%/0.45/5.6%
      256   True     2  16384     2.20/998/39.9%/0.24/3.0%    1.68/1308/52.3%/0.32/4.0%

Before — cutlass 4.4.2

FA4 peaks: Dense FWD 1886 TFLOPS, Dense BWD 1520, Varlen FWD 1749, Varlen BWD 1488 (peak across configs).

Before → after delta (GB300): no regression — flat within run-to-run noise: Dense FWD 1886→1882, Dense BWD 1520→1530, Varlen FWD 1749→1745, Varlen BWD 1488→1508 TFLOPS.

gb300_before — FA4 dense, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024   0.25/1086/43.5%/2.12/26.5%
       64  False    16   2048   0.47/1172/46.9%/1.14/14.3%
       64  False     8   4096    0.91/1210/48.4%/0.59/7.4%
       64  False     4   8192    1.78/1234/49.3%/0.30/3.8%
       64  False     2  16384    3.55/1240/49.6%/0.15/1.9%
       64   True    32   1024    0.25/543/21.7%/2.12/26.5%
       64   True    16   2048    0.35/786/31.4%/1.53/19.2%
       64   True     8   4096   0.54/1012/40.5%/0.99/12.4%
       64   True     4   8192    0.94/1167/46.7%/0.57/7.1%
       64   True     2  16384    1.76/1248/49.9%/0.30/3.8%
      128  False    32   1024   0.18/1538/61.5%/3.00/37.6%
      128  False    16   2048   0.32/1703/68.1%/1.66/20.8%
      128  False     8   4096   0.61/1806/72.2%/0.88/11.0%
      128  False     4   8192    1.19/1849/73.9%/0.45/5.6%
      128  False     2  16384    2.33/1886/75.5%/0.23/2.9%
      128   True    32   1024    0.17/829/33.1%/3.24/40.5%
      128   True    16   2048   0.24/1167/46.7%/2.28/28.5%
      128   True     8   4096   0.38/1444/57.8%/1.41/17.6%
      128   True     4   8192    0.69/1588/63.5%/0.78/9.7%
      128   True     2  16384    1.28/1718/68.7%/0.42/5.2%
      256  False    32   1024   0.27/1037/41.5%/2.02/25.3%
      256  False    16   2048   0.40/1363/54.5%/1.33/16.6%
      256  False     8   4096    0.70/1579/63.2%/0.77/9.6%
      256  False     4   8192    1.29/1707/68.3%/0.42/5.2%
      256  False     2  16384    2.42/1818/72.7%/0.22/2.8%
      256   True    32   1024    0.21/654/26.2%/2.56/31.9%
      256   True    16   2048    0.28/996/39.8%/1.94/24.3%
      256   True     8   4096   0.42/1318/52.7%/1.29/16.1%
      256   True     4   8192    0.71/1541/61.6%/0.75/9.4%
      256   True     2  16384    1.33/1649/66.0%/0.40/5.0%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    1.04/662/26.5%/0.90/11.3%
       64  False    16   2048     1.69/814/32.6%/0.56/7.0%
       64  False     8   4096     2.98/923/36.9%/0.32/3.9%
       64  False     4   8192     5.51/998/39.9%/0.17/2.1%
       64  False     2  16384   10.65/1032/41.3%/0.09/1.1%
       64   True    32   1024    0.77/447/17.9%/1.22/15.3%
       64   True    16   2048    1.15/599/24.0%/0.82/10.2%
       64   True     8   4096     1.91/719/28.8%/0.49/6.1%
       64   True     4   8192     3.44/800/32.0%/0.27/3.4%
       64   True     2  16384     6.47/850/34.0%/0.15/1.8%
      128  False    32   1024    0.73/941/37.7%/1.29/16.1%
      128  False    16   2048   1.14/1211/48.4%/0.83/10.3%
      128  False     8   4096    1.99/1378/55.1%/0.47/5.9%
      128  False     4   8192    3.70/1485/59.4%/0.25/3.2%
      128  False     2  16384    7.23/1520/60.8%/0.13/1.6%
      128   True    32   1024    0.56/613/24.5%/1.68/20.9%
      128   True    16   2048    0.76/906/36.3%/1.24/15.5%
      128   True     8   4096   1.17/1171/46.8%/0.80/10.0%
      128   True     4   8192    2.05/1342/53.7%/0.46/5.7%
      128   True     2  16384    3.76/1462/58.5%/0.25/3.1%
      256  False    32   1024    1.03/667/26.7%/0.91/11.4%
      256  False    16   2048     1.57/874/35.0%/0.60/7.5%
      256  False     8   4096    2.65/1037/41.5%/0.35/4.4%
      256  False     4   8192    4.83/1138/45.5%/0.19/2.4%
      256  False     2  16384    9.36/1174/47.0%/0.10/1.3%
      256   True    32   1024    0.85/404/16.1%/1.10/13.8%
      256   True    16   2048    1.14/603/24.1%/0.82/10.3%
      256   True     8   4096     1.71/804/32.2%/0.55/6.9%
      256   True     4   8192     2.90/948/37.9%/0.32/4.1%
      256   True     2  16384    5.41/1016/40.6%/0.17/2.2%
gb300_before — FA4 varlen, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    0.35/775/31.0%/1.51/18.9%
       64  False    16   2048    0.56/980/39.2%/0.96/12.0%
       64  False     8   4096    0.99/1109/44.4%/0.54/6.8%
       64  False     4   8192    1.86/1181/47.2%/0.29/3.6%
       64  False     2  16384    3.62/1215/48.6%/0.15/1.9%
       64   True    32   1024    0.30/457/18.3%/1.78/22.3%
       64   True    16   2048    0.39/696/27.8%/1.36/17.0%
       64   True     8   4096    0.61/902/36.1%/0.88/11.0%
       64   True     4   8192    1.03/1063/42.5%/0.52/6.5%
       64   True     2  16384    1.89/1164/46.6%/0.28/3.6%
      128  False    32   1024   0.24/1145/45.8%/2.24/28.0%
      128  False    16   2048   0.40/1385/55.4%/1.35/16.9%
      128  False     8   4096    0.71/1553/62.1%/0.76/9.5%
      128  False     4   8192    1.32/1665/66.6%/0.41/5.1%
      128  False     2  16384    2.51/1749/70.0%/0.21/2.7%
      128   True    32   1024    0.18/785/31.4%/3.07/38.3%
      128   True    16   2048   0.24/1142/45.7%/2.23/27.9%
      128   True     8   4096   0.39/1393/55.7%/1.36/17.0%
      128   True     4   8192    0.70/1580/63.2%/0.77/9.6%
      128   True     2  16384    1.29/1711/68.4%/0.42/5.2%
      256  False    32   1024      1.56/176/7.0%/0.34/4.3%
      256  False    16   2048     1.25/439/17.6%/0.43/5.4%
      256  False     8   4096     1.29/855/34.2%/0.42/5.2%
      256  False     4   8192    1.69/1299/52.0%/0.32/4.0%
      256  False     2  16384    2.85/1545/61.8%/0.19/2.4%
      256   True    32   1024       1.52/90/3.6%/0.35/4.4%
      256   True    16   2048      1.15/238/9.5%/0.47/5.8%
      256   True     8   4096     1.07/516/20.6%/0.50/6.3%
      256   True     4   8192     1.20/917/36.7%/0.45/5.6%
      256   True     2  16384    1.67/1315/52.6%/0.32/4.0%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    1.17/586/23.4%/0.80/10.0%
       64  False    16   2048     1.88/731/29.2%/0.50/6.2%
       64  False     8   4096     3.33/825/33.0%/0.28/3.5%
       64  False     4   8192     6.27/877/35.1%/0.15/1.9%
       64  False     2  16384    12.14/906/36.2%/0.08/1.0%
       64   True    32   1024    0.88/393/15.7%/1.07/13.4%
       64   True    16   2048     1.25/548/21.9%/0.75/9.4%
       64   True     8   4096     2.02/679/27.2%/0.46/5.8%
       64   True     4   8192     3.56/772/30.9%/0.26/3.3%
       64   True     2  16384     6.61/832/33.3%/0.14/1.8%
      128  False    32   1024    0.84/818/32.7%/1.12/14.0%
      128  False    16   2048    1.22/1127/45.1%/0.77/9.6%
      128  False     8   4096    2.07/1328/53.1%/0.45/5.7%
      128  False     4   8192    3.81/1444/57.8%/0.25/3.1%
      128  False     2  16384    7.39/1488/59.5%/0.13/1.6%
      128   True    32   1024    0.67/516/20.6%/1.41/17.6%
      128   True    16   2048    0.84/822/32.9%/1.12/14.1%
      128   True     8   4096    1.23/1117/44.7%/0.76/9.5%
      128   True     4   8192    2.06/1337/53.5%/0.46/5.7%
      128   True     2  16384    3.75/1467/58.7%/0.25/3.1%
      256  False    32   1024      16.07/43/1.7%/0.06/0.7%
      256  False    16   2048     13.67/101/4.0%/0.07/0.9%
      256  False     8   4096     12.78/215/8.6%/0.07/0.9%
      256  False     4   8192    13.27/414/16.6%/0.07/0.9%
      256  False     2  16384    14.96/735/29.4%/0.06/0.8%
      256   True    32   1024       4.46/77/3.1%/0.21/2.6%
      256   True    16   2048      2.91/237/9.5%/0.32/4.0%
      256   True     8   4096     2.59/530/21.2%/0.36/4.5%
      256   True     4   8192     3.32/827/33.1%/0.28/3.5%
      256   True     2  16384     5.59/984/39.3%/0.17/2.1%
gb300_before — FA4 vs cuDNN dense, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.28/980/39.2%/1.91/23.9%   0.27/1024/41.0%/2.00/25.0%
       64  False    16   2048   0.51/1074/43.0%/1.05/13.1%   0.50/1094/43.8%/1.07/13.4%
       64  False     8   4096    0.98/1116/44.7%/0.55/6.8%    0.98/1127/45.1%/0.55/6.9%
       64  False     4   8192    1.91/1149/45.9%/0.28/3.5%    1.92/1144/45.8%/0.28/3.5%
       64  False     2  16384    3.81/1155/46.2%/0.14/1.8%    3.81/1154/46.2%/0.14/1.8%
       64   True    32   1024    0.22/627/25.1%/2.45/30.6%    0.25/544/21.8%/2.13/26.6%
       64   True    16   2048    0.33/827/33.1%/1.61/20.2%    0.35/786/31.4%/1.53/19.2%
       64   True     8   4096    0.56/976/39.0%/0.95/11.9%   0.54/1013/40.5%/0.99/12.4%
       64   True     4   8192    1.03/1071/42.9%/0.52/6.5%    0.94/1170/46.8%/0.57/7.1%
       64   True     2  16384    1.96/1120/44.8%/0.27/3.4%    1.76/1252/50.1%/0.31/3.8%
      128  False    32   1024   0.19/1426/57.0%/2.79/34.8%   0.18/1540/61.6%/3.01/37.6%
      128  False    16   2048   0.34/1617/64.7%/1.58/19.7%   0.32/1715/68.6%/1.68/20.9%
      128  False     8   4096   0.65/1691/67.6%/0.83/10.3%   0.62/1784/71.3%/0.87/10.9%
      128  False     4   8192    1.24/1780/71.2%/0.43/5.4%    1.18/1858/74.3%/0.45/5.7%
      128  False     2  16384    2.45/1792/71.7%/0.22/2.7%    2.35/1870/74.8%/0.23/2.9%
      128   True    32   1024    0.14/963/38.5%/3.76/47.0%    0.17/832/33.3%/3.25/40.6%
      128   True    16   2048   0.22/1265/50.6%/2.47/30.9%   0.24/1159/46.4%/2.26/28.3%
      128   True     8   4096   0.38/1456/58.2%/1.42/17.8%   0.39/1421/56.9%/1.39/17.4%
      128   True     4   8192    0.68/1623/64.9%/0.79/9.9%    0.68/1621/64.8%/0.79/9.9%
      128   True     2  16384    1.30/1688/67.5%/0.41/5.2%    1.27/1734/69.3%/0.42/5.3%
      256  False    32   1024    0.31/884/35.4%/1.73/21.6%   0.26/1041/41.6%/2.03/25.4%
      256  False    16   2048    0.57/972/38.9%/0.95/11.9%   0.40/1373/54.9%/1.34/16.8%
      256  False     8   4096    1.10/1002/40.1%/0.49/6.1%    0.70/1571/62.8%/0.77/9.6%
      256  False     4   8192    2.17/1016/40.6%/0.25/3.1%    1.29/1710/68.4%/0.42/5.2%
      256  False     2  16384    4.20/1046/41.8%/0.13/1.6%    2.42/1816/72.6%/0.22/2.8%
      256   True    32   1024    0.22/631/25.3%/2.47/30.8%    0.21/655/26.2%/2.56/32.0%
      256   True    16   2048    0.35/781/31.3%/1.53/19.1%    0.28/998/39.9%/1.95/24.4%
      256   True     8   4096    0.61/902/36.1%/0.88/11.0%   0.42/1318/52.7%/1.29/16.1%
      256   True     4   8192     1.14/967/38.7%/0.47/5.9%    0.71/1545/61.8%/0.75/9.4%
      256   True     2  16384    2.19/1003/40.1%/0.24/3.1%    1.32/1662/66.5%/0.41/5.1%
gb300_before — FA4 vs cuDNN varlen, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.28/980/39.2%/1.91/23.9%    0.35/784/31.4%/1.53/19.1%
       64  False    16   2048   0.51/1072/42.9%/1.05/13.1%    0.56/989/39.6%/0.97/12.1%
       64  False     8   4096    0.98/1118/44.7%/0.55/6.8%    0.98/1120/44.8%/0.55/6.8%
       64  False     4   8192    1.92/1147/45.9%/0.28/3.5%    1.84/1193/47.7%/0.29/3.6%
       64  False     2  16384    3.80/1157/46.3%/0.14/1.8%    3.59/1225/49.0%/0.15/1.9%
       64   True    32   1024    0.22/627/25.1%/2.45/30.6%    0.30/463/18.5%/1.81/22.6%
       64   True    16   2048    0.33/826/33.1%/1.61/20.2%    0.39/706/28.2%/1.38/17.2%
       64   True     8   4096    0.56/977/39.1%/0.95/11.9%    0.60/917/36.7%/0.90/11.2%
       64   True     4   8192    1.03/1071/42.8%/0.52/6.5%    1.02/1079/43.2%/0.53/6.6%
       64   True     2  16384    1.97/1115/44.6%/0.27/3.4%    1.87/1176/47.0%/0.29/3.6%
      128  False    32   1024   0.19/1428/57.1%/2.79/34.9%   0.24/1150/46.0%/2.25/28.1%
      128  False    16   2048   0.34/1594/63.7%/1.56/19.5%   0.40/1391/55.7%/1.36/17.0%
      128  False     8   4096   0.64/1709/68.4%/0.83/10.4%    0.71/1557/62.3%/0.76/9.5%
      128  False     4   8192    1.27/1734/69.4%/0.42/5.3%    1.33/1653/66.1%/0.40/5.0%
      128  False     2  16384    2.45/1792/71.7%/0.22/2.7%    2.46/1790/71.6%/0.22/2.7%
      128   True    32   1024    0.14/959/38.4%/3.75/46.8%    0.17/788/31.5%/3.08/38.5%
      128   True    16   2048   0.22/1250/50.0%/2.44/30.5%   0.24/1135/45.4%/2.22/27.7%
      128   True     8   4096   0.37/1467/58.7%/1.43/17.9%   0.40/1382/55.3%/1.35/16.9%
      128   True     4   8192   0.67/1635/65.4%/0.80/10.0%    0.70/1580/63.2%/0.77/9.6%
      128   True     2  16384    1.31/1684/67.4%/0.41/5.1%    1.28/1715/68.6%/0.42/5.2%
      256  False    32   1024    0.31/888/35.5%/1.73/21.7%      1.56/177/7.1%/0.35/4.3%
      256  False    16   2048    0.57/968/38.7%/0.95/11.8%     1.25/440/17.6%/0.43/5.4%
      256  False     8   4096    1.09/1010/40.4%/0.49/6.2%     1.28/856/34.2%/0.42/5.2%
      256  False     4   8192    2.13/1030/41.2%/0.25/3.1%    1.69/1299/52.0%/0.32/4.0%
      256  False     2  16384    4.24/1038/41.5%/0.13/1.6%    2.83/1555/62.2%/0.19/2.4%
      256   True    32   1024    0.22/632/25.3%/2.47/30.9%       1.52/91/3.6%/0.35/4.4%
      256   True    16   2048    0.35/783/31.3%/1.53/19.1%      1.15/238/9.5%/0.47/5.8%
      256   True     8   4096    0.61/898/35.9%/0.88/11.0%     1.07/516/20.6%/0.50/6.3%
      256   True     4   8192     1.14/962/38.5%/0.47/5.9%     1.20/916/36.6%/0.45/5.6%
      256   True     2  16384    2.20/1000/40.0%/0.24/3.1%    1.67/1314/52.6%/0.32/4.0%

B200 (SM100)

After — cutlass 4.5.2

FA4 peaks: Dense FWD 1410 TFLOPS / 62.7% MFU (hd256), Dense BWD 1221 / 54.3% (hd128), Varlen FWD 1330 / 59.1% (hd128), Varlen BWD 1203 / 53.5% (hd128) — all non-causal at seqlen 16k. FA4 beats the cuDNN forward reference clearly at hd256 (~1.74× at long seqlens, e.g. dense nc 16k: 1395 vs 803 TFLOPS) and roughly matches cuDNN at hd128 non-causal. Weak spot (same as GB300): short-seqlen varlen hd256 — fwd 147 TFLOPS (6.5% MFU), bwd 38 TFLOPS (1.7% MFU, ~18 ms) at seqlen 1024, recovering only at long seqlens (small-tile/scheduling inefficiency; baseline characteristic).

b200_after — FA4 dense, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    0.39/702/31.2%/1.37/17.1%
       64  False    16   2048     0.75/733/32.6%/0.72/9.0%
       64  False     8   4096     1.47/749/33.3%/0.37/4.6%
       64  False     4   8192     2.90/758/33.7%/0.18/2.3%
       64  False     2  16384     5.77/762/33.9%/0.09/1.2%
       64   True    32   1024    0.32/423/18.8%/1.65/20.7%
       64   True    16   2048    0.47/579/25.7%/1.13/14.1%
       64   True     8   4096     0.77/713/31.7%/0.70/8.7%
       64   True     4   8192     1.39/792/35.2%/0.39/4.8%
       64   True     2  16384     2.64/834/37.1%/0.20/2.5%
      128  False    32   1024   0.23/1205/53.6%/2.35/29.4%
      128  False    16   2048   0.43/1279/56.8%/1.25/15.6%
      128  False     8   4096    0.83/1328/59.0%/0.65/8.1%
      128  False     4   8192    1.62/1358/60.4%/0.33/4.1%
      128  False     2  16384    3.23/1362/60.5%/0.17/2.1%
      128   True    32   1024    0.20/697/31.0%/2.72/34.0%
      128   True    16   2048    0.30/927/41.2%/1.81/22.6%
      128   True     8   4096   0.50/1090/48.4%/1.06/13.3%
      128   True     4   8192    0.91/1202/53.4%/0.59/7.3%
      128   True     2  16384    1.72/1276/56.7%/0.31/3.9%
      256  False    32   1024    0.30/921/41.0%/1.80/22.5%
      256  False    16   2048   0.49/1124/50.0%/1.10/13.7%
      256  False     8   4096    0.86/1275/56.6%/0.62/7.8%
      256  False     4   8192    1.63/1348/59.9%/0.33/4.1%
      256  False     2  16384    3.12/1410/62.7%/0.17/2.2%
      256   True    32   1024    0.24/583/25.9%/2.28/28.5%
      256   True    16   2048    0.33/837/37.2%/1.63/20.4%
      256   True     8   4096   0.52/1063/47.2%/1.04/13.0%
      256   True     4   8192    0.91/1206/53.6%/0.59/7.4%
      256   True     2  16384    1.75/1254/55.7%/0.31/3.8%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024     1.26/545/24.2%/0.75/9.3%
       64  False    16   2048     2.14/642/28.6%/0.44/5.5%
       64  False     8   4096     3.91/703/31.2%/0.24/3.0%
       64  False     4   8192     7.50/733/32.6%/0.13/1.6%
       64  False     2  16384    14.74/746/33.1%/0.06/0.8%
       64   True    32   1024    0.89/385/17.1%/1.05/13.2%
       64   True    16   2048     1.35/509/22.6%/0.70/8.7%
       64   True     8   4096     2.27/607/27.0%/0.41/5.2%
       64   True     4   8192     4.10/670/29.8%/0.23/2.9%
       64   True     2  16384     7.78/707/31.4%/0.12/1.5%
      128  False    32   1024    0.85/810/36.0%/1.11/13.8%
      128  False    16   2048     1.38/994/44.2%/0.68/8.5%
      128  False     8   4096    2.45/1124/50.0%/0.38/4.8%
      128  False     4   8192    4.66/1179/52.4%/0.20/2.5%
      128  False     2  16384    9.00/1221/54.3%/0.10/1.3%
      128   True    32   1024    0.63/543/24.1%/1.48/18.6%
      128   True    16   2048    0.92/751/33.4%/1.03/12.8%
      128   True     8   4096     1.45/945/42.0%/0.65/8.1%
      128   True     4   8192    2.49/1103/49.0%/0.38/4.7%
      128   True     2  16384    4.83/1139/50.6%/0.19/2.4%
      256  False    32   1024    1.13/609/27.1%/0.83/10.4%
      256  False    16   2048     1.84/745/33.1%/0.51/6.4%
      256  False     8   4096     3.29/836/37.2%/0.29/3.6%
      256  False     4   8192     6.14/895/39.8%/0.15/1.9%
      256  False     2  16384    11.82/931/41.4%/0.08/1.0%
      256   True    32   1024    0.88/388/17.3%/1.06/13.3%
      256   True    16   2048     1.26/547/24.3%/0.75/9.3%
      256   True     8   4096     2.00/686/30.5%/0.47/5.9%
      256   True     4   8192     3.56/772/34.3%/0.26/3.3%
      256   True     2  16384     6.78/810/36.0%/0.14/1.7%
b200_after — FA4 varlen, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    0.44/619/27.5%/1.21/15.1%
       64  False    16   2048     0.73/755/33.5%/0.74/9.2%
       64  False     8   4096     1.33/825/36.7%/0.40/5.0%
       64  False     4   8192     2.55/863/38.3%/0.21/2.6%
       64  False     2  16384     5.00/879/39.1%/0.11/1.3%
       64   True    32   1024    0.36/384/17.0%/1.50/18.7%
       64   True    16   2048    0.49/561/25.0%/1.10/13.7%
       64   True     8   4096     0.78/708/31.5%/0.69/8.6%
       64   True     4   8192     1.38/797/35.4%/0.39/4.9%
       64   True     2  16384     2.58/852/37.9%/0.21/2.6%
      128  False    32   1024    0.28/975/43.3%/1.90/23.8%
      128  False    16   2048   0.49/1131/50.3%/1.10/13.8%
      128  False     8   4096    0.88/1256/55.8%/0.61/7.7%
      128  False     4   8192    1.71/1289/57.3%/0.31/3.9%
      128  False     2  16384    3.31/1330/59.1%/0.16/2.0%
      128   True    32   1024    0.21/660/29.3%/2.58/32.2%
      128   True    16   2048    0.31/899/40.0%/1.76/22.0%
      128   True     8   4096   0.52/1062/47.2%/1.04/13.0%
      128   True     4   8192    0.92/1191/52.9%/0.58/7.3%
      128   True     2  16384    1.76/1252/55.6%/0.31/3.8%
      256  False    32   1024      1.87/147/6.5%/0.29/3.6%
      256  False    16   2048     1.45/379/16.8%/0.37/4.6%
      256  False     8   4096     1.44/763/33.9%/0.37/4.7%
      256  False     4   8192    2.08/1058/47.0%/0.26/3.2%
      256  False     2  16384    3.54/1243/55.3%/0.15/1.9%
      256   True    32   1024       1.83/75/3.3%/0.29/3.7%
      256   True    16   2048      1.34/205/9.1%/0.40/5.0%
      256   True     8   4096     1.21/456/20.3%/0.45/5.6%
      256   True     4   8192     1.35/816/36.3%/0.40/5.0%
      256   True     2  16384     2.09/1053/46.8%/0.26/3.2%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024     1.36/505/22.5%/0.69/8.6%
       64  False    16   2048     2.22/618/27.5%/0.42/5.3%
       64  False     8   4096     3.96/695/30.9%/0.24/3.0%
       64  False     4   8192     7.56/728/32.3%/0.12/1.6%
       64  False     2  16384    14.74/746/33.2%/0.06/0.8%
       64   True    32   1024    1.01/341/15.1%/0.93/11.6%
       64   True    16   2048     1.46/469/20.9%/0.64/8.0%
       64   True     8   4096     2.38/576/25.6%/0.39/4.9%
       64   True     4   8192     4.21/652/29.0%/0.22/2.8%
       64   True     2  16384     7.87/698/31.0%/0.12/1.5%
      128  False    32   1024    0.94/731/32.5%/1.00/12.5%
      128  False    16   2048     1.47/936/41.6%/0.64/8.0%
      128  False     8   4096    2.56/1073/47.7%/0.37/4.6%
      128  False     4   8192    4.74/1160/51.5%/0.20/2.5%
      128  False     2  16384    9.14/1203/53.5%/0.10/1.3%
      128   True    32   1024    0.75/456/20.3%/1.25/15.6%
      128   True    16   2048    1.01/678/30.2%/0.93/11.6%
      128   True     8   4096     1.58/868/38.6%/0.59/7.4%
      128   True     4   8192     2.77/993/44.1%/0.34/4.2%
      128   True     2  16384    5.06/1086/48.3%/0.19/2.3%
      256  False    32   1024      17.95/38/1.7%/0.05/0.7%
      256  False    16   2048      16.18/85/3.8%/0.06/0.7%
      256  False     8   4096     15.94/172/7.7%/0.06/0.7%
      256  False     4   8192    16.53/333/14.8%/0.06/0.7%
      256  False     2  16384    18.89/582/25.9%/0.05/0.6%
      256   True    32   1024       4.73/73/3.2%/0.20/2.5%
      256   True    16   2048      3.09/222/9.9%/0.30/3.8%
      256   True     8   4096     2.78/494/21.9%/0.34/4.2%
      256   True     4   8192     3.99/688/30.6%/0.24/2.9%
      256   True     2  16384     6.96/790/35.1%/0.13/1.7%
b200_after — FA4 vs cuDNN dense, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.32/846/37.6%/1.65/20.7%    0.39/703/31.2%/1.37/17.2%
       64  False    16   2048    0.64/856/38.0%/0.84/10.4%     0.75/734/32.6%/0.72/9.0%
       64  False     8   4096     1.24/889/39.5%/0.43/5.4%     1.47/750/33.3%/0.37/4.6%
       64  False     4   8192     2.44/901/40.0%/0.22/2.7%     2.90/758/33.7%/0.19/2.3%
       64  False     2  16384     4.80/916/40.7%/0.11/1.4%     5.77/763/33.9%/0.09/1.2%
       64   True    32   1024    0.27/501/22.3%/1.96/24.5%    0.32/427/19.0%/1.67/20.9%
       64   True    16   2048    0.42/658/29.2%/1.29/16.1%    0.47/585/26.0%/1.14/14.3%
       64   True     8   4096     0.71/769/34.2%/0.75/9.4%     0.77/716/31.8%/0.70/8.7%
       64   True     4   8192     1.32/833/37.0%/0.41/5.1%     1.38/798/35.5%/0.39/4.9%
       64   True     2  16384     2.53/870/38.7%/0.21/2.7%     2.61/842/37.4%/0.21/2.6%
      128  False    32   1024   0.24/1159/51.5%/2.26/28.3%   0.23/1205/53.5%/2.35/29.4%
      128  False    16   2048   0.45/1228/54.6%/1.20/15.0%   0.43/1283/57.0%/1.25/15.7%
      128  False     8   4096    0.84/1310/58.2%/0.64/8.0%    0.82/1342/59.6%/0.66/8.2%
      128  False     4   8192    1.64/1344/59.7%/0.33/4.1%    1.63/1349/60.0%/0.33/4.1%
      128  False     2  16384    3.25/1354/60.2%/0.17/2.1%    3.15/1396/62.0%/0.17/2.1%
      128   True    32   1024    0.18/753/33.5%/2.94/36.8%    0.20/705/31.3%/2.75/34.4%
      128   True    16   2048    0.29/959/42.6%/1.87/23.4%    0.29/947/42.1%/1.85/23.1%
      128   True     8   4096   0.48/1135/50.5%/1.11/13.9%   0.50/1100/48.9%/1.07/13.4%
      128   True     4   8192    0.91/1211/53.8%/0.59/7.4%    0.87/1258/55.9%/0.61/7.7%
      128   True     2  16384    1.71/1284/57.1%/0.31/3.9%    1.70/1291/57.4%/0.32/3.9%
      256  False    32   1024    0.37/740/32.9%/1.45/18.1%    0.30/923/41.0%/1.80/22.5%
      256  False    16   2048     0.71/780/34.6%/0.76/9.5%   0.48/1142/50.8%/1.12/13.9%
      256  False     8   4096     1.40/786/34.9%/0.38/4.8%    0.86/1281/56.9%/0.63/7.8%
      256  False     4   8192     2.78/792/35.2%/0.19/2.4%    1.62/1356/60.3%/0.33/4.1%
      256  False     2  16384     5.48/803/35.7%/0.10/1.2%    3.15/1395/62.0%/0.17/2.1%
      256   True    32   1024    0.27/517/23.0%/2.02/25.3%    0.24/583/25.9%/2.28/28.5%
      256   True    16   2048    0.43/632/28.1%/1.24/15.4%    0.33/837/37.2%/1.63/20.4%
      256   True     8   4096     0.78/706/31.4%/0.69/8.6%   0.52/1055/46.9%/1.03/12.9%
      256   True     4   8192     1.46/753/33.4%/0.37/4.6%    0.89/1238/55.0%/0.60/7.6%
      256   True     2  16384     2.86/769/34.2%/0.19/2.3%    1.66/1326/58.9%/0.32/4.0%
b200_after — FA4 vs cuDNN varlen, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.33/846/37.6%/1.65/20.6%    0.44/623/27.7%/1.22/15.2%
       64  False    16   2048    0.64/861/38.3%/0.84/10.5%     0.73/758/33.7%/0.74/9.3%
       64  False     8   4096     1.23/890/39.6%/0.43/5.4%     1.33/827/36.7%/0.40/5.0%
       64  False     4   8192     2.44/901/40.0%/0.22/2.7%     2.53/868/38.6%/0.21/2.6%
       64  False     2  16384     4.85/906/40.3%/0.11/1.4%     5.01/878/39.0%/0.11/1.3%
       64   True    32   1024    0.27/501/22.3%/1.96/24.5%    0.35/388/17.3%/1.52/19.0%
       64   True    16   2048    0.42/657/29.2%/1.28/16.0%    0.49/565/25.1%/1.10/13.8%
       64   True     8   4096     0.71/770/34.2%/0.75/9.4%     0.78/701/31.2%/0.68/8.6%
       64   True     4   8192     1.31/838/37.2%/0.41/5.1%     1.37/802/35.6%/0.39/4.9%
       64   True     2  16384     2.52/872/38.8%/0.21/2.7%     2.58/853/37.9%/0.21/2.6%
      128  False    32   1024   0.24/1156/51.4%/2.26/28.2%    0.28/972/43.2%/1.90/23.7%
      128  False    16   2048   0.44/1261/56.0%/1.23/15.4%   0.48/1137/50.5%/1.11/13.9%
      128  False     8   4096    0.84/1302/57.9%/0.64/7.9%    0.87/1259/56.0%/0.61/7.7%
      128  False     4   8192    1.65/1331/59.1%/0.32/4.1%    1.67/1316/58.5%/0.32/4.0%
      128  False     2  16384    3.27/1345/59.8%/0.16/2.1%    3.32/1326/58.9%/0.16/2.0%
      128   True    32   1024    0.18/756/33.6%/2.95/36.9%    0.21/665/29.6%/2.60/32.5%
      128   True    16   2048    0.28/965/42.9%/1.88/23.6%    0.30/902/40.1%/1.76/22.0%
      128   True     8   4096   0.49/1130/50.2%/1.10/13.8%   0.51/1069/47.5%/1.04/13.1%
      128   True     4   8192    0.90/1224/54.4%/0.60/7.5%    0.91/1214/53.9%/0.59/7.4%
      128   True     2  16384    1.69/1299/57.7%/0.32/4.0%    1.71/1283/57.0%/0.31/3.9%
      256  False    32   1024    0.38/727/32.3%/1.42/17.7%      1.88/146/6.5%/0.29/3.6%
      256  False    16   2048     0.71/774/34.4%/0.76/9.5%     1.45/379/16.8%/0.37/4.6%
      256  False     8   4096     1.39/792/35.2%/0.39/4.8%     1.44/762/33.9%/0.37/4.6%
      256  False     4   8192     2.79/787/35.0%/0.19/2.4%    2.06/1067/47.4%/0.26/3.3%
      256  False     2  16384     5.36/820/36.5%/0.10/1.3%    3.57/1233/54.8%/0.15/1.9%
      256   True    32   1024    0.26/522/23.2%/2.04/25.5%       1.83/75/3.3%/0.29/3.7%
      256   True    16   2048    0.44/624/27.8%/1.22/15.2%      1.34/204/9.1%/0.40/5.0%
      256   True     8   4096     0.77/716/31.8%/0.70/8.7%     1.21/456/20.3%/0.45/5.6%
      256   True     4   8192     1.46/755/33.6%/0.37/4.6%     1.35/814/36.2%/0.40/5.0%
      256   True     2  16384     2.86/769/34.2%/0.19/2.3%    2.09/1054/46.9%/0.26/3.2%

Before — cutlass 4.4.2

B200 numbers on cutlass 4.4.2 (commit c200609, immediately before the deprecated-API replacement in eaf806d), with quack 0.4.0. Same clock lock (1867 MHz), same --warmup 25 --rep 100, same shape grid. Tables sit alongside the after numbers above for direct side-by-side reading.

FA4 peaks: Dense FWD 1401 TFLOPS, Dense BWD 1215, Varlen FWD 1321, Varlen BWD 1205 (peak across configs).

Before → after delta (B200): no regression — flat within run-to-run noise: Dense FWD 1401→1410, Dense BWD 1215→1221, Varlen FWD 1321→1330, Varlen BWD 1205→1203 TFLOPS.

b200_before — FA4 dense, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    0.42/657/29.2%/1.28/16.0%
       64  False    16   2048     0.81/682/30.3%/0.67/8.3%
       64  False     8   4096     1.58/695/30.9%/0.34/4.2%
       64  False     4   8192     3.14/701/31.2%/0.17/2.1%
       64  False     2  16384     6.23/706/31.4%/0.09/1.1%
       64   True    32   1024    0.32/425/18.9%/1.66/20.8%
       64   True    16   2048    0.47/580/25.8%/1.13/14.2%
       64   True     8   4096     0.77/713/31.7%/0.70/8.7%
       64   True     4   8192     1.38/794/35.3%/0.39/4.8%
       64   True     2  16384     2.63/836/37.1%/0.20/2.6%
      128  False    32   1024   0.23/1196/53.2%/2.34/29.2%
      128  False    16   2048   0.43/1266/56.3%/1.24/15.5%
      128  False     8   4096    0.84/1309/58.2%/0.64/8.0%
      128  False     4   8192    1.62/1357/60.3%/0.33/4.1%
      128  False     2  16384    3.14/1401/62.3%/0.17/2.1%
      128   True    32   1024    0.20/704/31.3%/2.75/34.4%
      128   True    16   2048    0.29/944/42.0%/1.84/23.0%
      128   True     8   4096   0.49/1114/49.5%/1.09/13.6%
      128   True     4   8192    0.89/1231/54.7%/0.60/7.5%
      128   True     2  16384    1.70/1292/57.4%/0.32/3.9%
      256  False    32   1024    0.30/920/40.9%/1.80/22.5%
      256  False    16   2048   0.49/1123/49.9%/1.10/13.7%
      256  False     8   4096    0.87/1259/55.9%/0.61/7.7%
      256  False     4   8192    1.65/1336/59.4%/0.33/4.1%
      256  False     2  16384    3.19/1379/61.3%/0.17/2.1%
      256   True    32   1024    0.24/585/26.0%/2.28/28.6%
      256   True    16   2048    0.33/842/37.4%/1.64/20.6%
      256   True     8   4096   0.52/1060/47.1%/1.04/12.9%
      256   True     4   8192    0.90/1217/54.1%/0.59/7.4%
      256   True     2  16384    1.71/1285/57.1%/0.31/3.9%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024     1.26/545/24.2%/0.74/9.3%
       64  False    16   2048     2.14/642/28.5%/0.44/5.5%
       64  False     8   4096     3.91/703/31.2%/0.24/3.0%
       64  False     4   8192     7.52/731/32.5%/0.12/1.6%
       64  False     2  16384    14.73/746/33.2%/0.06/0.8%
       64   True    32   1024    0.89/384/17.1%/1.05/13.1%
       64   True    16   2048     1.35/509/22.6%/0.70/8.7%
       64   True     8   4096     2.27/606/26.9%/0.41/5.2%
       64   True     4   8192     4.10/670/29.8%/0.23/2.9%
       64   True     2  16384     7.76/708/31.5%/0.12/1.5%
      128  False    32   1024    0.85/809/36.0%/1.11/13.8%
      128  False    16   2048     1.38/994/44.2%/0.68/8.5%
      128  False     8   4096    2.47/1113/49.5%/0.38/4.8%
      128  False     4   8192    4.61/1191/52.9%/0.20/2.5%
      128  False     2  16384    9.05/1215/54.0%/0.10/1.3%
      128   True    32   1024    0.63/547/24.3%/1.50/18.7%
      128   True    16   2048    0.90/761/33.8%/1.04/13.0%
      128   True     8   4096     1.44/954/42.4%/0.65/8.2%
      128   True     4   8192    2.56/1075/47.8%/0.37/4.6%
      128   True     2  16384    4.66/1181/52.5%/0.20/2.5%
      256  False    32   1024    1.13/606/26.9%/0.83/10.4%
      256  False    16   2048     1.86/741/32.9%/0.51/6.3%
      256  False     8   4096     3.31/830/36.9%/0.28/3.5%
      256  False     4   8192     6.01/914/40.6%/0.16/2.0%
      256  False     2  16384    11.75/936/41.6%/0.08/1.0%
      256   True    32   1024    0.93/371/16.5%/1.01/12.7%
      256   True    16   2048     1.29/531/23.6%/0.73/9.1%
      256   True     8   4096     2.07/664/29.5%/0.45/5.7%
      256   True     4   8192     3.66/750/33.3%/0.26/3.2%
      256   True     2  16384     6.80/809/35.9%/0.14/1.7%
b200_before — FA4 varlen, forward + backward 
==========================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024    0.45/617/27.4%/1.21/15.1%
       64  False    16   2048     0.73/756/33.6%/0.74/9.2%
       64  False     8   4096     1.33/827/36.8%/0.40/5.0%
       64  False     4   8192     2.57/857/38.1%/0.21/2.6%
       64  False     2  16384     5.02/877/39.0%/0.11/1.3%
       64   True    32   1024    0.36/382/17.0%/1.49/18.7%
       64   True    16   2048    0.49/560/24.9%/1.09/13.7%
       64   True     8   4096     0.78/707/31.4%/0.69/8.6%
       64   True     4   8192     1.38/797/35.4%/0.39/4.9%
       64   True     2  16384     2.55/862/38.3%/0.21/2.6%
      128  False    32   1024    0.28/969/43.1%/1.89/23.7%
      128  False    16   2048   0.49/1132/50.3%/1.11/13.8%
      128  False     8   4096    0.90/1228/54.6%/0.60/7.5%
      128  False     4   8192    1.69/1302/57.9%/0.32/4.0%
      128  False     2  16384    3.33/1321/58.7%/0.16/2.0%
      128   True    32   1024    0.21/659/29.3%/2.58/32.2%
      128   True    16   2048    0.31/886/39.4%/1.73/21.6%
      128   True     8   4096   0.52/1061/47.2%/1.04/13.0%
      128   True     4   8192    0.92/1198/53.2%/0.58/7.3%
      128   True     2  16384    1.75/1260/56.0%/0.31/3.8%
      256  False    32   1024      1.87/147/6.5%/0.29/3.6%
      256  False    16   2048     1.45/379/16.9%/0.37/4.6%
      256  False     8   4096     1.44/763/33.9%/0.37/4.7%
      256  False     4   8192    2.07/1062/47.2%/0.26/3.2%
      256  False     2  16384    3.54/1241/55.2%/0.15/1.9%
      256   True    32   1024       1.83/75/3.3%/0.29/3.7%
      256   True    16   2048      1.34/205/9.1%/0.40/5.0%
      256   True     8   4096     1.20/457/20.3%/0.45/5.6%
      256   True     4   8192     1.35/814/36.2%/0.40/5.0%
      256   True     2  16384    2.08/1056/46.9%/0.26/3.2%
==========================================================
  BWD (ms / TFLOPS / MFU% / TBs / BW%)
==========================================================
     hdim causal batch seqlen                          FA4
----------------------------------------------------------
       64  False    32   1024     1.35/507/22.5%/0.69/8.7%
       64  False    16   2048     2.21/621/27.6%/0.42/5.3%
       64  False     8   4096     3.94/698/31.0%/0.24/3.0%
       64  False     4   8192     7.52/731/32.5%/0.12/1.6%
       64  False     2  16384    14.60/753/33.5%/0.06/0.8%
       64   True    32   1024    1.01/341/15.2%/0.93/11.7%
       64   True    16   2048     1.46/469/20.9%/0.64/8.0%
       64   True     8   4096     2.38/577/25.6%/0.39/4.9%
       64   True     4   8192     4.21/652/29.0%/0.22/2.8%
       64   True     2  16384     7.86/699/31.1%/0.12/1.5%
      128  False    32   1024    0.94/728/32.3%/0.99/12.4%
      128  False    16   2048     1.48/929/41.3%/0.64/7.9%
      128  False     8   4096    2.54/1081/48.0%/0.37/4.6%
      128  False     4   8192    4.83/1138/50.6%/0.19/2.4%
      128  False     2  16384    9.13/1205/53.5%/0.10/1.3%
      128   True    32   1024    0.75/457/20.3%/1.25/15.6%
      128   True    16   2048    1.02/676/30.0%/0.92/11.6%
      128   True     8   4096     1.60/861/38.3%/0.59/7.4%
      128   True     4   8192    2.69/1021/45.4%/0.35/4.4%
      128   True     2  16384    4.90/1123/49.9%/0.19/2.4%
      256  False    32   1024      17.98/38/1.7%/0.05/0.7%
      256  False    16   2048      16.25/85/3.8%/0.06/0.7%
      256  False     8   4096     15.87/173/7.7%/0.06/0.7%
      256  False     4   8192    16.53/333/14.8%/0.06/0.7%
      256  False     2  16384    18.79/585/26.0%/0.05/0.6%
      256   True    32   1024       4.80/72/3.2%/0.20/2.4%
      256   True    16   2048      3.16/218/9.7%/0.30/3.7%
      256   True     8   4096     2.85/482/21.4%/0.33/4.1%
      256   True     4   8192     4.01/685/30.4%/0.23/2.9%
      256   True     2  16384     6.98/788/35.0%/0.13/1.7%
b200_before — FA4 vs cuDNN dense, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.33/846/37.6%/1.65/20.6%    0.39/703/31.2%/1.37/17.2%
       64  False    16   2048    0.64/856/38.1%/0.84/10.5%     0.75/734/32.6%/0.72/9.0%
       64  False     8   4096     1.24/885/39.3%/0.43/5.4%     1.47/749/33.3%/0.37/4.6%
       64  False     4   8192     2.45/897/39.9%/0.22/2.7%     2.90/757/33.7%/0.18/2.3%
       64  False     2  16384     4.86/904/40.2%/0.11/1.4%     5.77/762/33.9%/0.09/1.2%
       64   True    32   1024    0.27/501/22.3%/1.96/24.5%    0.32/428/19.0%/1.67/20.9%
       64   True    16   2048    0.42/659/29.3%/1.29/16.1%    0.47/586/26.0%/1.14/14.3%
       64   True     8   4096     0.72/765/34.0%/0.75/9.3%     0.77/715/31.8%/0.70/8.7%
       64   True     4   8192     1.31/838/37.2%/0.41/5.1%     1.37/801/35.6%/0.39/4.9%
       64   True     2  16384     2.54/865/38.4%/0.21/2.6%     2.61/842/37.4%/0.21/2.6%
      128  False    32   1024   0.24/1156/51.4%/2.26/28.2%   0.23/1203/53.5%/2.35/29.4%
      128  False    16   2048   0.44/1257/55.9%/1.23/15.3%   0.43/1279/56.8%/1.25/15.6%
      128  False     8   4096    0.84/1312/58.3%/0.64/8.0%    0.84/1314/58.4%/0.64/8.0%
      128  False     4   8192    1.66/1323/58.8%/0.32/4.0%    1.62/1357/60.3%/0.33/4.1%
      128  False     2  16384    3.27/1344/59.7%/0.16/2.1%    3.16/1390/61.8%/0.17/2.1%
      128   True    32   1024    0.18/752/33.4%/2.94/36.7%    0.20/705/31.3%/2.75/34.4%
      128   True    16   2048    0.28/973/43.3%/1.90/23.8%    0.29/932/41.4%/1.82/22.8%
      128   True     8   4096   0.49/1129/50.2%/1.10/13.8%   0.50/1104/49.1%/1.08/13.5%
      128   True     4   8192    0.89/1234/54.8%/0.60/7.5%    0.90/1215/54.0%/0.59/7.4%
      128   True     2  16384    1.72/1281/56.9%/0.31/3.9%    1.70/1293/57.4%/0.32/3.9%
      256  False    32   1024    0.38/724/32.2%/1.41/17.7%    0.30/919/40.9%/1.80/22.4%
      256  False    16   2048     0.71/769/34.2%/0.75/9.4%   0.48/1136/50.5%/1.11/13.9%
      256  False     8   4096     1.39/788/35.0%/0.38/4.8%    0.88/1255/55.8%/0.61/7.7%
      256  False     4   8192     2.71/811/36.0%/0.20/2.5%    1.63/1352/60.1%/0.33/4.1%
      256  False     2  16384     5.48/803/35.7%/0.10/1.2%    3.17/1388/61.7%/0.17/2.1%
      256   True    32   1024    0.27/518/23.0%/2.02/25.3%    0.23/586/26.0%/2.29/28.6%
      256   True    16   2048    0.44/628/27.9%/1.23/15.3%    0.33/837/37.2%/1.63/20.4%
      256   True     8   4096     0.74/746/33.2%/0.73/9.1%   0.52/1055/46.9%/1.03/12.9%
      256   True     4   8192     1.46/755/33.6%/0.37/4.6%    0.91/1214/53.9%/0.59/7.4%
      256   True     2  16384     2.79/789/35.1%/0.19/2.4%    1.70/1291/57.4%/0.32/3.9%
b200_before — FA4 vs cuDNN varlen, forward (reference) 
=======================================================================================
  FWD (ms / TFLOPS / MFU% / TBs / BW%)
=======================================================================================
     hdim causal batch seqlen                        cuDNN                          FA4
---------------------------------------------------------------------------------------
       64  False    32   1024    0.33/844/37.5%/1.65/20.6%    0.44/623/27.7%/1.22/15.2%
       64  False    16   2048    0.64/862/38.3%/0.84/10.5%     0.72/760/33.8%/0.74/9.3%
       64  False     8   4096     1.23/891/39.6%/0.43/5.4%     1.33/828/36.8%/0.40/5.1%
       64  False     4   8192     2.45/898/39.9%/0.22/2.7%     2.53/869/38.6%/0.21/2.7%
       64  False     2  16384     4.86/905/40.2%/0.11/1.4%     5.01/878/39.0%/0.11/1.3%
       64   True    32   1024    0.27/501/22.3%/1.96/24.5%    0.35/387/17.2%/1.51/18.9%
       64   True    16   2048    0.42/658/29.2%/1.29/16.1%    0.49/565/25.1%/1.10/13.8%
       64   True     8   4096     0.72/769/34.2%/0.75/9.4%     0.78/705/31.4%/0.69/8.6%
       64   True     4   8192     1.31/837/37.2%/0.41/5.1%     1.37/800/35.5%/0.39/4.9%
       64   True     2  16384     2.53/869/38.6%/0.21/2.7%     2.57/856/38.0%/0.21/2.6%
      128  False    32   1024   0.24/1153/51.2%/2.25/28.1%    0.28/976/43.4%/1.91/23.8%
      128  False    16   2048   0.43/1278/56.8%/1.25/15.6%   0.48/1137/50.5%/1.11/13.9%
      128  False     8   4096    0.84/1305/58.0%/0.64/8.0%    0.89/1235/54.9%/0.60/7.5%
      128  False     4   8192    1.65/1332/59.2%/0.33/4.1%    1.68/1305/58.0%/0.32/4.0%
      128  False     2  16384    3.28/1339/59.5%/0.16/2.0%    3.34/1317/58.5%/0.16/2.0%
      128   True    32   1024    0.18/756/33.6%/2.95/36.9%    0.21/667/29.7%/2.61/32.6%
      128   True    16   2048    0.29/963/42.8%/1.88/23.5%    0.31/901/40.0%/1.76/22.0%
      128   True     8   4096   0.49/1113/49.5%/1.09/13.6%   0.51/1081/48.0%/1.06/13.2%
      128   True     4   8192    0.88/1246/55.4%/0.61/7.6%    0.91/1206/53.6%/0.59/7.4%
      128   True     2  16384    1.72/1280/56.9%/0.31/3.9%    1.71/1284/57.0%/0.31/3.9%
      256  False    32   1024    0.39/713/31.7%/1.39/17.4%      1.87/147/6.5%/0.29/3.6%
      256  False    16   2048     0.72/767/34.1%/0.75/9.4%     1.45/380/16.9%/0.37/4.6%
      256  False     8   4096     1.39/791/35.2%/0.39/4.8%     1.44/764/34.0%/0.37/4.7%
      256  False     4   8192     2.72/807/35.9%/0.20/2.5%    2.06/1070/47.5%/0.26/3.3%
      256  False     2  16384     5.48/803/35.7%/0.10/1.2%    3.56/1235/54.9%/0.15/1.9%
      256   True    32   1024    0.26/521/23.2%/2.04/25.5%       1.83/75/3.3%/0.29/3.7%
      256   True    16   2048    0.43/633/28.1%/1.24/15.5%      1.34/204/9.1%/0.40/5.0%
      256   True     8   4096     0.78/700/31.1%/0.68/8.5%     1.20/456/20.3%/0.45/5.6%
      256   True     4   8192     1.46/753/33.5%/0.37/4.6%     1.35/815/36.2%/0.40/5.0%
      256   True     2  16384     2.83/778/34.6%/0.19/2.4%    2.06/1066/47.4%/0.26/3.3%

Reproduce:

nvidia-smi -i 0 -lgc 1965,1965    # lock; reset afterward with: nvidia-smi -i 0 -rgc
CUDA_VISIBLE_DEVICES=0 FLASH_ATTENTION_CUTE_DSL_CACHE_ENABLED=1 \
python benchmarks/benchmark_attn.py --backend fa4 --fwd --bwd --causal both \
  --headdim 64,128,256 --seqlen 1024,2048,4096,8192,16384 --warmup 25 --rep 100
# add --varlen for the varlen baseline; use --backend fa4,cudnn --fwd for the reference

Conclusion

The nvidia-cutlass-dsl 4.4.2 → 4.5.2 / quack-kernels 0.4.0 → 0.5.0 bump is a clean upgrade — no correctness regression and no performance regression on either Blackwell arch (GB300, B200):

  • Correctness: identical pass/skip/fail (250 / 212 / 5) before and after, on both archs. All 5 failures are pre-existing test_score_mod.py tolerance cases unrelated to the bump.
  • Performance: peak FA4 TFLOPS move within ±1.3% (run-to-run noise) across dense/varlen, fwd/bwd — GB300 Dense FWD 1886→1882, BWD 1520→1530, Varlen FWD 1749→1745, BWD 1488→1508; B200 Dense FWD 1401→1410, BWD 1215→1221, Varlen FWD 1321→1330, BWD 1205→1203. The short-seqlen varlen-hd256 weak spot is present identically before and after, so it is a baseline characteristic, not introduced by this bump.

Safe to merge; the only deferred item is the unrelated OperandMajorMode deprecation below.

Notes / follow-up

  • New deprecation surfaced by 4.5.2 (not addressed here): tcgen05.OperandMajorMode → cute.nvgpu.OperandMajorMode. ~51 call sites across 7 SM100 kernel files (flash_bwd_sm100.py, flash_fwd_sm100.py, flash_fwd_mla_sm100.py, flash_bwd_postprocess.py, sm100_hd256_2cta_fmha_{forward,backward_dqkernel,backward_dkdvkernel}.py). Emits DeprecationWarning only — kernels compile and run correctly. Deferred to a separate PR because the new and deprecated enums are distinct objects (OperandMajorMode.K is not identical across the two), so equality comparisons against MMA-op major modes need careful review rather than a mechanical find-replace.

@Johnsonms
Bump nvidia-cutlass-dsl to >=4.5.2 and quack-kernels to >=0.5.0
fb712d2 
cutlass 4.5.2 is safe to update, and quack 0.5.0 has been published, so
bump the FA4 (flash_attn/cute) requirement floors to match. Updates the
dependencies and the cu13 extra in pyproject.toml, and the documented
versions in CLAUDE.md.

Verified on NVIDIA GB300 (SM100, CUDA 13.2): deps resolve cleanly
(nvidia-cutlass-dsl 4.5.2 base+cu13, quack-kernels 0.5.0), imports OK,
and a representative GPU sample of tests/cute/test_flash_attn.py passes
(6 passed / 6 skipped / 0 failed across hd 64/96/128/192, causal,
mha/gqa/mqa, fwd+bwd).
@Johnsonms
Johnsonms marked this pull request as ready for review May 30, 2026 04:28
@Johnsonms
Johnsonms requested a review from Copilot May 30, 2026 14:33
@Johnsonms
Johnsonms requested review from drisspg, jayhshah and tridao May 30, 2026 14:34
Copilot AI reviewed May 30, 2026
View reviewed changes

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Bumps FA4 Python runtime dependency floors: nvidia-cutlass-dsl to >=4.5.2 (base + cu13 extra) and quack-kernels to >=0.5.0, with matching documentation in CLAUDE.md.

Changes:

  • Raise nvidia-cutlass-dsl floor from 4.4.2 to 4.5.2 in both dependencies and the cu13 extra.
  • Raise quack-kernels floor from 0.4.0 to 0.5.0.
  • Sync CLAUDE.md dependency list to match.

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated no comments.

File Description
flash_attn/cute/pyproject.toml Bumps cutlass-dsl and quack-kernels version floors.
CLAUDE.md Updates documented dependency versions to match the manifest.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

drisspg
drisspg approved these changes May 30, 2026
View reviewed changes

@drisspg drisspg left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good, testing looks thorough

@Johnsonms

Johnsonms commented May 30, 2026

Copy link
Copy Markdown
Collaborator Author

Thanks @drisspg !

@Johnsonms
Johnsonms merged commit f82d0dc into Dao-AILab:main May 30, 2026
@Johnsonms
Johnsonms deleted the bump-cutlass-4.5.2-quack-0.5.0 branch May 30, 2026 15:32
@Johnsonms Johnsonms mentioned this pull request Jun 1, 2026
Closed
reubenconducts pushed a commit to reubenconducts/flash-attention that referenced this pull request Jun 2, 2026
@Johnsonms
Bump nvidia-cutlass-dsl to >=4.5.2 and quack-kernels to >=0.5.0 (Dao-…
6b33bbc 
…AILab#2605)

cutlass 4.5.2 is safe to update, and quack 0.5.0 has been published, so
bump the FA4 (flash_attn/cute) requirement floors to match. Updates the
dependencies and the cu13 extra in pyproject.toml, and the documented
versions in CLAUDE.md.

Verified on NVIDIA GB300 (SM100, CUDA 13.2): deps resolve cleanly
(nvidia-cutlass-dsl 4.5.2 base+cu13, quack-kernels 0.5.0), imports OK,
and a representative GPU sample of tests/cute/test_flash_attn.py passes
(6 passed / 6 skipped / 0 failed across hd 64/96/128/192, causal,
mha/gqa/mqa, fwd+bwd).
@MatthewBonanni MatthewBonanni mentioned this pull request Jun 16, 2026
Merged
@arpera arpera mentioned this pull request Jul 7, 2026
Merged
MatthewBonanni added a commit to vllm-project/flash-attention that referenced this pull request Jul 13, 2026
@MatthewBonanni @Luosuu
@brandon-yujie-sun @Johnsonms @jayhshah @ko3n1g @micmelesse @reubenconducts @drisspg @aryanputta @claude @sryap @yunweili3 @hhy3 @anakinxc @pashu-cohere @janeyx99 @oattia @bbuschkaemper @rocking5566 @DDEle @hyoon1 @yinli-systems Kevin-Li-2025 @johnnynunez
Sync upstream (#149)
3286005 
* [Fwd,Sm100] fix: decode↔prefill exp2 emulation consistency (Dao-AILab#2595)

apply_exp2_convert selected the exp2 implementation based on mask_fn
presence: hardware ex2.approx.ftz for causal-masked tiles, polynomial
emulation for unmasked tiles. Different q_stage values (1 for decode,
2 for prefill) compute different m_block for the same logical Q row,
shifting which tiles are processed with vs without mask_fn. The same
K tile could receive different exp2 methods across variants.

Fix: always pass self.ex2_emu_freq regardless of mask_fn presence.
Add regression test for decode↔prefill bitwise consistency on MLA
(192,128) shapes.

* replace deprecated apis (Dao-AILab#2602)

* Bump nvidia-cutlass-dsl to >=4.5.2 and quack-kernels to >=0.5.0 (Dao-AILab#2605)

cutlass 4.5.2 is safe to update, and quack 0.5.0 has been published, so
bump the FA4 (flash_attn/cute) requirement floors to match. Updates the
dependencies and the cu13 extra in pyproject.toml, and the documented
versions in CLAUDE.md.

Verified on NVIDIA GB300 (SM100, CUDA 13.2): deps resolve cleanly
(nvidia-cutlass-dsl 4.5.2 base+cu13, quack-kernels 0.5.0), imports OK,
and a representative GPU sample of tests/cute/test_flash_attn.py passes
(6 passed / 6 skipped / 0 failed across hd 64/96/128/192, causal,
mha/gqa/mqa, fwd+bwd).

* [CuTe,Fwd,Sm100] refactor mla sm100 forward and add page table (Dao-AILab#2558)

* refactor mla sm100 forward

* add benchmark; address deprecation warnings; tweak ptx gemm dispatch

* update interface and tests

* ci: bump Jimver/cuda-toolkit to v0.2.35 for CUDA 13.2 support (Dao-AILab#2617)

v0.2.30 only ships URLs up to CUDA 13.1.0; bumping to v0.2.35 adds
13.1.1, 13.2.0, and the matching aarch64 SBSA installers.

Signed-off-by: oliver könig <okoenig@nvidia.com>

* [ROCm] Bump Triton to >=3.6.0 and aiter submodule (Dao-AILab#2614)

* [Triton] Fix graph capture issues and env var (Dao-AILab#2620)

* graph capture fix

* rm env flag

* [CuTe,Bwd,Sm100] allow 2cta with score mod and mask mod in bwd (Dao-AILab#2557)

* [CuTe] Fix lint failures (Dao-AILab#2625)

stack-info: PR: Dao-AILab#2625, branch: drisspg/stack/42

* [CuTe] Fix lint failure in flash_bwd_sm100.py (Dao-AILab#2627)

ruff format flagged flash_attn/cute/flash_bwd_sm100.py (trailing
whitespace in a comment and an over-split call). It was missed by
the lint sweep in Dao-AILab#2625.

* fix: add weights_only=True to all torch.load call sites (Dao-AILab#2622)

Passing weights_only=False (the pre-2.4 default) to torch.load allows
arbitrary Python object deserialization from the checkpoint file.
A malicious .pt/.pth file can execute arbitrary code on the machine
loading it — a well-known PyTorch deserialization vector (CWE-502).

Four call sites updated:
  training/src/utils/checkpoint.py  load_checkpoint()
  training/src/eval.py               eval checkpoint loader
  flash_attn/utils/pretrained.py    partial(torch.load, ...) loader
  flash_attn/models/llama.py        state_dicts_from_checkpoint()

weights_only=True restricts deserialization to tensors, dicts, lists,
tuples, and other primitive types — no arbitrary Python objects.
Requires PyTorch >= 1.13; FA4's CuTeDSL dependency already requires
a modern PyTorch 2.x build, so no compatibility regression.

Fixes Dao-AILab#2583

* use correction warps if not tma store; remove outdated packgqa guard (Dao-AILab#2629)

* Add aux-scalars to interface to enable dynamic ints and floats in expressions (Dao-AILab#2616)

stack-info: PR: Dao-AILab#2616, branch: drisspg/stack/41

* fix: build and select cu13.2 prebuilt wheels (Dao-AILab#2618)

* ci: use 1 ninja job for cu13.2

Signed-off-by: oliver könig <okoenig@nvidia.com>

* fix(setup): request cu13 prebuilt wheels for CUDA 13 torch

get_wheel_url() binned every CUDA >= 12 to major '12', so under a CUDA 13
torch it requested cu12 wheels and never matched the published cu13
artifacts, falling back to a multi-hour source build. Add a CUDA 13
branch so the guessed wheel name uses cu13, matching WHEEL_CUDA_VERSION
in _build.yml.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: oliver könig <okoenig@nvidia.com>

---------

Signed-off-by: oliver könig <okoenig@nvidia.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* ci(fa4): enforce cutlass-dsl/quack dep floors and rebake cu130 image (Dao-AILab#2636)

* ci(fa4): assert cute dep floors in CI; fail loudly on a stale SIF

run_fa4_ci.py installs FA4 with --no-deps (to keep the SIF's baked
torch/cudnn), so the nvidia-cutlass-dsl>=4.5.2 / quack-kernels>=0.5.0 floors
in flash_attn/cute/pyproject.toml are not enforced at install time. A SIF
baked before a floor bump keeps a stale dep — e.g. cutlass-dsl 4.4.2, which
can't convert the AuxData JIT arg and dies with a cryptic DSLRuntimeError
deep in SM100 kernel launch (reproduced on B200).

Upgrading the dep in-place is not viable: the --writable-tmpfs overlay is
RAM-backed and too small for a cutlass-dsl reinstall (ENOSPC, and a partial
removal corrupts the baked torch). So instead of installing, add
assert_dsl_floor.py — it reads the floors from pyproject (no hardcoded
version to drift) and fails with an actionable "rebake the image" message
when the installed cutlass-dsl/quack are below them. Wired into run_step
right after the editable install.

The durable fix is to rebake the image at the current floors and bump the
digest in .github/workflows/ci.yml; this guard makes future drift fail fast
instead of silently.

* ci(fa4): bump cu130 image to 26.06.10 (cutlass-dsl 4.5.2 / quack 0.5.0)

* ci(fa4): fall back to tomli when tomllib is unavailable (Python 3.10)

* Fix SM100 FP8 fwd with cutlass-dsl >=4.5.2 (MmaF8F6F4Op) (Dao-AILab#2640)

cutlass-dsl >=4.5.2 changed make_trivial_tiled_mma to build plain FP8
MMAs as MmaF8F6F4Op (its _F8F6F4_TYPES branch) instead of the now-legacy
MmaFP8Op. The two are siblings under MmaOp, so _tcgen05_mma_kind's
isinstance(op, MmaFP8Op) check missed the new type and raised
"Unsupported tcgen05 MMA op kind: MmaF8F6F4Op", breaking the FP8 forward
path on Blackwell. Worked on 4.4.2.

Accept both ops in the f8f6f4 branch (both map to kind::f8f6f4).
mma_op_to_idesc only reads generic op attrs and is unaffected.

Validated on B200: FP8 fwd passes for all configs in the issue
(incl. hd=64) plus hd=128, causal and non-causal; mean abs err vs bf16
~0.002-0.01.

Fixes Dao-AILab#2639

* [cute] Fix int32 overflow in SM100 LPT tile scheduler for long context (Dao-AILab#2662)

The LPT tile scheduler sizes its L2 swizzle from

    seqlen_k * (headdim + headdim_v) * element_size

in int32. For long context this overflows once it exceeds 2**31
(seqlen_k > ~4M for hdim-128 bf16), making size_one_head negative.
That corrupts the swizzle and the L2 divmods, so get_current_work
decodes an out-of-bounds batch_idx and the kernel performs an illegal
memory access (cudaErrorIllegalAddress) on SM100.

Compute the byte size in int64. swizzle stays small and is cast back
to int32 for the device-side divmods, so there is no behavior or perf
change for non-overflowing shapes.

Fixes both SingleTileLPTScheduler (forward; selected for causal/local)
and SingleTileLPTBwdScheduler (backward; its extra
seqlen_k * headdim * 4 term overflows even sooner).

Repro on SM100 (e.g. GB200), causal forward at seqlen_k = 2**22:

    import torch
    from flash_attn.cute.interface import flash_attn_func

    sq, sk = 2048, 4_194_304  # seqlen_k = 2**22 -> int32 overflow
    q = torch.randn(1, sq, 8, 128, dtype=torch.bfloat16, device="cuda")
    k = torch.randn(1, sk, 1, 128, dtype=torch.bfloat16, device="cuda")
    v = torch.randn(1, sk, 1, 128, dtype=torch.bfloat16, device="cuda")
    out = flash_attn_func(q, k, v, causal=True)
    torch.cuda.synchronize()  # cudaErrorIllegalAddress here before the fix

Crashes before this change, runs clean after; seqlen_k = 2**22 - 128 is
clean both ways (the int32 boundary). Verified clean under
compute-sanitizer memcheck.

* [Fwd,Sm100] Tune FP8 causal hd128 ex2_emu_freq (8 vs inherited 16) (Dao-AILab#2642)

FP8 fwd is MUFU/ex2-bound on Blackwell, so the optimal exp2-emulation
frequency differs from bf16. The causal hd128 key (False,True,128,False)
had no FP8 entry and inherited bf16's freq=16; freq=8 offloads more exp
from the MUFU unit.

Thermally-matched back-to-back A/B on B200 (locked-ish clock, hot GPU,
median of 300 iters, nheads=16 = benchmark default) across the official
benchmark's causal hd128 shapes:

  b   s      f16 TFLOP  f8 TFLOP  delta
  32  512    500.6      516.3     +3.1%
  16  1024   796.5      832.5     +4.5%
  8   2048   1124.6     1175.1    +4.5%
  4   4096   1407.9     1481.3    +5.2%
  2   8192   1604.1     1661.7    +3.6%
  1   16384  1683.7     1726.0    +2.5%

Accuracy-neutral (FP8-vs-bf16 mean-abs-err unchanged; benchmark --check
passes 24/24). Keyed on is_causal=True only: freq=8 would regress
non-causal hd128 (0.94x), which keeps its existing freq=10.

* Make q_subtile_factor default to identity (Dao-AILab#2660)

* fix(hd256/sm100): make q/k/v contiguous before dedicated hd256 kernel (Dao-AILab#2666)

The BlackwellFusedMultiHeadAttentionForward kernel builds tensor layouts
with hardcoded contiguous strides computed from shape dimensions, so
non-contiguous inputs (e.g. from .transpose()) cause wrong memory accesses
and silently corrupt outputs on B200 (SM100) with head_dim=256.

maybe_contiguous() only guarantees stride(-1)==1; add explicit full
contiguity checks in both the forward and backward paths when the hd256
dedicated kernel is selected.

Fixes: Dao-AILab#2665

* [Cute,Bwd,Sm100] add sparse MLA (Deepseek v4) backward kernels (Dao-AILab#2621)

* add backward sparse mla kernels

* add dk gemm

* fix errors

* fix dq errors

* rename bwd kernels

* refactor interface

* fix predicate error in dq kernel

* update tests

* mla fwd fixes

* improve varlen fwd perf

* use cluster idx scheduling in fwd

* use packed scheduler for mqa 128

* fix int32 overflow in swizzle

* simplify bwd preprocess

* refactor bwd

* simplify preprocess

* update benchmark script

* add safety check

* remove test code

* ruff format

* ensure scale is 0 for masked out rows

* fix: sync callers with new _flash_attn_fwd 4-tuple return signature (Dao-AILab#2674)

* Fix compatibility issues with CuTe DSL 4.6.0+ (Dao-AILab#2648)

* Prepare for 4.6 release

* Bump version

* Update pyproject.toml

* Update nvidia-cutlass-dsl version in pyproject.toml

* Pass tmem scalar fields as .ptr to TmemAllocator on SM100 (Dao-AILab#2679)

The DSL now warns when a struct scalar is used directly as a pointer
("Use explicit struct.scalar.ptr for pointer instead"), so these fire
on every tmem_holding_buf / dealloc mbar access. Just pass .ptr like the
other SM100 kernels already do.

* Add FLASHATTENTION_DISABLE_SPLIT_ALIGNMENT (Dao-AILab#2680)

* ci: rebake cu130 image for cutlass-dsl 4.6.0.dev0 floor (Dao-AILab#2684)

PR Dao-AILab#2648 bumped the flash_attn/cute/pyproject.toml floor to
nvidia-cutlass-dsl==4.6.0.dev0, but the CI image (26.06.10) still ships
4.5.2. assert_dsl_floor.py correctly fails every push to main with
"installed 4.5.2 does not satisfy floor ==4.6.0.dev0", so FA4 CI has
been red since Dao-AILab#2648 landed.

- Dockerfile: add --prerelease=allow to the FA4 install. The dev-build
  floor pulls transitive pre-releases (nvidia-cutlass-dsl-libs-base==
  4.6.0.dev0 ...) that uv refuses without it; the old stable 4.5.2 floor
  didn't need it.
- ci.yml: bump fa4_image_cu130 to the rebaked 26.06.27 image
  (cutlass-dsl 4.6.0.dev0, quack-kernels 0.5.3, torch 2.12.1).

E2e verified on B200: assert_dsl_floor passes, compile + run + benchmark
all green (run_fa4_ci.py, exit 0).

* Update FA4 cute quack compatibility (Dao-AILab#2676)

* Update FA4 cute quack compatibility

* Use quack 0.5.3 make_smem_layout instead of vendored copy

Tri re-added the major_mode_size arg to quack.sm90_utils.make_smem_layout
in quack 0.5.3 (commit 68888e2), so FA4 no longer needs the local
sm90_layout helper. Revert the 4 backward call sites to quack's helper and
bump the floor to >=0.5.3 (0.5.2 lacks the arg).

---------

Co-authored-by: Johnsonms <lizhaofu@gmail.com>

* ci: install cutlass-dsl/quack at runtime to decouple from the baked image (Dao-AILab#2685)

* [Cute,Bwd,Sm100] Assume 16B stride divisibility for LSE/dPsum bulk-copy inputs (Dao-AILab#2686)

The SM100 backward stats (LSE, dPsum) are loaded via cp.async.bulk
(CopyBulkG2SOp), which - unlike cp.async.bulk.tensor - needs the source
pointer alignment provable at compile time. After slicing, the newer
cute-dsl can't deduce 16B alignment unless the input strides carry the
divisibility assumption, so the bulk copy fails to compile on real
tensors (the FakeTensor path masks it).

- flash_bwd_mla_sm100.py: add mdPsum to the new_stride divisibility list
  (it already covered ScaleP and the other stats; mdPsum was omitted).
- flash_bwd_sm100.py: the ordinary backward had no divisibility
  assumption at all; add it for both mLSE and mdPsum.

Only these two SM100 kernels use CopyBulkG2SOp; the SM90/SM80/SM120 and
MLA dK/dQ backward kernels use other copy paths and are unaffected.

Addresses the dPsum stride-divisibility finding (Finding 1) in Dao-AILab#2677.

* fix(hd256/sm100): forward reads actual input strides, drop .contiguous() patch (Dao-AILab#2670)

* follow up to Dao-AILab#2666: fixing the layouts in the sm100 hd256 kernels and removing the temporary fix of calling .contiguous everywhere

* respond to PR comments

* respond to PR comments-2: move to utils file

* Add tests

---------

Co-authored-by: drisspg <drisspguessous@gmail.com>

* ci: run MLA backward cases so CI exercises flash_bwd_mla_sm100.py (Dao-AILab#2690)

FA4_TEST_FILTER selected no MLA test, so the MLA backward kernels
(flash_bwd_mla_sm100.py + dq_dqv + dk) had zero CI coverage. Add four
small test_flash_attn_mla_absorbed cases covering the distinct backward
paths: sparse (kv_sparsity=True) non-causal and causal, dense
(kv_sparsity=False), and shared_kv=True. The ordinary SM100 backward is
already covered by the existing test_flash_attn_output cases.

Cold-cache cost on B200 (full 8-case filter): pass-1 compile ~4:54,
GPU run ~1:03 — well under the 60-min job timeout.

Stacked on Dao-AILab#2685 (runtime cutlass-dsl/quack install).

* Parallelize splitkv alignment templated kernels, remove flag (Dao-AILab#2683)

* [FA3] uv installation support (Dao-AILab#2458)

* Expose flash_attn_3 as package so imports work correctly.

* Add flash_attn_config package shim and fix uv packaging details

Builds on the flash_attn_3 package exposure so both import styles work
for downstream frameworks and uv/pyproject.toml installs:

- Add flash_attn_3/flash_attn_config.py re-export so
  `from flash_attn_3 import flash_attn_config` works (previously only the
  top-level module was importable), matching the interface shim.
- Un-ignore the committed shim in .gitignore; the bare `flash_attn_config.py`
  pattern (for the build-time generated top-level file) also matched the
  package shim and would have silently dropped it from the commit.
- Read flash_attn_3.__version__ from installed package metadata with a
  fallback, avoiding drift from setup.py's version source.
- README: move `dependencies` under `[project]` so the uv snippet is valid
  PEP 621.

Verified on H100 (SM90): editable `uv pip install -e .` now succeeds (fails
on main), both `import flash_attn_interface` and
`from flash_attn_3 import flash_attn_interface` resolve, `flash_attn_config`
imports both ways, and fp16 hdim128 forward matches a torch reference
(max_abs_err <= 2e-3). ruff check passes.

---------

Co-authored-by: Johnsonms <lizhaofu@gmail.com>

* [AMD ROCm] Enable RDNA backward and adopt CK unified workspace (Dao-AILab#2675)

* Add sink_ptr/d_sink_ptr to fmha_bwd_args to match updated CK submodule

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

* update submodule

* [CK_TILE] Use Unified Workspace for FMHA BWD (#182)

* [CK_TILE] Use Unified Workspace for FMHA BWD

Bump composable_kernel submodule to mono-split/users/yiding12/fmha-bwd-workspace
HEAD and adapt the FMHA BWD host wrappers to the new unified workspace API:

- Replace dq_acc tensor argument with workspace_ptr in get_ck_fmha_bwd_args
  / get_ck_fmha_varlen_bwd_args
- Drop dq_acc strides that have been removed from fmha_bwd_args
- In mha_bwd / mha_varlen_bwd, allocate the device workspace based on
  fmha_bwd_launcher::workspace_size and call launcher.prepare_workspace()
- Invoke launcher.run(args, stream_config) instead of fmha_bwd(...)

* Update CK pin as ROCm/rocm-libraries#6152 merged

* [CK_TILE] FMHA BWD: stream-async workspace prepare (#183)

* [CK_TILE] FMHA BWD: stream-async workspace prepare

Bump composable_kernel submodule to mono-split/users/yiding12/fmha-bwd-
async-prepare HEAD and adapt the FMHA BWD host wrappers to the new
async workspace prepare API (CK PR #7331):

- Replace launcher.prepare_workspace() with prepare_workspace_async(),
  which enqueues the full workspace setup (dq_acc zero, group-mode D2H
  of seqstart, host-side metadata pack via hipLaunchHostFunc, H2D back
  to device) on the caller's stream. No host-blocking sync remains in
  the BWD launch path.
- Pass a pinned_host_alloc lambda backed by PyTorch's CachingHostAllocator
  (torch::empty(..., pin_memory=true)). The launcher keeps the returned
  shared_ptr alive via a stream-tail hipLaunchHostFunc keepalive so the
  pinned buffer is not recycled while async copies are still in flight.
- mha_varlen_bwd: drop the cu_seqlens_q.cpu() / cu_seqlens_k.cpu() host
  copies; the launcher now reads device seqstart directly via async D2H.
  get_ck_fmha_varlen_bwd_traits no longer takes seqstart_qs/ks.

* [CK_TILE] FMHA BWD: bump CK submodule to develop tip (#7331 merged)

ROCm/rocm-libraries#7331 (async workspace prepare for FMHA BWD launcher)
landed on develop. Move csrc/composable_kernel from the pre-merge fork
tip ce838e19e5 to ROCm/composable_kernel develop tip 83566edb0f, which
is the split commit for #7331 (rocm-libraries 5692db0).

* [CK_TILE] FMHA BWD: explicit at::kCPU on pinned host TensorOptions

* Update CK and enable RDNA backward

---------

Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Co-authored-by: Yi DING <yi.ding@amd.com>
Co-authored-by: Hosang Yoon <hosang.yoon@amd.com>

* Fix CuTe SM120 compile-time argument handling (Dao-AILab#2671)

* Fix CuTe SM120 compile-time argument handling

* clean up

* guard empty SM120 local backward tiles

---------

Co-authored-by: Kevin-Li-2025 <2242139@qq.com>
Co-authored-by: drisspg <drisspguessous@gmail.com>

* [NVIDIA][CuTe,Fwd,sm120] Implement Pack-GQA on SM120 (+ graceful SplitKV fallback) (Dao-AILab#2656)

* [CuTe,Fwd,sm120] Fix use_tma_O crash on SM120 (issue Dao-AILab#2649)

On SM120 (Blackwell GeForce / RTX PRO 6000 / DGX Spark) the forward kernel set
`use_tma_O = self.arch >= Arch.sm_90`, enabling the TMA-based O-store epilogue.
But SM120 does not build the TMA store atom (tma_atom_O is None), so any forward
call crashes in cpasync.tma_partition with:

    AttributeError: 'NoneType' object has no attribute '_trait'

This makes the CuTe-DSL forward unusable on every SM120 GPU.

Restrict the TMA O-store to sm_90..sm_119, which is where the WGMMA-era epilogue
path is actually available:

    self.use_tma_O = Arch.sm_90 <= self.arch < Arch.sm_120

SM120 falls back to the non-TMA register->gmem O store (already used for the
SM80 path), which is correct and what the CpAsync SM120 kernel expects.

Verified on RTX PRO 6000 Blackwell (sm_120, cc 12.0), torch 2.12.0+cu130,
nvidia-cutlass-dsl 4.5.2: forward now runs and matches PyTorch SDPA reference
for hdim 64/96/128, causal and non-causal (max abs err <= 8e-3 in bf16). Before
this fix every SM120 forward call raised the AttributeError above.

* [CuTe,Fwd,sm120] Implement Pack-GQA on SM120; graceful SplitKV fallback

Pack-GQA was only half-wired in the SM80/SM120 CpAsync forward: the epilogue
referenced PackGQA.store_O/store_LSE, but the Q-load and head-indexing used the
plain (unpacked) path. So pack_gqa=True crashed in pack_gqa.store_O (crd2idx on a
packed (h_idx, m_idx) coordinate against an unpacked mO layout).

This implements Pack-GQA end to end on SM120 (and SM80), mirroring the SM90 path:
- Reshape mQ/mO (head_idx=2) and mLSE (head_idx=1) via pack_gqa_layout so
  qhead_per_kvhead folds into the seqlen mode ((qhead, seqlen)).
- Scheduler args use cute.size(mQ.shape[0]) (packed total rows) and seqlen_q_static
  = mQ.shape[0][1] (logical seqlen), so causal/mask q_idx stay correct.
- Kernel head-indexing: when pack_gqa, num_head from the scheduler already indexes
  the KV head (mQ/mK share nheads_kv); no division.
- Q-load: gather rows via PackGQA.load_Q (per-row (h_idx, m_idx) gmem pointers)
  instead of the contiguous local_tile path.

SplitKV (num_splits>1) is an SM100-only feature (SM80/SM90 also assert it
unsupported); SM120 has no forward+combine path. Fall back to num_splits=1, which
is numerically correct, instead of crashing in _check_type on the fp32 partials.

Verified on RTX PRO 6000 Blackwell (sm_120): pack_gqa=True matches PyTorch SDPA
GQA/MQA reference (err <= 8.4e-3 bf16) AND is bit-identical to the unpacked path
(max |packed - unpacked| = 0.0) across MHA/GQA/MQA, causal/non-causal, hd 64/128,
seqlen 512-2048. num_splits=3 falls back and matches reference (err 6.8e-4).
Stacked on the SM120 use_tma_O fix (Dao-AILab#2649).

* re-enable SM120 pack-gqa after rebase

* clean up SM120 pack-gqa split handling

* fix SM120 varlen pack-gqa offset

---------

Co-authored-by: drisspg <drisspguessous@gmail.com>

* Fix pre-commit

Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>

---------

Signed-off-by: oliver könig <okoenig@nvidia.com>
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: 鐘天楽 <tianle.zhong@bytedance.com>
Co-authored-by: brandonsun <brandons@nvidia.com>
Co-authored-by: Johnsonms <lizhaofu@gmail.com>
Co-authored-by: jayhshah <jayhshah@gmail.com>
Co-authored-by: oliver könig <okoenig@nvidia.com>
Co-authored-by: Michael Melesse <micmelesse@gmail.com>
Co-authored-by: Reuben Stern <107093092+reubenconducts@users.noreply.github.com>
Co-authored-by: Driss Guessous <32754868+drisspg@users.noreply.github.com>
Co-authored-by: aryan <aryansputta@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: sryap <17482891+sryap@users.noreply.github.com>
Co-authored-by: Yunwei Li <yunweili372423@gmail.com>
Co-authored-by: Zihao Wang <rekind133@outlook.com>
Co-authored-by: Anakin(Yancheng) Zheng <103552181+anakinxc@users.noreply.github.com>
Co-authored-by: Prashant Kumar <prashant.kumar@cohere.com>
Co-authored-by: Jane (Yuan) Xu <31798555+janeyx99@users.noreply.github.com>
Co-authored-by: Omar Attia <oy.attia@gmail.com>
Co-authored-by: drisspg <drisspguessous@gmail.com>
Co-authored-by: Björn Buschkämper <bjoern.buschkaemper@gmail.com>
Co-authored-by: rocking <ChunYu.Lai@amd.com>
Co-authored-by: Yi DING <yi.ding@amd.com>
Co-authored-by: Hosang Yoon <hosang.yoon@amd.com>
Co-authored-by: Yin Li <kxl474@student.bham.ac.uk>
Co-authored-by: Kevin-Li-2025 <2242139@qq.com>
Co-authored-by: Johnny <johnnynuca14@gmail.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

Copilot code review Copilot
Copilot review effort, defaults to Lite
Applies to this pull request for everyone.Learn more about Copilot code review.
Copilot left review comments
@drisspg drisspg drisspg approved these changes
@tridao tridao Awaiting requested review from tridao
@jayhshah jayhshah Awaiting requested review from jayhshah

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@Johnsonms @drisspg