Commit Graph

  • 72c7ba484d Fix Hopper tests (#2242) main Driss Guessous 2026-02-08 09:25:01 -08:00
  • 2a8d39c540 [DSL] warpgroup_reg_alloc -> setmaxregister_increase Tri Dao 2026-02-08 22:17:36 +07:00
  • 17d29436b8 [Bwd,Sm90] Have score_mod and score_mod_bwd as partial functions Tri Dao 2026-02-08 22:10:05 +07:00
  • deb183092b [Bwd,Sm100] Shorten PipelineTmaUmma create Tri Dao 2026-02-08 21:11:17 +07:00
  • c912a37d52 [Bwd,Sm90] Use quack.copy_utils Tri Dao 2026-02-08 18:57:20 +07:00
  • b9148cec6f [Layout] Use layout_utils.transpose_view and select from quack Tri Dao 2026-02-08 18:24:28 +07:00
  • 90f10faafd [DSL] Use cute.math.{exp2,log2,log} Tri Dao 2026-02-08 18:02:19 +07:00
  • 8dd8019cef [Layout] Use quack.layout_utils.mma_partition_C_vec Tri Dao 2026-02-08 17:05:48 +07:00
  • b735ef24c2 [Layout] Use reshape_acc_to_mn and reshape_acc_to_frgA from quack Tri Dao 2026-02-08 16:48:16 +07:00
  • 7edcf59c9e [DSL] Use cute.arch.warp_reduction_{max,sum} Tri Dao 2026-02-08 16:18:56 +07:00
  • 81f2c2dcdc [Sm90] Use functions from quack.sm90_utils Tri Dao 2026-02-08 15:52:49 +07:00
  • d39b6292bb [DSL] Remove coord_offset_i64, domain_offset_i64, elem_pointer_i64 Tri Dao 2026-02-08 11:11:43 +07:00
  • 5a66f2cca3 [DSL]Replace utils.{fma,mul,add}_packed_f32x2 with cute.arch version Tri Dao 2026-02-08 11:07:39 +07:00
  • a804a5a3ef [DSL] Replace old fence with cute.arch.fence_view_async_shared() Tri Dao 2026-02-08 10:48:54 +07:00
  • 48af662c53 pytest-dist round robin to gpus (#2241) Driss Guessous 2026-02-07 19:50:48 -08:00
  • abaa87875d [CUTE]Bump to Cutedsl (#2216) Driss Guessous 2026-02-07 18:32:44 -08:00
  • 912c6c4518 Add FLASH_ATTENTION_TRITON_AMD_CONFIG_JSON env var support (#2239) Alex Butler 2026-02-07 16:13:18 +00:00
  • f1284cff5d hdim 192 smem fix (#2235) jayhshah 2026-02-05 08:22:36 -08:00
  • e2743ab5b3 [FA3] Mark current main version as v3.0.0 stable (#2223) Luca Wehrstedt 2026-02-05 04:45:41 +01:00
  • 24445c0c17 short readme for flex flash (#2231) Markus Hoehnerbach 2026-02-04 18:30:56 -08:00
  • ef9e6a6441 Use TORCH_TARGET_VERSION over TORCH_STABLE_ONLY (#2155) Jane (Yuan) Xu 2026-02-04 16:58:49 -05:00
  • 188643b82d Fix shared-memory race (#2229) Driss Guessous 2026-02-04 12:33:20 -08:00
  • 514e63cc26 fix compute_block_sparsity usage in benchmark_mask_mod (#2221) zhuochen 2026-02-03 05:41:42 +08:00
  • 701ebe0578 [AMD] Triton Backend for ROCm #3 (#2178) Michael Melesse 2026-01-28 10:49:08 -05:00
  • 99589e5a66 [DSL] Optionally patch cute-dsl to use system's ptxas Tri Dao 2026-01-27 21:09:36 +07:00
  • 4f892461bb [Flex][SM100] Replay expand fix on sm100 (#2209) Driss Guessous 2026-01-26 12:22:14 -07:00
  • 438325c2c3 Update README to include 'psutil' package as build requirement (#2210) Wang Lecheng 2026-01-25 17:41:59 +08:00
  • 57cef6c2e7 ci: Use 1 ninja job for cu13 (#2195) oliver könig 2026-01-24 04:04:16 +01:00
  • 2580b5a488 [Cute][Flex] Allow q_offset 1 and add block-sizes to disambiguate edge cases (#2187) Driss Guessous 2026-01-21 17:25:32 -08:00
  • f15ccf5ff2 reduce chance of build oom (#2079) Qubitium-ModelCloud 2026-01-21 18:36:22 +08:00
  • 04e6ee1fb5 [Cute, SM90] fix fwd varlen Cute implementation bug for H100 (#2194) Kareem 2026-01-21 00:42:46 +07:00
  • a0f9f418fd [Cute][Flex] Fix expanded tensor bug (#2189) Driss Guessous 2026-01-16 20:47:10 -08:00
  • 2d6b146893 Merge pull request #2185 from henrylhtsang/test_local_r2p Markus Hoehnerbach 2026-01-16 16:06:04 -08:00
  • 137ad8e6e0 add comment Henry Tsang 2026-01-16 09:39:15 -08:00
  • 2e6ae05b4b added back clamp to avoid "OverflowError: Python int too large to convert to C long" Henry Tsang 2026-01-16 09:26:19 -08:00
  • e94012ac16 lint Henry Tsang 2026-01-15 20:40:24 -08:00
  • 94f034800e doc Henry Tsang 2026-01-15 20:38:42 -08:00
  • 08e65188b5 remove 24 clamp Henry Tsang 2026-01-15 20:27:15 -08:00
  • e34d84057d remove zero logic for right_s and left_s henrylhtsang 2026-01-15 19:43:44 -08:00
  • ac8885812e flip in_bound to out_bound Henry Tsang 2026-01-15 17:48:34 -08:00
  • e4ec1ad333 switch from xor to mask_right & ~ mask_left henrylhtsang 2026-01-15 17:10:03 -08:00
  • 7108d1c854 Add R2P dual bound masking for local attention henrylhtsang 2026-01-15 15:10:19 -08:00
  • 2020964fc8 remove benchmark result, undo changes to benchmark Henry Tsang 2026-01-15 14:55:59 -08:00
  • a512bd8c7c Add R2P dual bound masking for local attention henrylhtsang 2026-01-15 14:51:19 -08:00
  • fffabc3de1 [Cute,Fwd,Sm100] distributed offset calculation for paged KV (#2104) timmy-feng 2026-01-15 15:11:01 -05:00
  • 88067b00de baseline local flops henrylhtsang 2026-01-15 11:46:15 -08:00
  • 68649fb784 [Cute][Flex]Add pack-gqa divmod (#2180) Driss Guessous 2026-01-15 10:09:54 -08:00
  • 506441a3fc [Cute][Flex] add back in contig (#2177) Driss Guessous 2026-01-14 17:04:02 -08:00
  • 13696f2e5e [Cute] update row_max before safe overwrite for online_softmax (#2174) jayhshah 2026-01-13 13:09:43 -08:00
  • 4894657ee0 [Cute][Flex] Remove no longer needed contig (#2172) Driss Guessous 2026-01-12 14:08:33 -08:00
  • 4cb272ed75 [Cute,Flex] Add option to create and cache __cute_hash__ (#2171) Reuben Stern 2026-01-12 14:22:18 -05:00
  • dbf08eb8e1 Merge pull request #2156 from v0i0/v0i0/improve-paged-ldgsts Markus Hoehnerbach 2026-01-12 10:10:03 -08:00
  • ef7343b4cf [CUTE] Bump cutedsl to 4.3.5 (#2170) Driss Guessous 2026-01-12 09:51:09 -08:00
  • ea8f735063 [Cute,Bwd,Sm100] fix seqused in varlen bwd (#2167) jayhshah 2026-01-10 11:24:52 -08:00
  • 5d4c9537a1 [CUTE][SM90] GQA backward non deterministic (#2158) Driss Guessous 2026-01-09 17:42:22 -08:00
  • 8eff546d05 [CUTE][SM90]Enable pack-gqa with broadcasted maskmods (#2145) Driss Guessous 2026-01-09 17:40:52 -08:00
  • 26d4ee96a3 [CUTE][SM100] Fix backward gqa on sm100 post mask-mod semantic change (#2146) Driss Guessous 2026-01-09 16:58:58 -08:00
  • e317aa4978 [Cute] Clarify and fix subtle cachekey bug (#2143) Driss Guessous 2026-01-09 16:58:25 -08:00
  • 844b10fe36 score-mod backward SM90 (#2137) Driss Guessous 2026-01-09 16:56:47 -08:00
  • 27a3b54c45 block-sparse backward SM90 (#2136) Driss Guessous 2026-01-09 16:54:12 -08:00
  • ed6a82f050 [Cute, Bwd, Sm100] Add varlen for sm100 bwd (#2150) jayhshah 2026-01-09 15:24:29 -08:00
  • c15ffe3caa cleanup Markus Hoehnerbach 2026-01-08 22:16:15 +00:00
  • 6dd7e742df [Cute] Fix two tests that were failing (#2149) Henry Tsang 2026-01-08 10:42:26 -08:00
  • 3c8ca4e255 [Cute,Fwd,Sm100] Support q_stage=1 for inference (#1993) timmy-feng 2026-01-08 13:10:25 -05:00
  • 3e87e421f8 Update cutlass to fix undefined symbol: cuDriverGetVersion. (#2142) Qi Yuhang 2026-01-08 03:35:30 +08:00
  • f472175a9e Misc tests that should be xfailed for now (#2127) Driss Guessous 2026-01-05 15:09:06 -08:00
  • bb2efb3329 [Cute] Fix minor lint issue in shuffle_sync Tri Dao 2026-01-05 11:36:41 -05:00
  • f98d345f0a [Cute,Fwd] improved block sparsity (#2100) Reuben Stern 2026-01-05 11:35:13 -05:00
  • 9b6dbaceb6 Add pack-gqa support for blcoksparse impl w/ braodcasted H dim (#2098) Driss Guessous 2026-01-03 17:55:00 -08:00
  • f3423a8a29 [Cute,Fwd,Sm90] Change PipelineTMAAsync sublass to signal per warp Tri Dao 2025-12-31 19:44:13 -05:00
  • 4fd123e4b5 [Cute] Add quack as dependency Tri Dao 2025-12-31 18:15:26 -05:00
  • d234051316 Enable Thor (#2108) Johnny 2025-12-29 23:14:47 +01:00
  • 11b32fd20a improve paged cpasync Markus Hoehnerbach 2025-12-24 00:30:09 +00:00
  • 58fe37fba6 fix shuffle sync for pack gqa epilogue (#2097) jayhshah 2025-12-23 16:05:35 -08:00
  • 5663adffb0 README for AMD ROCm (#2068) seungrokj 2025-12-23 20:51:26 +09:00
  • ceb4110f7e Update to dsl 3.4.3 (#2092) Driss Guessous 2025-12-21 21:13:28 -05:00
  • bba578d439 Fix IMA in fwd on m boundary (#2091) Driss Guessous 2025-12-20 04:52:41 -08:00
  • eacbc560be Blocksparse impl (#2085) Driss Guessous 2025-12-18 14:45:35 -08:00
  • ac9b5f107f basics working (#2070) Driss Guessous 2025-12-16 09:19:20 -08:00
  • 0a5339f4cb [FIRST] Fix softcap scoremod kwargs typo. (#2072) Leo Dong 2025-12-15 19:30:14 -08:00
  • 179f793bbc [CUTE] Seeing if tvvm reduces cpu overhead (#2042) Driss Guessous 2025-12-15 15:40:43 -08:00
  • fd8d5eb363 [Cute,Fwd] Extend score_mod to variable sequence length (#2043) Reuben Stern 2025-12-15 17:16:30 -05:00
  • e240e0f7e4 [CUTE] Allow grads to be preallocated (#2065) Driss Guessous 2025-12-14 21:57:26 -05:00
  • bc0e4ac014 Fix use-after-free in FA3 deterministic mode. The pytorch caching allocator actually saves us here, but if you turn it off, then compute-sanitizer will detect this. (#2063) skarupke 2025-12-12 09:47:31 -05:00
  • c783ab2f7e fixing cute bwd func def (#2056) liangel-02 2025-12-09 16:34:16 -05:00
  • 6328432b8b [AMD ROCm] Update to latest composable_kernel to improve performance (#2052) rocking 2025-12-08 04:11:01 +08:00
  • 0d1ad61b7f Add hash attr to shortcut expensive check (#2048) Driss Guessous 2025-12-06 17:42:52 -08:00
  • 56fdf3e232 [Cute,Bwd,Sm100] Add local for sm100 bwd (#2046) jayhshah 2025-12-05 16:11:10 -08:00
  • 59df2f9bdf Merge pull request #1769 from guilhermeleobas/guilhermeleobas/fa3-compile Markus Hoehnerbach 2025-12-04 11:08:27 -08:00
  • 26ba559ee1 Add headdim64 tests (#2041) Driss Guessous 2025-12-02 15:18:59 -08:00
  • de6a6ad08b [Flash] Fix head dim 64 bwd (#2035) Driss Guessous 2025-12-02 13:58:14 -08:00
  • 91ba87d759 ruff all the smaller files (#2040) Driss Guessous 2025-12-02 13:43:24 -08:00
  • 672381f72c Bump pin (#2025) Driss Guessous 2025-11-25 12:38:30 -08:00
  • 92ca9da8d6 [Cute,Fwd] enable mask mod without blocksparsity (#2031) Reuben Stern 2025-11-25 00:43:48 -05:00
  • 63b66f2cd9 [Cute] Add authors Tri Dao 2025-11-24 23:45:34 -05:00
  • 5cc6fa48f9 Add LICENSE and AUTHORS to flash_attn/cute (#2032) Jean-Luc Duprat 2025-11-24 20:42:02 -08:00
  • 91942973d5 [NFC] Trivial fix to silence linter (#1928) Jean-Luc Duprat 2025-11-24 20:41:20 -08:00
  • 20cda05e6b [Cute,Bwd,Sm100] enable deterministic mode for sm100 bwd and fix race conditions (#2033) jayhshah 2025-11-24 17:33:08 -08:00
  • a986d0190e swap order of decorators (#2029) Anakin(Yancheng) Zheng 2025-11-24 09:51:17 +08:00
  • d063b333ba don't pass mask_fn to softmax_step generically (#2026) jayhshah 2025-11-21 18:33:53 -08:00