Metadata-Version: 2.1
Name: fastvideo-kernel
Version: 0.3.5
Summary: Unified CUDA kernels for FastVideo
Author-Email: Hao AI Lab <contact@haoailab.com>
License:                                  Apache License
                                    Version 2.0, January 2004
                                 http://www.apache.org/licenses/
         
            TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
         
            1. Definitions.
         
               "License" shall mean the terms and conditions for use, reproduction,
               and distribution as defined by Sections 1 through 9 of this document.
         
               "Licensor" shall mean the copyright owner or entity authorized by
               the copyright owner that is granting the License.
         
               "Legal Entity" shall mean the union of the acting entity and all
               other entities that control, are controlled by, or are under common
               control with that entity. For the purposes of this definition,
               "control" means (i) the power, direct or indirect, to cause the
               direction or management of such entity, whether by contract or
               otherwise, or (ii) ownership of fifty percent (50%) or more of the
               outstanding shares, or (iii) beneficial ownership of such entity.
         
               "You" (or "Your") shall mean an individual or Legal Entity
               exercising permissions granted by this License.
         
               "Source" form shall mean the preferred form for making modifications,
               including but not limited to software source code, documentation
               source, and configuration files.
         
               "Object" form shall mean any form resulting from mechanical
               transformation or translation of a Source form, including but
               not limited to compiled object code, generated documentation,
               and conversions to other media types.
         
               "Work" shall mean the work of authorship, whether in Source or
               Object form, made available under the License, as indicated by a
               copyright notice that is included in or attached to the work
               (an example is provided in the Appendix below).
         
               "Derivative Works" shall mean any work, whether in Source or Object
               form, that is based on (or derived from) the Work and for which the
               editorial revisions, annotations, elaborations, or other modifications
               represent, as a whole, an original work of authorship. For the purposes
               of this License, Derivative Works shall not include works that remain
               separable from, or merely link (or bind by name) to the interfaces of,
               the Work and Derivative Works thereof.
         
               "Contribution" shall mean any work of authorship, including
               the original version of the Work and any modifications or additions
               to that Work or Derivative Works thereof, that is intentionally
               submitted to Licensor for inclusion in the Work by the copyright owner
               or by an individual or Legal Entity authorized to submit on behalf of
               the copyright owner. For the purposes of this definition, "submitted"
               means any form of electronic, verbal, or written communication sent
               to the Licensor or its representatives, including but not limited to
               communication on electronic mailing lists, source code control systems,
               and issue tracking systems that are managed by, or on behalf of, the
               Licensor for the purpose of discussing and improving the Work, but
               excluding communication that is conspicuously marked or otherwise
               designated in writing by the copyright owner as "Not a Contribution."
         
               "Contributor" shall mean Licensor and any individual or Legal Entity
               on behalf of whom a Contribution has been received by Licensor and
               subsequently incorporated within the Work.
         
            2. Grant of Copyright License. Subject to the terms and conditions of
               this License, each Contributor hereby grants to You a perpetual,
               worldwide, non-exclusive, no-charge, royalty-free, irrevocable
               copyright license to reproduce, prepare Derivative Works of,
               publicly display, publicly perform, sublicense, and distribute the
               Work and such Derivative Works in Source or Object form.
         
            3. Grant of Patent License. Subject to the terms and conditions of
               this License, each Contributor hereby grants to You a perpetual,
               worldwide, non-exclusive, no-charge, royalty-free, irrevocable
               (except as stated in this section) patent license to make, have made,
               use, offer to sell, sell, import, and otherwise transfer the Work,
               where such license applies only to those patent claims licensable
               by such Contributor that are necessarily infringed by their
               Contribution(s) alone or by combination of their Contribution(s)
               with the Work to which such Contribution(s) was submitted. If You
               institute patent litigation against any entity (including a
               cross-claim or counterclaim in a lawsuit) alleging that the Work
               or a Contribution incorporated within the Work constitutes direct
               or contributory patent infringement, then any patent licenses
               granted to You under this License for that Work shall terminate
               as of the date such litigation is filed.
         
            4. Redistribution. You may reproduce and distribute copies of the
               Work or Derivative Works thereof in any medium, with or without
               modifications, and in Source or Object form, provided that You
               meet the following conditions:
         
               (a) You must give any other recipients of the Work or
                   Derivative Works a copy of this License; and
         
               (b) You must cause any modified files to carry prominent notices
                   stating that You changed the files; and
         
               (c) You must retain, in the Source form of any Derivative Works
                   that You distribute, all copyright, patent, trademark, and
                   attribution notices from the Source form of the Work,
                   excluding those notices that do not pertain to any part of
                   the Derivative Works; and
         
               (d) If the Work includes a "NOTICE" text file as part of its
                   distribution, then any Derivative Works that You distribute must
                   include a readable copy of the attribution notices contained
                   within such NOTICE file, excluding those notices that do not
                   pertain to any part of the Derivative Works, in at least one
                   of the following places: within a NOTICE text file distributed
                   as part of the Derivative Works; within the Source form or
                   documentation, if provided along with the Derivative Works; or,
                   within a display generated by the Derivative Works, if and
                   wherever such third-party notices normally appear. The contents
                   of the NOTICE file are for informational purposes only and
                   do not modify the License. You may add Your own attribution
                   notices within Derivative Works that You distribute, alongside
                   or as an addendum to the NOTICE text from the Work, provided
                   that such additional attribution notices cannot be construed
                   as modifying the License.
         
               You may add Your own copyright statement to Your modifications and
               may provide additional or different license terms and conditions
               for use, reproduction, or distribution of Your modifications, or
               for any such Derivative Works as a whole, provided Your use,
               reproduction, and distribution of the Work otherwise complies with
               the conditions stated in this License.
         
            5. Submission of Contributions. Unless You explicitly state otherwise,
               any Contribution intentionally submitted for inclusion in the Work
               by You to the Licensor shall be under the terms and conditions of
               this License, without any additional terms or conditions.
               Notwithstanding the above, nothing herein shall supersede or modify
               the terms of any separate license agreement you may have executed
               with Licensor regarding such Contributions.
         
            6. Trademarks. This License does not grant permission to use the trade
               names, trademarks, service marks, or product names of the Licensor,
               except as required for reasonable and customary use in describing the
               origin of the Work and reproducing the content of the NOTICE file.
         
            7. Disclaimer of Warranty. Unless required by applicable law or
               agreed to in writing, Licensor provides the Work (and each
               Contributor provides its Contributions) on an "AS IS" BASIS,
               WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
               implied, including, without limitation, any warranties or conditions
               of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
               PARTICULAR PURPOSE. You are solely responsible for determining the
               appropriateness of using or redistributing the Work and assume any
               risks associated with Your exercise of permissions under this License.
         
            8. Limitation of Liability. In no event and under no legal theory,
               whether in tort (including negligence), contract, or otherwise,
               unless required by applicable law (such as deliberate and grossly
               negligent acts) or agreed to in writing, shall any Contributor be
               liable to You for damages, including any direct, indirect, special,
               incidental, or consequential damages of any character arising as a
               result of this License or out of the use or inability to use the
               Work (including but not limited to damages for loss of goodwill,
               work stoppage, computer failure or malfunction, or any and all
               other commercial damages or losses), even if such Contributor
               has been advised of the possibility of such damages.
         
            9. Accepting Warranty or Additional Liability. While redistributing
               the Work or Derivative Works thereof, You may choose to offer,
               and charge a fee for, acceptance of support, warranty, indemnity,
               or other liability obligations and/or rights consistent with this
               License. However, in accepting such obligations, You may act only
               on Your own behalf and on Your sole responsibility, not on behalf
               of any other Contributor, and only if You agree to indemnify,
               defend, and hold each Contributor harmless for any liability
               incurred by, or claims asserted against, such Contributor by reason
               of your accepting any such warranty or additional liability.
         
            END OF TERMS AND CONDITIONS
         
            APPENDIX: How to apply the Apache License to your work.
         
               To apply the Apache License to your work, attach the following
               boilerplate notice, with the fields enclosed by brackets "[]"
               replaced with your own identifying information. (Don't include
               the brackets!)  The text should be enclosed in the appropriate
               comment syntax for the file format. We also recommend that a
               file or class name and description of purpose be included on the
               same "printed page" as the copyright notice for easier
               identification within third-party archives.
Classifier: Programming Language :: Python :: 3
Classifier: Environment :: GPU :: NVIDIA CUDA
Project-URL: Homepage, https://github.com/hao-ai-lab/FastVideo
Requires-Python: >=3.10
Requires-Dist: torch==2.12.0
Requires-Dist: triton>=2.0.0; sys_platform == "linux"
Description-Content-Type: text/markdown

# FastVideo Kernel

CUDA kernels for FastVideo video generation.

## Kernel inventory

Compiled CUDA extensions (CMake, see the build summary printed at the end of every configure):

| Extension | Kernels | Sources | GPU arch | Build gate |
|---|---|---|---|---|
| `fastvideo_kernel._C.fastvideo_kernel_ops` | TurboDiffusion INT8 GEMM, quant, RMSNorm, LayerNorm | `csrc/turbodiffusion/` | every arch in `TORCH_CUDA_ARCH_LIST` | always built |
| same extension, optional part | ThunderKittens sliding-tile attention (`sta_fwd`) and VSA block-sparse (`block_sparse_fwd/bwd`) | `csrc/attention/*_h100.cu` | Hopper `sm_90a` only | `FASTVIDEO_KERNEL_BUILD_TK` (AUTO = ON iff `9.0a` is in the arch list; always OFF on aarch64 hosts — TK headers don't compile there) |
| same extension, optional part | MiniMax-H3 block-sparse VSA forward (64- and 128-token blocks) | `csrc/attention/block_sparse*_sm100a.cu` | Blackwell `sm_100a` only | ON iff `10.0a` is in `TORCH_CUDA_ARCH_LIST` |
| same extension, optional part | fused NVLink Ulysses all-to-all | `csrc/comm/ulysses_all_to_all.cu` | CUDA | `FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A` (AUTO = ON with NCCL 2.29+ device headers and library; always OFF on ROCm) |
| `fp4attn_cuda`, `fp4quant_cuda` | FP4 attention + quantization ("attn_qat_infer", modified SageAttention3) | `attn_qat_infer/` | consumer Blackwell `sm_120a` only, CUDA ≥ 12.8 | `FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER` (AUTO = ON iff `12.0a` is in the arch list) |

Runtime-JIT kernels (no build step, ship in every wheel/image):

| Kernels | Where | Used when |
|---|---|---|
| Triton: STA, VSA block-sparse, SLA, fused compress+topk, FP4 QAT training, quant/norm utils | `python/fastvideo_kernel/triton_kernels/` | automatic fallback when the matching C++ op is absent (`ops.py`, `turbodiffusion_ops.py`) |
| FA4 CuTe-DSL block-sparse forward/backward (VSA-128/256 fastpath on `sm_100`) | `block_sparse_attn_cute_fwd.py` | optional `flash_attn.cute` dependency, see below |
| VMoBA `moba_attn_varlen` | `vmoba.py` | wraps flash-attn varlen |

## What gets built where, and when

| Surface | Trigger | Leg | `TORCH_CUDA_ARCH_LIST` | TK | Ulysses | FP4 |
|---|---|---|---|---|---|---|
| PyPI wheels (`.github/workflows/publish-kernel.yml`) | version bump in `fastvideo-kernel/pyproject.toml` on main, or manual dispatch | x86_64 cu126 | `9.0a` | ON | AUTO | — (CUDA < 12.8) |
| | | x86_64 cu130 | `9.0a;10.0a;12.0a` | ON | AUTO | ON |
| | | aarch64 cu130 | `10.0a;12.0a` | — | AUTO | ON |
| Docker images `ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev` (`.github/workflows/infra-build-image.yml`) | `docker/Dockerfile` changes on main, or manual dispatch | amd64 cuda12.6.3 + cuda13.0.0 | `9.0a` | ON | AUTO | — |
| | | arm64 cuda12.6.3 (GH200) | `9.0a` | — (aarch64) | AUTO | — |
| | | arm64 cuda13.0.0 (GB10 / DGX Spark) | `12.1` | — | AUTO | — |
| Local `./build.sh` | manual | probes the visible GPU via torch | detected | ON iff sm_90 (non-aarch64 host) | AUTO | ON iff sm_120 |

Notes:

- No Docker image ships the FP4 kernels; only the x86_64/aarch64 cu130 wheels do.
- Both cu130 PyPI wheels ship the MiniMax-H3 sm_100a VSA forward.
- On arm64 images (GH200 included) STA/VSA run on the Triton fallbacks, since TK never builds on aarch64.
- Ulysses AUTO builds only when CMake finds a NCCL library and device-API headers with the 2.29 initializers. Use
  `-DFASTVIDEO_KERNEL_BUILD_ULYSSES_A2A=ON` to require it or `OFF` to test the portable build.
- Kernel tests run on Buildkite GPU CI for PRs touching `fastvideo-kernel/**` (see `.buildkite/pipeline.yml`).

## Installation

### Standard Installation (Local Development)
This will automatically detect your GPU architecture. If an NVIDIA Hopper (H100/sm_90a) GPU is detected, ThunderKittens kernels will be enabled. Otherwise, they will be skipped, and the package will use Triton fallbacks at runtime.

Before installation, set CUDA toolchain paths:

```bash
export CUDA_HOME=/usr/local/cuda
export CUDACXX=$CUDA_HOME/bin/nvcc
```

```bash
git submodule update --init --recursive
cd fastvideo-kernel
./build.sh
```

### Rocm Build
If you are in a rocm environment without the compilation toolchaine of CUDA.

```bash
cd fastvideo-kernel
./build.sh --rocm
```

### Optional: FA4 CuTe block-sparse backend (VSA-128/256 fastpath)

The VSA-128/256 fastpaths (tile volume 128 or 256, on NVIDIA Blackwell / sm_100) route to the
FlashAttention-4 CuTe-DSL block-sparse kernel exposed as `flash_attn.cute`. This is
an **optional** dependency: it is imported lazily, and `video_sparse_attn`
transparently falls back to the Triton backend when it is absent (so the package is
fully usable without it).

The symbols the fastpath needs (`flash_attn.cute.block_sparsity.BlockSparseTensorsTorch`
and the public/private forward-backward bridges in `flash_attn.cute.interface`) are provided upstream by
[Dao-AILab/flash-attention](https://github.com/Dao-AILab/flash-attention). Pin to
commit `14c377950125c70b7a9dabf9c561fca53715ac7d`, the revision FastVideo pins as
the `flash-attn-4` source in the repo-root `pyproject.toml`; other revisions may
have incompatible block-sparse forward/backward interfaces.

Install it under its distribution name so its own runtime stack resolves with it.
Do **not** pre-install `nvidia-cutlass-dsl` by hand: this revision pins
`nvidia-cutlass-dsl==4.6.0.dev0` exactly, and a hand-installed 4.5.x floor either
gets silently upgraded or, if something else holds it back, leaves the CuTe
kernels broken.

```bash
pip install torchvision
pip install "flash-attn-4 @ git+https://github.com/Dao-AILab/flash-attention.git@14c377950125c70b7a9dabf9c561fca53715ac7d#subdirectory=flash_attn/cute"
```

That resolves `nvidia-cutlass-dsl` to 4.6.0.dev0 and `quack-kernels` to 0.5.3, a
combination this revision works with. A mismatched CuTe DSL only surfaces when the
kernel JIT-compiles, so the error points at CuTe internals rather than at the
install:

| Error on first VSA-128/256 CuTe call | Cause |
|---|---|
| `TypeError: fmax() missing 1 required positional argument: 'b'` | `nvidia-cutlass-dsl` 4.5.x |
| `AttributeError: module 'cutlass.cute.core' has no attribute 'ThrMma'` | `quack-kernels` older than 0.5.1 |
| `ImportError: cannot import name 'alloc_reserved_mbarrier'` | `quack-kernels` 0.6.2 or newer |

An environment whose `flash_attn.cute` came from a prebuilt flash-attn wheel rather
than from this pin hits the first row; that is what the overlay step in
`docker/Dockerfile` works around.

The CuTe kernels JIT-compile on first use. Forward and backward are verified on
Blackwell (sm_100) against `tests/test_vsa128_*.py` and `tests/test_vsa256_*.py`.

## Usage

### Sliding Tile Attention (STA) & Video Sparse Attention (VSA)

For detailed usage, please check the [Attention Documentation](../docs/attention/index.md).

```python
from fastvideo_kernel import sliding_tile_attention, video_sparse_attn, moba_attn_varlen

# Example: Sliding Tile Attention
out = sliding_tile_attention(q, k, v, window_sizes, text_len)

# Example: Video Sparse Attention (with Triton fallback)
out = video_sparse_attn(q, k, v, block_sizes, block_sizes, topk=5)

# Example: VMoBA
out = moba_attn_varlen(q, k, v, cu_seqlens_q, cu_seqlens_k, ...)
```

## Benchmark

### Attn-QAT training

The default shape matches one sequence-parallel rank of the 4-GPU
Wan2.1-T2V-1.3B MixKit recipe (`B=1, H=3, L=31200, D=128`):

```bash
cd fastvideo-kernel
python benchmarks/benchmark_attn_qat_train.py
```

The benchmark reports both conventional attention FLOPs and the extra matrix
multiplications executed by the QAT straight-through path. Override
`--peak-tflops` when running on a GPU other than RTX 5090.

The QAT kernel is entirely Triton and routes by architecture at runtime. SM100
uses a large-tile forward and split 64x64 backward for the production
non-causal, head-dimension-128 configuration with a 16-aligned KV length. SM120
(including RTX 5090) keeps the previous tiling but joins the quantized and STE
P@V operations and uses a shallower backward pipeline for long sequences. Set
`FASTVIDEO_ATTN_QAT_SM120_JOIN_QAT_PV=0` to compare against the split P@V path.
Unsupported configurations retain the previous implementation. Set
`FASTVIDEO_ATTN_QAT_SM100_OPTIMIZED=0` to benchmark that previous path on SM100. Forward tuning is available through
`FASTVIDEO_ATTN_QAT_FWD_MODE=fast|balanced|reference`; exact reference-order
softmax statistics are controlled by `FASTVIDEO_ATTN_QAT_FWD_EXACT_M` and are
disabled by default for maximum throughput. Set it to `1` for reference-order
statistics and bitwise-compatible `dV`.

### VSA (block-sparse) TFLOPs

After building/installing `fastvideo-kernel`, run:

```bash
cd fastvideo-kernel
python benchmarks/bench_vsa.py --batch_size 1 --num_heads 16 --head_dim 128 --q_seq_lens 49152 --topk 64

# VSA-256 FA4 CuTe forward/backward on Blackwell
python benchmarks/bench_vsa.py --block_size 256 --use_cute \
  --batch_size 1 --num_heads 12 --head_dim 128 --q_seq_lens 39936 --topk 20

# VSA-128 FA4 CuTe forward/backward on Blackwell
python benchmarks/bench_vsa.py --block_size 128 --use_cute \
  --batch_size 1 --num_heads 12 --head_dim 128 --q_seq_lens 39936 --topk 40
```

### TurboDiffusion Kernels

This package also includes kernels from [TurboDiffusion](https://github.com/thu-ml/TurboDiffusion), including INT8 GEMM, Quantization, RMSNorm and LayerNorm.

## Requirements

- **Runtime**:
  - NVIDIA H100 (sm_90a) for C++ optimized kernels.
  - Any CUDA GPU for Triton-based fallbacks.
- **Build**:
  - CUDA Toolkit 12.3+
  - `CUDA_HOME` must be set (for example, `/usr/local/cuda`)
  - `CUDACXX` must be set (for example, `$CUDA_HOME/bin/nvcc`)
  - C++20 compatible compiler (GCC 10+, Clang 11+)

## Acknowledgement

This package structure and build system are based on [sgl-kernel](https://github.com/sgl-project/sglang/tree/main/sgl-kernel) from the SGLang project.

The implementation of `turbodiffusion` kernels is adapted from [TurboDiffusion](https://github.com/thu-ml/TurboDiffusion). If you use these kernels, please cite:

```bibtex
@article{zhang2025turbodiffusion,
  title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
  author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
  journal={arXiv preprint arXiv:2512.16093},
  year={2025}
}
```
