ckpt

[bugfix]: _compile_conditions regression (#1077 )
[misc] readme small fix (#1076 )
2026-02-09 05:50:48 +00:00 · 2026-02-06 18:57:43 -08:00 · 2026-02-06 16:57:48 -08:00 · 2026-02-06 19:43:18 -05:00 · 2026-02-06 13:38:35 -08:00 · 2026-02-05 20:33:08 -08:00
308 changed files with 37628 additions and 3195 deletions
@@ -61,7 +61,7 @@ steps:
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
-                command: "timeout 60m .buildkite/scripts/pr_test.sh"
+                command: "timeout 90m .buildkite/scripts/pr_test.sh"
                label: "SSIM Tests"
                env:
                  - TEST_TYPE=ssim
@@ -156,8 +156,23 @@ jobs:
          
          # Fix the wheel to be manylinux compliant
          pip install auditwheel
+          # Point auditwheel at torch libs, but do not vendor them into the wheel.
+          TORCH_LIB_DIR=$(python - <<'PY'
+          import os
+          import torch
+
+          print(os.path.join(os.path.dirname(torch.__file__), "lib"))
+          PY
+          )
+          export LD_LIBRARY_PATH="${TORCH_LIB_DIR}:${LD_LIBRARY_PATH}"
          # Target manylinux_2_35 (Ubuntu 22.04 native)
-          auditwheel repair dist/*.whl --plat manylinux_2_35_x86_64 -w fixed_dist
+          auditwheel repair dist/*.whl --plat manylinux_2_35_x86_64 -w fixed_dist \
+            --exclude libtorch_cuda.so \
+            --exclude libtorch_cpu.so \
+            --exclude libtorch.so \
+            --exclude libc10.so \
+            --exclude libc10_cuda.so \
+            --exclude libtorch_python.so
          # Move fixed wheels back to dist for upload consistency
          rm dist/*.whl
          mv fixed_dist/*.whl dist/
@@ -68,7 +68,7 @@ repos:
    entry: bash
    args:
      - -c
-      - 'git ls-files | grep -v "^fastvideo/tests/ssim/" | grep -v "^fastvideo/tests/inference/lora/L40S_reference_videos/" | grep " " && echo "Filenames should not contain spaces!" && exit 1 || exit 0'
+      - 'git ls-files | grep -v "^\"*fastvideo/tests/ssim/" | grep -v "^\"*fastvideo/tests/inference/lora/L40S_reference_videos/" | grep " " && echo "Filenames should not contain spaces!" && exit 1 || exit 0'
    language: system
    always_run: true
    pass_filenames: false
@@ -1,35 +1,29 @@
-<div align="center">
-<img src=assets/logos/logo.svg width="30%"/>
-</div>

-<p align="center">
-     | <a href="https://hao-ai-lab.github.io/FastVideo"><b>Documentation</b></a> | <a href="https://hao-ai-lab.github.io/FastVideo/inference/inference_quick_start/"><b> Quick Start</b></a> | <a href="https://github.com/hao-ai-lab/FastVideo/discussions/982"  target="_blank"><b>Weekly Dev Meeting</b></a>  | 🟣💬 <a href="https://join.slack.com/t/fastvideo/shared_invite/zt-3f4lao1uq-u~Ipx6Lt4J27AlD2y~IdLQ" target="_blank"> <b>Slack</b> </a> |  🟣💬 <a href="https://ibb.co/c7g1qdD" target="_blank"> <b> WeChat </b> </a> |
-</p>
+| **[Documentation](https://hao-ai-lab.github.io/FastVideo)** | **[Quick Start](https://hao-ai-lab.github.io/FastVideo/inference/inference_quick_start/)** | **[Weekly Dev Meeting](https://github.com/hao-ai-lab/FastVideo/discussions/982)** | 🟣💬 **[Slack](https://join.slack.com/t/fastvideo/shared_invite/zt-3f4lao1uq-u~Ipx6Lt4J27AlD2y~IdLQ)** |

 **FastVideo is a unified post-training and inference framework for accelerated video generation.**

 ## NEWS
- ```2025/11/19```: Release [CausalWan2.2 I2V A14B Preview](https://huggingface.co/FastVideo/CausalWan2.2-I2V-A14B-Preview-Diffusers) models, [Blog](https://hao-ai-lab.github.io/blogs/fastvideo_causalwan_preview/) and [Inference Code!](https://github.com/hao-ai-lab/FastVideo/blob/main/examples/inference/basic/basic_self_forcing_causal_wan2_2_i2v.py)
- ```2025/08/04```: Release [FastWan](https://hao-ai-lab.github.io/FastVideo/distillation/dmd) models and [Sparse-Distillation](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/).

-<details>
-<summary>More</summary>
+- `2025/11/19`: Release [CausalWan2.2 I2V A14B Preview](https://huggingface.co/FastVideo/CausalWan2.2-I2V-A14B-Preview-Diffusers) models, [Blog](https://hao-ai-lab.github.io/blogs/fastvideo_causalwan_preview/) and [Inference Code!](https://github.com/hao-ai-lab/FastVideo/blob/main/examples/inference/basic/basic_self_forcing_causal_wan2_2_i2v.py)
+- `2025/08/04`: Release [FastWan](https://hao-ai-lab.github.io/FastVideo/distillation/dmd) models and [Sparse-Distillation](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/).

- ```2025/06/14```: Release finetuning and inference code for [VSA](https://arxiv.org/pdf/2505.13389)
- ```2025/04/24```: [FastVideo V1](https://hao-ai-lab.github.io/blogs/fastvideo/) is released!
- ```2025/02/18```: Release the inference code for [Sliding Tile Attention](https://hao-ai-lab.github.io/blogs/sta/).
+### More News

-</details>
+- `2025/06/14`: Release finetuning and inference code for [VSA](https://arxiv.org/pdf/2505.13389)
+- `2025/04/24`: [FastVideo V1](https://hao-ai-lab.github.io/blogs/fastvideo/) is released!
+- `2025/02/18`: Release the inference code for [Sliding Tile Attention](https://hao-ai-lab.github.io/blogs/sta/).

 ## Key Features

 FastVideo has the following features:
+
 - End-to-end post-training support for bidirectional and autoregressive models:
  - Support full finetuning and LoRA finetuning for state-of-the-art open video DiTs
  - Data preprocessing pipeline for video, image, and text data
  - Distribution Matching Distillation (DMD2) stepwise distillation.
  - Sparse attention with [Video Sparse Attention](https://arxiv.org/pdf/2505.13389)
-  - [Sparse distillation](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/) to achineve >50x denoising speedup
+  - [Sparse distillation](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/) to achieve >50x denoising speedup
  - Scalable training with FSDP2, sequence parallelism, and selective activation checkpointing.
  - Causal distillation through Self-Forcing
  - See this [page](https://hao-ai-lab.github.io/FastVideo/training/overview/) for full list of supported models and recipes.
@@ -44,6 +38,7 @@ FastVideo has the following features:
  - See this [page](https://hao-ai-lab.github.io/FastVideo/inference/hardware_support/) for full list of supported hardware and OS.

 ## Getting Started
+
 We recommend using an environment manager such as `Conda` to create a clean environment:

 ```bash
@@ -58,18 +53,20 @@ pip install fastvideo
 Please see our [docs](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/) for more detailed installation instructions.

 ## Sparse Distillation
+
 For our sparse distillation techniques, please see our [distillation docs](https://hao-ai-lab.github.io/FastVideo/distillation/dmd/) and check out our [blog](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/).

 See below for recipes and datasets:

-|                                            Model                                              |                                               Sparse Distillation                                                 |                                                  Dataset                                                  |
-|:-------------------------------------------------------------------------------------------:  |:---------------------------------------------------------------------------------------------------------------:  |:--------------------------------------------------------------------------------------------------------: |
-| [FastWan2.1-T2V-1.3B](https://huggingface.co/FastVideo/FastWan2.1-T2V-1.3B-Diffusers)         |    [Recipe](https://github.com/hao-ai-lab/FastVideo/tree/main/examples/distill/Wan2.1-T2V/Wan-Syn-Data-480P)      | [FastVideo Synthetic Wan2.1 480P](https://huggingface.co/datasets/FastVideo/Wan-Syn_77x448x832_600k)      |
-| [FastWan2.1-T2V-14B-Preview](https://huggingface.co/FastVideo/FastWan2.1-T2V-14B-Diffusers)   |                                                   Coming soon!                                                    |   [FastVideo Synthetic Wan2.1 720P](https://huggingface.co/datasets/FastVideo/Wan-Syn_77x768x1280_250k)   |
-| [FastWan2.2-TI2V-5B](https://huggingface.co/FastVideo/FastWan2.2-TI2V-5B-Diffusers)           | [Recipe](https://github.com/hao-ai-lab/FastVideo/tree/main/examples/distill/Wan2.2-TI2V-5B-Diffusers/Data-free)   | [FastVideo Synthetic Wan2.2 720P](https://huggingface.co/datasets/FastVideo/Wan2.2-Syn-121x704x1280_32k)  |
+| Model                                                                                 | Sparse Distillation                                                                                             | Dataset                                                                                                  |
+| ------------------------------------------------------------------------------------- | --------------------------------------------------------------------------------------------------------------- | -------------------------------------------------------------------------------------------------------- |
+| [FastWan2.1-T2V-1.3B](https://huggingface.co/FastVideo/FastWan2.1-T2V-1.3B-Diffusers) | [Recipe](https://github.com/hao-ai-lab/FastVideo/tree/main/examples/distill/Wan2.1-T2V/Wan-Syn-Data-480P)       | [FastVideo Synthetic Wan2.1 480P](https://huggingface.co/datasets/FastVideo/Wan-Syn_77x448x832_600k)     |
+| [FastWan2.2-TI2V-5B](https://huggingface.co/FastVideo/FastWan2.2-TI2V-5B-Diffusers)   | [Recipe](https://github.com/hao-ai-lab/FastVideo/tree/main/examples/distill/Wan2.2-TI2V-5B-Diffusers/Data-free) | [FastVideo Synthetic Wan2.2 720P](https://huggingface.co/datasets/FastVideo/Wan2.2-Syn-121x704x1280_32k) |

 ## Inference
+
 ### Generating Your First Video
+
 Here's a minimal example to generate a video using the default settings. Make sure VSA kernels are [installed](https://hao-ai-lab.github.io/FastVideo/video_sparse_attention/installation/). Create a file called `example.py` with the following code:

 ```python
@@ -108,55 +105,36 @@ python example.py

 For a more detailed guide, please see our [inference quick start](https://hao-ai-lab.github.io/FastVideo/inference/inference_quick_start/).

-### Other docs:
+## More Guides

 - [Design Overview](https://hao-ai-lab.github.io/FastVideo/design/overview/)
- [Contribution Guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/)
-
-## Distillation and Finetuning
 - [Distillation Guide](https://hao-ai-lab.github.io/FastVideo/distillation/dmd/)
-<!-- - [Finetuning Guide](https://hao-ai-lab.github.io/FastVideo/training/finetune.html) -->
+- [Contribution Guide](https://hao-ai-lab.github.io/FastVideo/contributing/overview/)

 ## Awesome work using FastVideo or our research projects

- [SGLang](https://github.com/sgl-project/sglang/tree/main/python/sglang/multimodal_gen): SGLang's diffusion inference functionality is based  on a fork of FastVideo on Sept. 24, 2025. [![Star](https://img.shields.io/github/stars/sgl-project/sglang.svg?style=social&label=Star)](https://github.com/sgl-project/sglang)
-
- [DanceGRPO](https://github.com/XueZeyue/DanceGRPO): A  unified framework to adapt Group Relative Policy Optimization (GRPO) to visual generation paradigms. Code based on FastVideo. [![Star](https://img.shields.io/github/stars/XueZeyue/DanceGRPO.svg?style=social&label=Star)](https://github.com/XueZeyue/DanceGRPO)
- [SRPO](https://github.com/Tencent-Hunyuan/SRPO): A method to directly align the full diffusion trajectory with fine-grained human preference. Code based on FastVideo. [![Star](https://img.shields.io/github/stars/Tencent-Hunyuan/SRPO.svg?style=social&label=Star)](https://github.com/Tencent-Hunyuan/SRPO)
- [DCM](https://github.com/Vchitect/DCM): Dual-expert consistency model for efficient and high-quality video generation. Code based on FastVideo. [![Star](https://img.shields.io/github/stars/Vchitect/DCM.svg?style=social&label=Star)](https://github.com/Vchitect/DCM)
- [Hunyuan Video 1.5](https://github.com/Tencent-Hunyuan/HunyuanVideo-1.5): A leading lightweight video generation model, where they proposed SSTA based on Sliding Tile Attention. [![Star](https://img.shields.io/github/stars/Tencent-Hunyuan/HunyuanVideo-1.5.svg?style=social&label=Star)](https://github.com/Tencent-Hunyuan/HunyuanVideo-1.5)
- [Kandinsky-5.0](https://github.com/kandinskylab/kandinsky-5): A family of diffusion models for video & image generation, where their NABLA attention includes a Sliding Tile Attention branch. [![Star](https://img.shields.io/github/stars/kandinskylab/kandinsky-5.svg?style=social&label=Star)](https://github.com/kandinskylab/kandinsky-5)
- [LongCat Video](https://github.com/meituan-longcat/LongCat-Video): A foundational video generation model with 13.6B parameters with block-sparse attention similar to Video Sparse Attention. [![Star](https://img.shields.io/github/stars/meituan-longcat/LongCat-Video.svg?style=social&label=Star)](https://github.com/meituan-longcat/LongCat-Video)
+- [SGLang](https://github.com/sgl-project/sglang/tree/main/python/sglang/multimodal_gen): SGLang's diffusion inference functionality is based on a fork of FastVideo on Sept. 24, 2025.
+- [DanceGRPO](https://github.com/XueZeyue/DanceGRPO): A unified framework to adapt Group Relative Policy Optimization (GRPO) to visual generation paradigms. Code based on FastVideo.
+- [SRPO](https://github.com/Tencent-Hunyuan/SRPO): A method to directly align the full diffusion trajectory with fine-grained human preference. Code based on FastVideo.
+- [DCM](https://github.com/Vchitect/DCM): Dual-expert consistency model for efficient and high-quality video generation. Code based on FastVideo.
+- [Hunyuan Video 1.5](https://github.com/Tencent-Hunyuan/HunyuanVideo-1.5): A leading lightweight video generation model, where they proposed SSTA based on Sliding Tile Attention.
+- [Kandinsky-5.0](https://github.com/kandinskylab/kandinsky-5): A family of diffusion models for video & image generation, where their NABLA attention includes a Sliding Tile Attention branch.
+- [LongCat Video](https://github.com/meituan-longcat/LongCat-Video): A foundational video generation model with 13.6B parameters with block-sparse attention similar to Video Sparse Attention.

 ## 🤝 Contributing

 We welcome all contributions. Please check out our guide [here](https://hao-ai-lab.github.io/FastVideo/contributing/overview/).
 See details in [development roadmap](https://github.com/hao-ai-lab/FastVideo/issues/899).
-## Acknowledgement
-We learned and reused code from the following projects:
- [Wan-Video](https://github.com/Wan-Video)
- [ThunderKittens](https://github.com/HazyResearch/ThunderKittens)
- [Triton](https://github.com/triton-lang/triton)
- [DMD2](https://github.com/tianweiy/DMD2)
- [diffusers](https://github.com/huggingface/diffusers)
- [xDiT](https://github.com/xdit-project/xDiT)
- [vLLM](https://github.com/vllm-project/vllm)
- [SGLang](https://github.com/sgl-project/sglang)

-We thank [MBZUAI](https://ifm.mbzuai.ac.ae/), [Anyscale](https://www.anyscale.com/), and [GMI Cloud](https://www.gmicloud.ai/) for their support throughout this project.
+## Acknowledgement
+
+We learned the design and reused code from the following projects: [Wan-Video](https://github.com/Wan-Video), [ThunderKittens](https://github.com/HazyResearch/ThunderKittens), [DMD2](https://github.com/tianweiy/DMD2), [diffusers](https://github.com/huggingface/diffusers), [xDiT](https://github.com/xdit-project/xDiT), [vLLM](https://github.com/vllm-project/vllm), [SGLang](https://github.com/sgl-project/sglang). We thank [MBZUAI](https://ifm.mbzuai.ac.ae/), [Anyscale](https://www.anyscale.com/), and [GMI Cloud](https://www.gmicloud.ai/) for their support throughout this project.

 ## Citation
-If you find FastVideo useful, please considering citing our work:
+
+If you find FastVideo useful, please consider citing our research work:

 ```bibtex
-@software{fastvideo2024,
-  title        = {FastVideo: A Unified Framework for Accelerated Video Generation},
-  author       = {The FastVideo Team},
-  url          = {https://github.com/hao-ai-lab/FastVideo},
-  month        = apr,
-  year         = {2024},
-}
-
@article{zhang2025vsa,
  title={Vsa: Faster video diffusion with trainable sparse attention},
  author={Zhang, Peiyuan and Chen, Yongqi and Huang, Haofeng and Lin, Will and Liu, Zhengzhong and Stoica, Ion and Xing, Eric and Zhang, Hao},
@@ -43,7 +43,7 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    uv pip install --no-cache-dir --upgrade pip && \
    uv pip install --no-cache-dir .[dev] && \
-    uv pip install --no-cache-dir https://github.com/mjun0812/flash-attention-prebuild-wheels/releases/download/v0.5.4/flash_attn-2.8.3%2Bcu128torch2.9-cp310-cp310-linux_x86_64.whl
+    uv pip install --no-cache-dir https://github.com/mjun0812/flash-attention-prebuild-wheels/releases/download/v0.7.16/flash_attn-2.8.3+cu128torch2.10-cp310-cp310-linux_x86_64.whl

 COPY . .

@@ -43,7 +43,7 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    uv pip install --no-cache-dir --upgrade pip && \
    uv pip install --no-cache-dir .[dev] && \
-    uv pip install --no-cache-dir https://github.com/mjun0812/flash-attention-prebuild-wheels/releases/download/v0.5.4/flash_attn-2.8.3%2Bcu128torch2.9-cp311-cp311-linux_x86_64.whl
+    uv pip install --no-cache-dir https://github.com/mjun0812/flash-attention-prebuild-wheels/releases/download/v0.7.16/flash_attn-2.8.3+cu128torch2.10-cp311-cp311-linux_x86_64.whl

 COPY . .

@@ -43,7 +43,7 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    uv pip install --no-cache-dir --upgrade pip && \
    uv pip install --no-cache-dir .[dev] && \
-    uv pip install --no-cache-dir https://github.com/mjun0812/flash-attention-prebuild-wheels/releases/download/v0.5.4/flash_attn-2.8.3%2Bcu128torch2.9-cp312-cp312-linux_x86_64.whl
+    uv pip install --no-cache-dir https://github.com/mjun0812/flash-attention-prebuild-wheels/releases/download/v0.7.16/flash_attn-2.8.3+cu128torch2.10-cp312-cp312-linux_x86_64.whl

 COPY . .

@@ -121,6 +121,15 @@ This page contains the complete API reference for the FastVideo library.
      show_root_toc_entry: true
      heading_level: 4

+## fastvideo.registry
+
+::: fastvideo.registry
+    options:
+      show_source: true
+      show_root_heading: true
+      show_root_toc_entry: true
+      heading_level: 3
+
 ## fastvideo.pipelines

 ::: fastvideo.pipelines
@@ -6,6 +6,8 @@ FastVideo provides highly optimized custom attention kernels to accelerate video

 * **[Video Sparse Attention (VSA)](vsa/index.md)**: Sparse attention mechanism selecting top-k blocks.
 * **[Sliding Tile Attention (STA)](sta/index.md)**: Optimized attention for window-based video generation.
+* **Backend development guide**: See the developer guide at
+  [Attention Backend Development](../contributing/attention_backend.md).

 ## General Build Instructions

@@ -41,7 +43,7 @@ Clone the repository and build the kernel:

 ```bash
 # Clone recursively to get ThunderKittens submodule
-git clone --recursive https://github.com/hao-ai-lab/FastVideo.git
+git clone https://github.com/hao-ai-lab/FastVideo.git
 cd FastVideo/fastvideo-kernel

 # Build and install
@@ -13,11 +13,13 @@ from fastvideo_kernel import video_sparse_attn

 # q, k, v: [batch_size, num_heads, seq_len, head_dim]
 # variable_block_sizes: Number of valid tokens per block
+# q_variable_block_sizes: Number of valid tokens per q block (can differ from KV for q/k of different lengths)
 # topk: Number of blocks to attend

 output = video_sparse_attn(
    q, k, v, 
-    variable_block_sizes=block_sizes, 
+    block_sizes,
+    block_sizes,
    topk=32
 )
 ```
@@ -0,0 +1,176 @@
+# Attention Backend Development
+
+This guide is for contributors adding a new attention backend (kernel or
+implementation) to FastVideo. If you just want to use existing kernels or build
+fastvideo-kernel, see [Attention overview](../attention/index.md).
+
+## When you need this guide
+
+Use this guide when you are:
+
+- Adding a new attention kernel or algorithm.
+- Wiring an existing kernel into FastVideo's attention selection.
+- Extending attention support to a new platform.
+
+## 0) Choose a backend name and scope
+
+Pick a backend name in `UPPER_SNAKE_CASE` and decide where it should run.
+Example: `MY_NEW_ATTN`.
+
+You will use this name in:
+
+- `AttentionBackendEnum` (global list of backends).
+- `get_name()` in your backend class (match the enum name).
+- Platform selectors (CUDA/ROCm/MPS/NPU) to return your backend class.
+
+## 1) Add enum + platform selection
+
+1) Add your backend to `fastvideo/platforms/interface.py`:
+
+```python
+class AttentionBackendEnum(enum.Enum):
+    ...
+    MY_NEW_ATTN = enum.auto()
+```
+
+1) Register it in platform selection (example: CUDA). Update
+`fastvideo/platforms/cuda.py` inside `get_attn_backend_cls`:
+
+```python
+elif selected_backend == AttentionBackendEnum.MY_NEW_ATTN:
+    try:
+        from fastvideo.attention.backends.my_new_attn import MyNewAttnBackend
+        return "fastvideo.attention.backends.my_new_attn.MyNewAttnBackend"
+    except ImportError as e:
+        logger.error("Failed to import MY_NEW_ATTN backend: %s", str(e))
+        raise
+```
+
+If you want support on other platforms, add a similar branch in
+`fastvideo/platforms/rocm.py`, `fastvideo/platforms/mps.py`, or `fastvideo/platforms/npu.py`.
+
+## 2) Implement the backend
+
+Create `fastvideo/attention/backends/my_new_attn.py` and implement the required
+classes.
+
+Minimal skeleton (no custom metadata):
+
+```python
+import torch
+from dataclasses import dataclass
+from fastvideo.attention.backends.abstract import (
+    AttentionBackend,
+    AttentionImpl,
+    AttentionMetadata,
+    AttentionMetadataBuilder,
+)
+
+class MyNewAttnBackend(AttentionBackend):
+    @staticmethod
+    def get_name() -> str:
+        return "MY_NEW_ATTN"
+
+    @staticmethod
+    def get_impl_cls() -> type["MyNewAttnImpl"]:
+        return MyNewAttnImpl
+
+    @staticmethod
+    def get_metadata_cls() -> type["AttentionMetadata"]:
+        return AttentionMetadata
+
+    @staticmethod
+    def get_builder_cls() -> type["AttentionMetadataBuilder"]:
+        return MyNewAttnMetadataBuilder
+
+
+@dataclass
+class MyNewAttnMetadata(AttentionMetadata):
+    current_timestep: int
+
+
+class MyNewAttnMetadataBuilder(AttentionMetadataBuilder):
+    def __init__(self) -> None:
+        pass
+
+    def prepare(self) -> None:
+        pass
+
+    def build(self, current_timestep: int, **kwargs):
+        return MyNewAttnMetadata(current_timestep=current_timestep)
+
+
+class MyNewAttnImpl(AttentionImpl):
+    def __init__(
+        self,
+        num_heads: int,
+        head_size: int,
+        softmax_scale: float,
+        causal: bool = False,
+        num_kv_heads: int | None = None,
+        prefix: str = "",
+        **extra_impl_args,
+    ) -> None:
+        self.softmax_scale = softmax_scale
+        self.causal = causal
+
+    def forward(
+        self,
+        query: torch.Tensor,
+        key: torch.Tensor,
+        value: torch.Tensor,
+        attn_metadata: MyNewAttnMetadata,
+    ) -> torch.Tensor:
+        # Implement attention
+        return torch.nn.functional.scaled_dot_product_attention(
+            query.transpose(1, 2),
+            key.transpose(1, 2),
+            value.transpose(1, 2),
+            is_causal=self.causal,
+            scale=self.softmax_scale,
+        ).transpose(1, 2)
+```
+
+Optional:
+
+- Implement `preprocess_qkv` / `postprocess_output` if your kernel needs tiling
+  or reshaping.
+- Use `fastvideo.forward_context.get_forward_context()` if you need dynamic
+  per-step data (e.g., window sizes).
+- Set `accept_output_buffer = True` if your backend writes into a provided
+  output buffer.
+
+## 3) Wire into attention layers
+
+Backends are used by `LocalAttention` and `DistributedAttention`. These layers
+accept a `supported_attention_backends` tuple. If your backend should be
+eligible, update the call sites that construct these layers (search for
+`supported_attention_backends=`).
+
+## 4) Add compiled kernels (optional)
+
+If you have a custom CUDA kernel:
+
+1) Add sources in `fastvideo-kernel/csrc/attention/`.
+2) Register bindings in `fastvideo-kernel/csrc/common_extension.cpp`.
+3) Add to `fastvideo-kernel/CMakeLists.txt` (and any feature flags).
+4) Expose in `fastvideo-kernel/python/fastvideo_kernel/ops.py`.
+5) Export in `fastvideo-kernel/python/fastvideo_kernel/__init__.py`.
+
+Keep a Python/Triton fallback so the backend runs even when the kernel is not
+available.
+
+## 5) Testing and debugging
+
+- Add a small parity test or microbenchmark comparing to SDPA.
+- Force your backend with the env var:
+  `FASTVIDEO_ATTENTION_BACKEND=MY_NEW_ATTN`.
+- Check logs from `fastvideo/attention/selector.py` to confirm selection.
+
+## Checklist
+
+- [ ] Added enum entry in `fastvideo/platforms/interface.py`.
+- [ ] Implemented backend in `fastvideo/attention/backends/`.
+- [ ] Registered selection in platform(s).
+- [ ] Updated layer call sites to include the backend where appropriate.
+- [ ] Added tests and documentation.
@@ -0,0 +1,480 @@
+# FastVideo + Coding Agents
+
+Coding agents are now strong at navigating large codebases and iterating fast
+with parity tests and examples. This guide shows how to use them to add new
+model pipelines and ship PRs in a production-grade video diffusion framework.
+
+FastVideo is a great project to contribute to, with production-grade
+infrastructure, active collaborations (including NVIDIA), and a pipeline design
+and inference architecture that has been forked by [SGLang’s
+multimodal generation stack](https://github.com/sgl-project/sglang/tree/main/python/sglang/multimodal_gen).
+
+Goal: run the new pipeline with a minimal script like
+`examples/inference/basic/basic.py`. In production, FastVideo can download
+models automatically via `HF_HOME`; for development, use local directories so
+agents can run scripts and tests deterministically. We standardize local paths
+as:
+
+- `official_weights/<model_name>/` for official checkpoints
+- `converted_weights/<model_name>/` if conversion is required
+
+## Tips when prompting the agent
+
+When prompting the agent, include:
+
+- This guide and the [FastVideo design overview](../design/overview.md).
+- Exact file paths to edit.
+- A closest reference example file in FastVideo.
+- Expected behavior and acceptance criteria.
+- Repro steps (command, inputs, logs).
+- Constraints (performance, memory, compatibility).
+- Local paths (e.g., `official_weights/<model_name>/` or
+  `converted_weights/<model_name>/`) for parity tests.
+
+## FastVideo structure at a glance
+
+Before diving in, scan these references:
+
+- [Contributing overview](overview.md) for environment/setup context.
+- [FastVideo design overview](../design/overview.md) for pipeline architecture, configs, and HF layout.
+
+FastVideo maps a Diffusers-style repo into a pipeline like:
+
+- `fastvideo/models/*`: model implementations (DiT, VAE, encoders, upsamplers).
+- `fastvideo/configs/models/*`: arch configs and `param_names_mapping` for
+  weight name translation.
+- `fastvideo/configs/pipelines/*`: pipeline wiring (component classes + names).
+- `fastvideo/configs/sample/*`: default runtime sampling parameters.
+- `fastvideo/pipelines/basic/*`: end-to-end pipeline logic built from stages.
+- `model_index.json`: the HF repo entrypoint that maps component names to
+  classes and weight files.
+- Component loading happens in `VideoGenerator.from_pretrained`, which reads
+  `model_index.json`, resolves configs, and loads weights.
+
+Minimal usage example (based on `examples/inference/basic/basic.py`):
+
+```python
+from fastvideo import VideoGenerator
+from fastvideo.configs.sample import SamplingParam
+
+model_id = "Wan-AI/Wan2.1-T2V-1.3B-Diffusers"  # or official_weights/<model_name>/
+generator = VideoGenerator.from_pretrained(model_id, num_gpus=1)
+
+sampling = SamplingParam.from_pretrained(model_id)
+sampling.num_frames = 45
+video = generator.generate_video(
+    "A vibrant city street at sunset.",
+    sampling_param=sampling,
+    output_path="video_samples",
+    save_video=True,
+)
+```
+
+## Some questions to ask yourself before starting
+
+Answering these upfront clarifies the work and speeds up implementation.
+
+### Is the model already supported by SGLang's multimodal generation stack?
+If yes, you can port many components from SGLang. It is a FastVideo fork, so
+interfaces line up, but you still need to swap layers/modules to match
+FastVideo's architecture and attention stack.
+
+If not, implement the model directly in FastVideo.
+
+### Is there an official implementation of the model you are adding?
+
+If yes, use it as the numerical reference. For example, LTX‑2 has an official
+implementation here: https://github.com/Lightricks/LTX-2. Prefer official code
+even if Diffusers also has one.
+
+### Is there a HuggingFace repo for the model you are adding? Is it in Diffusers format?
+
+If yes, load it directly in FastVideo after setting tensor mapping rules in the
+config. Otherwise, convert the weights to Diffusers format. See [Weights and
+Diffusers format](../design/overview.md#weights-and-diffusers-format) for details.
+
+### Do I have official weights + local paths ready?
+
+Standardize local paths as:
+
+- `official_weights/<model_name>/` for official checkpoints
+- `converted_weights/<model_name>/` if conversion is required (can be created later)
+
+### What pipeline components are required for the model you are adding?
+
+Usually you need a transformer (DiT), VAE, text encoder, and tokenizer. Some
+models add extra components.
+
+### What tasks does the model support?
+
+Usually a video diffusion model supports text‑to‑video (T2V),
+image‑to‑video (I2V), and video‑to‑video (V2V). Some add extra tasks (two‑stage
+generation, keyframe interpolation), which require extra components.
+
+It's usually easiest to start with a T2V pipeline and add the other tasks later.
+
+You can refer to the [Pipeline system](../design/overview.md#pipeline-system)
+section for more details.
+
+### Am I able to generate videos with the official implementation?
+
+These videos and prompts are your reference. Once the FastVideo pipeline works,
+compare outputs to the official implementation. Due to seeding and other
+factors, outputs may not match exactly, but they should be comparable.
+
+## Workflow: adding a full pipeline
+
+This is an example workflow for adding a full model pipeline (model + configs +
+examples + tests). This guide is in active development; feedback is welcome.
+
+!!! note
+    If you get stuck, refer to existing models/pipelines in FastVideo or ask in Slack.
+
+### 0) Fetch official model's code and weights
+
+Purpose:
+
+- Keep official checkpoints and source code local so conversion, parity tests,
+  and reference runs are reproducible.
+- Clone the official repo so you can use it as a numerical reference.
+
+Action:
+
+- Download official weights into `official_weights/<model_name>/`
+  (Diffusers format or not).
+- Clone the official repo under the project root (e.g., `FastVideo/LTX-2/`).
+- If a Diffusers-format HF repo already exists, you can skip manual weight
+  handling and download it directly with
+  `scripts/huggingface/download_hf.py`.
+
+!!! note
+    This step is best done manually because large downloads can time out.
+    Example:
+    ```bash
+    python scripts/huggingface/download_hf.py \
+      --repo_id Wan-AI/Wan2.1-T2V-1.3B-Diffusers \
+      --local_dir official_weights/Wan2.1-T2V-1.3B-Diffusers \
+      --repo_type model
+    ```
+  
+### 1) Implement the model + config mapping
+
+Purpose:
+
+- Model weights are a dictionary of named tensors (`state_dict`). If the names
+  don’t line up with FastVideo’s module names, weights won’t load correctly (or
+  will silently load into the wrong layer).
+- Official checkpoints often use different prefixes or module layouts than
+  FastVideo, so we translate names via the mapping (during load or conversion).
+- Mapping aligns three things:
+  1. the official implementation’s module names,
+  2. the checkpoint `state_dict` keys,
+  3. FastVideo’s model classes and layer naming conventions.
+- If names don’t align, weights won’t load; implement the FastVideo model and
+  define mapping rules first.
+
+Action:
+
+- Implement the FastVideo model + config mapping.
+  - Add/extend the model in `fastvideo/models/...` and config in
+    `fastvideo/configs/models/...` (including `param_names_mapping`).
+  - Reuse existing FastVideo layers/modules where possible.
+  - Use FastVideo’s attention layers:
+    - `DistributedAttention` only for full‑sequence self‑attention in the DiT.
+    - `LocalAttention` for cross‑attention and other attention layers.
+  - See the “Configuration System” and “Weights and Diffusers format” sections
+    in `docs/design/overview.md` for how these pieces connect.
+  - If you are using an agent, ask it to implement the model, config mapping,
+    and a parity test together so you can validate numerics immediately.
+
+!!! note
+    After the first component is aligned and parity‑tested, open a **DRAFT PR**
+    on FastVideo so the rest of the pipeline work can build on top of it.
+
+!!! note
+    If a Diffusers-format HF repo already exists and loads correctly, you can
+    skip conversion entirely (no conversion script needed) and just download it
+    with `scripts/huggingface/download_hf.py`. Otherwise, you may need a
+    conversion script + a `converted_weights/<model>/` staging directory.
+
+Example (key renaming via arch config mapping, Wan2.1‑style):
+
+```python
+# Official model (simplified) in the upstream repo.
+class OfficialWanTransformer(torch.nn.Module):
+    def __init__(self):
+        super().__init__()
+        self.patch_embedding = torch.nn.Conv3d(16, 1536, kernel_size=2, padding=0)
+
+    def forward(self, x):
+        return self.patch_embedding(x)
+
+# FastVideo model (simplified) in fastvideo/models/dits/wanvideo.py
+class PatchEmbed(torch.nn.Module):
+    def __init__(self):
+        super().__init__()
+        self.proj = torch.nn.Conv3d(16, 1536, kernel_size=2, padding=0)
+
+    def forward(self, x):
+        return self.proj(x)
+
+class WanTransformer3DModel(torch.nn.Module):
+    def __init__(self):
+        super().__init__()
+        self.patch_embedding = PatchEmbed()
+
+    def forward(self, x):
+        return self.patch_embedding(x)
+
+# Mapping defined in a config (simplified; see the real mapping in
+# fastvideo/configs/models/dits/wanvideo.py)
+param_names_mapping = {
+    r"^patch_embedding\.(.*)$": r"patch_embedding.proj.\1",
+    r"^blocks\.(\d+)\.attn1\.to_q\.(.*)$": r"blocks.\1.to_q.\2",
+}
+
+def apply_regex_map(state_dict, mapping):
+    # Pseudocode: apply regex substitutions in order
+    ...
+
+# Official checkpoint keys (example)
+official = {
+    "patch_embedding.weight": ...,
+    "blocks.0.attn1.to_q.weight": ...,
+}
+
+# Apply mapping so keys match FastVideo modules
+converted = apply_regex_map(official, param_names_mapping)
+
+```
+
+Optional helper (print a few checkpoint keys quickly):
+
+```bash
+python - <<'PY'
+import safetensors.torch as st
+keys = list(st.load_file("official_weights/<model>/transformer/diffusion_pytorch_model.safetensors").keys())
+print(keys[:20])
+PY
+```
+
+Example agent prompt (task request):
+
+```
+Please add the Wan2.1 T2V 1.3B Diffusers pipeline to FastVideo:
+- Add a FastVideo native Wan2.1 DiT implementation + config mapping.
+- Make sure to use the existing FastVideo layers and attention modules where possible.
+- Add a parity test that loads the official model alongside the FastVideo model and compares outputs numerically with fixed seeds and inputs.
+
+Paths:
+  - Official repo: Wan-AI/Wan2.1-T2V-1.3B-Diffusers
+  - Local download: official_weights/Wan2.1-T2V-1.3B-Diffusers
+Mapping steps:
+  - Load the official DiT weights from
+    official_weights/Wan2.1-T2V-1.3B-Diffusers/transformer/diffusion_pytorch_model.safetensors.
+  - Instantiate the FastVideo DiT (`WanTransformer3DModel`) and compare
+    its `state_dict().keys()` to the official keys.
+  - Update `param_names_mapping` in
+    fastvideo/configs/models/dits/wanvideo.py to resolve missing/unexpected keys.
+  - Use `load_state_dict(strict=False)` during iteration to surface mismatches.
+```
+
+External examples of the same pattern:
+- SGLang uses prefix-based routing in its weight loader to map checkpoint keys
+  into internal submodules (e.g., stripping a top-level prefix before delegating).
+- vLLM includes model-specific renamers for certain checkpoints that adjust
+  key prefixes so weights match its internal naming.
+
+### 2) Test numerical alignment with the official implementation
+
+Purpose:
+
+- Verify that the FastVideo component is numerically aligned with the official
+  implementation.
+
+Action:
+
+- Add or reuse a numerical parity test that loads the official model and the
+  FastVideo model and compares outputs.
+- See examples in `tests/local_tests/` (e.g., `tests/local_tests/upsamplers/`)
+  and the commands in `tests/local_tests/README.md`.
+- If there are discrepancies, add opt‑in logging to both models and compare
+  activation summaries (layer output sums, per‑stage logs).
+- First align the loaded weights (validate `param_names_mapping`).
+- Then align forward outputs using fixed seeds and inputs.
+  - Start with `atol=1e-4, rtol=1e-4` in `assert_close`.
+  - Keep dtype consistent (bf16 if available; otherwise fp32).
+  - If attention parity is unstable, align backends (e.g.,
+    `FASTVIDEO_ATTENTION_BACKEND=TORCH_SDPA`).
+
+### 3) Repeat the process for each component
+
+If the model requires additional components, repeat Steps 1–2 for each one.
+For example, implement the VAE in `fastvideo/models/vaes/` and its config in
+`fastvideo/configs/models/vaes/`, then add parity coverage for it.
+
+### 4) Add a pipeline config + sample defaults
+
+Purpose:
+
+- `fastvideo/configs/pipelines/` describes pipeline wiring and model module
+  names.
+- `fastvideo/configs/sample/` defines default runtime parameters.
+
+Action:
+
+- Add a new pipeline config + sampling params.
+- Register them in `fastvideo/registry.py` using explicit
+  `register_configs(...)` blocks (this file is the single source of truth now).
+
+### 5) Wire pipeline stages
+
+Purpose:
+
+- `fastvideo/pipelines/basic/<pipeline>/` contains the actual pipeline logic.
+- `fastvideo/pipelines/stages/` holds reusable, testable stages.
+
+Action:
+
+- Build the pipeline using stages; keep new stages isolated and documented.
+- Prefer opt‑in flags for expensive or optional steps.
+
+### 6) Add pipeline‑level tests
+
+Purpose:
+
+- Ensure the end‑to‑end pipeline works and stays aligned as pieces evolve.
+
+Action:
+
+- Add a pipeline parity test under `tests/local_tests/pipelines/`.
+- See the [Testing Guide](testing.md) for test conventions.
+
+### 7) Add user‑facing examples
+
+Purpose:
+
+- `examples/inference/basic/` is the entry point for simple, runnable scripts.
+
+Action:
+
+- Provide a minimal “hello world” example plus advanced variations.
+- Use fixed seeds and stable prompts.
+- Run the example locally to confirm end‑to‑end behavior.
+
+### 8) Add SSIM tests for CI checks
+
+Purpose:
+
+- Ensure visual similarity stays within expected bounds for regression testing.
+- SSIM tests act as a higher‑level guardrail beyond unit/parity tests.
+
+Action:
+
+- Add SSIM tests under `fastvideo/tests/ssim/` and include reference videos
+  (see the structure in the Testing Guide).
+- Use stable prompts/seeds and document any GPU‑specific requirements.
+- Follow the [Testing Guide](testing.md) for reference video placement and
+  execution details.
+
+### 9) Document it
+
+Purpose:
+
+- `docs/` is where users find the new pipeline usage and limitations.
+
+Action:
+
+- Add a short doc page or update an existing one.
+- Mention any caveats (memory, speed, constraints).
+
+## Common pitfalls when porting models
+
+- **Attention backend mismatch**: parity can fail if the official model uses a
+  different attention backend (e.g., SDPA vs custom). Align backends before
+  debugging deeper issues.
+- **Patchifier shape mistakes**: wrong patchification or reshape lengths can
+  silently corrupt outputs. Validate patch shapes early.
+- **Mask handling**: attention masks must match the official behavior (padding,
+  causal masks, and broadcast shapes).
+- **Scheduler / sigma schedule mismatch**: even small differences in schedules
+  or timestep shapes can cause noticeable drift.
+
+## Diffusers vs manual conversion
+
+If a model already ships in Diffusers format (with a proper `model_index.json`),
+prefer downloading it directly and loading it via FastVideo. In that case:
+
+- You usually **do not need** a conversion script.
+- You still need a correct `param_names_mapping` if the internal module names
+  differ from FastVideo’s implementation.
+
+If the model does **not** have a Diffusers-format repo:
+
+- You will need a conversion script to rewrite `state_dict` keys into FastVideo
+  naming and stage the result (e.g., under `converted_weights/<model>/`).
+- You may still use the official repo for reference parity and debugging.
+
+In both cases, parity testing is required to validate correctness.
+
+If you want to publish a Diffusers‑style repo after conversion, use
+`scripts/checkpoint_conversion/create_hf_repo.py` to assemble a HuggingFace‑ready
+directory before uploading.
+
+## FAQ
+
+**Q: Why do we implement the FastVideo model before conversion?**  
+A: You can’t define the key‑mapping rules until the FastVideo module names are
+known. The implementation determines the target `state_dict` schema.
+
+**Q: Do we always need a conversion script?**  
+A: No. If a Diffusers‑format repo exists and loads correctly, download it and
+skip conversion.
+
+**Q: How do I figure out `param_names_mapping` quickly?**  
+A: Load the official weights, instantiate the FastVideo model, and diff
+`state_dict().keys()` on both sides. Add regex rules until missing/unexpected
+keys are resolved. Agents can help you with this.
+
+**Q: What if parity fails even after mapping?**  
+A: Align attention backends, sigma schedules, and timestep shapes first. Then
+add opt‑in activation logging to locate the first divergent layer.
+
+## Case study: LTX‑2 port (from PLAN.md)
+
+The LTX‑2 port in `PLAN.md` shows the real sequence of steps and backtracking
+that happened during integration. Use it as a reference for how parity work
+actually unfolds:
+
+- Ported components first (transformer, VAE, audio, text encoder).
+- Added parity tests per component; used SDPA for reference parity.
+- Added debug logging to compare per‑block activations and isolate divergence.
+- Fixed cross‑attention reshape and patch grid bounds issues after logging.
+- Aligned sigma schedule and masking behavior to match the official pipeline.
+
+Recommendation:
+
+- Keep raw step‑by‑step logs in your own local `PLAN.md` for large ports.
+
+## Worked example: Wan2.1 T2V 1.3B pipeline
+
+The Wan2.1 T2V 1.3B Diffusers pipeline is a good “standard” example for
+FastVideo integration.
+
+1. Verify model config + mapping.
+   - DiT mapping: `fastvideo/configs/models/dits/wanvideo.py`
+   - VAE: `fastvideo/models/vaes/wanvae.py`
+   - Text encoder: `fastvideo/models/encoders/t5.py`
+
+2. Parity test the core components.
+   - Example tests: `fastvideo/tests/transformers/test_wanvideo.py`,
+     `fastvideo/tests/vaes/test_wan_vae.py`,
+     `fastvideo/tests/encoders/test_t5_encoder.py`
+
+3. Pipeline wiring.
+   - Pipeline: `fastvideo/pipelines/basic/wan/wan_pipeline.py`
+   - Pipeline config: `fastvideo/configs/pipelines/wan.py`
+   - Sampling defaults: `fastvideo/configs/sample/wan.py`
+
+4. Minimal example.
+   - Script: `examples/inference/basic/basic.py`
@@ -1,16 +1,55 @@

 # 📦 Developing FastVideo on RunPod

-You can easily use the FastVideo Docker image as a custom container on [RunPod](https://www.runpod.io) for development or experimentation.
+You can easily use the FastVideo Pod Template on [RunPod](https://www.runpod.io) for development or experimentation.

 ## Creating a new pod

-Choose a GPU that supports CUDA 12.8
-
-Pick 1 or 2 L40S GPU(s)
+- Make sure you are using the correct RunPod account.
+![RunPod Account Selection](../../assets/images/runpod_account.png)

+- Use "Additional Filters" to select CUDA 12.8.
 ![RunPod CUDA selection](../../assets/images/runpod_cuda.png)

+- Click "Deploy" and Pick a single A40 or RTX 4090 GPU.
+![RunPod GPU Selection](../../assets/images/runpod_deploy.png)
+
+- Select the "FastVideo" or "fastvideo-dev" Pod Template.
+![RunPod Pod Template Selection](../../assets/images/runpod_create.png)
+
+- Set the Pod name to "`<name>-<FastVideo>-<date>`".
+
+- Finally, once the pod is deployed (will take a few minutes as the image is being pulled), you can SSH into it using "SSH exposed over TCP". You'll need to use the matching private ssh key you provided.
+![RunPod SSH](../../assets/images/runpod_ssh.png)
+
+## Working with the pod
+
+After SSH'ing into your pod, you'll find the correct `uv` environment already activated and you should be in /FastVideo/ directory. Make sure to use /FastVideo/ for all your work.
+
+To pull in the latest changes from the GitHub repo:
+
+```bash
+cd /FastVideo
+git pull
+```
+
+Run your development workflows as usual:
+
+```bash
+# Run linters
+pre-commit run --all-files
+
+# Run tests
+pytest tests/
+```
+
+Make sure to push your changes back to the GitHub repo as nothing will be saved to the pod when it is terminated.
+
+After you are done with your work, you can terminate the pod by clicking the "Terminate" and "Delete" buttons. Remember if the pod is not completely deleted, Runpod will keep charging you for it.
+
+## Extra Information:
+If you need to customize the pod template this section has some useful information. For the most part you can leave the defaults of the FastVideo Pod Template.
+
 When creating your pod template, use this image:

 ```
@@ -24,30 +63,3 @@ bash -c "apt update;DEBIAN_FRONTEND=noninteractive apt-get install openssh-serve
 ```

 ![RunPod template configuration](../../assets/images/runpod_template.png)
-
-After deploying, the pod will take a few minutes to pull the image and start the SSH service.
-
-![RunPod ssh](../../assets/images/runpod_ssh.png)
-
-## Working with the pod
-
-After SSH'ing into your pod, you'll find the `fastvideo-dev` Conda environment already activated.
-
-To pull in the latest changes from the GitHub repo:
-
-```bash
-cd /FastVideo
-git pull
-```
-
-`If you have a persistent volume and want to keep your code changes, you can move /FastVideo to /workspace/FastVideo, or simply clone the repository there.`
-
-Run your development workflows as usual:
-
-```bash
-# Run linters
-pre-commit run --all-files
-
-# Run tests
-pytest tests/
-```
@@ -1,76 +1,84 @@

 # 🛠️ Contributing to FastVideo

-Thank you for your interest in contributing to FastVideo. We want to make the process as smooth for you as possible and this is a guide to help get you started!
+Thank you for your interest in contributing to FastVideo. We want the process
+to be smooth and beginner‑friendly, whether you are adding a new pipeline,
+improving performance, or fixing a bug.

-Our community is open to everyone and welcomes any contributions no matter how large or small.
+## Quick prerequisites

-# Developer Environment:
-Do make sure you have CUDA 12.4 installed and supported. FastVideo currently only supports Linux and CUDA GPUs, but we hope to support other platforms in the future.
+- **OS**: Linux is the primary development target (WSL can work).
+- **GPU**: NVIDIA GPU recommended for inference and training workflows.
+- **CUDA**: Use a recent CUDA 12.x toolchain (see the installation guide for
+  the current recommendation).

-We recommend using a fresh Python 3.10 Conda environment to develop FastVideo:
+For a full install checklist, see `docs/getting_started/installation/gpu.md`.
+
+## Local development (Conda + editable install)

 Install Miniconda:

-```
+```bash
 wget https://repo.anaconda.com/miniconda/Miniconda3-latest-Linux-x86_64.sh
 bash Miniconda3-latest-Linux-x86_64.sh
 source ~/.bashrc
 ```

-Create and activate a Conda environment for FastVideo:
+Create and activate a Conda environment:

-```
+```bash
 conda create -n fastvideo python=3.12 -y
 conda activate fastvideo
 ```

 Install `uv` (optional, but recommended):

-From instructions on [uv](https://astral.sh/uv/):
-
-```
+```bash
 curl -LsSf https://astral.sh/uv/install.sh | sh
-# or 
+# or
 wget -qO- https://astral.sh/uv/install.sh | sh
 ```

-Clone the FastVideo repository and go to the FastVideo directory:
+Clone the repo:

-```
+```bash
 git clone https://github.com/hao-ai-lab/FastVideo.git && cd FastVideo
-
 ```

-Now you can install FastVideo and setup git hooks for running linting. By using `pre-commit`, the linters will run and have to pass before you'll be able to make a commit.
+Install FastVideo in editable mode and set up hooks:

 ```bash
 uv pip install -e .[dev]

-# Can also install flash-attn (optional)
-uv pip install flash-attn --no-build-isolation 
+# Optional: FlashAttention (builds native kernels)
+uv pip install flash-attn --no-build-isolation

-# Linting, formatting and static type checking
+# Linting, formatting, static typing
 pre-commit install --hook-type pre-commit --hook-type commit-msg
-
-# You can manually run pre-commit with
 pre-commit run --all-files

 # Unit tests
 pytest tests/
 ```

-If you are on a Hopper GPU, you should also install [FA3](https://github.com/Dao-AILab/flash-attention) for much better performance:
+If you are on a Hopper GPU, installing FlashAttention 3 can improve
+performance (see `docs/inference/optimizations.md`).

-```
-git clone https://github.com/Dao-AILab/flash-attention.git && cd flash-attention/hopper
+## Docker development (optional)

-# make sure you have ninja installed
-uv pip install ninja
-
-python setup.py install
-```
+If you prefer a containerized environment, use the dev image documented in
+`docs/contributing/developer_env/docker.md`.

 ## Testing

-Please refer to the [Testing Guide](testing.md) for more information on how to add and run tests in FastVideo.
+See the [Testing Guide](testing.md) for how to add and run tests in FastVideo.
+
+## Attention backend development
+
+If you are adding a new attention kernel or backend, follow
+[Attention Backend Development](attention_backend.md).
+
+## Contributing with coding agents
+
+For a step‑by‑step workflow on adding pipelines or components with coding
+agents, see `docs/contributing/coding_agents.md`.
@@ -1,424 +1,179 @@
-# 🔍 FastVideo Overview
+# FastVideo Architecture Overview

-This document outlines FastVideo's architecture for developers interested in framework internals or contributions. It serves as an onboarding guide for new contributors by providing an overview of the most important directories and files within the `fastvideo/` codebase.
+This document summarizes how FastVideo is structured and how a Diffusers-style
+model repo maps into a runnable pipeline. It is intended for contributors who
+need the high-level layout and key entrypoints, not every internal detail.

-## Table of Contents - Directory Structure and Files
+## FastVideo structure at a glance

- [`fastvideo/pipelines/`](#pipeline-system) - Core diffusion pipeline components
- [`fastvideo/models/`](#model-components) - Model implementations
-  - [`dits/`](#transformer-models) - Transformer-based diffusion models
-  - [`vaes/`](#vae-variational-auto-encoder) - Variational autoencoders
-  - [`encoders/`](#text-and-image-encoders) - Text and image encoders
-  - [`schedulers/`](#schedulers) - Diffusion schedulers
- [`fastvideo/attention/`](#optimized-attention) - Optimized attention implementations
- [`fastvideo/distributed/`](#distributed-processing) - Distributed computing utilities
- [`fastvideo/layers/`](#tensor-parallelism) - Custom neural network layers
- [`fastvideo/platforms/`](#platforms) - Hardware platform abstractions
- [`fastvideo/worker/`](#executor-and-worker-system) - Multi-GPU process management
- [`fastvideo/fastvideo_args.py`](#fastvideoargs) - Argument handling
- [`fastvideo/forward_context.py`](#forward-context-management) - Forward pass context management
- `fastvideo/utils.py` - Utility functions
- [`fastvideo/logger.py`](#logger) - Logging infrastructure
+FastVideo maps a Diffusers-style repo into a pipeline like this:

-## Core Architecture
+- `fastvideo/models/*`: model implementations (DiT, VAE, encoders, upsamplers).
+- `fastvideo/configs/models/*`: arch configs and `param_names_mapping` for
+  weight name translation.
+- `fastvideo/configs/pipelines/*`: pipeline wiring (component classes + names).
+- `fastvideo/configs/sample/*`: default runtime sampling parameters.
+- `fastvideo/pipelines/basic/*`: end-to-end pipelines.
+- `fastvideo/pipelines/stages/*`: reusable pipeline stages.
+- `fastvideo/models/loader/*`: component loaders for Diffusers-style repos.
+- `model_index.json`: HF repo entrypoint mapping component names to classes.

-FastVideo separates model components from execution logic with these principles:
+Flow:
+`model_index.json` -> component loaders -> model modules -> pipeline stages ->
+sampling params.

- **Component Isolation**: Models (encoders, VAEs, transformers) are isolated from execution (pipelines, stages, distributed processing)
- **Modular Design**: Components can be independently replaced
- **Distributed Execution**: Supports various parallelism strategies (Tensor, Sequence)
- **Custom Attention Backends**: Components can support and use different Attention implementations
- **Pipeline Abstraction**: Consistent interface across diffusion models
-
-## FastVideoArgs
-
-The `FastVideoArgs` class in `fastvideo/fastvideo_args.py` serves as the central configuration system for FastVideo. It contains all parameters needed to control model loading, inference configuration, performance optimization settings, and more.
-
-Key features include:
-
- **Command-line Interface**: Automatic conversion between CLI arguments and dataclass fields
- **Configuration Groups**: Organized by functional areas (model loading, video params, optimization settings)
- **Context Management**: Global access to current settings via `get_current_fastvideo_args()`
- **Parameter Validation**: Ensures valid combinations of settings
-
-Common configuration areas:
-
- **Model paths and loading options**: `model_path`, `trust_remote_code`, `revision`
- **Distributed execution settings**: `num_gpus`, `tp_size`, `sp_size`
- **Video generation parameters**: `height`, `width`, `num_frames`, `num_inference_steps`
- **Precision settings**: Control computation precision for different components
-
-Example usage:
+Minimal usage (from `examples/inference/basic/basic.py`):

 ```python
-# Load arguments from command line
-fastvideo_args = prepare_fastvideo_args(sys.argv[1:])
+from fastvideo import VideoGenerator
+from fastvideo.configs.sample import SamplingParam

-# Access parameters
-model = load_model(fastvideo_args.model_path)
+model_id = "Wan-AI/Wan2.1-T2V-1.3B-Diffusers"  # or official_weights/<model_name>/
+generator = VideoGenerator.from_pretrained(model_id, num_gpus=1)

-# Set as global context
-with set_current_fastvideo_args(fastvideo_args):
-    # Code that requires access to these arguments
-    result = generate_video()
-```
-
-## Pipeline System
-
-### `ComposedPipelineBase`
-
-This foundational class provides:
-
- **Model Loading**: Automatically loads components from HuggingFace-Diffusers-compatible model directories
- **Stage Management**: Creates and orchestrates processing stages
- **Data Flow Coordination**: Ensures proper state flow between stages
-
-```python
-class MyCustomPipeline(ComposedPipelineBase):
-    _required_config_modules = [
-        "text_encoder", "tokenizer", "vae", "transformer", "scheduler"
-    ]
-    
-    def initialize_pipeline(self, fastvideo_args: FastVideoArgs):
-        # Pipeline-specific initialization
-        pass
-        
-    def create_pipeline_stages(self, fastvideo_args: FastVideoArgs):
-        self.add_stage("input_validation_stage", InputValidationStage())
-        self.add_stage("text_encoding_stage", CLIPTextEncodingStage(
-            text_encoder=self.get_module("text_encoder"),
-            tokenizer=self.get_module("tokenizer")
-        ))
-        # Additional stages...
-```
-
-### Pipeline Stages
-
-Each stage handles a specific diffusion process component:
-
- **Input Validation**: Parameter verification
- **Text Encoding**: CLIP, LLaMA, or T5-based encoding
- **Image Encoding**: Image input processing
- **Timestep & Latent Preparation**: Setup for diffusion
- **Denoising**: Core diffusion loop
- **Decoding**: Latent-to-pixel conversion
-
-Each stage implements a standard interface:
-
-```python
-def forward(self, batch: ForwardBatch, fastvideo_args: FastVideoArgs) -> ForwardBatch:
-    # Process batch and update state
-    return batch
-```
-
-![Pipeline execution and data flow](../assets/images/pipeline.png)
-
-### ForwardBatch
-
-Defined in `fastvideo/pipelines/pipeline_batch_info.py`, `ForwardBatch` encapsulates the data payload passed between pipeline stages. It typically holds:
-
- **Input Data**: Prompts, images, generation parameters
- **Intermediate State**: Embeddings, latents, timesteps, accumulated during stage execution
- **Output Storage**: Generated results and metadata
- **Configuration**: Sampling parameters, precision settings
-
-This structure facilitates clear state transitions between stages.
-
-## Model Components
-
-The `fastvideo/models/` directory contains implementations of the core neural network models used in video diffusion:
-
-### Transformer Models
-
-Transformer networks perform the actual denoising during diffusion:
-
- **Location**: `fastvideo/models/dits/`
- **Examples**:
-  - `WanTransformer3DModel`
-  - `HunyuanVideoTransformer3DModel`
-
-Features include:
-
- Text/image conditioning
- Standardized interface for model-specific optimizations
-
-```python
-def forward(
-    self, 
-    latents,                    # [B, T, C, H, W]
-    encoder_hidden_states,      # Text embeddings
-    timestep,                   # Current diffusion timestep
-    encoder_hidden_states_image=None,  # Optional image embeddings
-    **kwargs
-):
-    # Perform denoising computation
-    return noise_pred  # Predicted noise residual
-```
-
-### VAE (Variational Auto-Encoder)
-
-VAEs handle conversion between pixel space and latent space:
-
- **Location**: `fastvideo/models/vaes/`
- **Examples**:
-  - `AutoencoderKLWan`
-  - `AutoencoderKLHunyuanVideo`
-
-These models compress image/video data to a more efficient latent representation (typically 4x-8x smaller in each dimension).
-
-FastVideo's VAE implementations include:
-
- Efficient video batch processing
- Memory optimization
- Optional tiling for large frames
- Distributed weight support
-
-### Text and Image Encoders
-
-Encoders process conditioning inputs into embeddings:
-
- **Location**: `fastvideo/models/encoders/`
- **Text Encoders**:
-  - `CLIPTextModel`
-  - `LlamaModel`
-  - `UMT5EncoderModel`
- **Image Encoders**:
-  - `CLIPVisionModel`
-
-FastVideo implements optimizations such as:
-
- Vocab parallelism for distributed processing
- Caching for common prompts
- Precision-tuned computation
-
-### Schedulers
-
-Schedulers manage the diffusion sampling process:
-
- **Location**: `fastvideo/models/schedulers/`
- **Examples**:
-  - `UniPCMultistepScheduler`
-  - `FlowMatchEulerDiscreteScheduler`
-
-These components control:
-
- Diffusion timestep sequences
- Noise prediction to latent update conversions
- Quality/speed trade-offs
-
-```python
-def step(
-    self, 
-    model_output: torch.Tensor,
-    timestep: torch.LongTensor,
-    sample: torch.Tensor,
-    **kwargs
-) -> torch.Tensor:
-    # Process model output and update latents
-    # Return updated latents
-    return prev_sample
-```
-
-This diagram shows how models are discovered, validated, and loaded across entrypoints, executors, pipelines, and model loaders.
-
-![Model loading flow](../assets/images/load_models.png)
-
-## Optimized Attention
-
-The `fastvideo/attention/` directory contains optimized attention implementations crucial for efficient video diffusion:
-
-### Attention Backends
-
-Multiple implementations with automatic selection:
-
- **FLASH_ATTN**: Optimized for supporting hardware
- **TORCH_SDPA**: Built-in PyTorch scaled dot-product attention
- **SLIDING_TILE_ATTN**: For very long sequences
-
-```python
-# Configure available attention backends for this layer
-self.attn = LocalAttention(
-    num_heads=num_heads,
-    head_size=head_dim,
-    causal=False,
-    supported_attention_backends=(_Backend.FLASH_ATTN, _Backend.TORCH_SDPA)
-)
-
-# Override via environment variable
-# export FASTVIDEO_ATTENTION_BACKEND=FLASH_ATTN
-```
-
-![Attention backend selector design](../assets/images/attention_backend.png)
-
-### Attention Patterns
-
-Supports various patterns with memory optimization techniques:
-
- **Cross/Self/Temporal/Global-Local Attention**
- Chunking, progressive computation, optimized masking
-
-## Distributed Processing
-
-The `fastvideo/distributed/` directory contains implementations for distributed model execution:
-
-### Tensor Parallelism
-
-Tensor parallelism splits model weights across devices:
-
- **Implementation**: Through `RowParallelLinear` and `ColumnParallelLinear` layers
- **Use cases**: Will be used by encoder models as their sequence lengths are shorter and enables efficient sharding.
-
-```python
-# Tensor-parallel layers in a transformer block
-from fastvideo.layers.linear import ColumnParallelLinear, RowParallelLinear
-
-# Split along output dimension
-self.qkv_proj = ColumnParallelLinear(
-    input_size=hidden_size,
-    output_size=3 * hidden_size,
-    bias=True,
-    gather_output=False
-)
-
-# Split along input dimension
-self.out_proj = RowParallelLinear(
-    input_size=hidden_size,
-    output_size=hidden_size,
-    bias=True,
-    input_is_parallel=True
+sampling = SamplingParam.from_pretrained(model_id)
+sampling.num_frames = 45
+video = generator.generate_video(
+    "A vibrant city street at sunset.",
+    sampling_param=sampling,
+    output_path="video_samples",
+    save_video=True,
 )
 ```

-### Sequence Parallelism
+## Configuration system

-Sequence parallelism splits sequences across devices:
+FastVideo uses typed configs to keep model definitions, pipeline wiring, and
+runtime parameters consistent:

- **Implementation**: Through `DistributedAttention` and sequence splitting
- **Use cases**: Long video sequences or high-resolution processing. Used by DiT models.
+- `fastvideo/configs/models/`: architecture definitions, layer shapes, and
+  `param_names_mapping` rules for key renaming.
+- `fastvideo/configs/pipelines/`: pipeline wiring and required components.
+- `fastvideo/configs/sample/`: default sampling parameters (steps, frames,
+  guidance scale, resolution, fps).
+- `fastvideo/registry.py`: unified registry for pipeline config + sampling
+  defaults and model metadata resolution, defined via explicit
+  `register_configs(...)` blocks (no separate dict registries).

-```python
-# Distributed attention for long sequences
-from fastvideo.attention import DistributedAttention
+`FastVideoArgs` (in `fastvideo/fastvideo_args.py`) provides runtime settings and
+is passed into pipeline construction and stages.

-self.attn = DistributedAttention(
-    num_heads=num_heads,
-    head_size=head_dim,
-    causal=False,
-    supported_attention_backends=(_Backend.SLIDING_TILE_ATTN, _Backend.FLASH_ATTN)
-)
+## Weights and Diffusers format
+
+FastVideo follows the HuggingFace Diffusers repo layout. This keeps loaders
+compatible with HF repos and makes it easy to add new components.
+
+Typical Diffusers repo:
+
+```
+<model-repo>/
+  model_index.json
+  scheduler/
+    scheduler_config.json
+  transformer/               # or unet/ for image models
+    config.json
+    diffusion_pytorch_model.safetensors
+  vae/
+    config.json
+    diffusion_pytorch_model.safetensors
+  text_encoder/
+    config.json
+    model.safetensors
+  tokenizer/
+    tokenizer_config.json
+    tokenizer.json
 ```

-### Communication Primitives
+Key points:

-Efficient distributed operations via AllGather, AllReduce, and synchronization mechanisms.
+- `model_index.json` is the root map that tells FastVideo which components to
+  load and which classes implement them.
+- Each component lives in its own folder with a `config.json` and weights.
+- Weights are usually in `diffusion_pytorch_model.safetensors`.

-Efficient communication primitives minimize distributed overhead:
+Note on tensor names:

- **Sequence-Parallel AllGather**: Collects sequence chunks
- **Tensor-Parallel AllReduce**: Combines partial results
- **Distributed Synchronization**: Coordinates execution
+Official checkpoints often use different `state_dict` names than FastVideo's
+module layout. We translate tensor names via the DiT arch config mapping
+(`param_names_mapping` under `fastvideo/configs/models/dits/`). This is similar
+in spirit to name-translation layers used in systems like vLLM and SGLang.

-## Forward Context Management
+Example HF repo (Wan 2.1 T2V 1.3B Diffusers):

-### ForwardContext
-
-Defined in `fastvideo/forward_context.py`, `ForwardContext` manages execution-specific state *within* a forward pass, particularly for low-level optimizations. It is accessed via `get_forward_context()`.
-
- **Attention Metadata**: Configuration for optimized attention kernels (`attn_metadata`)
- **Profiling Data**: Potential hooks for performance metrics collection
-
-This context-based approach enables:
-
- Dynamic optimization based on execution state (e.g., attention backend selection)
- Step-specific customizations within model components
-
-Usage example:
-
-```python
-with set_forward_context(current_timestep, attn_metadata, fastvideo_args):
-    # During this forward pass, components can access context
-    # through get_forward_context()
-    output = model(inputs)
+```
+https://huggingface.co/Wan-AI/Wan2.1-T2V-1.3B-Diffusers/tree/main
 ```

-## Executor and Worker System
+Example `model_index.json` from that repo:

-The `fastvideo/worker/` directory contains the distributed execution framework:
-
-### Executor Abstraction
-
-FastVideo implements a flexible execution model for distributed processing:
-
- **Executor Base Class**: An abstract base class defining the interface for all executors
- **MultiProcExecutor**: Primary implementation that spawns and manages worker processes
- **GPU Workers**: Handle actual model execution on individual GPUs
-
-The MultiProcExecutor implementation:
-
-1. Spawns worker processes for each GPU
-2. Establishes communication channels via pipes
-3. Coordinates distributed operations across workers
-4. Handles graceful startup and shutdown of the process group
-
-Each GPU worker:
-
-1. Initializes the distributed environment
-2. Builds the pipeline for the specified model
-3. Executes requested operations on its assigned GPU
-4. Manages local resources and communicates results back to the executor
-
-This design allows FastVideo to efficiently utilize multiple GPUs while providing a simple, unified interface for model execution.
-
-## Platforms
-
-The `fastvideo/platforms/` directory provides hardware platform abstractions that enable FastVideo to run efficiently on different hardware configurations:
-
-### Platform Abstraction
-
-FastVideo's platform abstraction layer enables:
-
- **Hardware Detection**: Automatic detection of available hardware
- **Backend Selection**: Appropriate selection of compute kernels
- **Memory Management**: Efficient utilization of hardware-specific memory features
-
-The primary components include:
-
- **Platform Interface**: Defines the common API for all platform implementations
- **CUDA Platform**: Optimized implementation for NVIDIA GPUs
- **Backend Enum**: Used throughout the codebase for feature selection
-
-Usage example:
-
-```python
-from fastvideo.platforms import current_platform, _Backend
-
-# Check hardware capabilities
-if current_platform.supports_backend(_Backend.FLASH_ATTN):
-    # Use FlashAttention implementation
-else:
-    # Fall back to standard implementation
+```json
+{
+  "_class_name": "WanPipeline",
+  "_diffusers_version": "0.33.0.dev0",
+  "scheduler": [
+    "diffusers",
+    "UniPCMultistepScheduler"
+  ],
+  "text_encoder": [
+    "transformers",
+    "UMT5EncoderModel"
+  ],
+  "tokenizer": [
+    "transformers",
+    "T5TokenizerFast"
+  ],
+  "transformer": [
+    "diffusers",
+    "WanTransformer3DModel"
+  ],
+  "vae": [
+    "diffusers",
+    "AutoencoderKLWan"
+  ]
+}
 ```

-The platform system is designed to be extensible for future hardware targets.
+How this maps to FastVideo:

-## Logger
+- `WanPipeline` -> `fastvideo/pipelines/basic/wan/wan_pipeline.py`
+- `WanTransformer3DModel` -> `fastvideo/models/dits/wanvideo.py`
+- `AutoencoderKLWan` -> `fastvideo/models/vaes/wanvae.py`
+- `UMT5EncoderModel` -> `fastvideo/models/encoders/t5.py`
+- `T5TokenizerFast` -> loaded via HF in `fastvideo/models/loader/`
+- `UniPCMultistepScheduler` -> loaded via Diffusers scheduler utilities
+- Pipeline defaults -> `fastvideo/configs/pipelines/wan.py`
+- Sampling defaults -> `fastvideo/configs/sample/wan.py`

-See [PR](https://github.com/hao-ai-lab/FastVideo/pull/356)
+## Pipeline system

-*TODO*: (help wanted) Add an environment variable that disables process-aware logging.
+- `fastvideo/pipelines/basic/*` contains end-to-end pipelines for each model
+  family.
+- `fastvideo/pipelines/stages/*` contains reusable, testable stages.
+- Pipelines subclass `ComposedPipelineBase` and declare required components via
+  `_required_config_modules`.
+- `ForwardBatch` (in `fastvideo/pipelines/pipeline_batch_info.py`) carries
+  prompts, latents, timesteps, and intermediate state across stages.

-## Contributing to FastVideo
+## Model components

-If you're a new contributor, here are some common areas to explore:
+- DiT models: `fastvideo/models/dits/`
+- VAEs: `fastvideo/models/vaes/`
+- Text/image encoders: `fastvideo/models/encoders/`
+- Schedulers: `fastvideo/models/schedulers/`
+- Upsamplers: `fastvideo/models/upsamplers/`
+- Optional audio models: `fastvideo/models/audio/`

-1. **Adding a new model**: Implement new model types in the appropriate subdirectory of `fastvideo/models/`
-2. **Optimizing performance**: Look at attention implementations or memory management
-3. **Adding a new pipeline**: Create a new pipeline subclass in `fastvideo/pipelines/`
-4. **Hardware support**: Extend the `platforms` module for new hardware targets
+## Attention and distributed execution

-When adding code, follow these practices:
+- Attention backends live in `fastvideo/attention/` and can be selected via
+  `FASTVIDEO_ATTENTION_BACKEND`.
+- `LocalAttention` is used for cross-attention and most attention layers.
+- `DistributedAttention` is used for full-sequence self-attention in the DiT.
+- Tensor-parallel layers live in `fastvideo/layers/`.
+- Sequence/tensor parallel utilities live in `fastvideo/distributed/`.

- Use type hints for better code readability
- Add appropriate docstrings
- Maintain the separation between model components and execution logic
- Follow existing patterns for distributed processing
+## Related docs
+
+- [Contributing overview](../contributing/overview.md)
+- [Coding agents workflow](../contributing/coding_agents.md)
+- [Testing guide](../contributing/testing.md)
@@ -11,15 +11,13 @@ FastVideo supports the following hardware platforms:
 ### Using pip

 ```bash
+# Create and activate a new conda environment
+conda create -n fastvideo python=3.12
+conda activate fastvideo
+
 pip install fastvideo
 ```

-### Using conda
-
-```bash
-conda install -c conda-forge fastvideo
-```
-
 ### From source

 ```bash
@@ -28,6 +26,12 @@ cd FastVideo
 pip install -e .
 ```

+Also optionally install flash-attn:
+
+```bash
+pip install flash-attn --no-build-isolation
+```
+
 ## Hardware Requirements

 - **NVIDIA GPUs**: CUDA 11.8+ with compute capability 7.0+
@@ -15,6 +15,12 @@ conda activate fastvideo
 pip install fastvideo
 ```

+Also optionally install flash-attn:
+
+```bash
+pip install flash-attn --no-build-isolation
+```
+
 ## Basic Usage

 ### Text-to-Video Generation
@@ -131,17 +131,41 @@ self.attn = DistributedAttention(

 ### Registering Models

-Register implemented modules in the model registry:
+Register implemented modules for auto‑discovery by adding `EntryClass` in each
+model module (the registry scans for it):

 ```python
-# In fastvideo/models/registry.py
-_TEXT_TO_VIDEO_DIT_MODELS = {
-    "YourTransformerModel": ("dits", "yourmodule", "YourTransformerClass"),
-}
+# In fastvideo/models/dits/your_module.py
+class YourTransformerModel(...):
+    ...

-_VAE_MODELS = {
-    "YourVAEModel": ("vaes", "yourvae", "YourVAEClass"),
-}
+# Entry point for model registry
+EntryClass = YourTransformerModel
+```
+
+```python
+# In fastvideo/models/vaes/your_vae.py
+class YourVAEModel(...):
+    ...
+
+# Entry point for model registry
+EntryClass = YourVAEModel
+```
+
+Register pipeline config + sampling defaults in the unified registry:
+
+```python
+# In fastvideo/registry.py
+register_configs(
+    sampling_param_cls=YourSamplingParam,
+    pipeline_config_cls=YourPipelineConfig,
+    hf_model_paths=[
+        "org/your-model-id",
+    ],
+    model_detectors=[
+        lambda path: "your-model" in path.lower(),
+    ],
+)
 ```

 ## Step 2: Directory Structure
@@ -106,6 +106,8 @@ If you encounter CUDA out of memory errors:
 - Enable memory optimization with `enable_model_cpu_offload`
 - Try a smaller model or use distilled versions
 - Use `num_gpus` > 1 if multiple GPUs are available
+- Try enabling FSDP inference with `use_fsdp_inference=True` (may slow down generation)
+- Try enabling DiT layerwise offload with `dit_layerwise_offload=True` (now only a few models support this, but may introduce less overhead than FSDP)

 ### Slow Generation

@@ -0,0 +1,140 @@
+# Offloading
+
+This page describes how to use offloading techniques for inference to reduce GPU memory usage while maintaining acceptable performance.
+
+## Default Behavior
+
+```python
+dit_cpu_offload: bool = True
+use_fsdp_inference: bool = False
+dit_layerwise_offload: bool = True
+text_encoder_cpu_offload: bool = True
+image_encoder_cpu_offload: bool = True
+vae_cpu_offload: bool = True
+pin_cpu_memory: bool = True
+```
+
+## Behavior Explanation
+
+!!! note
+    For CLI usage, replace underscores (`_`) with hyphens (`-`).
+
+### `use_fsdp_inference`
+
+Enables [FSDP](https://docs.pytorch.org/tutorials/intermediate/FSDP_tutorial.html) for inference. The model weights are sharded across multiple GPUs to reduce memory usage per GPU, and weights are broadcast to all GPUs layer by layer during inference.
+
+#### Performance Impact
+
+FSDP inference introduces negligible performance overhead due to weight prefetching. Performance overhead may be visible when GPU interconnect is slow (e.g., multiple consumer-level GPUs connected by slow PCIe without GPU P2P support).
+
+#### Usage Recommendation
+
+We recommend enabling this option when multiple GPUs are available.
+
+### `dit_cpu_offload`
+
+Enables CPU offloading for FSDP inference. When enabled, the model weights are offloaded to CPU memory, and the weight of each layer is moved to GPU memory only when that layer is being computed.
+
+#### Performance Impact
+
+The PyTorch FSDP implementation does not overlap computation and data transfer perfectly for inference, so enabling this option will harm performance.
+
+#### Usage Recommendation
+
+This option only takes effect when FSDP is enabled. For single GPU usage, we recommend using `dit_layerwise_offload` instead.
+
+### `dit_layerwise_offload`
+
+This option is similar to `dit_cpu_offload`, but with two key differences:
+
+1. It overlaps computation and PCIe data transfer
+2. It only works for single GPU inference
+
+#### Performance Impact
+
+This option introduces negligible performance overhead.
+
+#### Usage Recommendation
+
+We recommend enabling this option for single GPU usage. This option is not compatible with FSDP.
+
+### `text_encoder_cpu_offload`
+
+When enabled, the text encoder model weights are offloaded to CPU memory, and text encoding is computed on CPU.
+
+#### Performance Impact
+
+This option significantly slows down text encoding computation, but text encoding is usually not the bottleneck.
+
+#### Usage Recommendation
+
+We recommend enabling this option only when OOM happens.
+
+### `image_encoder_cpu_offload` and `vae_cpu_offload`
+
+When enabled, the weights are stored in CPU memory and moved to GPU memory when the corresponding module is being computed. After computation, the weights are moved back to CPU memory.
+
+#### Performance Impact
+
+These options introduce performance overhead due to PCIe data transfer.
+
+#### Usage Recommendation
+
+We recommend enabling these options when OOM happens.
+
+## General Recommendations
+
+### Single GPU Inference
+
+We recommend enabling `dit_layerwise_offload`. If OOM happens, also enable `image_encoder_cpu_offload` and `vae_cpu_offload`. If OOM still happens, consider enabling `text_encoder_cpu_offload`.
+
+### Multi-GPU Inference
+
+We recommend enabling `use_fsdp_inference` and disabling both `dit_layerwise_offload` and `dit_cpu_offload`. If OOM happens, consider enabling `text_encoder_cpu_offload`, `image_encoder_cpu_offload`, and `vae_cpu_offload`. If OOM still happens, consider enabling `dit_cpu_offload`.
+
+## Examples
+
+### Single GPU with Layerwise Offloading
+
+```python
+from fastvideo import VideoGenerator
+
+generator = VideoGenerator.from_pretrained(
+    "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
+    num_gpus=1,
+    # Recommended for single GPU
+    dit_layerwise_offload=True,
+    # Enable if OOM happens
+    vae_cpu_offload=True,
+    image_encoder_cpu_offload=True,
+    text_encoder_cpu_offload=True,
+    # Speeds up CPU-GPU transfer
+    pin_cpu_memory=True,
+)
+
+prompt = "A curious raccoon peers through a vibrant field of yellow sunflowers."
+video = generator.generate_video(prompt, output_path="output/", save_video=True)
+```
+
+### Multi-GPU with FSDP
+
+```python
+from fastvideo import VideoGenerator
+
+generator = VideoGenerator.from_pretrained(
+    "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
+    num_gpus=2,
+    # Recommended for multi-GPU
+    use_fsdp_inference=True,
+    dit_layerwise_offload=False,
+    dit_cpu_offload=False,
+    # Enable if OOM happens
+    vae_cpu_offload=True,
+    image_encoder_cpu_offload=True,
+    text_encoder_cpu_offload=True,
+    pin_cpu_memory=True,
+)
+
+prompt = "A majestic lion strides across the golden savanna."
+video = generator.generate_video(prompt, output_path="output/", save_video=True)
+```
@@ -103,7 +103,7 @@ python setup.py install  # or pip install -e .

 **`SAGE_ATTN_THREE`**

-[SageAttention 3](https://huggingface.co/jt-zhang/SageAttention3) is an advanced attention mechanism that leverages FP4 quantization and Blackwell GPU Tensor Cores for significant performance improvements.
+[SageAttention 3](https://github.com/thu-ml/SageAttention/tree/main/sageattention3_blackwell) is an advanced attention mechanism that leverages FP4 quantization and Blackwell GPU Tensor Cores for significant performance improvements.

 #### Hardware Requirements

@@ -113,11 +113,7 @@ python setup.py install  # or pip install -e .

 Note that Sage Attention 3 requires `python>=3.13`, `torch>=2.8.0`, `CUDA >=12.8`. If you are using `uv` and using `torch==2.8.0` make sure that `sentencepiece==0.2.1` in the pyproject.toml file.

-To use Sage Attention 3 in FastVideo, first get access to the SageAttention3 code, then move `sageattn/` and `setup.py` to the directory `fastvideo/attention/backends`, then install from using:
-
-```bash
-python setup.py install
-```
+To use Sage Attention 3 in FastVideo, follow the `README.md` in the linked repository to install the package from source.

 ## Teacache

@@ -0,0 +1,71 @@
+# LoRA Extraction and Merging
+
+Tools for extracting and merging LoRA adapters for FastVideo models.
+
+## Extract LoRA Adapter
+
+```bash
+python scripts/lora_extraction/extract_lora.py \
+  --base Wan-AI/Wan2.2-TI2V-5B-Diffusers \
+  --ft FastVideo/FastWan2.2-TI2V-5B-FullAttn-Diffusers \
+  --out adapter_r32.safetensors \
+  --rank 32
+```
+
+**Options:**
+
+- `--base`: Base model (HuggingFace ID or local path)
+- `--ft`: Fine-tuned model (HuggingFace ID or local path)
+- `--out`: Output adapter file
+- `--rank`: LoRA rank (16, 32, 64, 128)
+- `--full-rank`: Extract full-rank adapter (optional)
+
+## Merge Adapter
+
+```bash
+python scripts/lora_extraction/merge_lora.py \
+  --base Wan-AI/Wan2.2-TI2V-5B-Diffusers \
+  --adapter adapter_r32.safetensors \
+  --ft FastVideo/FastWan2.2-TI2V-5B-FullAttn-Diffusers \
+  --output merged_model
+```
+
+**Options:**
+
+- `--base`: Base model (HuggingFace ID or local path)
+- `--adapter`: LoRA adapter file (.safetensors)
+- `--ft`: Fine-tuned model (for configuration)
+- `--output`: Output directory
+
+## Validate Quality (Optional)
+
+```bash
+python scripts/lora_extraction/lora_inference_comparison.py \
+  --base merged_model \
+  --ft FastVideo/FastWan2.2-TI2V-5B-FullAttn-Diffusers \
+  --adapter NONE \
+  --output-dir results \
+  --prompt "A cat sitting on a windowsill" \
+  --seed 42 \
+  --height 480 \
+  --width 480 \
+  --num-frames 49 \
+  --num-inference-steps 32 \
+  --compute-ssim \
+  --compute-lpips
+```
+
+**Options:**
+
+- `--base`: Merged model or base model path
+- `--ft`: Fine-tuned model (reference)
+- `--adapter`: Path to adapter or NONE
+- `--output-dir`: Output directory
+- `--prompt`: Text prompt (default: "A cat sitting on a windowsill")
+- `--seed`: Random seed (default: 42)
+- `--height`: Video height (default: 480)
+- `--width`: Video width (default: 832)
+- `--num-frames`: Number of frames (default: 49)
+- `--num-inference-steps`: Inference steps (default: 32)
+- `--compute-ssim`: Compute SSIM metric
+- `--compute-lpips`: Compute LPIPS metric
@@ -12,7 +12,7 @@ def main():
        "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=False,
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -0,0 +1,51 @@
+# SPDX-License-Identifier: Apache-2.0
+from fastvideo import VideoGenerator
+from fastvideo.configs.sample import SamplingParam
+
+
+def main():
+    # Point this to your local diffusers model dir (or replace with a HF model ID).
+    model_path = "KyleShao/Cosmos-Predict2.5-2B-Diffusers"
+
+    generator = VideoGenerator.from_pretrained(
+        model_path,
+        num_gpus=1,
+        use_fsdp_inference=False,  # set True if GPU is out of memory
+        dit_cpu_offload=False,
+        vae_cpu_offload=False,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=True,
+    )
+
+    sampling_param = SamplingParam.from_pretrained(model_path)
+
+    # image2world example from official repo
+    image_path = "images/bus_terminal.jpg"
+
+    prompt = (
+        "A nighttime city bus terminal gradually shifts from stillness to subtle movement. "
+        "At first, multiple double-decker buses are parked under the glow of overhead lights, "
+        "with a central bus labeled '87D' facing forward and stationary. "
+        "As the video progresses, the bus in the middle moves ahead slowly, its headlights brightening the surrounding area "
+        "and casting reflections onto adjacent vehicles. "
+        "The motion creates space in the lineup, signaling activity within the otherwise quiet station. "
+        "It then comes to a smooth stop, resuming its position in line. "
+        "Overhead signage in Chinese characters remains illuminated, enhancing the vibrant, urban night scene."
+    )
+
+    generator.generate_video(
+        prompt,
+        sampling_param=sampling_param,
+        image_path=str(image_path),
+        num_cond_frames=1,
+        output_path="outputs_video/cosmos2_5_i2w.mp4",
+        save_video=True,
+    )
+
+    generator.shutdown()
+
+
+if __name__ == "__main__":
+    main()
+
+
@@ -0,0 +1,51 @@
+from fastvideo import VideoGenerator
+from fastvideo.configs.sample import SamplingParam
+
+
+def main():
+    # Point this to your local diffusers model dir (or replace with a HF model ID).
+    model_path = "KyleShao/Cosmos-Predict2.5-2B-Diffusers"
+
+    generator = VideoGenerator.from_pretrained(
+        model_path,
+        num_gpus=1,
+        use_fsdp_inference=False,  # set True if GPU is out of memory
+        dit_cpu_offload=False,
+        vae_cpu_offload=False,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=True,
+    )
+
+    # Load default sampling parameters (negative_prompt, resolution, steps, etc.)
+    sampling_param = SamplingParam.from_pretrained(model_path)
+
+    prompt = (
+        "A high-definition video captures the precision of robotic welding in an industrial setting. "
+        "The first frame showcases a robotic arm, equipped with a welding torch, positioned over a large metal structure. "
+        "The welding process is in full swing, with bright sparks and intense light illuminating the scene, "
+        "creating a vivid display of blue and white hues. "
+        "A significant amount of smoke billows around the welding area, partially obscuring the view but emphasizing the heat and activity. "
+        "The background reveals parts of the workshop environment, including a ventilation system and various pieces of machinery, "
+        "indicating a busy and functional industrial workspace. "
+        "As the video progresses, the robotic arm maintains its steady position, continuing the welding process and moving to its left. "
+        "The welding torch consistently emits sparks and light, and the smoke continues to rise, diffusing slightly as it moves upward. "
+        "The metal surface beneath the torch shows ongoing signs of heating and melting. "
+        "The scene retains its industrial ambiance, with the welding sparks and smoke dominating the visual field, "
+        "underscoring the ongoing nature of the welding operation."
+    )
+
+    generator.generate_video(
+        prompt,
+        sampling_param=sampling_param,
+        output_path="outputs_video/cosmos2_5_t2w.mp4",
+        save_video=True,
+    )
+
+    generator.shutdown()
+
+
+if __name__ == "__main__":
+    main()
+
+
+
@@ -0,0 +1,54 @@
+# SPDX-License-Identifier: Apache-2.0
+from fastvideo import VideoGenerator
+from fastvideo.configs.sample import SamplingParam
+
+
+def main():
+    # Point this to your local diffusers model dir (or replace with a HF model ID).
+    model_path = "KyleShao/Cosmos-Predict2.5-2B-Diffusers"
+
+    generator = VideoGenerator.from_pretrained(
+        model_path,
+        num_gpus=1,
+        use_fsdp_inference=False,  # set True if GPU is out of memory
+        dit_cpu_offload=False,
+        vae_cpu_offload=False,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=True,
+    )
+
+    sampling_param = SamplingParam.from_pretrained(model_path)
+
+    # video2world example from official repo
+    video_path = "videos/robot_pouring.mp4"
+
+    prompt = (
+        "A robotic arm, primarily white with black joints and cables, is shown in a clean, modern indoor setting with a white tabletop. "
+        "The arm, equipped with a gripper holding a small, light green pitcher, is positioned above a clear glass containing a reddish-brown liquid and a spoon. "
+        "The robotic arm is in the process of pouring a transparent liquid into the glass. "
+        "To the left of the pitcher, there is an opened jar with a similar reddish-brown substance visible through its transparent body. "
+        "In the background, a vase with white flowers and a brown couch are partially visible, adding to the contemporary ambiance. "
+        "The lighting is bright, casting soft shadows on the table. "
+        "The robotic arm's movements are smooth and controlled, demonstrating precision in its task. "
+        "As the video progresses, the robotic arm completes the pour, leaving the glass half-filled with the reddish-brown liquid. "
+        "The jar remains untouched throughout the sequence, and the spoon inside the glass remains stationary. "
+        "The other robotic arm on the right side also stays stationary throughout the video. "
+        "The final frame captures the robotic arm with the pitcher finishing the pour, with the glass now filled to a higher level, while the pitcher is slightly tilted but still held securely by the gripper."
+    )
+
+    generator.generate_video(
+        prompt,
+        sampling_param=sampling_param,
+        video_path=str(video_path),
+        num_cond_frames=1,
+        output_path="outputs_video/cosmos2_5_v2w.mp4",
+        save_video=True,
+    )
+
+    generator.shutdown()
+
+
+if __name__ == "__main__":
+    main()
+
+
@@ -14,7 +14,7 @@ def main():
        model_name,
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        # Adjust these offload parameters if you have < 32GB of VRAM
        text_encoder_cpu_offload=True,
        pin_cpu_memory=True, # set to false if low CPU RAM or hit obscure "CUDA error: Invalid argument"
@@ -12,7 +12,7 @@ def main():
        "hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True,
        vae_cpu_offload=True,
        text_encoder_cpu_offload=True,
@@ -39,4 +39,4 @@ def main():


 if __name__ == "__main__":
-    main()
+    main()
@@ -0,0 +1,43 @@
+from fastvideo import VideoGenerator
+import json
+# from fastvideo.configs.sample import SamplingParam
+
+OUTPUT_PATH = "video_samples_hy15_1080p"
+def main():
+    # FastVideo will automatically use the optimal default arguments for the
+    # model.
+    # If a local path is provided, FastVideo will make a best effort
+    # attempt to identify the optimal arguments.
+    generator = VideoGenerator.from_pretrained(
+        "weizhou03/HunyuanVideo-1.5-Diffusers-1080p-2SR", # 480p -> 720p -> 1080p
+        # or "weizhou03/HunyuanVideo-1.5-Diffusers-1080p" # 720p -> 1080p
+        # FastVideo will automatically handle distributed setup
+        num_gpus=1,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
+        dit_cpu_offload=True,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=True, # set to false if low CPU RAM or hit obscure "CUDA error: Invalid argument"
+        # image_encoder_cpu_offload=False,
+    )
+
+    prompt = (
+        "A curious raccoon peers through a vibrant field of yellow sunflowers, its eyes "
+        "wide with interest. The playful yet serene atmosphere is complemented by soft "
+        "natural light filtering through the petals. Mid-shot, warm and cheerful tones."
+    )
+
+    video = generator.generate_video(prompt, output_path=OUTPUT_PATH, save_video=True, negative_prompt="")
+
+    prompt2 = (
+        "A majestic lion strides across the golden savanna, its powerful frame "
+        "glistening under the warm afternoon sun. The tall grass ripples gently in "
+        "the breeze, enhancing the lion's commanding presence. The tone is vibrant, "
+        "embodying the raw energy of the wild. Low angle, steady tracking shot, "
+        "cinematic.")
+
+    video2 = generator.generate_video(prompt2, output_path=OUTPUT_PATH, save_video=True, negative_prompt="")
+
+
+if __name__ == "__main__":
+    main()
@@ -0,0 +1,66 @@
+from fastvideo import VideoGenerator
+from fastvideo.models.dits.hyworld.resolution_utils import get_resolution_from_image
+
+# Default prompt from HY-WorldPlay run.sh
+DEFAULT_PROMPT = 'A paved pathway leads towards a stone arch bridge spanning a calm body of water.  Lush green trees and foliage line the path and the far bank of the water. A traditional-style pavilion with a tiered, reddish-brown roof sits on the far shore. The water reflects the surrounding greenery and the sky.  The scene is bathed in soft, natural light, creating a tranquil and serene atmosphere. The pathway is composed of large, rectangular stones, and the bridge is constructed of light gray stone.  The overall composition emphasizes the peaceful and harmonious nature of the landscape.'
+DEFAULT_IMAGE = 'https://raw.githubusercontent.com/Tencent-Hunyuan/HY-WorldPlay/main/assets/img/test.png'
+
+OUTPUT_PATH = "video_samples_hyworld"
+def main():
+    import argparse
+
+    # pose: (a, w, s, d) - (15, 31)
+    # num_frames: (61, 125)
+    parser = argparse.ArgumentParser(description="HYWorld video generation with FastVideo")
+    parser.add_argument("--prompt", type=str, default=DEFAULT_PROMPT, help="Text prompt for video generation")
+    parser.add_argument("--image", type=str, default=DEFAULT_IMAGE, help="Path or URL to input image")
+    parser.add_argument("--pose", type=str, default='w-31', help="Pose string (e.g., 'a-31', 'w-31', 's-31', 'd-31')")
+    parser.add_argument("--output_path", type=str, default=OUTPUT_PATH, help="Output video path")
+    parser.add_argument("--num-frames", type=int, default=125, help="Number of frames")
+    parser.add_argument("--seed", type=int, default=1, help="Random seed")
+    parser.add_argument("--resolution", type=str, default="480p", help="Only support 480p for now")
+    args = parser.parse_args()
+
+    # Automatically determine resolution from input image
+    HEIGHT, WIDTH = get_resolution_from_image(args.image, args.resolution)
+    print(f"Image: {args.image}")
+    print(f"Pose: {args.pose}")
+    print(f"Resolution: {HEIGHT}x{WIDTH} (from {args.resolution} buckets)")
+    print(f"Num frames: {args.num_frames}")
+    print(f"Output path: {args.output_path}")
+
+    # Initialize generator
+    print("\nInitializing VideoGenerator for HYWorld...")
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/HY-WorldPlay-Bidirectional-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=True,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=True,
+        image_encoder_cpu_offload=True,
+    )
+
+    # Generate video
+    # The pose string is automatically converted to camera matrices by the pipeline
+    print("\nGenerating video...")
+    generator.generate_video(
+        prompt=args.prompt,
+        image_path=args.image,
+        pose=args.pose,  # Camera trajectory control
+        output_path=args.output_path,
+        save_video=True,
+        negative_prompt="",
+        num_frames=args.num_frames,
+        fps=24,
+        height=HEIGHT,
+        width=WIDTH,
+        seed=args.seed,
+    )
+
+    print(f"\nVideo saved to: {args.output_path}")
+
+
+if __name__ == "__main__":
+    main()
@@ -0,0 +1,210 @@
+"""
+LongCat Image-to-Video (I2V) Example Script
+
+This script demonstrates LongCat I2V inference using the FastVideo Python API.
+LongCat I2V takes an input image and generates a video from it.
+
+It runs both basic generation (50 steps) and distill+refine generation
+(16 steps distill + 50 steps refinement to 720p with BSA).
+
+Usage:
+    python examples/inference/basic/basic_longcat_i2v.py
+
+Note:
+    Refinement uses 768x768 dimensions where latent (48x48) is divisible by 8,
+    compatible with BSA chunks [4, 4, 8].
+"""
+
+import glob
+import os
+
+from fastvideo import VideoGenerator
+
+# Common prompts and settings matching the shell script examples
+PROMPT = (
+    "A woman sits at a wooden table by the window in a cozy café. She reaches out "
+    "with her right hand, picks up the white coffee cup from the saucer, and gently "
+    "brings it to her lips to take a sip. After drinking, she places the cup back on "
+    "the table and looks out the window, enjoying the peaceful atmosphere."
+)
+
+NEGATIVE_PROMPT = (
+    "Bright tones, overexposed, static, blurred details, subtitles, style, works, "
+    "paintings, images, static, overall gray, worst quality, low quality, JPEG compression "
+    "residue, ugly, incomplete, extra fingers, poorly drawn hands, poorly drawn faces, "
+    "deformed, disfigured, misshapen limbs, fused fingers, still picture, messy background, "
+    "three legs, many people in the background, walking backwards"
+)
+
+# Input image path
+IMAGE_PATH = "assets/girl.png"
+
+SEED = 42
+
+
+def basic_generation():
+    """
+    Run basic LongCat I2V generation (50 steps at 480p).
+    
+    This uses the full 50-step denoising process for highest quality.
+    """
+    print("=" * 60)
+    print("LongCat I2V: Basic Generation (50 steps, 480p)")
+    print("=" * 60)
+    
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-I2V-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
+        dit_cpu_offload=False,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=False,
+    )
+    
+    output_path = "outputs_video/longcat_i2v_basic"
+    
+    generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        image_path=IMAGE_PATH,
+        output_path=output_path,
+        save_video=True,
+        height=480,
+        width=480,  # Square
+        num_frames=93,
+        num_inference_steps=50,
+        fps=15,
+        guidance_scale=4.0,
+        seed=SEED,
+    )
+    
+    print(f"\nBasic generation complete! Video saved to: {output_path}")
+    generator.shutdown()
+
+
+def distill_refine_generation():
+    """
+    Run LongCat I2V with distill+refine pipeline (16 steps + refinement to 768p).
+    
+    This uses the distilled LoRA for fast 480p generation (16 steps),
+    then refines to 768p using the refinement LoRA with BSA enabled.
+    """
+    print("\n" + "=" * 60)
+    print("LongCat I2V: Distill + Refine Pipeline")
+    print("=" * 60)
+    
+    # Stage 1: Distilled generation (16 steps at 480p)
+    print("\n[Stage 1] Distilled generation (16 steps, 480p)")
+    print("-" * 40)
+    
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-I2V-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=False,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=False,
+        lora_path="FastVideo/LongCat-Video-T2V-Distilled-LoRA",
+        lora_nickname="distilled",
+    )
+    
+    distill_output_path = "outputs_video/longcat_i2v_distill"
+    
+    generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        image_path=IMAGE_PATH,
+        output_path=distill_output_path,
+        save_video=True,
+        height=480,
+        width=480,  # Square
+        num_frames=93,
+        num_inference_steps=16,
+        fps=15,
+        guidance_scale=1.0,
+        seed=SEED,
+    )
+    
+    print(f"Distilled generation complete! Video saved to: {distill_output_path}")
+    generator.shutdown()
+    
+    # Stage 2: Refinement (480p -> 768p)
+    print("\n[Stage 2] Refinement (480p -> 768p with BSA)")
+    print("-" * 40)
+    
+    # Find the actual saved video file from stage 1
+    video_files = glob.glob(os.path.join(distill_output_path, "*.mp4"))
+    if not video_files:
+        raise FileNotFoundError(f"No video file found in {distill_output_path}")
+    # Use the most recently created video file
+    distill_video_path = max(video_files, key=os.path.getmtime)
+    print(f"Using stage 1 video: {distill_video_path}")
+    
+    # Create a new generator with refinement LoRA and BSA enabled
+    # Note: Refinement uses the T2V model (not I2V) since it's upscaling the generated video
+    # For BSA [4, 4, 8]: latent must be divisible by 8
+    # 768x768: latent 48x48, 48%8=0 ✓
+    refine_generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-T2V-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=True,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=True,
+        bsa_sparsity=0.875,
+        bsa_chunk_q=[4, 4, 4],
+        bsa_chunk_k=[4, 4, 4],
+        lora_path="FastVideo/LongCat-Video-T2V-Refinement-LoRA",
+        lora_nickname="refinement",
+    )
+    
+    refine_output_path = "outputs_video/longcat_i2v_refine_720p"
+    
+    refine_generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        output_path=refine_output_path,
+        save_video=True,
+        refine_from=distill_video_path,
+        t_thresh=0.5,
+        spatial_refine_only=False,
+        num_cond_frames=0,
+        height=720,
+        width=720,
+        num_inference_steps=50,
+        fps=30,
+        guidance_scale=1.0,
+        seed=SEED,
+    )
+    
+    print(f"Refinement complete! Video saved to: {refine_output_path}")
+    refine_generator.shutdown()
+
+
+def main():
+    """Run both basic and distill+refine generation pipelines."""
+    print("\n" + "=" * 60)
+    print("LongCat Image-to-Video Example")
+    print("=" * 60 + "\n")
+    
+    # Run basic generation
+    basic_generation()
+    
+    # Run distill+refine pipeline
+    distill_refine_generation()
+    
+    print("\n" + "=" * 60)
+    print("All generations complete!")
+    print("=" * 60)
+
+
+if __name__ == "__main__":
+    main()
+
+
@@ -0,0 +1,198 @@
+"""
+LongCat Text-to-Video (T2V) Example Script
+
+This script demonstrates LongCat T2V inference using the FastVideo Python API.
+It runs both basic generation (50 steps) and distill+refine generation
+(16 steps distill + 50 steps refinement to 720p).
+
+Usage:
+    python examples/inference/basic/basic_longcat_t2v.py
+"""
+
+import glob
+import os
+
+from fastvideo import VideoGenerator
+
+# Common prompts and settings matching the shell script examples
+PROMPT = (
+    "In a realistic photography style, a white boy around seven or eight years old "
+    "sits on a park bench, wearing a light blue T-shirt, denim shorts, and white sneakers. "
+    "He holds an ice cream cone with vanilla and chocolate flavors, and beside him is a "
+    "medium-sized golden Labrador. Smiling, the boy offers the ice cream to the dog, "
+    "who eagerly licks it with its tongue. The sun is shining brightly, and the background "
+    "features a green lawn and several tall trees, creating a warm and loving scene."
+)
+
+NEGATIVE_PROMPT = (
+    "Bright tones, overexposed, static, blurred details, subtitles, style, works, "
+    "paintings, images, static, overall gray, worst quality, low quality, JPEG compression "
+    "residue, ugly, incomplete, extra fingers, poorly drawn hands, poorly drawn faces, "
+    "deformed, disfigured, misshapen limbs, fused fingers, still picture, messy background, "
+    "three legs, many people in the background, walking backwards"
+)
+
+SEED = 42
+
+
+def basic_generation():
+    """
+    Run basic LongCat T2V generation (50 steps at 480p).
+    
+    This uses the full 50-step denoising process for highest quality.
+    """
+    print("=" * 60)
+    print("LongCat T2V: Basic Generation (50 steps, 480p)")
+    print("=" * 60)
+    
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-T2V-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
+        dit_cpu_offload=False,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=False,
+    )
+    
+    output_path = "outputs_video/longcat_t2v_basic"
+    
+    generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        output_path=output_path,
+        save_video=True,
+        height=480,
+        width=832,
+        num_frames=93,
+        num_inference_steps=50,
+        fps=15,
+        guidance_scale=4.0,
+        seed=SEED,
+    )
+    
+    print(f"\nBasic generation complete! Video saved to: {output_path}")
+    generator.shutdown()
+
+
+def distill_refine_generation():
+    """
+    Run LongCat T2V with distill+refine pipeline (16 steps + refinement to 720p).
+    
+    This uses the distilled LoRA for fast 480p generation (16 steps),
+    then refines to 720p using the refinement LoRA with BSA enabled.
+    """
+    print("\n" + "=" * 60)
+    print("LongCat T2V: Distill + Refine Pipeline")
+    print("=" * 60)
+    
+    # Stage 1: Distilled generation (16 steps at 480p)
+    print("\n[Stage 1] Distilled generation (16 steps, 480p)")
+    print("-" * 40)
+    
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-T2V-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=False,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=False,
+        lora_path="FastVideo/LongCat-Video-T2V-Distilled-LoRA",
+        lora_nickname="distilled",
+    )
+    
+    distill_output_path = "outputs_video/longcat_t2v_distill"
+    
+    generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        output_path=distill_output_path,
+        save_video=True,
+        height=480,
+        width=832,
+        num_frames=93,
+        num_inference_steps=16,
+        fps=15,
+        guidance_scale=1.0,
+        seed=SEED,
+    )
+    
+    print(f"Distilled generation complete! Video saved to: {distill_output_path}")
+    generator.shutdown()
+    
+    # Stage 2: Refinement (480p -> 720p)
+    print("\n[Stage 2] Refinement (480p -> 720p with BSA)")
+    print("-" * 40)
+    
+    # Find the actual saved video file from stage 1
+    video_files = glob.glob(os.path.join(distill_output_path, "*.mp4"))
+    if not video_files:
+        raise FileNotFoundError(f"No video file found in {distill_output_path}")
+    # Use the most recently created video file
+    distill_video_path = max(video_files, key=os.path.getmtime)
+    print(f"Using stage 1 video: {distill_video_path}")
+    
+    # Create a new generator with refinement LoRA and BSA enabled
+    refine_generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-T2V-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=True,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=True,
+        bsa_sparsity=0.875,
+        bsa_chunk_q=[4, 4, 8],
+        bsa_chunk_k=[4, 4, 8],
+        lora_path="FastVideo/LongCat-Video-T2V-Refinement-LoRA",
+        lora_nickname="refinement",
+    )
+    
+    refine_output_path = "outputs_video/longcat_t2v_refine_720p"
+    
+    refine_generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        output_path=refine_output_path,
+        save_video=True,
+        refine_from=distill_video_path,
+        t_thresh=0.5,
+        spatial_refine_only=False,
+        num_cond_frames=0,
+        height=720,
+        width=1280,
+        num_inference_steps=50,
+        fps=30,
+        guidance_scale=1.0,
+        seed=SEED,
+    )
+    
+    print(f"Refinement complete! Video saved to: {refine_output_path}")
+    refine_generator.shutdown()
+
+
+def main():
+    """Run both basic and distill+refine generation pipelines."""
+    print("\n" + "=" * 60)
+    print("LongCat Text-to-Video Example")
+    print("=" * 60 + "\n")
+    
+    # Run basic generation
+    basic_generation()
+    
+    # Run distill+refine pipeline
+    distill_refine_generation()
+    
+    print("\n" + "=" * 60)
+    print("All generations complete!")
+    print("=" * 60)
+
+
+if __name__ == "__main__":
+    main()
+
+
@@ -0,0 +1,228 @@
+"""
+LongCat Video Continuation (VC) Example Script
+
+This script demonstrates LongCat VC inference using the FastVideo Python API.
+LongCat VC takes an input video and generates a continuation of it.
+
+It runs both basic generation (50 steps) and distill+refine generation
+(16 steps distill + 50 steps refinement to 720p).
+
+Usage:
+    python examples/inference/basic/basic_longcat_vc.py
+
+Prerequisites:
+    - Ensure the input video exists at assets/motorcycle.mp4
+      (or provide your own video)
+"""
+
+import glob
+import os
+
+from fastvideo import VideoGenerator
+
+# Common prompts and settings matching the shell script examples
+PROMPT = (
+    "A person rides a motorcycle along a long, straight road that stretches between "
+    "a body of water and a forested hillside. The rider steadily accelerates, keeping "
+    "the motorcycle centered between the guardrails, while the scenery passes by on "
+    "both sides. The video captures the journey from the rider's perspective, emphasizing "
+    "the sense of motion and adventure."
+)
+
+NEGATIVE_PROMPT = (
+    "Bright tones, overexposed, static, blurred details, subtitles, style, works, "
+    "paintings, images, static, overall gray, worst quality, low quality, JPEG compression "
+    "residue, ugly, incomplete, extra fingers, poorly drawn hands, poorly drawn faces, "
+    "deformed, disfigured, misshapen limbs, fused fingers, still picture, messy background, "
+    "three legs, many people in the background, walking backwards"
+)
+
+# Input video path
+VIDEO_PATH = "assets/motorcycle.mp4"
+
+# Number of conditioning frames from the input video
+NUM_COND_FRAMES = 13
+
+SEED = 42
+
+
+def basic_generation():
+    """
+    Run basic LongCat VC generation (50 steps at 480p).
+    
+    This uses the full 50-step denoising process for highest quality.
+    """
+    print("=" * 60)
+    print("LongCat VC: Basic Generation (50 steps, 480p)")
+    print("=" * 60)
+    
+    # Check if video exists
+    if not os.path.exists(VIDEO_PATH):
+        raise FileNotFoundError(
+            f"Video not found at {VIDEO_PATH}. "
+            "Please provide a valid video path."
+        )
+    
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-VC-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
+        dit_cpu_offload=False,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=False,
+    )
+    
+    output_path = "outputs_video/longcat_vc_basic"
+    
+    generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        video_path=VIDEO_PATH,
+        num_cond_frames=NUM_COND_FRAMES,
+        output_path=output_path,
+        save_video=True,
+        height=480,
+        width=832,
+        num_frames=93,
+        num_inference_steps=50,
+        fps=15,
+        guidance_scale=4.0,
+        seed=SEED,
+    )
+    
+    print(f"\nBasic generation complete! Video saved to: {output_path}")
+    generator.shutdown()
+
+
+def distill_refine_generation():
+    """
+    Run LongCat VC with distill+refine pipeline (16 steps + refinement to 720p).
+    
+    This uses the distilled LoRA for fast 480p generation (16 steps),
+    then refines to 720p using the refinement LoRA with BSA enabled.
+    """
+    print("\n" + "=" * 60)
+    print("LongCat VC: Distill + Refine Pipeline")
+    print("=" * 60)
+    
+    # Check if video exists
+    if not os.path.exists(VIDEO_PATH):
+        raise FileNotFoundError(
+            f"Video not found at {VIDEO_PATH}. "
+            "Please provide a valid video path."
+        )
+    
+    # Stage 1: Distilled generation (16 steps at 480p)
+    print("\n[Stage 1] Distilled generation (16 steps, 480p)")
+    print("-" * 40)
+    
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-VC-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=False,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=False,
+        lora_path="FastVideo/LongCat-Video-T2V-Distilled-LoRA",
+        lora_nickname="distilled",
+    )
+    
+    distill_output_path = "outputs_video/longcat_vc_distill"
+    
+    generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        video_path=VIDEO_PATH,
+        num_cond_frames=NUM_COND_FRAMES,
+        output_path=distill_output_path,
+        save_video=True,
+        height=480,
+        width=832,
+        num_frames=93,
+        num_inference_steps=16,
+        fps=15,
+        guidance_scale=1.0,
+        seed=SEED,
+    )
+    
+    print(f"Distilled generation complete! Video saved to: {distill_output_path}")
+    generator.shutdown()
+    
+    # Stage 2: Refinement (480p -> 720p)
+    print("\n[Stage 2] Refinement (480p -> 720p with BSA)")
+    print("-" * 40)
+    
+    # Find the actual saved video file from stage 1
+    video_files = glob.glob(os.path.join(distill_output_path, "*.mp4"))
+    if not video_files:
+        raise FileNotFoundError(f"No video file found in {distill_output_path}")
+    # Use the most recently created video file
+    distill_video_path = max(video_files, key=os.path.getmtime)
+    print(f"Using stage 1 video: {distill_video_path}")
+    
+    # Create a new generator with refinement LoRA and BSA enabled
+    # Note: Refinement uses the T2V model (not VC) since it's upscaling the generated video
+    refine_generator = VideoGenerator.from_pretrained(
+        "FastVideo/LongCat-Video-T2V-Diffusers",
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=True,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+        enable_bsa=True,
+        bsa_sparsity=0.875,
+        bsa_chunk_q=[4, 4, 8],
+        bsa_chunk_k=[4, 4, 8],
+        lora_path="FastVideo/LongCat-Video-T2V-Refinement-LoRA",
+        lora_nickname="refinement",
+    )
+    
+    refine_output_path = "outputs_video/longcat_vc_refine_720p"
+    
+    refine_generator.generate_video(
+        prompt=PROMPT,
+        negative_prompt=NEGATIVE_PROMPT,
+        output_path=refine_output_path,
+        save_video=True,
+        refine_from=distill_video_path,
+        t_thresh=0.5,
+        spatial_refine_only=False,
+        num_cond_frames=0,  # For refinement, no conditioning frames
+        height=720,
+        width=1280,
+        num_inference_steps=50,
+        fps=30,
+        guidance_scale=1.0,
+        seed=SEED,
+    )
+    
+    print(f"Refinement complete! Video saved to: {refine_output_path}")
+    refine_generator.shutdown()
+
+
+def main():
+    """Run both basic and distill+refine generation pipelines."""
+    print("\n" + "=" * 60)
+    print("LongCat Video Continuation Example")
+    print("=" * 60 + "\n")
+    
+    # Run basic generation
+    basic_generation()
+    
+    # Run distill+refine pipeline
+    distill_refine_generation()
+    
+    print("\n" + "=" * 60)
+    print("All generations complete!")
+    print("=" * 60)
+
+
+if __name__ == "__main__":
+    main()
+
+
@@ -0,0 +1,34 @@
+from fastvideo import VideoGenerator
+
+PROMPT = (
+    "A warm sunny backyard. The camera starts in a tight cinematic close-up "
+    "of a woman and a man in their 30s, facing each other with serious "
+    "expressions. The woman, emotional and dramatic, says softly, \"That's "
+    "it... Dad's lost it. And we've lost Dad.\" The man exhales, slightly "
+    "annoyed: \"Stop being so dramatic, Jess.\" A beat. He glances aside, "
+    "then mutters defensively, \"He's just having fun.\" The camera slowly "
+    "pans right, revealing the grandfather in the garden wearing enormous "
+    "butterfly wings, waving his arms in the air like he's trying to take "
+    "off. He shouts, \"Wheeeew!\" as he flaps his wings with full commitment. "
+    "The woman covers her face, on the verge of tears. The tone is deadpan, "
+    "absurd, and quietly tragic."
+)
+
+
+def main() -> None:
+    generator = VideoGenerator.from_pretrained(
+        "FastVideo/LTX2-Distilled-Diffusers",
+        num_gpus=1,
+    )
+
+    output_path = "outputs_video/ltx2_basic/output_ltx2_distilled_t2v.mp4"
+    generator.generate_video(
+        prompt=PROMPT,
+        output_path=output_path,
+        save_video=True,
+    )
+    generator.shutdown()
+
+
+if __name__ == "__main__":
+    main()
@@ -1,6 +1,5 @@
 from fastvideo import VideoGenerator
-from fastvideo.configs.pipelines.wan import MatrixGameI2V480PConfig
-from fastvideo.models.dits.matrix_game.utils import create_action_presets
+from fastvideo.models.dits.matrixgame.utils import create_action_presets

 import torch

@@ -43,7 +42,7 @@ def main():
        config["model_path"],
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -1,5 +1,5 @@
 from fastvideo.entrypoints.streaming_generator import StreamingVideoGenerator
-from fastvideo.models.dits.matrix_game.utils import get_current_action_async, expand_action_to_frames
+from fastvideo.models.dits.matrixgame.utils import get_current_action_async, expand_action_to_frames

 import torch
 import asyncio
@@ -46,7 +46,7 @@ async def main():
        config["model_path"],
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -13,7 +13,7 @@ def main():
        model_name,
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        text_encoder_cpu_offload=False,
        dit_cpu_offload=False,
    )
@@ -14,7 +14,7 @@ def main():
        "FastVideo/SFWan2.2-I2V-A14B-Preview-Diffusers",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        dit_precision="fp32",
        vae_cpu_offload=False,
@@ -14,7 +14,7 @@ def main():
        "rand0nmr/SFWan2.2-T2V-A14B-Diffusers",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -15,12 +15,14 @@ def main() -> None:
        "loayrashid/TurboWan2.1-T2V-1.3B-Diffusers",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        # TurboDiffusion uses a custom pipeline with RCM scheduler
-        override_pipeline_cls_name="TurboDiffusionPipeline",
+        use_fsdp_inference=False, # set to True if GPU is out of memory
+
+        # set to false if using RTX 4090 
+        # pin_cpu_memory=False,
    )

    # Generate videos with the same simple API, regardless of GPU count
-    # TurboDiffusion uses guidance_scale=1.0 (no CFG) and only 4 steps
+    # TurboDiffusion defaults: guidance_scale=1.0 and num_inference_steps=4 (from config)
    prompt = (
        "A curious raccoon peers through a vibrant field of yellow sunflowers, its eyes "
        "wide with interest. The playful yet serene atmosphere is complemented by soft "
@@ -30,9 +32,7 @@ def main() -> None:
        prompt,
        output_path=OUTPUT_PATH,
        save_video=True,
-        num_inference_steps=4,
        seed=42,
-        guidance_scale=1.0,
    )

    # Generate another video with a different prompt, without reloading the model!
@@ -47,9 +47,7 @@ def main() -> None:
        prompt2,
        output_path=OUTPUT_PATH,
        save_video=True,
-        num_inference_steps=4,
        seed=42,
-        guidance_scale=1.0,
    )


@@ -15,8 +15,6 @@ def main() -> None:
        "loayrashid/TurboWan2.1-T2V-14B-Diffusers",
        # 14B model needs more GPUs
        num_gpus=2,
-        # TurboDiffusion uses a custom pipeline with RCM scheduler
-        override_pipeline_cls_name="TurboDiffusionPipeline",
    )

    prompt = (
@@ -28,9 +26,7 @@ def main() -> None:
        prompt,
        output_path=OUTPUT_PATH,
        save_video=True,
-        num_inference_steps=4,
        seed=42,
-        guidance_scale=1.0,
    )

    # Generate another video with a different prompt, without reloading the model!
@@ -45,9 +41,7 @@ def main() -> None:
        prompt2,
        output_path=OUTPUT_PATH,
        save_video=True,
-        num_inference_steps=4,
        seed=42,
-        guidance_scale=1.0,
    )


@@ -0,0 +1,37 @@
+import os
+
+# Set SLA attention backend BEFORE fastvideo imports
+os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "SLA_ATTN"
+
+from fastvideo import VideoGenerator
+
+# Use local model path
+MODEL_PATH = "loayrashid/TurboWan2.2-I2V-A14B-Diffusers"
+OUTPUT_PATH = "video_samples_turbodiffusion_i2v"
+
+
+def main() -> None:
+    # TurboDiffusion I2V: 1-4 step image-to-video generation
+    generator = VideoGenerator.from_pretrained(
+        MODEL_PATH,
+        num_gpus=2,
+    )
+
+    # Example prompt and image for I2V
+    prompt = ("Summer beach vacation style, a white cat wearing sunglasses sits on a surfboard. The fluffy-furred feline gazes directly at the camera with a relaxed expression. Blurred beach scenery forms the background featuring crystal-clear waters, distant green hills, and a blue sky dotted with white clouds. The cat assumes a naturally relaxed posture, as if savoring the sea breeze and warm sunlight. A close-up shot highlights the feline's intricate details and the refreshing atmosphere of the seaside.")
+
+    # Use an example image path
+    image_path = "https://huggingface.co/datasets/YiYiXu/testing-images/resolve/main/wan_i2v_input.JPG"
+
+
+    video = generator.generate_video(
+        prompt,
+        image_path=image_path,
+        output_path=OUTPUT_PATH,
+        save_video=True,
+        seed=42,
+    )
+
+
+if __name__ == "__main__":
+    main()
@@ -12,7 +12,7 @@ def main():
        "Wan-AI/Wan2.2-T2V-A14B-Diffusers",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory 
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -14,7 +14,7 @@ def main():
        # "alibaba-pai/Wan2.2-Fun-A14B-Control",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -12,7 +12,7 @@ def main():
        "Wan-AI/Wan2.2-I2V-A14B-Diffusers",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -11,7 +11,7 @@ def main():
        model_name,
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=True,
+        use_fsdp_inference=False, # set to True if GPU is out of memory
        dit_cpu_offload=True,
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -10,7 +10,7 @@ from fastapi import FastAPI, Request, HTTPException
 from fastapi.responses import HTMLResponse, FileResponse

 from fastvideo.entrypoints.streaming_generator import StreamingVideoGenerator
-from fastvideo.models.dits.matrix_game.utils import expand_action_to_frames
+from fastvideo.models.dits.matrixgame.utils import expand_action_to_frames


 VARIANT_CONFIG = {
@@ -4,7 +4,7 @@ export WANDB_BASE_URL="https://api.wandb.ai"
 export WANDB_MODE=online
 export TOKENIZERS_PARALLELISM=false

-MODEL_PATH="Wan-AI/Wan2.1-T2V-1.3B-Diffusers"
+MODEL_PATH="wlsaidhi/SFWan2.1-T2V-1.3B-Diffusers"
 DATA_DIR="data/crush-smol_processed_t2v_1_3b_ode_init/"
 VALIDATION_DATASET_FILE="$(dirname "$0")/validation.json"
 NUM_GPUS=1
@@ -14,7 +14,6 @@ NUM_GPUS=1
 training_args=(
  --tracker_project_name "wan_ode_init"
  --output_dir "wan_ode_init_crush_smol"
-  --override_transformer_cls_name "CausalWanTransformer3DModel"
  --wandb_run_name "wan_ode_init_crush_smol"
  --max_train_steps 6000
  --train_batch_size 1
@@ -34,7 +33,7 @@ parallel_args=(
  --sp_size 1
  --tp_size 1
  --hsdp_replicate_dim 1
-  --hsdp_shard_dim 1
+  --hsdp_shard_dim $NUM_GPUS
 )

 # Model arguments
@@ -51,20 +50,17 @@ dataset_args=(

 # Validation arguments
 validation_args=(
-  --log_validation
-  --validation_dataset_file "$VALIDATION_DATASET_FILE"
-  --validation_steps 50
-  --validation_sampling_steps "50"
-  --validation_guidance_scale "6.0"
+  --log-visualization
+  --visualization-steps 100
 )

 # Optimizer arguments
 optimizer_args=(
-  --learning_rate 6e-6
+  --learning_rate 1e-5
  --mixed_precision "bf16"
  --weight_only_checkpointing_steps 1000
  --training_state_checkpointing_steps 1000
-  --weight_decay 1e-4
+  --weight_decay 0.01
  --max_grad_norm 1.0
 )

@@ -0,0 +1,93 @@
+#!/bin/bash
+
+export WANDB_BASE_URL="https://api.wandb.ai"
+export WANDB_MODE=online
+export TOKENIZERS_PARALLELISM=false
+# export FASTVIDEO_ATTENTION_BACKEND=TORCH_SDPA
+
+MODEL_PATH="FastVideo/Matrix-Game-2.0-Foundation-Diffusers"
+DATA_DIR="footsies-dataset/preprocessed/combined_parquet_dataset"
+VALIDATION_DATASET_FILE="$(dirname "$0")/validation.json"
+NUM_GPUS=8
+# export CUDA_VISIBLE_DEVICES=4,5
+# IP=[MASTER NODE IP]
+
+# Training arguments
+training_args=(
+  --tracker_project_name "matrixgame_finetune"
+  --output_dir "checkpoints/matrixgame_finetune"
+  --max_train_steps 1500
+  --train_batch_size 1
+  --train_sp_batch_size 1
+  --gradient_accumulation_steps 4
+  --num_latent_t 20
+  --num_height 352
+  --num_width 640
+  --num_frames 77
+  --enable_gradient_checkpointing_type "full"
+)
+
+# Parallel arguments
+parallel_args=(
+  --num_gpus $NUM_GPUS
+  --sp_size 2
+  --tp_size 1
+  --hsdp_replicate_dim 4
+  --hsdp_shard_dim 2
+)
+
+# Model arguments
+model_args=(
+  --model_path $MODEL_PATH
+  --pretrained_model_name_or_path $MODEL_PATH
+)
+
+# Dataset arguments
+dataset_args=(
+  --data_path "$DATA_DIR"
+  --dataloader_num_workers 1
+)
+
+# Validation arguments
+validation_args=(
+  --log_validation
+  --validation_dataset_file "$VALIDATION_DATASET_FILE"
+  --validation_steps 100
+  --validation_sampling_steps "40"
+  --validation_guidance_scale "6.0"
+)
+
+# Optimizer arguments
+optimizer_args=(
+  --learning_rate 2e-5
+  --mixed_precision "bf16"
+  --weight_only_checkpointing_steps 100
+  --training_state_checkpointing_steps 100
+  --weight_decay 1e-4
+  --max_grad_norm 1.0
+)
+
+# Miscellaneous arguments
+miscellaneous_args=(
+  --inference_mode False
+  --checkpoints_total_limit 3
+  --training_cfg_rate 0.1
+  --multi_phased_distill_schedule "4000-1"
+  --not_apply_cfg_solver
+  --dit_precision "fp32"
+  --num_euler_timesteps 50
+  --ema_start_step 0
+)
+
+# If you do not have 32 GPUs and to fit in memory, you can: 1. increase sp_size. 2. reduce num_latent_t
+torchrun \
+  --nnodes 1 \
+  --nproc_per_node $NUM_GPUS \
+    fastvideo/training/matrixgame_training_pipeline.py \
+    "${parallel_args[@]}" \
+    "${model_args[@]}" \
+    "${dataset_args[@]}" \
+    "${training_args[@]}" \
+    "${optimizer_args[@]}" \
+    "${validation_args[@]}" \
+    "${miscellaneous_args[@]}"
@@ -0,0 +1,27 @@
+#!/bin/bash
+
+GPU_NUM=1 # 2,4,8
+MODEL_PATH="Matrix-Game-2.0-Foundation-Diffusers"
+DATA_MERGE_PATH="footsies-dataset/merge.txt"
+OUTPUT_DIR="footsies-dataset/preprocessed/"
+
+# export CUDA_VISIBLE_DEVICES=0
+export MASTER_ADDR=localhost
+export MASTER_PORT=29500
+export RANK=0
+export WORLD_SIZE=1
+
+python fastvideo/pipelines/preprocess/v1_preprocess.py \
+    --model_path $MODEL_PATH \
+    --data_merge_path $DATA_MERGE_PATH \
+    --preprocess_video_batch_size 4 \
+    --seed 42 \
+    --max_height 352 \
+    --max_width 640 \
+    --num_frames 77 \
+    --dataloader_num_workers 0 \
+    --output_dir=$OUTPUT_DIR \
+    --samples_per_file 4 \
+    --train_fps 25 \
+    --flush_frequency 4 \
+    --preprocess_task matrixgame
@@ -10,6 +10,8 @@ if(GPU_BACKEND STREQUAL "ROCM")
    enable_language(HIP)
 else()
    enable_language(CUDA)
+    # Ensure CUDA toolkit targets (CUDA::cudart, CUDA::cuda_driver, etc.) are available.
+    find_package(CUDAToolkit REQUIRED)
 endif()

 # Import common utils if needed, but we keep it simple for now
@@ -153,6 +155,30 @@ if(BUILD_CXX_KERNELS)
        $<$<COMPILE_LANGUAGE:CUDA>:${CUDA_FLAGS}>
    )

+    # Link against Torch libraries to avoid undefined symbols at import time
+    # (e.g., torch::autograd vtables) when loading the extension module.
+    target_link_libraries(fastvideo_kernel_ops PRIVATE ${TORCH_LIBRARIES})
+
+    # Also link against libtorch_python to satisfy Python-binding symbols
+    # (e.g., torch::PyWarningHandler) required by torch/extension.h.
+    execute_process(
+        COMMAND "${Python_EXECUTABLE}" -c "import torch; from pathlib import Path; p=Path(torch.__file__).parent/'lib'; m=sorted(p.glob('libtorch_python*')); print(str(m[0]) if m else '')"
+        OUTPUT_VARIABLE TORCH_PYTHON_LIBRARY_PATH
+        OUTPUT_STRIP_TRAILING_WHITESPACE
+        ERROR_QUIET
+    )
+    if(TORCH_PYTHON_LIBRARY_PATH)
+        message(STATUS "TORCH_PYTHON_LIBRARY_PATH: ${TORCH_PYTHON_LIBRARY_PATH}")
+        target_link_libraries(fastvideo_kernel_ops PRIVATE "${TORCH_PYTHON_LIBRARY_PATH}")
+    else()
+        message(WARNING "Could not locate libtorch_python; fastvideo_kernel_ops may fail to import.")
+    endif()
+
+    # Link CUDA runtime + driver explicitly (fixes missing symbols like cuGetErrorString at import time)
+    if(NOT GPU_BACKEND STREQUAL "ROCM")
+        target_link_libraries(fastvideo_kernel_ops PRIVATE CUDA::cudart CUDA::cuda_driver)
+    endif()
+
    # We install it to fastvideo_kernel/_C so we can load it to register the ops
    install(TARGETS fastvideo_kernel_ops LIBRARY DESTINATION fastvideo_kernel/_C)
 endif()
@@ -34,12 +34,23 @@ from fastvideo_kernel import sliding_tile_attention, video_sparse_attn, moba_att
 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, topk=5)
+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
+
+### 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
+```
+
 ### TurboDiffusion Kernels

 This package also includes kernels from [TurboDiffusion](https://github.com/thu-ml/TurboDiffusion), including INT8 GEMM, Quantization, RMSNorm and LayerNorm.
@@ -0,0 +1,166 @@
+#!/usr/bin/env python3
+"""
+Benchmark VSA *wrapper* performance (forward + backward) and report TFLOPs.
+
+This script benchmarks the autograd-enabled wrapper:
+  - fastvideo_kernel.block_sparse_attn.block_sparse_attn
+
+So measured time includes wrapper overhead (map->index conversion, dispatch) plus kernel time.
+"""
+
+from __future__ import annotations
+
+import argparse
+import os
+import random
+from typing import Tuple, Callable
+
+import numpy as np
+import torch
+
+try:
+    from triton.testing import do_bench
+except Exception as e:  # pragma: no cover
+    raise ImportError("This benchmark requires triton (for triton.testing.do_bench).") from e
+
+
+BLOCK_M = 64
+BLOCK_N = 64
+
+
+def set_seed(seed: int = 42) -> None:
+    random.seed(seed)
+    np.random.seed(seed)
+    torch.manual_seed(seed)
+    torch.cuda.manual_seed(seed)
+    torch.cuda.manual_seed_all(seed)
+
+
+def parse_arguments() -> argparse.Namespace:
+    p = argparse.ArgumentParser(description="Benchmark FastVideo VSA block-sparse attention")
+    p.add_argument("--batch_size", type=int, default=1)
+    p.add_argument("--num_heads", type=int, default=12)
+    p.add_argument("--head_dim", type=int, default=128, choices=[64, 128])
+    p.add_argument("--topk", type=int, default=None, help="KV blocks per Q block (default: ~90%% sparsity)")
+    p.add_argument("--q_seq_lens", type=int, nargs="+", default=[49152], help="Q sequence lengths (must be /64)")
+    p.add_argument("--kv_seq_lens", type=int, nargs="+", default=None, help="KV sequence lengths (defaults to q_seq_len)")
+    p.add_argument("--warmup", type=int, default=5)
+    p.add_argument("--rep", type=int, default=20)
+    p.add_argument("--seed", type=int, default=42)
+    p.add_argument("--dtype", type=str, default="bf16", choices=["bf16", "fp16"])
+    p.add_argument("--force_triton", action="store_true", help="Force wrapper to use Triton path (if supported by shapes).")
+    return p.parse_args()
+
+
+def create_qkv(batch: int, heads: int, q_len: int, kv_len: int, d: int, dtype: torch.dtype) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    q = torch.randn(batch, heads, q_len, d, dtype=dtype, device="cuda")
+    k = torch.randn(batch, heads, kv_len, d, dtype=dtype, device="cuda")
+    v = torch.randn(batch, heads, kv_len, d, dtype=dtype, device="cuda")
+    return q, k, v
+
+
+def make_block_map(bs: int, h: int, num_q_blocks: int, num_kv_blocks: int, topk: int) -> torch.Tensor:
+    # block_map: [bs, h, num_q_blocks, num_kv_blocks] bool
+    scores = torch.rand(bs, h, num_q_blocks, num_kv_blocks, device="cuda")
+    topk = min(max(1, topk), num_kv_blocks)
+    idx = torch.topk(scores, topk, dim=-1).indices
+    block_map = torch.zeros(bs, h, num_q_blocks, num_kv_blocks, dtype=torch.bool, device="cuda")
+    block_map.scatter_(-1, idx, True)
+    return block_map
+
+
+def flops_sparse_attention(bs: int, h: int, d: int, q_len: int, topk_blocks: int, block_n: int) -> float:
+    # Approx: QK^T + PV, each is ~2*bs*h*q_len*(topk_blocks*block_n)*d
+    return 4.0 * bs * h * d * q_len * (topk_blocks * block_n)
+
+
+def bench_ms(fn: Callable[[], object], warmup: int, rep: int) -> float:
+    return do_bench(fn, warmup=warmup, rep=rep, quantiles=None)
+
+
+def main() -> None:
+    args = parse_arguments()
+    set_seed(args.seed)
+
+    dtype = torch.bfloat16 if args.dtype == "bf16" else torch.float16
+
+    if args.force_triton:
+        os.environ["FASTVIDEO_KERNEL_VSA_FORCE_TRITON"] = "1"
+
+    from fastvideo_kernel.block_sparse_attn import block_sparse_attn
+
+    bs, h, d = args.batch_size, args.num_heads, args.head_dim
+    kv_seq_lens = args.kv_seq_lens
+    if kv_seq_lens is None:
+        kv_seq_lens = args.q_seq_lens
+    if len(kv_seq_lens) != len(args.q_seq_lens):
+        raise ValueError("kv_seq_lens must have the same number of entries as q_seq_lens (or be omitted).")
+
+    print("VSA Block-Sparse Attention Benchmark (WRAPPER)")
+    print(f"device: {torch.cuda.get_device_name(0)}")
+    print(f"batch={bs}, heads={h}, head_dim={d}, dtype={args.dtype}")
+    print(f"BLOCK_M={BLOCK_M}, BLOCK_N={BLOCK_N}")
+    print("NOTE: timings include wrapper overhead (map->index + dispatch).")
+    if args.force_triton:
+        print("dispatch: forced Triton (FASTVIDEO_KERNEL_VSA_FORCE_TRITON=1)")
+    else:
+        print("dispatch: SM90 if available, else Triton")
+
+    for q_len, kv_len in zip(args.q_seq_lens, kv_seq_lens):
+        if q_len % BLOCK_M != 0 or kv_len % BLOCK_N != 0:
+            print(f"[skip] q_len={q_len}, kv_len={kv_len} must be divisible by 64")
+            continue
+
+        num_q_blocks = q_len // BLOCK_M
+        num_kv_blocks = kv_len // BLOCK_N
+        topk = args.topk if args.topk is not None else max(1, num_kv_blocks // 10)
+        topk = min(topk, num_kv_blocks)
+
+        print("\n" + "=" * 80)
+        print(f"q_len={q_len}, kv_len={kv_len}, num_q_blocks={num_q_blocks}, num_kv_blocks={num_kv_blocks}, topk={topk}")
+
+        q, k, v = create_qkv(bs, h, q_len, kv_len, d, dtype)
+        block_map = make_block_map(bs, h, num_q_blocks, num_kv_blocks, topk)
+
+        # Variable block sizes: default full blocks (64 tokens per KV block)
+        variable_block_sizes = torch.full((num_kv_blocks,), BLOCK_N, dtype=torch.int32, device="cuda")
+
+        def _fwd():
+            return block_sparse_attn(q, k, v, block_map, variable_block_sizes)
+
+        fwd_ms = bench_ms(_fwd, warmup=args.warmup, rep=args.rep)
+
+        # Backward benchmark (wrapper autograd). We build the graph once, then repeatedly run backward
+        # on the retained graph so bwd timing excludes the forward compute.
+        q_ = q.detach().requires_grad_(True)
+        k_ = k.detach().requires_grad_(True)
+        v_ = v.detach().requires_grad_(True)
+        o_, _aux_ = block_sparse_attn(q_, k_, v_, block_map, variable_block_sizes)
+        og = torch.randn_like(o_)
+        loss = (o_ * og).sum()
+
+        for _ in range(max(1, args.warmup // 2)):
+            torch.autograd.grad(loss, (q_, k_, v_), retain_graph=True)
+        torch.cuda.synchronize()
+
+        bwd_ms = bench_ms(
+            lambda: torch.autograd.grad(loss, (q_, k_, v_), retain_graph=True),
+            warmup=0,
+            rep=max(5, args.rep // 2),
+        )
+
+        flops = flops_sparse_attention(bs, h, d, q_len, topk, BLOCK_N)
+        fwd_tflops = flops / fwd_ms * 1e-12 * 1e3
+        # Rough backward multiplier (attention backward typically ~2-3x forward)
+        bwd_tflops = (2.5 * flops) / bwd_ms * 1e-12 * 1e3
+
+        print(f"fwd(wrapper): {fwd_ms:.3f} ms  | {fwd_tflops:.2f} TFLOPs (approx)")
+        print(f"bwd(wrapper): {bwd_ms:.3f} ms  | {bwd_tflops:.2f} TFLOPs (approx)")
+
+
+if __name__ == "__main__":
+    if not torch.cuda.is_available():
+        raise RuntimeError("CUDA is required for this benchmark.")
+    main()
+
+
@@ -639,7 +639,8 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {

    // store kq and vq

-    // ensuring all writes are finished
+    // ! the following two line seems unnecessary.
+    // tma::store_async_wait(); // ensure qg is finished
    __syncthreads();

    warpgroup::store(kg_smem[0], kg_reg);
@@ -660,145 +661,6 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {
    tma::store_async_wait(); 
 }

-
-template<int D>
-void block_sparse_attention_forward_impl(
-    bf16* d_q, bf16* d_k, bf16* d_v, float* d_l, bf16* d_o,
-    int batch, int qo_heads, int kv_heads, int seq_len, int hr,
-    int max_kv_blocks_per_q, 
-    int32_t* q2k_block_sparse_index_ptr,
-    int32_t* q2k_block_sparse_num_ptr,
-    int32_t* block_size_ptr,
-    cudaStream_t stream
-) {
-    using K = fwd_attend_ker_tile_dims<D>;
-    using q_tile    =         st_bf<K::qo_height, K::tile_width>;
-    using k_tile    =         st_bf<K::kv_height, K::tile_width>;
-    using v_tile    =         st_bf<K::kv_height, K::tile_width>;
-    using l_col_vec = col_vec<st_fl<K::qo_height, K::tile_width>>;
-    using o_tile    =         st_bf<K::qo_height, K::tile_width>;
-
-    using q_global = gl<bf16,  -1, -1, -1, -1, q_tile>;
-    using k_global = gl<bf16,  -1, -1, -1, -1, k_tile>;
-    using v_global = gl<bf16,  -1, -1, -1, -1, v_tile>;
-    using l_global = gl<float, -1, -1, -1, -1, l_col_vec>;
-    using o_global = gl<bf16,  -1, -1, -1, -1, o_tile>;
-
-    using globals      = fwd_globals<D>;
-
-    q_global qg_arg{d_q, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    k_global kg_arg{d_k, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    v_global vg_arg{d_v, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    l_global lg_arg{d_l, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,   static_cast<unsigned int>(seq_len)};
-    o_global og_arg{d_o, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-
-    globals g{
-        qg_arg, kg_arg, vg_arg, lg_arg, og_arg, 
-        static_cast<int>(seq_len), static_cast<int>(hr), static_cast<int>(max_kv_blocks_per_q), 
-        q2k_block_sparse_index_ptr, q2k_block_sparse_num_ptr, block_size_ptr
-    };
-
-    // Shared memory size for the kernel
-    // 54000 bytes is calibrated for H100 shared memory constraints for these tile sizes
-    constexpr int mem_size = 54000;
-
-    dim3 grid(seq_len/(64), qo_heads, batch);
-
-    cudaFuncSetAttribute(
-        fwd_attend_ker<D>,
-        cudaFuncAttributeMaxDynamicSharedMemorySize,
-        mem_size
-    );
-
-    fwd_attend_ker<D><<<grid, (128), mem_size, stream>>>(g);
-}
-
-template<int D>
-void block_sparse_attention_backward_impl(
-    bf16* d_q, bf16* d_k, bf16* d_v, bf16* d_o, bf16* d_og, float* d_l, float* d_d, float* d_qg, float* d_kg, float* d_vg,
-    int batch, int qo_heads, int kv_heads, int seq_len, int hr, int max_q_blocks_per_kv,
-    int32_t* k2q_block_sparse_index_ptr,
-    int32_t* k2q_block_sparse_num_ptr,
-    int32_t* block_size_ptr,
-    cudaStream_t stream
-) {
-    using G = bwd_attend_ker_tile_dims<D>;
-    using og_tile = st_bf<4*16, D>;
-    using o_tile  = st_bf<4*16, D>;
-    using d_tile  = col_vec<st_fl<4*16, D>>;
-
-    using og_global = gl<bf16,  -1, -1, -1, -1, og_tile>;
-    using o_global  = gl<bf16,  -1, -1, -1, -1, o_tile>;
-    using d_global  = gl<float, -1, -1, -1, -1, d_tile>;
-
-    using prep_globals = bwd_prep_globals<D>;
-
-    constexpr int mem_size_prep = kittens::MAX_SHARED_MEMORY;
-    int threads_prep = PREP_NUM_WARPS * kittens::WARP_THREADS;
-    dim3 grid_bwd_prep(seq_len/(PREP_NUM_WARPS*kittens::TILE_ROW_DIM<bf16>*4), qo_heads, batch);
-
-    og_global prep_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    o_global  prep_o_arg {d_o,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    d_global  prep_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
-
-    prep_globals bwd_g{prep_og_arg, prep_o_arg, prep_d_arg};
-
-    cudaFuncSetAttribute(
-        bwd_attend_prep_ker<D>,
-        cudaFuncAttributeMaxDynamicSharedMemorySize,
-        mem_size_prep
-    );
-    bwd_attend_prep_ker<D><<<grid_bwd_prep, threads_prep, mem_size_prep, stream>>>(bwd_g);
-
-    using bwd_q_tile    =         st_bf<G::tile_h_qo, G::tile_width>;
-    using bwd_k_tile    =         st_bf<G::tile_h,    G::tile_width>;
-    using bwd_v_tile    =         st_bf<G::tile_h,    G::tile_width>;
-    using bwd_og_tile   =         st_bf<G::tile_h_qo, G::tile_width>;
-    using bwd_qg_tile   =         st_fl<G::tile_h_qo, G::tile_width>;
-    using bwd_kg_tile   =         st_fl<G::tile_h,    G::tile_width>;
-    using bwd_vg_tile   =         st_fl<G::tile_h,    G::tile_width>;
-    using bwd_l_tile    = row_vec<st_fl<G::tile_h_qo, G::tile_h>>;
-    using bwd_d_tile    = row_vec<st_fl<G::tile_h_qo, G::tile_h>>;
-
-    using bwd_q_global  = gl<bf16,  -1, -1, -1, -1, bwd_q_tile>;
-    using bwd_k_global  = gl<bf16,  -1, -1, -1, -1, bwd_k_tile>;
-    using bwd_v_global  = gl<bf16,  -1, -1, -1, -1, bwd_v_tile>;
-    using bwd_og_global = gl<bf16,  -1, -1, -1, -1, bwd_og_tile>;
-    using bwd_qg_global = gl<float, -1, -1, -1, -1, bwd_qg_tile>;
-    using bwd_kg_global = gl<float, -1, -1, -1, -1, bwd_kg_tile>;
-    using bwd_vg_global = gl<float, -1, -1, -1, -1, bwd_vg_tile>;
-    using bwd_l_global  = gl<float, -1, -1, -1, -1, bwd_l_tile>;
-    using bwd_d_global  = gl<float, -1, -1, -1, -1, bwd_d_tile>;
-
-    using bwd_global_args = bwd_globals<D>;
-
-    bwd_q_global  bwd_q_arg {d_q,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    bwd_k_global  bwd_k_arg {d_k,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    bwd_v_global  bwd_v_arg {d_v,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    bwd_og_global bwd_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    bwd_qg_global bwd_qg_arg{d_qg, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    bwd_kg_global bwd_kg_arg{d_kg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    bwd_vg_global bwd_vg_arg{d_vg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), static_cast<uint32_t>(D)};
-    bwd_l_global  bwd_l_arg {d_l,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
-    bwd_d_global  bwd_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
-
-    bwd_global_args bwd_global{bwd_q_arg, bwd_k_arg, bwd_v_arg, bwd_og_arg, bwd_qg_arg, bwd_kg_arg, bwd_vg_arg, bwd_l_arg, bwd_d_arg, 
-                        static_cast<int>(seq_len), static_cast<int>(hr), static_cast<int>(max_q_blocks_per_kv),
-                        k2q_block_sparse_index_ptr, k2q_block_sparse_num_ptr, block_size_ptr};
-
-    dim3 grid_bwd_main(seq_len/64, qo_heads, batch);
-    int threads_main = 128;
-    // Calibrated shared memory sizes for different head dimensions
-    int bwd_mem_size = (D == 64) ? 72000 : 113000;
-
-    cudaFuncSetAttribute(
-        bwd_attend_ker<D>,
-        cudaFuncAttributeMaxDynamicSharedMemorySize,
-        bwd_mem_size
-    );
-    bwd_attend_ker<D><<<grid_bwd_main, threads_main, bwd_mem_size, stream>>>(bwd_global);
-}
-
 #include "pyutils/torch_helpers.cuh"
 #include <ATen/cuda/CUDAContext.h>
 #include <iostream>
@@ -810,23 +672,32 @@ block_sparse_attention_forward(
    torch::Tensor v, 
    torch::Tensor q2k_block_sparse_index, 
    torch::Tensor q2k_block_sparse_num, 
-    torch::Tensor block_size
+    torch::Tensor kv_block_size
 )
 {
    CHECK_INPUT(q);
    CHECK_INPUT(k);
    CHECK_INPUT(v);

+    // q shape: (batch, qo_heads, q_seq_len, head_dim)
+    // k shape: (batch, kv_heads, kv_seq_len, head_dim)
+    // v shape: (batch, kv_heads, kv_seq_len, head_dim)
+    // q2k_block_sparse_index shape: (batch, qo_heads, num_q_blocks, max_kv_blocks_per_q)
+    // q2k_block_sparse_num shape: (batch, qo_heads, num_q_blocks)
+    // kv_block_size shape: (num_kv_blocks) This does not need other dimensions because across all batch/heads the padding is the same.
+    
    auto batch    = q.size(0);
-    auto seq_len  = q.size(2); 
+    auto q_seq_len  = q.size(2);
+    auto kv_seq_len = k.size(2);
    auto head_dim = q.size(3); 
    auto qo_heads = q.size(1);
    auto kv_heads = k.size(1);
    auto max_kv_blocks_per_q = q2k_block_sparse_index.size(3);
-    auto num_q_blocks = block_size.size(0);
+    auto num_q_blocks = q2k_block_sparse_index.size(2);
+    auto num_kv_blocks = kv_block_size.size(0);
    TORCH_CHECK(batch==1, "Batch size dim will be removed in the future, please set batch to 1");
-    TORCH_CHECK(num_q_blocks * 64 == seq_len, "This kernel supports variable block size, but it assumes the input sequence is properly padded.");
-    TORCH_CHECK(num_q_blocks == q2k_block_sparse_index.size(2), "Number of Q blocks does not match between q2k_block_sparse_index and block_size");
+    TORCH_CHECK(num_q_blocks * BLOCK_M == q_seq_len, "This kernel supports variable q block size, but it assumes the input sequence is properly padded.");
+    TORCH_CHECK(num_kv_blocks * BLOCK_M == kv_seq_len, "This kernel supports variable kv block size, but it assumes the input sequence is properly padded.");
    // check to see that these dimensions match for all inputs
    TORCH_CHECK(q.size(0) == batch, "Q batch dimension - idx 0 - must match for all inputs");
    TORCH_CHECK(k.size(0) == batch, "K batch dimension - idx 0 - must match for all inputs");
@@ -834,11 +705,9 @@ block_sparse_attention_forward(
    TORCH_CHECK(q2k_block_sparse_index.size(0) == batch, "q2k_block_sparse_index batch dimension - idx 0 - must match for all inputs");
    TORCH_CHECK(q2k_block_sparse_num.size(0) == batch, "q2k_block_sparse_num batch dimension - idx 0 - must match for all inputs");

-    TORCH_CHECK(q.size(2) == seq_len, "Q sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(k.size(2) == seq_len, "K sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(v.size(2) == seq_len, "V sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(q2k_block_sparse_index.size(2) == seq_len / BLOCK_M, "q2k_block_sparse_index idx 2 - must match seq_len / BLOCK_M");
-    TORCH_CHECK(q2k_block_sparse_num.size(2) == seq_len / BLOCK_M, "q2k_block_sparse_num idx 2 - must match seq_len / BLOCK_M");
+    TORCH_CHECK(v.size(2) == kv_seq_len, "V sequence length dimension - idx 2 - must match K inputs");
+    TORCH_CHECK(q2k_block_sparse_num.size(2) == num_q_blocks, "q2k_block_sparse_num idx 2 - must match num_q_blocks");
+

    TORCH_CHECK(q.size(3) == head_dim, "Q head dimension - idx 3 - must match for all non-vector inputs");
    TORCH_CHECK(k.size(3) == head_dim, "K head dimension - idx 3 - must match for all non-vector inputs");
@@ -864,12 +733,12 @@ block_sparse_attention_forward(
    // for the returned outputs
    torch::Tensor o     = torch::empty({static_cast<const uint>(batch), 
                                        static_cast<const uint>(qo_heads), 
-                                        static_cast<const uint>(seq_len), 
+                                        static_cast<const uint>(q_seq_len), 
                                        static_cast<const uint>(head_dim)}, v.options());
    
    torch::Tensor l_vec = torch::empty({static_cast<const uint>(batch), 
                                        static_cast<const uint>(qo_heads), 
-                                        static_cast<const uint>(seq_len), 
+                                        static_cast<const uint>(q_seq_len), 
                                        static_cast<const uint>(1)}, 
                                        torch::TensorOptions().dtype(torch::kFloat).device(q.device()).memory_format(at::MemoryFormat::Contiguous));
        
@@ -880,32 +749,110 @@ block_sparse_attention_forward(
    float* l_ptr = reinterpret_cast<float*>(l_vec.data_ptr<float>());
    float* d_l   = reinterpret_cast<float*>(l_ptr);

+    //cudadevicesynchronize();
    const c10::cuda::OptionalCUDAGuard device_guard(q.device());    
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream(); 

-    // Temporated implementation to avoid code duplication between head_dim=64 and 128
    if (head_dim == 64) {
-        block_sparse_attention_forward_impl<64>(
-            d_q, d_k, d_v, d_l, d_o,
-            batch, qo_heads, kv_heads, seq_len, hr,
-            max_kv_blocks_per_q,
-            reinterpret_cast<int32_t*>(q2k_block_sparse_index.data_ptr()),
+        using q_tile    =         st_bf<fwd_attend_ker_tile_dims<64>::qo_height, fwd_attend_ker_tile_dims<64>::tile_width>;
+        using k_tile    =         st_bf<fwd_attend_ker_tile_dims<64>::kv_height, fwd_attend_ker_tile_dims<64>::tile_width>;
+        using v_tile    =         st_bf<fwd_attend_ker_tile_dims<64>::kv_height, fwd_attend_ker_tile_dims<64>::tile_width>;
+        using l_col_vec = col_vec<st_fl<fwd_attend_ker_tile_dims<64>::qo_height, fwd_attend_ker_tile_dims<64>::tile_width>>;
+        using o_tile    =         st_bf<fwd_attend_ker_tile_dims<64>::qo_height, fwd_attend_ker_tile_dims<64>::tile_width>;
+
+        using q_global = gl<bf16,  -1, -1, -1, -1, q_tile>;
+        using k_global = gl<bf16,  -1, -1, -1, -1, k_tile>;
+        using v_global = gl<bf16,  -1, -1, -1, -1, v_tile>;
+        using l_global = gl<float, -1, -1, -1, -1, l_col_vec>;
+        using o_global = gl<bf16,  -1, -1, -1, -1, o_tile>;
+
+        using globals      = fwd_globals<64>;
+
+        q_global qg_arg{d_q, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 64U};
+        k_global kg_arg{d_k, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 64U};
+        v_global vg_arg{d_v, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 64U};
+        l_global lg_arg{d_l, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,   static_cast<unsigned int>(q_seq_len)};
+        o_global og_arg{d_o, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 64U};
+
+        globals g{
+            qg_arg, 
+            kg_arg, 
+            vg_arg, 
+            lg_arg, 
+            og_arg, 
+            static_cast<int>(q_seq_len), 
+            static_cast<int>(hr), 
+            static_cast<int>(max_kv_blocks_per_q), 
+            reinterpret_cast<int32_t*>(q2k_block_sparse_index.data_ptr()), 
            reinterpret_cast<int32_t*>(q2k_block_sparse_num.data_ptr()),
-            reinterpret_cast<int32_t*>(block_size.data_ptr()),
-            stream
+            reinterpret_cast<int32_t*>(kv_block_size.data_ptr())
+        };
+
+        constexpr int mem_size = 54000;
+
+        dim3 grid(q_seq_len/(BLOCK_M), qo_heads, batch);
+
+        cudaFuncSetAttribute(
+            fwd_attend_ker<64>,
+            cudaFuncAttributeMaxDynamicSharedMemorySize,
+            mem_size
        );
-    } else if (head_dim == 128) {
-        block_sparse_attention_forward_impl<128>(
-            d_q, d_k, d_v, d_l, d_o,
-            batch, qo_heads, kv_heads, seq_len, hr,
-            max_kv_blocks_per_q,
-            reinterpret_cast<int32_t*>(q2k_block_sparse_index.data_ptr()),
+
+        fwd_attend_ker<64><<<grid, (128), mem_size, stream>>>(g);
+
+        CHECK_CUDA_ERROR(cudaGetLastError());
+        // cudaStreamSynchronize(stream);
+    }
+
+    if (head_dim == 128) {
+        using q_tile    =         st_bf<fwd_attend_ker_tile_dims<128>::qo_height, fwd_attend_ker_tile_dims<128>::tile_width>;
+        using k_tile    =         st_bf<fwd_attend_ker_tile_dims<128>::kv_height, fwd_attend_ker_tile_dims<128>::tile_width>;
+        using v_tile    =         st_bf<fwd_attend_ker_tile_dims<128>::kv_height, fwd_attend_ker_tile_dims<128>::tile_width>;
+        using l_col_vec = col_vec<st_fl<fwd_attend_ker_tile_dims<128>::qo_height, fwd_attend_ker_tile_dims<128>::tile_width>>;
+        using o_tile    =         st_bf<fwd_attend_ker_tile_dims<128>::qo_height, fwd_attend_ker_tile_dims<128>::tile_width>;
+
+        using q_global = gl<bf16,  -1, -1, -1, -1, q_tile>;
+        using k_global = gl<bf16,  -1, -1, -1, -1, k_tile>;
+        using v_global = gl<bf16,  -1, -1, -1, -1, v_tile>;
+        using l_global = gl<float, -1, -1, -1, -1, l_col_vec>;
+        using o_global = gl<bf16,  -1, -1, -1, -1, o_tile>;
+
+        using globals      = fwd_globals<128>;
+
+        q_global qg_arg{d_q, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 128U};
+        k_global kg_arg{d_k, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 128U};
+        v_global vg_arg{d_v, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 128U};
+        l_global lg_arg{d_l, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,   static_cast<unsigned int>(q_seq_len)};
+        o_global og_arg{d_o, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 128U};
+
+        globals g{
+            qg_arg, 
+            kg_arg, 
+            vg_arg, 
+            lg_arg, 
+            og_arg, 
+            static_cast<int>(q_seq_len), 
+            static_cast<int>(hr), 
+            static_cast<int>(max_kv_blocks_per_q), 
+            reinterpret_cast<int32_t*>(q2k_block_sparse_index.data_ptr()), 
            reinterpret_cast<int32_t*>(q2k_block_sparse_num.data_ptr()),
-            reinterpret_cast<int32_t*>(block_size.data_ptr()),
-            stream
+            reinterpret_cast<int32_t*>(kv_block_size.data_ptr())
+        };
+
+        constexpr int mem_size = 54000;
+
+        dim3 grid(q_seq_len/(BLOCK_M), qo_heads, batch);
+
+        cudaFuncSetAttribute(
+            fwd_attend_ker<128>,
+            cudaFuncAttributeMaxDynamicSharedMemorySize,
+            mem_size
        );
-    } else {
-        TORCH_CHECK(false, "Unsupported head_dim: ", head_dim, ". Only 64 and 128 are supported.");
+
+        fwd_attend_ker<128><<<grid, (128), mem_size, stream>>>(g);
+
+        CHECK_CUDA_ERROR(cudaGetLastError());
+        // cudaStreamSynchronize(stream);
    }

    return {o, l_vec};
@@ -921,7 +868,7 @@ block_sparse_attention_backward(torch::Tensor q,
                   torch::Tensor og,
                   torch::Tensor k2q_block_sparse_index,
                   torch::Tensor k2q_block_sparse_num,
-                   torch::Tensor block_size)
+                   torch::Tensor kv_block_size)
 {
    CHECK_INPUT(q);
    CHECK_INPUT(k);
@@ -930,11 +877,23 @@ block_sparse_attention_backward(torch::Tensor q,
    CHECK_INPUT(o);
    CHECK_INPUT(og);

+    // q: [batch, qo_heads, q_seq_len, head_dim]
+    // k: [batch, kv_heads, kv_seq_len, head_dim]
+    // v: [batch, kv_heads, kv_seq_len, head_dim]
+    // o: [batch, qo_heads, q_seq_len, head_dim]
+    // l_vec: [batch, qo_heads, q_seq_len, 1]
+    // og: [batch, qo_heads, q_seq_len, head_dim]
+    // k2q_block_sparse_index: [batch, kv_heads, num_kv_blocks, max_num_q_blocks]
+    // k2q_block_sparse_num: [batch, kv_heads, num_kv_blocks]
+    // kv_block_size: [num_kv_blocks]
+
    auto batch    = q.size(0);
-    auto seq_len  = q.size(2);
+    auto q_seq_len  = q.size(2);
+    auto kv_seq_len = k.size(2);
    auto head_dim = q.size(3);
    auto max_q_blocks_per_kv = k2q_block_sparse_index.size(3);
-    TORCH_CHECK(k2q_block_sparse_index.size(2) == block_size.size(0), "k2q_block_sparse_index.size(2) must match block_size.size(0)");
+    auto num_kv_blocks = kv_block_size.size(0);
+    TORCH_CHECK(k2q_block_sparse_index.size(2) == num_kv_blocks, "k2q_block_sparse_index.size(2) must match num_kv_blocks (kv_block_size.size(0))");
    // check to see that these dimensions match for all inputs
    TORCH_CHECK(q.size(0)     == batch, "Q  batch dimension - idx 0 - must match for all inputs");
    TORCH_CHECK(k.size(0)     == batch, "K  batch dimension - idx 0 - must match for all inputs");
@@ -945,23 +904,18 @@ block_sparse_attention_backward(torch::Tensor q,
    TORCH_CHECK(k2q_block_sparse_index.size(0) == batch, "k2q_block_sparse_index batch dimension - idx 0 - must match for all inputs");
    TORCH_CHECK(k2q_block_sparse_num.size(0) == batch, "k2q_block_sparse_num batch dimension - idx 0 - must match for all inputs");

-
-    TORCH_CHECK(q.size(2)     == seq_len, "Q  sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(k.size(2)     == seq_len, "K  sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(v.size(2)     == seq_len, "V  sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(l_vec.size(2) == seq_len, "L  sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(o.size(2)     == seq_len, "O  sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(og.size(2)    == seq_len, "OG sequence length dimension - idx 2 - must match for all inputs");
-    TORCH_CHECK(k2q_block_sparse_index.size(2) == seq_len / BLOCK_N, "k2q_block_sparse_index idx 2 - must match seq_len / BLOCK_N");
-    TORCH_CHECK(k2q_block_sparse_num.size(2) == seq_len / BLOCK_N, "k2q_block_sparse_num idx 2 - must match seq_len / BLOCK_N");
+    TORCH_CHECK(v.size(2)     == kv_seq_len, "V  sequence length dimension - idx 2 - must match K sequence length");
+    TORCH_CHECK(l_vec.size(2) == q_seq_len, "L  sequence length dimension - idx 2 - must match Q sequence length");
+    TORCH_CHECK(o.size(2)     == q_seq_len, "O  sequence length dimension - idx 2 - must match Q sequence length");
+    TORCH_CHECK(og.size(2)    == q_seq_len, "OG sequence length dimension - idx 2 - must match Q sequence length");
+    TORCH_CHECK(k2q_block_sparse_index.size(2) == num_kv_blocks, "k2q_block_sparse_index idx 2 - must match num_kv_blocks (kv_block_size.size(0))");
+    TORCH_CHECK(k2q_block_sparse_num.size(2) == num_kv_blocks, "k2q_block_sparse_num idx 2 - must match num_kv_blocks (kv_block_size.size(0))");

    TORCH_CHECK(q.size(3)  == head_dim, "Q  head dimension - idx 3 - must match for all non-vector inputs");
    TORCH_CHECK(k.size(3)  == head_dim, "K  head dimension - idx 3 - must match for all non-vector inputs");
    TORCH_CHECK(v.size(3)  == head_dim, "V  head dimension - idx 3 - must match for all non-vector inputs");
    TORCH_CHECK(o.size(3)  == head_dim, "O  head dimension - idx 3 - must match for all non-vector inputs");
    TORCH_CHECK(og.size(3) == head_dim, "OG head dimension - idx 3 - must match for all non-vector inputs");
-    
-

    auto qo_heads = q.size(1);
    auto kv_heads = k.size(1);
@@ -988,20 +942,20 @@ block_sparse_attention_backward(torch::Tensor q,

    torch::Tensor qg = torch::zeros({static_cast<const uint>(batch), 
                                     static_cast<const uint>(qo_heads), 
-                                     static_cast<const uint>(seq_len), 
+                                     static_cast<const uint>(q_seq_len), 
                                     static_cast<const uint>(head_dim)},   l_vec.options());
    torch::Tensor kg = torch::zeros({static_cast<const uint>(batch), 
                                     static_cast<const uint>(kv_heads), 
-                                     static_cast<const uint>(seq_len), 
+                                     static_cast<const uint>(kv_seq_len), 
                                     static_cast<const uint>(head_dim)},   l_vec.options());
    torch::Tensor vg = torch::zeros({static_cast<const uint>(batch), 
                                     static_cast<const uint>(kv_heads), 
-                                     static_cast<const uint>(seq_len), 
+                                     static_cast<const uint>(kv_seq_len), 
                                     static_cast<const uint>(head_dim)},   l_vec.options());
    
    torch::Tensor d_vec = torch::empty({static_cast<const uint>(batch), 
                                        static_cast<const uint>(qo_heads), 
-                                        static_cast<const uint>(seq_len), 
+                                        static_cast<const uint>(q_seq_len), 
                                        static_cast<const uint>(1)},       l_vec.options());

    float*         qg_ptr = qg.data_ptr<float>();
@@ -1030,7 +984,7 @@ block_sparse_attention_backward(torch::Tensor q,
    //  cudaStreamSynchronize(stream);

    // TORCH_CHECK(seq_len % (4*kittens::TILE_DIM*4) == 0, "sequence length must be divisible by 256");
-    dim3 grid_bwd(seq_len/(PREP_NUM_WARPS*kittens::TILE_ROW_DIM<bf16>*4), qo_heads, batch);
+    dim3 grid_bwd(q_seq_len/(PREP_NUM_WARPS*kittens::TILE_ROW_DIM<bf16>*4), qo_heads, batch);

    if (head_dim == 64)  {
        using og_tile = st_bf<4*16, 64>;
@@ -1043,9 +997,9 @@ block_sparse_attention_backward(torch::Tensor q,

        using bwd_prep_globals = bwd_prep_globals<64>;

-        og_global prep_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 64U};
-        o_global  prep_o_arg {d_o,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 64U};
-        d_global  prep_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
+        og_global prep_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 64U};
+        o_global  prep_o_arg {d_o,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 64U};
+        d_global  prep_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(q_seq_len)};

        bwd_prep_globals bwd_g{prep_og_arg, prep_o_arg, prep_d_arg};

@@ -1082,15 +1036,15 @@ block_sparse_attention_backward(torch::Tensor q,

        using bwd_global_args = bwd_globals<64>;

-        bwd_q_global  bwd_q_arg {d_q,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 64U};
-        bwd_k_global  bwd_k_arg {d_k,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 64U};
-        bwd_v_global  bwd_v_arg {d_v,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 64U};
-        bwd_og_global bwd_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 64U};
-        bwd_qg_global bwd_qg_arg{d_qg, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 64U};
-        bwd_kg_global bwd_kg_arg{d_kg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 64U};
-        bwd_vg_global bwd_vg_arg{d_vg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 64U};
-        bwd_l_global  bwd_l_arg {d_l,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
-        bwd_d_global  bwd_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
+        bwd_q_global  bwd_q_arg {d_q,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 64U};
+        bwd_k_global  bwd_k_arg {d_k,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 64U};
+        bwd_v_global  bwd_v_arg {d_v,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 64U};
+        bwd_og_global bwd_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 64U};
+        bwd_qg_global bwd_qg_arg{d_qg, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 64U};
+        bwd_kg_global bwd_kg_arg{d_kg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 64U};
+        bwd_vg_global bwd_vg_arg{d_vg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 64U};
+        bwd_l_global  bwd_l_arg {d_l,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(q_seq_len)};
+        bwd_d_global  bwd_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(q_seq_len)};

        bwd_global_args bwd_global{bwd_q_arg, 
                        bwd_k_arg, 
@@ -1101,14 +1055,14 @@ block_sparse_attention_backward(torch::Tensor q,
                        bwd_vg_arg, 
                        bwd_l_arg, 
                        bwd_d_arg, 
-                        static_cast<int>(seq_len), 
+                        static_cast<int>(kv_seq_len), // N is not used in the kernel
                        static_cast<int>(hr),
                        static_cast<int>(max_q_blocks_per_kv),
                        reinterpret_cast<int32_t*>(k2q_block_sparse_index.data_ptr()),
                        reinterpret_cast<int32_t*>(k2q_block_sparse_num.data_ptr()),
-                        reinterpret_cast<int32_t*>(block_size.data_ptr())};
+                        reinterpret_cast<int32_t*>(kv_block_size.data_ptr())};

-        dim3 grid_bwd_2(seq_len/64, qo_heads, batch);
+        dim3 grid_bwd_2(kv_seq_len/BLOCK_N, qo_heads, batch);
        threads = 128;

        //cudadevicesynchronize();
@@ -1147,9 +1101,9 @@ block_sparse_attention_backward(torch::Tensor q,

        using bwd_prep_globals = bwd_prep_globals<128>;

-        og_global prep_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 128U};
-        o_global  prep_o_arg {d_o,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 128U};
-        d_global  prep_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
+        og_global prep_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 128U};
+        o_global  prep_o_arg {d_o,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 128U};
+        d_global  prep_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(q_seq_len)};

        bwd_prep_globals bwd_g{prep_og_arg, prep_o_arg, prep_d_arg};

@@ -1186,15 +1140,15 @@ block_sparse_attention_backward(torch::Tensor q,

        using bwd_global_args = bwd_globals<128>;

-        bwd_q_global  bwd_q_arg {d_q,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 128U};
-        bwd_k_global  bwd_k_arg {d_k,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 128U};
-        bwd_v_global  bwd_v_arg {d_v,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 128U};
-        bwd_og_global bwd_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 128U};
-        bwd_qg_global bwd_qg_arg{d_qg, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(seq_len), 128U};
-        bwd_kg_global bwd_kg_arg{d_kg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 128U};
-        bwd_vg_global bwd_vg_arg{d_vg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(seq_len), 128U};
-        bwd_l_global  bwd_l_arg {d_l,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
-        bwd_d_global  bwd_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(seq_len)};
+        bwd_q_global  bwd_q_arg {d_q,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 128U};
+        bwd_k_global  bwd_k_arg {d_k,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 128U};
+        bwd_v_global  bwd_v_arg {d_v,  static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 128U};
+        bwd_og_global bwd_og_arg{d_og, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 128U};
+        bwd_qg_global bwd_qg_arg{d_qg, static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), static_cast<unsigned int>(q_seq_len), 128U};
+        bwd_kg_global bwd_kg_arg{d_kg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 128U};
+        bwd_vg_global bwd_vg_arg{d_vg, static_cast<unsigned int>(batch), static_cast<unsigned int>(kv_heads), static_cast<unsigned int>(kv_seq_len), 128U};
+        bwd_l_global  bwd_l_arg {d_l,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(q_seq_len)};
+        bwd_d_global  bwd_d_arg {d_d,  static_cast<unsigned int>(batch), static_cast<unsigned int>(qo_heads), 1U,  static_cast<unsigned int>(q_seq_len)};

        bwd_global_args bwd_global{bwd_q_arg, 
                        bwd_k_arg, 
@@ -1205,14 +1159,14 @@ block_sparse_attention_backward(torch::Tensor q,
                        bwd_vg_arg, 
                        bwd_l_arg, 
                        bwd_d_arg, 
-                        static_cast<int>(seq_len), 
+                        static_cast<int>(kv_seq_len), // N is not used in the kernel
                        static_cast<int>(hr),
                        static_cast<int>(max_q_blocks_per_kv),
                        reinterpret_cast<int32_t*>(k2q_block_sparse_index.data_ptr()),
                        reinterpret_cast<int32_t*>(k2q_block_sparse_num.data_ptr()),
-                        reinterpret_cast<int32_t*>(block_size.data_ptr())};
+                        reinterpret_cast<int32_t*>(kv_block_size.data_ptr())};

-        dim3 grid_bwd_2(seq_len/64, qo_heads, batch);
+        dim3 grid_bwd_2(kv_seq_len/BLOCK_N, qo_heads, batch);
        threads = 128;

        //cudadevicesynchronize();
@@ -1233,4 +1187,4 @@ block_sparse_attention_backward(torch::Tensor q,

    return {qg, kg, vg};
    //cudadevicesynchronize();
-}
+}
@@ -4,6 +4,7 @@
 #include <torch/all.h>
 #include <torch/python.h>
 #include <cutlass/cutlass.h>
+#include <cutlass/numeric_types.h>
 #include "common/common.hpp"
 #include "norm/layernorm.hpp"

@@ -14,10 +15,6 @@ auto layer_norm(
  std::optional<at::Tensor const> const B,
  std::optional<at::Tensor> Output
 ) {
-  using ElementIn = float;
-  using ElementOut = float;
-  using ElementWeight = float;
-
  int64_t const m = Input.size(0);
  int64_t const n = Input.size(1);
  torch::Device const input_device = Input.device();
@@ -26,31 +23,70 @@ auto layer_norm(
    Output.emplace(
      torch::empty(
        {m, n},
-        torch::TensorOptions().device(input_device).dtype(torch::kFloat32)
+        torch::TensorOptions().device(input_device).dtype(Input.scalar_type())
      )
    );
  }

+  TORCH_CHECK(Output.value().scalar_type() == Input.scalar_type(),
+              "Output dtype must match Input dtype. Got Output=",
+              Output.value().scalar_type(), ", Input=", Input.scalar_type());
+  if (W.has_value()) {
+    TORCH_CHECK(W.value().scalar_type() == Input.scalar_type(),
+                "W dtype must match Input dtype. Got W=",
+                W.value().scalar_type(), ", Input=", Input.scalar_type());
+  }
+  if (B.has_value()) {
+    TORCH_CHECK(B.value().scalar_type() == Input.scalar_type(),
+                "B dtype must match Input dtype. Got B=",
+                B.value().scalar_type(), ", Input=", Input.scalar_type());
+  }

  void *Iptr = Input.data_ptr();
  void *Wptr = W.has_value() ? W.value().data_ptr() : nullptr;
  void *Bptr = B.has_value() ? B.value().data_ptr() : nullptr;
  void *Optr = Output.value().data_ptr();

-  BOOL_SWITCH(B.has_value(), BIAS, [&]{
-    BOOL_SWITCH(W.has_value(), AFFINE, [&]{
-    CONFIG_SWITCH(n, [&]{
-    layernorm<
-    ElementIn, ElementOut, ElementWeight,
-    AFFINE, BIAS,
-    MAX_HIDDEN_SIZE, NUM_THR_PER_CTA> (
-      Iptr, Wptr, Bptr,
-      Optr, eps, m, n,
-      at::cuda::getCurrentCUDAStream().stream()
-        );
+  auto stream = at::cuda::getCurrentCUDAStream().stream();
+  if (Input.scalar_type() == at::kHalf) {
+    using ElementIn = cutlass::half_t;
+    using ElementOut = cutlass::half_t;
+    using ElementWeight = cutlass::half_t;
+    BOOL_SWITCH(B.has_value(), BIAS, [&]{
+      BOOL_SWITCH(W.has_value(), AFFINE, [&]{
+        CONFIG_SWITCH(n, [&]{
+          layernorm<ElementIn, ElementOut, ElementWeight, AFFINE, BIAS, MAX_HIDDEN_SIZE, NUM_THR_PER_CTA>(
+              Iptr, Wptr, Bptr, Optr, eps, m, n, stream);
+        });
      });
    });
-  });
+  } else if (Input.scalar_type() == at::kBFloat16) {
+    using ElementIn = cutlass::bfloat16_t;
+    using ElementOut = cutlass::bfloat16_t;
+    using ElementWeight = cutlass::bfloat16_t;
+    BOOL_SWITCH(B.has_value(), BIAS, [&]{
+      BOOL_SWITCH(W.has_value(), AFFINE, [&]{
+        CONFIG_SWITCH(n, [&]{
+          layernorm<ElementIn, ElementOut, ElementWeight, AFFINE, BIAS, MAX_HIDDEN_SIZE, NUM_THR_PER_CTA>(
+              Iptr, Wptr, Bptr, Optr, eps, m, n, stream);
+        });
+      });
+    });
+  } else if (Input.scalar_type() == at::kFloat) {
+    using ElementIn = float;
+    using ElementOut = float;
+    using ElementWeight = float;
+    BOOL_SWITCH(B.has_value(), BIAS, [&]{
+      BOOL_SWITCH(W.has_value(), AFFINE, [&]{
+        CONFIG_SWITCH(n, [&]{
+          layernorm<ElementIn, ElementOut, ElementWeight, AFFINE, BIAS, MAX_HIDDEN_SIZE, NUM_THR_PER_CTA>(
+              Iptr, Wptr, Bptr, Optr, eps, m, n, stream);
+        });
+      });
+    });
+  } else {
+    TORCH_CHECK(false, "Unsupported dtype for layer_norm_cuda: ", Input.scalar_type());
+  }
  
    

@@ -68,9 +68,22 @@ public:
    // mean reduction
    float u = _reduce_sum(x, shared_data) / params.n;

+    // IMPORTANT:
+    // Loader pads out-of-range lanes with 0. That is OK for the sum, but after
+    // subtracting mean, those padded lanes become -u and would incorrectly
+    // contribute to the variance. Mask them back to 0 before variance reduction.
+    // We launch exactly NumThrPerCta threads for a 1xMaxHiddenSize tile,
+    // so each thread is responsible for a contiguous chunk in N.
+    int thr_n_offset = tidx * NumElementPerThread;
    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      x[i] -= u;
+    for (int i = 0; i < NumElementPerThread; ++i) {
+      int idx = thr_n_offset + i;
+      if (idx < params.n) {
+        x[i] -= u;
+      } else {
+        x[i] = 0.f;
+      }
+    }

    __syncthreads();
    // var reduction
@@ -4,6 +4,7 @@
 #include <torch/all.h>
 #include <torch/python.h>
 #include <cutlass/cutlass.h>
+#include <cutlass/numeric_types.h>
 #include <pybind11/pybind11.h>

 #include "common/common.hpp"
@@ -16,10 +17,6 @@ auto rms_norm(
  std::optional<at::Tensor>& Output
 ) {
  
-  using ElementIn = float;
-  using ElementOut = float;
-  using ElementWeight = float;
-  
  int64_t const m = Input.size(0);
  int64_t const n = Input.size(1);
  torch::Device const input_device = Input.device();
@@ -28,27 +25,51 @@ auto rms_norm(
    Output.emplace(
      torch::empty(
        {m, n},
-        torch::TensorOptions().device(input_device).dtype(torch::kFloat32)
+        torch::TensorOptions().device(input_device).dtype(Input.scalar_type())
      )
    );
  }

+  TORCH_CHECK(Output.value().scalar_type() == Input.scalar_type(),
+              "Output dtype must match Input dtype. Got Output=",
+              Output.value().scalar_type(), ", Input=", Input.scalar_type());
+  if (Weight.has_value()) {
+    TORCH_CHECK(Weight.value().scalar_type() == Input.scalar_type(),
+                "Weight dtype must match Input dtype. Got Weight=",
+                Weight.value().scalar_type(), ", Input=", Input.scalar_type());
+  }
+
  void *Iptr = Input.data_ptr();
  void *Wptr = Weight.has_value() ? Weight.value().data_ptr() : nullptr;
  void *Optr = Output.value().data_ptr();

-
-  CONFIG_SWITCH(n, [&]{
-    rmsnorm<
-    ElementIn, ElementOut, ElementWeight,
-    MAX_HIDDEN_SIZE, NUM_THR_PER_CTA
-    > (
-      Iptr, Wptr,
-      Optr, 
-      eps, m, n, 
-      at::cuda::getCurrentCUDAStream().stream()
-    );
-});
+  if (Input.scalar_type() == at::kHalf) {
+    using ElementIn = cutlass::half_t;
+    using ElementOut = cutlass::half_t;
+    using ElementWeight = cutlass::half_t;
+    CONFIG_SWITCH(n, [&]{
+      rmsnorm<ElementIn, ElementOut, ElementWeight, MAX_HIDDEN_SIZE, NUM_THR_PER_CTA>(
+          Iptr, Wptr, Optr, eps, m, n, at::cuda::getCurrentCUDAStream().stream());
+    });
+  } else if (Input.scalar_type() == at::kBFloat16) {
+    using ElementIn = cutlass::bfloat16_t;
+    using ElementOut = cutlass::bfloat16_t;
+    using ElementWeight = cutlass::bfloat16_t;
+    CONFIG_SWITCH(n, [&]{
+      rmsnorm<ElementIn, ElementOut, ElementWeight, MAX_HIDDEN_SIZE, NUM_THR_PER_CTA>(
+          Iptr, Wptr, Optr, eps, m, n, at::cuda::getCurrentCUDAStream().stream());
+    });
+  } else if (Input.scalar_type() == at::kFloat) {
+    using ElementIn = float;
+    using ElementOut = float;
+    using ElementWeight = float;
+    CONFIG_SWITCH(n, [&]{
+      rmsnorm<ElementIn, ElementOut, ElementWeight, MAX_HIDDEN_SIZE, NUM_THR_PER_CTA>(
+          Iptr, Wptr, Optr, eps, m, n, at::cuda::getCurrentCUDAStream().stream());
+    });
+  } else {
+    TORCH_CHECK(false, "Unsupported dtype for rms_norm_cuda: ", Input.scalar_type());
+  }
  

  return Output;
@@ -9,7 +9,7 @@ build-backend = "scikit_build_core.build"

 [project]
 name = "fastvideo-kernel"
-version = "0.2.1"
+version = "0.2.5"
 description = "Unified CUDA kernels for FastVideo"
 readme = "README.md"
 requires-python = ">=3.10"
@@ -0,0 +1,295 @@
+from __future__ import annotations
+
+import os
+from typing import Tuple
+
+import torch
+
+
+def _get_sm90_ops():
+    try:
+        from fastvideo_kernel._C import fastvideo_kernel_ops  # type: ignore
+    except Exception:
+        return None, None
+    return (
+        getattr(fastvideo_kernel_ops, "block_sparse_fwd", None),
+        getattr(fastvideo_kernel_ops, "block_sparse_bwd", None),
+    )
+
+
+def _is_sm90() -> bool:
+    if not torch.cuda.is_available():
+        return False
+    major, minor = torch.cuda.get_device_capability(0)
+    return major == 9 and minor == 0
+
+
+def _force_triton() -> bool:
+    # Force Triton even on SM90 and even if the compiled extension is available.
+    # Useful for CI / debugging / parity testing.
+    return os.environ.get("FASTVIDEO_KERNEL_VSA_FORCE_TRITON", "0") == "1"
+
+
+def _map_to_index(block_map: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
+    """
+    Preferred map->index conversion used by the wrapper.
+
+    This wrapper **requires** the Triton implementation.
+    If Triton (or the Triton map_to_index module) is not available, it raises.
+    """
+    if block_map.dim() == 3:
+        block_map = block_map.unsqueeze(0)
+    if block_map.dim() != 4:
+        raise ValueError(f"block_map must be [B,H,Q,KV] (or [H,Q,KV]), got shape={tuple(block_map.shape)}")
+    if block_map.dtype != torch.bool:
+        block_map = block_map.to(torch.bool)
+
+    if not block_map.is_cuda:
+        raise RuntimeError("block_map must be a CUDA tensor (Triton map_to_index required).")
+
+    try:
+        from fastvideo_kernel.triton_kernels.index import map_to_index as triton_map_to_index  # local import
+    except Exception as e:
+        raise ImportError(
+            "Triton map_to_index is required but not available. "
+            "Ensure Triton is installed and fastvideo_kernel.triton_kernels.index is importable."
+        ) from e
+    return triton_map_to_index(block_map)
+
+
+@torch.library.custom_op(
+    "fastvideo_kernel::block_sparse_attn_triton",
+    mutates_args=(),
+    device_types="cuda",
+)
+def block_sparse_attn_triton(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    q = q.contiguous()
+    k = k.contiguous()
+    v = v.contiguous()
+    block_map = block_map.to(torch.bool)
+    q2k_idx, q2k_num = _map_to_index(block_map)
+
+    from fastvideo_kernel.triton_kernels.block_sparse_attn_triton import (  # local import
+        triton_block_sparse_attn_forward,
+    )
+
+    o, M = triton_block_sparse_attn_forward(q, k, v, q2k_idx, q2k_num, variable_block_sizes)
+    return o, M
+
+
+
+@torch.library.register_fake("fastvideo_kernel::block_sparse_attn_triton")
+def _block_sparse_attn_triton_fake(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    o = torch.empty_like(q)
+    M = torch.empty((q.shape[0], q.shape[1], q.shape[2]), device=q.device, dtype=torch.float32)
+    return o, M
+
+
+@torch.library.custom_op(
+    "fastvideo_kernel::block_sparse_attn_backward_triton",
+    mutates_args=(),
+    device_types="cuda",
+)
+def block_sparse_attn_backward_triton(
+    grad_output: torch.Tensor,
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    o: torch.Tensor,
+    M: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    grad_output = grad_output.contiguous()
+    block_map = block_map.to(torch.bool)
+    q2k_idx, q2k_num = _map_to_index(block_map)
+    k2q_idx, k2q_num = _map_to_index(block_map.transpose(-1, -2).contiguous())
+
+    from fastvideo_kernel.triton_kernels.block_sparse_attn_triton import (  # local import
+        triton_block_sparse_attn_backward,
+    )
+
+    dq, dk, dv = triton_block_sparse_attn_backward(
+        grad_output, q, k, v, o, M, q2k_idx, q2k_num, k2q_idx, k2q_num, variable_block_sizes
+    )
+    return dq, dk, dv
+
+
+@torch.library.register_fake("fastvideo_kernel::block_sparse_attn_backward_triton")
+def _block_sparse_attn_backward_triton_fake(
+    grad_output: torch.Tensor,
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    o: torch.Tensor,
+    M: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    dq = torch.empty_like(q)
+    dk = torch.empty_like(k)
+    dv = torch.empty_like(v)
+    return dq, dk, dv
+
+
+def _backward_triton(ctx, grad_o, grad_M):
+    q, k, v, o, M, block_map, variable_block_sizes = ctx.saved_tensors
+    dq, dk, dv = block_sparse_attn_backward_triton(grad_o, q, k, v, o, M, block_map, variable_block_sizes)
+    return dq, dk, dv, None, None
+
+
+def _setup_context_triton(ctx, inputs, output):
+    q, k, v, block_map, variable_block_sizes = inputs
+    o, M = output
+    ctx.save_for_backward(q, k, v, o, M, block_map, variable_block_sizes)
+
+
+block_sparse_attn_triton.register_autograd(_backward_triton, setup_context=_setup_context_triton)
+
+
+@torch.library.custom_op(
+    "fastvideo_kernel::block_sparse_attn_sm90",
+    mutates_args=(),
+    device_types="cuda",
+)
+def block_sparse_attn_sm90(
+    q_padded: torch.Tensor,
+    k_padded: torch.Tensor,
+    v_padded: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    block_sparse_fwd, _ = _get_sm90_ops()
+    if block_sparse_fwd is None:
+        raise ImportError("fastvideo_kernel_ops.block_sparse_fwd is not available")
+
+    q_padded = q_padded.contiguous()
+    k_padded = k_padded.contiguous()
+    v_padded = v_padded.contiguous()
+    block_map = block_map.to(torch.bool)
+    q2k_idx, q2k_num = _map_to_index(block_map)
+
+    o_padded, lse_padded = block_sparse_fwd(
+        q_padded, k_padded, v_padded, q2k_idx, q2k_num, variable_block_sizes.int()
+    )
+    return o_padded, lse_padded
+
+
+@torch.library.register_fake("fastvideo_kernel::block_sparse_attn_sm90")
+def _block_sparse_attn_sm90_fake(
+    q_padded: torch.Tensor,
+    k_padded: torch.Tensor,
+    v_padded: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    o = torch.empty_like(q_padded)
+    lse = torch.empty((q_padded.shape[0], q_padded.shape[1], q_padded.shape[2], 1), device=q_padded.device, dtype=torch.float32)
+    return o, lse
+
+
+@torch.library.custom_op(
+    "fastvideo_kernel::block_sparse_attn_backward_sm90",
+    mutates_args=(),
+    device_types="cuda",
+)
+def block_sparse_attn_backward_sm90(
+    grad_output_padded: torch.Tensor,
+    q_padded: torch.Tensor,
+    k_padded: torch.Tensor,
+    v_padded: torch.Tensor,
+    o_padded: torch.Tensor,
+    lse_padded: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    _, block_sparse_bwd = _get_sm90_ops()
+    if block_sparse_bwd is None:
+        raise ImportError("fastvideo_kernel_ops.block_sparse_bwd is not available")
+
+    grad_output_padded = grad_output_padded.contiguous()
+    block_map = block_map.to(torch.bool)
+    k2q_idx, k2q_num = _map_to_index(block_map.transpose(-1, -2).contiguous())
+
+    dq, dk, dv = block_sparse_bwd(
+        q_padded,
+        k_padded,
+        v_padded,
+        o_padded,
+        lse_padded,
+        grad_output_padded,
+        k2q_idx,
+        k2q_num,
+        variable_block_sizes.int(),
+    )
+    # C++ kernel returns fp32 grads; cast back to match PyTorch convention if needed
+    return dq.to(grad_output_padded.dtype), dk.to(grad_output_padded.dtype), dv.to(grad_output_padded.dtype)
+
+
+@torch.library.register_fake("fastvideo_kernel::block_sparse_attn_backward_sm90")
+def _block_sparse_attn_backward_sm90_fake(
+    grad_output_padded: torch.Tensor,
+    q_padded: torch.Tensor,
+    k_padded: torch.Tensor,
+    v_padded: torch.Tensor,
+    o_padded: torch.Tensor,
+    lse_padded: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    dq = torch.empty_like(q_padded)
+    dk = torch.empty_like(k_padded)
+    dv = torch.empty_like(v_padded)
+    return dq, dk, dv
+
+
+def _backward_sm90(ctx, grad_o, grad_lse):
+    q, k, v, o, lse, block_map, variable_block_sizes = ctx.saved_tensors
+    dq, dk, dv = block_sparse_attn_backward_sm90(
+        grad_o, q, k, v, o, lse, block_map, variable_block_sizes
+    )
+    return dq, dk, dv, None, None
+
+
+def _setup_context_sm90(ctx, inputs, output):
+    q, k, v, block_map, variable_block_sizes = inputs
+    o, lse = output
+    ctx.save_for_backward(q, k, v, o, lse, block_map, variable_block_sizes)
+
+
+block_sparse_attn_sm90.register_autograd(_backward_sm90, setup_context=_setup_context_sm90)
+
+
+def block_sparse_attn(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    """
+    Unified block-sparse attention op with autograd support.
+    - On SM90 with compiled extension present: uses fastvideo_kernel_ops.block_sparse_fwd/bwd.
+    - Otherwise: uses Triton implementation (requires q/k/v to have same padded length today).
+    """
+    block_sparse_fwd, block_sparse_bwd = _get_sm90_ops()
+    if (not _force_triton()) and _is_sm90() and (block_sparse_fwd is not None) and (block_sparse_bwd is not None):
+        return block_sparse_attn_sm90(q, k, v, block_map, variable_block_sizes)
+    # Triton path: generally assumes q/k/v share the same padded length
+    if q.shape[2] != k.shape[2] or q.shape[2] != v.shape[2]:
+        raise RuntimeError("Triton fallback requires q/k/v to have the same padded length.")
+    return block_sparse_attn_triton(q, k, v, block_map, variable_block_sizes)
+
+
@@ -1,5 +1,6 @@
 import math
 import torch
+from .block_sparse_attn import block_sparse_attn
 from .triton_kernels.block_sparse_attn_triton import triton_block_sparse_attn_forward
 from .triton_kernels.st_attn_triton import sliding_tile_attention_triton
 from .triton_kernels.index import map_to_index
@@ -45,14 +46,22 @@ def sliding_tile_attention(
    flag = shape_map[seq_shape]

    for head_idx, (t, h, w) in enumerate(window_size):
+        # Per-head slices are not contiguous in the batch dimension when batch>1
+        # (they keep the original head-stride). The TK kernel assumes contiguous
+        # [B, H, S, D] layout, so we materialize a contiguous [B,1,S,D] view.
+        q_h = q[:, head_idx:head_idx + 1].contiguous()
+        k_h = k[:, head_idx:head_idx + 1].contiguous()
+        v_h = v[:, head_idx:head_idx + 1].contiguous()
+        o_h = torch.empty_like(q_h)
        sta_fwd(
-            q[:, head_idx:head_idx + 1], k[:, head_idx:head_idx + 1],
-            v[:, head_idx:head_idx + 1], output[:, head_idx:head_idx + 1],
+            q_h, k_h,
+            v_h, o_h,
            t, h, w, text_length, False, has_text, flag
        )
+        output[:, head_idx:head_idx + 1] = o_h

    if has_text:
-        sta_fwd(q, k, v, output, 3, 3, 3, text_length, True, True, flag)
+        sta_fwd(q.contiguous(), k.contiguous(), v.contiguous(), output, 3, 3, 3, text_length, True, True, flag)

    return output[:, :, :seq_length]

@@ -62,6 +71,7 @@ def video_sparse_attn(
    k: torch.Tensor,
    v: torch.Tensor,
    variable_block_sizes: torch.Tensor,
+    q_variable_block_sizes: torch.Tensor,
    topk: int,
    block_size: int | tuple = 64,
    compress_attn_weight: torch.Tensor = None,
@@ -70,14 +80,42 @@ def video_sparse_attn(
        block_size = (block_size, block_size, block_size)

    block_elements = block_size[0] * block_size[1] * block_size[2]
-    batch, heads, seq_len, dim = q.shape
+    batch, heads, q_seq_len, dim = q.shape
+    kv_seq_len = k.shape[2]
+    if v.shape[2] != kv_seq_len:
+        raise ValueError(
+            f"Expected k and v to have the same sequence length, got "
+            f"k.shape[2]={kv_seq_len}, v.shape[2]={v.shape[2]}"
+        )
+    if k.shape[0] != batch or v.shape[0] != batch or k.shape[1] != heads or v.shape[1] != heads:
+        raise ValueError("Expected q/k/v to have the same batch and head dimensions.")
+
+    if q_seq_len % block_elements != 0 or kv_seq_len % block_elements != 0:
+        raise ValueError(
+            f"q_seq_len and kv_seq_len must be divisible by block_elements={block_elements}, "
+            f"got q_seq_len={q_seq_len}, kv_seq_len={kv_seq_len}"
+        )
+    q_num_blocks = q_seq_len // block_elements
+    kv_num_blocks = kv_seq_len // block_elements
+
+    if variable_block_sizes.numel() != kv_num_blocks:
+        raise ValueError(
+            f"variable_block_sizes must have length kv_num_blocks={kv_num_blocks}, "
+            f"got {variable_block_sizes.numel()}"
+        )
+
+    if q_variable_block_sizes.numel() != q_num_blocks:
+        raise ValueError(
+            f"q_variable_block_sizes must have length q_num_blocks={q_num_blocks}, "
+            f"got {q_variable_block_sizes.numel()}"
+        )

    # Compression branch
-    q_c = q.view(batch, heads, seq_len // block_elements, block_elements, dim)
-    k_c = k.view(batch, heads, seq_len // block_elements, block_elements, dim)
-    v_c = v.view(batch, heads, seq_len // block_elements, block_elements, dim)
+    q_c = q.view(batch, heads, q_num_blocks, block_elements, dim)
+    k_c = k.view(batch, heads, kv_num_blocks, block_elements, dim)
+    v_c = v.view(batch, heads, kv_num_blocks, block_elements, dim)

-    q_c = (q_c.float().sum(dim=3) / variable_block_sizes.view(1, 1, -1, 1)).to(
+    q_c = (q_c.float().sum(dim=3) / q_variable_block_sizes.view(1, 1, -1, 1)).to(
        q.dtype)
    k_c = (k_c.float().sum(dim=3) / variable_block_sizes.view(1, 1, -1, 1)).to(
        k.dtype)
@@ -88,9 +126,9 @@ def video_sparse_attn(
    attn = torch.softmax(scores, dim=-1)
    out_c = torch.matmul(attn, v_c)

-    out_c = out_c.view(batch, heads, seq_len // block_elements, 1, dim)
+    out_c = out_c.view(batch, heads, q_num_blocks, 1, dim)
    out_c = out_c.repeat(1, 1, 1, block_elements,
-                         1).view(batch, heads, seq_len, dim)
+                         1).view(batch, heads, q_seq_len, dim)

    # Sparse branch
    topk_idx = torch.topk(scores, topk, dim=-1).indices
@@ -100,12 +138,17 @@ def video_sparse_attn(
    idx, num = map_to_index(mask)
    
    if block_sparse_fwd is not None:
-        out_s = block_sparse_fwd(
-            q, k, v, idx, num, variable_block_sizes.int()
-        )[0] # block_sparse_fwd returns vector<Tensor>
+        # Use autograd-enabled wrapper so backward works (and still uses SM90 kernel when available)
+        out_s = block_sparse_attn(q, k, v, mask, variable_block_sizes)[0]
    else:
-        out_s, _ = triton_block_sparse_attn_forward(q, k, v, idx, num,
-                                                    variable_block_sizes)
+        if q_seq_len != kv_seq_len:
+            raise RuntimeError(
+                "q/k have different lengths, but the compiled CUDA kernel (block_sparse_fwd) "
+                "is not available. The Triton fallback currently requires q and k/v to have "
+                "the same padded length."
+            )
+        # Triton-only forward (kept for environments without the wrapper deps)
+        out_s, _ = triton_block_sparse_attn_forward(q, k, v, idx, num, variable_block_sizes)

    if compress_attn_weight is not None:
        return out_c * compress_attn_weight + out_s
@@ -1 +1 @@
-__version__ = "0.2.1"
+__version__ = "0.2.5"
@@ -42,37 +42,13 @@ def block_sparse_kernel_test(Q, K, V, block_sparse_mask, variable_block_sizes, q
    q_padded = vsa_pad(Q, q_non_pad_index, q_num_blocks, BLOCK_M)
    k_padded = vsa_pad(K, kv_non_pad_index, kv_num_blocks, BLOCK_M)
    v_padded = vsa_pad(V, kv_non_pad_index, kv_num_blocks, BLOCK_M)
-    # Use raw kernel or triton
-    try:
-        from fastvideo_kernel._C import fastvideo_kernel_ops
-        raw_kernel = getattr(fastvideo_kernel_ops, "block_sparse_fwd", None)
-    except ImportError:
-        raw_kernel = None
+    # Use autograd-enabled wrapper (internally dispatches to SM90 kernel or Triton)
+    from fastvideo_kernel.block_sparse_attn import block_sparse_attn
+    output_padded, _aux = block_sparse_attn(
+        q_padded, k_padded, v_padded, block_sparse_mask, variable_block_sizes
+    )

-    from fastvideo_kernel.triton_kernels.index import map_to_index
-
-    
-    # Convert mask to indices
-    # block_sparse_mask is [H, M, N] bool
-    # We need to map it to index.
-    # block_sparse_mask needs to be expanded/reshaped? 
-    # generate_block_sparse_mask_for_function returns [H, NumBlocksQ, NumBlocksKV]
-    
-    # Ops.py logic:
-    # mask = torch.zeros_like(scores, dtype=torch.bool).scatter_(-1, topk_idx, True)
-    # idx, num = map_to_index(mask)
-    
-    idx, num = map_to_index(block_sparse_mask.unsqueeze(0)) # Add batch dim [1, H, M, N]
-    
-    if raw_kernel:
-        out_s = raw_kernel(q_padded, k_padded, v_padded, idx, num, variable_block_sizes.int())
-        output = out_s[0]
-    else:
-        # Fallback to triton testing if C++ not available
-        from fastvideo_kernel.triton_kernels.block_sparse_attn_triton import triton_block_sparse_attn_forward
-        output, _ = triton_block_sparse_attn_forward(q_padded, k_padded, v_padded, idx, num, variable_block_sizes)
-
-    output = output[:, :, q_non_pad_index, :]
+    output = output_padded[:, :, q_non_pad_index, :]
    output.backward(dO)
    return output, Q.grad, K.grad, V.grad

@@ -264,7 +240,6 @@ def generate_error_graphs_qkdiff(h, d, error_mode='all'):

    print("-" * 150)

-@pytest.mark.skip()
 def test_video_sparse_attention_backward():
    if not torch.cuda.is_available():
        return
@@ -3,6 +3,7 @@ import sys
 from typing import Tuple

 import torch
+import pytest

 from .utils import (
    generate_block_sparse_mask_for_function,
@@ -57,23 +58,14 @@ def block_sparse_forward_test(
    k_padded = ref.vsa_pad(K, kv_non_pad_index, kv_num_blocks, BLOCK_M)
    v_padded = ref.vsa_pad(V, kv_non_pad_index, kv_num_blocks, BLOCK_M)

-    # Use raw kernel or triton
+    # Use autograd-enabled wrapper (internally dispatches SM90 C++ vs Triton)
+    from fastvideo_kernel.block_sparse_attn import block_sparse_attn
    try:
-        from fastvideo_kernel._C import fastvideo_kernel_ops
-        raw_kernel = getattr(fastvideo_kernel_ops, "block_sparse_fwd", None)
-    except ImportError:
-        raw_kernel = None
-        
-    from fastvideo_kernel.triton_kernels.index import map_to_index
-    idx, num = map_to_index(block_sparse_mask)
-    
-    if raw_kernel:
-        out_padded = raw_kernel(q_padded, k_padded, v_padded, idx, num, variable_block_sizes.int())[0]
-    else:
-        from fastvideo_kernel.triton_kernels.block_sparse_attn_triton import triton_block_sparse_attn_forward
-        out_padded, _ = triton_block_sparse_attn_forward(
-            q_padded, k_padded, v_padded, idx, num, variable_block_sizes
+        out_padded, _aux = block_sparse_attn(
+            q_padded, k_padded, v_padded, block_sparse_mask, variable_block_sizes
        )
+    except RuntimeError as e:
+        pytest.skip(str(e))

    # Remove padding on the query side
    out = out_padded[:, :, q_non_pad_index, :]
@@ -156,11 +148,6 @@ def run_forward_qk_diff(
 ) -> Tuple[float, float]:
    """
    Forward-only correctness test for the case S_q != S_kv.
-
-    NOTE:
-      - The Triton backend supports different Q/KV logical lengths via padding.
-      - The SM90 (H100) CUDA backend currently assumes the same number of blocks
-        for Q and KV, so we skip this test there.
    """
    assert torch.cuda.is_available(), "VSA kernels require CUDA"

@@ -1,7 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0

 import torch
-from fastvideo.attention.backends.sageattn.api import sageattn_blackwell
+from sageattn3 import sageattn3_blackwell

 from fastvideo.attention.backends.abstract import (AttentionBackend,
                                                   AttentionImpl,
@@ -67,6 +67,6 @@ class SageAttention3Impl(AttentionImpl):
        query = query.transpose(1, 2)
        key = key.transpose(1, 2)
        value = value.transpose(1, 2)
-        output = sageattn_blackwell(query, key, value, is_causal=self.causal)
+        output = sageattn3_blackwell(query, key, value, is_causal=self.causal)
        output = output.transpose(1, 2)
        return output
@@ -276,8 +276,9 @@ class VideoSparseAttentionImpl(AttentionImpl):
            query,
            key,
            value,
-            variable_block_sizes=attn_metadata.variable_block_sizes,
-            topk=cur_topk,
+            attn_metadata.variable_block_sizes,
+            attn_metadata.variable_block_sizes,
+            cur_topk,
            block_size=VSA_TILE_SIZE,
            compress_attn_weight=gate_compress).transpose(1, 2)

@@ -2,5 +2,18 @@ from fastvideo.configs.models.base import ModelConfig
 from fastvideo.configs.models.dits.base import DiTConfig
 from fastvideo.configs.models.encoders.base import EncoderConfig
 from fastvideo.configs.models.vaes.base import VAEConfig
+from fastvideo.configs.models.audio import (LTX2AudioDecoderConfig,
+                                            LTX2AudioEncoderConfig,
+                                            LTX2VocoderConfig)
+from fastvideo.configs.models.upsamplers.base import UpsamplerConfig

-__all__ = ["ModelConfig", "VAEConfig", "DiTConfig", "EncoderConfig"]
+__all__ = [
+    "ModelConfig",
+    "VAEConfig",
+    "DiTConfig",
+    "EncoderConfig",
+    "LTX2AudioEncoderConfig",
+    "LTX2AudioDecoderConfig",
+    "LTX2VocoderConfig",
+    "UpsamplerConfig",
+]
@@ -0,0 +1,13 @@
+# SPDX-License-Identifier: Apache-2.0
+
+from fastvideo.configs.models.audio.ltx2_audio_vae import (
+    LTX2AudioDecoderConfig,
+    LTX2AudioEncoderConfig,
+    LTX2VocoderConfig,
+)
+
+__all__ = [
+    "LTX2AudioEncoderConfig",
+    "LTX2AudioDecoderConfig",
+    "LTX2VocoderConfig",
+]
@@ -0,0 +1,31 @@
+# SPDX-License-Identifier: Apache-2.0
+"""
+LTX-2 audio VAE and vocoder configuration.
+"""
+
+from dataclasses import dataclass, field
+
+from fastvideo.configs.models.base import ArchConfig, ModelConfig
+
+
+@dataclass
+class LTX2AudioArchConfig(ArchConfig):
+    architectures: list[str] = field(default_factory=list)
+
+
+@dataclass
+class LTX2AudioEncoderConfig(ModelConfig):
+    arch_config: ArchConfig = field(default_factory=lambda: LTX2AudioArchConfig(
+        architectures=["LTX2AudioEncoder"]))
+
+
+@dataclass
+class LTX2AudioDecoderConfig(ModelConfig):
+    arch_config: ArchConfig = field(default_factory=lambda: LTX2AudioArchConfig(
+        architectures=["LTX2AudioDecoder"]))
+
+
+@dataclass
+class LTX2VocoderConfig(ModelConfig):
+    arch_config: ArchConfig = field(default_factory=lambda: LTX2AudioArchConfig(
+        architectures=["LTX2Vocoder"]))
@@ -3,11 +3,13 @@ from fastvideo.configs.models.dits.cosmos2_5 import Cosmos25VideoConfig
 from fastvideo.configs.models.dits.hunyuanvideo import HunyuanVideoConfig
 from fastvideo.configs.models.dits.hunyuanvideo15 import HunyuanVideo15Config
 from fastvideo.configs.models.dits.longcat import LongCatVideoConfig
+from fastvideo.configs.models.dits.ltx2 import LTX2VideoConfig
 from fastvideo.configs.models.dits.stepvideo import StepVideoConfig
 from fastvideo.configs.models.dits.wanvideo import WanVideoConfig
+from fastvideo.configs.models.dits.hyworld import HYWorldConfig

 __all__ = [
    "HunyuanVideoConfig", "HunyuanVideo15Config", "WanVideoConfig",
    "StepVideoConfig", "CosmosVideoConfig", "Cosmos25VideoConfig",
-    "LongCatVideoConfig"
+    "LongCatVideoConfig", "LTX2VideoConfig", "HYWorldConfig"
 ]
@@ -55,6 +55,8 @@ class HunyuanVideo15ArchConfig(DiTArchConfig):
            r"txt_in.refiner_blocks.\1.mlp.fc_out.\2",
            r"^context_embedder\.token_refiner\.refiner_blocks\.(\d+)\.norm_out\.linear\.(.*)$":
            r"txt_in.refiner_blocks.\1.adaLN_modulation.linear.\2",
+            r"^context_embedder\.token_refiner\.refiner_blocks\.(\d+)\.self_attn_qkv\.(.*)$":
+            r"txt_in.refiner_blocks.\1.self_attn_qkv.\2",

            # 2. txt_in_2 mapping:
            r"^context_embedder_2\.(.*)$":
@@ -0,0 +1,202 @@
+# SPDX-License-Identifier: Apache-2.0
+from dataclasses import dataclass, field
+
+from fastvideo.configs.models.dits.base import DiTArchConfig, DiTConfig
+
+
+def is_double_block(n: str, m) -> bool:
+    return "double" in n and str.isdigit(n.split(".")[-1])
+
+
+def is_refiner_block(n: str, m) -> bool:
+    return "refiner" in n and str.isdigit(n.split(".")[-1])
+
+
+def is_txt_in(n: str, m) -> bool:
+    return n.split(".")[-1] == "txt_in"
+
+
+@dataclass
+class HYWorldArchConfig(DiTArchConfig):
+    _fsdp_shard_conditions: list = field(
+        default_factory=lambda: [is_double_block, is_refiner_block])
+
+    _compile_conditions: list = field(
+        default_factory=lambda: [is_double_block, is_refiner_block, is_txt_in])
+
+    param_names_mapping: dict = field(
+        default_factory=lambda: {
+            # 1. txt_in submodules (text embedder, refiner blocks):
+            r"^txt_in\.t_embedder\.mlp\.0\.(.*)$":
+            r"txt_in.t_embedder.mlp.fc_in.\1",
+            r"^txt_in\.t_embedder\.mlp\.2\.(.*)$":
+            r"txt_in.t_embedder.mlp.fc_out.\1",
+            r"^txt_in\.c_embedder\.linear_1\.(.*)$":
+            r"txt_in.c_embedder.fc_in.\1",
+            r"^txt_in\.c_embedder\.linear_2\.(.*)$":
+            r"txt_in.c_embedder.fc_out.\1",
+            r"^txt_in\.individual_token_refiner\.blocks\.(\d+)\.norm1\.(.*)$":
+            r"txt_in.refiner_blocks.\1.norm1.\2",
+            r"^txt_in\.individual_token_refiner\.blocks\.(\d+)\.norm2\.(.*)$":
+            r"txt_in.refiner_blocks.\1.norm2.\2",
+            r"^txt_in\.individual_token_refiner\.blocks\.(\d+)\.self_attn_qkv\.(.*)$":
+            r"txt_in.refiner_blocks.\1.self_attn_qkv.\2",
+            r"^txt_in\.individual_token_refiner\.blocks\.(\d+)\.self_attn_proj\.(.*)$":
+            r"txt_in.refiner_blocks.\1.self_attn_proj.\2",
+            r"^txt_in\.individual_token_refiner\.blocks\.(\d+)\.mlp\.fc1\.(.*)$":
+            r"txt_in.refiner_blocks.\1.mlp.fc_in.\2",
+            r"^txt_in\.individual_token_refiner\.blocks\.(\d+)\.mlp\.fc2\.(.*)$":
+            r"txt_in.refiner_blocks.\1.mlp.fc_out.\2",
+            r"^txt_in\.individual_token_refiner\.blocks\.(\d+)\.adaLN_modulation\.1\.(.*)$":
+            r"txt_in.refiner_blocks.\1.adaLN_modulation.linear.\2",
+
+            # 2. time_in mappings:
+            r"^time_in\.mlp\.0\.(.*)$":
+            r"time_in.timestep_embedder.mlp.fc_in.\1",
+            r"^time_in\.mlp\.2\.(.*)$":
+            r"time_in.timestep_embedder.mlp.fc_out.\1",
+
+            # 3. action_in mappings:
+            r"^action_in\.mlp\.0\.(.*)$":
+            r"action_in.mlp.fc_in.\1",
+            r"^action_in\.mlp\.2\.(.*)$":
+            r"action_in.mlp.fc_out.\1",
+
+            # 4. byt5_in -> txt_in_2 mappings:
+            r"^byt5_in\.layernorm\.(.*)$":
+            r"txt_in_2.norm.\1",
+            r"^byt5_in\.fc1\.(.*)$":
+            r"txt_in_2.linear_1.\1",
+            r"^byt5_in\.fc2\.(.*)$":
+            r"txt_in_2.linear_2.\1",
+            r"^byt5_in\.fc3\.(.*)$":
+            r"txt_in_2.linear_3.\1",
+
+            # 5. cond_type_embedding -> cond_type_embed:
+            r"^cond_type_embedding\.(.*)$":
+            r"cond_type_embed.\1",
+
+            # 6. vision_in -> image_embedder mappings:
+            r"^vision_in\.proj\.0\.(.*)$":
+            r"image_embedder.norm_in.\1",
+            r"^vision_in\.proj\.1\.(.*)$":
+            r"image_embedder.linear_1.\1",
+            r"^vision_in\.proj\.3\.(.*)$":
+            r"image_embedder.linear_2.\1",
+            r"^vision_in\.proj\.4\.(.*)$":
+            r"image_embedder.norm_out.\1",
+
+            # 7. double_blocks mapping:
+            r"^double_blocks\.(\d+)\.img_attn_q\.(.*)$":
+            (r"double_blocks.\1.img_attn_qkv.\2", 0, 3),
+            r"^double_blocks\.(\d+)\.img_attn_k\.(.*)$":
+            (r"double_blocks.\1.img_attn_qkv.\2", 1, 3),
+            r"^double_blocks\.(\d+)\.img_attn_v\.(.*)$":
+            (r"double_blocks.\1.img_attn_qkv.\2", 2, 3),
+            r"^double_blocks\.(\d+)\.txt_attn_q\.(.*)$":
+            (r"double_blocks.\1.txt_attn_qkv.\2", 0, 3),
+            r"^double_blocks\.(\d+)\.txt_attn_k\.(.*)$":
+            (r"double_blocks.\1.txt_attn_qkv.\2", 1, 3),
+            r"^double_blocks\.(\d+)\.txt_attn_v\.(.*)$":
+            (r"double_blocks.\1.txt_attn_qkv.\2", 2, 3),
+            r"^double_blocks\.(\d+)\.img_mlp\.fc1\.(.*)$":
+            r"double_blocks.\1.img_mlp.fc_in.\2",
+            r"^double_blocks\.(\d+)\.img_mlp\.fc2\.(.*)$":
+            r"double_blocks.\1.img_mlp.fc_out.\2",
+            r"^double_blocks\.(\d+)\.txt_mlp\.fc1\.(.*)$":
+            r"double_blocks.\1.txt_mlp.fc_in.\2",
+            r"^double_blocks\.(\d+)\.txt_mlp\.fc2\.(.*)$":
+            r"double_blocks.\1.txt_mlp.fc_out.\2",
+
+            # 8. Final layer mapping:
+            r"^final_layer\.adaLN_modulation\.1\.(.*)$":
+            r"final_layer.adaLN_modulation.linear.\1",
+        })
+
+    # Reverse mapping for saving checkpoints: custom -> hf
+    reverse_param_names_mapping: dict = field(default_factory=lambda: {})
+
+    # Parameters from HY-WorldPlay config.json (loaded from checkpoint)
+    patch_size: list | tuple | int = field(default_factory=lambda: [1, 1, 1])
+    # Base latent channels - will be expanded in __post_init__ if concat_condition=True
+    in_channels: int = 32
+    concat_condition: bool = True
+    out_channels: int = 32
+    hidden_size: int = 2048
+    heads_num: int = 16
+    mlp_width_ratio: float = 4.0
+    mlp_act_type: str = "gelu_tanh"
+    mm_double_blocks_depth: int = 54
+    mm_single_blocks_depth: int = 0
+    rope_dim_list: list | tuple = field(default_factory=lambda: [16, 56, 56])
+    qkv_bias: bool = True
+    qk_norm: bool | str = True
+    qk_norm_type: str = "rms"
+    guidance_embed: bool = False
+    use_meanflow: bool = False
+    text_projection: str = "single_refiner"
+    use_attention_mask: bool = True
+    text_states_dim: int = 3584
+    text_states_dim_2: int | None = None
+    text_pool_type: str | None = None
+    rope_theta: float = 256.0
+    attn_mode: str = "flash"
+    attn_param: str | None = None
+    glyph_byT5_v2: bool = True
+    vision_projection: str = "linear"
+    vision_states_dim: int = 1152
+    is_reshape_temporal_channels: bool = False
+    use_cond_type_embedding: bool = True
+    ideal_resolution: str = "480p"
+    ideal_task: str = "i2v"
+    task_type: str = "i2v"
+    exclude_lora_layers: list[str] = field(
+        default_factory=lambda: ["img_in", "txt_in", "time_in", "vector_in"])
+
+    def __post_init__(self):
+        super().__post_init__()
+        # Convert HY-WorldPlay naming to FastVideo naming conventions
+        self.num_attention_heads: int = self.heads_num
+        self.attention_head_dim: int = self.hidden_size // self.heads_num
+        self.num_layers: int = self.mm_double_blocks_depth
+        self.num_single_layers: int = self.mm_single_blocks_depth
+        self.num_refiner_layers: int = 2  # Default for HYWorld
+        self.mlp_ratio: float = float(self.mlp_width_ratio)
+        self.text_embed_dim: int = self.text_states_dim
+        self.text_embed_2_dim: int = self.text_states_dim_2 if self.text_states_dim_2 else 1472
+        self.image_embed_dim: int = self.vision_states_dim
+        self.rope_axes_dim: tuple[int, ...] = tuple(self.rope_dim_list)
+        self.num_channels_latents: int = self.out_channels
+        self.target_size: int = 640
+
+        # Handle concat_condition: when True, actual in_channels = base * 2 + 1
+        # (base latent + condition latent + mask channel)
+        # config.json has base in_channels (32), but img_in needs full (65)
+        if self.concat_condition and self.in_channels == 32:
+            if self.is_reshape_temporal_channels:
+                self.in_channels = self.in_channels + self.in_channels // 2 + 1
+            else:
+                self.in_channels = self.in_channels * 2 + 1  # 32 * 2 + 1 = 65
+
+        # Handle patch_size (can be list/tuple or int)
+        if isinstance(self.patch_size, list | tuple):
+            self.patch_size_t: int = self.patch_size[0]
+            # assume square patch size for height and width
+            patch_size_hw: int = self.patch_size[1]
+            object.__setattr__(self, 'patch_size', patch_size_hw)
+        else:
+            self.patch_size_t = 1
+
+        # Convert qk_norm to string format
+        if isinstance(self.qk_norm, bool):
+            if self.qk_norm:
+                self.qk_norm = "rms_norm" if self.qk_norm_type == "rms" else self.qk_norm_type
+            else:
+                self.qk_norm = "none"
+
+
+@dataclass
+class HYWorldConfig(DiTConfig):
+    arch_config: DiTArchConfig = field(default_factory=HYWorldArchConfig)
+
+    prefix: str = "HYWorld"
@@ -26,7 +26,6 @@ class LongCatVideoArchConfig(DiTArchConfig):
        default_factory=lambda: [is_longcat_blocks])

    # Parameter name mapping for weight conversion
-    # Maps original LongCat third_party names -> native FastVideo names
    param_names_mapping: dict = field(
        default_factory=lambda: {
            # Embedders
@@ -0,0 +1,84 @@
+# SPDX-License-Identifier: Apache-2.0
+"""
+LTX-2 Transformer configuration for native FastVideo integration.
+"""
+
+from dataclasses import dataclass, field
+
+from fastvideo.configs.models.dits.base import DiTArchConfig, DiTConfig
+
+
+def is_ltx2_blocks(name: str, _module) -> bool:
+    """FSDP shard condition for LTX-2 transformer blocks."""
+    return "transformer_blocks" in name
+
+
+@dataclass
+class LTX2VideoArchConfig(DiTArchConfig):
+    """Architecture configuration for LTX-2 video transformer."""
+
+    _fsdp_shard_conditions: list = field(
+        default_factory=lambda: [is_ltx2_blocks])
+    _compile_conditions: list = field(default_factory=lambda: [is_ltx2_blocks])
+
+    # Parameter name mapping for weight conversion (hf/comfy -> FastVideo)
+    param_names_mapping: dict = field(
+        default_factory=lambda: {
+            r"^model\.diffusion_model\.(.*)$": r"model.\1",
+            r"^diffusion_model\.(.*)$": r"model.\1",
+            r"^model\.(.*)$": r"model.\1",
+            r"^(.*)$": r"model.\1",
+        })
+
+    reverse_param_names_mapping: dict = field(default_factory=lambda: {})
+    lora_param_names_mapping: dict = field(default_factory=lambda: {})
+
+    # Core transformer settings (defaults from LTX-2 metadata)
+    num_attention_heads: int = 32
+    attention_head_dim: int = 128
+    num_layers: int = 48
+    cross_attention_dim: int = 4096
+    caption_channels: int = 3840
+    norm_eps: float = 1e-6
+    attention_type: str = "default"
+    rope_type: str = "split"
+    double_precision_rope: bool = True
+
+    positional_embedding_theta: float = 10000.0
+    positional_embedding_max_pos: list[int] = field(
+        default_factory=lambda: [20, 2048, 2048])
+    timestep_scale_multiplier: int = 1000
+    use_middle_indices_grid: bool = True
+
+    # Patchification (video-only path)
+    patch_size: tuple[int, int, int] = (1, 1, 1)
+    num_channels_latents: int = 128
+    in_channels: int | None = None
+    out_channels: int | None = None
+
+    # Audio defaults (reserved for joint AV ports)
+    audio_num_attention_heads: int = 32
+    audio_attention_head_dim: int = 64
+    audio_in_channels: int = 128
+    audio_out_channels: int = 128
+    audio_cross_attention_dim: int = 2048
+    audio_positional_embedding_max_pos: list[int] = field(
+        default_factory=lambda: [20])
+    av_ca_timestep_scale_multiplier: int = 1
+
+    def __post_init__(self):
+        super().__post_init__()
+        patch_volume = self.patch_size[0] * self.patch_size[
+            1] * self.patch_size[2]
+        if self.in_channels is None:
+            self.in_channels = self.num_channels_latents * patch_volume
+        if self.out_channels is None:
+            self.out_channels = self.in_channels
+
+
+@dataclass
+class LTX2VideoConfig(DiTConfig):
+    """Main configuration for LTX-2 transformer."""
+
+    arch_config: DiTArchConfig = field(default_factory=LTX2VideoArchConfig)
+    prefix: str = "ltx2"
@@ -1,4 +1,6 @@
 from dataclasses import dataclass, field
+
+import torch
 from fastvideo.configs.models.dits.wanvideo import WanVideoArchConfig, WanVideoConfig


@@ -8,8 +10,8 @@ class MatrixGameWanVideoArchConfig(WanVideoArchConfig):
    # because MatrixGame checkpoints already have patch_embedding.proj format
    param_names_mapping: dict = field(
        default_factory=lambda: {
-            # Removed: r"^patch_embedding\.(.*)$": r"patch_embedding.proj.\1"
-            # because checkpoint already has correct format
+            r"^patch_embedding\.(?!proj\.)(.*)$":
+            r"patch_embedding.proj.\1",
            r"^condition_embedder\.text_embedder\.linear_1\.(.*)$":
            r"condition_embedder.text_embedder.fc_in.\1",
            r"^condition_embedder\.text_embedder\.linear_2\.(.*)$":
@@ -76,8 +78,14 @@ class MatrixGameWanVideoArchConfig(WanVideoArchConfig):
    image_dim: int = 1280


+def _is_transformer_block(param_name: str, module: torch.nn.Module) -> bool:
+    return bool("blocks" in param_name and param_name.split(".")[-1].isdigit())
+
+
@dataclass
 class MatrixGameWanVideoConfig(WanVideoConfig):
    arch_config: MatrixGameWanVideoArchConfig = field(
        default_factory=MatrixGameWanVideoArchConfig)
    prefix: str = "Wan"
+    _compile_conditions: list = field(
+        default_factory=lambda: [_is_transformer_block])
@@ -7,10 +7,14 @@ from fastvideo.configs.models.encoders.clip import (
 from fastvideo.configs.models.encoders.llama import LlamaConfig
 from fastvideo.configs.models.encoders.t5 import T5Config, T5LargeConfig
 from fastvideo.configs.models.encoders.qwen2_5 import Qwen2_5_VLConfig
+from fastvideo.configs.models.encoders.siglip import SiglipVisionConfig
+from fastvideo.configs.models.encoders.reason1 import Reason1ArchConfig, Reason1Config
+from fastvideo.configs.models.encoders.gemma import LTX2GemmaConfig

 __all__ = [
    "EncoderConfig", "TextEncoderConfig", "ImageEncoderConfig",
    "BaseEncoderOutput", "CLIPTextConfig", "CLIPVisionConfig",
    "WAN2_1ControlCLIPVisionConfig", "LlamaConfig", "T5Config", "T5LargeConfig",
-    "Qwen2_5_VLConfig"
+    "Qwen2_5_VLConfig", "Reason1ArchConfig", "Reason1Config", "LTX2GemmaConfig",
+    "SiglipVisionConfig"
 ]
@@ -87,10 +87,10 @@ class CLIPVisionConfig(ImageEncoderConfig):
    arch_config: ImageEncoderArchConfig = field(
        default_factory=CLIPVisionArchConfig)

-    num_hidden_layers_override: int | None = None
+    num_hidden_layers_override: int | None = 31
    require_post_norm: bool | None = None
-    enable_scale: bool = True
-    is_causal: bool = True
+    enable_scale: bool = False
+    is_causal: bool = False
    prefix: str = "clip"


@@ -0,0 +1,48 @@
+# SPDX-License-Identifier: Apache-2.0
+from dataclasses import dataclass, field
+
+from fastvideo.configs.models.encoders.base import (
+    TextEncoderArchConfig,
+    TextEncoderConfig,
+)
+
+
+@dataclass
+class LTX2GemmaArchConfig(TextEncoderArchConfig):
+    architectures: list[str] = field(
+        default_factory=lambda: ["LTX2GemmaTextEncoderModel"])
+    hidden_size: int = 3840
+    num_hidden_layers: int = 48
+    num_attention_heads: int = 30
+    text_len: int = 1024
+    pad_token_id: int = 0
+    eos_token_id: int = 2
+
+    gemma_model_path: str = ""
+    gemma_dtype: str = "bfloat16"
+    padding_side: str = "left"
+
+    feature_extractor_in_features: int = 3840 * 49
+    feature_extractor_out_features: int = 3840
+
+    connector_num_attention_heads: int = 30
+    connector_attention_head_dim: int = 128
+    connector_num_layers: int = 2
+    connector_positional_embedding_theta: float = 10000.0
+    connector_positional_embedding_max_pos: list[int] = field(
+        default_factory=lambda: [4096])
+    connector_rope_type: str = "split"
+    connector_double_precision_rope: bool = False
+    connector_num_learnable_registers: int | None = 128
+
+    def __post_init__(self) -> None:
+        super().__post_init__()
+        self.tokenizer_kwargs["padding"] = "max_length"
+
+
+@dataclass
+class LTX2GemmaConfig(TextEncoderConfig):
+    arch_config: TextEncoderArchConfig = field(
+        default_factory=LTX2GemmaArchConfig)
+
+    prefix: str = "ltx2_gemma"
@@ -0,0 +1,72 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Config for Reason1 (Qwen2.5-VL) text encoder."""
+
+from dataclasses import dataclass, field
+from typing import Any
+
+from fastvideo.configs.models.encoders.base import TextEncoderArchConfig, TextEncoderConfig
+
+
+@dataclass
+class Reason1ArchConfig(TextEncoderArchConfig):
+    """Architecture settings (defaults match Qwen2.5-VL-7B-Instruct)."""
+
+    architectures: list[str] = field(
+        default_factory=lambda: ["Qwen2_5_VLForConditionalGeneration"])
+    model_type: str = "qwen2_5_vl"
+
+    vocab_size: int = 152064
+    hidden_size: int = 3584
+    num_hidden_layers: int = 28
+    num_attention_heads: int = 28
+    num_key_value_heads: int = 4
+    intermediate_size: int = 18944
+
+    text_len: int = 512
+    hidden_state_skip_layer: int = 0
+    bos_token_id: int = 151643
+    pad_token_id: int = 151643
+    eos_token_id: int = 151645
+
+    image_token_id: int = 151655
+    video_token_id: int = 151656
+    vision_token_id: int = 151654
+    vision_start_token_id: int = 151652
+    vision_end_token_id: int = 151653
+
+    vision_config: dict[str, Any] | None = None
+
+    rope_theta: float = 1000000.0
+    rope_scaling: dict[str, Any] | None = field(default_factory=lambda: {
+        "type": "mrope",
+        "mrope_section": [16, 24, 24]
+    })
+    max_position_embeddings: int = 128000
+    max_window_layers: int = 28
+
+    embedding_concat_strategy: str = "mean_pooling"
+    n_layers_per_group: int = 5
+    num_embedding_padding_tokens: int = 512
+
+    attention_dropout: float = 0.0
+    hidden_act: str = "silu"
+    initializer_range: float = 0.02
+    rms_norm_eps: float = 1e-6
+
+    use_sliding_window: bool = False
+    sliding_window: int = 32768
+
+    tie_word_embeddings: bool = False
+    use_cache: bool = False
+    output_hidden_states: bool = True
+
+    torch_dtype: str = "bfloat16"
+    _attn_implementation: str = "flash_attention_2"
+
+
+@dataclass
+class Reason1Config(TextEncoderConfig):
+    """Reason1 text encoder config."""
+
+    arch_config: Reason1ArchConfig = field(default_factory=Reason1ArchConfig)
+    tokenizer_type: str = "Qwen/Qwen2.5-VL-7B-Instruct"
@@ -0,0 +1,53 @@
+# SPDX-License-Identifier: Apache-2.0
+"""SigLIP vision encoder configuration for FastVideo."""
+
+from dataclasses import dataclass, field
+
+from fastvideo.configs.models.encoders.base import (ImageEncoderArchConfig,
+                                                    ImageEncoderConfig)
+
+
+@dataclass
+class SiglipVisionArchConfig(ImageEncoderArchConfig):
+    """Architecture configuration for SigLIP vision encoder.
+    
+    Fields match the config.json from HuggingFace SigLIP checkpoints.
+    """
+
+    # From config.json
+    architectures: list[str] = field(
+        default_factory=lambda: ["SiglipVisionModel"])
+    attention_dropout: float = 0.0
+    dtype: str | None = None
+    hidden_act: str = "gelu_pytorch_tanh"
+    hidden_size: int = 1152
+    image_size: int = 384
+    intermediate_size: int = 4304
+    layer_norm_eps: float = 1e-6
+    model_type: str = "siglip_vision_model"
+    num_attention_heads: int = 16
+    num_channels: int = 3
+    num_hidden_layers: int = 27
+    patch_size: int = 14
+
+    # FastVideo specific - QKV fusion mapping
+    stacked_params_mapping: list = field(default_factory=lambda: [
+        ("qkv_proj", "q_proj", "q"),
+        ("qkv_proj", "k_proj", "k"),
+        ("qkv_proj", "v_proj", "v"),
+    ])
+
+
+@dataclass
+class SiglipVisionConfig(ImageEncoderConfig):
+    """Configuration for SigLIP vision encoder."""
+
+    arch_config: ImageEncoderArchConfig = field(
+        default_factory=SiglipVisionArchConfig)
+
+    # FastVideo specific
+    num_hidden_layers_override: int | None = None
+    require_post_norm: bool | None = None
+    enable_scale: bool = True
+    is_causal: bool = False
+    prefix: str = "siglip"
@@ -0,0 +1,6 @@
+from fastvideo.configs.models.upsamplers.hunyuan15 import SRTo720pUpsamplerConfig, SRTo1080pUpsamplerConfig
+from fastvideo.configs.models.upsamplers.base import UpsamplerConfig
+
+__all__ = [
+    "SRTo720pUpsamplerConfig", "SRTo1080pUpsamplerConfig", "UpsamplerConfig"
+]
@@ -0,0 +1,7 @@
+from dataclasses import dataclass
+from fastvideo.configs.models.base import ModelConfig
+
+
+@dataclass
+class UpsamplerConfig(ModelConfig):
+    pass
@@ -0,0 +1,20 @@
+from dataclasses import dataclass
+from fastvideo.configs.models.upsamplers.base import UpsamplerConfig
+
+
+@dataclass
+class SRTo720pUpsamplerConfig(UpsamplerConfig):
+    in_channels: int = 0
+    out_channels: int = 0
+    hidden_channels: int = 64
+    num_blocks: int = 6
+    global_residual: bool = False
+
+
+@dataclass
+class SRTo1080pUpsamplerConfig(UpsamplerConfig):
+    z_channels: int = 0
+    out_channels: int = 0
+    block_out_channels: tuple[int, ...] = (0, 0)
+    num_res_blocks: int = 2
+    is_residual: bool = False
@@ -1,6 +1,8 @@
 from fastvideo.configs.models.vaes.cosmosvae import CosmosVAEConfig
+from fastvideo.configs.models.vaes.cosmos2_5vae import Cosmos25VAEConfig
 from fastvideo.configs.models.vaes.hunyuanvae import HunyuanVAEConfig
 from fastvideo.configs.models.vaes.hunyuan15vae import Hunyuan15VAEConfig
+from fastvideo.configs.models.vaes.ltx2vae import LTX2VAEConfig
 from fastvideo.configs.models.vaes.stepvideovae import StepVideoVAEConfig
 from fastvideo.configs.models.vaes.wanvae import WanVAEConfig

@@ -9,5 +11,7 @@ __all__ = [
    "WanVAEConfig",
    "StepVideoVAEConfig",
    "CosmosVAEConfig",
+    "Cosmos25VAEConfig",
    "Hunyuan15VAEConfig",
+    "LTX2VAEConfig",
 ]
@@ -0,0 +1,223 @@
+"""Cosmos 2.5 (Wan2.1-style) VAE config and checkpoint-key mapping."""
+
+from __future__ import annotations
+
+import re
+from dataclasses import dataclass, field
+
+import torch
+
+from fastvideo.configs.models.vaes.base import VAEArchConfig, VAEConfig
+
+
+@dataclass
+class Cosmos25VAEArchConfig(VAEArchConfig):
+    _name_or_path: str = ""
+    base_dim: int = 96
+    decoder_base_dim: int | None = None
+    z_dim: int = 16
+    dim_mult: tuple[int, ...] = (1, 2, 4, 4)
+    num_res_blocks: int = 2
+    attn_scales: tuple[float, ...] = ()
+    temperal_downsample: tuple[bool, ...] = (False, True, True)
+    dropout: float = 0.0
+    is_residual: bool = False
+    in_channels: int = 3
+    out_channels: int = 3
+    patch_size: int | None = None
+    scale_factor_temporal: int = 4
+    scale_factor_spatial: int = 8
+    clip_output: bool = True
+
+    latents_mean: tuple[float, ...] = (
+        -0.7571,
+        -0.7089,
+        -0.9113,
+        0.1075,
+        -0.1745,
+        0.9653,
+        -0.1517,
+        1.5508,
+        0.4134,
+        -0.0715,
+        0.5517,
+        -0.3632,
+        -0.1922,
+        -0.9497,
+        0.2503,
+        -0.2921,
+    )
+    latents_std: tuple[float, ...] = (
+        2.8184,
+        1.4541,
+        2.3275,
+        2.6558,
+        1.2196,
+        1.7708,
+        2.6052,
+        2.0743,
+        3.2687,
+        2.1526,
+        2.8652,
+        1.5579,
+        1.6382,
+        1.1253,
+        2.8251,
+        1.9160,
+    )
+
+    # Simple 1:1 renames. More complex decoder remapping is handled by
+    # `map_official_key()`.
+    param_names_mapping: dict[str, str] = field(
+        default_factory=lambda: {
+            r"^conv1\.(.*)$": r"quant_conv.\1",
+            r"^conv2\.(.*)$": r"post_quant_conv.\1",
+            r"^encoder\.conv1\.(.*)$": r"encoder.conv_in.\1",
+            r"^decoder\.conv1\.(.*)$": r"decoder.conv_in.\1",
+            r"^encoder\.head\.0\.gamma$": r"encoder.norm_out.gamma",
+            r"^encoder\.head\.2\.(.*)$": r"encoder.conv_out.\1",
+            r"^decoder\.head\.0\.gamma$": r"decoder.norm_out.gamma",
+            r"^decoder\.head\.2\.(.*)$": r"decoder.conv_out.\1",
+        })
+
+    @staticmethod
+    def map_official_key(key: str) -> str | None:
+        """Map a single official checkpoint key into FastVideo key space."""
+
+        def map_residual_subkey(prefix: str, sub: str) -> str | None:
+            if re.match(r"^residual\.0\.gamma$", sub):
+                return f"{prefix}.norm1.gamma"
+            m = re.match(r"^residual\.2\.(weight|bias)$", sub)
+            if m:
+                return f"{prefix}.conv1.{m.group(1)}"
+            if re.match(r"^residual\.3\.gamma$", sub):
+                return f"{prefix}.norm2.gamma"
+            m = re.match(r"^residual\.6\.(weight|bias)$", sub)
+            if m:
+                return f"{prefix}.conv2.{m.group(1)}"
+            m = re.match(r"^shortcut\.(weight|bias)$", sub)
+            if m:
+                return f"{prefix}.conv_shortcut.{m.group(1)}"
+            return None
+
+        def map_attn_subkey(prefix: str, sub: str) -> str | None:
+            if re.match(r"^norm\.gamma$", sub):
+                return f"{prefix}.norm.gamma"
+            m = re.match(r"^to_qkv\.(weight|bias)$", sub)
+            if m:
+                return f"{prefix}.to_qkv.{m.group(1)}"
+            m = re.match(r"^proj\.(weight|bias)$", sub)
+            if m:
+                return f"{prefix}.proj.{m.group(1)}"
+            return None
+
+        def map_resample_subkey(prefix: str, sub: str) -> str | None:
+            m = re.match(r"^resample\.1\.(weight|bias)$", sub)
+            if m:
+                return f"{prefix}.resample.1.{m.group(1)}"
+            m = re.match(r"^time_conv\.(weight|bias)$", sub)
+            if m:
+                return f"{prefix}.time_conv.{m.group(1)}"
+            return None
+
+        m = re.match(r"^conv1\.(weight|bias)$", key)
+        if m:
+            return f"quant_conv.{m.group(1)}"
+        m = re.match(r"^conv2\.(weight|bias)$", key)
+        if m:
+            return f"post_quant_conv.{m.group(1)}"
+        m = re.match(r"^(encoder|decoder)\.conv1\.(weight|bias)$", key)
+        if m:
+            return f"{m.group(1)}.conv_in.{m.group(2)}"
+        m = re.match(r"^(encoder|decoder)\.head\.0\.gamma$", key)
+        if m:
+            return f"{m.group(1)}.norm_out.gamma"
+        m = re.match(r"^(encoder|decoder)\.head\.2\.(weight|bias)$", key)
+        if m:
+            return f"{m.group(1)}.conv_out.{m.group(2)}"
+
+        m = re.match(r"^(encoder|decoder)\.middle\.0\.(.*)$", key)
+        if m:
+            return map_residual_subkey(f"{m.group(1)}.mid_block.resnets.0",
+                                       m.group(2))
+        m = re.match(r"^(encoder|decoder)\.middle\.1\.(.*)$", key)
+        if m:
+            return map_attn_subkey(f"{m.group(1)}.mid_block.attentions.0",
+                                   m.group(2))
+        m = re.match(r"^(encoder|decoder)\.middle\.2\.(.*)$", key)
+        if m:
+            return map_residual_subkey(f"{m.group(1)}.mid_block.resnets.1",
+                                       m.group(2))
+
+        m = re.match(r"^encoder\.downsamples\.(\d+)\.(.*)$", key)
+        if m:
+            idx = int(m.group(1))
+            sub = m.group(2)
+            if sub.startswith("residual.") or sub.startswith("shortcut."):
+                return map_residual_subkey(f"encoder.down_blocks.{idx}", sub)
+            if sub.startswith("resample.") or sub.startswith("time_conv."):
+                return map_resample_subkey(f"encoder.down_blocks.{idx}", sub)
+            return None
+
+        m = re.match(r"^decoder\.upsamples\.(\d+)\.(.*)$", key)
+        if m:
+            uidx = int(m.group(1))
+            sub = m.group(2)
+
+            if uidx in (0, 1, 2):
+                block_i, res_i = 0, uidx
+            elif uidx == 3:
+                block_i, res_i = 0, None
+            elif uidx in (4, 5, 6):
+                block_i, res_i = 1, uidx - 4
+            elif uidx == 7:
+                block_i, res_i = 1, None
+            elif uidx in (8, 9, 10):
+                block_i, res_i = 2, uidx - 8
+            elif uidx == 11:
+                block_i, res_i = 2, None
+            elif uidx in (12, 13, 14):
+                block_i, res_i = 3, uidx - 12
+            else:
+                return None
+
+            if res_i is None:
+                return map_resample_subkey(
+                    f"decoder.up_blocks.{block_i}.upsamplers.0",
+                    sub,
+                )
+
+            return map_residual_subkey(
+                f"decoder.up_blocks.{block_i}.resnets.{res_i}",
+                sub,
+            )
+
+        return None
+
+    temporal_compression_ratio: int = 4
+    spatial_compression_ratio: int = 8
+
+    def __post_init__(self):
+        self.scaling_factor: torch.Tensor = 1.0 / torch.tensor(
+            self.latents_std).view(1, self.z_dim, 1, 1, 1)
+        self.shift_factor: torch.Tensor = torch.tensor(self.latents_mean).view(
+            1, self.z_dim, 1, 1, 1)
+        self.temporal_compression_ratio = self.scale_factor_temporal
+        self.spatial_compression_ratio = self.scale_factor_spatial
+
+
+@dataclass
+class Cosmos25VAEConfig(VAEConfig):
+    """Cosmos2.5 VAE config."""
+
+    arch_config: Cosmos25VAEArchConfig = field(
+        default_factory=Cosmos25VAEArchConfig)
+
+    use_feature_cache: bool = True
+    use_tiling: bool = False
+    use_temporal_tiling: bool = False
+    use_parallel_tiling: bool = False
+
+    def __post_init__(self):
+        self.blend_num_frames = (self.tile_sample_min_num_frames -
+                                 self.tile_sample_stride_num_frames) * 2
@@ -0,0 +1,45 @@
+# SPDX-License-Identifier: Apache-2.0
+"""
+LTX-2 VAE configuration.
+"""
+
+from dataclasses import dataclass, field
+
+from fastvideo.configs.models.vaes.base import VAEArchConfig, VAEConfig
+
+
+@dataclass
+class LTX2VAEArchConfig(VAEArchConfig):
+    # Mirrors LTX-2 safetensors metadata config under "vae"
+    _class_name: str = "CausalVideoAutoencoder"
+    dims: int = 3
+    in_channels: int = 3
+    out_channels: int = 3
+    latent_channels: int = 128
+    encoder_blocks: list = field(default_factory=list)
+    decoder_blocks: list = field(default_factory=list)
+    patch_size: int = 4
+    norm_layer: str = "pixel_norm"
+    latent_log_var: str = "uniform"
+    encoder_spatial_padding_mode: str = "zeros"
+    decoder_spatial_padding_mode: str = "reflect"
+    causal_decoder: bool = False
+    timestep_conditioning: bool = True
+    use_quant_conv: bool = False
+    scaling_factor: float = 1.0
+    normalize_latent_channels: bool = False
+
+    # Match FastVideo naming for compression ratios (LTX-2 default)
+    temporal_compression_ratio: int = 8
+    spatial_compression_ratio: int = 32
+
+
+@dataclass
+class LTX2VAEConfig(VAEConfig):
+    arch_config: VAEArchConfig = field(default_factory=LTX2VAEArchConfig)
+
+    # LTX-2 tiling defaults (match ltx_core.video_vae.TilingConfig.default()).
+    ltx2_spatial_tile_size_in_pixels: int = 512
+    ltx2_spatial_tile_overlap_in_pixels: int = 64
+    ltx2_temporal_tile_size_in_frames: int = 64
+    ltx2_temporal_tile_overlap_in_frames: int = 24
@@ -1,10 +1,12 @@
 from fastvideo.configs.pipelines.base import (PipelineConfig,
                                              SlidingTileAttnConfig)
 from fastvideo.configs.pipelines.cosmos import CosmosConfig
+from fastvideo.configs.pipelines.cosmos2_5 import Cosmos25Config
 from fastvideo.configs.pipelines.hunyuan import FastHunyuanConfig, HunyuanConfig
 from fastvideo.configs.pipelines.hunyuan15 import Hunyuan15T2V480PConfig, Hunyuan15T2V720PConfig
-from fastvideo.configs.pipelines.registry import (
-    get_pipeline_config_cls_from_name)
+from fastvideo.configs.pipelines.hyworld import HYWorldConfig
+from fastvideo.configs.pipelines.ltx2 import LTX2T2VConfig
+from fastvideo.registry import get_pipeline_config_cls_from_name
 from fastvideo.configs.pipelines.stepvideo import StepVideoT2VConfig
 from fastvideo.configs.pipelines.wan import (SelfForcingWanT2V480PConfig,
                                             WanI2V480PConfig, WanI2V720PConfig,
@@ -15,5 +17,6 @@ __all__ = [
    "Hunyuan15T2V480PConfig", "Hunyuan15T2V720PConfig", "SlidingTileAttnConfig",
    "WanT2V480PConfig", "WanI2V480PConfig", "WanT2V720PConfig",
    "WanI2V720PConfig", "StepVideoT2VConfig", "SelfForcingWanT2V480PConfig",
-    "CosmosConfig", "get_pipeline_config_cls_from_name"
+    "CosmosConfig", "Cosmos25Config", "LTX2T2VConfig", "HYWorldConfig",
+    "get_pipeline_config_cls_from_name"
 ]
@@ -8,7 +8,7 @@ from typing import Any, cast
 import torch

 from fastvideo.configs.models import (DiTConfig, EncoderConfig, ModelConfig,
-                                      VAEConfig)
+                                      VAEConfig, UpsamplerConfig)
 from fastvideo.configs.models.encoders import BaseEncoderOutput
 from fastvideo.configs.utils import update_config_from_args
 from fastvideo.logger import init_logger
@@ -44,12 +44,15 @@ class PipelineConfig:
    # Video generation parameters
    embedded_cfg_scale: float = 6.0
    flow_shift: float | None = None
+    flow_shift_sr: float | None = None
    disable_autocast: bool = False
    is_causal: bool = False

    # Model configuration
    dit_config: DiTConfig = field(default_factory=DiTConfig)
    dit_precision: str = "bf16"
+    upsampler_config: UpsamplerConfig = field(default_factory=UpsamplerConfig)
+    upsampler_precision: str = "fp32"

    # VAE configuration
    vae_config: VAEConfig = field(default_factory=VAEConfig)
@@ -216,6 +219,24 @@ class PipelineConfig:
            "Comma-separated list of denoising steps (e.g., '1000,757,522')",
        )

+        # STA (Sliding Tile Attention) parameters
+        parser.add_argument(
+            f"--{prefix_with_dot}STA-mode",
+            type=str,
+            dest=f"{prefix_with_dot.replace('-', '_')}STA_mode",
+            default=PipelineConfig.STA_mode.value,
+            choices=[mode.value for mode in STA_Mode],
+            help=
+            "STA mode: STA_inference, STA_searching, STA_tuning, STA_tuning_cfg, None",
+        )
+        parser.add_argument(
+            f"--{prefix_with_dot}skip-time-steps",
+            type=int,
+            dest=f"{prefix_with_dot.replace('-', '_')}skip_time_steps",
+            default=PipelineConfig.skip_time_steps,
+            help="Number of time steps to warmup (full attention) for STA",
+        )
+
        # Add VAE configuration arguments
        from fastvideo.configs.models.vaes.base import VAEConfig
        VAEConfig.add_cli_args(parser, prefix=f"{prefix_with_dot}vae-config")
@@ -245,8 +266,7 @@ class PipelineConfig:
        """
        use the pipeline class setting from model_path to match the pipeline config
        """
-        from fastvideo.configs.pipelines.registry import (
-            get_pipeline_config_cls_from_name)
+        from fastvideo.registry import get_pipeline_config_cls_from_name
        pipeline_config_cls = get_pipeline_config_cls_from_name(model_path)

        return cast(PipelineConfig, pipeline_config_cls(model_path=model_path))
@@ -260,8 +280,7 @@ class PipelineConfig:
        kwargs: dictionary of kwargs
        config_cli_prefix: prefix of CLI arguments for this PipelineConfig instance
        """
-        from fastvideo.configs.pipelines.registry import (
-            get_pipeline_config_cls_from_name)
+        from fastvideo.registry import get_pipeline_config_cls_from_name

        prefix_with_dot = f"{config_cli_prefix}." if (config_cli_prefix.strip()
                                                      != "") else ""
@@ -298,6 +317,12 @@ class PipelineConfig:
        # 4. Update PipelineConfig from CLI arguments if provided
        kwargs[prefix_with_dot + 'model_path'] = model_path
        pipeline_config.update_config_from_dict(kwargs, config_cli_prefix)
+
+        # Convert STA_mode string to enum if necessary
+        if isinstance(pipeline_config.STA_mode, str) and not isinstance(
+                pipeline_config.STA_mode, STA_Mode):
+            pipeline_config.STA_mode = STA_Mode(pipeline_config.STA_mode)
+
        return pipeline_config

    def check_pipeline_config(self) -> None:
@@ -0,0 +1,89 @@
+# SPDX-License-Identifier: Apache-2.0
+from collections.abc import Callable
+from dataclasses import dataclass, field
+
+import torch
+
+from fastvideo.configs.models import DiTConfig, EncoderConfig, VAEConfig
+from fastvideo.configs.models.dits import Cosmos25VideoConfig
+from fastvideo.configs.models.dits.cosmos2_5 import Cosmos25ArchConfig
+from fastvideo.configs.models.encoders import BaseEncoderOutput
+from fastvideo.configs.models.encoders.reason1 import Reason1Config, Reason1ArchConfig
+from fastvideo.configs.models.vaes import Cosmos25VAEConfig
+from fastvideo.configs.pipelines.base import PipelineConfig, STA_Mode
+
+
+def _identity_preprocess_text(prompt: str) -> str:
+    return prompt
+
+
+def reason1_postprocess_text(outputs: BaseEncoderOutput) -> torch.Tensor:
+    hidden_states = getattr(outputs, "hidden_states", None)
+    if hidden_states is None:
+        raise ValueError("Reason1 postprocess requires outputs.hidden_states")
+
+    hs = list(hidden_states)[1:]
+    normed = []
+    for h in hs:
+        h = h.float()
+        h = (h - h.mean(dim=-1, keepdim=True)) / (h.std(dim=-1, keepdim=True) +
+                                                  1e-8)
+        normed.append(h)
+    return torch.cat(normed, dim=-1).to(hidden_states[0].dtype)
+
+
+@dataclass
+class Cosmos25Config(PipelineConfig):
+    """Configuration for Cosmos 2.5 (Predict2.5) video generation pipeline."""
+
+    dit_config: DiTConfig = field(default_factory=lambda: Cosmos25VideoConfig(
+        arch_config=Cosmos25ArchConfig(
+            num_attention_heads=16,
+            attention_head_dim=128,
+            in_channels=16,
+            out_channels=16,
+            num_layers=28,
+            patch_size=[1, 2, 2],
+            max_size=[128, 240, 240],
+            rope_scale=[1.0, 3.0, 3.0],
+            text_embed_dim=1024,
+            mlp_ratio=4.0,
+            adaln_lora_dim=256,
+            use_adaln_lora=True,
+            concat_padding_mask=True,
+            extra_pos_embed_type=None,
+            use_crossattn_projection=True,
+            rope_enable_fps_modulation=False,
+            qk_norm="rms_norm",
+        )))
+
+    vae_config: VAEConfig = field(default_factory=Cosmos25VAEConfig)
+
+    text_encoder_configs: tuple[EncoderConfig, ...] = field(
+        default_factory=lambda: (Reason1Config(arch_config=Reason1ArchConfig(
+            embedding_concat_strategy="full_concat")), ))
+
+    preprocess_text_funcs: tuple[Callable[[str], str], ...] = field(
+        default_factory=lambda: (_identity_preprocess_text, ))
+    postprocess_text_funcs: tuple[Callable[[BaseEncoderOutput], torch.Tensor],
+                                  ...] = field(default_factory=lambda:
+                                               (reason1_postprocess_text, ))
+
+    dit_precision: str = "bf16"
+    vae_precision: str = "bf16"
+    text_encoder_precisions: tuple[str, ...] = field(
+        default_factory=lambda: ("bf16", ))
+
+    embedded_cfg_scale: float = 0.0
+    flow_shift: float = 5.0
+
+    vae_tiling: bool = False
+    vae_sp: bool = False
+
+    STA_mode: STA_Mode = STA_Mode.NONE
+    skip_time_steps: int = 0
+
+    def __post_init__(self):
+        self.vae_config.load_encoder = True
+        self.vae_config.load_decoder = True
+        self._vae_latent_dim = 16
@@ -11,7 +11,8 @@ from fastvideo.configs.models.dits import HunyuanVideo15Config
 from fastvideo.configs.models.encoders import (BaseEncoderOutput,
                                               Qwen2_5_VLConfig, T5Config)
 from fastvideo.configs.models.vaes import Hunyuan15VAEConfig
-from fastvideo.configs.pipelines.base import PipelineConfig
+from fastvideo.configs.models.upsamplers import SRTo720pUpsamplerConfig, SRTo1080pUpsamplerConfig
+from fastvideo.configs.pipelines.base import PipelineConfig, UpsamplerConfig

 PROMPT_TEMPLATE_TOKEN_LENGTH = 108

@@ -131,9 +132,35 @@ class Hunyuan15T2V480PConfig(PipelineConfig):
        self.vae_config.load_decoder = True


+@dataclass
+class Hunyuan15I2V480PStepDistilledConfig(Hunyuan15T2V480PConfig):
+    flow_shift: int = 7
+
+
@dataclass
 class Hunyuan15T2V720PConfig(Hunyuan15T2V480PConfig):
    """Base configuration for HunYuan pipeline architecture."""

    # HunyuanConfig-specific parameters with defaults
    flow_shift: int = 9
+
+
+@dataclass
+class Hunyuan15I2V720PConfig(Hunyuan15T2V720PConfig):
+    """Base configuration for HunYuan pipeline architecture."""
+
+    # HunyuanConfig-specific parameters with defaults
+    flow_shift: int = 7
+
+
+@dataclass
+class Hunyuan15SR1080PConfig(Hunyuan15T2V720PConfig):
+    """Base configuration for HunYuan pipeline architecture."""
+
+    # HunyuanConfig-specific parameters with defaults
+    flow_shift: int = 7
+    flow_shift_sr: int = 2
+    upsampler_config: tuple[UpsamplerConfig, ...] = field(
+        default_factory=lambda:
+        (SRTo720pUpsamplerConfig(), SRTo1080pUpsamplerConfig()))
+    upsampler_precision: str = "fp32"
@@ -0,0 +1,29 @@
+# SPDX-License-Identifier: Apache-2.0
+from dataclasses import dataclass, field
+
+from fastvideo.configs.models import DiTConfig, EncoderConfig
+from fastvideo.configs.models.dits import HYWorldConfig as HYWorldDiTConfig
+from fastvideo.configs.models.encoders import SiglipVisionConfig
+from fastvideo.configs.pipelines.hunyuan15 import Hunyuan15T2V480PConfig
+
+
+@dataclass
+class HYWorldConfig(Hunyuan15T2V480PConfig):
+    """Base configuration for HYWorld pipeline architecture."""
+
+    # HYWorldConfig-specific parameters with defaults
+    dit_config: DiTConfig = field(default_factory=HYWorldDiTConfig)
+
+    # SigLIP image encoder for I2V
+    image_encoder_config: EncoderConfig = field(
+        default_factory=SiglipVisionConfig)
+    image_encoder_precision: str = "fp16"
+    # vae_precision: str = "fp32"
+
+    # Text encoding
+    text_encoder_precisions: tuple[str, ...] = field(
+        default_factory=lambda: ("fp16", "fp32"))
+
+    def __post_init__(self):
+        self.vae_config.load_encoder = True
+        self.vae_config.load_decoder = True
@@ -17,11 +17,7 @@ from fastvideo.configs.pipelines.base import PipelineConfig

@dataclass
 class LongCatDiTArchConfig(DiTArchConfig):
-    """Extended DiTArchConfig with LongCat-specific fields.
-    
-    NOTE: This is for Phase 1 wrapper compatibility. For native model (Phase 2),
-    use LongCatVideoConfig from fastvideo.configs.models.dits.longcat instead.
-    """
+    """Extended DiTArchConfig with LongCat-specific fields."""
    # LongCat-specific architecture parameters
    adaln_tembed_dim: int = 512
    caption_channels: int = 4096
@@ -88,20 +84,16 @@ def umt5_postprocess_text(outputs: BaseEncoderOutput) -> torch.Tensor:

@dataclass
 class LongCatT2V480PConfig(PipelineConfig):
-    """Configuration for LongCat pipeline (480p) aligned to LongCat-Video modules.
+    """Configuration for LongCat pipeline (480p).

    Components expected by loaders:
      - tokenizer: AutoTokenizer
      - text_encoder: UMT5EncoderModel
-      - transformer: LongCatVideoTransformer3DModel (Phase 1 wrapper)
-                  OR LongCatTransformer3DModel (Phase 2 native)
+      - transformer: LongCatTransformer3DModel
      - vae: AutoencoderKLWan (Wan VAE, 4x8 compression)
      - scheduler: FlowMatchEulerDiscreteScheduler
    """

-    # DiT config with LongCat-specific arch_config
-    # NOTE: For Phase 1 wrapper, uses LongCatDiTArchConfig
-    # For Phase 2 native model, can use LongCatVideoConfig directly
    dit_config: DiTConfig = field(
        default_factory=lambda: DiTConfig(arch_config=LongCatDiTArchConfig()))

@@ -0,0 +1,50 @@
+# SPDX-License-Identifier: Apache-2.0
+from collections.abc import Callable
+from dataclasses import dataclass, field
+
+import torch
+
+from fastvideo.configs.models import (DiTConfig, EncoderConfig, ModelConfig,
+                                      LTX2AudioDecoderConfig, LTX2VocoderConfig,
+                                      VAEConfig)
+from fastvideo.configs.models.dits import LTX2VideoConfig
+from fastvideo.configs.models.encoders import BaseEncoderOutput, LTX2GemmaConfig
+from fastvideo.configs.models.vaes import LTX2VAEConfig
+from fastvideo.configs.pipelines.base import PipelineConfig, preprocess_text
+
+
+def ltx2_postprocess_text(outputs: BaseEncoderOutput) -> torch.Tensor:
+    return outputs.last_hidden_state
+
+
+@dataclass
+class LTX2T2VConfig(PipelineConfig):
+    """Configuration for LTX-2 T2V pipeline."""
+
+    dit_config: DiTConfig = field(default_factory=LTX2VideoConfig)
+    vae_config: VAEConfig = field(default_factory=LTX2VAEConfig)
+    vae_tiling: bool = True
+    vae_sp: bool = False
+
+    text_encoder_configs: tuple[EncoderConfig, ...] = field(
+        default_factory=lambda: (LTX2GemmaConfig(), ))
+    preprocess_text_funcs: tuple[Callable[[str], str], ...] = field(
+        default_factory=lambda: (preprocess_text, ))
+    postprocess_text_funcs: tuple[Callable[[BaseEncoderOutput], torch.Tensor],
+                                  ...] = field(default_factory=lambda:
+                                               (ltx2_postprocess_text, ))
+
+    dit_precision: str = "bf16"
+    vae_precision: str = "bf16"
+    text_encoder_precisions: tuple[str, ...] = field(
+        default_factory=lambda: ("bf16", ))
+
+    audio_decoder_config: ModelConfig = field(
+        default_factory=LTX2AudioDecoderConfig)
+    vocoder_config: ModelConfig = field(default_factory=LTX2VocoderConfig)
+    audio_decoder_precision: str = "bf16"
+    vocoder_precision: str = "bf16"
+
+    def __post_init__(self) -> None:
+        self.vae_config.load_encoder = False
+        self.vae_config.load_decoder = True
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
JerryZhou54	8cb2ae9d27	ckpt	2026-02-09 05:50:48 +00:00
XOR-op	f76efe798e	[bugfix]: `_compile_conditions` regression (#1077 )	2026-02-06 18:57:43 -08:00
Zhang Peiyuan	09f455233e	[misc] readme small fix (#1076 )	2026-02-06 16:57:48 -08:00
XOR-op	aea300f690	[perf]: use CUDA IPC in multiproc executor to avoid serialization overhead (#1061 )	2026-02-06 19:43:18 -05:00
Hao Zhang	b92219f6a6	more fix and relocate STA arguments to pipeline config (#1073 )	2026-02-06 13:38:35 -08:00
Jinzhe Pan	a321b95a8a	[Fix] remove video ratio limitation (#1069 )	2026-02-05 20:33:08 -08:00
Wei Zhou	98308db7e0	[Feature] [Hy1.5] Support HY1.5 super-resolution pipeline for 1080p videos (#1046 )	2026-02-05 16:39:20 -08:00
Hao Zhang	c1e18f6722	Some minor fixes (#1068 )	2026-02-05 16:28:00 -08:00
William Lin	75e193a2c9	[core] Refactor and centralize our registry for models, pipelines, and sampling params (#1066 )	2026-02-05 14:40:30 -08:00
XOR-op	d6e0a7d0dd	[refactor] Action module (#1065 )	2026-02-05 13:56:35 -08:00
William Linandgemini-code-assist[bot]	7fc5f241da	[misc] Fix naming instruction in runpod.md (#1067 ) Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-02-05 13:45:49 -08:00
Mingjia Huo	aae48a7e90	[feat] HYworld VAE with cache (#1057 )	2026-02-05 04:18:17 -08:00
William Lin	88f38eb0f4	[misc] upgrade torch to 2.10 (#1048 )	2026-02-05 04:15:22 -08:00
Shao Duan	d750b463dc	Added Sequence Parallelism for LTX-2 Distilled (#1036 )	2026-02-04 15:25:30 -08:00
KyleShaoandWill Lin	e10b26a3d8	[feat] Add Cosmos 2.5 I2W/V2W support (staged pipeline + examples) (#1021 ) Co-authored-by: Will Lin <wlsaidhi@gmail.com>	2026-02-04 14:19:30 -08:00
XOR-op	74636ba246	[chore]: use higher precision timestamp in logging (#1062 )	2026-02-04 13:22:59 -08:00
Wei Zhou	7e2f3f14e7	[Bugfix] [Wan I2V] Fix CLIP Image encoder config (#1063 )	2026-02-04 13:20:19 -08:00
Kaiqin Kong	caa1c402ba	[bugfix] Double Normalization in Preprocessing Dataset (#1055 )	2026-01-31 11:11:59 -08:00
XOR-op	38a6bd93d3	[chore]: update sageattn3 installation instructions (#1050 )	2026-01-29 15:45:52 -08:00
alexzms	b867ef7e7c	[SP Sharding] Fix SP loss sharding on token axis (thw) with padding; add distributed correctness tests (#1045 ) Fixes the sequence parallel sharding on t, now SP shards on thw	2026-01-27 22:53:04 -08:00
William Lin	3ae58c277a	[docs] Update design overview and add agents tutorial (#1044 )	2026-01-27 15:56:58 -08:00
Kaiqin Kong	0c6862ca55	[feature] Add Matrix Game 2.0 training (#1017 ) The CI tests are quite unstable, but since multiple CI tests indicates that each individual tests are passed, I think we can merge this.	2026-01-26 19:21:43 -08:00
XOR-opandWill Lin	e8c854bcf1	[docs] Offloading instruction (#1022 ) Co-authored-by: Will Lin <wlsaidhi@gmail.com>	2026-01-26 13:40:28 -08:00
William Lin	06860e96fe	[docs] Update runpod instructions (#1043 )	2026-01-26 13:13:54 -08:00
Matthew Noto	1b503554d1	[bugfix] fix torchvision import (#1039 )	2026-01-24 22:37:14 -08:00
Shreejith SGandgemini-code-assist[bot]	351ceb7c59	[bugfix]: handle architectural differences while lora extraction (#1035 ) Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-01-24 15:42:28 -08:00
KyleShao	10875e0d7b	[bugfix] Fix NCCL all_gather contiguity + correct ParallelTiledVAE decode tiling threshold (#1037 )	2026-01-24 15:37:35 -08:00
alexzms	1eaae8a10b	[ci] Increase ci test error threshold (#1038 )	2026-01-24 15:36:10 -08:00
Mingjia Huo	59e00f6164	[feat] Add HY-World1.5-Bidirectional-480P-I2V (#1027 ) VAE requires further improvement, will raise PR in near future.	2026-01-23 14:18:04 -08:00
Junda (David) Su Davids048 gemini-code-assist[bot]	745cc05b10	[bugfix] Allow update timesteps for hy1.5 model. (#1033 ) Co-authored-by: Davids048 <jundasu@ucsd.edu> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-01-22 22:04:53 -08:00
William Lin	c5dc244871	[bugfix] add omegaconf as dep. (#1032 )	2026-01-22 11:59:28 -08:00
alexzms	dbf3917bf4	[fastvideo-kernel] replace map to index with Triton implementation + add vsa benchmark (#1029 )	2026-01-22 11:35:02 -08:00
XOR-op	050f189c95	fix: SP for hunyuanvideo 1.5 (#1026 )	2026-01-21 14:40:06 -08:00
Shao DuanandWill Lin	029216029f	Added LTX-2 Distilled T2V Generation (#1016 ) Co-authored-by: Will Lin <wlsaidhi@gmail.com>	2026-01-21 14:11:39 -08:00
alexzmsandWilliam Lin	31f44110b5	[kernel] [bugfix] [ci] bump v0.2.4. Fix STA output handling, TurboDiffusion CUDA norm dtypes for fastvideo-kernel unit tests. (#1020 ) Co-authored-by: William Lin <SolitaryThinker@users.noreply.github.com>	2026-01-19 17:42:42 -08:00
William Lin	21f3ce6577	[kernel] Fix fastvideo-kernel release workflow (#1019 )	2026-01-17 15:57:46 -08:00
XOR-op	785d123e36	[feat] Hooks API and layerwise offloading for all DiTs (#1006 )	2026-01-17 11:22:02 -08:00
William Lin	d58c551c11	[chore] release fastvideo-kernel 0.2.3 (#1018 )	2026-01-17 02:24:23 -08:00
alexzms	560628709c	[Bug Fix] Add autograd wrapper for block-sparse attention in fastvideo-kernel + fix CMake extension linking (#1015 )	2026-01-16 21:16:43 -08:00
William Lin	0f53b51e6c	[CI] Fix OOM issues in ssim tests (#1011 )	2026-01-16 21:15:20 -08:00
alexzmsandWill Lin	06093a9c4e	[CI] SSIM tests optimization: load all model weights from Modal persistent Volume (#958 ) Co-authored-by: Will Lin <wlsaidhi@gmail.com>	2026-01-16 11:50:43 -08:00
KyleShao	dbddfab6d2	[feat] Introduce Cosmos 2.5 Text2World pipeline (#974 )	2026-01-15 15:09:05 -08:00
William Lin	7188170277	[misc] [bugfix] unpin 'av' in pyproject (#1009 )	2026-01-13 15:40:46 -08:00
XOR-op	b7f69c2c1d	[feat!] Disable FSDP inference by default (#1001 )	2026-01-13 14:20:05 -08:00
Loay Rashid	23a4531491	[CI] Fixed Turbodiffusion I2V CI (#1002 )	2026-01-13 01:08:58 -08:00
William Linandgemini-code-assist[bot]	7d52ad0118	[ci] temporarily disable turbodiffusion ssim test (#1000 ) Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-01-08 15:43:50 -08:00
Will Lin	4d7bf35fa3	Revert "dit" This reverts commit `a6a9c9ca07`.	2026-01-07 03:22:48 -08:00
Will Lin	a6a9c9ca07	dit	2026-01-07 03:18:48 -08:00
Loay Rashid gemini-code-assist[bot]Will Lin	f4704847c2	[bugfix] Add configs for TurboDiffusion T2V/I2V models (#993 ) Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> Co-authored-by: Will Lin <wlsaidhi@gmail.com>	2026-01-06 16:36:45 -06:00
Shreejith SGandWill Lin	d9c996310b	[docs]: add LoRA extraction utilities documentation (#992 ) Co-authored-by: Will Lin <wlsaidhi@gmail.com>	2026-01-06 16:36:26 -06:00
Shao Duan	d6651afd2e	[examples] Added longcat-video python api examples (#994 )	2026-01-06 15:03:42 -06:00
William Lin	cf67618cad	[chore] release 0.1.7 (real) (#980 )	2026-01-05 15:47:05 -06:00
William Lin	2f0a2b3c57	[misc] add pin_cpu_memory false for RTX 4090 (#990 )	2026-01-05 15:45:35 -06:00
Loay Rashid	e7748d9952	[feat] add Turbodiffusion I2V pipeline (#984 )	2026-01-05 15:41:23 -06:00
William Lin	8eb3140b2f	[misc] pin fastvideo-kernel in .toml file (#989 )	2026-01-05 13:42:32 -06:00
Shao Duan	d6ddcea682	Add LongCat-Video I2V and Video Continuation (Base, Distillation and Refinement) Support to FastVideo (#953 )	2026-01-04 22:20:09 -06:00
William Lin	3559ba2377	[chore] update wechat QR code (#988 )	2026-01-04 21:59:38 -06:00
William Lin	61e63ea0d7	[chore] release fastvideo-kernel 0.2.2 (#986 )	2026-01-04 21:21:06 -06:00