Wenbo Ji 嵇文博
ef63de1f60
[perf] Fuse VSA64 tile layout for SM100 inference ( #1900 )
...
For Wan VSA inference with 64-token tiles, each attention layer scattered Q, K,
V and the gate into a tile buffer and then made four transposed copies. On
SM100, one Triton kernel (tile_to_bhsd) now does the scatter and the transposes
in one pass. Training, 256-token tiles, other GPUs and builds without the kernel
use the previous path. FASTVIDEO_DISABLE_VSA64_FUSED_LAYOUT=1 turns the fused
path off.
On GB200 the layout and attention step goes from 15.7 ms to 8.6 ms. FastWan 1.3B
frames are bit-exact with the previous path at sequence parallel 1 and 2.
2026-10-07 15:31:58 -07:00
Wenbo Ji 嵇文博
b7bd7775ed
[perf] Default SM100 long-sequence Attn-QAT tiling ( #1899 )
...
On SM100, the Attn-QAT training kernel now uses a faster route for non-causal
BF16 attention with equal query and key lengths of 16384 tokens or more. The
forward pass processes full KV tiles without bounds masks, and a separate masked
step handles the partial tail tile. The backward pass uses 64x128 dQ and dK/dV
tiles. Outputs, softmax statistics and all gradients are bitwise identical to the
previous route. FASTVIDEO_ATTN_QAT_SM100_OPTIMIZED=0 turns off the SM100 route,
and FASTVIDEO_ATTN_QAT_SM100_WIDE_BWD=0 keeps the 64x64 backward tiles.
On GB200 with 6 heads and 31200 tokens, the forward kernel time goes from 9.14 ms
to 6.96 ms and the backward kernel time from 17.41 ms to 15.85 ms.
2026-10-07 15:14:38 -07:00
f82c190cbb
[kernel] Build + allow attn_qat_infer FP4 attention on sm_121a (DGX Spark) ( #1598 )
...
Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com >
Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com >
2026-10-07 04:09:39 -07:00
bc0191d775
[perf] Harden Ulysses agreement and tune bounded H3 transfers ( #1807 )
...
Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com >
Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com >
2026-10-07 04:01:53 -07:00
72795d0af8
[perf] avoid full-sequence materialization in VSA coarse/sparse combine ( #1813 )
...
Co-authored-by: Claude Opus 5 <noreply@anthropic.com >
Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com >
2026-10-07 01:54:35 -07:00
8b51380466
[feat]: FastH3 on single RTX GPUs ( #1919 )
...
Co-authored-by: Aryan Kumar <aryan5v@users.noreply.github.com >
Co-authored-by: Aryan Kumar <aryank@Aryans-Mac-Studio.local >
Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com >
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com >
2026-10-06 11:37:01 -07:00
lpc0220 and Claude Opus 5.5
0cc41a22dc
[bugfix] sm_100a VSA: fence the CLC response read before the slot release; forward PV descriptor lifetime ( #1904 )
...
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com >
2026-10-02 12:37:34 -07:00
e1f3904799
[kernel] sm_100a CUDA backward for VSA block-sparse attention, 128-token blocks ( #1866 )
...
Co-authored-by: Pengcheng Li <pengchengl@ptyche0203.ptyche.clusters.nvidia.com >
Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com >
2026-09-29 16:57:22 -07:00
Junda Su
e3b88bb12a
[feat] Add a typed environment-variable registry, policy doc, and contract test ( #1896 )
2026-09-29 12:50:45 -07:00
Yogya Mehrotra
017ce6602d
[bugfix] Fix dead NVFP4 QAT plumbing: bench imports, int8 error string, smooth_q crash ( #1723 )
2026-09-24 17:09:33 -07:00
Yogya Mehrotra
a575055eec
[perf] benchmark_attn_qat_train: device peak-TFLOPS table instead of silent RTX 5090 default ( #1724 )
2026-09-24 17:08:03 -07:00
Keith
f47ad3f5b7
[bugfix] fastvideo-kernel: install the Python + Triton package when the HIP toolchain cannot be configured ( #1871 )
2026-09-21 11:59:13 -07:00
lpc0220 and Claude Fable 5.1
aef4e9b3b1
[kernel] VSA kernel: one sm_100a / sm_103a image per listed arch; un-gate backward on sm_103a ( #1833 )
...
Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com >
2026-09-12 16:49:36 -07:00
lpc0220
c993d7393e
[kernel] sm_100a CUDA backward for VSA block-sparse attention (blk64) ( #1819 )
2026-09-06 21:00:18 -07:00
SYLAR
7bb76b5ec9
[feat] Add native SM103a VSA support ( #1812 )
...
Signed-off-by: lishunyang12 <lishunyang12@163.com >
2026-09-03 15:00:28 -07:00
William Lin
1aed667377
[chore] release fastvideo-kernel 0.3.5 ( #1775 )
2026-08-27 23:18:59 -07:00
William Lin
c1612ff397
[bugfix]: pin fastvideo-kernel to Torch 2.12.0 ( #1774 )
2026-08-27 23:13:54 -07:00
Raghav K
b2062556a9
[perf] VSA Triton: widen the autotune num_stages range (the optimum was outside it) ( #1706 )
2026-08-26 12:30:56 -07:00
lpc0220
7a4285189f
[kernel] Route block-sparse VSA to the sm_100a forward behind FASTVIDEO_VSA_SM100A (opt-in) ( #1754 )
2026-08-24 15:26:56 -07:00
William Lin
f9e3680f11
[perf] Align FastH3 optimized inference profile ( #1748 )
2026-08-23 02:01:03 -07:00
Shao Duan
c041318f2c
[perf] Add fused NVLink all-to-all for Ulysses ( #1740 )
2026-08-22 23:06:24 -07:00
William Lin and Claude Fable 5
56d4a6074f
[bugfix] fastvideo-kernel: fix Triton block-sparse backward logit scaling (bf16 K pre-scaling) ( #1730 )
...
Co-authored-by: Claude Fable 5 <noreply@anthropic.com >
2026-08-21 04:50:06 -05:00
lpc0220
907f2100ec
[kernel] sm_100a CUDA block-sparse VSA forward (Blackwell), 64- and 128-token blocks ( #1719 )
2026-08-20 23:16:25 -05:00
00338aa9ca
[perf] Add FA4 CuTe backward support for VSA-256 ( #1639 )
...
Co-authored-by: Hyunsung Lee <hyunsungl@sizigistudios.com >
Co-authored-by: alexzms <3036648523@qq.com >
2026-08-19 11:46:49 -07:00
William Lin
9028953625
[misc] yapf pass under CI's interpreter (3.12) + pin hook language_version ( #1702 )
2026-08-08 22:24:17 -07:00
Mac Lee
7a592ff09a
[ci]: cache FastVideo kernel builds in Modal ( #1562 )
2026-07-26 04:21:09 -07:00
9709b7513b
[feat] Port NVFP4 QAT/QAD to modular train framework ( #1619 )
...
Co-authored-by: Peiyuan Zhang <email>
Co-authored-by: Peiyuan <a>
2026-07-18 15:01:09 -07:00
William Lin
98be9b3da2
[bugfix]: address the three remaining #1447 review findings ( #1549 )
2026-07-05 06:43:27 -07:00
William Lin
40a8bd2d3b
[bugfix]: skip ThunderKittens kernels on aarch64 and document the kernel build matrix ( #1548 )
2026-07-05 06:13:11 -07:00
William Lin
0d626ef2d1
[kernel] Bump fastvideo-kernel pin to 0.3.1 and version the FA4 tile_mn port as 0.3.2 ( #1539 )
2026-07-03 14:11:08 -07:00
William Lin
b36d0ef085
[infra] Add DGX Spark and multi-architecture CUDA support ( #1447 )
2026-07-03 13:41:43 -07:00
William Lin
31e26abec4
[chore] release fastvideo-kernel 0.3.1 ( #1520 )
2026-06-30 12:25:27 -07:00
sumyyyyy and SolitaryThinker
9e83ba630c
[kernel] Extract VSA utility functions into fastvideo_kernel ( #1408 )
...
Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com >
2026-06-30 12:19:25 -07:00
William Lin
3d36160fc4
[bugfix] -fsigned-char so ThunderKittens compiles on aarch64 (Grace Hopper) ( #1515 )
2026-06-29 14:05:38 -07:00
William Lin
2656d6530c
[bugfix] compile FP4 (attn_qat_infer) kernels for sm_120a only via per-arch split ( #1508 )
2026-06-29 11:07:27 -07:00
Kevin Lin
e3f54e7169
[bugfix] Fix causal attention mask for Blackwell FP4 MMA column layout ( #1506 )
2026-06-28 20:36:40 -07:00
William Lin
4ba5681307
[bugfix] guard ThunderKittens Hopper kernels for non-sm_90a device passes ( #1507 )
2026-06-28 20:07:43 -07:00
Kevin Lin
719a1879bd
[bugfix] Remove incorrect V-row permutation in scaled_fp4_quant_trans_kernel ( #1493 )
2026-06-25 03:05:12 -07:00
xsank and mergify[bot]
31b719ae49
[perf] optimize compress & topk kernel ( #1421 )
...
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-23 18:51:22 +00:00
William Lin
b1d89eba1f
[chore] release fastvideo-kernel 0.3.0 ( #1478 )
2026-06-23 10:02:27 -07:00
b2ade71467
[Bugfix] QAD 5090: Torch.compile and other optimizations (15/12) ( #1466 )
...
Co-authored-by: alexzms <3036648523@qq.com >
Co-authored-by: Kaiqin Kong <k1kong@ucsd.edu >
Co-authored-by: William Lin <SolitaryThinker@users.noreply.github.com >
Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com >
2026-06-23 01:23:57 -07:00
sumyyyyy
87f98c9b8b
[kernel] Add varlen support for block-sparse attention ( #1319 )
2026-06-18 01:06:00 +00:00
eed9c4bfbf
[kernel] QAD 5090: Add Attn-QAT training Triton kernels (11/12) ( #1460 )
...
Co-authored-by: William Lin <SolitaryThinker@users.noreply.github.com >
Co-authored-by: Loay Rashid <42599591+loaydatrain@users.noreply.github.com >
Co-authored-by: Kaiqin Kong <k1kong@ucsd.edu >
2026-06-17 13:35:12 -07:00
88e753f281
[feat] QAD 5090: Wire the Attn-QAT inference attention backend (8/12) ( #1457 )
...
Co-authored-by: William Lin <SolitaryThinker@users.noreply.github.com >
Co-authored-by: Loay Rashid <42599591+loaydatrain@users.noreply.github.com >
Co-authored-by: Kaiqin Kong <k1kong@ucsd.edu >
Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com >
Co-authored-by: alexzms <26690162+alexzms@users.noreply.github.com >
2026-06-16 13:12:33 -07:00
77832059cc
[kernel] QAD 5090: Add modified SageAttention3 FP4 inference kernels (7/12) ( #1455 )
...
Co-authored-by: William Lin <SolitaryThinker@users.noreply.github.com >
Co-authored-by: Loay Rashid <42599591+loaydatrain@users.noreply.github.com >
Co-authored-by: Kaiqin Kong <k1kong@ucsd.edu >
Co-authored-by: Edenzzzz <wtan45@wisc.edu >
2026-06-16 12:02:30 -07:00
William Lin
1c627a3f98
[bugfix]: build fastvideo-kernel on GPU-less Docker runners ( #1437 )
2026-06-06 17:52:45 -07:00
alexzms
be548a78ea
[feat] VSA-256 fastpath on Blackwell via FA4 CuTe block-sparse attention ( #1354 )
2026-05-26 12:58:29 -07:00
William Lin
140bd1a6cf
[misc]: standardize install instructions on uv pip install ( #1279 )
2026-05-02 12:45:50 -07:00
Mook
7b872cc41e
[Perf] Skip bool-mask round-trip in block-sparse VSA attention ( #1243 )
2026-04-26 15:14:37 -07:00
William Lin
24ced500f5
[test] add LTX-2 distilled T2V SSIM regression test ( #1240 )
2026-04-21 12:03:38 -07:00
Wenbo Ji 嵇文博