[chore] release fastvideo-kernel 0.3.0

[SolitaryThinker]

fastvideo-kernel 0.3.0

The largest fastvideo-kernel release since the 0.2.x line. It adds FP4 quantized
attention for both inference and training on NVIDIA Blackwell, a VSA-256 Blackwell
fastpath built on FlashAttention-4 CuTe, variable-length block-sparse attention,
an index-native block-sparse API, and a round of build/packaging improvements
(GPU-less Docker builds, uv-based install).

Diff base: everything merged into fastvideo-kernel/ since 0.2.6 (PR #1094).

Highlights

• SageAttention3-style FP4 inference attention for Blackwell consumer GPUs (sm_120a / RTX 5090). New importable attn_qat_infer package. ([kernel] QAD 5090: Add modified SageAttention3 FP4 inference kernels (7/12) #1455, [feat] QAD 5090: Wire the Attn-QAT inference attention backend (8/12) #1457)
• Attn-QAT FP4 training Triton kernels — fused forward + backward with straight-through estimation and MXFP4 fake-quant. ([kernel] QAD 5090: Add Attn-QAT training Triton kernels (11/12) #1460)
• VSA-256 fastpath on Blackwell via FlashAttention-4 CuTe block-sparse attention, opt-in with Triton fallback. ([feat] VSA-256 fastpath on Blackwell via FA4 CuTe block-sparse attention #1354)
• Variable-length (varlen) block-sparse attention with Q/KV length mismatch support. ([kernel] Add varlen support for block-sparse attention #1319)
• Index-native block-sparse API (block_sparse_attn_from_indices) that skips the bool-mask round-trip. ([Perf] Skip bool-mask round-trip in block-sparse VSA attention #1243)
• GPU-less / Docker build support, conda-toolchain neutralization, and uv migration. ([bugfix]: build fastvideo-kernel on GPU-less Docker runners #1437, [Bugfix] QAD 5090: Torch.compile and other optimizations (15/12) #1466, migrate uv #1127, [misc]: standardize install instructions on uv pip install #1279)

---

New: FP4 quantized attention — inference (Blackwell)

A new top-level, importable package attn_qat_infer (shipped in the wheel
alongside fastvideo_kernel) provides modified SageAttention3 FP4 inference kernels
for Blackwell. Q/K/V are quantized to FP4 with FP8 (e4m3) scale factors during the
forward pass; outputs stay in BF16/FP16. (#1455, #1457)

• Entry point: attn_qat_infer.sageattn_blackwell(q, k, v, is_causal=False, per_block_mean=True, single_level_p_quant=True, sm_scale=None, ...) — inputs [B, H, L, D] in BF16/FP16.
• Quantization: single-level (default) or two-level softmax-matrix quantization; optional per-block query smoothing for stability.
• Built CUDA modules: fp4attn_cuda (block-scaled FP4 attention) and fp4quant_cuda (FP4 quantization), using CuTe/CUTLASS and Blackwell TMA.
• FastVideo integration: the main package wires these kernels in as a selectable attention backend.

Requirements / caveats

• NVIDIA Blackwell sm_120a (e.g. RTX 5090); CUDA Toolkit ≥ 12.8.
• Head dim ∈ {64, 128} (other sizes fall back to PyTorch SDPA). No dropout. GQA/MQA not yet supported (Q/K/V head counts must match). Sequences are padded to the 128-token block boundary internally and trimmed on output. ROCm not supported.

New: Attn-QAT FP4 training kernels (Triton)

Triton kernels for quantization-aware attention training, mirroring the inference
quantization scheme so models can be trained to match FP4 inference numerics. (#1460)

• Entry point: fastvideo_kernel.triton_kernels.attn_qat_train.attention (autograd Function) — fused forward + backward.
• Scheme: MXFP4 (nvfp4) fake-quant (quantize→dequantize) of attention weights P, with optional Q/K/V quantization; straight-through estimator on the backward pass; high-precision (BF16) accumulation. SageAttn-style K-smoothing on by default.
• Configurable: single- vs two-level P quantization, global vs per-block scale factors, Q/K/V backward quant toggle, warp specialization.
• Tested: causal / non-causal / cross-attention, fwd-only and fwd+bwd, batch 1–2, 4–40 heads, seq lengths up to ~9.4k, head dims 64/128, including a WAN-shaped [1, 40, 9360, 128] case (BF16).
• Note: currently an internal API — not re-exported at the fastvideo_kernel top level.

Block-sparse / VSA attention improvements

VSA-256 Blackwell fastpath (#1354). A FlashAttention-4 CuTe block-sparse kernel for
256-token KV blocks on Blackwell, exposed transparently through video_sparse_attn()
plus a new BSHD-native video_sparse_attn_bshd() entry point that avoids layout
round-trips on the CuTe path.

• Opt-in: set FASTVIDEO_VSA_CUTEDSL=1; transparently falls back to the Triton backend when the optional deps are absent.
• Optional deps: nvidia-cutlass-dsl>=4.5.0, torchvision, and FlashAttention CuTe pinned to flash-attention@c19cd20e (flash_attn/cute) for signature compatibility.
• Coverage: forward, cross-attention, and variable-block-size (VBS) variants, with a Triton-parity test.

Varlen block-sparse attention (#1319). New block_sparse_attn_varlen() handles
packed variable-length sequences (with mismatched Q/KV lengths) in a single launch via
cu_seqlens_q / cu_seqlens_kv, instead of one kernel call per sequence. Supports
forward and backward.

Perf: skip the bool-mask round-trip (#1243). Block-sparse kernels now operate on
compact index tensors natively.

• New preferred API block_sparse_attn_from_indices(q, k, v, q2k_idx, q2k_num, variable_block_sizes) avoids the bool-mask → index → bool conversions on both the forward and backward paths.
• Legacy block_sparse_attn(..., block_map, ...) is retained as a bool-mask wrapper for backward compatibility; both the SM90 and Triton backends were refactored to accept indices directly.

Build, packaging & install

• GPU-less / Docker builds ([bugfix]: build fastvideo-kernel on GPU-less Docker runners #1437, [Bugfix] QAD 5090: Torch.compile and other optimizations (15/12) #1466). build.sh no longer requires a live GPU — set TORCH_CUDA_ARCH_LIST (e.g. 9.0a, 12.0a) and the GPU probe is skipped. sm_120 is auto-detected to 12.0a so the CMake AUTO gate compiles the attn_qat_infer FP4 kernels instead of silently skipping them.
• Architecture gating (CMake). New FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER option (AUTO/ON/OFF; AUTO requires CUDA ≥ 12.8 + sm_120a), replacing the older ..._BUILD_MODIFIED_SAGE3. sm_90 keeps auto-enabling ThunderKittens (90a); other arches build to their detected arch.
• Conda toolchain neutralization ([feat] [3/n] Improve API: extend support to cli #1226). build.sh clears conda-injected CC/CXX/CFLAGS/CMAKE_ARGS and discovers the system compiler, avoiding nvcc conflicts with mismatched conda toolchains.
• Packaging. MANIFEST.in and wheel.packages now include the attn_qat_infer/ tree (*.py/*.cu/*.cuh/*.cpp/*.h) so the new kernels ship in the sdist/wheel. Triton is now gated to Linux: triton>=2.0.0; sys_platform == 'linux'.
• uv migration (migrate uv #1127, [misc]: standardize install instructions on uv pip install #1279). README and build.sh standardize on uv pip install; README documents required CUDA_HOME/CUDACXX and the optional FA4 CuTe backend for the VSA-256 fastpath.

New / changed public API

Symbol	Module	Status
sageattn_blackwell	attn_qat_infer	New — FP4 inference attention (Blackwell)
video_sparse_attn_bshd	fastvideo_kernel	New — BSHD-native VSA entry point
block_sparse_attn_from_indices	fastvideo_kernel	New — index-native block-sparse attention
block_sparse_attn_varlen	fastvideo_kernel	New — varlen block-sparse attention
block_sparse_attn	fastvideo_kernel	Retained (bool-mask wrapper; internals refactored)
attn_qat_train.attention	fastvideo_kernel.triton_kernels	New — QAT training kernel (internal)

Requirements & compatibility

• FP4 inference kernels: NVIDIA Blackwell sm_120a, CUDA Toolkit ≥ 12.8, Linux.
• QAT training kernels and the VSA Triton path: Triton (Linux).
• VSA-256 fastpath: optional nvidia-cutlass-dsl>=4.5.0 + FlashAttention CuTe (@c19cd20e); otherwise Triton fallback.
• Published wheel: Python 3.12, torch 2.12 / CUDA 13 (cu130) — includes the Blackwell sm_120a FP4 kernels. A torch 2.12 / CUDA 12.6 (cu126) build is also produced (Hopper-only, no FP4) and kept as a build/release artifact.

Install

uv pip install fastvideo-kernel==0.3.0

PRs since 0.2.6

• [Bugfix] QAD 5090: Torch.compile and other optimizations (15/12) #1466 — QAD 5090: Torch.compile and other optimizations (build/arch detection) (15/12)
• [kernel] QAD 5090: Add Attn-QAT training Triton kernels (11/12) #1460 — QAD 5090: Attn-QAT training Triton kernels (11/12)
• [feat] QAD 5090: Wire the Attn-QAT inference attention backend (8/12) #1457 — QAD 5090: Wire the Attn-QAT inference attention backend (8/12)
• [kernel] QAD 5090: Add modified SageAttention3 FP4 inference kernels (7/12) #1455 — QAD 5090: Add modified SageAttention3 FP4 inference kernels (7/12)
• [bugfix]: build fastvideo-kernel on GPU-less Docker runners #1437 — Build fastvideo-kernel on GPU-less Docker runners
• [feat] VSA-256 fastpath on Blackwell via FA4 CuTe block-sparse attention #1354 — VSA-256 fastpath on Blackwell via FA4 CuTe block-sparse attention
• [kernel] Add varlen support for block-sparse attention #1319 — Add varlen support for block-sparse attention
• [misc]: standardize install instructions on uv pip install #1279 — Standardize install instructions on uv pip install
• [Perf] Skip bool-mask round-trip in block-sparse VSA attention #1243 — Perf: skip bool-mask round-trip in block-sparse VSA attention
• [feat] [3/n] Improve API: extend support to cli #1226 — Improve API: extend support to CLI (kernel build.sh conda fix)
• migrate uv #1127 — Docs/kernel: migrate to uv
• Packaging: Triton gated to Linux via sys_platform ([feat] Job Runner UI #1172/[feat] Job Runner UI #1189); ops.py fallback cleanup

[gemini-code-assist]

Code Review

This pull request updates the version of fastvideo-kernel from 0.2.6 to 0.3.0 across several files, including configuration files, dependency lists, and documentation. There are no review comments, and I have no feedback to provide.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

[mergify]

Merge Protections

Your pull request matches the following merge protections and will not be merged until they are valid.

🔴 PR merge requirements

Waiting for

• #approved-reviews-by>=1
• check-success=fastcheck-passed
• check-success=full-suite-passed

This rule is failing.

• #approved-reviews-by>=1
• check-success=fastcheck-passed
• check-success=full-suite-passed
• check-success~=pre-commit
• title~=(?i)^\[(feat|feature|bugfix|fix|refactor|perf|ci|doc|docs|misc|chore|kernel|new.?model|skill|skills|infra)\]