Revert "[chore] release v0.1.7 (#955 )"

This reverts commit 9cd6a86b95.
2025-12-27 05:05:42 -06:00
313 changed files with 2272 additions and 37532 deletions
@@ -22,7 +22,7 @@ steps:
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
-                command: "timeout 20m .buildkite/scripts/pr_test.sh"
+                command: "timeout 15m .buildkite/scripts/pr_test.sh"
                label: "Encoder Tests"
                env:
                  - TEST_TYPE=encoder
@@ -35,7 +35,7 @@ steps:
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
-                command: "timeout 20m .buildkite/scripts/pr_test.sh"
+                command: "timeout 15m .buildkite/scripts/pr_test.sh"
                label: "VAE Tests"
                env:
                  - TEST_TYPE=vae
@@ -61,7 +61,7 @@ steps:
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
-                command: "timeout 90m .buildkite/scripts/pr_test.sh"
+                command: "timeout 45m .buildkite/scripts/pr_test.sh"
                label: "SSIM Tests"
                env:
                  - TEST_TYPE=ssim
@@ -76,7 +76,7 @@ steps:
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
-                command: "timeout 20m .buildkite/scripts/pr_test.sh"
+                command: "timeout 15m .buildkite/scripts/pr_test.sh"
                label: "LoRA Inference Tests"
                env:
                  - TEST_TYPE=inference_lora
@@ -157,9 +157,33 @@ steps:
                - "docker/Dockerfile.python3.12"
              config:
                command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Kernel Tests"
+                label: "Precision Tests STA"
                env:
-                  - TEST_TYPE=kernel_tests
+                  - TEST_TYPE=precision_sta
+                agents:
+                  queue: "default"
+            - path:
+                - "fastvideo-kernel/**"
+                - "pyproject.toml"
+                - "docker/Dockerfile.python3.12"
+              config:
+                command: "timeout 15m .buildkite/scripts/pr_test.sh"
+                label: "Precision Tests VSA"
+                env:
+                  - TEST_TYPE=precision_vsa
+                agents:
+                  queue: "default"
+            - path:
+                - "fastvideo-kernel/**"
+                - "pyproject.toml"
+                - "docker/Dockerfile.python3.12"
+              config:
+                command: "timeout 15m .buildkite/scripts/pr_test.sh"
+                label: "Precision Tests VMoBA"
+                env:
+                  - TEST_TYPE=precision_vmoba
+                agents:
+                  queue: "default"
            - path:
                - "fastvideo-kernel/**"
                - "fastvideo/attention/backends/vmoba.py"
@@ -183,14 +207,14 @@ steps:
                  - TEST_TYPE=unit_test
                agents:
                  queue: "default"
-            # - path:
-            #     - "scripts/lora_extraction/**"
-            #     - "pyproject.toml"
-            #     - "docker/Dockerfile.python3.12"
-            #   config:
-            #     command: "timeout 90m .buildkite/scripts/pr_test.sh"
-            #     label: "LoRA Extraction Tests"
-            #     env:
-            #       - TEST_TYPE=lora_extraction
-            #     agents:
-            #       queue: "default"
+            - path:
+                - "scripts/lora_extraction/**"
+                - "pyproject.toml"
+                - "docker/Dockerfile.python3.12"
+              config:
+                command: "timeout 90m .buildkite/scripts/pr_test.sh"
+                label: "LoRA Extraction Tests"
+                env:
+                  - TEST_TYPE=lora_extraction
+                agents:
+                  queue: "default"
@@ -93,9 +93,13 @@ case "$TEST_TYPE" in
        log "Running inference STA tests..."
        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_inference_tests_STA"
        ;;
-    "kernel_tests")
-        log "Running kernel tests..."
-        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_kernel_tests"
+    "precision_sta")
+        log "Running precision STA tests..."
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_precision_tests_STA"
+        ;;
+    "precision_vsa")
+        log "Running precision VSA tests..."
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_precision_tests_VSA"
        ;;
    "inference_lora")
        log "Running LoRA tests..."
@@ -114,6 +118,10 @@ case "$TEST_TYPE" in
        log "Running V-MoBA inference tests..."
        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_inference_tests_vmoba"
        ;;
+    "precision_vmoba")
+        log "Running V-MoBA precision tests..."
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_precision_tests_vmoba"
+        ;;
    "unit_test")
        log "Running unit tests..."
        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_unit_test"
@@ -156,23 +156,8 @@ 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 \
-            --exclude libtorch_cuda.so \
-            --exclude libtorch_cpu.so \
-            --exclude libtorch.so \
-            --exclude libc10.so \
-            --exclude libc10_cuda.so \
-            --exclude libtorch_python.so
+          auditwheel repair dist/*.whl --plat manylinux_2_35_x86_64 -w fixed_dist
          # Move fixed wheels back to dist for upload consistency
          rm dist/*.whl
          mv fixed_dist/*.whl dist/
@@ -1,6 +1,3 @@
 [submodule "fastvideo-kernel/include/tk"]
 	path = fastvideo-kernel/include/tk
 	url = https://github.com/HazyResearch/ThunderKittens.git
-[submodule "fastvideo-kernel/include/cutlass"]
-	path = fastvideo-kernel/include/cutlass
-	url = https://github.com/NVIDIA/cutlass.git
@@ -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,47 +1,41 @@
 <div align="center">
 <img src=assets/logos/logo.svg width="30%"/>
 </div>
+**FastVideo is a unified post-training and inference framework for accelerated video generation.**
+
+FastVideo features an end-to-end unified pipeline for accelerating diffusion models, starting from data preprocessing to model training, finetuning, distillation, and inference. FastVideo is designed to be modular and extensible, allowing users to easily add new optimizations and techniques. Whether it is training-free optimizations or post-training optimizations, FastVideo has you covered.

 <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/sv3MMKyv" target="_blank"> <b> WeChat </b> </a> |
+    | 🕹️ <a href="https://fastwan.fastvideo.org/"<b>Online Demo</b></a> | <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://huggingface.co/collections/FastVideo/fastwan-6886a305d9799c8cd1496408"  target="_blank"><b>FastWan</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>

-**FastVideo is a unified post-training and inference framework for accelerated video generation.**
+<div align="center">
+<img src=assets/fastwan.png width="90%"/>
+</div>

 ## 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/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/).

-</details>
-
 ## Key Features

 FastVideo has the following features:
- End-to-end post-training support for bidirectional and autoregressive models:
+- End-to-end post-training support:
+  - [Sparse distillation](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/) for Wan2.1 and Wan2.2 to achineve >50x denoising speedup
+  - Data preprocessing pipeline for video data
  - 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
-  - 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.
+  - Scalable training with FSDP2, sequence parallelism, and selective activation checkpointing, with near linear scaling to 64 GPUs
 - State-of-the-art performance optimizations for inference
-  - Sequence Parallelism for distributed inference
-  - Multiple state-of-the-art attention backends
-  - User-friendly CLI and Python API
-  - See this [page](https://hao-ai-lab.github.io/FastVideo/inference/optimizations/) for full list of supported optimizations.
+  - [Video Sparse Attention](https://arxiv.org/pdf/2505.13389)
+  - [Sliding Tile Attention](https://arxiv.org/pdf/2502.04507)
+  - [TeaCache](https://arxiv.org/pdf/2411.19108)
+  - [Sage Attention](https://arxiv.org/abs/2410.02367)
 - Diverse hardware and OS support
  - Support H100, A100, 4090
  - Support Linux, Windows, MacOS
-  - 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:
@@ -38,8 +38,7 @@ python -m benchmarks.fvd.cli \
    --num-frames 32 \
    --clip-strategy random \
    --batch-size 32 \
-    --seed 42 \
-    --extractor clip
+    --seed 42
 ```

 **Standard protocols:**
@@ -52,8 +51,6 @@ python -m benchmarks.fvd.cli \
    --protocol fvd2048_16f  # or fvd2048_128f, quick_test, etc.
 ```

-This would use i3d model by default as the feature extractor
-
 **Feature caching** (speed up repeated evaluations):

 ```bash
@@ -61,7 +58,7 @@ python -m benchmarks.fvd.cli \
    --real-path data/real/ \
    --gen-path outputs/gen/ \
    --protocol fvd2048_16f \
-    --cache-real-features fvd-cache/extractor_name  # Directory path (will save/load fvd-cache/extractor_name/extractor-name_real_features.pkl)
+    --cache-real-features cache/real  # Directory path (will save/load cache/real/real_features.pkl)
 ```

 Run `python -m benchmarks.fvd.cli --help` for all options.
@@ -86,7 +83,6 @@ batch_size=32,                # GPU batch size
 device='cuda',                # cuda|cpu
 cache_real_features=None,     # Cache path for speed
 seed=42,                      # Reproducibility
-extractor='i3d',              # i3d|clip|videomae
 ```

 ## Programmatic Usage
@@ -101,6 +97,7 @@ print(f"FVD: {results['fvd']:.2f}")

 ## Notes

+- I3D model auto-downloads from Hugging Face on first run
 - Requires minimum 10 frames per clip
 - Supports both video files (.mp4, .avi, etc.) and frame directories
 - `--cache-real-features` expects a **directory path** (e.g., `cache/real`), it will automatically create/load `real_features.pkl` inside that directory
@@ -13,8 +13,7 @@ from .fvd import (
    compute_statistics,
    FVDConfig,
 )
-from .feature_extractors import (BaseFeatureExtractor, I3DFeatureExtractor,
-                                 load_extractor)
+from .i3d_model import I3DFeatureExtractor
 from .video_utils import (
    load_video_auto,
    sample_clips_from_video,
@@ -28,9 +27,7 @@ __all__ = [
    'compute_frechet_distance',
    'compute_statistics',
    'FVDConfig',
-    'BaseFeatureExtractor',
    'I3DFeatureExtractor',
-    'load_extractor',
    'load_video_auto',
    'sample_clips_from_video',
    'load_video_clips_streaming',
@@ -1,64 +1,122 @@
 import argparse
+import json
 import sys
-import traceback
+from pathlib import Path
+
 from .fvd import compute_fvd_with_config, FVDConfig


 def main() -> int:
    parser = argparse.ArgumentParser(
-        description='Compute Fréchet Video Distance (FVD)')
+        description='Compute Fréchet Video Distance (FVD)',
+        formatter_class=argparse.RawDescriptionHelpFormatter,
+        epilog="""
+Examples:
+  # Standard FVD2048_16f protocol
+  python -m fastvideo.benchmarks.fvd.cli \\
+      --real-path data/real/ \\
+      --gen-path outputs/gen/ \\
+      --protocol fvd2048_16f
+
+  # Custom configuration
+  python -m fastvideo.benchmarks.fvd.cli \\
+      --real-path data/real/ \\
+      --gen-path outputs/gen/ \\
+      --num-videos 1024 \\
+      --num-frames 32 \\
+      --clip-strategy random \\
+      --frame-stride 2
+        """)

    # Required arguments
    parser.add_argument('--real-path',
                        type=str,
                        required=True,
-                        help='Path to real videos')
+                        help='Path to real videos directory')
    parser.add_argument('--gen-path',
                        type=str,
                        required=True,
-                        help='Path to generated videos')
+                        help='Path to generated videos directory')

-    # Extractor selection
-    parser.add_argument('--extractor',
-                        type=str,
-                        default='i3d',
-                        choices=['i3d', 'clip', 'videomae'],
-                        help='Feature extractor model to use (default: i3d)')
+    # Reproducibility
+    parser.add_argument(
+        '--seed',
+        type=int,
+        default=None,
+        help='Random seed for reproducibility (np.random, random, torch)')

-    # Standard args
-    parser.add_argument('--seed',
-                        type=int,
-                        default=None,
-                        help='Random seed for reproducibility')
+    # Protocol presets
    parser.add_argument('--protocol',
                        type=str,
                        default=None,
-                        choices=['fvd2048_16f', 'fvd2048_128f', 'quick_test'],
+                        choices=[
+                            'fvd2048_16f', 'fvd2048_128f',
+                            'fvd2048_128f_subsample8', 'quick_test'
+                        ],
                        help='Use standard protocol (overrides other settings)')
+
+    # Video selection
    parser.add_argument('--num-videos',
                        type=int,
                        default=2048,
-                        help='Number of videos to use')
+                        help='Number of videos to use (default: 2048)')
+
+    # Clip sampling
    parser.add_argument('--num-frames',
                        type=int,
                        default=16,
-                        help='Number of frames per clip')
-    parser.add_argument('--clip-strategy',
-                        type=str,
-                        default='beginning',
-                        help='Clip sampling strategy')
+                        help='Number of frames per clip (default: 16)')
+    parser.add_argument('--num-clips',
+                        type=int,
+                        default=1,
+                        help='Number of clips per video (default: 1)')
+    parser.add_argument(
+        '--clip-strategy',
+        type=str,
+        default='beginning',
+        choices=['beginning', 'random', 'uniform', 'middle', 'sliding', 'all'],
+        help='Clip sampling strategy (default: beginning)')
+    parser.add_argument(
+        '--frame-stride',
+        type=int,
+        default=1,
+        help='Frame stride for FPS subsampling (default: 1, no subsampling)')
+    parser.add_argument('--temporal-stride',
+                        type=int,
+                        default=1,
+                        help='Temporal stride for sliding window (default: 1)')
+
+    # Data processing
+    parser.add_argument('--no-frame-dirs',
+                        action='store_true',
+                        help='Disable frame directory support')
+
+    # Computation
    parser.add_argument('--batch-size',
                        type=int,
                        default=32,
-                        help='Batch size for feature extraction')
+                        help='Batch size for feature extraction (default: 32)')
    parser.add_argument('--device',
                        type=str,
                        default='cuda',
-                        help='Device to use (cuda or cpu)')
+                        choices=['cuda', 'cpu'],
+                        help='Device to use (default: cuda)')
+
+    # Caching
    parser.add_argument('--cache-real-features',
                        type=str,
                        default=None,
                        help='Path to cache real video features')
+    parser.add_argument('--i3d-model-path',
+                        type=str,
+                        default=None,
+                        help='Custom cache path for I3D model')
+
+    # Output
+    parser.add_argument('--output',
+                        type=str,
+                        default='fvd_results.json',
+                        help='Output JSON file (default: fvd_results.json)')
    parser.add_argument('--quiet',
                        action='store_true',
                        help='Suppress progress output')
@@ -70,36 +128,56 @@ def main() -> int:
        protocol_map = {
            'fvd2048_16f': FVDConfig.fvd2048_16f,
            'fvd2048_128f': FVDConfig.fvd2048_128f,
+            'fvd2048_128f_subsample8': FVDConfig.fvd2048_128f_subsample8,
            'quick_test': FVDConfig.quick_test,
        }
        config = protocol_map[args.protocol]()
-        # Apply overrides
+
+        # Override device and caching from args
        config.device = args.device
        config.cache_real_features = args.cache_real_features
-        config.extractor_model = args.extractor  # Apply extractor arg
+        config.i3d_model_path = args.i3d_model_path
+        config.batch_size = args.batch_size
+        config.seed = args.seed
    else:
-        config = FVDConfig(
-            num_videos=args.num_videos,
-            num_frames_per_clip=args.num_frames,
-            extractor_model=args.extractor,  # Apply extractor arg
-            clip_strategy=args.clip_strategy,
-            batch_size=args.batch_size,
-            device=args.device,
-            cache_real_features=args.cache_real_features,
-            seed=args.seed)
+        # Custom config from args
+        config = FVDConfig(num_videos=args.num_videos,
+                           num_frames_per_clip=args.num_frames,
+                           num_clips_per_video=args.num_clips,
+                           clip_strategy=args.clip_strategy,
+                           frame_stride=args.frame_stride,
+                           temporal_stride=args.temporal_stride,
+                           support_frame_dirs=not args.no_frame_dirs,
+                           batch_size=args.batch_size,
+                           device=args.device,
+                           cache_real_features=args.cache_real_features,
+                           i3d_model_path=args.i3d_model_path,
+                           seed=args.seed)

+    # Compute FVD
    try:
-        _ = compute_fvd_with_config(
-            args.real_path,  # noqa: F841
-            args.gen_path,
-            config,
-            verbose=not args.quiet)
+        results = compute_fvd_with_config(real_videos=args.real_path,
+                                          gen_videos=args.gen_path,
+                                          config=config,
+                                          verbose=not args.quiet)
+
+        # Save results
+        output_path = Path(args.output)
+        output_path.parent.mkdir(parents=True, exist_ok=True)
+
+        with open(output_path, 'w') as f:
+            json.dump(results, f, indent=2)
+
+        print(f"\nResults saved to {output_path}")
+        print(f"FVD: {results['fvd']:.2f}")
+        print(f"Protocol: {results['protocol']}")

        return 0

    except Exception as e:
        print(f"Error: {e}", file=sys.stderr)
-        traceback.print_exc(file=sys.stderr)
+        import traceback
+        traceback.print_exc()
        return 1


@@ -1,264 +0,0 @@
-"""
-Pluggable Feature Extractors for FVD Computation.
-Supports I3D (standard), CLIP, and VideoMAE via a common interface.
-"""
-
-import torch
-import torch.nn as nn
-import torch.nn.functional as F
-from abc import ABC, abstractmethod
-from huggingface_hub import hf_hub_download
-from tqdm import tqdm
-
-try:
-    from transformers import CLIPModel, CLIPProcessor, VideoMAEModel
-    TRANSFORMERS_AVAILABLE = True
-except ImportError:
-    TRANSFORMERS_AVAILABLE = False
-
-
-class BaseFeatureExtractor(ABC, nn.Module):
-    """Abstract base class for all video feature extractors."""
-
-    def __init__(self, device: str = 'cuda'):
-        super().__init__()
-        self.device = torch.device(
-            device if torch.cuda.is_available() else 'cpu')
-
-    @property
-    @abstractmethod
-    def feature_dim(self) -> int:
-        """Dimension of the output feature vector."""
-        pass
-
-    @abstractmethod
-    def preprocess(self, videos: torch.Tensor) -> torch.Tensor:
-        """
-        Args:
-            videos: [B, T, C, H, W] in [0, 255] range.
-        Returns:
-            Preprocessed tensor ready for the model.
-        """
-        pass
-
-    @abstractmethod
-    def extract_features_batch(self, videos: torch.Tensor) -> torch.Tensor:
-        """
-        Extract features for a single batch.
-        Args:
-            videos: [B, T, C, H, W] (raw input)
-        Returns:
-            Features: [B, feature_dim]
-        """
-        pass
-
-    @torch.no_grad()
-    def extract_features(self,
-                         videos: torch.Tensor,
-                         batch_size: int = 32,
-                         verbose: bool = True) -> torch.Tensor:
-        """
-        Extract features for a large tensor of videos by batching.
-        """
-        N = len(videos)
-        all_features = []
-
-        iterator = range(0, N, batch_size)
-        if verbose:
-            iterator = tqdm(
-                iterator,
-                desc=f"Extracting features ({self.__class__.__name__})")
-
-        for i in iterator:
-            batch = videos[i:i + batch_size].to(self.device)
-            features = self.extract_features_batch(batch)
-            all_features.append(features.cpu())
-
-        return torch.cat(all_features, dim=0)
-
-
-# 1. I3D Extractor (The Standard FVD Metric)
-class I3DFeatureExtractor(BaseFeatureExtractor):
-    REPO_ID = 'flateon/FVD-I3D-torchscript'
-    MODEL_FILENAME = 'i3d_torchscript.pt'
-
-    def __init__(self, device: str = 'cuda', cache_dir: str | None = None):
-        super().__init__(device)
-        self.cache_dir = cache_dir
-        self.model = self._load_model()
-        self.model.eval()
-        self.model.to(self.device)
-
-    @property
-    def feature_dim(self) -> int:
-        return 400
-
-    def _load_model(self) -> torch.nn.Module:
-        try:
-            model_path = hf_hub_download(repo_id=self.REPO_ID,
-                                         filename=self.MODEL_FILENAME,
-                                         cache_dir=self.cache_dir)
-            return torch.jit.load(model_path, map_location=self.device)
-        except Exception as e:
-            raise RuntimeError(f"Failed to load I3D model: {e}") from e
-
-    def preprocess(self, videos: torch.Tensor) -> torch.Tensor:
-        """Standard I3D preprocessing: Resize to 224, Norm to [-1, 1]."""
-        B, T, C, H, W = videos.shape
-
-        if T < 10:
-            raise ValueError(f"I3D requires at least 10 frames, got {T}")
-
-        # Normalize to [0, 1]
-        if videos.max() > 1.0:
-            videos = videos / 255.0
-
-        # Scale to [-1, 1]
-        videos = videos * 2.0 - 1.0
-
-        # Resize to 224x224
-        if H != 224 or W != 224:
-            videos = videos.reshape(B * T, C, H, W)
-            videos = F.interpolate(videos,
-                                   size=(224, 224),
-                                   mode='bilinear',
-                                   align_corners=False)
-            videos = videos.reshape(B, T, C, 224, 224)
-
-        # [B, T, C, H, W] -> [B, C, T, H, W]
-        return videos.permute(0, 2, 1, 3, 4).contiguous()
-
-    def extract_features_batch(self, videos: torch.Tensor) -> torch.Tensor:
-        batch = self.preprocess(videos)
-        # TorchScript I3D returns raw logits when return_features=True
-        return self.model(batch,
-                          rescale=False,
-                          resize=False,
-                          return_features=True)
-
-
-# 2. CLIP Extractor (Semantic/Content Quality)
-class CLIPFeatureExtractor(BaseFeatureExtractor):
-
-    def __init__(self,
-                 device: str = 'cuda',
-                 model_name: str = "openai/clip-vit-base-patch32"):
-        if not TRANSFORMERS_AVAILABLE:
-            raise ImportError(
-                "Please install transformers: pip install transformers")
-        super().__init__(device)
-        self.processor = CLIPProcessor.from_pretrained(model_name)
-        self.model = CLIPModel.from_pretrained(model_name).to(self.device)
-        self.model.eval()
-        self._feature_dim = self.model.config.projection_dim
-
-    @property
-    def feature_dim(self) -> int:
-        return self._feature_dim
-
-    def preprocess(self, videos: torch.Tensor) -> torch.Tensor:
-        # Ensure values are [0, 255]
-        if videos.max() <= 1.0:
-            videos = videos * 255.0
-
-        return videos.to(torch.uint8)
-
-    def extract_features_batch(self, videos: torch.Tensor) -> torch.Tensor:
-        # Input: [B, T, C, H, W]
-        B, T, C, H, W = videos.shape
-        videos = self.preprocess(videos)
-
-        # Flatten B*T to treat frames as images
-        images = videos.view(B * T, C, H, W)
-
-        # HF Processor
-        inputs = self.processor(images=images,
-                                return_tensors="pt",
-                                padding=True)
-        inputs = {k: v.to(self.device) for k, v in inputs.items()}
-
-        # Extract features [B*T, Dim]
-        outputs = self.model.get_image_features(**inputs)
-
-        # Reshape [B, T, Dim] and Average Pooling over time
-        outputs = outputs.view(B, T, -1)
-        return outputs.mean(dim=1)
-
-
-# 3. VideoMAE Extractor (Structure/Motion Quality)
-class VideoMAEFeatureExtractor(BaseFeatureExtractor):
-
-    def __init__(self,
-                 device: str = 'cuda',
-                 model_name: str = "MCG-NJU/videomae-base"):
-        if not TRANSFORMERS_AVAILABLE:
-            raise ImportError(
-                "Please install transformers: pip install transformers")
-        super().__init__(device)
-        self.model = VideoMAEModel.from_pretrained(model_name).to(self.device)
-        self.model.eval()
-
-        self.register_buffer(
-            'mean',
-            torch.tensor([0.485, 0.456, 0.406],
-                         device=self.device).view(1, 1, 3, 1, 1))
-        self.register_buffer(
-            'std',
-            torch.tensor([0.229, 0.224, 0.225],
-                         device=self.device).view(1, 1, 3, 1, 1))
-
-    @property
-    def feature_dim(self) -> int:
-        return self.model.config.hidden_size
-
-    def preprocess(self, videos: torch.Tensor) -> torch.Tensor:
-        """
-        Efficient GPU-based preprocessing.
-        Input: [B, T, C, H, W] in range [0, 255]
-        """
-        B, T, C, H, W = videos.shape
-
-        # 1. Resize to 224x224
-        if H != 224 or W != 224:
-            videos = videos.view(B * T, C, H, W)
-            videos = F.interpolate(videos,
-                                   size=(224, 224),
-                                   mode='bilinear',
-                                   align_corners=False)
-            videos = videos.view(B, T, C, 224, 224)
-
-        # 2. Normalize to [0, 1]
-        if videos.dtype != torch.float32:
-            videos = videos.float()
-
-        if videos.max() > 1.0:
-            videos = videos / 255.0
-
-        # 3. Apply ImageNet Mean/Std
-        return (videos - self.mean) / self.std
-
-    def extract_features_batch(self, videos: torch.Tensor) -> torch.Tensor:
-        # Input: [B, T, C, H, W]
-
-        # Fast GPU Preprocessing
-        pixel_values = self.preprocess(videos)
-
-        # Forward pass
-        outputs = self.model(pixel_values)
-
-        # Global Average Pooling of last hidden state [B, T_patches, 768] -> [B, 768]
-        return outputs.last_hidden_state.mean(dim=1)
-
-
-# Factory
-def load_extractor(name: str, device: str = 'cuda') -> BaseFeatureExtractor:
-    name = name.lower()
-    if name == 'i3d':
-        return I3DFeatureExtractor(device)
-    elif name == 'clip':
-        return CLIPFeatureExtractor(device)
-    elif name == 'videomae':
-        return VideoMAEFeatureExtractor(device)
-    else:
-        raise ValueError(
-            f"Unknown extractor: {name}. Options: i3d, clip, videomae")
@@ -5,7 +5,8 @@ from pathlib import Path
 from collections.abc import Iterator
 import pickle
 from dataclasses import dataclass, field
-from .feature_extractors import BaseFeatureExtractor, load_extractor
+
+from .i3d_model import I3DFeatureExtractor
 from .video_utils import ClipSamplingStrategy, load_video_clips_streaming


@@ -54,9 +55,6 @@ class FVDConfig:
    # Video selection
    num_videos: int = 2048

-    # Feature Extractor Selection
-    extractor_model: str = 'i3d'  # Options: 'i3d', 'clip', 'videomae'
-
    # Clip sampling
    num_frames_per_clip: int = 16
    num_clips_per_video: int = 1
@@ -87,7 +85,11 @@ class FVDConfig:

    @classmethod
    def fvd2048_16f(cls) -> 'FVDConfig':
-        """Standard FVD protocol: 2048 videos, 16 frames, beginning clip."""
+        """
+        Standard FVD protocol: 2048 videos, 16 frames, beginning clip.
+        
+        most common FVD configuration used in papers
+        """
        return cls(num_videos=2048,
                   num_frames_per_clip=16,
                   clip_strategy='beginning',
@@ -101,6 +103,18 @@ class FVDConfig:
                   clip_strategy='beginning',
                   use_streaming=True)

+    @classmethod
+    def fvd2048_128f_subsample8(cls) -> 'FVDConfig':
+        """
+        Long video with FPS subsampling: 2048 videos, 128 frames (every 8th).
+        Used for very long videos - samples every 8th frame
+        """
+        return cls(num_videos=2048,
+                   num_frames_per_clip=16,
+                   frame_stride=8,
+                   clip_strategy='beginning',
+                   use_streaming=True)
+
    @classmethod
    def quick_test(cls) -> 'FVDConfig':
        """Quick test config: 100 videos, 16 frames."""
@@ -110,13 +124,22 @@ class FVDConfig:

    def to_dict(self) -> dict:
        """Export config to dict for logging"""
-        d = self.__dict__.copy()
-        d['clip_strategy'] = str(self.clip_strategy)
-        return d
+        return {
+            'num_videos': self.num_videos,
+            'num_frames_per_clip': self.num_frames_per_clip,
+            'num_clips_per_video': self.num_clips_per_video,
+            'clip_strategy': str(self.clip_strategy),
+            'frame_stride': self.frame_stride,
+            'temporal_stride': self.temporal_stride,
+            'batch_size': self.batch_size,
+            'device': self.device,
+            'seed': self.seed,
+            'use_streaming': self.use_streaming,
+        }

    def __str__(self) -> str:
        """Human-readable protocol name"""
-        desc = f"FVD_{self.extractor_model.upper()}_{self.num_videos}_{self.num_frames_per_clip}f"
+        desc = f"FVD{self.num_videos}_{self.num_frames_per_clip}f"
        if self.frame_stride > 1:
            desc += f"_subsample{self.frame_stride}"
        if self.num_clips_per_video > 1:
@@ -127,42 +150,57 @@ class FVDConfig:


 def extract_features_streaming(video_generator: Iterator[torch.Tensor],
-                               extractor: BaseFeatureExtractor,
+                               extractor: I3DFeatureExtractor,
                               batch_size: int = 32,
                               max_clips: int | None = None,
                               verbose: bool = True) -> np.ndarray:
    """
    Extract features from a video clip generator using streaming.
+
+    Args:
+        video_generator: Iterator yielding clips [T, C, H, W]
+        extractor: I3D feature extractor
+        batch_size: Batch size for processing
+        max_clips: Maximum clips to process (for validation)
+        verbose: Show progress
+    
+    Returns:
+        features: [N, 400] numpy array
    """
    all_features = []
    batch = []
+    clip_count = 0

    if verbose:
        print(f"Extracting features with batch_size={batch_size}...")

-    with torch.no_grad():
-        for clip_count, clip in enumerate(video_generator):
-            batch.append(clip)
+    for clip_count, clip in enumerate(video_generator):
+        batch.append(clip)

-            # Process batch when full
-            if len(batch) == batch_size:
-                batch_tensor = torch.stack(batch).to(extractor.device)
-                features = extractor.extract_features_batch(batch_tensor)
-
-                all_features.append(features.detach().cpu().numpy())
-                batch = []
-
-                if verbose and clip_count % (batch_size * 10) == 0:
-                    print(f"Processed {clip_count} clips...")
-
-            if max_clips is not None and clip_count >= max_clips:
-                break
-
-        # Process remaining clips
-        if len(batch) > 0:
+        # Process batch when full
+        if len(batch) == batch_size:
            batch_tensor = torch.stack(batch).to(extractor.device)
-            features = extractor.extract_features_batch(batch_tensor)
-            all_features.append(features.detach().cpu().numpy())
+            features = extractor.extract_features(batch_tensor,
+                                                  batch_size=batch_size,
+                                                  verbose=False)
+            all_features.append(features.cpu().numpy())
+
+            batch = []  # Clear batch
+
+            if verbose and clip_count % (batch_size * 10) == 0:
+                print(f"Processed {clip_count} clips...")
+
+        # Stop if we've reached max_clips
+        if max_clips is not None and clip_count >= max_clips:
+            break
+
+    # Process remaining clips
+    if len(batch) > 0:
+        batch_tensor = torch.stack(batch).to(extractor.device)
+        features = extractor.extract_features(batch_tensor,
+                                              batch_size=len(batch),
+                                              verbose=False)
+        all_features.append(features.cpu().numpy())

    if len(all_features) == 0:
        raise RuntimeError("No features extracted - check video loading")
@@ -176,17 +214,14 @@ def extract_features_streaming(video_generator: Iterator[torch.Tensor],


 def load_or_compute_features(videos: str | Path | torch.Tensor,
-                             extractor: BaseFeatureExtractor,
+                             extractor: I3DFeatureExtractor,
                             config: FVDConfig,
                             cache_path: str | None = None,
                             cache_name: str = "real_features") -> np.ndarray:
    """Load features from cache or compute (with streaming support)"""

    if cache_path is not None:
-        script_dir = Path(__file__).parent
-        cache_dir = script_dir / cache_path
-        cache_file = cache_dir / f"{config.extractor_model}_{cache_name}.pkl"
-
+        cache_file = Path(cache_path) / f"{cache_name}.pkl"
        if cache_file.exists():
            print(f"Loading cached features from {cache_file}")
            with open(cache_file, 'rb') as f:
@@ -194,25 +229,19 @@ def load_or_compute_features(videos: str | Path | torch.Tensor,

            # Validate and limit based on config
            max_features = config.num_videos * config.num_clips_per_video
-
            if len(features) < max_features:
                print(
                    f"WARNING: Cache has {len(features)} features but need {max_features}"
                )
-                print("Cached features insufficient - will recompute...")
+                print("Recomputing features...")
            elif len(features) > max_features:
-                print(
-                    f"Using {max_features} features from cache (truncated from {len(features)})"
-                )
                features = features[:max_features]
                return features
            else:
-                print(f"Using all {len(features)} cached features")
                return features

-    print("Computing features from scratch...")
-
-    if isinstance(videos, (str | Path)):
+    # Compute features
+    if isinstance(videos, str | Path):
        target_size = (224, 224) if config.resize_before_extraction else None

        video_generator = load_video_clips_streaming(
@@ -233,7 +262,9 @@ def load_or_compute_features(videos: str | Path | torch.Tensor,
                                              batch_size=config.batch_size,
                                              max_clips=max_clips,
                                              verbose=True)
+
    else:
+        # Already a tensor
        print(f"Extracting features from {len(videos)} video tensors...")
        features = extractor.extract_features(videos,
                                              batch_size=config.batch_size,
@@ -252,10 +283,9 @@ def load_or_compute_features(videos: str | Path | torch.Tensor,

    # Cache features if requested
    if cache_path is not None:
-        script_dir = Path(__file__).parent
-        cache_dir = script_dir / cache_path
+        cache_dir = Path(cache_path)
        cache_dir.mkdir(parents=True, exist_ok=True)
-        cache_file = cache_dir / f"{config.extractor_model}_{cache_name}.pkl"
+        cache_file = cache_dir / f"{cache_name}.pkl"
        print(f"Caching features to {cache_file}")
        with open(cache_file, 'wb') as f:
            pickle.dump(features, f)
@@ -263,34 +293,79 @@ def load_or_compute_features(videos: str | Path | torch.Tensor,
    return features


+def compute_fvd(real_videos: str | Path | torch.Tensor,
+                gen_videos: str | Path | torch.Tensor,
+                num_frames: int = 16,
+                batch_size: int = 32,
+                device: str = 'cuda',
+                num_videos: int | None = 2048,
+                cache_real_features: str | None = None,
+                i3d_model_path: str | None = None,
+                seed: int | None = None,
+                verbose: bool = True) -> float:
+    """
+    Compute Fréchet Video Distance (FVD)
+    
+    For advanced control, use compute_fvd_with_config() instead.
+    
+    Args:
+        real_videos: Path to real videos or tensor [N, T, C, H, W]
+        gen_videos: Path to generated videos or tensor [N, T, C, H, W]
+        num_frames: Frames per video (default: 16)
+        batch_size: Batch size (default: 32)
+        device: 'cuda' or 'cpu' (default: 'cuda')
+        num_videos: Max videos (default: 2048)
+        cache_real_features: Cache path for real features
+        i3d_model_path: Custom I3D model cache path
+        seed: Random seed for reproducibility
+        verbose: Print progress
+    
+    Returns:
+        FVD score (float). Lower is better.
+    """
+    num_videos = num_videos if num_videos is not None else 2048
+
+    config = FVDConfig(
+        num_videos=num_videos,
+        num_frames_per_clip=num_frames,
+        batch_size=batch_size,
+        device=device,
+        cache_real_features=cache_real_features,
+        i3d_model_path=i3d_model_path,
+        seed=seed,
+    )
+
+    result = compute_fvd_with_config(real_videos, gen_videos, config, verbose)
+    return result['fvd']
+
+
 def compute_fvd_with_config(real_videos: str | Path | torch.Tensor,
                            gen_videos: str | Path | torch.Tensor,
                            config: FVDConfig,
                            verbose: bool = True) -> dict:
    """
-        Compute FVD using a standardized configuration.
-
-        This is the recommended way to compute FVD for reproducibility.
-
-        Args:
-            real_videos: Path or tensors
-            gen_videos: Path or tensors
-            config: FVDConfig specifying protocol
-            verbose: Print progress
-
-        Returns:
-            results: Dictionary with:
-                - 'fvd': FVD score (float)
-                - 'protocol': Protocol name (str)
-                - 'model': Feature extractor model name (str)
-                - 'config': Configuration dict
-
-        Example:
-            >>> config = FVDConfig.fvd2048_16f()
-            >>> results = compute_fvd_with_config('data/real/', 'outputs/gen/', config)
-            >>> print(f"FVD: {results['fvd']:.2f}")
+    Compute FVD using a standardized configuration.
+    
+    This is the recommended way to compute FVD for reproducibility.
+    
+    Args:
+        real_videos: Path or tensors
+        gen_videos: Path or tensors
+        config: FVDConfig specifying protocol
+        verbose: Print progress
+    
+    Returns:
+        results: Dictionary with:
+            - 'fvd': FVD score (float)
+            - 'protocol': Protocol name (str)
+            - 'config': Configuration dict
+    
+    Example:
+        >>> config = FVDConfig.fvd2048_16f()
+        >>> results = compute_fvd_with_config('data/real/', 'outputs/gen/', config)
+        >>> print(f"FVD: {results['fvd']:.2f}")
+        >>> print(f"Protocol: {results['protocol']}")  # "FVD2048_16f"
    """
-
    # Seed for reproducibility
    if config.seed is not None:
        import random as _rnd
@@ -303,20 +378,18 @@ def compute_fvd_with_config(real_videos: str | Path | torch.Tensor,
    if verbose:
        print("=" * 70)
        print(f"Computing FVD with protocol: {config}")
-        print(f"Model: {config.extractor_model.upper()}")
        print("=" * 70)
        print("\nConfiguration:")
        for key, value in config.to_dict().items():
            print(f"  {key}: {value}")
        print()

-    # Initialize Extractor using Factory
+    # Initialize I3D
    if verbose:
-        print(
-            f"\nInitializing {config.extractor_model.upper()} model on {config.device}..."
-        )
+        print(f"\nInitializing I3D model on {config.device}...")

-    extractor = load_extractor(config.extractor_model, device=config.device)
+    extractor = I3DFeatureExtractor(device=config.device,
+                                    cache_dir=config.i3d_model_path)

    # Extract features
    if verbose:
@@ -361,45 +434,14 @@ def compute_fvd_with_config(real_videos: str | Path | torch.Tensor,

    if verbose:
        print(f"\n{'='*70}")
-        print(f"FVD Score ({config.extractor_model.upper()}): {fvd:.4f}")
+        print(f"FVD Score: {fvd:.4f}")
        print(f"Protocol: {config}")
        print(f"{'='*70}\n")

    results = {
        'fvd': fvd,
        'protocol': str(config),
-        'model': config.extractor_model,
        'config': config.to_dict(),
    }

    return results
-
-
-def compute_fvd(real_videos: str | Path | torch.Tensor,
-                gen_videos: str | Path | torch.Tensor,
-                num_frames: int = 16,
-                batch_size: int = 32,
-                device: str = 'cuda',
-                num_videos: int | None = 2048,
-                cache_real_features: str | None = None,
-                i3d_model_path: str | None = None,
-                seed: int | None = None,
-                verbose: bool = True) -> float:
-    """
-    Backward compatibility wrapper for computing FVD (defaults to I3D).
-    """
-    num_videos = num_videos if num_videos is not None else 2048
-
-    config = FVDConfig(
-        num_videos=num_videos,
-        num_frames_per_clip=num_frames,
-        extractor_model='i3d',  # Default to I3D
-        batch_size=batch_size,
-        device=device,
-        cache_real_features=cache_real_features,
-        i3d_model_path=i3d_model_path,
-        seed=seed,
-    )
-
-    result = compute_fvd_with_config(real_videos, gen_videos, config, verbose)
-    return result['fvd']
@@ -1,53 +1,33 @@
 import sys
 from pathlib import Path
+from benchmarks.fvd.fvd import FVDConfig, compute_fvd_with_config

 root_dir = Path(__file__).parent.parent.parent
 sys.path.insert(0, str(root_dir))

-from benchmarks.fvd.fvd import FVDConfig, compute_fvd_with_config  # noqa: E402
-

 def main() -> None:
+    # Get script directory
    script_dir = Path(__file__).parent.resolve()

-    # Define directories
+    clip_strategy = 'beginning'  # Options: 'uniform', 'random', 'beginning', 'end', 'all'
+    cfg = FVDConfig(
+        num_videos=650,
+        num_frames_per_clip=16,
+        num_clips_per_video=1,
+        clip_strategy=clip_strategy,
+        frame_stride=1,
+        batch_size=32,
+        device='cuda',
+        seed=42,
+        cache_real_features=str(script_dir / f'fvd-cache/{clip_strategy}'),
+    )
+
    real_dir = "benchmarks/data/real_videos"
    gen_dir = "benchmarks/data/generated_videos"

-    # Compare all 3 models
-    models_to_test = ['i3d', 'clip', 'videomae']
-
-    print(f"\n{'='*60}")
-    print("STARTING COMPARISON BENCHMARK")
-    print(f"{'='*60}")
-
-    for model_name in models_to_test:
-        print(f"\n>>> Running evaluation with {model_name.upper()}...")
-
-        try:
-            cfg = FVDConfig(
-                num_videos=650,
-                num_frames_per_clip=16,
-                extractor_model=model_name,
-                clip_strategy='beginning',
-                device='cuda',
-                seed=42,
-                # Use separate cache folders for each model to avoid conflicts
-                cache_real_features=str(script_dir / f'fvd-cache/{model_name}'),
-            )
-
-            results = compute_fvd_with_config(real_dir,
-                                              gen_dir,
-                                              cfg,
-                                              verbose=False)
-            print(f"FVD: {results['fvd']}\nModel: {results['model']}")
-
-        except Exception as e:
-            print(f"{model_name.upper()} Failed: {e}")
-
-    print(f"\n{'='*60}")
-    print("BENCHMARK COMPLETE")
-    print(f"{'='*60}")
+    results = compute_fvd_with_config(real_dir, gen_dir, cfg, verbose=True)
+    print(f"FVD = {results['fvd']:.2f}")


 if __name__ == '__main__':
@@ -1,7 +1,7 @@
 #!/bin/bash

 # 1. Install missing dependency
-pip install -q opencv-python-headless transformers huggingface_hub
+pip install -q opencv-python-headless

 # 2. Run FVD script
 python benchmarks/fvd/run_fvd.py
@@ -60,7 +60,7 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    cd fastvideo-kernel && \
    git submodule update --init --recursive && \
-    ./build.sh
+    ./build.sh --release


 EXPOSE 22
@@ -60,7 +60,7 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    cd fastvideo-kernel && \
    git submodule update --init --recursive && \
-    ./build.sh
+    ./build.sh --release


 EXPOSE 22
@@ -60,6 +60,6 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    cd fastvideo-kernel && \
    git submodule update --init --recursive && \
-    ./build.sh
+    ./build.sh --release

 EXPOSE 22
@@ -60,7 +60,7 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    cd fastvideo-kernel && \
    git submodule update --init --recursive && \
-    ./build.sh
+    ./build.sh --release


 EXPOSE 22
@@ -36,7 +36,7 @@ RUN echo "# Placeholder" > README.md
 RUN source $HOME/.local/bin/env && \
    uv venv --python 3.10 --seed /opt/venv && \
    source /opt/venv/bin/activate && \
-    uv pip install --no-cache-dir --upgrade pip
+    uv pip install --no-cache-dir --upgrade pip 

 COPY . .

@@ -53,6 +53,6 @@ RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
    cd fastvideo-kernel && \
    git submodule update --init --recursive && \
-    ./build.sh --rocm
+    ./build.sh --release

 EXPOSE 22
@@ -50,25 +50,21 @@ class MyNewAttnBackend(AttentionBackend):
 FastVideo uses a `ForwardContext` to pass global metadata (like current timestep, batch info, or custom attention configurations) to attention backends without changing the `forward` signature of every layer. **This is optional and only required if your backend needs dynamic per-step information.**

 To use this:
-
 1. **Set Context**: In your pipeline or generation loop, use the `set_forward_context` context manager.
 2. **Access Context**: Inside your attention backend, use `get_forward_context()`.

-See [`docs/attention/sta/index.md`](../sta/index.md) (Sliding Tile Attention) for an example of how complex configuration (window sizes) is passed this way.
+See `docs/attention/sta/index.md` (Sliding Tile Attention) for an example of how complex configuration (window sizes) is passed this way.

 ## 3. Adding Compiled Kernels (C++/CUDA)

 If your backend requires custom CUDA kernels, you need to add them to the `fastvideo-kernel` package.

 ### A. Add Source Files
-
 Place your kernel implementation files in `fastvideo-kernel/csrc/attention/`.
-
 * `mynew_attn.cu` (CUDA implementation)
 * `mynew_attn.h` (Optional headers)

 ### B. Register in Extension
-
 Update `fastvideo-kernel/csrc/common_extension.cpp` to expose your function to Python.

 ```cpp
@@ -88,7 +84,6 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
 ```

 ### C. Update CMakeLists.txt
-
 Update `fastvideo-kernel/CMakeLists.txt` to compile your new files.

 **Case 1: General CUDA Kernel (Runs on all GPUs)**
@@ -116,7 +111,6 @@ endif()
 ```

 ### D. Expose in Python Ops
-
 Update `fastvideo-kernel/python/fastvideo_kernel/ops.py` to make the function importable and handle fallbacks gracefully.

 ```python
@@ -139,7 +133,6 @@ def my_compiled_attn_func(q, k, v):
 ```

 ### E. Expose in Package Init
-
 Update `fastvideo-kernel/python/fastvideo_kernel/__init__.py` to export the function.

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

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

 # Build and install
@@ -13,13 +13,11 @@ 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, 
-    block_sizes,
-    block_sizes,
+    variable_block_sizes=block_sizes, 
    topk=32
 )
 ```
@@ -4,26 +4,25 @@ This document outlines FastVideo's architecture for developers interested in fra

 ## Table of Contents - Directory Structure and Files

- [`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/pipelines/`](#design-pipeline-system) - Core diffusion pipeline components
+- [`fastvideo/models/`](#design-model-components) - Model implementations
+  - [`dits/`](#design-transformer-models) - Transformer-based diffusion models
+  - [`vaes/`](#design-vae-variational-auto-encoder) - Variational autoencoders
+  - [`encoders/`](#design-text-and-image-encoders) - Text and image encoders
+  - [`schedulers/`](#design-schedulers) - Diffusion schedulers
+- [`fastvideo/attention/`](#design-optimized-attention) - Optimized attention implementations
+- [`fastvideo/distributed/`](#design-distributed-processing) - Distributed computing utilities
+- [`fastvideo/layers/`](#design-tensor-parallelism) - Custom neural network layers
+- [`fastvideo/platforms/`](#design-platforms) - Hardware platform abstractions
+- [`fastvideo/worker/`](#design-executor-and-worker-abstractions) - Multi-GPU process management
+- [`fastvideo/fastvideo_args.py`](#design-fastvideo-args) - Argument handling
+- [`fastvideo/forward_context.py`](#design-forwardcontext) - Forward pass context management
 - `fastvideo/utils.py` - Utility functions
- [`fastvideo/logger.py`](#logger) - Logging infrastructure
+- [`fastvideo/logger.py`](#design-logger) - Logging infrastructure

 ## Core Architecture

 FastVideo separates model components from execution logic with these principles:
-
 - **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)
@@ -35,14 +34,12 @@ FastVideo separates model components from execution logic with these principles:
 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`
@@ -93,9 +90,7 @@ class MyCustomPipeline(ComposedPipelineBase):
 ```

 ### 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
@@ -138,7 +133,6 @@ Transformer networks perform the actual denoising during diffusion:
  - `HunyuanVideoTransformer3DModel`

 Features include:
-
 - Text/image conditioning
 - Standardized interface for model-specific optimizations

@@ -167,7 +161,6 @@ VAEs handle conversion between pixel space and latent space:
 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
@@ -186,7 +179,6 @@ Encoders process conditioning inputs into embeddings:
  - `CLIPVisionModel`

 FastVideo implements optimizations such as:
-
 - Vocab parallelism for distributed processing
 - Caching for common prompts
 - Precision-tuned computation
@@ -201,7 +193,6 @@ Schedulers manage the diffusion sampling process:
  - `FlowMatchEulerDiscreteScheduler`

 These components control:
-
 - Diffusion timestep sequences
 - Noise prediction to latent update conversions
 - Quality/speed trade-offs
@@ -228,9 +219,7 @@ This diagram shows how models are discovered, validated, and loaded across entry
 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
@@ -251,9 +240,7 @@ self.attn = LocalAttention(
 ![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

@@ -309,7 +296,6 @@ self.attn = DistributedAttention(
 ```

 ### Communication Primitives
-
 Efficient distributed operations via AllGather, AllReduce, and synchronization mechanisms.

 Efficient communication primitives minimize distributed overhead:
@@ -328,7 +314,6 @@ Defined in `fastvideo/forward_context.py`, `ForwardContext` manages execution-sp
 - **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

@@ -354,14 +339,12 @@ FastVideo implements a flexible execution model for distributed processing:
 - **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
@@ -376,13 +359,11 @@ The `fastvideo/platforms/` directory provides hardware platform abstractions tha
 ### 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
@@ -402,7 +383,6 @@ else:
 The platform system is designed to be extensible for future hardware targets.

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

 *TODO*: (help wanted) Add an environment variable that disables process-aware logging.
@@ -417,7 +397,6 @@ If you're a new contributor, here are some common areas to explore:
 4. **Hardware support**: Extend the `platforms` module for new hardware targets

 When adding code, follow these practices:
-
 - Use type hints for better code readability
 - Add appropriate docstrings
 - Maintain the separation between model components and execution logic
@@ -13,7 +13,7 @@ We provide two distilled models:
 Both models are trained on **61×448×832** resolution but support generating videos with **any resolution** (1.3B  model mainly support 480P, 14B model support 480P and 720P, quality may degrade for different resolutions).

 ## ⚙️ Inference
-First install [VSA](../attention/vsa/index.md). Set `MODEL_BASE` to your own model path and run:
+First install [VSA](https://hao-ai-lab.github.io/FastVideo/video_sparse_attention/installation). Set `MODEL_BASE` to your own model path and run:

 ```bash
 bash scripts/inference/v1_inference_wan_dmd.sh
@@ -11,13 +11,15 @@ 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
@@ -26,12 +28,6 @@ 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+
@@ -42,4 +38,4 @@ pip install flash-attn --no-build-isolation

 - [Quick Start Guide](quick_start.md) - Get started with your first video generation
 - [Configuration](../inference/configuration.md) - Learn about configuration options
- [Examples](../inference/examples/examples_inference_index.md) - Explore example scripts and notebooks
+- [Examples](../inference/examples/) - Explore example scripts and notebooks
@@ -84,12 +84,12 @@ pip install flash-attn --no-build-isolation

 ## Set up using Docker
 We also have prebuilt docker images with FastVideo dependencies pre-installed:
-[Docker Images](../../contributing/developer_env/docker.md)
+[Docker Images](#docker)

 ## Development Environment Setup

 If you're planning to contribute to FastVideo please see the following page:
-[Contributor Guide](../../contributing/overview.md)
+[Contributor Guide](#developer-overview)

 ## Hardware Requirements

@@ -78,7 +78,7 @@ uv pip install -e .
 ## Development Environment Setup

 If you're planning to contribute to FastVideo please see the following page:
-[Contributor Guide](../../contributing/overview.md)
+[Contributor Guide](#developer-overview)

 ## Hardware Requirements

@@ -15,12 +15,6 @@ 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
@@ -45,7 +45,6 @@ FastVideo uses the Hugging Face Diffusers format for model organization:
 ### Implementing Modules

 Place new modules in the appropriate directories:
-
 - Encoders: `fastvideo/models/encoders/`
 - VAEs: `fastvideo/models/vaes/`
 - Transformer models: `fastvideo/models/dits/`
@@ -54,15 +53,12 @@ Place new modules in the appropriate directories:
 ### Adapting Model Layers

 #### Layer Replacements
-
 Replace standard PyTorch layers with FastVideo optimized versions:
-
 - nn.LayerNorm → fastvideo.layers.layernorm.RMSNorm
 - Embedding layers → fastvideo.layers.vocab_parallel_embedding modules
 - Activation functions → versions from fastvideo.layers.activation

 #### Distributed Linear Layers
-
 Use appropriate parallel layers for distribution:

 ```python
@@ -95,7 +91,6 @@ self.out_proj = RowParallelLinear(
 ```

 ### Attention Layers
-
 Replace standard attention with FastVideo's optimized attention:

 ```python
@@ -309,7 +304,6 @@ EntryClass = [MyCustomPipeline, MyOtherPipeline]
 ```

 The registry will automatically:
-
 1. Scan all packages under `fastvideo/pipelines/`
 2. Look for `EntryClass` variables
 3. Register pipelines using their class names as identifiers
@@ -1,7 +1,7 @@
 # FastVideo CLI Inference

 The FastVideo CLI provides a quick way to access the FastVideo inference pipeline for video generation. For more advanced usage,
-see the Python interface [here](examples/basic.md).
+see the Python interface [here](https://hao-ai-lab.github.io/FastVideo/inference/examples/basic.html).

 ## Basic Usage

@@ -74,4 +74,4 @@ if __name__ == '__main__':

 ## Performance Optimization

-For configuring optimizations, please see our [optimizations guide](optimizations.md)
+For configuring optimizations, please see our [optimizations guide](#inference-optimizations)
@@ -3,7 +3,6 @@
 This page contains step-by-step instructions to get you quickly started with video generation using FastVideo.

 ## Requirements
-
 - **OS**: Linux (Tested on Ubuntu 22.04+)
 - **Python**: 3.10-3.12
 - **CUDA**: 12.8
@@ -22,10 +21,9 @@ conda activate fastvideo
 pip install fastvideo
 ```

-For advanced installation options, see the [Installation Guide](../getting_started/installation.md).
+For advanced installation options, see the [Installation Guide](installation.md).

 ## Generating Your First Video
-
 Here's a minimal example to generate a video using the default settings. Create a file called `example.py` with the following code:

 ```python
@@ -62,10 +60,9 @@ python example.py
 The generated video will be saved in the current directory under `my_videos/`  

 More inference example scripts can be found in `scripts/inference/`
-
 ## Available Models

-Please see the [support matrix](support_matrix.md) for the list of supported models and their available optimizations.
+Please see the [support matrix](#support-matrix) for the list of supported models and their available optimizations.

 ## Image-to-Video Generation

@@ -99,28 +96,20 @@ if __name__ == '__main__':
 Common issues and their solutions:

 ### Out of Memory Errors
-
 If you encounter CUDA out of memory errors:
-
 - Reduce `num_frames` or video resolution
 - 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
-
 To speed up generation:
-
 - Reduce `num_inference_steps` (20-30 is usually sufficient)
 - Use half precision (`fp16`) for the VAE
 - Use multiple GPUs if available

 ### Unexpected Results
-
 If the generated video doesn't match your prompt:
-
 - Try increasing `guidance_scale` (7.0-9.0 works well)
 - Make your prompt more detailed and specific
 - Experiment with different random seeds
@@ -128,8 +117,8 @@ If the generated video doesn't match your prompt:

 ## Next Steps

- Learn about [Advanced Inference Configurations](configuration.md)
- Learn about using [Optimizations](optimizations.md)
- See [Examples](examples/examples_inference_index.md) for more usage scenarios
+- Learn about [Advanced Inference Configurations](#inference-configuration)
+- Learn about using [Optimizations](#inference-optimizations)
+- See [Examples](../examples/examples_inference_index.md) for more usage scenarios
 - Join our [Community Discord](https://discord.gg/JA7cksDz86).
 - Join our [Community Slack](https://join.slack.com/t/fastvideo/shared_invite/zt-38u6p1jqe-yDI1QJOCEnbtkLoaI5bjZQ).
@@ -7,13 +7,13 @@ This page describes the various options for speeding up generation times in Fast

 - Optimized Attention Backends

-  - [Flash Attention](#flash-attention)
-  - [Sliding Tile Attention](#sliding-tile-attention)
-  - [Sage Attention](#sage-attention)
-  - [Sage Attention 3](#sage-attention-3)
+  - [Flash Attention](#optimizations-flash)
+  - [Sliding Tile Attention](#optimizations-sta)
+  - [Sage Attention](#optimizations-sage)
+  - [Sage Attention 3](#optimizations-sage3)

 - Caching Techniques
-  - [TeaCache](#teacache)
+  - [TeaCache](#optimizations-teacache)

 ## Attention Backends

@@ -74,7 +74,7 @@ python setup.py install
 pip install st_attn==0.0.4
 ```

-Please see [this page](../attention/sta/index.md) for more installation instructions.
+Please see [this page](#sta-installation) for more installation instructions.

 ### Video Sparse Attention

@@ -85,7 +85,7 @@ git submodule update --init --recursive
 python setup_vsa.py install
 ```

-Please see [this page](../attention/vsa/index.md) for more installation instructions.
+Please see [this page](#vsa-installation) for more installation instructions.

 ### Sage Attention

@@ -40,26 +40,20 @@ The `HuggingFace Model ID` can be directly pass to `from_pretrained()` methods a
  }
 </style>

-| Model Name | HuggingFace Model ID | Resolutions | TeaCache | Sliding Tile Attn | Sage Attn | VSA | BSA |
-|------------|---------------------|-------------|----------|-------------------|-----------|-----|-----|
-| FastWan2.1 T2V 1.3B | `FastVideo/FastWan2.1-T2V-1.3B-Diffusers` | 480P | ⭕ | ⭕ | ⭕ | ✅ | ⭕ |
-| FastWan2.2 TI2V 5B Full Attn* | `FastVideo/FastWan2.2-TI2V-5B-FullAttn-Diffusers` | 720P | ⭕ | ⭕ | ⭕ | ✅ | ⭕ |
-| Wan2.2 TI2V 5B | `Wan-AI/Wan2.2-TI2V-5B-Diffusers` | 720P | ⭕ | ⭕ | ✅ | ⭕ | ⭕ |
-| Wan2.2 T2V A14B | `Wan-AI/Wan2.2-T2V-A14B-Diffusers` | 480P<br>720P | ❌ | ❌ | ✅ | ⭕ | ⭕ |
-| Wan2.2 I2V A14B | `Wan-AI/Wan2.2-I2V-A14B-Diffusers` | 480P<br>720P | ❌ | ❌ | ✅ | ⭕ | ⭕ |
-| HunyuanVideo | `hunyuanvideo-community/HunyuanVideo` | 720px1280p<br>544px960p | ❌ | ✅ | ✅ | ⭕ | ⭕ |
-| FastHunyuan | `FastVideo/FastHunyuan-diffusers` | 720px1280p<br>544px960p | ❌ | ✅ | ✅ | ⭕ | ⭕ |
-| Wan2.1 T2V 1.3B | `Wan-AI/Wan2.1-T2V-1.3B-Diffusers` | 480P | ✅ | ✅* | ✅ | ⭕ | ⭕ |
-| Wan2.1 T2V 14B | `Wan-AI/Wan2.1-T2V-14B-Diffusers` | 480P, 720P | ✅ | ✅* | ✅ | ⭕ | ⭕ |
-| Wan2.1 I2V 480P | `Wan-AI/Wan2.1-I2V-14B-480P-Diffusers` | 480P | ✅ | ✅* | ✅ | ⭕ | ⭕ |
-| Wan2.1 I2V 720P | `Wan-AI/Wan2.1-I2V-14B-720P-Diffusers` | 720P | ✅ | ✅ | ✅ | ⭕ | ⭕ |
-| StepVideo T2V | `FastVideo/stepvideo-t2v-diffusers` | 768px768px204f<br>544px992px204f<br>544px992px136f | ❌ | ❌ | ✅ | ⭕ | ⭕ |
-| TurboWan2.1 T2V 1.3B | `loayrashid/TurboWan2.1-T2V-1.3B-Diffusers` | 480P | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
-| TurboWan2.1 T2V 14B | `loayrashid/TurboWan2.1-T2V-14B-Diffusers` | 480P, 720P | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
-| LongCat T2V 13.6B | See note** | 480P<br>720P | ❌ | ❌ | ❌ | ⭕ | ✅ |
-| Matrix Game 2.0 Base | `FastVideo/Matrix-Game-2.0-Base-Diffusers` | 352x640 | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
-| Matrix Game 2.0 GTA | `FastVideo/Matrix-Game-2.0-GTA-Diffusers` | 352x640 | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
-| Matrix Game 2.0 TempleRun | `FastVideo/Matrix-Game-2.0-TempleRun-Diffusers` | 352x640 | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
+| Model Name | HuggingFace Model ID | Resolutions | TeaCache | Sliding Tile Attn | Sage Attn | VSA |
+|------------|---------------------|-------------|----------|-------------------|-----------|-----|
+| FastWan2.1 T2V 1.3B | `FastVideo/FastWan2.1-T2V-1.3B-Diffusers` | 480P | ⭕ | ⭕ | ⭕ | ✅ |
+| FastWan2.2 TI2V 5B Full Attn* | `FastVideo/FastWan2.2-TI2V-5B-FullAttn-Diffusers` | 720P | ⭕ | ⭕ | ⭕ | ✅ |
+| Wan2.2 TI2V 5B | `Wan-AI/Wan2.2-TI2V-5B-Diffusers` | 720P | ⭕ | ⭕ | ✅ | ⭕ |
+| Wan2.2 T2V A14B | `Wan-AI/Wan2.2-T2V-A14B-Diffusers` | 480P<br>720P | ❌ | ❌ | ✅ | ⭕ |
+| Wan2.2 I2V A14B | `Wan-AI/Wan2.2-I2V-A14B-Diffusers` | 480P<br>720P | ❌ | ❌ | ✅ | ⭕ |
+| HunyuanVideo | `hunyuanvideo-community/HunyuanVideo` | 720px1280p<br>544px960p | ❌ | ✅ | ✅ | ⭕ |
+| FastHunyuan | `FastVideo/FastHunyuan-diffusers` | 720px1280p<br>544px960p | ❌ | ✅ | ✅ | ⭕ |
+| Wan2.1 T2V 1.3B | `Wan-AI/Wan2.1-T2V-1.3B-Diffusers` | 480P | ✅ | ✅* | ✅ | ⭕ |
+| Wan2.1 T2V 14B | `Wan-AI/Wan2.1-T2V-14B-Diffusers` | 480P, 720P | ✅ | ✅* | ✅ | ⭕ |
+| Wan2.1 I2V 480P | `Wan-AI/Wan2.1-I2V-14B-480P-Diffusers` | 480P | ✅ | ✅* | ✅ | ⭕ |
+| Wan2.1 I2V 720P | `Wan-AI/Wan2.1-I2V-14B-720P-Diffusers` | 720P | ✅ | ✅ | ✅ | ⭕ |
+| StepVideo T2V | `FastVideo/stepvideo-t2v-diffusers` | 768px768px204f<br>544px992px204f<br>544px992px136f | ❌ | ❌ | ✅ | ⭕ |

 **Note**: Wan2.2 TI2V 5B has some quality issues when performing I2V generation. We are working on fixing this issue.

@@ -70,13 +64,3 @@ The `HuggingFace Model ID` can be directly pass to `from_pretrained()` methods a

 ### Sliding Tile Attention
 - Currently only Hopper GPUs (H100s) are supported.
-
-### TurboWan2.1 (TurboDiffusion)
- Uses TurboDiffusionPipeline with RCM scheduler for 1-4 step generation
- Requires SLA attention backend: `export FASTVIDEO_ATTENTION_BACKEND=SLA_ATTN`
- Uses `guidance_scale=1.0` (no classifier-free guidance)
-
-### Matrix Game 2.0
- Image-to-video game world models with keyboard/mouse control input
- Three variants available: Base (universal), GTA, and TempleRun
- Each variant has different keyboard dimensions for control inputs
@@ -1,130 +1,45 @@
-# 🧱 Data Preprocessing

-To save GPU memory during training, FastVideo precomputes text embeddings and VAE latents. This eliminates the need to load the text encoder and VAE during training.
+# 🧱 Data Preprocess

-## Quick Start
+To save GPU memory, we precompute text embeddings and VAE latents to eliminate the need to load the text encoder and VAE during training.

-Download the sample dataset and run preprocessing:
+We provide a sample dataset to help you get started. Download the source media using the following command:

 ```bash
-# Download the crush-smol dataset
-python scripts/huggingface/download_hf.py \
-    --repo_id "wlsaidhi/crush-smol-merged" \
-    --local_dir "data/crush-smol" \
-    --repo_type "dataset"
-
-# Run preprocessing
-bash examples/training/finetune/wan_t2v_1.3B/crush_smol/preprocess_wan_data_t2v_new.sh
+python scripts/huggingface/download_hf.py --repo_id=FastVideo/mini_i2v_dataset --local_dir=data/mini_i2v_dataset --repo_type=dataset
 ```

-## Preprocessing Pipeline
+The folder `crush-smol_raw/` contains raw videos and captions for testing preprocessing, while `crush-smol_preprocessed/` contains latents prepared for testing training.

-The new preprocessing pipeline supports multiple dataset formats and video loaders:
-
-```bash
-GPU_NUM=2
-MODEL_PATH="Wan-AI/Wan2.1-T2V-1.3B-Diffusers"
-DATASET_PATH="data/crush-smol/"
-OUTPUT_DIR="data/crush-smol_processed_t2v/"
-
-torchrun --nproc_per_node=$GPU_NUM \
-    -m fastvideo.pipelines.preprocess.v1_preprocessing_new \
-    --model_path $MODEL_PATH \
-    --mode preprocess \
-    --workload_type t2v \
-    --preprocess.video_loader_type torchvision \
-    --preprocess.dataset_type merged \
-    --preprocess.dataset_path $DATASET_PATH \
-    --preprocess.dataset_output_dir $OUTPUT_DIR \
-    --preprocess.preprocess_video_batch_size 2 \
-    --preprocess.dataloader_num_workers 0 \
-    --preprocess.max_height 480 \
-    --preprocess.max_width 832 \
-    --preprocess.num_frames 77 \
-    --preprocess.train_fps 16 \
-    --preprocess.samples_per_file 8 \
-    --preprocess.flush_frequency 8 \
-    --preprocess.video_length_tolerance_range 5
-```
-
-### Key Parameters
-
-| Parameter | Description |
-|-----------|-------------|
-| `--workload_type` | Task type: `t2v` (text-to-video) or `i2v` (image-to-video) |
-| `--preprocess.dataset_type` | Input format: `hf` (HuggingFace) or `merged` (local folder) |
-| `--preprocess.dataset_path` | Path to dataset (HF repo ID or local folder) |
-| `--preprocess.dataset_output_dir` | Output directory for Parquet files |
-| `--preprocess.video_loader_type` | Video decoder: `torchcodec` or `torchvision` |
-| `--preprocess.max_height` / `max_width` | Target resolution for videos |
-| `--preprocess.num_frames` | Number of frames to extract per video |
-| `--preprocess.train_fps` | Target FPS for frame extraction |
-
-## Dataset Formats
-
-### Merged Dataset (Local Folder)
-
-Structure your dataset as follows:
+To preprocess the dataset for fine-tuning or distillation, run:

 ```
-your_dataset/
-├── videos/
-│   ├── video_001.mp4
-│   ├── video_002.mp4
-│   └── ...
-└── videos2caption.json
+bash scripts/preprocess/v1_preprocess_wan_data_t2v # for wan
 ```

-The `videos2caption.json` maps video filenames to captions:
+## Process your own dataset

-```json
-[
-  {"path": "video_001.mp4", "cap": "A cat playing with yarn..."},
-  {"path": "video_002.mp4", "cap": "Ocean waves at sunset..."}
-]
-```
-
-### HuggingFace Dataset
-
-Use `--preprocess.dataset_type hf` and point `--preprocess.dataset_path` to a HuggingFace dataset with `video` and `caption` columns.
-
-## Creating Your Own Dataset
-
-If you have raw videos and captions in separate files, generate the `videos2caption.json`:
-
-```bash
-python scripts/dataset_preparation/prepare_json_file.py \
-    --data_folder path/to/your_raw_data/ \
-    --output path/to/output_folder
-```
-
-Your raw data folder should contain:
+If you wish to create your own dataset for finetuning or distillation, please refer `mini_i2v_dataset/crush-smol_raw/` to structure you video dataset in the following format:

 ```
-your_raw_data/
+path_to_your_dataset_folder/
 ├── videos/
 │   ├── 0.mp4
 │   ├── 1.mp4
-│   └── ...
-├── videos.txt    # list of video filenames
-└── prompt.txt    # corresponding captions (one per line)
+├── videos.txt
+└── prompt.txt
 ```

-## Output Format
+To generate the `videos2caption.json` and `merge.txt`, run

-Preprocessing outputs Parquet files in the `combined_parquet_dataset/` subdirectory containing:
+``` python
+python scripts/dataset_preparation/prepare_json_file.py --data_folder mini_i2v_dataset/crush-smol_raw/ --output your_output_folder
+```

- `vae_latent_bytes` — VAE-encoded video latent
- `text_embedding_bytes` — text encoder output
- `clip_feature_bytes` — CLIP image features (I2V only)
- `first_frame_latent_bytes` — first frame latent (I2V only)
- Metadata: shapes, dtypes, and sample identifiers
+Adjust the `DATA_MERGE_PATH` and `OUTPUT_DIR` in `scripts/preprocess/v1_preprocess_****.sh` accordingly and run:

-## Examples
+```
+bash scripts/preprocess/v1_preprocess_****.sh
+```

-See ready-to-run preprocessing scripts in the training examples:
-
- **T2V**: `examples/training/finetune/wan_t2v_1.3B/crush_smol/preprocess_wan_data_t2v_new.sh`
- **I2V**: `examples/training/finetune/wan_i2v_14B_480p/crush_smol/preprocess_wan_data_i2v_new.sh`
-
-**→ [Browse all training examples](examples/examples_training_index.md)**
+The preprocessed data will be put into the `OUTPUT_DIR` and the `videos2caption.json` can be used in finetune and distill scripts.
@@ -1,176 +1,78 @@
-# 🧠 Finetuning

-This guide covers finetuning video diffusion models with FastVideo, including full finetuning and LoRA.
-
-## Training Arguments
-
-FastVideo training scripts use several argument groups:
-
-### Training Arguments
-
-| Argument | Description |
-|----------|-------------|
-| `--max_train_steps` | Total training steps |
-| `--train_batch_size` | Batch size per GPU |
-| `--gradient_accumulation_steps` | Steps to accumulate before optimizer update |
-| `--num_latent_t` | Temporal latent dimension (reduce to save memory) |
-| `--num_height` / `--num_width` | Video resolution |
-| `--num_frames` | Number of frames per video |
-| `--output_dir` | Directory for checkpoints |
-
-### Parallelism Arguments
-
-| Argument | Description |
-|----------|-------------|
-| `--num_gpus` | Total number of GPUs |
-| `--sp_size` | Sequence parallel size (increase to reduce memory per GPU) |
-| `--tp_size` | Tensor parallel size |
-| `--hsdp_replicate_dim` | HSDP replication dimension |
-| `--hsdp_shard_dim` | HSDP sharding dimension |
-
-### Optimizer Arguments
-
-| Argument | Description |
-|----------|-------------|
-| `--learning_rate` | Base learning rate |
-| `--mixed_precision` | Precision mode (`bf16` recommended) |
-| `--weight_decay` | Weight decay for regularization |
-| `--max_grad_norm` | Gradient clipping threshold |
-
-### Validation Arguments
-
-| Argument | Description |
-|----------|-------------|
-| `--log_validation` | Enable validation logging |
-| `--validation_dataset_file` | JSON file with validation prompts |
-| `--validation_steps` | Run validation every N steps |
-| `--validation_sampling_steps` | Inference steps for validation |
-| `--validation_guidance_scale` | CFG scale for validation |
-
-## Full Finetuning
-
-Full finetuning updates all model weights. This provides the best quality but requires more GPU memory.
+# 🧠 Finetune
+## ⚡ Full Finetune
+Ensure your data is prepared and preprocessed in the format specified in [data_preprocess.md](#v0-data-preprocess). For convenience, we also provide a mochi preprocessed Black Myth Wukong data that can be downloaded directly:

 ```bash
-# Example: Wan2.1 T2V 1.3B full finetune (4 GPUs)
-bash examples/training/finetune/wan_t2v_1.3B/crush_smol/finetune_t2v.sh
+python scripts/huggingface/download_hf.py --repo_id=FastVideo/Mochi-Black-Myth --local_dir=data/Mochi-Black-Myth --repo_type=dataset
 ```

-**Typical settings:**
+Download the original model weights as specified in the [Distillation Section](../distillation/dmd.md):

- Learning rate: `1e-5` to `5e-5`
- Gradient checkpointing: `--enable_gradient_checkpointing_type "full"`
- Memory scaling: Increase `--sp_size` or reduce `--num_latent_t` to fit in memory
+Then you can run the finetune with:

-## LoRA Finetuning
+```
+bash scripts/finetune/finetune_mochi.sh # for mochi
+```

-LoRA (Low-Rank Adaptation) trains lightweight adapters while keeping the base model frozen. This significantly reduces memory usage and training time.
-
-### LoRA-Specific Arguments
-
-| Argument | Description |
-|----------|-------------|
-| `--lora_training True` | Enable LoRA mode |
-| `--lora_rank` | Rank of LoRA adapters (16, 32, 64, 128) |
-
-### Learning Rate for LoRA
-
-**Important:** LoRA typically requires a **10–20× higher learning rate** than full finetuning because only the low-rank adapters are being trained while the base model is frozen.
-
-| Training Mode | Recommended Learning Rate |
-|---------------|---------------------------|
-| Full finetune | `1e-5` to `5e-5` |
-| LoRA | `1e-4` to `2e-4` |
-
-### Example LoRA Training
+**Note that for finetuning, we did not tune the hyperparameters in the provided script.**
+## ⚡ Finetune with VSA
+Follow [data_preprocess.md](#v0-data-preprocess) to get parquet files for preproccessed latent, and then run:

 ```bash
-# Example: Wan2.1 T2V 1.3B LoRA finetune (1 GPU)
-bash examples/training/finetune/wan_t2v_1.3B/crush_smol/finetune_t2v_lora.sh
+bash scripts/finetune/finetune_v1_VSA.sh
 ```

-Key differences from full finetune:
+## ⚡ Lora Finetune

- Add `--lora_training True --lora_rank 32`
- Use higher learning rate (10–20× full finetune)
- Can run on fewer GPUs (even single GPU)
- Outputs adapter weights instead of full model
-
-## LoRA Extraction and Merging
-
-FastVideo provides tools to extract LoRA adapters from finetuned models and merge them back.
-
-### Extract LoRA Adapter
-
-Extract a LoRA adapter by comparing a finetuned model to its base:
+Hunyuan supports Lora fine-tuning of videos up to 720p. Demos and prompts of Black-Myth-Wukong can be found in [here](https://huggingface.co/FastVideo/Hunyuan-Black-Myth-Wukong-lora-weight). You can download the Lora weight through:

 ```bash
-python scripts/lora_extraction/extract_lora.py \
-  --base Wan-AI/Wan2.1-T2V-1.3B-Diffusers \
-  --ft path/to/your/finetuned_model \
-  --out adapter_r32.safetensors \
-  --rank 32
+python scripts/huggingface/download_hf.py --repo_id=FastVideo/Hunyuan-Black-Myth-Wukong-lora-weight --local_dir=data/Hunyuan-Black-Myth-Wukong-lora-weight --repo_type=model
 ```

-| Argument | Description |
-|----------|-------------|
-| `--base` | Base model (HuggingFace ID or local path) |
-| `--ft` | Finetuned model path |
-| `--out` | Output adapter file (.safetensors) |
-| `--rank` | LoRA rank (16, 32, 64, 128) |
-| `--full-rank` | Extract full-rank adapter (optional) |
+### Minimum Hardware Requirement
+- 40 GB GPU memory each for 2 GPUs with lora.
+- 30 GB GPU memory each for 2 GPUs with CPU offload and lora.

-### Merge LoRA Adapter
+Currently, both Mochi and Hunyuan models support Lora finetuning through diffusers. To generate personalized videos from your own dataset, you'll need to follow three main steps: dataset preparation, finetuning, and inference.

-Merge an adapter back into a base model:
+### Dataset Preparation
+We provide scripts to better help you get started to train on your own characters!
+You can run this to organize your dataset to get the videos2caption.json before preprocess. Specify your video folder and corresponding caption folder (caption files should be .txt files and have the same name with its video):
+
+```
+python scripts/dataset_preparation/prepare_json_file.py --video_dir data/input_videos/ --prompt_dir data/captions/ --output_path data/output_folder/videos2caption.json --verbose
+```
+
+Also, we provide script to resize your videos:
+
+```
+python scripts/data_preprocess/resize_videos.py
+```
+
+### Finetuning
+After basic dataset preparation and preprocess, you can start to finetune your model using Lora:
+
+```
+bash scripts/finetune/finetune_hunyuan_hf_lora.sh
+```
+
+### Inference
+For inference with Lora checkpoint, you can run the following scripts with additional parameter `--lora_checkpoint_dir`:
+
+```
+bash scripts/inference/inference_hunyuan_hf.sh
+```
+
+**We also provide scripts for Mochi in the same directory.**
+
+### Finetune with Both Image and Video
+Our codebase support finetuning with both image and video.

 ```bash
-python scripts/lora_extraction/merge_lora.py \
-  --base Wan-AI/Wan2.1-T2V-1.3B-Diffusers \
-  --adapter adapter_r32.safetensors \
-  --ft path/to/your/finetuned_model \
-  --output merged_model
+bash scripts/finetune/finetune_hunyuan.sh
+bash scripts/finetune/finetune_mochi_lora_mix.sh
 ```

-| Argument | Description |
-|----------|-------------|
-| `--base` | Base model path |
-| `--adapter` | LoRA adapter file |
-| `--ft` | Finetuned model (for config reference) |
-| `--output` | Output directory for merged model |
-
-### Validate Merged Model
-
-Compare the merged model against the original finetuned model:
-
-```bash
-python scripts/lora_extraction/lora_inference_comparison.py \
-  --base merged_model \
-  --ft path/to/your/finetuned_model \
-  --adapter NONE \
-  --output-dir results \
-  --prompt "A cat sitting on a windowsill" \
-  --compute-ssim \
-  --compute-lpips
-```
-
-## Training Examples
-
-Ready-to-run training scripts are available for multiple models:
-
-**→ [Browse all training examples](examples/examples_training_index.md)**
-
-| Model | Type | Example |
-|-------|------|---------|
-| Wan2.1 T2V 1.3B | T2V | `examples/training/finetune/wan_t2v_1.3B/crush_smol/` |
-| Wan2.1 I2V 14B | I2V | `examples/training/finetune/wan_i2v_14B_480p/crush_smol/` |
-| Wan2.1-Fun 1.3B InP | I2V | `examples/training/finetune/Wan2.1-Fun-1.3B-InP/crush_smol/` |
-| Wan2.1 VSA | T2V/I2V | `examples/training/finetune/Wan2.1-VSA/Wan-Syn-Data/` |
-
-Each example includes:
-
- `download_dataset.sh` — download sample data
- `preprocess_*.sh` — run preprocessing
- `finetune_*.sh` — full finetune launcher
- `finetune_*_lora.sh` — LoRA finetune launcher
- `validation.json` — validation prompts
+For Image-Video Mixture Fine-tuning, make sure to enable the `--group_frame` option in your script.
@@ -1,67 +0,0 @@
-# Training Overview
-
-FastVideo supports finetuning video diffusion models on custom datasets. This page explains what data you need and how to get started.
-
-## Data Requirements
-
-To save GPU memory during training, FastVideo precomputes embeddings and latents ahead of time. This eliminates the need to load the text encoder and VAE during training, significantly reducing memory usage.
-
-### Text-to-Video (T2V) Finetuning
-
-For T2V models, you need:
-
-| Component | Description |
-|-----------|-------------|
-| **Text embeddings** | Precomputed embeddings from the model's text encoder (e.g., T5 or LLaMA). Stored as numpy arrays in Parquet files. |
-| **Video latents** | VAE-encoded representations of your training videos. Each video is encoded into a compressed latent tensor. |
-
-### Image-to-Video (I2V) Finetuning
-
-For I2V models, you need everything from T2V plus additional image conditioning. Note that not all I2V architectures require encoded images—this depends on how the model conditions on the input frame. Wan2.1 and Wan2.2 A14B I2V models do require these additional components:
-
-| Component | Description |
-|-----------|-------------|
-| **Text embeddings** | Same as T2V—precomputed from the text encoder. |
-| **Video latents** | Same as T2V—VAE-encoded video representations. |
-| **First frame latent** | VAE-encoded representation of the first frame, used as the conditioning image. |
-| **CLIP features** | Image embeddings from a CLIP vision encoder for the conditioning frame. |
-
-## Preprocessing
-
-Before training, you need to preprocess your raw videos and captions into Parquet files containing precomputed latents and embeddings.
-
-FastVideo supports two input formats:
-
- **HuggingFace datasets** — load directly from HF Hub or local HF datasets
- **Merged datasets** — local folder with videos and a `videos2caption.json` metadata file
-
-**→ See [Data Preprocessing](data_preprocess.md) for full details and examples.**
-
-## Training Examples
-
-Ready-to-run examples with preprocessing scripts, training launchers, and validation configs are available for multiple models and datasets:
-
-**→ [Browse all training examples](examples/examples_training_index.md)**
-
-Each example includes:
-
- `download_dataset.sh` — download sample data
- `preprocess_*.sh` — run preprocessing
- `finetune_*.sh` — launch training (full finetune or LoRA)
- `validation.json` — validation prompts for checkpoints
-
-## Training Methods
-
-FastVideo supports several training approaches:
-
-| Method | Use Case |
-|--------|----------|
-| **Full finetune** | Adapt entire model to a new domain or style |
-| **LoRA finetune** | Lightweight adaptation with frozen base weights |
-| **VSA finetune** | Finetune with Variable Sparse Attention for efficiency |
-
-## Next Steps
-
-1. **Get started**: Pick an example from the [training examples index](examples/examples_training_index.md)
-2. **Prepare data**: Follow [data preprocessing](data_preprocess.md) for your own dataset
-3. **Run inference**: After training, see [inference examples](../inference/examples/examples_inference_index.md)
@@ -1,71 +0,0 @@
-# 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
@@ -1,169 +0,0 @@
-#!/bin/bash
-#SBATCH --job-name=dmd_3333
-#SBATCH --partition=main
-#SBATCH --nodes=4
-#SBATCH --ntasks=4
-#SBATCH --ntasks-per-node=1
-#SBATCH --gres=gpu:8
-#SBATCH --cpus-per-task=128
-#SBATCH --mem=1440G
-#SBATCH --output=dmd_3333_output/dmd_3333_%j.out
-#SBATCH --error=dmd_3333_output/dmd_3333_%j.err
-#SBATCH --exclusive
-
-# Basic Info
-export NCCL_P2P_DISABLE=1
-export TORCH_NCCL_ENABLE_MONITORING=0
-export NCCL_DEBUG_SUBSYS=INIT,NET
-# different cache dir for different processes
-export TRITON_CACHE_DIR=/tmp/triton_cache_${SLURM_PROCID}
-export MASTER_PORT=29500
-export NODE_RANK=$SLURM_PROCID
-nodes=( $(scontrol show hostnames $SLURM_JOB_NODELIST) )
-export MASTER_ADDR=${nodes[0]}
-export TOKENIZERS_PARALLELISM=false
-# export WANDB_API_KEY="8d9f4b39abd68eb4e29f6fc010b7ee71a2207cde"
-export WANDB_API_KEY="50632ebd88ffd970521cec9ab4a1a2d7e85bfc45"
-# export WANDB_API_KEY='your_wandb_api_key_here'
-export WANDB_BASE_URL="https://api.wandb.ai"
-export WANDB_MODE=online
-export FASTVIDEO_ATTENTION_BACKEND=FLASH_ATTN
-
-source ~/conda/miniconda/bin/activate
-conda activate wei-fv-distill
-export HOME="/mnt/weka/home/hao.zhang/wei"
-
-# Configs
-NUM_GPUS=8
-NUM_TOTAL_GPUS=32
-
-# Model paths for Self-Forcing DMD distillation with Hunyuan1.5:
-GENERATOR_MODEL_PATH="weizhou03/SFHunyuanVideo-1.5-Diffusers-480p_t2v"  # Updated to Wan2.2
-REAL_SCORE_MODEL_PATH="hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v" # Teacher model
-FAKE_SCORE_MODEL_PATH="hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v"
-# REAL_SCORE_MODEL_PATH="hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_i2v_step_distilled"
-# FAKE_SCORE_MODEL_PATH="hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_i2v_step_distilled" # Critic model
-
-# DATA_DIR="data/test-text-preprocessing/Node_0_GPU_1_File_1/combined_parquet_dataset/"
-# DATA_DIR="/mnt/weka/home/hao.zhang/matthew/FastVideo/data/test-text-preprocessing"
-DATA_DIR="/mnt/weka/home/hao.zhang/wl/sharefs/wl/release/vidprom_16k_text_embed"
-DATA_DIR_2="/mnt/weka/home/hao.zhang/wl/sharefs/wl/release/fv-ode-preprocessing-16k-hy15-121"
-# DATA_DIR=data/crush-smol_processed_t2v/combined_parquet_dataset
-# DATA_DIR="/mnt/sharefs/users/hao.zhang/Vchitect-2M/Wan-Syn-upload/latents_i2v/train/"
-# VALIDATION_DATASET_FILE="data/crush-smol-single_processed_t2v/validation.json"
-VALIDATION_DATASET_FILE="/mnt/weka/home/hao.zhang/wl/FastVideo/examples/distill/Wan2.2-TI2V-5B-Diffusers/Data-free/validation_64.json"
-
-# export CUDA_VISIBLE_DEVICES=4,5
-# IP=[MASTER NODE IP]
-
-training_args=(
-  --tracker_project_name SFhy1.5_t2v_distill_self_forcing_dmd  # Updated for Wan2.2
-  --output_dir "/mnt/weka/home/hao.zhang/wei/SFhy1.5_distill_3333_1e-5_1e-5_cfg6_corrected_scheduler"
-  --wandb_run_name "self_forcing_3333_context_forcing"
-  --max_train_steps 4000
-  --train_batch_size 1
-  --train_sp_batch_size 1
-  --gradient_accumulation_steps 1
-  --num_latent_t 21
-  --num_height 480
-  --num_width 848
-  --enable_gradient_checkpointing_type "full"
-  --simulate_generator_forward
-  --log_visualization
-  --num_frames 81
-  --num_frame_per_block 3  # Frame generation block size for self-forcing
-  # --resume-from-checkpoint "/mnt/weka/home/hao.zhang/wei/SFhy1.5_distill_1333/checkpoint-300"
-  --init_weights_from_safetensors /mnt/weka/home/hao.zhang/wei/hy15_worldplay_df_init_3333/checkpoint-2000/transformer/diffusion_pytorch_model.safetensors
-  # --init_weights_from_safetensors /mnt/weka/home/hao.zhang/wei/hy15_distilled_ode_init_3333/checkpoint-2400/transformer/diffusion_pytorch_model.safetensors
-)
-
-parallel_args=(
-  --num_gpus $NUM_TOTAL_GPUS # 64
-  --sp_size 1
-  --tp_size 1
-  --hsdp_replicate_dim 1 # 64
-  --hsdp_shard_dim $NUM_TOTAL_GPUS
-)
-
-model_args=(
-  --model_path $GENERATOR_MODEL_PATH
-  --pretrained_model_name_or_path $GENERATOR_MODEL_PATH
-  --real_score_model_path $REAL_SCORE_MODEL_PATH
-  --fake_score_model_path $FAKE_SCORE_MODEL_PATH
-)
-
-dataset_args=(
-  --data_path "$DATA_DIR"
-  # --data_path_2 "$DATA_DIR_2"
-  --dataloader_num_workers 4
-)
-
-validation_args=(
-  --log_validation
-  --validation_dataset_file "$VALIDATION_DATASET_FILE"
-  --validation_steps 50
-  --validation_sampling_steps "8"
-  --validation_guidance_scale "1.0" # not used for dmd inference
-  --text-encoder-cpu-offload
-  # --vae_cpu_offload True
-)
-
-optimizer_args=(
-  --learning_rate 1e-5
-  --mixed_precision "bf16"
-  --training_state_checkpointing_steps 100
-  --weight_only_checkpointing_steps 500
-  --weight_decay 0.01
-  --betas '0.0,0.999'
-  --max_grad_norm 1.0
-)
-
-miscellaneous_args=(
-  --inference_mode False
-  --checkpoints_total_limit 3
-  --training_cfg_rate 0.0
-  --dit_precision "fp32"
-  --flow_shift 5
-  --seed 1000
-  --use_ema True
-  --ema_decay 0.99
-  --ema_start_step 200
-)
-
-dmd_args=(
-  --warp_denoising_step
-  --dmd_denoising_steps '1000,875,750,625,500,375,250,125'
-  # --dmd_denoising_steps '1000,760,520,280'
-  # --dmd_denoising_steps '1000,750,500,250'
-  --min_timestep_ratio 0.02
-  --max_timestep_ratio 0.98
-  --dfake_gen_update_ratio 5
-  --real_score_guidance_scale 3.5
-  --fake_score_learning_rate 5e-6
-  --fake_score_betas '0.0,0.999'
-)
-
-self_forcing_args=(
-  --independent_first_frame False  # Whether to treat first frame independently
-  --same_step_across_blocks True  # Whether to use same denoising step across all blocks
-  --last_step_only False  # Whether to only use the last denoising step
-  --context_noise 0  # Amount of noise to add during context caching (0 = no noise)
-  # --use-context-forcing True
-)
-
-srun torchrun \
--nnodes $SLURM_JOB_NUM_NODES \
--nproc_per_node $NUM_GPUS \
--node_rank $SLURM_PROCID \
--rdzv_backend=c10d \
--rdzv_endpoint="$MASTER_ADDR:$MASTER_PORT" \
-    fastvideo/training/hy15_self_forcing_distillation_pipeline.py \
-    "${parallel_args[@]}" \
-    "${model_args[@]}" \
-    "${dataset_args[@]}" \
-    "${training_args[@]}" \
-    "${optimizer_args[@]}" \
-    "${validation_args[@]}" \
-    "${miscellaneous_args[@]}" \
-    "${dmd_args[@]}" \
-    "${self_forcing_args[@]}"
@@ -1,19 +0,0 @@
-# Self-Forcing Distillation for SFWan2.1 T2V 1.3B
-
-These scripts demonstrate self-forcing distillation (SFwan) for the causal Wan2.1 T2V 1.3B model. The workflow mirrors DMD2 while injecting self-forcing blocks so the student can autoregressively refine later frames.
-
-## Run the recipe
-1. Download the preprocessed text-video dataset:
-   ```bash
-   bash examples/distill/SFWan2.1-T2V/download_dataset.sh
-   ```
-2. (Optional) Regenerate parquet shards locally:
-   ```bash
-   bash examples/distill/SFWan2.1-T2V/preprocess_data.sh
-   ```
-3. Launch self-forcing distillation with your cluster settings:
-   ```bash
-   sbatch examples/distill/SFWan2.1-T2V/distill_dmd_t2v_1.3B.sh
-   ```
-
-Update the dataset paths and wandb credentials inside the script before running on your environment.
@@ -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=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        dit_cpu_offload=False,
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -1,42 +0,0 @@
-from fastvideo import VideoGenerator
-
-
-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,
-    )
-
-    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."
-    )
-
-    video = generator.generate_video(
-        prompt,
-        negative_prompt="",
-        height=704,
-        width=1280,
-        num_frames=77,
-        num_inference_steps=35,
-        guidance_scale=7.0,
-        fps=24,
-        output_path="outputs_video/cosmos2_5_t2w.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=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        # 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"
@@ -2,51 +2,31 @@ from fastvideo import VideoGenerator
 import json
 # from fastvideo.configs.sample import SamplingParam

-OUTPUT_PATH = "video_samples_hy15_t2v_distilled"
+OUTPUT_PATH = "video_samples_hy15"
 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(
-        "hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_i2v_step_distilled",
+        "hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v",
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        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,
-        # init_weights_from_safetensors="/mnt/weka/home/hao.zhang/wei/SFhy1.5_distill_1333_5e-5_2e-6_cfg3.5/checkpoint-500/ema/generator_ema.safetensors"
    )

-    # json_path = "/mnt/weka/home/hao.zhang/wei/FastVideo/data/mixkit_i2v_full_720p.json"
-    # with open(json_path, 'r') as f:
-    #     data_list = json.load(f)["data"]
-    
-    # # Now you can index into data_list however you like
-    # # For example: data_list[0], data_list[1:3], etc.
-    # for data in data_list:
-    #     prompt = data["prompt"]
-    #     image_path = data["image_path"]
-    #     generator.generate_video(prompt, output_path=OUTPUT_PATH, save_video=True, image_path=image_path, num_frames=121, fps=24)
-    # return
-
-    # prompt = (
-    #     "In a brightly lit studio, a photographer wearing a denim jacket focuses intently, capturing shots with a professional camera. Facing him, a model stands gracefully, adjusting her long, flowing hair with delicate movements. The scene is characterized by strong contrasts; the model's soft pink attire and gentle gestures complement the rugged, precise demeanor of the photographer. Positioned against a minimalist backdrop, the pair work seamlessly, with the camera\u2019s lens pointed directly at the model, capturing her elegance. The soft, diffused lighting casts a gentle glow on both subjects, creating an airy and ethereal atmosphere perfect for a high-fashion photo shoot."
-    # )
-
-    # video = generator.generate_video(prompt, output_path=OUTPUT_PATH, save_video=True, negative_prompt="", num_frames=121, fps=24, image_path="data/1.png")
-    # return
-
    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)
+    video = generator.generate_video(prompt, output_path=OUTPUT_PATH, save_video=True, negative_prompt="", num_frames=81, fps=16)

    prompt2 = (
        "A majestic lion strides across the golden savanna, its powerful frame "
@@ -55,7 +35,7 @@ def main():
        "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)
+    video2 = generator.generate_video(prompt2, output_path=OUTPUT_PATH, save_video=True, negative_prompt="", num_frames=81, fps=16)


 if __name__ == "__main__":
@@ -1,224 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-"""
-Basic example for HYWorld (HY-WorldPlay) video generation using FastVideo.
-
-This example replicates the same functionality as HY-WorldPlay/run.sh,
-demonstrating image-to-video generation with camera trajectory control.
-"""
-
-import time
-import math
-import numpy as np
-import imageio
-import torchvision
-from einops import rearrange
-
-from fastvideo import VideoGenerator
-from fastvideo.pipelines import ForwardBatch
-from fastvideo.utils import shallow_asdict, align_to
-from fastvideo.logger import init_logger
-
-from fastvideo.models.dits.hyworld.pose import pose_to_input, compute_latent_num
-from fastvideo.models.dits.hyworld.resolution_utils import get_resolution_from_image
-
-logger = init_logger(__name__)
-
-class HYWorldVideoGenerator(VideoGenerator):
-    """Extended VideoGenerator that adds HYWorld-specific parameters to batch.extra."""
-
-    def _generate_single_video(self, prompt: str, sampling_param=None, **kwargs):
-        """Override to add viewmats, Ks, and action to batch.extra."""
-        fastvideo_args = self.fastvideo_args
-        pipeline_config = fastvideo_args.pipeline_config
-
-        if sampling_param is None:
-            from fastvideo.configs.sample import SamplingParam
-            sampling_param = SamplingParam.from_pretrained(fastvideo_args.model_path)
-
-        # Update sampling param with kwargs
-        if kwargs:
-            for key, value in kwargs.items():
-                if hasattr(sampling_param, key):
-                    setattr(sampling_param, key, value)
-
-        # Get pose string from sampling_param or kwargs
-        pose = kwargs.get('pose', getattr(sampling_param, 'POSE', 'w-31'))
-        num_frames = kwargs.get('num_frames', getattr(sampling_param, 'num_frames', 125))
-
-        # Calculate number of latents
-        latent_num = compute_latent_num(num_frames)
-
-        # Convert pose to viewmats, Ks, and action
-        viewmats, Ks, action = pose_to_input(pose, latent_num)
-
-        # Convert to tensors and add batch dimension
-        viewmats = viewmats.unsqueeze(0)  # (1, T, 4, 4)
-        Ks = Ks.unsqueeze(0)  # (1, T, 3, 3)
-        action = action.unsqueeze(0)  # (1, T)
-
-        # Validate inputs
-        prompt = prompt.strip()
-        sampling_param = sampling_param.__class__(**shallow_asdict(sampling_param))
-        output_path = kwargs.get("output_path", sampling_param.output_path)
-        sampling_param.prompt = prompt
-
-        if sampling_param.negative_prompt is not None:
-            sampling_param.negative_prompt = sampling_param.negative_prompt.strip()
-
-        # Validate dimensions
-        if (sampling_param.height <= 0 or sampling_param.width <= 0 or
-            sampling_param.num_frames <= 0):
-            raise ValueError(
-                f"Height, width, and num_frames must be positive integers")
-
-        temporal_scale_factor = pipeline_config.vae_config.arch_config.temporal_compression_ratio
-        num_frames = sampling_param.num_frames
-        num_gpus = fastvideo_args.num_gpus
-        use_temporal_scaling_frames = pipeline_config.vae_config.use_temporal_scaling_frames
-
-        # Adjust number of frames based on number of GPUs
-        if use_temporal_scaling_frames:
-            orig_latent_num_frames = (num_frames - 1) // temporal_scale_factor + 1
-        else:
-            orig_latent_num_frames = sampling_param.num_frames // 17 * 3
-
-        if orig_latent_num_frames % fastvideo_args.num_gpus != 0:
-            if use_temporal_scaling_frames:
-                new_num_frames = (orig_latent_num_frames - 1) * temporal_scale_factor + 1
-            else:
-                divisor = math.lcm(3, num_gpus)
-                orig_latent_num_frames = (
-                    (orig_latent_num_frames + divisor - 1) // divisor) * divisor
-                new_num_frames = orig_latent_num_frames // 3 * 17
-
-            logger.info(
-                "Adjusting number of frames from %s to %s based on number of GPUs (%s)",
-                sampling_param.num_frames, new_num_frames, fastvideo_args.num_gpus)
-            sampling_param.num_frames = new_num_frames
-
-        # Calculate sizes
-        target_height = align_to(sampling_param.height, 16)
-        target_width = align_to(sampling_param.width, 16)
-
-        # Calculate latent sizes
-        latents_size = [(sampling_param.num_frames - 1) // 4 + 1,
-                        sampling_param.height // 8, sampling_param.width // 8]
-        n_tokens = latents_size[0] * latents_size[1] * latents_size[2]
-
-        # Prepare batch
-        batch = ForwardBatch(
-            **shallow_asdict(sampling_param),
-            eta=0.0,
-            n_tokens=n_tokens,
-            VSA_sparsity=fastvideo_args.VSA_sparsity,
-        )
-
-        # Add HYWorld-specific parameters to batch.extra
-        batch.extra['viewmats'] = viewmats
-        batch.extra['Ks'] = Ks
-        batch.extra['action'] = action
-        batch.extra['chunk_latent_frames'] = 16  # For bidirectional model
-
-        # Run inference
-        start_time = time.perf_counter()
-        output_batch = self.executor.execute_forward(batch, fastvideo_args)
-        samples = output_batch.output
-        logging_info = output_batch.logging_info
-
-        gen_time = time.perf_counter() - start_time
-        logger.info("Generated successfully in %.2f seconds", gen_time)
-
-        # Process outputs
-        videos = rearrange(samples, "b c t h w -> t b c h w")
-        frames = []
-        for x in videos:
-            x = torchvision.utils.make_grid(x, nrow=6)
-            x = x.transpose(0, 1).transpose(1, 2).squeeze(-1)
-            frames.append((x * 255).numpy().astype(np.uint8))
-
-        # Save video if requested
-        if batch.save_video:
-            imageio.mimsave(output_path, frames, fps=batch.fps, format="mp4")
-            logger.info("Saved video to %s", output_path)
-
-        if batch.return_frames:
-            return frames
-        else:
-            return {
-                "samples": samples,
-                "frames": frames,
-                "prompts": prompt,
-                "size": (target_height, target_width, batch.num_frames),
-                "generation_time": gen_time,
-                "logging_info": logging_info,
-                "trajectory": output_batch.trajectory_latents,
-                "trajectory_timesteps": output_batch.trajectory_timesteps,
-                "trajectory_decoded": output_batch.trajectory_decoded,
-            }
-
-# 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'
-
-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='video_samples_hyworld', 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 = HYWorldVideoGenerator.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
-    print("\nGenerating video...")
-    start_time = time.time()
-    video = generator.generate_video(
-        prompt=args.prompt,
-        image_path=args.image,
-        output_path=args.output_path,
-        save_video=True,
-        negative_prompt="",
-        num_frames=args.num_frames,
-        fps=24,
-        height=HEIGHT,
-        width=WIDTH,
-        seed=args.seed,
-        pose=args.pose,
-    )
-    elapsed = time.time() - start_time
-
-    print(f"\nVideo generated successfully!")
-    print(f"Saved to: {args.output_path}")
-    print(f"Time: {elapsed:.2f}s")
-
-
-if __name__ == "__main__":
-    main()
@@ -1,210 +0,0 @@
-"""
-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()
-
-
@@ -1,198 +0,0 @@
-"""
-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()
-
-
@@ -1,228 +0,0 @@
-"""
-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()
-
-
@@ -1,34 +0,0 @@
-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()
@@ -43,7 +43,7 @@ def main():
        config["model_path"],
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -1,98 +0,0 @@
-from fastvideo.entrypoints.streaming_generator import StreamingVideoGenerator
-from fastvideo.models.dits.matrix_game.utils import get_current_action_async, expand_action_to_frames
-
-import torch
-import asyncio
-
-# Available variants: "base_distilled_model", "gta_distilled_model", "templerun_distilled_model"
-# Each variant has different keyboard_dim:
-#   - base_distilled_model: keyboard_dim=4
-#   - gta_distilled_model: keyboard_dim=2
-#   - templerun_distilled_model: keyboard_dim=7 (keyboard only, no mouse)
-MODEL_VARIANT = "base_distilled_model"
-
-# Variant-specific settings
-VARIANT_CONFIG = {
-    "base_distilled_model": {
-        "model_path": "FastVideo/Matrix-Game-2.0-Base-Diffusers",
-        "keyboard_dim": 4,
-        "mode": "universal",
-        "image_url": "https://raw.githubusercontent.com/SkyworkAI/Matrix-Game/main/Matrix-Game-2/demo_images/universal/0000.png",
-    },
-    "gta_distilled_model": {
-        "model_path": "FastVideo/Matrix-Game-2.0-GTA-Diffusers",
-        "keyboard_dim": 2,
-        "mode": "gta_drive",
-        "image_url": "https://raw.githubusercontent.com/SkyworkAI/Matrix-Game/main/Matrix-Game-2/demo_images/gta_drive/0000.png",
-    },
-    "templerun_distilled_model": {
-        "model_path": "FastVideo/Matrix-Game-2.0-TempleRun-Diffusers",
-        "keyboard_dim": 7,
-        "mode": "templerun",
-        "image_url": "https://raw.githubusercontent.com/SkyworkAI/Matrix-Game/main/Matrix-Game-2/demo_images/temple_run/0000.png",
-    },
-}
-
-
-OUTPUT_PATH = "video_samples_matrixgame2"
-async 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.
-    config = VARIANT_CONFIG[MODEL_VARIANT]
-
-    generator = StreamingVideoGenerator.from_pretrained(
-        config["model_path"],
-        # 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, # DiT need to be offloaded for MoE
-        vae_cpu_offload=False,
-        text_encoder_cpu_offload=True,
-        # Set pin_cpu_memory to false if CPU RAM is limited and there're no frequent CPU-GPU transfer
-        pin_cpu_memory=True,
-        # image_encoder_cpu_offload=False,
-    )
-
-    max_blocks = 50
-    num_frames = 597    
-    actions = {
-        "keyboard": torch.zeros((num_frames, config["keyboard_dim"])),
-        "mouse": torch.zeros((num_frames, 2))
-    }
-    grid_sizes = torch.tensor([150, 44, 80])
-    mode = config["mode"]
-
-    generator.reset(
-        prompt="",
-        image_path=config["image_url"],
-        mouse_cond=actions["mouse"].unsqueeze(0),
-        keyboard_cond=actions["keyboard"].unsqueeze(0),
-        grid_sizes=grid_sizes,
-        num_frames=num_frames,
-        height=352,
-        width=640,
-        num_inference_steps=50,
-        output_path=OUTPUT_PATH,
-        save_video=True,
-    )
-    print("Initialization complete.")
-
-    for block_id in range(max_blocks):
-        print(f"\n=== Block {block_id + 1}/{max_blocks} ===")
-        
-        action = await get_current_action_async(mode)
-        keyboard_cond, mouse_cond = expand_action_to_frames(action, 12)
-        await generator.step_async(keyboard_cond, mouse_cond)
-        
-        if (await asyncio.to_thread(input, "\nContinue? (y/n): ")).lower() == 'n':
-            break
-
-    # Save final video
-    generator.finalize()
-    generator.shutdown()
-
-
-if __name__ == "__main__":
-    asyncio.run(main())
@@ -13,7 +13,7 @@ def main():
        model_name,
        # FastVideo will automatically handle distributed setup
        num_gpus=1,
-        use_fsdp_inference=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        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=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        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=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        dit_cpu_offload=True, # DiT need to be offloaded for MoE
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -1,55 +0,0 @@
-import os
-
-# Set SLA attention backend BEFORE fastvideo imports
-os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "SLA_ATTN"
-
-from fastvideo import VideoGenerator
-
-OUTPUT_PATH = "video_samples_turbodiffusion"
-
-
-def main() -> None:
-    # TurboDiffusion: 1-4 step video generation using RCM scheduler + SLA attention
-    # FastVideo will automatically use TurboDiffusionPipeline when specified
-    generator = VideoGenerator.from_pretrained(
-        "loayrashid/TurboWan2.1-T2V-1.3B-Diffusers",
-        # FastVideo will automatically handle distributed setup
-        num_gpus=1,
-        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 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 "
-        "natural light filtering through the petals. Mid-shot, warm and cheerful tones."
-    )
-    video = generator.generate_video(
-        prompt,
-        output_path=OUTPUT_PATH,
-        save_video=True,
-        seed=42,
-    )
-
-    # Generate another video with a different prompt, without reloading the model!
-    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,
-        seed=42,
-    )
-
-
-if __name__ == "__main__":
-    main()
@@ -1,49 +0,0 @@
-import os
-
-# Set SLA attention backend BEFORE fastvideo imports
-os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "SLA_ATTN"
-
-from fastvideo import VideoGenerator
-
-OUTPUT_PATH = "video_samples_turbodiffusion_14B"
-
-
-def main() -> None:
-    # TurboDiffusion 14B: 1-4 step video generation using RCM scheduler + SLA attention
-    # FastVideo will automatically use TurboDiffusionPipeline when specified
-    generator = VideoGenerator.from_pretrained(
-        "loayrashid/TurboWan2.1-T2V-14B-Diffusers",
-        # 14B model needs more GPUs
-        num_gpus=2,
-    )
-
-    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,
-        seed=42,
-    )
-
-    # Generate another video with a different prompt, without reloading the model!
-    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,
-        seed=42,
-    )
-
-
-if __name__ == "__main__":
-    main()
@@ -1,37 +0,0 @@
-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=False, # set to True if GPU is out of memory 
+        use_fsdp_inference=True,
        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=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        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=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        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=False, # set to True if GPU is out of memory
+        use_fsdp_inference=True,
        dit_cpu_offload=True,
        vae_cpu_offload=False,
        text_encoder_cpu_offload=True,
@@ -1,684 +0,0 @@
-import argparse
-import asyncio
-import os
-import time
-
-import gradio as gr
-import torch
-import uvicorn
-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
-
-
-VARIANT_CONFIG = {
-    "Matrix-Game-2.0-Base": {
-        "model_path": "FastVideo/Matrix-Game-2.0-Base-Diffusers",
-        "keyboard_dim": 4,
-        "mode": "universal",
-        "image_url": "https://raw.githubusercontent.com/SkyworkAI/Matrix-Game/main/Matrix-Game-2/demo_images/universal/0000.png",
-    },
-    "Matrix-Game-2.0-GTA": {
-        "model_path": "FastVideo/Matrix-Game-2.0-GTA-Diffusers",
-        "keyboard_dim": 2,
-        "mode": "gta_drive",
-        "image_url": "https://raw.githubusercontent.com/SkyworkAI/Matrix-Game/main/Matrix-Game-2/demo_images/gta_drive/0000.png",
-    },
-    "Matrix-Game-2.0-TempleRun": {
-        "model_path": "FastVideo/Matrix-Game-2.0-TempleRun-Diffusers",
-        "keyboard_dim": 7,
-        "mode": "templerun",
-        "image_url": "https://raw.githubusercontent.com/SkyworkAI/Matrix-Game/main/Matrix-Game-2/demo_images/temple_run/0000.png",
-    },
-}
-
-MODEL_PATH_MAPPING = {
-    name: config["model_path"] for name, config in VARIANT_CONFIG.items()
-}
-
-
-CAM_VALUE = 0.1
-KEYBOARD_MAP_UNIVERSAL = {
-    "W (Forward)": [1, 0, 0, 0],
-    "S (Back)": [0, 1, 0, 0],
-    "A (Left)": [0, 0, 1, 0],
-    "D (Right)": [0, 0, 0, 1],
-    "Q (Stop)": [0, 0, 0, 0],
-}
-KEYBOARD_MAP_GTA = {
-    "W (Forward)": [1, 0],
-    "S (Back)": [0, 1],
-    "Q (Stop)": [0, 0],
-}
-KEYBOARD_MAP_TEMPLERUN = {
-    "Q (Run)": [1, 0, 0, 0, 0, 0, 0],
-    "W (Jump)": [0, 1, 0, 0, 0, 0, 0],
-    "S (Slide)": [0, 0, 1, 0, 0, 0, 0],
-    "Z (Turn Left)": [0, 0, 0, 1, 0, 0, 0],
-    "C (Turn Right)": [0, 0, 0, 0, 1, 0, 0],
-    "A (Left)": [0, 0, 0, 0, 0, 1, 0],
-    "D (Right)": [0, 0, 0, 0, 0, 0, 1],
-}
-
-
-CAMERA_MAP_UNIVERSAL = {
-    "U (Center)": [0, 0],
-    "I (Up)": [CAM_VALUE, 0],
-    "K (Down)": [-CAM_VALUE, 0],
-    "J (Left)": [0, -CAM_VALUE],
-    "L (Right)": [0, CAM_VALUE],
-}
-CAMERA_MAP_GTA = {
-    "Q (Straight)": [0, 0],
-    "A (Steer Left)": [0, -CAM_VALUE],
-    "D (Steer Right)": [0, CAM_VALUE],
-}
-
-def setup_model_environment(model_path: str) -> None:
-    # if "fullattn" in model_path.lower():
-    #     os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "FLASH_ATTN"
-    # else:
-    #     os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "VIDEO_SPARSE_ATTN"
-    os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "FLASH_ATTN"
-    os.environ["FASTVIDEO_STAGE_LOGGING"] = "1"
-
-def create_timing_display(inference_time, total_time, stage_execution_times, num_frames):
-    dit_denoising_time = f"{stage_execution_times[5]:.2f}s" if len(stage_execution_times) > 5 else "N/A"
-    
-    timing_html = f"""
-    <div style="margin: 10px 0;">
-        <h3 style="text-align: center; margin-bottom: 10px;">⏱️ Timing Breakdown</h3>
-        <div style="display: grid; grid-template-columns: repeat(5, 1fr); gap: 10px; margin-bottom: 10px;">
-            <div class="timing-card timing-card-highlight">
-                <div style="font-size: 20px;">🚀</div>
-                <div style="font-weight: bold; margin: 3px 0; font-size: 14px;">DiT Denoising</div>
-                <div style="font-size: 16px; color: #ffa200; font-weight: bold;">{dit_denoising_time}</div>
-            </div>
-            <div class="timing-card">
-                <div style="font-size: 20px;">🧠</div>
-                <div style="font-weight: bold; margin: 3px 0; font-size: 14px;">E2E (w. vae/text encoder)</div>
-                <div style="font-size: 16px; color: #2563eb;">{inference_time:.2f}s</div>
-            </div>
-            <div class="timing-card">
-                <div style="font-size: 20px;">🎬</div>
-                <div style="font-weight: bold; margin: 3px 0; font-size: 14px;">Video Encoding</div>
-                <div style="font-size: 16px; color: #dc2626;">N/A</div>
-            </div>
-            <div class="timing-card">
-                <div style="font-size: 20px;">🌐</div>
-                <div style="font-weight: bold; margin: 3px 0; font-size: 14px;">Network Transfer</div>
-                <div style="font-size: 16px; color: #059669;">N/A</div>
-            </div>
-            <div class="timing-card">
-                <div style="font-size: 20px;">📊</div>
-                <div style="font-weight: bold; margin: 3px 0; font-size: 14px;">Total Processing</div>
-                <div style="font-size: 18px; color: #0277bd;">{total_time:.2f}s</div>
-            </div>
-        </div>"""
-    
-    if inference_time > 0:
-        fps = num_frames / inference_time
-        timing_html += f"""
-        <div class="performance-card" style="margin-top: 15px;">
-            <span style="font-weight: bold;">Generation Speed: </span>
-            <span style="font-size: 18px; color: #6366f1; font-weight: bold;">{fps:.1f} frames/second</span>
-        </div>"""
-    
-    return timing_html + "</div>"
-
-def get_action_tensors(mode: str, keyboard_key: str, mouse_key: str | None):
-    if mode == "universal":
-        keyboard = torch.tensor(KEYBOARD_MAP_UNIVERSAL.get(keyboard_key, [0, 0, 0, 0])).cuda()
-        mouse = torch.tensor(CAMERA_MAP_UNIVERSAL.get(mouse_key, [0, 0])).cuda()
-    elif mode == "gta_drive":
-        keyboard = torch.tensor(KEYBOARD_MAP_GTA.get(keyboard_key, [0, 0])).cuda()
-        mouse = torch.tensor(CAMERA_MAP_GTA.get(mouse_key, [0, 0])).cuda()
-    elif mode == "templerun":
-        keyboard = torch.tensor(KEYBOARD_MAP_TEMPLERUN.get(keyboard_key, [1, 0, 0, 0, 0, 0, 0])).cuda()
-        mouse = None
-    else:
-        raise ValueError(f"Unknown mode: {mode}")
-    
-    return {"keyboard": keyboard, "mouse": mouse}
-
-def create_gradio_interface(generators: dict[str, StreamingVideoGenerator], loaded_model_name: str):
-    initial_config = VARIANT_CONFIG.get(loaded_model_name, VARIANT_CONFIG["Matrix-Game-2.0-Base"])
-    initial_mode = initial_config["mode"]
-    
-    if initial_mode == "universal":
-        initial_kb_choices = list(KEYBOARD_MAP_UNIVERSAL.keys())
-        initial_mouse_choices = list(CAMERA_MAP_UNIVERSAL.keys())
-        initial_mouse_visible = True
-    elif initial_mode == "gta_drive":
-        initial_kb_choices = list(KEYBOARD_MAP_GTA.keys())
-        initial_mouse_choices = list(CAMERA_MAP_GTA.keys())
-        initial_mouse_visible = True
-    else:  # templerun
-        initial_kb_choices = list(KEYBOARD_MAP_TEMPLERUN.keys())
-        initial_mouse_choices = []
-        initial_mouse_visible = False
-    
-    theme = gr.themes.Base().set(
-        button_primary_background_fill="#2563eb",
-        button_primary_background_fill_hover="#1d4ed8",
-        button_primary_text_color="white",
-        slider_color="#2563eb",
-        checkbox_background_color_selected="#2563eb",
-    )
-    
-    with gr.Blocks(title="FastVideo - Matrix Game 2.0", theme=theme) as demo:
-        game_state = gr.State({
-            "initialized": False,
-            "current_model": None,
-            "block_idx": 0,
-            "max_blocks": 50,
-        })
-        
-        # Header
-        gr.Image("assets/full.svg", show_label=False, container=False, height=80)
-        
-        gr.HTML("""
-        <div style="text-align: center; margin-bottom: 10px;">
-            <p style="font-size: 18px;"> Make Video Generation Go Blurrrrrrr </p>
-            <p style="font-size: 18px;"> <a href="https://github.com/hao-ai-lab/FastVideo/tree/main" target="_blank">Code</a> | <a href="https://hao-ai-lab.github.io/blogs/fastvideo_post_training/" target="_blank">Blog</a> | <a href="https://hao-ai-lab.github.io/FastVideo/" target="_blank">Docs</a>  </p>
-        </div>
-        """)
-        
-        with gr.Accordion("🎥 What Is FastVideo?", open=False):
-            gr.HTML("""
-            <div style="padding: 20px; line-height: 1.6;">
-                <p style="font-size: 16px; margin-bottom: 15px;">
-                    FastVideo is an inference and post-training framework for diffusion models. It features an end-to-end unified pipeline for accelerating diffusion models, starting from data preprocessing to model training, finetuning, distillation, and inference. FastVideo is designed to be modular and extensible, allowing users to easily add new optimizations and techniques. Whether it is training-free optimizations or post-training optimizations, FastVideo has you covered.
-                </p>
-            </div>
-            """)
-        
-        # Model Selection
-        with gr.Row():
-            model_selection = gr.Dropdown(
-                choices=[loaded_model_name],
-                value=loaded_model_name,
-                label="Select Model",
-                interactive=False
-            )
-        
-
-        # Main Layout
-        with gr.Row(equal_height=True, elem_classes="main-content-row"):
-            with gr.Column(scale=1, elem_classes="advanced-options-column"):
-                with gr.Group():
-                    gr.HTML("<div style='margin: 0 0 15px 0; text-align: center; font-size: 16px;'>Game Controls</div>")
-                    
-                    with gr.Group():
-                        gr.HTML("<div style='font-size: 14px; margin-bottom: 5px; font-weight: bold;'>🎮 Keyboard Control</div>")
-                        keyboard_action = gr.Radio(
-                            choices=initial_kb_choices,
-                            value=initial_kb_choices[0] if initial_kb_choices else None,
-                            label="Movement",
-                            show_label=False,
-                            interactive=True
-                        )
-                    
-                    with gr.Group(visible=initial_mouse_visible) as mouse_group:
-                        gr.HTML("<div style='font-size: 14px; margin-bottom: 5px; font-weight: bold;'>🖱️ Mouse/Camera Control</div>")
-                        mouse_action = gr.Radio(
-                            choices=initial_mouse_choices if initial_mouse_visible else [],
-                            value=initial_mouse_choices[0] if initial_mouse_choices else None,
-                            label="Camera",
-                            show_label=False,
-                            interactive=True
-                        )
-                    
-                    with gr.Row():
-                        action_btn = gr.Button("Start", variant="primary")
-                        stop_btn = gr.Button("Stop", variant="stop")
-                    
-                    gr.HTML("<div style='margin-top: 15px;'></div>")
-                    
-                    seed = gr.Slider(
-                        label="Seed",
-                        minimum=0,
-                        maximum=1000000,
-                        step=1,
-                        value=1024,
-                    )
-                    randomize_seed = gr.Checkbox(label="Randomize seed", value=False)
-                    seed_output = gr.Number(label="Used Seed")
-                    
-                    block_counter = gr.Textbox(label="Progress", value="Block: 0 / 50", interactive=False, lines=1)
-
-
-            # Right Column: Video Output
-            with gr.Column(scale=1, elem_classes="video-column"):
-                video_output = gr.Video(
-                    label="Generated Video",
-                    show_label=True,
-                    height=466,
-                    width=600,
-                    container=True,
-                    elem_classes="video-component",
-                    autoplay=True
-                )
-
-        # Styles
-        gr.HTML("""
-        <style>
-        .center-button {
-            display: flex !important;
-            justify-content: center !important;
-            height: 100% !important;
-            padding-top: 1.4em !important;
-        }
-        
-        .gradio-container {
-            max-width: 1200px !important;
-            margin: 0 auto !important;
-        }
-        
-        .main {
-            max-width: 1200px !important;
-            margin: 0 auto !important;
-        }
-        
-        .gr-form, .gr-box, .gr-group {
-            max-width: 1200px !important;
-        }
-        
-        .gr-video {
-            max-width: 500px !important;
-            margin: 0 auto !important;
-        }
-        
-        .main-content-row {
-            display: flex !important;
-            align-items: flex-start !important;
-            min-height: 500px !important;
-            gap: 20px !important;
-        }
-        
-        .advanced-options-column,
-        .video-column {
-            display: flex !important;
-            flex-direction: column !important;
-            flex: 1 !important;
-            min-height: 400px !important;
-            align-items: stretch !important;
-        }
-        
-        .video-column > * {
-            margin-top: 0 !important;
-        }
-        
-        .video-column .gr-video,
-        .video-component {
-            margin-top: 0 !important;
-            padding-top: 0 !important;
-        }
-        
-        .video-column .gr-video .gr-form {
-            margin-top: 0 !important;
-        }
-        
-        .advanced-options-column .gr-group,
-        .video-column .gr-video {
-            margin-top: 0 !important;
-            vertical-align: top !important;
-        }
-        
-        .advanced-options-column > *:last-child,
-        .video-column > *:last-child {
-            flex-grow: 0 !important;
-        }
-        
-        @media (max-width: 1400px) {
-            .main-content-row {
-                min-height: 600px !important;
-            }
-            
-            .advanced-options-column,
-            .video-column {
-                min-height: 600px !important;
-            }
-        }
-        
-        @media (max-width: 1200px) {
-            .main-content-row {
-                flex-direction: column !important;
-                align-items: stretch !important;
-            }
-            
-            .advanced-options-column,
-            .video-column {
-                min-height: auto !important;
-                width: 100% !important;
-            }
-        }
-        
-        .timing-card {
-            background: var(--background-fill-secondary) !important;
-            border: 1px solid var(--border-color-primary) !important;
-            color: var(--body-text-color) !important;
-            padding: 10px;
-            border-radius: 8px;
-            text-align: center;
-            min-height: 80px;
-            display: flex;
-            flex-direction: column;
-            justify-content: center;
-        }
-        
-        .timing-card-highlight {
-            background: var(--background-fill-primary) !important;
-            border: 2px solid var(--color-accent) !important;
-        }
-        
-        .performance-card {
-            background: var(--background-fill-secondary) !important;
-            border: 1px solid var(--border-color-primary) !important;
-            color: var(--body-text-color) !important;
-            padding: 10px;
-            border-radius: 6px;
-            text-align: center;
-        }
-
-        .gr-number input[readonly] {
-            background-color: var(--background-fill-secondary) !important;
-            border: 1px solid var(--border-color-primary) !important;
-            color: var(--body-text-color-subdued) !important;
-            cursor: default !important;
-            text-align: center !important;
-            font-weight: 500 !important;
-        }
-        </style>
-        """)
-        
-        # UI update based on model selection
-        def on_model_change(model_name):
-            config = VARIANT_CONFIG.get(model_name, VARIANT_CONFIG["Matrix-Game-2.0-Base"])
-            mode = config["mode"]
-            
-            if mode == "universal":
-                kb_choices = list(KEYBOARD_MAP_UNIVERSAL.keys())
-                mouse_choices = list(CAMERA_MAP_UNIVERSAL.keys())
-                mouse_visible = True
-            elif mode == "gta_drive":
-                kb_choices = list(KEYBOARD_MAP_GTA.keys())
-                mouse_choices = list(CAMERA_MAP_GTA.keys())
-                mouse_visible = True
-            else:  # templerun
-                kb_choices = list(KEYBOARD_MAP_TEMPLERUN.keys())
-                mouse_choices = []
-                mouse_visible = False
-            
-            return (
-                gr.update(choices=kb_choices, value=kb_choices[0] if kb_choices else None),
-                gr.update(choices=mouse_choices, value=mouse_choices[0] if mouse_choices else None, visible=mouse_visible),
-                gr.update(visible=mouse_visible),
-            )
-        
-        model_selection.change(
-            fn=on_model_change,
-            inputs=model_selection,
-            outputs=[keyboard_action, mouse_action, mouse_group]
-        )
-
-        def start_game(model_name, seed_val, randomize, state):
-            if randomize:
-                seed_val = torch.randint(0, 1000000, (1,)).item()
-            
-            config = VARIANT_CONFIG.get(model_name)
-            if not config:
-                return state, seed_val, "Block: 0 / 50", None, "", gr.update(), gr.update()
-            
-            generator = generators.get(config["model_path"])
-            if not generator:
-                return state, seed_val, "Block: 0 / 50", None, "", gr.update(), gr.update()
-            
-            # If already initialized, clean up first
-            if state.get("initialized"):
-                try:
-                    # Clear accumulated frames without saving
-                    generator.accumulated_frames = []
-                    generator.executor.execute_streaming_clear()
-                except Exception as e:
-                    print(f"Warning: cleanup error: {e}")
-            
-            # Streaming parameters
-            num_latent_frames_per_block = 3
-            max_blocks = 50
-            total_latent_frames = num_latent_frames_per_block * max_blocks
-            num_frames = (total_latent_frames - 1) * 4 + 1
-            
-            actions = {
-                "keyboard": torch.zeros((num_frames, config["keyboard_dim"])),
-                "mouse": torch.zeros((num_frames, 2))
-            }
-            grid_sizes = torch.tensor([150, 44, 80])
-            
-            output_dir = os.path.abspath("outputs/matrixgame")
-            os.makedirs(output_dir, exist_ok=True)
-            video_path = os.path.join(output_dir, f"video_{int(time.time())}.mp4")
-            
-            generator.reset(
-                prompt="",
-                image_path=config["image_url"],
-                mouse_cond=actions["mouse"].unsqueeze(0),
-                keyboard_cond=actions["keyboard"].unsqueeze(0),
-                grid_sizes=grid_sizes,
-                num_frames=num_frames,
-                height=352,
-                width=640,
-                num_inference_steps=50,
-                output_path=video_path,
-            )
-            
-            new_state = {
-                "initialized": True,
-                "current_model": model_name,
-                "block_idx": 0,
-                "max_blocks": max_blocks,
-                "video_path": video_path,
-                "frames_per_block": num_latent_frames_per_block * 4,
-                "mode": config["mode"],
-                "seed": seed_val,
-            }
-            
-            return new_state, seed_val, "Block: 0 / 50", None, gr.update(value="Step"), gr.update(interactive=True)
-        
-        async def step_game(keyboard_key, mouse_key, model_name, state):
-            if not state.get("initialized"):
-                return state, state.get("seed", 0), "Block: 0 / 50", None, gr.update(), gr.update()
-            
-            # total_start_time = time.time()
-            config = VARIANT_CONFIG.get(model_name)
-            generator = generators.get(config["model_path"])
-            mode = state["mode"]
-            frames_per_block = state["frames_per_block"]
-
-            # Parse inputs to tensors
-            action = get_action_tensors(mode, keyboard_key, mouse_key)
-            keyboard_cond, mouse_cond = expand_action_to_frames(action, frames_per_block)
-            
-            # run step async
-            # inference_start_time = time.time()
-            frames, block_future = await generator.step_async(keyboard_cond, mouse_cond)
-            # inference_time = time.time() - inference_start_time
-            
-            # wait for block file to be written
-            block_path = await asyncio.to_thread(block_future.result) if block_future else None
-            state["block_idx"] = generator.block_idx
-            block_str = f"Block: {state['block_idx']} / {state['max_blocks']}"
-            
-            # total_time = time.time() - total_start_time
-            
-            # Timing breakdown
-            # timing_html = create_timing_display(inference_time, total_time, [], frames_per_block)
-
-            return state, state.get("seed", 0), block_str, block_path, gr.update(), gr.update()
-
-        def stop_game(model_name, state):
-            if not state.get("initialized"):
-                return {"initialized": False}, 0, "Block: 0 / 50", None, gr.update(value="Start"), gr.update(interactive=False)
-            
-            config = VARIANT_CONFIG.get(model_name)
-            generator = generators.get(config["model_path"])
-            
-            final_path = state.get("video_path")
-            generator.finalize(final_path)
-            
-            return {"initialized": False}, state.get("seed", 0), "Block: 0 / 50", final_path, gr.update(value="Start"), gr.update(interactive=False)
-        
-        async def handle_action(keyboard_key, mouse_key, model_name, seed_val, randomize, state):
-            if not state.get("initialized"):
-                return start_game(model_name, seed_val, randomize, state)
-            else:
-                return await step_game(keyboard_key, mouse_key, model_name, state)
-        
-        action_btn.click(
-            fn=handle_action,
-            inputs=[keyboard_action, mouse_action, model_selection, seed, randomize_seed, game_state],
-            outputs=[game_state, seed_output, block_counter, video_output, action_btn, stop_btn]
-        )
-        
-        stop_btn.click(
-            fn=stop_game,
-            inputs=[model_selection, game_state],
-            outputs=[game_state, seed_output, block_counter, video_output, action_btn, stop_btn]
-        )
-        
-        gr.HTML("""
-        <div style="text-align: center; margin-top: 10px; margin-bottom: 15px;">
-            <p style="font-size: 16px; margin: 0;">Note that this demo is meant to showcase Matrix Game's quality and that under a large number of requests, generation speed may be affected.</p>
-        </div>
-        """)
-    
-    return demo
-
-
-def main():
-    parser = argparse.ArgumentParser(description="Matrix Game Gradio Demo")
-    parser.add_argument("--model", type=str, default="Matrix-Game-2.0-Base",
-                        choices=list(VARIANT_CONFIG.keys()),
-                        help="Model variant to load")
-    parser.add_argument("--host", type=str, default="0.0.0.0")
-    parser.add_argument("--port", type=int, default=7860)
-    args = parser.parse_args()
-    
-    # Load the selected model
-    config = VARIANT_CONFIG[args.model]
-    model_path = config["model_path"]
-    
-    print(f"Loading model: {model_path}")
-    setup_model_environment(model_path)
-    generator = StreamingVideoGenerator.from_pretrained(
-        model_path,
-        num_gpus=1,
-        use_fsdp_inference=True,
-        dit_cpu_offload=True,
-        vae_cpu_offload=False,
-        text_encoder_cpu_offload=True,
-        pin_cpu_memory=True,
-    )
-    
-    generators = {model_path: generator}
-    
-    demo = create_gradio_interface(generators, args.model)
-    
-    print(f"Starting Gradio at http://{args.host}:{args.port}")
-    
-    # FastAPI Wrapper
-    app = FastAPI()
-    
-    @app.get("/logo.png")
-    def get_logo():
-        return FileResponse(
-            "assets/full.svg",
-            media_type="image/svg+xml",
-            headers={
-                "Cache-Control": "public, max-age=3600",
-                "Access-Control-Allow-Origin": "*"
-            }
-        )
-
-    @app.get("/favicon.ico")
-    def get_favicon():
-        favicon_path = "assets/icon-simple.svg"
-        
-        if os.path.exists(favicon_path):
-            return FileResponse(
-                favicon_path, 
-                media_type="image/svg+xml",
-                headers={
-                    "Cache-Control": "public, max-age=3600",
-                    "Access-Control-Allow-Origin": "*"
-                }
-            )
-        else:
-            raise HTTPException(status_code=404, detail="Favicon not found")
-    
-    @app.get("/", response_class=HTMLResponse)
-    def index(request: Request):
-        base_url = str(request.base_url).rstrip('/')
-        return f"""
-        <!DOCTYPE html>
-        <html lang="en">
-        <head>
-            <meta charset="UTF-8" />
-            <meta name="viewport" content="width=device-width, initial-scale=1.0" />
-            
-            <title>FastVideo - Matrix Game 2.0</title>
-            <meta name="title" content="MatrixGame2.0">
-            <meta name="description" content="Make video generation go blurrrrrrr">
-            <meta name="keywords" content="FastVideo, video generation, AI, machine learning, Matrix Game 2.0">
-            
-            <meta property="og:type" content="website">
-            <meta property="og:url" content="{base_url}/">
-            <meta property="og:title" content="FastVideo - Matrix Game 2.0">
-            <meta property="og:description" content="Make video generation go blurrrrrrr">
-            <meta property="og:image" content="{base_url}/logo.png">
-            <meta property="og:image:width" content="1200">
-            <meta property="og:image:height" content="630">
-            <meta property="og:site_name" content="MatrixGame2.0">
-            
-            <meta property="twitter:card" content="summary_large_image">
-            <meta property="twitter:url" content="{base_url}/">
-            <meta property="twitter:title" content="MatrixGame2.0">
-            <meta property="twitter:description" content="Make video generation go blurrrrrrr">
-            <meta property="twitter:image" content="{base_url}/logo.png">
-            <link rel="icon" type="image/png" sizes="32x32" href="/favicon.ico">
-            <link rel="icon" type="image/png" sizes="16x16" href="/favicon.ico">
-            <link rel="apple-touch-icon" href="/favicon.ico">
-            <style>
-                body, html {{
-                    margin: 0;
-                    padding: 0;
-                    height: 100%;
-                    overflow: hidden;
-                }}
-                iframe {{
-                    width: 100%;
-                    height: 100vh;
-                    border: none;
-                }}
-            </style>
-        </head>
-        <body>
-            <iframe src="/gradio" width="100%" height="100%" style="border: none;"></iframe>
-        </body>
-        </html>
-        """
-
-    app = gr.mount_gradio_app(
-        app, 
-        demo, 
-        path="/gradio",
-        allowed_paths=[os.path.abspath("outputs"), os.path.abspath("fastvideo-logos")]
-    )
-    
-    uvicorn.run(app, host=args.host, port=args.port)
-
-
-if __name__ == "__main__":
-    main()
@@ -1,46 +0,0 @@
-from fastvideo import VideoGenerator
-import argparse
-
-OUTPUT_PATH = "video_samples_wan2_2_5B_ti2v"
-
-
-def main(text_encoder_path: str):
-    # 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.
-    model_name = "Wan-AI/Wan2.2-TI2V-5B-Diffusers"
-    generator = VideoGenerator.from_pretrained(
-        model_name,
-        # FastVideo will automatically handle distributed setup
-        num_gpus=1,
-        use_fsdp_inference=True,
-        dit_cpu_offload=True,
-        vae_cpu_offload=False,
-        text_encoder_cpu_offload=False,
-        # AbsMaxFP8 is the quantization method used by ComfyUI; 
-        # check fastvideo/layers/quantization/* for more quantization methods
-        override_text_encoder_quant="AbsMaxFP8",
-        # for Wan 2.2, this is the path to "umt5_xxl_fp8_e4m3fn_scaled.safetensors"
-        override_text_encoder_safetensors=text_encoder_path,
-        pin_cpu_memory=True,  # set to false if low CPU RAM or hit obscure "CUDA error: Invalid argument"
-    )
-
-    # I2V is triggered just by passing in an image_path argument
-    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."
-    image_path = "https://huggingface.co/datasets/YiYiXu/testing-images/resolve/main/wan_i2v_input.JPG"
-    video = generator.generate_video(
-        prompt, output_path=OUTPUT_PATH, save_video=True, image_path=image_path
-    )
-
-
-if __name__ == "__main__":
-    parser = argparse.ArgumentParser()
-    parser.add_argument(
-        "--text_encoder_path",
-        type=str,
-        required=True,
-        help="Path to the quantized text encoder safetensors file.",
-    )
-    args = parser.parse_args()
-    main(args.text_encoder_path)
@@ -1,38 +1,36 @@
 #!/bin/bash

 export WANDB_BASE_URL="https://api.wandb.ai"
-export WANDB_MODE=offline
+export WANDB_MODE=online
 export TOKENIZERS_PARALLELISM=false

-MODEL_PATH="hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v"
-DATA_DIR="data/ode-preprocessing-hy15-test/"
-VALIDATION_DATASET_FILE="data/validation_64.json"
+MODEL_PATH="Wan-AI/Wan2.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
 # IP=[MASTER NODE IP]

 # Training arguments
 training_args=(
  --tracker_project_name "wan_ode_init"
-  --output_dir "ode_init_hy15_test"
-  --wandb_run_name "vidprom_bz128_1e-5"
-  # --resume_from_checkpoint "ode_init_diffusers/"
-  # --warp_denoising_step
-  # --log_visualization
-  --max_train_steps 6001
+  --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
  --train_sp_batch_size 1
  --gradient_accumulation_steps 1
-  --num_latent_t 19
+  --num_latent_t 21
  --num_height 480
  --num_width 832
-  --num_frames 73
-  --dmd_denoising_steps "1000,750,500,250"
+  --num_frames 77
+  --warp_denoising_step
  --enable_gradient_checkpointing_type "full"
 )

 # Parallel arguments
 parallel_args=(
-  --num_gpus 1
+  --num_gpus $NUM_GPUS
  --sp_size 1
  --tp_size 1
  --hsdp_replicate_dim 1
@@ -53,17 +51,18 @@ dataset_args=(

 # Validation arguments
 validation_args=(
-  # --log_validation
+  --log_validation
  --validation_dataset_file "$VALIDATION_DATASET_FILE"
-  --validation_steps 200
+  --validation_steps 50
  --validation_sampling_steps "50"
  --validation_guidance_scale "6.0"
 )

 # Optimizer arguments
 optimizer_args=(
-  --learning_rate 1e-5
+  --learning_rate 6e-6
  --mixed_precision "bf16"
+  --weight_only_checkpointing_steps 1000
  --training_state_checkpointing_steps 1000
  --weight_decay 1e-4
  --max_grad_norm 1.0
@@ -79,7 +78,6 @@ miscellaneous_args=(
  --dit_precision "fp32"
  --num_euler_timesteps 50
  --ema_start_step 0
-  # --enable_gradient_checkpointing_type "full"
 )

 # If you do not have 32 GPUs and to fit in memory, you can: 1. increase sp_size. 2. reduce num_latent_t
@@ -1,132 +0,0 @@
-#!/bin/bash
-#SBATCH --job-name=hy15_ode_3333
-#SBATCH --partition=main
-#SBATCH --nodes=8
-#SBATCH --ntasks=8
-#SBATCH --ntasks-per-node=1
-#SBATCH --gres=gpu:8
-#SBATCH --cpus-per-task=128
-#SBATCH --mem=1440G
-#SBATCH --output=hy15_ode_3333_output/hy15_ode_3333_%j.out
-#SBATCH --error=hy15_ode_3333_output/hy15_ode_3333_%j.err
-#SBATCH --exclusive
-
-# Basic Info
-export NCCL_P2P_DISABLE=1
-export TORCH_NCCL_ENABLE_MONITORING=0
-export NCCL_DEBUG_SUBSYS=INIT,NET
-# different cache dir for different processes
-export TRITON_CACHE_DIR=/tmp/triton_cache_${SLURM_PROCID}
-export MASTER_PORT=29500
-export NODE_RANK=$SLURM_PROCID
-nodes=( $(scontrol show hostnames $SLURM_JOB_NODELIST) )
-export MASTER_ADDR=${nodes[0]}
-export TOKENIZERS_PARALLELISM=false
-# export WANDB_API_KEY="8d9f4b39abd68eb4e29f6fc010b7ee71a2207cde"
-export WANDB_API_KEY="50632ebd88ffd970521cec9ab4a1a2d7e85bfc45"
-# export WANDB_API_KEY='your_wandb_api_key_here'
-export WANDB_BASE_URL="https://api.wandb.ai"
-export WANDB_MODE=online
-export FASTVIDEO_ATTENTION_BACKEND=FLASH_ATTN
-
-source ~/conda/miniconda/bin/activate
-conda activate wei-fv-distill
-export HOME="/mnt/weka/home/hao.zhang/wei"
-
-# Configs
-NUM_GPUS=8
-NUM_TOTAL_GPUS=64
-
-MODEL_PATH="weizhou03/SFHunyuanVideo-1.5-Diffusers-480p_t2v"
-DATA_DIR="/mnt/weka/home/hao.zhang/wl/sharefs/wl/release/fv-ode-preprocessing-16k-hy15-121"
-VALIDATION_DATASET_FILE="/mnt/weka/home/hao.zhang/wl/FastVideo/examples/distill/Wan2.2-TI2V-5B-Diffusers/Data-free/validation_64.json"
-# IP=[MASTER NODE IP]
-
-# Training arguments
-training_args=(
-  --tracker_project_name "hy15_ode_init"
-  --output_dir "/mnt/weka/home/hao.zhang/wei/hy15_ode_init_1333_new"
-  --wandb_run_name "hy15_ode_init_1333_new"
-  --warp_denoising_step
-  --dmd_denoising_steps '1000,760,520,280,0'
-  --log_visualization
-  --visualization_steps 100
-  --max_train_steps 4000
-  --train_batch_size 1
-  --train_sp_batch_size 1
-  --gradient_accumulation_steps 2
-  --num_latent_t 31
-  --num_height 480
-  --num_width 848
-  --num_frames 121
-  --enable_gradient_checkpointing_type "full"
-)
-
-# Parallel arguments
-parallel_args=(
-  --num_gpus $NUM_TOTAL_GPUS # 64
-  --sp_size 1
-  --tp_size 1
-  --hsdp_replicate_dim 1 # 64
-  --hsdp_shard_dim $NUM_TOTAL_GPUS
-)
-
-# Model arguments
-model_args=(
-  --model_path $MODEL_PATH
-  --pretrained_model_name_or_path $MODEL_PATH
-  --text-encoder-cpu-offload
-)
-
-# Dataset arguments
-dataset_args=(
-  --data_path "$DATA_DIR" 
-  --dataloader_num_workers 4
-)
-
-# Optimizer arguments
-optimizer_args=(
-  --learning_rate 1e-5
-  --mixed_precision "bf16"
-  --training_state_checkpointing_steps 400
-  --weight_decay 0.01
-  --max_grad_norm 1.0
-)
-
-validation_args=(
-  # --log_validation
-  --validation_dataset_file "$VALIDATION_DATASET_FILE"
-  --validation_steps 10
-  --validation_sampling_steps "4"
-  --validation_guidance_scale "1.0" # not used for dmd inference
-  # --vae_cpu_offload True
-)
-
-# Miscellaneous arguments
-miscellaneous_args=(
-  --inference_mode False
-  --checkpoints_total_limit 6
-  --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
-  # --enable_gradient_checkpointing_type "full"
-)
-
-# If you do not have 32 GPUs and to fit in memory, you can: 1. increase sp_size. 2. reduce num_latent_t
-srun torchrun \
--nnodes $SLURM_JOB_NUM_NODES \
--nproc_per_node $NUM_GPUS \
--node_rank $SLURM_PROCID \
--rdzv_backend=c10d \
--rdzv_endpoint="$MASTER_ADDR:$MASTER_PORT" \
-    fastvideo/training/ode_causal_pipeline.py \
-    "${parallel_args[@]}" \
-    "${model_args[@]}" \
-    "${dataset_args[@]}" \
-    "${training_args[@]}" \
-    "${optimizer_args[@]}" \
-    "${validation_args[@]}" \
-    "${miscellaneous_args[@]}"
@@ -1,18 +1,5 @@
 cmake_minimum_required(VERSION 3.26 FATAL_ERROR)
-project(fastvideo-kernel LANGUAGES CXX)
-
-# Prefer environment variable (used by CI or pip install git+repo_addr) if CMake var is not explicitly set.
-if(NOT DEFINED GPU_BACKEND AND DEFINED ENV{GPU_BACKEND})
-    set(GPU_BACKEND "$ENV{GPU_BACKEND}")
-endif()
-
-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()
+project(fastvideo-kernel LANGUAGES CXX CUDA)

 # Import common utils if needed, but we keep it simple for now

@@ -33,11 +20,9 @@ find_package(Torch REQUIRED)
 # Include directories
 include_directories(
    ${CMAKE_SOURCE_DIR}/include
-    ${CMAKE_SOURCE_DIR}/include/cutlass/include
    ${CMAKE_SOURCE_DIR}/include/tk/include
    ${CMAKE_SOURCE_DIR}/include/tk/prototype
    ${CMAKE_SOURCE_DIR}/csrc
-    ${CMAKE_SOURCE_DIR}/csrc/turbodiffusion
    ${TORCH_INCLUDE_DIRS}
 )

@@ -74,7 +59,7 @@ else()
    else()
        # Best-effort local detection (works when a CUDA device is visible).
        execute_process(
-            COMMAND "${Python_EXECUTABLE}" -c "import torch; import sys; \nprint('1' if (torch.cuda.is_available() and torch.version.cuda and torch.cuda.get_device_capability()[0] >= 9) else '0')"
+            COMMAND "${Python_EXECUTABLE}" -c "import torch; import sys; \nprint('1' if (torch.cuda.is_available() and torch.cuda.get_device_capability()[0] >= 9) else '0')"
            OUTPUT_VARIABLE _LOCAL_HAS_HOPPER
            OUTPUT_STRIP_TRAILING_WHITESPACE
            ERROR_QUIET
@@ -121,17 +106,11 @@ endif()

 if(BUILD_CXX_KERNELS)
    # Source files
-    set(EXTENSION_SOURCES 
-        csrc/common_extension.cpp
-        csrc/turbodiffusion/gemm/gemm.cu
-        csrc/turbodiffusion/norm/rmsnorm.cu
-        csrc/turbodiffusion/norm/layernorm.cu
-        csrc/turbodiffusion/quant/quant.cu
-    )
+    set(EXTENSION_SOURCES csrc/common_extension.cpp)
    
    # Conditionally add TK kernels
    if(ENABLE_TK_KERNELS)
-        list(APPEND EXTENSION_SOURCES
+        list(APPEND EXTENSION_SOURCES 
            csrc/attention/st_attn_h100.cu
            csrc/attention/block_sparse_h100.cu
        )
@@ -142,7 +121,7 @@ if(BUILD_CXX_KERNELS)
    Python_add_library(fastvideo_kernel_ops MODULE USE_SABI ${SKBUILD_SABI_VERSION} WITH_SOABI
        ${EXTENSION_SOURCES}
    )
-
+    
    # Build compile definitions list
    set(COMPILE_DEFS TORCH_EXTENSION_NAME=fastvideo_kernel_ops)
    if(ENABLE_TK_KERNELS)
@@ -151,34 +130,10 @@ if(BUILD_CXX_KERNELS)

    target_compile_definitions(fastvideo_kernel_ops PRIVATE ${COMPILE_DEFS})

-    target_compile_options(fastvideo_kernel_ops PRIVATE
+    target_compile_options(fastvideo_kernel_ops PRIVATE 
        $<$<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()
@@ -13,20 +13,17 @@ cd fastvideo-kernel
 ./build.sh
 ```

-### Rocm Build
-If you are in a rocm environment without the compilation toolchaine of CUDA.
+### Release Build (Force Enable Kernels)
+If you are building a release wheel or docker image on a machine without a GPU (e.g., CI/CD), you can force-enable the compilation of Hopper-specific ThunderKittens kernels.

 ```bash
 cd fastvideo-kernel
-./build.sh --rocm
+./build.sh --release
 ```
+*Note: The resulting wheel will contain kernels that require an H100 GPU to run, but can be built on any machine with CUDA 12.3+ toolchain.*

 ## Usage

-### Sliding Tile Attention (STA) & Video Sparse Attention (VSA)
-
-For detailed usage, please check the [Attention Documentation](../docs/attention/index.md).
-
 ```python
 from fastvideo_kernel import sliding_tile_attention, video_sparse_attn, moba_attn_varlen

@@ -34,27 +31,12 @@ 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, block_sizes, topk=5)
+out = video_sparse_attn(q, k, v, 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.
-
 ## Requirements

 - **Runtime**:
@@ -67,14 +49,3 @@ This package also includes kernels from [TurboDiffusion](https://github.com/thu-
 ## Acknowledgement

 This package structure and build system are based on [sgl-kernel](https://github.com/sgl-project/sglang/tree/main/sgl-kernel) from the SGLang project.
-
-The implementation of `turbodiffusion` kernels is adapted from [TurboDiffusion](https://github.com/thu-ml/TurboDiffusion). If you use these kernels, please cite:
-
-```bibtex
-@article{zhang2025turbodiffusion,
-  title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
-  author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
-  journal={arXiv preprint arXiv:2512.16093},
-  year={2025}
-}
-```
@@ -1,166 +0,0 @@
-#!/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()
-
-
@@ -15,24 +15,19 @@ git submodule update --init --recursive
 pip install scikit-build-core cmake ninja

 RELEASE=0
-GPU_BACKEND=CUDA
-for arg in "$@"; do
-    case "$arg" in
-        --rocm)
-            GPU_BACKEND=ROCM
-            ;;
-    esac
-done
+if [ "${1:-}" = "--release" ] || [ "${1:-}" = "-r" ]; then
+    RELEASE=1
+fi

-# Force-enable ThunderKittens kernels and compile for Hopper.
-export TORCH_CUDA_ARCH_LIST="9.0a"
-export CMAKE_ARGS="${CMAKE_ARGS:-} -DFASTVIDEO_KERNEL_BUILD_TK=ON -DCMAKE_CUDA_ARCHITECTURES=90a"
-
-export CMAKE_ARGS="${CMAKE_ARGS:-} -DGPU_BACKEND=${GPU_BACKEND}"
+if [ "$RELEASE" -eq 1 ]; then
+    # Force-enable ThunderKittens kernels and compile for Hopper.
+    # Intended for producing release wheels/images on machines without a GPU.
+    export TORCH_CUDA_ARCH_LIST="9.0a"
+    export CMAKE_ARGS="${CMAKE_ARGS:-} -DFASTVIDEO_KERNEL_BUILD_TK=ON -DCMAKE_CUDA_ARCHITECTURES=90a"
+fi

 echo "TORCH_CUDA_ARCH_LIST: ${TORCH_CUDA_ARCH_LIST:-<unset>}"
 echo "CMAKE_ARGS: ${CMAKE_ARGS:-<unset>}"
-echo "GPU_BACKEND: ${GPU_BACKEND:-<unset>}"
 # Build and install
 # Use -v for verbose output
 pip install . -v --no-build-isolation
@@ -639,8 +639,7 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {

    // store kq and vq

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

    warpgroup::store(kg_smem[0], kg_reg);
@@ -661,6 +660,145 @@ 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>
@@ -672,32 +810,23 @@ block_sparse_attention_forward(
    torch::Tensor v, 
    torch::Tensor q2k_block_sparse_index, 
    torch::Tensor q2k_block_sparse_num, 
-    torch::Tensor kv_block_size
+    torch::Tensor 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 q_seq_len  = q.size(2);
-    auto kv_seq_len = k.size(2);
+    auto seq_len  = q.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 = q2k_block_sparse_index.size(2);
-    auto num_kv_blocks = kv_block_size.size(0);
+    auto num_q_blocks = 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 * 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.");
+    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");
    // 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");
@@ -705,9 +834,11 @@ 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(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(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(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");
@@ -733,12 +864,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>(q_seq_len), 
+                                        static_cast<const uint>(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>(q_seq_len), 
+                                        static_cast<const uint>(seq_len), 
                                        static_cast<const uint>(1)}, 
                                        torch::TensorOptions().dtype(torch::kFloat).device(q.device()).memory_format(at::MemoryFormat::Contiguous));
        
@@ -749,110 +880,32 @@ 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) {
-        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()), 
+        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()),
            reinterpret_cast<int32_t*>(q2k_block_sparse_num.data_ptr()),
-            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
+            reinterpret_cast<int32_t*>(block_size.data_ptr()),
+            stream
        );
-
-        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()), 
+    } 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()),
            reinterpret_cast<int32_t*>(q2k_block_sparse_num.data_ptr()),
-            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
+            reinterpret_cast<int32_t*>(block_size.data_ptr()),
+            stream
        );
-
-        fwd_attend_ker<128><<<grid, (128), mem_size, stream>>>(g);
-
-        CHECK_CUDA_ERROR(cudaGetLastError());
-        // cudaStreamSynchronize(stream);
+    } else {
+        TORCH_CHECK(false, "Unsupported head_dim: ", head_dim, ". Only 64 and 128 are supported.");
    }

    return {o, l_vec};
@@ -868,7 +921,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 kv_block_size)
+                   torch::Tensor block_size)
 {
    CHECK_INPUT(q);
    CHECK_INPUT(k);
@@ -877,23 +930,11 @@ 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 q_seq_len  = q.size(2);
-    auto kv_seq_len = k.size(2);
+    auto seq_len  = q.size(2);
    auto head_dim = q.size(3);
    auto max_q_blocks_per_kv = k2q_block_sparse_index.size(3);
-    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))");
+    TORCH_CHECK(k2q_block_sparse_index.size(2) == block_size.size(0), "k2q_block_sparse_index.size(2) must match 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");
@@ -904,18 +945,23 @@ 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(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(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(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);
@@ -942,20 +988,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>(q_seq_len), 
+                                     static_cast<const uint>(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>(kv_seq_len), 
+                                     static_cast<const uint>(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>(kv_seq_len), 
+                                     static_cast<const uint>(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>(q_seq_len), 
+                                        static_cast<const uint>(seq_len), 
                                        static_cast<const uint>(1)},       l_vec.options());

    float*         qg_ptr = qg.data_ptr<float>();
@@ -984,7 +1030,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(q_seq_len/(PREP_NUM_WARPS*kittens::TILE_ROW_DIM<bf16>*4), qo_heads, batch);
+    dim3 grid_bwd(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>;
@@ -997,9 +1043,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>(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)};
+        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)};

        bwd_prep_globals bwd_g{prep_og_arg, prep_o_arg, prep_d_arg};

@@ -1036,15 +1082,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>(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_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_global_args bwd_global{bwd_q_arg, 
                        bwd_k_arg, 
@@ -1055,14 +1101,14 @@ block_sparse_attention_backward(torch::Tensor q,
                        bwd_vg_arg, 
                        bwd_l_arg, 
                        bwd_d_arg, 
-                        static_cast<int>(kv_seq_len), // N is not used in the kernel
+                        static_cast<int>(seq_len), 
                        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*>(kv_block_size.data_ptr())};
+                        reinterpret_cast<int32_t*>(block_size.data_ptr())};

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

        //cudadevicesynchronize();
@@ -1101,9 +1147,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>(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)};
+        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)};

        bwd_prep_globals bwd_g{prep_og_arg, prep_o_arg, prep_d_arg};

@@ -1140,15 +1186,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>(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_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_global_args bwd_global{bwd_q_arg, 
                        bwd_k_arg, 
@@ -1159,14 +1205,14 @@ block_sparse_attention_backward(torch::Tensor q,
                        bwd_vg_arg, 
                        bwd_l_arg, 
                        bwd_d_arg, 
-                        static_cast<int>(kv_seq_len), // N is not used in the kernel
+                        static_cast<int>(seq_len), 
                        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*>(kv_block_size.data_ptr())};
+                        reinterpret_cast<int32_t*>(block_size.data_ptr())};

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

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

    return {qg, kg, vg};
    //cudadevicesynchronize();
-}
+}
@@ -22,12 +22,6 @@ extern std::vector<torch::Tensor> block_sparse_attention_backward(
 );
 #endif

-// TurboDiffusion kernels
-void register_quant(pybind11::module_ &);
-void register_rms_norm(pybind11::module_ &);
-void register_layer_norm(pybind11::module_ &);
-void register_gemm(pybind11::module_ &);
-
 PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.doc() = "FastVideo CUDA Kernels";

@@ -39,10 +33,4 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("block_sparse_fwd", torch::wrap_pybind_function(block_sparse_attention_forward), "block sparse attention forward (Hopper)");
    m.def("block_sparse_bwd", torch::wrap_pybind_function(block_sparse_attention_backward), "block sparse attention backward (Hopper)");
 #endif
-
-    // TurboDiffusion
-    register_quant(m);
-    register_rms_norm(m);
-    register_layer_norm(m);
-    register_gemm(m);
 }
@@ -1,85 +0,0 @@
-#pragma once
-
-#include <torch/torch.h>
-#include <torch/extension.h>
-
-CUTLASS_HOST_DEVICE int64_t cdiv(int64_t const& a, int64_t const &b) {
-  return (a + b - 1) / b;
-}
-
-template <class T>
-CUTLASS_HOST_DEVICE T max(T a, T b) { return a > b ? a : b; }
-
-template <class T>
-CUTLASS_HOST_DEVICE T min(T a, T b) { return a > b ? b : a; }
-
-#define MIN(a, b) ((a) > (b) ? (b) : (a))
-#define MAX(a, b) ((a) > (b) ? (a) : (b))
-
-#define BOOL_SWITCH(COND, CONST_NAME, ...)    \
-[&] {                                         \
-  if (COND) {                                 \
-    static constexpr bool CONST_NAME = true;  \
-    return (__VA_ARGS__)();                   \
-  } else {                                    \
-    static constexpr bool CONST_NAME = false; \
-    return (__VA_ARGS__)();                   \
-  }                                           \
-}() 
-
-#define CUDA_CHECK(call)    \
-{                           \
-  cudaError_t err = call;   \
-  if (err != cudaSuccess) { \
-    fprintf(stderr, "CUDA Error at %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
-    exit(err);              \
-  }                         \
-}
-
-#define CONFIG_SWITCH(N, ...)                                   \
-[&] {                                                           \
-  if (N <= 1024) {                                              \
-    constexpr int NUM_THR_PER_CTA = 128;                        \
-    constexpr int MAX_HIDDEN_SIZE = 1024;                       \
-    return (__VA_ARGS__)();                                     \
-  } else if (N <= 2048) {                                       \
-    constexpr int NUM_THR_PER_CTA = 128;                        \
-    constexpr int MAX_HIDDEN_SIZE = 2048;                       \
-    return (__VA_ARGS__)();                                     \
-  } else if (N <= 4096) {                                       \
-    constexpr int NUM_THR_PER_CTA = 128;                        \
-    constexpr int MAX_HIDDEN_SIZE = 4096;                       \
-    return (__VA_ARGS__)();                                     \
-  } else if (N <= 8192) {                                       \
-    constexpr int NUM_THR_PER_CTA = 256;                        \
-    constexpr int MAX_HIDDEN_SIZE = 8192;                       \
-    return (__VA_ARGS__)();                                     \
-  } else {                                                      \
-    constexpr int NUM_THR_PER_CTA = 256;                        \
-    constexpr int MAX_HIDDEN_SIZE = 16384;                      \
-    return (__VA_ARGS__)();                                     \
-  }                                                             \
-}()
-
-
-template <int BlockSize>
-  void create_tensor(
-    torch::Device const &device,
-    std::optional<at::Tensor> &output,
-    std::optional<at::Tensor> &scale,
-    int m, int n
-  ) {
-    int num_block_m = cdiv(m, BlockSize);
-    int num_block_n = cdiv(n, BlockSize);
-    if (!output.has_value()) {
-      output.emplace(torch::empty(
-        {m, n},
-        torch::TensorOptions().device(device).dtype(torch::kInt8)
-      ));
-      scale.emplace(torch::empty(
-        {num_block_m, num_block_n},
-        torch::TensorOptions().device(device).dtype(torch::kFloat32)
-      ));
-    }
-  }
-
@@ -1,44 +0,0 @@
-#pragma once
-
-#include <cuda.h>
-#include <cuda_runtime.h>
-
-
-template <class Kernel>
-__global__ void device_kernel(
-  __grid_constant__ typename Kernel::Params const params
-) {
-  extern __shared__ char smem[];
-  Kernel op;
-  op(params, smem);
-}
-
-template <class Kernel>
-__global__ __launch_bounds__(Kernel::MaxThreadsPerBlock, Kernel::MinBlocksPerMultiprocessor)
-void device_kernel_with_launch_bounds(
-  __grid_constant__ typename Kernel::Params const params
-) {
-  extern __shared__ char smem[];
-  Kernel op;
-  op(params, smem);
-}
-
-template <class Kernel>
-void launch_kernel(
-  typename Kernel::Params const &params,
-  dim3 grid_shape,
-  dim3 cta_shape,
-  size_t ShmSize,
-  cudaStream_t stream = nullptr
-) {
-  auto func = device_kernel<Kernel>;
-  if (ShmSize >= 48 * 1024) {
-    CUDA_CHECK(cudaFuncSetAttribute(
-      func,
-      cudaFuncAttributeMaxDynamicSharedMemorySize,
-      ShmSize
-    ));
-  }
-  func<<<grid_shape, cta_shape, ShmSize, stream>>>(params);
-  CUDA_CHECK(cudaGetLastError());
-}
@@ -1,75 +0,0 @@
-#pragma once
-
-template <
-  class InputDtype_,
-  int TileM_,
-  int TileN_,
-  int NumThrPerCta_,
-  bool IsEvenM,
-  bool IsEvenN
->
-class Loader {
-public: 
-  using InputDtype = InputDtype_;
-  static constexpr int TileM = TileM_;
-  static constexpr int TileN = TileN_;
-  static constexpr int NumThrPerCta = NumThrPerCta_;
-  static constexpr int NumElementPerThread = TileM * TileN / NumThrPerCta;
-  static constexpr int NumThrPerRow = TileN / NumElementPerThread;
-
-  static_assert(NumThrPerCta % TileM == 0);
-  static_assert(TileM * TileN % NumThrPerCta == 0);
-
-
-  CUTLASS_DEVICE void
-  load(void const *input_ptr, void *thr_output_reg, int64_t m, int64_t n, int blk_m, int blk_n, int tid) {
-    int n_alignment = (n & 31) * sizeof(InputDtype);
-    int thr_m_offset = tid / NumThrPerRow;
-    int thr_n_offset = (tid % NumThrPerRow) * NumElementPerThread;
-    void const *cta_input_ptr = (void*)((InputDtype*)input_ptr + blk_m * TileM * n + blk_n * TileN);
-    void const *thr_input_ptr = (void*)((InputDtype*)cta_input_ptr + thr_m_offset * n + thr_n_offset);
-    InputDtype tmp_reg[NumElementPerThread];
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      tmp_reg[i] = InputDtype(0.f);
-    bool pred = IsEvenM ? true : thr_m_offset + blk_m * TileM < m;
-    int limit = IsEvenN ? NumElementPerThread : MIN(NumElementPerThread, n - (blk_n * TileN + thr_n_offset));
-    if (n_alignment % 128 == 0)
-      _load<int4, IsEvenN>(thr_input_ptr, (void*)tmp_reg, limit, pred);
-    else if (n_alignment % 64 == 0)
-      _load<int2, IsEvenN>(thr_input_ptr, (void*)tmp_reg, limit, pred);
-    else
-      _load<InputDtype, IsEvenN>(thr_input_ptr, (void*)tmp_reg, limit, pred);
-
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      *((float*)thr_output_reg + i) = static_cast<float>(reinterpret_cast<InputDtype const&>(tmp_reg[i]));
-  }
-
-private:
-  template <class LoadDataType, bool IsEven>
-  CUTLASS_DEVICE void
-  _load(void const *thr_input_ptr, void *thr_output_reg, int limit, bool pred) {
-    static constexpr int NumElementPerLoad = sizeof(LoadDataType) / sizeof(InputDtype);
-    if (pred) {
-      if constexpr (IsEven) {
-        CUTLASS_PRAGMA_UNROLL
-        for (int i = 0; i < NumElementPerThread; i += NumElementPerLoad) {
-          *(LoadDataType*)((InputDtype*)thr_output_reg + i) = *(LoadDataType*)((InputDtype*)thr_input_ptr + i);
-        }
-      } else {
-        CUTLASS_PRAGMA_UNROLL
-        for (int i = 0; i < limit; i += NumElementPerLoad) {
-          if (limit - i > NumElementPerLoad)
-            *(LoadDataType*)((InputDtype*)thr_output_reg + i) = *(LoadDataType*)((InputDtype*)thr_input_ptr + i);
-          else {
-            for (int j = 0; j < NumElementPerLoad; ++j) {
-              if (i + j < limit)
-                *((InputDtype*)thr_output_reg + i + j) = *((InputDtype*)thr_input_ptr + i + j);
-            }
-          }
-        }
-      }
-    }
-  }
-};
@@ -1,77 +0,0 @@
-#pragma once
-#include "common/common.hpp"
-
-
-template <
-  class OutputDtype_,
-  int TileM_,
-  int TileN_,
-  int NumThrPerCta_,
-  bool IsEvenM,
-  bool IsEvenN,
-  bool Round = true,
-  bool SaveScale = true
->
-class Saver {
-public:
-  using OutputDtype = OutputDtype_;
-
-  static constexpr int TileM = TileM_;
-  static constexpr int TileN = TileN_;
-  static constexpr int NumThrPerCta = NumThrPerCta_;
-  static constexpr int NumElementPerThread = TileM * TileN / NumThrPerCta;
-  static constexpr int NumThrPerRow = TileN / NumElementPerThread;
-
-  static_assert(TileM * TileN % NumThrPerCta == 0);
-  static_assert(NumThrPerCta % TileM == 0);
-
-  CUTLASS_DEVICE void
-  store(void *Optr, void *OSptr, void *reg, float scale_inv, int64_t m, int64_t n, int blk_m, int blk_n, int tid) {
-    int n_alignment = (n & 31) * sizeof(OutputDtype);
-    int thr_m_offset = tid / NumThrPerRow;
-    int thr_n_offset = (tid % NumThrPerRow) * NumElementPerThread;
-    void *cta_output_ptr = (void*)((OutputDtype*)Optr + blk_m * TileM * (Round ? cdiv(n, TileN) * TileN : n) + blk_n * TileN);
-    void *thr_output_ptr = (void*)((OutputDtype*)cta_output_ptr + thr_m_offset * (Round ? cdiv(n, TileN) * TileN : n) + thr_n_offset);
-    bool pred = IsEvenM ? true : thr_m_offset + blk_m * TileM < m;
-    int limit = IsEvenN ? NumElementPerThread : MIN(NumElementPerThread, n - (blk_n * TileN + thr_n_offset));
-    if (n_alignment % 128 == 0)
-      _store<int4, IsEvenN>(thr_output_ptr, reg, limit, pred);
-    else if (n_alignment % 64 == 0)
-      _store<int2, IsEvenN>(thr_output_ptr, reg, limit, pred);
-    else
-      _store<OutputDtype, IsEvenN>(thr_output_ptr, reg, limit, pred);
-
-    if constexpr (SaveScale) {
-      if (tid == 0) {
-        *((float*)OSptr + blk_m * cdiv(n, TileN)+ blk_n) = scale_inv;
-      }
-    }
-  }
-
-private:
-  template <class StoreDataType, bool IsEven>
-  CUTLASS_DEVICE void
-  _store(void *thr_output_ptr, void *reg, int limit, bool pred) {
-    static constexpr int NumElementPerStore = sizeof(StoreDataType) / sizeof(OutputDtype);
-    if (pred) {
-      if constexpr (IsEven) {
-        CUTLASS_PRAGMA_UNROLL
-        for (int i = 0; i < NumElementPerThread; i += NumElementPerStore) {
-          *(StoreDataType*)((OutputDtype*)thr_output_ptr + i) = *(StoreDataType*)((OutputDtype*)reg + i);
-        }
-      } else {
-        CUTLASS_PRAGMA_UNROLL
-        for (int i = 0; i < limit; i += NumElementPerStore) {
-          if (limit - i > NumElementPerStore)
-            *(StoreDataType*)((OutputDtype*)thr_output_ptr + i) = *(StoreDataType*)((OutputDtype*)reg + i);
-          else {
-            for (int j = 0; j < limit - i; ++j) {
-              *((OutputDtype*)thr_output_ptr + i + j) = *((OutputDtype*)reg + i + j);
-            }
-          }
-        }
-      }
-    }
-  }
-
-};
@@ -1,77 +0,0 @@
-/* 
- * Copyright (c) 2025 by TurboDiffusion team.
- * 
- * Licensed under the Apache License, Version 2.0 (the "License");
- * 
- * Citation (please cite if you use this code):
- * 
- * @article{zhang2025turbodiffusion,
- *   title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
- *   author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
- *   journal={arXiv preprint arXiv:2512.16093},
- *   year={2025}
- * }
- */
-
-#include <pybind11/pybind11.h>
-#include <torch/extension.h>
-#include <ATen/cuda/CUDAContext.h>
-#include <torch/all.h>
-#include <torch/python.h>
-#include <cutlass/cutlass.h>
-#include <cutlass/numeric_types.h>
-
-#include "common/common.hpp"
-#include "gemm/launch.hpp"
-
-void int8_gemm(
-  at::Tensor const& A, at::Tensor const& A_S, 
-  at::Tensor const& B, at::Tensor const& B_S, 
-  torch::Tensor& C
-) {
-
-  
-  static constexpr int swizzle_dir = 1;
-  static constexpr int swizzle_size_log = 5;
-
-  int k = B.size(1);
-  int m = A.size(0);
-  int n = B.size(0);
-
-  switch (C.scalar_type()) {
-    case torch::kHalf:{
-        int8_gemm_<cutlass::half_t> (
-            (int8_t*)A.data_ptr(), A_S.data_ptr<float>(),
-            (int8_t*)B.data_ptr(), B_S.data_ptr<float>(), 
-            (cutlass::half_t*)C.data_ptr(), 
-            m, n, k, swizzle_dir, swizzle_size_log, at::cuda::getCurrentCUDAStream().stream()
-        );
-        break;
-    }
-
-    case torch::kBFloat16:{
-        int8_gemm_<cutlass::bfloat16_t> (
-                (int8_t*)A.data_ptr(), A_S.data_ptr<float>(),
-                (int8_t*)B.data_ptr(), B_S.data_ptr<float>(), 
-                (cutlass::bfloat16_t*)C.data_ptr(), 
-                m, n, k, swizzle_dir, swizzle_size_log, at::cuda::getCurrentCUDAStream().stream()
-            );
-        break;
-    }
-
-    default: {
-      std::cerr << "Observing: " << C.scalar_type() << " for the output datatype which is invalid";
-      throw std::runtime_error("Unsupported output data type for int8 gemm.");
-    }
-  }
-
-}
-
-void register_gemm(pybind11::module_ &m) {
-    m.def("gemm_cuda", &int8_gemm);
-}
-
-
-  
-
-
@@ -1,522 +0,0 @@
-/* 
- * Copyright (c) 2025 by TurboDiffusion team.
- * 
- * Licensed under the Apache License, Version 2.0 (the "License");
- * 
- * Citation (please cite if you use this code):
- * 
- * @article{zhang2025turbodiffusion,
- *   title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
- *   author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
- *   journal={arXiv preprint arXiv:2512.16093},
- *   year={2025}
- * }
- */
-
-#pragma once
-
-#include <cuda.h>
-#include "cute/tensor.hpp"
-
-#include "common/common.hpp"
-#include "gemm/utils.hpp"
-
-using namespace cute;
-
-template <
-  class OutputDtype_,
-  bool IsEvenM,
-  bool IsEvenN
->
-struct GemmKernel {
-  using ElementA = int8_t;
-  using ElementB = int8_t;
-  using OutputDtype = OutputDtype_;
-  using AccumulatorDtype = int32_t;
-  static constexpr int BlockSize = 128;
-  static constexpr int TileM = 128;
-  static constexpr int TileN = 128;
-  static constexpr int TileK = 128;
-  static constexpr int Stage = 3;
-  static constexpr int EpiStage = 2;
-  
-  static_assert(
-       BlockSize % TileM == 0
-    && BlockSize % TileN == 0
-    && BlockSize % TileK == 0
-  );
-
-  static constexpr int NumTilePerBlock = BlockSize / TileK;
-
-  using SmemLayoutAtom = decltype(
-    composition(
-      Swizzle<3, 4, 3>{},
-      make_layout(
-        make_shape(Int<8>{}, Int<TileK>{}),
-        make_stride(Int<TileK>{}, Int<1>{})
-      )
-    )
-  );
-
-  using SmemLayoutA = decltype(
-    tile_to_shape(
-      SmemLayoutAtom{},
-      make_shape(Int<TileM>{}, Int<TileK>{}, Int<Stage>{})
-    )
-  );
-
-  using SmemLayoutB = decltype(
-    tile_to_shape(
-      SmemLayoutAtom{},
-      make_shape(Int<TileN>{}, Int<TileK>{}, Int<Stage>{})
-    )
-  );
-
-  using MmaOP = cute::SM80_16x8x32_S32S8S8S32_TN;
-  using TiledMma = decltype(
-    make_tiled_mma(
-      MMA_Atom<MMA_Traits<MmaOP>>{},
-      make_layout(make_shape(
-        _4{}, _2{}, _1{}
-      )),
-      make_tile(Int<64>{}, Int<32>{}, Int<32>{})
-    )
-  );
-
-  using G2SCopyAtomA = Copy_Atom<Copy_Traits<SM80_CP_ASYNC_CACHEGLOBAL<cute::uint128_t>>, ElementA>;
-  using G2SCopyAtomB = Copy_Atom<Copy_Traits<SM80_CP_ASYNC_CACHEGLOBAL<cute::uint128_t>>, ElementB>;
-  using G2STiledCopyA = decltype(
-    make_tiled_copy(
-      G2SCopyAtomA{},
-      make_layout(
-        make_shape(Int<64>{}, Int<4>{}),
-        make_stride(Int<4>{}, Int<1>{})
-      ),
-      make_layout(make_shape(Int<1>{}, Int<16>{}))
-    )
-  );
-  using G2STiledCopyB = decltype(
-    make_tiled_copy(
-      G2SCopyAtomB{},
-      make_layout(
-        make_shape(Int<64>{}, Int<4>{}),
-        make_stride(Int<4>{}, Int<1>{})
-      ),
-      make_layout(make_shape(Int<1>{}, Int<16>{}))
-    )
-  );
-
-  using S2RCopyAtomA = Copy_Atom<Copy_Traits<SM75_U32x4_LDSM_N>, ElementA>;
-  using S2RCopyAtomB = Copy_Atom<Copy_Traits<SM75_U32x4_LDSM_N>, ElementB>;
-  using S2RTiledCopyA = decltype(make_tiled_copy_A(S2RCopyAtomA{}, TiledMma{}));
-  using S2RTiledCopyB = decltype(make_tiled_copy_B(S2RCopyAtomB{}, TiledMma{}));
-
-  // epilogue
-  using SmemLayoutAtomD = decltype(
-    composition(
-      Swizzle<2, 3, 3>{},
-      make_layout(
-        make_shape(Int<32>{}, Int<32>{}),
-        LayoutRight{}
-      )
-    )
-  );
-
-  using SmemLayoutD = decltype(
-    tile_to_shape(
-      SmemLayoutAtomD{},
-      make_shape(Int<64>{}, Int<32>{}, Int<EpiStage>{})
-    )
-  );
-
-  using R2SCopyAtomD = Copy_Atom<UniversalCopy<std::conditional_t<sizeof(OutputDtype) == 4, int32_t, int16_t>>, OutputDtype>;
-  using R2STiledCopyD = decltype(make_tiled_copy_C(R2SCopyAtomD{}, TiledMma{}));
-
-  using S2GCopyAtomD = Copy_Atom<UniversalCopy<uint128_t>, OutputDtype>;
-  using S2GCopyD = decltype(make_tiled_copy(
-    S2GCopyAtomD{},
-    make_layout(Shape<_64, _4>{}),
-    make_layout(Shape<_1, _8>{})
-  ));
-
-  using TileShape = decltype(make_shape(Int<TileM>{}, Int<TileN>{}, Int<TileK>{}));
-
-  struct SharedStorageAB: cute::aligned_struct<128> {
-    array_aligned<typename TiledMma::ValTypeA, cosize_v<SmemLayoutA>, 128> smem_A;
-    array_aligned<typename TiledMma::ValTypeB, cosize_v<SmemLayoutB>, 128> smem_B;
-    array_aligned<float, 1> smem_AS;
-    array_aligned<float, 1> smem_BS;
-    array_aligned<int32_t, 1> smem_AF;
-  };
-
-  struct SharedStorageD: cute::aligned_struct<128> {
-    array_aligned<OutputDtype, cosize_v<SmemLayoutD>> smem_D;
-  };
-
-  union SharedStorage {
-    SharedStorageAB storage_AB;
-    SharedStorageD storage_D;
-  };
-
-
-  struct Params {
-    void const* Aptr;
-    void const* ASptr;
-    void const* Bptr;
-    void const* BSptr;
-    void* Dptr;
-    int64_t const m;
-    int64_t const n;
-    int64_t const k;
-    int const swizzle_dir;
-    int const swizzle_size;
-  };
-
-  using Arguments = Params;
-
-  static constexpr int ThreadNum = size(TiledMma{});
-  static constexpr int ShmSize = sizeof(SharedStorage);
-  static constexpr bool FastInt2Float = false;
-
-  static bool can_implement(int64_t m, int64_t n, int64_t k) {
-    if (k % BlockSize != 0) return false;
-    if ((n * sizeof(OutputDtype)) % 16 != 0)
-      return false;
-    return true;
-  }
-
-  static Params to_underlying_arguments(Arguments const& args) {
-    return args;
-  }
-
-  static dim3 get_grid_size(int64_t m, int64_t n) {
-    return dim3(cdiv(m, TileM) * cdiv(n, TileN));
-  }
-
-  CUTLASS_HOST_DEVICE 
-  static auto get_block_coord(
-    int64_t m_blocks, 
-    int64_t n_blocks, 
-    int const swizzle_dir,
-    int64_t const swizzle_size_log
-  ) {
-    int64_t blk_m;
-    int64_t blk_n;
-
-    if (swizzle_dir == 1)
-      std::swap(m_blocks, n_blocks);
-
-    if (swizzle_size_log == 0) {
-      blk_m = blockIdx.x % m_blocks;
-      blk_n = blockIdx.x / m_blocks;
-    } else {
-      int64_t group_size = n_blocks << swizzle_size_log;
-      int64_t num_groups = m_blocks >> swizzle_size_log;
-      int64_t group_idx = blockIdx.x / group_size;
-      int64_t local_idx = blockIdx.x % group_size;
-      if (group_idx == num_groups) {
-        blk_m = (num_groups << swizzle_size_log) + local_idx % (m_blocks - (num_groups << swizzle_size_log));
-        blk_n = local_idx / (m_blocks - (num_groups << swizzle_size_log));
-      } else {
-        blk_m = (local_idx & ((1LL << swizzle_size_log) - 1)) + (group_idx << swizzle_size_log);
-        blk_n = local_idx >> swizzle_size_log;
-      }
-    }
-
-    if (swizzle_dir == 1)
-      std::swap(blk_m, blk_n);
-
-    return make_coord(blk_m, blk_n);
-  }
-
-  CUTLASS_DEVICE
-  void operator()(
-    Params const& params, char* smem_data
-  ) {
-
-    SharedStorage& shared_storage = *reinterpret_cast<SharedStorage*>(smem_data);
-
-    auto t_idx = threadIdx.x;
-
-    int64_t const m = params.m;
-    int64_t const n = params.n;
-    int64_t const k = params.k;
-    int const swizzle_dir = params.swizzle_dir;
-    int const swizzle_size = params.swizzle_size;
-
-    Tensor A = make_tensor(
-      make_gmem_ptr<ElementA>(params.Aptr),
-      make_shape(m, k),
-      make_stride(k, _1{})
-    );
-    Tensor B = make_tensor(
-      make_gmem_ptr<ElementB>(params.Bptr),
-      make_shape(m, k),
-      make_stride(k, _1{})
-    );
-    Tensor AS = make_tensor(
-      make_gmem_ptr<float>(params.ASptr),
-      make_shape(cdiv(m, BlockSize), cdiv(k, BlockSize)),
-      make_stride(cdiv(k, BlockSize), _1{})
-    );
-    Tensor BS = make_tensor(
-      make_gmem_ptr<float>(params.BSptr),
-      make_shape(cdiv(n, BlockSize), cdiv(k, BlockSize)),
-      make_stride(cdiv(k, BlockSize), _1{})
-    );
-    Tensor D = make_tensor(
-      make_gmem_ptr<OutputDtype>(params.Dptr),
-      make_shape(m, n),
-      LayoutRight{}
-    );
-
-    auto [m_coord, n_coord] = get_block_coord(
-      cdiv(m, size<0>(TileShape{})),
-      cdiv(n, size<1>(TileShape{})),
-      swizzle_dir, swizzle_size
-    );
-
-    int32_t blk_m_coord = m_coord / (BlockSize / TileM);
-    int32_t blk_n_coord = n_coord / (BlockSize / TileN);
-
-    // local tile
-    auto gA = local_tile(A, TileShape{}, make_coord(m_coord, n_coord, _), Step<_1, X, _1>{});
-    auto gB = local_tile(B, TileShape{}, make_coord(m_coord, n_coord, _), Step<X, _1, _1>{});
-    auto gD = local_tile(D, TileShape{}, make_coord(m_coord, n_coord, _), Step<_1, _1, X>{});
-
-    // shared memory
-    Tensor sA = make_tensor(
-      make_smem_ptr<ElementA>(shared_storage.storage_AB.smem_A.data()),
-      SmemLayoutA{}
-    );
-
-    Tensor sB = make_tensor(
-      make_smem_ptr<ElementB>(shared_storage.storage_AB.smem_B.data()),
-      SmemLayoutB{}
-    );
-
-    // register
-    TiledMma tiled_mma;
-    auto thr_mma = tiled_mma.get_slice(t_idx);
-    auto tCrA = thr_mma.partition_fragment_A(gA(_, _, 0));
-    auto tCrB = thr_mma.partition_fragment_B(gB(_, _, 0));
-    auto tDrC = thr_mma.partition_fragment_C(gD); // mma accumulator
-    auto tDrD = make_tensor_like<float>(tDrC); // float accumulator
-
-    if constexpr (FastInt2Float) {
-      CUTLASS_PRAGMA_UNROLL
-      for (int i = 0; i < size(tDrC); ++i)
-        tDrC(i) = 0x4B400000;
-    } else {
-      clear(tDrC);
-    }
-
-    clear(tDrD);
-
-
-    // global to shared copy
-    G2STiledCopyA g2s_tiled_copy_a;
-    auto g2s_thr_copy_a = g2s_tiled_copy_a.get_slice(t_idx);
-    auto tAgA = g2s_thr_copy_a.partition_S(gA);
-    auto tAsA = g2s_thr_copy_a.partition_D(sA);
-    auto cA = make_identity_tensor(make_shape(size<0>(sA), size<1>(sA)));
-    auto tAcA = g2s_thr_copy_a.partition_S(cA);
-    int const m_limit = m - TileM * m_coord;
-    int const n_limit = n - TileN * n_coord;
-
-    G2STiledCopyB g2s_tiled_copy_b;
-    auto g2s_thr_copy_b = g2s_tiled_copy_b.get_slice(t_idx);
-    auto tBgB = g2s_thr_copy_b.partition_S(gB);
-    auto tBsB = g2s_thr_copy_b.partition_D(sB);
-    auto cB = make_identity_tensor(make_shape(size<0>(sB), size<1>(sB)));
-    auto tBcB = g2s_thr_copy_a.partition_S(cB);
-
-
-    // shared to register copy
-    S2RTiledCopyA s2r_tiled_copy_a;
-    auto s2r_thr_copy_a = s2r_tiled_copy_a.get_slice(t_idx);
-    auto tCsA = s2r_thr_copy_a.partition_S(sA);
-    auto tCrA_view = s2r_thr_copy_a.retile_D(tCrA);
-
-    S2RTiledCopyB s2r_tiled_copy_b;
-    auto s2r_thr_copy_b = s2r_tiled_copy_b.get_slice(t_idx);
-    auto tCsB = s2r_thr_copy_b.partition_S(sB);
-    auto tCrB_view = s2r_thr_copy_b.retile_D(tCrB);
-
-    // pipeline status
-    int64_t g2s_a_tile = 0;
-    int64_t g2s_b_tile = 0;
-    int g2s_a_smem = 0;
-    int g2s_b_smem = 0;
-
-    int g2s_tile_in_block = 0;
-    int g2s_block = 0; // b block idx
-
-    int s2r_a_smem = 0;
-    int s2r_b_smem = 0;
-    int s2r_tile_in_block = 0;
-
-    int mma_block_a = 0;
-    int mma_block_b = 0;
-
-    int ntile = k / TileK;
-    // load scale and fallback
-    // we assume all ptrs are 128bit aligned
-    // auto smem_fallback_A = raw_pointer_cast(make_smem_ptr<int32_t>(shared_storage.storage_AB.smem_AF.data()));
-    // auto smem_scale_A = raw_pointer_cast(make_smem_ptr<float>(shared_storage.storage_AB.smem_AS.data()));
-    // auto smem_scale_B = raw_pointer_cast(make_smem_ptr<float>(shared_storage.storage_AB.smem_BS.data()));
-    __syncthreads();
-
-
-    int32_t fallbackA_load = 0;
-    int32_t fallbackA_mma = 0;
-
-    // copy first Stage - 1 tile
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0, _i = min(Stage - 1, ntile); i < _i; ++i) {
-      if (g2s_b_tile < ntile) {
-        g2s_tile_in_block = (g2s_tile_in_block + 1) % NumTilePerBlock;
-        copy_AB<IsEvenM>(g2s_tiled_copy_a, tAgA, tAsA, tAcA, g2s_a_tile, g2s_a_smem, m_limit);
-        copy_AB<IsEvenN>(g2s_tiled_copy_b, tBgB, tBsB, tBcB, g2s_b_tile, g2s_b_smem, n_limit);
-        ++g2s_b_tile;
-        ++g2s_b_smem;
-        ++g2s_block;
-        g2s_a_tile = g2s_block * NumTilePerBlock;
-        ++g2s_a_smem;
-      }
-      cp_async_fence();
-    }
-
-    constexpr int nk = size<2>(tCrA);
-    float scale_a = AS(blk_m_coord, 0);
-    float scale_b = BS(blk_n_coord, 0);
-
-    CUTLASS_PRAGMA_NO_UNROLL
-    for (int64_t mma_b_tile = 0; mma_b_tile < ntile; ++mma_b_tile) {
-      s2r_tile_in_block = (s2r_tile_in_block + 1) % NumTilePerBlock;
-      cp_async_wait<Stage - 2>();
-      __syncthreads();
-
-      // do mma first
-      CUTLASS_PRAGMA_UNROLL
-      for (int ik = 0; ik < nk; ++ik) {
-        cute::copy(s2r_tiled_copy_a, tCsA(_, _, ik, s2r_a_smem),
-                    tCrA_view(_, _, ik));
-        cute::copy(s2r_tiled_copy_b, tCsB(_, _, ik, s2r_b_smem),
-                    tCrB_view(_, _, ik));
-        cute::gemm(tiled_mma, tDrC, tCrA(_, _, ik), tCrB(_, _, ik), tDrC);
-      }
-
-      // a s2r increase anyway
-      s2r_a_smem = (s2r_a_smem + 1) % Stage;
-
-      // get next s2r b tile int64_t
-        // end of a block
-
-        // dequant first
-      dequant<AccumulatorDtype, TileM * TileN / ThreadNum, FastInt2Float>(
-        tDrC.data(), tDrD.data(), scale_a * scale_b
-      );
-
-      s2r_b_smem = (s2r_b_smem + 1) % Stage;
-      // b advance
-      ++mma_block_b;
-      if (mma_block_b < size<1>(BS)) scale_b = BS(blk_n_coord, mma_block_b);
-      mma_block_a = mma_block_b;
-      if (mma_block_a < size<1>(AS)) scale_a = AS(blk_m_coord, mma_block_a);
-
-      // load next stage
-      if (g2s_b_tile < ntile) {
-        g2s_tile_in_block = (g2s_tile_in_block + 1) % NumTilePerBlock;
-        copy_AB<IsEvenM>(g2s_tiled_copy_a, tAgA, tAsA, tAcA, g2s_a_tile, g2s_a_smem, m_limit);
-        copy_AB<IsEvenN>(g2s_tiled_copy_b, tBgB, tBsB, tBcB, g2s_b_tile, g2s_b_smem, n_limit);
-        ++g2s_b_tile;
-        g2s_b_smem = (g2s_b_smem + 1) % Stage;
-        ++g2s_block;
-        g2s_a_tile = g2s_block * NumTilePerBlock;
-        g2s_a_smem = (g2s_a_smem + 1) % Stage;
-      }
-      cp_async_fence();
-    }
-
-
-    // epilogue
-
-    Tensor sD = make_tensor(
-      make_smem_ptr<OutputDtype>(shared_storage.storage_D.smem_D.data()),
-      SmemLayoutD{}
-    );
-
-    R2STiledCopyD r2s_tiled_copy_d;
-    auto r2s_thr_copy_d = r2s_tiled_copy_d.get_slice(t_idx);
-    auto tDrD_r2s = r2s_thr_copy_d.retile_S(tDrD);
-    auto tDsD_r2s = r2s_thr_copy_d.partition_D(sD);
-
-    S2GCopyD s2g_tiled_copy_d;
-    auto s2g_thr_copy_d = s2g_tiled_copy_d.get_slice(t_idx);
-    auto tDsD_s2g = s2g_thr_copy_d.partition_S(sD);
-    auto tDgD_s2g = s2g_thr_copy_d.partition_D(gD);
-    Tensor cD = make_identity_tensor(make_shape(Int<TileM>{}, Int<TileN>{}));
-    auto tDcD_s2g = s2g_thr_copy_d.partition_D(cD);
-
-    auto tDgD_s2gx = group_modes<1, 3>(tDgD_s2g);  // (CPY_, CPY_MN)
-    auto tDrD_r2sx = group_modes<1, 3>(tDrD_r2s);  // (CPY_, CPY_MN)
-    auto tDcD_s2gx = group_modes<1, 3>(tDcD_s2g);
-    
-    int32_t step = size<3>(tDsD_r2s);  // pipe
-    CUTLASS_PRAGMA_UNROLL
-    for (int32_t i = 0; i < size<1>(tDrD_r2sx); i += step) {
-      CUTLASS_PRAGMA_UNROLL
-      for (int32_t j = 0; j < step; ++j) {
-        if constexpr (std::is_same<OutputDtype, float>::value) {
-            cute::copy(r2s_tiled_copy_d, tDrD_r2sx(_, i + j), tDsD_r2s(_, 0, 0, j));
-        } else {
-            auto t = make_tensor_like<OutputDtype>(tDrD_r2sx(_, i + j));
-            cute::copy(tDrD_r2sx(_, i + j), t);
-            cute::copy(r2s_tiled_copy_d, t, tDsD_r2s(_, 0, 0, j));
-        }
-      }
-
-      __syncthreads();
-
-      // shm -> global
-      if constexpr (IsEvenM && IsEvenN) {
-        CUTLASS_PRAGMA_UNROLL
-        for (int32_t j = 0; j < step; ++j)
-          cute::copy(s2g_tiled_copy_d, tDsD_s2g(_, 0, 0, j), tDgD_s2gx(_, i + j));
-      } else if constexpr (IsEvenN) {
-          CUTLASS_PRAGMA_UNROLL
-          for (int32_t j = 0; j < step; ++j) {
-            if (get<0>(tDcD_s2gx(0, i + j)) < m_limit)
-              cute::copy(s2g_tiled_copy_d, tDsD_s2g(_, 0, 0, j), tDgD_s2gx(_, i + j));
-          }
-      } else if constexpr (IsEvenM) {
-          CUTLASS_PRAGMA_UNROLL
-          for (int32_t j = 0; j < step; ++j)
-            if (get<1>(tDcD_s2gx(size<0>(tDsD_s2g) - 1, i + j)) < n_limit) {
-              cute::copy(s2g_tiled_copy_d, tDsD_s2g(_, 0, 0, j), tDgD_s2gx(_, i + j));
-            } else {
-              CUTLASS_PRAGMA_UNROLL
-              for (int k = 0; k < size<0>(tDsD_s2g); ++k)
-                if (get<1>(tDcD_s2gx(k, i + j)) < n_limit)
-                  tDgD_s2gx(k, i + j) = tDsD_s2g(k, 0, 0, j);
-            }
-      } else {
-          CUTLASS_PRAGMA_UNROLL
-          for (int32_t j = 0; j < step; ++j)
-            if (get<0>(tDcD_s2gx(0, i + j)) < m_limit) {
-              if (get<1>(tDcD_s2gx(size<0>(tDsD_s2g) - 1, i + j)) < n_limit) {
-                cute::copy(s2g_tiled_copy_d, tDsD_s2g(_, 0, 0, j), tDgD_s2gx(_, i + j));
-              } else {
-                for (int32_t k = 0; k < size<0>(tDsD_s2g); ++k)
-                  if (get<1>(tDcD_s2gx(k, i + j)) < n_limit)
-                    tDgD_s2gx(k, i + j) = tDsD_s2g(k, 0, 0, j);
-              }
-            }
-      }
-      __syncthreads();
-    }
-  }
-
-};
-  
@@ -1,64 +0,0 @@
-/* 
- * Copyright (c) 2025 by TurboDiffusion team.
- * 
- * Licensed under the Apache License, Version 2.0 (the "License");
- * 
- * Citation (please cite if you use this code):
- * 
- * @article{zhang2025turbodiffusion,
- *   title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
- *   author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
- *   journal={arXiv preprint arXiv:2512.16093},
- *   year={2025}
- * }
- */
-
-#pragma once
-
-#include "common/common.hpp"
-#include "common/launch.hpp"
-#include "gemm/kernel.hpp"
-
-
-template <class OutputDtype>
-bool int8_gemm_(
-  int8_t const *Aptr, float const *ASptr,
-  int8_t const *Bptr, float const *BSptr,
-  OutputDtype* Dptr, int64_t m, int64_t n, int64_t k,
-  int swizzle_dir = 1, int swizzle_size_log = 0,
-  cudaStream_t stream = nullptr
-) {
-  BOOL_SWITCH(m % 128 == 0, IsEvenM, [&] {
-    BOOL_SWITCH(n % 128 == 0, IsEvenN, [&] {
-      using Kernel = GemmKernel<OutputDtype, IsEvenM, IsEvenN>;
-      if (!Kernel::can_implement(m, n, k))
-        return false;
-      using Args = typename Kernel::Arguments;
-      Args args {
-        (void*)Aptr, (void*)ASptr,
-        (void*)Bptr, (void*)BSptr, (void*)Dptr,
-        m, n, k, swizzle_dir,
-        swizzle_size_log
-      };
-
-      auto params = Kernel::to_underlying_arguments(args);
-
-      static constexpr size_t ShmSize = Kernel::ShmSize;
-      dim3 grid_shape = Kernel::get_grid_size(m, n);
-      dim3 block_shape = dim3(Kernel::ThreadNum);
-      auto func = device_kernel<Kernel>;
-      if (ShmSize >= 48 * 1024) {
-        cudaFuncSetAttribute(
-          func,
-          cudaFuncAttributeMaxDynamicSharedMemorySize,
-          ShmSize
-        );
-      }
-      func<<<grid_shape, block_shape, ShmSize, stream>>>(
-        params
-      );
-      return true;
-    });
-  });
-  return true;
-}
@@ -1,129 +0,0 @@
-/* 
- * Copyright (c) 2025 by TurboDiffusion team.
- * 
- * Licensed under the Apache License, Version 2.0 (the "License");
- * 
- * Citation (please cite if you use this code):
- * 
- * @article{zhang2025turbodiffusion,
- *   title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
- *   author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
- *   journal={arXiv preprint arXiv:2512.16093},
- *   year={2025}
- * }
- */
-
-#pragma once
-
-#include <cuda.h>
-#include "cute/tensor.hpp"
-
-template <
-  bool IsEven, 
-  class TiledCopy,
-  class SrcTensor,
-  class DstTensor,
-  class PrdTensor
->
-CUTLASS_DEVICE void
-copy_AB(
-    TiledCopy const& _copy, 
-    SrcTensor const &S, 
-    DstTensor &D, 
-    PrdTensor const &ID, 
-    const int64_t &i_read, 
-    const int64_t &i_write, 
-    const int64_t &limit
-) {
-  using namespace cute;
-  if constexpr (IsEven)
-    cute::copy(_copy, S(_, _, _, i_read), D(_, _, _, i_write));
-  else {
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < size<1>(ID); ++i)
-      if (get<0>(ID(0, i, 0)) < limit)
-        cute::copy(_copy, S(_, i, _, i_read), D(_, i, _, i_write));
-  }
-}
-
-template <int N>
-CUTLASS_DEVICE void copy_async(
-  void const* gmem_src,
-  void* smem_dst
-) {
-  uint32_t smem_int_ptr = static_cast<uint32_t>(__cvta_generic_to_shared(smem_dst));;
-  asm volatile("cp.async.ca.shared.global.L2::128B [%0], [%1], %2;\n"
-      :: "r"(smem_int_ptr),
-          "l"(gmem_src),
-          "n"(N));
-}
-
-template<class LoadType, class T, int NumThreads>
-CUTLASS_DEVICE void copy_aligned(const void* src, void* dst, size_t N, int64_t thread_idx) {
-  static constexpr int NumElementPerLoad = sizeof(LoadType) / sizeof(T);
-  for (int64_t i = thread_idx * NumElementPerLoad; i < N; i += NumElementPerLoad * NumThreads) {
-    if (i + NumElementPerLoad <= N) {
-      copy_async<sizeof(LoadType)>(
-        (void*)((T*)src + i),
-        (void*)((T*)dst + i)
-      );
-    } else {
-      for (int64_t j = 0; j < N - i; ++j)
-        copy_async<sizeof(T)>(
-          (void*)((T*)src + i + j),
-          (void*)((T*)dst + i + j)
-        );
-    }
-  }
-}
-
-template<class T, int NumThreads, bool Wait = true, bool Commit = true>
-CUTLASS_DEVICE void g2s_vector_copy(const void* src, void* dst, size_t N, int64_t thread_idx) {
-  
-  uintptr_t src_addr = reinterpret_cast<uintptr_t>(src);
-
-  if (src_addr % 16 == 0) {
-      copy_aligned<int4, T, NumThreads>(src, dst, N, thread_idx);
-  } else if (src_addr % 8 == 0) {
-      copy_aligned<int2, T, NumThreads>(src, dst, N, thread_idx);
-  } else if (src_addr % 4 == 0) {
-      copy_aligned<int, T, NumThreads>(src, dst, N, thread_idx);
-  } else {
-    assert(0);
-  }
-  if constexpr (Commit) {
-    asm volatile("cp.async.commit_group;\n" ::);
-  }
-  if constexpr (Wait) {
-    asm volatile("cp.async.wait_all;\n" ::);
-  }
-}
-
-template<class T, int N, bool FastInt2Float>
-CUTLASS_DEVICE
-static void dequant(
-  T* mma_accum_ptr,
-  float* float_accum_ptr,
-  float scale
-) {
-  static int const ic = 0x4B400000;
-  if constexpr (FastInt2Float && std::is_same_v<T, int32_t>) {
-    CUTLASS_PRAGMA_UNROLL
-    for (size_t i = 0; i < N; ++i) {
-      *(float_accum_ptr + i) += (__int_as_float(*(mma_accum_ptr + i)) - __int_as_float(ic)) * scale;
-      *(mma_accum_ptr + i) = ic;
-    }
-  } else if constexpr (std::is_same_v<T, int32_t>) {
-    CUTLASS_PRAGMA_UNROLL
-    for (size_t i = 0; i < N; ++i) {
-      *(float_accum_ptr + i) += __int2float_rn(*(mma_accum_ptr + i)) * scale;
-      *(mma_accum_ptr + i) = 0;
-    }
-  } else {
-    CUTLASS_PRAGMA_UNROLL
-    for (size_t i = 0; i < N; ++i) {
-      *(float_accum_ptr + i) += (*(mma_accum_ptr + i)) * scale;
-      *(mma_accum_ptr + i) = 0;
-    }
-  }
-}
@@ -1,99 +0,0 @@
-#include <pybind11/pybind11.h>
-#include <torch/extension.h>
-#include <ATen/cuda/CUDAContext.h>
-#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"
-
-auto layer_norm(
-  at::Tensor const Input, 
-  float eps,
-  std::optional<at::Tensor const> W,
-  std::optional<at::Tensor const> const B,
-  std::optional<at::Tensor> Output
-) {
-  int64_t const m = Input.size(0);
-  int64_t const n = Input.size(1);
-  torch::Device const input_device = Input.device();
-
-  if (!Output.has_value()) {
-    Output.emplace(
-      torch::empty(
-        {m, n},
-        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();
-
-  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());
-  }
-  
-    
-
-  return Output;
-}
-
-void register_layer_norm(pybind11::module_ &m) {
-    m.def("layer_norm_cuda", &layer_norm);
-}
-
@@ -1,215 +0,0 @@
-#pragma once
-
-#include "common/load.hpp"
-#include "common/store.hpp"
-#include "common/launch.hpp"
-
-
-template <
-  class InputDtype_,
-  class OutputDtype_,
-  class WeightDtype_,
-  bool Affine_,
-  bool Bias_,
-  int MaxHiddenSize_,
-  int NumThrPerCta_,
-  bool IsEven
->
-class LayerNorm {
-public:
-  using InputDtype = InputDtype_;
-  using OutputDtype = OutputDtype_;
-  using WeightDtype = WeightDtype_;
-  static constexpr int NumThrPerCta = NumThrPerCta_;
-  static constexpr int MaxHiddenSize = MaxHiddenSize_;
-  static constexpr bool Affine = Affine_;
-  static constexpr bool Bias = Bias_;
-
-  static constexpr size_t ShmSize = 32;
-  static constexpr int NumElementPerThread = MaxHiddenSize / NumThrPerCta;
-
-  static_assert(MaxHiddenSize % NumThrPerCta == 0);
-
-  struct Params {
-    void const *Iptr;
-    void const *Wptr;
-    void const *Bptr;
-    void *Optr;
-    float eps;
-    int64_t m;
-    int64_t n;
-  };
-
-  using Arguments = Params;
-
-  static Params to_underlying_arguments(Arguments const& args) {
-    return args;
-  }
-
-  static dim3 get_grid_size(int64_t m, int64_t n) {
-    return dim3(m);
-  }
-
-  static dim3 get_cta_size(int64_t m, int64_t n) {
-    return dim3(NumThrPerCta, 1, 1);
-  }
-
-  CUTLASS_DEVICE
-  void operator()(Params const& params, char *shared_data) {
-    int const blk_m = blockIdx.x;
-    int const blk_n = 1;
-    int tidx = threadIdx.x;
-    float x[NumElementPerThread];
-
-    // load
-    Loader<InputDtype, 1, MaxHiddenSize, NumThrPerCta, true, IsEven> loader;
-    loader.load(params.Iptr, x, params.m, params.n, blk_m, 0, tidx);
-
-    // 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) {
-      int idx = thr_n_offset + i;
-      if (idx < params.n) {
-        x[i] -= u;
-      } else {
-        x[i] = 0.f;
-      }
-    }
-
-    __syncthreads();
-    // var reduction
-    float v = sqrtf(_reduce_square(x, shared_data) / params.n + params.eps);
-
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      x[i] /= v;
-
-    if constexpr (Affine) {
-      // load weight
-      Loader<WeightDtype, 1, MaxHiddenSize, NumThrPerCta, true, IsEven> weight_loader;
-      float w[NumElementPerThread];
-      weight_loader.load(params.Wptr, w, 1, params.n, 0, 0, tidx);
-      if constexpr (Bias) {
-        float b[NumElementPerThread];
-        weight_loader.load(params.Bptr, b, 1, params.n, 0, 0, tidx);
-        CUTLASS_PRAGMA_UNROLL
-        for (int i = 0; i < NumElementPerThread; ++i)
-          x[i] = x[i] * w[i] + b[i];
-      } else {
-        CUTLASS_PRAGMA_UNROLL
-        for (int i = 0; i < NumElementPerThread; ++i)
-          x[i] = x[i] * w[i];
-      }
-    }
-
-    // save y
-    {
-      Saver<OutputDtype, 1, MaxHiddenSize, NumThrPerCta, true, IsEven, false, false> saver;
-      if constexpr (std::is_same_v<OutputDtype, float>) {
-        saver.store(params.Optr, nullptr, x, 0, params.m, params.n, blk_m, 0, tidx);
-      } else {
-        OutputDtype tmp[NumElementPerThread];
-        for (int i = 0; i < NumElementPerThread; ++i)
-          tmp[i] = OutputDtype(x[i]);
-        saver.store(params.Optr, nullptr, tmp, 0, params.m, params.n, blk_m, 0, tidx);
-      }
-    }
-    
-  }
-
-private:
-  CUTLASS_DEVICE
-  float _reduce_square(float *reg, char *shared_data) {
-    // thread
-    float sum_square = 0;
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      sum_square += reg[i] * reg[i];
-
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 16; i >= 1; i >>= 1) {
-      sum_square += __shfl_down_sync(0xFFFFFFFF, sum_square, i);
-    }
-    if (threadIdx.x == 0) {
-      *(float*)shared_data = 0;
-    }
-    __syncthreads();
-
-    if (threadIdx.x % 32 == 0) {
-      atomicAdd((float*)shared_data, sum_square);
-    }
-
-    __syncthreads();
-    sum_square = *(float*)shared_data;
-    return sum_square;
-  }
-
-  CUTLASS_DEVICE
-  float _reduce_sum(float *reg, char *shared_data) {
-    // thread
-    float sum = 0;
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      sum += reg[i];
-
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 16; i >= 1; i >>= 1) {
-      sum += __shfl_down_sync(0xFFFFFFFF, sum, i);
-    }
-    if (threadIdx.x == 0) {
-      *(float*)shared_data = 0;
-    }
-    __syncthreads();
-
-    if (threadIdx.x % 32 == 0) {
-      atomicAdd((float*)shared_data, sum);
-    }
-
-    __syncthreads();
-    sum = *(float*)shared_data;
-    return sum;
-  }
-};
-
-template <
-  class InputDtype,
-  class OutputDtype,
-  class WeightDtype,
-  bool Affine,
-  bool Bias,
-  int MaxHiddenSize,
-  int NumThrPerCta
->
-bool layernorm(
-  void const *Iptr, void const *Wptr, void const *Bptr,
-  void *Optr, float eps, int64_t m, int64_t n, 
-  cudaStream_t stream = nullptr
-) {
-  BOOL_SWITCH(n % MaxHiddenSize == 0, IsEven, [&] {
-    using Kernel = LayerNorm<
-      InputDtype, OutputDtype, WeightDtype,
-      Affine, Bias,
-      MaxHiddenSize, NumThrPerCta,
-      IsEven>;
-    using Arguments = typename Kernel::Arguments;
-    Arguments args = {
-      Iptr, Wptr, Bptr, Optr, 
-      eps, m, n
-    };
-    auto params = Kernel::to_underlying_arguments(args);
-    auto grid_shape = Kernel::get_grid_size(m, n);
-    auto cta_shape = Kernel::get_cta_size(m, n);
-    static constexpr size_t ShmSize = Kernel::ShmSize;
-    launch_kernel<Kernel>(params, grid_shape, cta_shape, ShmSize, stream);
-  });
-  return true;
-}
@@ -1,80 +0,0 @@
-#include <pybind11/pybind11.h>
-#include <torch/extension.h>
-#include <ATen/cuda/CUDAContext.h>
-#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"
-#include "norm/rmsnorm.hpp"
-
-auto rms_norm(
-  at::Tensor const& Input, 
-  float eps,
-  const std::optional<at::Tensor>& Weight,
-  std::optional<at::Tensor>& Output
-) {
-  
-  int64_t const m = Input.size(0);
-  int64_t const n = Input.size(1);
-  torch::Device const input_device = Input.device();
-  
-  if (!Output.has_value()) {
-    Output.emplace(
-      torch::empty(
-        {m, n},
-        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();
-
-  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;
-}
-
-void register_rms_norm(pybind11::module_ &m) {
-    m.def("rms_norm_cuda", &rms_norm);
-}
@@ -1,147 +0,0 @@
-#pragma once
-
-#include "common/load.hpp"
-#include "common/store.hpp"
-#include "common/launch.hpp"
-
-
-template <
-  class InputDtype_,
-  class OutputDtype_,
-  class WeightDtype_,
-  int MaxHiddenSize_,
-  int NumThrPerCta_,
-  bool IsEven
->
-class RMSNorm {
-public:
-  using InputDtype = InputDtype_;
-  using OutputDtype = OutputDtype_;
-  using WeightDtype = WeightDtype_;
-  static constexpr int NumThrPerCta = NumThrPerCta_;
-  static constexpr int MaxHiddenSize = MaxHiddenSize_;
-
-  static constexpr size_t ShmSize = 32;
-  static constexpr int NumElementPerThread = MaxHiddenSize / NumThrPerCta;
-
-  static_assert(MaxHiddenSize % NumThrPerCta == 0);
-
-  struct Params {
-    void const *Iptr;
-    void const *Wptr;
-    void *Optr;
-    float eps;
-    int64_t m;
-    int64_t n;
-  };
-
-  using Arguments = Params;
-
-  static Params to_underlying_arguments(Arguments const& args) {
-    return args;
-  }
-
-  static dim3 get_grid_size(int64_t m, int64_t n) {
-    return dim3(m);
-  }
-
-  static dim3 get_cta_size(int64_t m, int64_t n) {
-    return dim3(NumThrPerCta, 1, 1);
-  }
-
-  CUTLASS_DEVICE
-  void operator()(Params const& params, char *shared_data) {
-    int const blk_m = blockIdx.x;
-    int const blk_n = 1;
-    int tidx = threadIdx.x;
-    float x[NumElementPerThread];
-
-    // load
-    Loader<InputDtype, 1, MaxHiddenSize, NumThrPerCta, true, IsEven> loader;
-    loader.load(params.Iptr, x, params.m, params.n, blk_m, 0, tidx);
-
-    // rms reduction
-    float rms = sqrtf(_reduce_square(x, shared_data) / params.n + params.eps);
-
-    // load weight
-    Loader<WeightDtype, 1, MaxHiddenSize, NumThrPerCta, true, IsEven> weight_loader;
-    float w[NumElementPerThread];
-    loader.load(params.Wptr, w, 1, params.n, 0, 0, tidx);
-
-    // norm
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      x[i] = w[i] * x[i] / rms ;
-
-    // save y
-    OutputDtype *output_reg = (OutputDtype*)x;
-    if constexpr (!std::is_same_v<OutputDtype, float>) {
-      output_reg = (OutputDtype*)w;
-      CUTLASS_PRAGMA_UNROLL
-      for (int i = 0; i < NumElementPerThread; ++i)
-        output_reg[i] = OutputDtype(x[i]);
-    }
-    Saver<OutputDtype, 1, MaxHiddenSize, NumThrPerCta, true, IsEven, false, false> saver;
-    saver.store(params.Optr, nullptr, output_reg, 0, params.m, params.n, blk_m, 0, tidx);
-    
-  }
-
-private:
-  CUTLASS_DEVICE
-  float _reduce_square(float *reg, char *shared_data) {
-    // thread
-    float sum_square = 0;
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      sum_square += reg[i] * reg[i];
-
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 16; i >= 1; i >>= 1) {
-      sum_square += __shfl_down_sync(0xFFFFFFFF, sum_square, i);
-    }
-    if (threadIdx.x == 0) {
-      *(float*)shared_data = 0;
-    }
-    __syncthreads();
-
-    if (threadIdx.x % 32 == 0) {
-      atomicAdd((float*)shared_data, sum_square);
-    }
-
-    __syncthreads();
-    sum_square = *(float*)shared_data;
-    return sum_square;
-  }
-};
-
-
-template <
-  class InputDtype,
-  class OutputDtype,
-  class WeightDtype,
-  int MaxHiddenSize,
-  int NumThrPerCta
->
-bool rmsnorm(
-  void const *Iptr, void const *Wptr, 
-  void *Optr, float eps, 
-  int64_t m, int64_t n, 
-  cudaStream_t stream = nullptr
-) {
-  BOOL_SWITCH(n % MaxHiddenSize == 0, IsEven, [&] {
-    using Kernel = RMSNorm<
-      InputDtype, OutputDtype, WeightDtype,
-      MaxHiddenSize, NumThrPerCta,
-      IsEven>;
-    using Arguments = typename Kernel::Arguments;
-    Arguments args = {
-      Iptr, Wptr, Optr, eps, m, n
-    };
-    auto params = Kernel::to_underlying_arguments(args);
-    auto grid_shape = Kernel::get_grid_size(m, n);
-    auto cta_shape = Kernel::get_cta_size(m, n);
-    static constexpr size_t ShmSize = Kernel::ShmSize;
-    launch_kernel<Kernel>(params, grid_shape, cta_shape, ShmSize, stream);
-  });
-  return true;
-}
@@ -1,75 +0,0 @@
-/* 
- * Copyright (c) 2025 by TurboDiffusion team.
- * 
- * Licensed under the Apache License, Version 2.0 (the "License");
- * 
- * Citation (please cite if you use this code):
- * 
- * @article{zhang2025turbodiffusion,
- *   title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
- *   author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
- *   journal={arXiv preprint arXiv:2512.16093},
- *   year={2025}
- * }
- */
-
-#include <pybind11/pybind11.h>
-#include <torch/extension.h>
-#include <ATen/cuda/CUDAContext.h>
-#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"
-#include "quant/quant.hpp"
-
-auto quant(
-  torch::Tensor const& Input, 
-  std::optional<torch::Tensor>& Output, 
-  std::optional<torch::Tensor>& Output_S
-) {
-
-  using ElementOut = int8_t;
-  static constexpr int BlockSize = 128;
-  static constexpr int NumThrPerCta = 256;
-
-  int64_t m = Input.size(0);
-  int64_t n = Input.size(1);
-  torch::Device const input_device = Input.device();
-
-  create_tensor<BlockSize>(input_device, Output, Output_S, m, n);
-  
-  ElementOut *Optr = (ElementOut*)Output.value().data_ptr();
-  float *OSptr = Output_S.value().data_ptr<float>();
-
-  switch (Input.scalar_type()) {
-    case torch::kHalf:{
-        cutlass::half_t *Iptr = (cutlass::half_t*)Input.data_ptr();
-        quantization<cutlass::half_t, BlockSize, NumThrPerCta> (
-                    Iptr, Optr, OSptr, m, n, at::cuda::getCurrentCUDAStream().stream()
-                  );
-        break;
-    }
-
-    case torch::kBFloat16:{
-      cutlass::bfloat16_t *Iptr = (cutlass::bfloat16_t*)Input.data_ptr();
-      quantization<cutlass::bfloat16_t, BlockSize, NumThrPerCta> (
-                    Iptr, Optr, OSptr, m, n, at::cuda::getCurrentCUDAStream().stream()
-                  );
-      break;
-    }
-
-    default: {
-      std::cerr << "Observing: " << Input.scalar_type() << " for the input datatype which is invalid";
-      throw std::runtime_error("Unsupported input data type for quantize_to_fp4.");
-    }
-  }
-  
-  return std::make_tuple(Output, Output_S);
-}
-
-void register_quant(pybind11::module_ &m) {
-    m.def("quant_cuda", &quant);
-}
@@ -1,195 +0,0 @@
-/* 
- * Copyright (c) 2025 by TurboDiffusion team.
- * 
- * Licensed under the Apache License, Version 2.0 (the "License");
- * 
- * Citation (please cite if you use this code):
- * 
- * @article{zhang2025turbodiffusion,
- *   title={TurboDiffusion: Accelerating Video Diffusion Models by 100-200 Times},
- *   author={Zhang, Jintao and Zheng, Kaiwen and Jiang, Kai and Wang, Haoxu and Stoica, Ion and Gonzalez, Joseph E and Chen, Jianfei and Zhu, Jun},
- *   journal={arXiv preprint arXiv:2512.16093},
- *   year={2025}
- * }
- */
-
-#pragma once
-
-#include <cuda.h>
-#include <cuda_runtime.h>
-#include "cutlass/numeric_conversion.h"
-
-#include "common/load.hpp"
-#include "common/store.hpp"
-#include "common/launch.hpp"
-
-template <
-  class InputDtype_, 
-  int NumThrPerCta_,
-  bool IsEvenM,
-  bool IsEvenN
->
-class Quantization {
-public:
-  using InputDtype = InputDtype_;
-  using OutputDtype = int8_t;
-  using FPConverter = cutlass::NumericConverter<int8_t, float, cutlass::FloatRoundStyle::round_to_nearest>;
-
-  static constexpr int BlockSize = 128;
-  static constexpr int NumThrPerCta = NumThrPerCta_;
-  static constexpr int NumElementPerThread = BlockSize * BlockSize / NumThrPerCta;
-  static constexpr int NumThrPerRow = BlockSize / NumElementPerThread;
-
-  static_assert(BlockSize * BlockSize % NumThrPerCta == 0);
-  static_assert(NumThrPerCta % BlockSize == 0);
-
-  static constexpr size_t ShmSize = 32;
-
-  static constexpr float int8_max = 128.f;
-
-  struct Params {
-    void const *Iptr;
-    void *Optr;
-    void *OSptr;
-    int64_t const m;
-    int64_t const n;
-  };
-
-  using Arguments = Params;
-
-  static Params to_underlying_arguments(Arguments const& args) {
-    return args;
-  }
-
-  static dim3 get_grid_size(int64_t m, int64_t n) {
-    return dim3(
-      cdiv(n, BlockSize),
-      cdiv(m, BlockSize) 
-    );
-  }
-
-  static dim3 get_cta_size(int64_t m, int64_t n) {
-    return dim3(
-      NumThrPerCta, 1, 1
-    );
-  }
-  
-  CUTLASS_DEVICE
-  void quantization(
-    float *float_reg, 
-    void *Optr, void *OSptr, 
-    int64_t const m, int64_t const n, 
-    int blk_m, int blk_n, int tidx,
-    char *shared_data
-  ) {
-
-    OutputDtype output_reg[NumElementPerThread];
-
-
-    Saver<OutputDtype, BlockSize, BlockSize, NumThrPerCta, IsEvenM, IsEvenN> saver;
-
-    float amax = _reduce_amax(float_reg, (float*)shared_data);
-    
-    
-    _quantization(float_reg, output_reg, int8_max / amax);
-    
-    float scale_inv = amax / int8_max;
-
-    saver.store(Optr, OSptr, output_reg, scale_inv, m, n, blk_m, blk_n, tidx);
-
-    __syncthreads();
-    }
-  
-
-  CUTLASS_DEVICE 
-  void operator()(Params const& params, char *shared_data) {
-    int blk_m = blockIdx.y;
-    int blk_n = blockIdx.x;
-    int tidx = threadIdx.x;
-
-    float float_reg[NumElementPerThread];
-
-    // load float32 data
-    Loader<InputDtype, BlockSize, BlockSize, NumThrPerCta, IsEvenM, IsEvenN> loader;
-    loader.load(params.Iptr, float_reg, params.m, params.n, blk_m, blk_n, tidx);
-    quantization(
-      float_reg, params.Optr, params.OSptr, params.m, params.n, blk_m, blk_n, tidx, shared_data
-    );
-  }
-
-private: 
-
-  CUTLASS_DEVICE float 
-  _reduce_amax(float *reg, float *smem_ptr) {
-    float amax = 1e-8;
-    // thread reduction
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i)
-      amax = max(amax, fabs(reg[i]));
-
-    __syncwarp();
-
-    // warp reduction
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 16; i >= 1; i /= 2) {
-      amax = max(
-        __shfl_xor_sync(0xffffffff, amax, i, 32),
-        amax
-      );
-    }
-
-    // cta reduction
-    if (threadIdx.x == 0) {
-      *smem_ptr = 0;
-    }
-    __syncthreads();
-
-    atomicMax((uint32_t*)smem_ptr, reinterpret_cast<const uint32_t&>(amax));
-
-    __syncthreads();
-
-    amax = *smem_ptr;
-
-    return amax;
-  }
-
-  CUTLASS_DEVICE void
-  _quantization(float *float_reg, OutputDtype *out_reg, float scale) {
-    FPConverter converter;
-    CUTLASS_PRAGMA_UNROLL
-    for (int i = 0; i < NumElementPerThread; ++i) {
-        out_reg[i] = converter(float_reg[i] * scale);
-    }
-
-  }
-};
-
-template <
-  class InputDtype,
-  int BlockSize,
-  int NumThrPerCta
->
-bool quantization(
-  void const *Iptr, void *Optr, void *OSptr,
-  int64_t m, int64_t n, 
-  cudaStream_t stream = nullptr
-) {
-  BOOL_SWITCH(m % BlockSize == 0, IsEvenM, [&] {
-    BOOL_SWITCH(n % BlockSize == 0, IsEvenN, [&] {
-      using Kernel = Quantization<
-        InputDtype, NumThrPerCta, IsEvenM, IsEvenN>;
-      using Arguments = typename Kernel::Arguments;
-      Arguments args = {
-        Iptr, Optr, OSptr,
-        m, n
-      };
-      auto params = Kernel::to_underlying_arguments(args);
-      auto grid_shape = Kernel::get_grid_size(m, n);
-      auto cta_shape = Kernel::get_cta_size(m, n);
-      static constexpr size_t ShmSize = Kernel::ShmSize;
-      launch_kernel<Kernel>(params, grid_shape, cta_shape, ShmSize, stream);
-    });
-  });
-
-  return true;
-}
@@ -9,7 +9,7 @@ build-backend = "scikit_build_core.build"

 [project]
 name = "fastvideo-kernel"
-version = "0.2.5"
+version = "0.2.1"
 description = "Unified CUDA kernels for FastVideo"
 readme = "README.md"
 requires-python = ">=3.10"
@@ -11,24 +11,11 @@ from fastvideo_kernel.vmoba import (
    process_moba_output,
 )

-from fastvideo_kernel.turbodiffusion_ops import (
-    Int8Linear,
-    FastRMSNorm,
-    FastLayerNorm,
-    int8_linear,
-    int8_quant,
-)
-
 __all__ = [
    "sliding_tile_attention",
    "video_sparse_attn",
    "moba_attn_varlen",
    "process_moba_input",
    "process_moba_output",
-    "Int8Linear",
-    "FastRMSNorm",
-    "FastLayerNorm",
-    "int8_linear",
-    "int8_quant",
    "__version__",
 ]
@@ -1,295 +0,0 @@
-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,6 +1,5 @@
 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
@@ -46,22 +45,14 @@ 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_h, k_h,
-            v_h, o_h,
+            q[:, head_idx:head_idx + 1], k[:, head_idx:head_idx + 1],
+            v[:, head_idx:head_idx + 1], output[:, head_idx:head_idx + 1],
            t, h, w, text_length, False, has_text, flag
        )
-        output[:, head_idx:head_idx + 1] = o_h

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

    return output[:, :, :seq_length]

@@ -71,7 +62,6 @@ 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,
@@ -80,42 +70,14 @@ 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, 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()}"
-        )
+    batch, heads, seq_len, dim = q.shape

    # Compression branch
-    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.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_c.float().sum(dim=3) / q_variable_block_sizes.view(1, 1, -1, 1)).to(
+    q_c = (q_c.float().sum(dim=3) / 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)
@@ -126,9 +88,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, q_num_blocks, 1, dim)
+    out_c = out_c.view(batch, heads, seq_len // block_elements, 1, dim)
    out_c = out_c.repeat(1, 1, 1, block_elements,
-                         1).view(batch, heads, q_seq_len, dim)
+                         1).view(batch, heads, seq_len, dim)

    # Sparse branch
    topk_idx = torch.topk(scores, topk, dim=-1).indices
@@ -138,17 +100,12 @@ def video_sparse_attn(
    idx, num = map_to_index(mask)
    
    if block_sparse_fwd is not None:
-        # 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]
+        out_s = block_sparse_fwd(
+            q, k, v, idx, num, variable_block_sizes.int()
+        )[0] # block_sparse_fwd returns vector<Tensor>
    else:
-        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)
+        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,311 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# Adapted from TurboDiffusion SLA implementation
-# Copyright (c) 2025 by SLA team.
-#
-# Citation:
-# @article{zhang2025sla,
-#   title={SLA: Beyond Sparsity in Diffusion Transformers via Fine-Tunable Sparse-Linear Attention},
-#   author={Jintao Zhang and Haoxu Wang and Kai Jiang and Shuo Yang and Kaiwen Zheng and
-#           Haocheng Xi and Ziteng Wang and Hongzhou Zhu and Min Zhao and Ion Stoica and
-#           Joseph E. Gonzalez and Jun Zhu and Jianfei Chen},
-#   journal={arXiv preprint arXiv:2509.24006},
-#   year={2025}
-# }
-
-import torch
-import triton
-import triton.language as tl
-
-
-@triton.jit
-def _attn_fwd(
-    Q, K, V,
-    qk_scale: tl.constexpr,
-    topk: tl.constexpr,
-    LUT, LSE, OS,
-    L: tl.constexpr,
-    M_BLOCKS: tl.constexpr,
-    D: tl.constexpr,
-    BLOCK_M: tl.constexpr,
-    BLOCK_N: tl.constexpr,
-):
-    idx_m = tl.program_id(0).to(tl.int64)
-    idx_bh = tl.program_id(1).to(tl.int64)
-
-    qkv_offset = idx_bh * L * D
-    lut_offset = (idx_bh * M_BLOCKS + idx_m) * topk
-    lse_offset = idx_bh * L
-    offs_m = idx_m * BLOCK_M + tl.arange(0, BLOCK_M)
-    offs_n = tl.arange(0, BLOCK_N)
-    offs_d = tl.arange(0, D)
-
-    Q_ptrs = Q + qkv_offset + offs_m[:, None] * D + offs_d[None, :]
-    K_ptrs = K + qkv_offset + offs_n[None, :] * D + offs_d[:, None]
-    V_ptrs = V + qkv_offset + offs_n[:, None] * D + offs_d[None, :]
-    OS_ptrs = OS + qkv_offset + offs_m[:, None] * D + offs_d[None, :]
-    LUT_ptr = LUT + lut_offset
-    LSE_ptrs = LSE + lse_offset + offs_m
-    
-    m_i = tl.full([BLOCK_M], -float('inf'), dtype=tl.float32)
-    l_i = tl.zeros([BLOCK_M], dtype=tl.float32)
-    o_s = tl.zeros([BLOCK_M, D], dtype=tl.float32)
-
-    q = tl.load(Q_ptrs, mask=offs_m[:, None] < L)
-    for block_idx in tl.range(topk):
-        idx_n = tl.load(LUT_ptr + block_idx)
-        n_mask = offs_n < L - idx_n * BLOCK_N
-        
-        k = tl.load(K_ptrs + idx_n * BLOCK_N * D, mask=n_mask[None, :])
-        qk = tl.dot(q, k) * (qk_scale * 1.4426950408889634)  # = 1 / ln(2)
-        if L - idx_n * BLOCK_N < BLOCK_N:
-            qk = tl.where(n_mask[None, :], qk, float("-inf"))
-
-        v = tl.load(V_ptrs + idx_n * BLOCK_N * D, mask=n_mask[:, None])
-        local_m = tl.max(qk, 1)
-        new_m = tl.maximum(m_i, local_m)
-        qk = qk - new_m[:, None]
-
-        p = tl.math.exp2(qk)
-        l_ij = tl.sum(p, 1)
-        alpha = tl.math.exp2(m_i - new_m)
-        o_s = o_s * alpha[:, None]
-        o_s += tl.dot(p.to(v.dtype), v)
-
-        l_i = l_i * alpha + l_ij
-        m_i = new_m
-
-    o_s = o_s / l_i[:, None]
-    tl.store(OS_ptrs, o_s.to(OS.type.element_ty), mask=offs_m[:, None] < L)
-    
-    m_i += tl.math.log2(l_i)
-    tl.store(LSE_ptrs, m_i, mask=offs_m < L)
-
-
-@triton.jit
-def _attn_bwd_preprocess(
-    OS, DOS, DELTAS,
-    L,
-    D: tl.constexpr,
-    BLOCK_M: tl.constexpr,
-):
-    idx_m = tl.program_id(0).to(tl.int64)
-    idx_bh = tl.program_id(1).to(tl.int64)
-
-    OS += idx_bh * L * D
-    DOS += idx_bh * L * D
-    DELTAS += idx_bh * L
-
-    offs_m = idx_m * BLOCK_M + tl.arange(0, BLOCK_M)
-    offs_d = tl.arange(0, D)
-
-    o_s = tl.load(OS + offs_m[:, None] * D + offs_d[None, :], mask=offs_m[:, None] < L)
-    do_s = tl.load(DOS + offs_m[:, None] * D + offs_d[None, :], mask=offs_m[:, None] < L)
-    
-    delta_s = tl.sum(o_s * do_s, axis=1).to(DELTAS.type.element_ty)
-    tl.store(DELTAS + offs_m, delta_s, mask=offs_m < L)
-
-
-@triton.jit
-def _attn_bwd_dq(
-    Q, K, V, LSE, DELTAS,
-    DOS, DQ, LUT,
-    qk_scale: tl.constexpr,
-    topk: tl.constexpr,
-    L: tl.constexpr,
-    M_BLOCKS: tl.constexpr,
-    D: tl.constexpr,
-    BLOCK_M: tl.constexpr,
-    BLOCK_N: tl.constexpr,
-):
-    idx_m = tl.program_id(0).to(tl.int64)
-    idx_bh = tl.program_id(1).to(tl.int64)
-
-    offs_m = idx_m * BLOCK_M + tl.arange(0, BLOCK_M)
-    offs_n = tl.arange(0, BLOCK_N)
-    offs_d = tl.arange(0, D)
-
-    qkv_offset = idx_bh * L * D
-    lse_offset = idx_bh * L
-    lut_offset = (idx_bh * M_BLOCKS + idx_m) * topk
-
-    Q_ptrs = Q + qkv_offset + offs_m[:, None] * D + offs_d[None, :]
-    K_ptrs = K + qkv_offset + offs_n[:, None] * D + offs_d[None, :]
-    V_ptrs = V + qkv_offset + offs_n[:, None] * D + offs_d[None, :]
-    DQ_ptrs = DQ + qkv_offset + offs_m[:, None] * D + offs_d[None, :]
-    DOS_ptrs = DOS + qkv_offset + offs_m[:, None] * D + offs_d[None, :]
-    LSE_ptrs = LSE + lse_offset + offs_m
-    DELTAS_ptrs = DELTAS + lse_offset + offs_m
-    LUT_ptr = LUT + lut_offset
-
-    q = tl.load(Q_ptrs, mask=offs_m[:, None] < L)
-    do_s = tl.load(DOS_ptrs, mask=offs_m[:, None] < L)
-    delta_s = tl.load(DELTAS_ptrs, mask=offs_m < L)
-    lse = tl.load(LSE_ptrs, mask=offs_m < L, other=float("inf"))
-    
-    dq = tl.zeros([BLOCK_M, D], dtype=tl.float32)
-    for block_idx in tl.range(topk, num_stages=2):
-        idx_n = tl.load(LUT_ptr + block_idx)
-        n_mask = offs_n < L - idx_n * BLOCK_N
-        
-        k = tl.load(K_ptrs + idx_n * BLOCK_N * D, mask=n_mask[:, None])
-        v = tl.load(V_ptrs + idx_n * BLOCK_N * D, mask=n_mask[:, None])
-        qk = tl.dot(q, k.T) * (qk_scale * 1.4426950408889634)
-        p = tl.math.exp2(qk - lse[:, None])
-        p = tl.where(n_mask[None, :], p, 0.0)
-        
-        dp = tl.dot(do_s, v.T).to(tl.float32)
-        ds = p * (dp - delta_s[:, None])
-        dq += tl.dot(ds.to(k.dtype), k)
-    tl.store(DQ_ptrs, dq * qk_scale, mask=offs_m[:, None] < L)
-    
-
-@triton.jit
-def _attn_bwd_dkdv(
-    Q, K, V, DOS, DK, DV,
-    qk_scale, KBID, LSE, DELTAS,
-    L: tl.constexpr,
-    M_BLOCKS: tl.constexpr,
-    N_BLOCKS: tl.constexpr,
-    D: tl.constexpr,
-    BLOCK_M: tl.constexpr,
-    BLOCK_N: tl.constexpr,
-    BLOCK_SLICE_FACTOR: tl.constexpr,
-):
-    BLOCK_M2: tl.constexpr = BLOCK_M // BLOCK_SLICE_FACTOR
-
-    idx_n = tl.program_id(0).to(tl.int64)
-    idx_bh = tl.program_id(1).to(tl.int64)
-
-    offs_n = idx_n * BLOCK_N + tl.arange(0, BLOCK_N)
-    offs_m = tl.arange(0, BLOCK_M2)
-    offs_d = tl.arange(0, D)
-
-    qkv_offset = idx_bh * L * D
-    kbid_offset = idx_bh * M_BLOCKS * N_BLOCKS
-    lse_offset = idx_bh * L
-
-    Q_ptrs = Q + qkv_offset + offs_m[:, None] * D + offs_d[None, :]
-    K_ptrs = K + qkv_offset + offs_n[:, None] * D + offs_d[None, :]
-    V_ptrs = V + qkv_offset + offs_n[:, None] * D + offs_d[None, :]
-    DOS_ptrs = DOS + qkv_offset + offs_m[:, None] * D + offs_d[None, :]
-    DK_ptrs = DK + qkv_offset + offs_n[:, None] * D + offs_d[None, :]
-    DV_ptrs = DV + qkv_offset + offs_n[:, None] * D + offs_d[None, :]
-    LSE_ptrs = LSE + lse_offset + offs_m
-    DELTAS_ptrs = DELTAS + lse_offset + offs_m
-    KBID_ptr = KBID + kbid_offset + idx_n
-
-    k = tl.load(K_ptrs, mask=offs_n[:, None] < L)
-    v = tl.load(V_ptrs, mask=offs_n[:, None] < L)
-        
-    dk = tl.zeros([BLOCK_N, D], dtype=tl.float32)
-    dv = tl.zeros([BLOCK_N, D], dtype=tl.float32)
-    for idx_m in tl.range(0, L, BLOCK_M2):
-        kbid = tl.load(KBID_ptr)
-        if kbid == 1:
-            m_mask = offs_m < L - idx_m
-            q = tl.load(Q_ptrs, mask=m_mask[:, None])
-            lse = tl.load(LSE_ptrs, mask=m_mask, other=float("inf"))
-            qkT = tl.dot(k, q.T) * (qk_scale * 1.4426950408889634)
-            pT = tl.math.exp2(qkT - lse[None, :])
-            pT = tl.where(offs_n[:, None] < L, pT, 0.0)
-
-            do = tl.load(DOS_ptrs, mask=m_mask[:, None])
-            dv += tl.dot(pT.to(do.dtype), do)
-            delta = tl.load(DELTAS_ptrs, mask=m_mask)
-            dpT = tl.dot(v, tl.trans(do))
-            dsT = pT * (dpT - delta[None, :])
-            dk += tl.dot(dsT.to(q.dtype), q)
-        
-        Q_ptrs += BLOCK_M2 * D
-        DOS_ptrs += BLOCK_M2 * D
-        LSE_ptrs += BLOCK_M2
-        DELTAS_ptrs += BLOCK_M2
-        if (idx_m + BLOCK_M2) % BLOCK_M == 0:
-            KBID_ptr += N_BLOCKS
-
-    tl.store(DK_ptrs, dk * qk_scale, mask=offs_n[:, None] < L)
-    tl.store(DV_ptrs, dv, mask=offs_n[:, None] < L)
-    
-
-class _attention(torch.autograd.Function):
-    """Sparse attention forward/backward with autograd support."""
-    
-    @staticmethod
-    def forward(ctx, q, k, v, k_block_id, lut, topk, BLOCK_M, BLOCK_N, qk_scale=None):
-        assert q.is_contiguous() and k.is_contiguous() and v.is_contiguous()
-        assert k_block_id.is_contiguous() and lut.is_contiguous()
-
-        assert BLOCK_M == 64 or BLOCK_M == 128
-        assert BLOCK_N == 64
-
-        B, H, L, D = q.shape
-        if qk_scale is None:
-            qk_scale = D**-0.5
-
-        M_BLOCKS = triton.cdiv(L, BLOCK_M)
-
-        o_s = torch.empty_like(v)
-        lse = torch.empty(q.shape[:-1], device=q.device, dtype=torch.float32)
-
-        grid = (M_BLOCKS, B * H)
-        _attn_fwd[grid](
-            q, k, v, qk_scale, topk,
-            lut, lse, o_s,
-            L, M_BLOCKS,
-            D, BLOCK_M, BLOCK_N,
-            num_warps=4 if q.shape[-1] == 64 else 8,
-            num_stages=3
-        )
-        
-        ctx.save_for_backward(q, k, v, k_block_id, lut, lse, o_s)
-        ctx.qk_scale = qk_scale
-        ctx.topk = topk
-        ctx.BLOCK_M = BLOCK_M
-        ctx.BLOCK_N = BLOCK_N
-        return o_s
-
-    @staticmethod
-    def backward(ctx, do_s):
-        q, k, v, k_block_id, lut, lse, o_s = ctx.saved_tensors
-        do_s = do_s.contiguous()
-
-        BLOCK_M, BLOCK_N = ctx.BLOCK_M, ctx.BLOCK_N
-        B, H, L, D = q.shape
-
-        M_BLOCKS = triton.cdiv(L, BLOCK_M)
-        N_BLOCKS = triton.cdiv(L, BLOCK_N)
-
-        dq = torch.empty_like(q)
-        dk = torch.empty_like(k)
-        dv = torch.empty_like(v)
-        delta_s = torch.empty_like(lse)
-
-        grid = (M_BLOCKS, B * H)
-        _attn_bwd_preprocess[grid](
-            o_s, do_s, delta_s,
-            L, D, BLOCK_M,
-        )
-
-        grid = (M_BLOCKS, B * H)
-        _attn_bwd_dq[grid](
-            q, k, v, lse, delta_s,
-            do_s, dq, lut,
-            ctx.qk_scale, ctx.topk,
-            L, M_BLOCKS,
-            D, BLOCK_M, BLOCK_N,
-            num_warps=4 if q.shape[-1] == 64 else 8,
-            num_stages=4 if q.shape[-1] == 64 else 5
-        )
-
-        grid = (N_BLOCKS, B * H)
-        _attn_bwd_dkdv[grid](
-            q, k, v, do_s, dk, dv,
-            ctx.qk_scale, k_block_id, lse, delta_s,
-            L, M_BLOCKS, N_BLOCKS,
-            D, BLOCK_M, BLOCK_N,
-            BLOCK_SLICE_FACTOR=BLOCK_M // 64,
-            num_warps=4 if q.shape[-1] == 64 else 8,
-            num_stages=4 if q.shape[-1] == 64 else 5
-        )
-
-        return dq, dk, dv, None, None, None, None, None, None
@@ -11,20 +11,12 @@ def is_hip():
    target = triton.runtime.driver.active.get_current_target()
    return target.backend == 'hip'

-
-def is_cdna3_cdna4():
-    target = triton.runtime.driver.active.get_current_target()
-    return (target.arch == 'gfx1201' or target.arch == 'gfx1101' or target.arch == 'gfx1100' or target.arch == 'gfx1030')
-
-
 def get_common_autotune_config():
-    # cdna arch does not support a 4-stage software pipeline, see https://github.com/ROCm/triton/issues/916
-    supported_num_staged = [1, 2] if is_cdna3_cdna4() else [1, 2, 3, 4]
    configs = [
        triton.Config({'BLOCK_Q': BLOCK_Q, 'BLOCK_KV': BLOCK_KV}, num_stages=s, num_warps=w) \
        for BLOCK_Q in [32, 64, 128]\
        for BLOCK_KV in [32, 64, 128]\
-        for s in supported_num_staged\
+        for s in [1, 2, 3, 4]\
        for w in [4, 8]\
    ]
    return configs
@@ -1,716 +0,0 @@
-from __future__ import annotations
-
-from typing import Optional, Tuple
-import torch
-import torch.nn as nn
-import torch.nn.functional as F
-import triton
-import triton.language as tl
-
-# Try to load the C++ extension
-try:
-    from fastvideo_kernel._C import fastvideo_kernel_ops
-    quant_cuda = getattr(fastvideo_kernel_ops, "quant_cuda", None)
-    gemm_cuda = getattr(fastvideo_kernel_ops, "gemm_cuda", None)
-    rms_norm_cuda = getattr(fastvideo_kernel_ops, "rms_norm_cuda", None)
-    layer_norm_cuda = getattr(fastvideo_kernel_ops, "layer_norm_cuda", None)
-except ImportError:
-    quant_cuda = None
-    gemm_cuda = None
-    rms_norm_cuda = None
-    layer_norm_cuda = None
-
-def int8_quant(x: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
-    """
-    Quantize a floating-point tensor to int8 using a custom CUDA kernel.
-
-    Args:
-        x (torch.Tensor): Input tensor of type float16/bfloat16.
-
-    Returns:
-        Tuple[torch.Tensor, torch.Tensor]:
-            - x_q: Quantized int8 tensor.
-            - x_scale: Per-block scale tensor used for quantization.
-    """
-    x_q, x_scale = quant_cuda(x, None, None)
-    return x_q, x_scale
-
-
-def int8_linear(
-    x: torch.Tensor,
-    w_q: torch.Tensor,
-    w_s: torch.Tensor,
-    **kwargs,
-) -> torch.Tensor:
-    """
-    Perform an int8 GEMM (matrix multiplication) using quantized weights and a
-    quantized version of the input. The underlying compute is performed by a
-    custom CUDA kernel.
-
-    Args:
-        x (torch.Tensor): Input activation of shape (M, K) in float32.
-        w_q (torch.Tensor): Quantized int8 weight tensor of shape (N, K).
-        w_s (torch.Tensor): Scale tensor associated with w_q.
-        **kwargs: Additional options (reserved for future use).
-
-    Returns:
-        torch.Tensor: Output tensor of shape (M, N) in float32.
-    """
-    assert w_q.dtype == torch.int8, "Weight tensor must be int8."
-    shape = x.shape
-    x = x.reshape(-1, shape[-1])
-    m = x.shape[0]
-    n = w_q.shape[0]
-    y = torch.zeros(m, n, dtype=x.dtype, device=x.device)
-
-    x_q, x_s = int8_quant(x)
-    gemm_cuda(x_q, x_s, w_q, w_s, y)
-    return y.reshape(*shape[:-1], n)
-
-def flatten_if_batched(*tensors):
-    """
-    Flattens all input tensors from (B, N, D_i) to (B * N, D_i) if they are batched (3D).
-
-    Args:
-        *tensors: Any number of input tensors, each must have shape (B, N, D_i) or (N, D_i)
-
-    Returns:
-        flat_tensors: List of flattened tensors
-        batched: Boolean flag indicating whether inputs were batched
-        batch_size: Batch size if batched, else None
-    """
-    if not tensors:
-        raise ValueError("At least one tensor must be provided.")
-
-    first = tensors[0]
-    assert len(first.shape) in [
-        2,
-        3,
-    ], "Input tensors must be batched (3D) or not batched (2D)"
-
-    if len(first.shape) == 3:  # batched
-        batched = True
-        batch_size = first.shape[0]
-        assert all(t.shape[0] == batch_size for t in tensors), "All input tensors must have the same batch size"
-        assert all(
-            t.shape[1] == first.shape[1] for t in tensors
-        ), "All input tensors must have the same sequence length"
-        flat_tensors = [t.reshape(-1, t.shape[-1]) for t in tensors]
-    else:
-        batched = False
-        batch_size = None
-        flat_tensors = list(tensors)
-
-    return flat_tensors, batched, batch_size
-
-
-@triton.jit
-def _rms_norm_fwd_fused(
-    X,
-    Y,
-    W,
-    Rstd,
-    x_stride,
-    y_stride,
-    N: tl.constexpr,  # number of columns in X,
-    N2: tl.constexpr,
-    eps,  # epsilon to avoid division by zero
-    BLOCK_M: tl.constexpr,
-):
-    # Map the program id to the row of X and Y it should compute.
-    pid = tl.program_id(0)
-    rows = pid * BLOCK_M + tl.arange(0, BLOCK_M)
-    cols = tl.arange(0, N2)
-    mask = cols < N
-
-    x_ptr = X + rows[:, None] * x_stride + cols[None, :]
-    y_ptr = Y + rows[:, None] * y_stride + cols[None, :]
-
-    x = tl.load(x_ptr, mask=mask[None, :], other=0.0).to(tl.float32)
-
-    # Compute variance
-    _var = x * x
-    var = tl.sum(_var, axis=1) / N
-    rstd = 1 / tl.sqrt(var + eps)
-
-    # Write mean / rstd
-    tl.store(Rstd + rows, rstd)
-    rstd = tl.reshape(rstd, (BLOCK_M, 1))
-
-    # Normalize and apply linear transformation
-    w = tl.load(W + cols)
-    x_hat = x * rstd
-    y = x_hat * w
-
-    # Write output
-    y = y.to(Y.type.element_ty)
-    tl.store(y_ptr, y, mask=mask[None, :])
-
-
-def rmsnorm(x, w, eps):
-    """
-    Forward pass of the RMSNorm.
-
-    Args:
-        x (torch.Tensor): Input tensor, High precision.
-        w (torch.Tensor): RMSNorm weight tensor.
-        eps (float): RMSNorm epsilon value.
-
-    Returns:
-        y (torch.Tensor): Output tensor, High precision.
-        rstd (torch.Tensor): Inverse standard deviation, needed for backward.
-    """
-    assert x.is_contiguous(), "Input must be contiguous"
-    # Change batched 3D input to 2D
-    [x], batched, BS = flatten_if_batched(x)
-
-    # allocate output
-    M, N = x.shape
-    y = torch.empty_like(x, dtype=x.dtype)
-    rstd = torch.empty((M,), dtype=torch.float32, device=x.device)
-
-    # heuristics for number of warps
-    num_warps = 8
-    
-    # Avoid illegal memory access
-    N2 = triton.next_power_of_2(N)
-    
-    if N <= 512:
-        BLOCK_M = 32
-    else:
-        BLOCK_M = 1
-
-    # Call the triton kernel
-    _rms_norm_fwd_fused[(triton.cdiv(M, BLOCK_M),)](  #
-        x,
-        y,
-        w,
-        rstd,  #
-        x.stride(0),
-        y.stride(0),
-        N,
-        N2,
-        eps,
-        num_warps=num_warps,
-        BLOCK_M=BLOCK_M,
-    )
-
-    # Recover 2D to 3D
-    if batched:
-        y = y.reshape(BS, -1, y.shape[-1])
-
-    return y, rstd
-
-@triton.jit
-def _layer_norm_param_fwd_fused(
-    X,  # pointer to the input
-    Y,  # pointer to the output
-    W,  # pointer to the weights
-    B,  # pointer to the biases
-    Mean,  # pointer to the mean
-    Rstd,  # pointer to the 1/std
-    x_stride,  # how much to increase the pointer when moving by 1 row
-    y_stride,  # how much to increase the pointer when moving by 1 row
-    N: tl.constexpr,  # number of columns in X,
-    N2: tl.constexpr,  # number of columns in X,
-    eps,  # epsilon to avoid division by zero
-    BLOCK_M: tl.constexpr,
-):
-    # Map the program id to the row of X and Y it should compute.
-    pid = tl.program_id(0)
-    rows = pid * BLOCK_M + tl.arange(0, BLOCK_M)
-    cols = tl.arange(0, N2)
-    mask = cols < N
-
-    x_ptr = X + rows[:, None] * x_stride + cols[None, :]
-    y_ptr = Y + rows[:, None] * y_stride + cols[None, :]
-
-    x = tl.load(x_ptr, mask=mask[None, :], other=0.0).to(tl.float32)
-
-    # Compute mean and Variance
-    mean = tl.sum(x, axis=1, keep_dims=True) / N
-    # Compute variance
-    _var = (x - mean) * (x - mean)
-    var = tl.sum(_var, axis=1, keep_dims=True) / N
-    rstd = 1 / tl.sqrt(var + eps)
-
-    # Write mean / rstd
-    _mean = tl.reshape(mean, (BLOCK_M))
-    _rstd = tl.reshape(rstd, (BLOCK_M))
-    tl.store(Mean + rows, _mean)
-    tl.store(Rstd + rows, _rstd)
-
-    # Normalize and apply linear transformation
-    x_hat = (x - mean) * rstd
-
-    w = tl.load(W + cols)
-    b = tl.load(B + cols)
-    
-    x_hat = x_hat * w + b
-
-    # Write output
-    x_hat = x_hat.to(Y.type.element_ty)
-    tl.store(y_ptr, x_hat, mask=mask[None, :])
-
-
-def layernorm_param(x, w, b, eps):
-    # Change batched 3D input to 2D
-    [x], batched, BS = flatten_if_batched(x)
-
-    # allocate output
-    M, N = x.shape
-    y = torch.empty_like(x, dtype=torch.float32)
-    mean = torch.empty((M,), dtype=torch.float32, device=x.device)
-    rstd = torch.empty((M,), dtype=torch.float32, device=x.device)
-    # heuristics for number of warps
-    num_warps = 8
-
-    N2 = triton.next_power_of_2(N)
-    
-    if N <= 512:
-        BLOCK_M = 32
-    else:
-        BLOCK_M = 1
-
-    # enqueue kernel
-    _layer_norm_param_fwd_fused[(triton.cdiv(M, BLOCK_M),)](  #
-        x,
-        y,
-        w,
-        b,
-        mean,
-        rstd,  #
-        x.stride(0),
-        y.stride(0),
-        N,
-        N2,
-        eps,
-        num_warps=num_warps,
-        BLOCK_M=BLOCK_M,
-    )
-
-    # Recover 2D to 3D
-    if batched:
-        y = y.reshape(BS, -1, y.shape[-1])
-
-    return y, mean, rstd
-
-
-########################################################
-# Elementwise_affine=False
-########################################################
-
-
-@triton.jit
-def _layer_norm_noparam_fwd_fused(
-    X,  # pointer to the input
-    Y,  # pointer to the output
-    Mean,  # pointer to the mean
-    Rstd,  # pointer to the 1/std
-    x_stride,  # how much to increase the pointer when moving by 1 row
-    y_stride,  # how much to increase the pointer when moving by 1 row
-    N: tl.constexpr,  # number of columns in X,
-    N2: tl.constexpr,  # number of columns in X,
-    eps,  # epsilon to avoid division by zero
-    BLOCK_M: tl.constexpr,
-):
-    # Map the program id to the row of X and Y it should compute.
-    pid = tl.program_id(0)
-    rows = pid * BLOCK_M + tl.arange(0, BLOCK_M)
-    cols = tl.arange(0, N2)
-    mask = cols < N
-
-    x_ptr = X + rows[:, None] * x_stride + cols[None, :]
-    y_ptr = Y + rows[:, None] * y_stride + cols[None, :]
-
-    x = tl.load(x_ptr, mask=mask[None, :], other=0.0).to(tl.float32)
-
-    # Compute mean and Variance
-    mean = tl.sum(x, axis=1, keep_dims=True) / N
-    # Compute variance
-    _var = (x - mean) * (x - mean)
-    var = tl.sum(_var, axis=1, keep_dims=True) / N
-    rstd = 1 / tl.sqrt(var + eps)
-
-    # Write mean / rstd
-    _mean = tl.reshape(mean, (BLOCK_M))
-    _rstd = tl.reshape(rstd, (BLOCK_M))
-    tl.store(Mean + rows, _mean)
-    tl.store(Rstd + rows, _rstd)
-
-    # Normalize and apply linear transformation
-    x_hat = (x - mean) * rstd
-
-    # Write output
-    x_hat = x_hat.to(Y.type.element_ty)
-    tl.store(y_ptr, x_hat, mask=mask[None, :])
-
-
-def layernorm_noparam(x, eps):
-    assert x.is_contiguous(), "Input must be contiguous"
-
-    # Change batched 3D input to 2D
-    [x], batched, BS = flatten_if_batched(x)
-
-    # allocate output
-    M, N = x.shape
-    y = torch.empty_like(x, dtype=torch.float32)
-    mean = torch.empty((M,), dtype=torch.float32, device=x.device)
-    rstd = torch.empty((M,), dtype=torch.float32, device=x.device)
-    # heuristics for number of warps
-    num_warps = 8
-
-    N2 = triton.next_power_of_2(N)
-    
-    if N <= 512:
-        BLOCK_M = 32
-    else:
-        BLOCK_M = 1
-
-    # enqueue kernel
-    _layer_norm_noparam_fwd_fused[(triton.cdiv(M, BLOCK_M),)](  #
-        x,
-        y,
-        mean,
-        rstd,  #
-        x.stride(0),
-        y.stride(0),
-        N,
-        N2,
-        eps,
-        num_warps=num_warps,
-        BLOCK_M=BLOCK_M,
-    )
-
-    # Recover 2D to 3D
-    if batched:
-        y = y.reshape(BS, -1, y.shape[-1])
-
-    return y, mean, rstd
-
-def layernorm(x, w, b, eps, elementwise_affine=True):
-    if elementwise_affine:
-        assert w is not None and b is not None
-        return layernorm_param(x, w, b, eps)
-    else:
-        assert w is None and b is None
-        return layernorm_noparam(x, eps)
-
-def cdiv(a: int, b: int):
-    return (a + b - 1) // b
-
-def dequantize_weight(w_q, w_s, dtype):
-    # w_q: (N, K) int8
-    # w_s: (NB, KB) float32
-    # Block size 128
-    BLOCK = 128
-    N, K = w_q.shape
-    # Expand w_s
-    # Repeat interleave
-    w_s_exp = w_s.repeat_interleave(BLOCK, dim=0).repeat_interleave(BLOCK, dim=1)
-    # Crop
-    w_s_exp = w_s_exp[:N, :K]
-    return w_q.to(dtype) * w_s_exp.to(dtype)
-
-class Int8LinearFunction(torch.autograd.Function):
-    @staticmethod
-    def forward(ctx, x, w_q, w_s, bias=None):
-        ctx.save_for_backward(x, w_q, w_s, bias)
-        ctx.bias_requires_grad = bias.requires_grad if bias is not None else False
-        
-        out = int8_linear(x, w_q, w_s)
-        if bias is not None:
-            out = out + bias
-        return out
-
-    @staticmethod
-    def backward(ctx, grad_output):
-        x, w_q, w_s, bias = ctx.saved_tensors
-        
-        grad_input = None
-        grad_weight = None 
-        grad_scale = None
-        grad_bias = None
-        
-        if ctx.needs_input_grad[0]: 
-            # grad_input = grad_output @ W
-            w_float = dequantize_weight(w_q, w_s, grad_output.dtype)
-            grad_input = torch.matmul(grad_output, w_float)
-            
-        if ctx.bias_requires_grad:
-            dim_to_sum = list(range(grad_output.dim() - 1))
-            grad_bias = grad_output.sum(dim=dim_to_sum)
-            
-        return grad_input, grad_weight, grad_scale, grad_bias
-
-class RMSNormFunction(torch.autograd.Function):
-    @staticmethod
-    def forward(ctx, x, w, eps):
-        y, rstd = rmsnorm(x, w, eps)
-        ctx.save_for_backward(x, w, rstd)
-        ctx.eps = eps
-        return y
-
-    @staticmethod
-    def backward(ctx, grad_output):
-        # x, w, rstd are saved
-        # grad_output is dL/dy
-        # We can implement backward using torch ops for correctness
-        x, w, rstd = ctx.saved_tensors
-        eps = ctx.eps
-        N = x.shape[-1]
-        
-        # dL/dy * w
-        dx = grad_output * w
-        
-        # Expand rstd
-        rstd = rstd.unsqueeze(-1) # (B, 1) or (M, 1)
-        if x.dim() == 3:
-            # Flatten if necessary or handle dimensions. 
-            # forward flattens x to (M, N) internally but saved x might be 3D?
-            # rmsnorm takes x, flattens it. But does it modify x in place? No.
-            # But ctx.save_for_backward saves the original x (3D if input was 3D).
-            # rstd is (M,).
-            # We should flatten x and grad_output to match logic
-            x_flat = x.reshape(-1, N)
-            grad_output_flat = grad_output.reshape(-1, N)
-            dx = dx.reshape(-1, N)
-        else:
-            x_flat = x
-            grad_output_flat = grad_output
-            dx = dx
-            
-        # Standard RMSNorm backward
-        # dy = grad_output_flat
-        # x_hat = x * rstd
-        # w * dy
-        # c1 = mean(dy * w * x) * rstd^2
-        # dx = (dy * w - c1 * x) * rstd
-        
-        # More precise:
-        # y = x * rstd * w
-        # dy = grad_output
-        # dw = sum(dy * x * rstd, dim=0)
-        
-        grad_w = None
-        if ctx.needs_input_grad[1]: # w
-             # grad_w = (grad_output_flat * x_flat * rstd).sum(dim=0)
-             # More accurate gradient for weight when considering rstd was computed from x
-             # Actually, for Affine part (w), it is just sum(dL/dy * x_hat).
-             # y = x_hat * w
-             # dL/dw = sum(dL/dy * x_hat)
-             x_hat = x_flat * rstd
-             grad_w = (grad_output_flat * x_hat).sum(dim=0)
-
-        grad_input = None
-        if ctx.needs_input_grad[0]: # x
-            # x_hat = x * rstd
-            # dx = rstd * (w * dy - mean(w * dy * x_hat) * x_hat) 
-            # but for RMSNorm:
-            # dx = rstd * (w * dy - (x * rstd^2) * mean(w * dy * x)) -> check formula
-            
-            # Using PyTorch autograd for reference logic:
-            # y = x / sqrt(mean(x^2) + eps) * w
-            # Let sigma = sqrt(...)
-            # dL/dx = dL/dy * w * (1/sigma) + dL/dsigma * dsigma/dx
-            # dsigma/dx = 1/(2*sigma) * 2x/N = x / (N * sigma)
-            # dL/dsigma = sum(dL/dy * w * x * (-1/sigma^2))
-            # dL/dx = (dL/dy * w)/sigma - sum(dL/dy * w * x) * x / (N * sigma^3)
-            #       = (1/sigma) * [ (dL/dy * w) - x * sum(dL/dy * w * x) / (N * sigma^2) ]
-            #       = rstd * [ (dL/dy * w) - x * rstd^2 * mean(dL/dy * w * x) ] # Wait, sum/N is mean
-            
-            # Let's compute:
-            dy_w = grad_output_flat * w
-            # term2 = (dy_w * x_flat).sum(dim=-1, keepdim=True) / N # mean(dy*w*x)
-            # grad_input = rstd * (dy_w - x_flat * (rstd ** 2) * term2)
-            
-            # Re-derivation:
-            # x_hat = x * rstd
-            # y = x_hat * w
-            # dL/dx = dL/dy * dy/dx
-            # dy/dx = w * dx_hat/dx
-            # dx_hat/dx = rstd * (I - x * x^T * rstd^2 / N) ?? No.
-            # dx_hat_i / dx_j = rstd * (delta_ij - x_i * x_j * rstd^2 / N)
-            
-            # dL/dx_hat = dL/dy * w
-            # dL/dx = rstd * (dL/dx_hat - x * rstd^2 * mean(dL/dx_hat * x))
-            
-            dx_hat = grad_output_flat * w
-            dx_hat_x_mean = (dx_hat * x_flat).mean(dim=-1, keepdim=True)
-            grad_input = rstd * (dx_hat - x_flat * (rstd ** 2) * dx_hat_x_mean)
-            
-            if x.dim() == 3:
-                grad_input = grad_input.reshape(x.shape)
-            
-        return grad_input, grad_w, None
-
-class LayerNormFunction(torch.autograd.Function):
-    @staticmethod
-    def forward(ctx, x, w, b, eps, elementwise_affine):
-        if elementwise_affine:
-            y, mean, rstd = layernorm_param(x, w, b, eps)
-            ctx.save_for_backward(x, w, b, mean, rstd)
-        else:
-            y, mean, rstd = layernorm_noparam(x, eps)
-            ctx.save_for_backward(x, None, None, mean, rstd)
-            
-        ctx.eps = eps
-        ctx.elementwise_affine = elementwise_affine
-        return y
-
-    @staticmethod
-    def backward(ctx, grad_output):
-        x, w, b, mean, rstd = ctx.saved_tensors
-        N = x.shape[-1]
-        
-        if x.dim() == 3:
-            x_flat = x.reshape(-1, N)
-            grad_output_flat = grad_output.reshape(-1, N)
-        else:
-            x_flat = x
-            grad_output_flat = grad_output
-            
-        if w is None:
-             w_eff = 1.0
-        else:
-             w_eff = w
-        
-        # dx calculation
-        # x_hat = (x - mean) * rstd
-        # y = x_hat * w + b
-        # dL/dx_hat = dL/dy * w
-        
-        dy = grad_output_flat
-        dx_hat = dy * w_eff
-        
-        # dL/dvar = sum(dL/dx_hat * (x-mean) * (-0.5) * (var+eps)^(-1.5))
-        #         = -0.5 * rstd^3 * sum(dx_hat * (x-mean))
-        # dL/dmean = sum(dL/dx_hat * (-rstd)) + dL/dvar * (-2/N) * sum(x-mean)
-        # term sum(x-mean) is 0. So second part vanishes.
-        # dL/dmean = -rstd * sum(dx_hat)
-        
-        # dL/dx = dL/dx_hat * rstd + dL/dvar * 2(x-mean)/N + dL/dmean * 1/N
-        #       = dx_hat * rstd + (-0.5 * rstd^3 * sum(dx_hat * (x-mean))) * 2(x-mean)/N + (-rstd * sum(dx_hat))/N
-        #       = rstd * [ dx_hat - mean(dx_hat) - (x-mean)*rstd^2 * mean(dx_hat * (x-mean)) ]
-        
-        x_centered = x_flat - mean.unsqueeze(-1)
-        dx_hat_mean = dx_hat.mean(dim=-1, keepdim=True)
-        dx_hat_x_centered_mean = (dx_hat * x_centered).mean(dim=-1, keepdim=True)
-        
-        grad_input = rstd.unsqueeze(-1) * (dx_hat - dx_hat_mean - x_centered * (rstd.unsqueeze(-1)**2) * dx_hat_x_centered_mean)
-        
-        if x.dim() == 3:
-            grad_input = grad_input.reshape(x.shape)
-
-        grad_w = None
-        grad_b = None
-        
-        if ctx.elementwise_affine:
-            if ctx.needs_input_grad[1]:
-                 x_hat = x_centered * rstd.unsqueeze(-1)
-                 grad_w = (dy * x_hat).sum(dim=0)
-            if ctx.needs_input_grad[2]:
-                 grad_b = dy.sum(dim=0)
-                 
-        return grad_input, grad_w, grad_b, None, None
-
-
-class Int8Linear(nn.Module):
-    def __init__(self, in_features, out_features, bias=True, dtype=torch.bfloat16):
-        super().__init__()
-        self.in_features = in_features
-        self.out_features = out_features
-
-        row_blocks = cdiv(out_features, b=128)
-        col_blocks = cdiv(in_features, b=128)
-        
-        self.register_buffer("int8_weight", torch.empty((out_features, in_features), dtype=torch.int8))
-        self.register_buffer("scale", torch.empty((row_blocks, col_blocks), dtype=torch.float32))
-        if bias:
-            self.register_buffer("bias", torch.empty(out_features, dtype=dtype))
-        else:
-            self.bias = None
-        
-
-    def forward(self, x):
-        return Int8LinearFunction.apply(x, self.int8_weight, self.scale, self.bias)
-
-    @classmethod
-    def from_linear(cls, original_linear: nn.Linear, quantize: bool = True):
-    
-        int8_layer = cls(
-            original_linear.in_features,
-            original_linear.out_features,
-            bias=original_linear.bias is not None,
-            dtype=original_linear.weight.dtype
-        )
-        if quantize:
-            w_data = original_linear.weight.data.cuda()
-            int8_w, scale = int8_quant(w_data)
-
-            int8_layer.int8_weight.copy_(int8_w)
-            int8_layer.scale.copy_(scale)
-            if original_linear.bias is not None:
-                int8_layer.bias.data.copy_(original_linear.bias.data.cuda())
-            
-        return int8_layer
-    
-class FastRMSNorm(nn.Module):
-    def __init__(self, dim: int, eps: float = 1e-5):
-        super().__init__()
-        self.dim = dim
-        self.eps = eps
-        self.register_buffer("weight", torch.ones(dim))
-
-    def forward(self, x):
-        return RMSNormFunction.apply(x.float(), self.weight, self.eps).to(x.dtype)
-
-    @classmethod
-    def from_rmsnorm(cls, original_rmsnorm):
-        rmsnorm_layer = cls(
-            dim=original_rmsnorm.dim,
-            eps=original_rmsnorm.eps
-        )
-        if original_rmsnorm.weight.device != torch.device('meta'):
-            rmsnorm_layer.weight.data.copy_(original_rmsnorm.weight.float().data)
-        return rmsnorm_layer
-    
-class FastLayerNorm(nn.Module):
-
-    def __init__(
-        self,
-        dim: int,
-        eps: float = 1e-5,
-        elementwise_affine: bool = False,
-        bias: bool = True
-    ) :
-        super().__init__()
-        self.dim = dim  # type: ignore[arg-type]
-        self.eps = eps
-        self.elementwise_affine = elementwise_affine
-        if self.elementwise_affine:
-            self.register_buffer("weight", torch.empty(self.dim))
-            if bias:
-                self.register_buffer("bias", torch.empty(self.dim))
-            else:
-                self.bias = None
-        else:
-            self.register_parameter("weight", None)
-            self.register_parameter("bias", None)
-
-    def forward(self, x):
-        return LayerNormFunction.apply(x.float(), self.weight, self.bias, self.eps, self.elementwise_affine).to(x.dtype)
-    
-    @classmethod
-    def from_layernorm(cls, original_layernorm):
-        layernorm_layer = cls(
-            dim=original_layernorm.normalized_shape[0],
-            eps=original_layernorm.eps,
-            elementwise_affine=False if original_layernorm.weight is None else True,
-            bias=original_layernorm.bias is not None
-        )
-        if original_layernorm.weight is not None and original_layernorm.weight.device != torch.device('meta'):
-            layernorm_layer.weight.data.copy_(original_layernorm.weight.data)
-        if original_layernorm.bias is not None and original_layernorm.bias.device != torch.device('meta'):
-            layernorm_layer.bias.data.copy_(original_layernorm.bias.data)
-        return layernorm_layer
@@ -1 +1 @@
-__version__ = "0.2.5"
+__version__ = "0.2.1"
@@ -0,0 +1,63 @@
+import torch
+import sys
+import os
+from tqdm import tqdm
+
+# Local support import
+from .support_flex_sta import get_sliding_tile_attention_mask
+
+# USE OUR NEW PACKAGE!
+from fastvideo_kernel import sliding_tile_attention
+from torch.nn.attention.flex_attention import flex_attention
+
+flex_attention = torch.compile(flex_attention, dynamic=False)
+
+def flex_test(Q, K, V, kernel_size):
+    mask = get_sliding_tile_attention_mask(kernel_size, (6, 8, 8), (18, 48, 80), 0, 'cuda', 0)
+    output = flex_attention(Q, K, V, block_mask=mask)
+    return output
+
+def h100_fwd_kernel_test(Q, K, V, kernel_size):
+    # Using the same parameters as the original test
+    o = sliding_tile_attention(Q, K, V, [kernel_size] * 24, 0, False, '18x48x80')
+    return o
+
+def generate_tensor(shape, mean, std, dtype, device):
+    tensor = torch.randn(shape, dtype=dtype, device=device)
+    magnitude = torch.norm(tensor, dim=-1, keepdim=True)
+    scaled_tensor = tensor * (torch.randn(magnitude.shape, dtype=dtype, device=device) * std + mean) / magnitude
+    return scaled_tensor.contiguous()
+
+def check_correctness(b, h, n, d, causal, mean, std, num_iterations=2):
+    print(f"Running correctness check: batch={b}, heads={h}, seq_len={n}, dim={d}")
+    kernel_size_ls = [(3, 3, 5), (3, 1, 10)]
+    
+    for kernel_size in kernel_size_ls:
+        print(f"Testing kernel_size: {kernel_size}")
+        for xi in tqdm(range(num_iterations)):
+            torch.manual_seed(xi)
+            Q = generate_tensor((b, h, n, d), mean, std, torch.bfloat16, 'cuda')
+            K = generate_tensor((b, h, n, d), mean, std, torch.bfloat16, 'cuda')
+            V = generate_tensor((b, h, n, d), mean, std, torch.bfloat16, 'cuda')
+            
+            tk_o = h100_fwd_kernel_test(Q, K, V, kernel_size)
+            pt_o = flex_test(Q, K, V, kernel_size)
+
+            diff = pt_o - tk_o
+            abs_diff = torch.abs(diff)
+            max_d = torch.max(abs_diff).item()
+            avg_d = torch.sum(abs_diff).item() / (b * h * n * d)
+            
+            if max_d > 0.1:
+                print(f"Warning: Large diff detected! max={max_d}, avg={avg_d}")
+
+    print("\n✅ TEST COMPLETE: New package matches FlexAttention behavior.")
+
+if __name__ == "__main__":
+    b, h, d = 2, 24, 128
+    n = 69120 
+    causal = False
+    mean = 1e-1
+    std = 10
+    
+    check_correctness(b, h, n, d, causal, mean, std, num_iterations=2)
@@ -1,95 +0,0 @@
-import torch
-from .support_flex_sta import get_sliding_tile_attention_mask
-from fastvideo_kernel import sliding_tile_attention
-from torch.nn.attention.flex_attention import flex_attention
-# from flash_attn_interface import flash_attn_func
-from tqdm import tqdm
-
-flex_attention = torch.compile(flex_attention, dynamic=False)
-
-
-def flex_test(Q, K, V, kernel_size):
-    mask = get_sliding_tile_attention_mask(kernel_size, (6, 8, 8), (18, 48, 80), 0, 'cuda', 0)
-    output = flex_attention(Q, K, V, block_mask=mask)
-
-    return output
-
-
-def h100_fwd_kernel_test(Q, K, V, kernel_size):
-    o = sliding_tile_attention(Q, K, V, [kernel_size] * 24, 0, False, '18x48x80')
-    return o
-
-
-def generate_tensor(shape, mean, std, dtype, device):
-    tensor = torch.randn(shape, dtype=dtype, device=device)
-
-    magnitude = torch.norm(tensor, dim=-1, keepdim=True)
-    scaled_tensor = tensor * (torch.randn(magnitude.shape, dtype=dtype, device=device) * std + mean) / magnitude
-
-    return scaled_tensor.contiguous()
-
-
-def check_correctness(b, h, n, d, causal, mean, std, num_iterations=50, error_mode='all'):
-    results = {
-        'TK vs FLEX': {
-            'sum_diff': 0,
-            'sum_abs': 0,
-            'max_diff': 0
-        },
-    }
-    kernel_size_ls = [(3, 3, 5), (3, 1, 10)]
-    from tqdm import tqdm
-    for kernel_size in tqdm(kernel_size_ls):
-        for _ in range(num_iterations):
-            torch.manual_seed(0)
-
-            Q = generate_tensor((b, h, n, d), mean, std, torch.bfloat16, 'cuda')
-            K = generate_tensor((b, h, n, d), mean, std, torch.bfloat16, 'cuda')
-            V = generate_tensor((b, h, n, d), mean, std, torch.bfloat16, 'cuda')
-            tk_o = h100_fwd_kernel_test(Q, K, V, kernel_size)
-            pt_o = flex_test(Q, K, V, kernel_size)
-
-            diff = pt_o - tk_o
-            abs_diff = torch.abs(diff)
-            results['TK vs FLEX']['sum_diff'] += torch.sum(abs_diff).item()
-            results['TK vs FLEX']['max_diff'] = max(results['TK vs FLEX']['max_diff'], torch.max(abs_diff).item())
-
-            torch.cuda.empty_cache()
-        print("kernel_size", kernel_size)
-        print("max_diff", torch.max(abs_diff).item())
-        print(
-            "avg_diff",
-            torch.sum(abs_diff).item() / (b * h * n * d *
-                                          (1 if error_mode == 'output' else 3 if error_mode == 'backward' else 4)))
-
-    total_elements = b * h * n * d * num_iterations * (1 if error_mode == 'output' else
-                                                       3 if error_mode == 'backward' else 4) * len(kernel_size_ls)
-    for name, data in results.items():
-        avg_diff = data['sum_diff'] / total_elements
-        max_diff = data['max_diff']
-        results[name] = {'avg_diff': avg_diff, 'max_diff': max_diff}
-
-    return results
-
-
-# Example usage
-def test_sliding_tile_attention():
-    if not torch.cuda.is_available():
-        return
-
-    b, h, d = 2, 24, 128
-    n = 69120  # Sequence length
-    causal = False
-    mean = 1e-1
-    std = 10
-
-    # Run correctness check directly
-    results = check_correctness(b, h, n, d, causal, mean, std, error_mode='output')
-    assert results['TK vs FLEX']['avg_diff'] < 3e-6, f"Average difference: {results['TK vs FLEX']['avg_diff']} is too large"
-    assert results['TK vs FLEX']['max_diff'] < 4e-2, f"Maximum difference: {results['TK vs FLEX']['max_diff']} is too large"
-    print(f"Average difference: {results['TK vs FLEX']['avg_diff']}")
-    print(f"Maximum difference: {results['TK vs FLEX']['max_diff']}")
-
-if __name__ == "__main__":
-    test_sliding_tile_attention()
-
@@ -1,294 +0,0 @@
-import torch
-import pytest
-import sys
-import os
-
-# Ensure local package is imported
-# sys.path.insert(0, os.path.abspath(os.path.join(os.path.dirname(__file__), "../python")))
-
-try:
-    from fastvideo_kernel._C import fastvideo_kernel_ops
-except ImportError:
-    fastvideo_kernel_ops = None
-
-from fastvideo_kernel import turbodiffusion_ops
-
-# Helper for RMS Norm reference
-def rms_norm_ref(x, w, eps=1e-6):
-    dtype = x.dtype
-    x = x.float()
-    variance = x.pow(2).mean(-1, keepdim=True)
-    x = x * torch.rsqrt(variance + eps)
-    return (x * w.float()).to(dtype)
-
-@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA not available")
-class TestTurboDiffusion:
-    
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    @pytest.mark.parametrize("shape", [(16, 128), (32, 256), (1, 1024)])
-    def test_quant_correctness(self, dtype, shape):
-        if turbodiffusion_ops.quant_cuda is None:
-            pytest.skip("quant_cuda not available")
-            
-        x = torch.randn(shape, dtype=dtype, device="cuda")
-        x_q, x_scale = turbodiffusion_ops.int8_quant(x)
-        
-        assert x_q.dtype == torch.int8
-        assert x_scale.dtype == torch.float32
-        
-        # Simple check: dequantize and compute error
-        # Note: The quantization scheme details matter here (per block? per tensor?).
-        # Looking at quant.cu, it seems to be block-based but the output scale shape isn't immediately obvious from python signature 
-        # without looking at C++ code deeper.
-        # But let's check shapes at least.
-        
-        # If we can't easily dequantize without knowing block size logic in python, 
-        # checking that it runs and produces valid shapes is a good start.
-        assert x_q.shape == shape
-        # x_scale shape depends on block size, usually smaller than x
-        
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    def test_gemm_correctness(self, dtype):
-        if turbodiffusion_ops.gemm_cuda is None:
-            pytest.skip("gemm_cuda not available")
-            
-        M, N, K = 32, 64, 128
-        x = torch.randn(M, K, dtype=dtype, device="cuda")
-        
-        # Create weights
-        # For simplicity in testing, let's create random int8 weights and scales
-        w_q = torch.randint(-127, 127, (N, K), dtype=torch.int8, device="cuda")
-        
-        # Scale shape: The Int8Linear class uses:
-        # row_blocks = cdiv(out_features, b=128)
-        # col_blocks = cdiv(in_features, b=128)
-        # scale shape: (row_blocks, col_blocks)
-        
-        row_blocks = (N + 127) // 128
-        col_blocks = (K + 127) // 128
-        w_s = torch.randn(row_blocks, col_blocks, dtype=torch.float32, device="cuda").abs()
-        
-        # Run int8_linear
-        output = turbodiffusion_ops.int8_linear(x, w_q, w_s)
-        
-        assert output.shape == (M, N)
-        assert output.dtype == dtype
-
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    def test_gemm_backward(self, dtype):
-        if turbodiffusion_ops.gemm_cuda is None:
-            pytest.skip("gemm_cuda not available")
-            
-        M, N, K = 32, 64, 128
-        x = torch.randn(M, K, dtype=dtype, device="cuda", requires_grad=True)
-        
-        # Weights (frozen)
-        w_q = torch.randint(-127, 127, (N, K), dtype=torch.int8, device="cuda")
-        row_blocks = (N + 127) // 128
-        col_blocks = (K + 127) // 128
-        w_s = torch.randn(row_blocks, col_blocks, dtype=torch.float32, device="cuda").abs()
-        
-        bias = torch.randn(N, dtype=dtype, device="cuda", requires_grad=True)
-        
-        # Use Int8LinearFunction
-        output = turbodiffusion_ops.Int8LinearFunction.apply(x, w_q, w_s, bias)
-        loss = output.sum()
-        loss.backward()
-        
-        assert x.grad is not None
-        assert bias.grad is not None
-        assert x.grad.shape == (M, K)
-        
-        # Check correctness against dequantized weight
-        w_float = turbodiffusion_ops.dequantize_weight(w_q, w_s, dtype)
-        
-        # Reference
-        x_ref = x.detach().clone().requires_grad_()
-        bias_ref = bias.detach().clone().requires_grad_()
-        
-        # Note: Int8Linear forward uses quantized input, so output won't match exactly reference with float input.
-        # But backward gradients should be consistent with the logic we implemented:
-        # grad_input = grad_output @ w_float
-        
-        # Let's verify our backward logic matches standard matmul backward with dequantized weights
-        # We manually compute expected gradient given grad_output = ones
-        grad_output = torch.ones_like(output)
-        expected_x_grad = grad_output @ w_float
-        expected_bias_grad = grad_output.sum(0)
-        
-        torch.testing.assert_close(x.grad, expected_x_grad, atol=1e-2, rtol=1e-2)
-        torch.testing.assert_close(bias.grad, expected_bias_grad, atol=1e-2, rtol=1e-2)
-
-        
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    @pytest.mark.parametrize("shape", [(2, 16, 128), (4, 32, 256)])
-    def test_rms_norm_triton(self, dtype, shape):
-        x = torch.randn(shape, dtype=dtype, device="cuda")
-        dim = shape[-1]
-        w = torch.randn(dim, dtype=dtype, device="cuda")
-        eps = 1e-5
-        
-        # Triton implementation
-        # Note: rmsnorm returns tuple now
-        res = turbodiffusion_ops.rmsnorm(x, w, eps)
-        if isinstance(res, tuple):
-            out_triton = res[0]
-        else:
-            out_triton = res
-        
-        # Reference
-        out_ref = rms_norm_ref(x, w, eps)
-        
-        torch.testing.assert_close(out_triton, out_ref, atol=1e-2, rtol=1e-2)
-
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    @pytest.mark.parametrize("shape", [(2, 16, 128)])
-    def test_rms_norm_backward(self, dtype, shape):
-        x = torch.randn(shape, dtype=dtype, device="cuda", requires_grad=True)
-        dim = shape[-1]
-        w = torch.randn(dim, dtype=dtype, device="cuda", requires_grad=True)
-        eps = 1e-5
-        
-        # Forward via Function
-        y = turbodiffusion_ops.RMSNormFunction.apply(x, w, eps)
-        loss = y.sum()
-        loss.backward()
-        
-        x_grad = x.grad
-        w_grad = w.grad
-        
-        # Reference
-        x_ref = x.detach().clone().requires_grad_()
-        w_ref = w.detach().clone().requires_grad_()
-        # Custom RMSNorm ref in pytorch
-        def rms_norm_ref_grad(x, w, eps):
-             x_float = x.float()
-             var = x_float.pow(2).mean(-1, keepdim=True)
-             rstd = torch.rsqrt(var + eps)
-             return (x_float * rstd * w.float()).to(x.dtype)
-             
-        y_ref = rms_norm_ref_grad(x_ref, w_ref, eps)
-        loss_ref = y_ref.sum()
-        loss_ref.backward()
-        
-        torch.testing.assert_close(x_grad, x_ref.grad, atol=1e-2, rtol=1e-2)
-        torch.testing.assert_close(w_grad, w_ref.grad, atol=2e-2, rtol=2e-2)
-
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    def test_rms_norm_cuda(self, dtype):
-        if fastvideo_kernel_ops is None or not hasattr(fastvideo_kernel_ops, "rms_norm_cuda"):
-            pytest.skip("rms_norm_cuda not available")
-            
-        shape = (16, 128)
-        x = torch.randn(shape, dtype=dtype, device="cuda")
-        dim = shape[-1]
-        w = torch.randn(dim, dtype=dtype, device="cuda")
-        eps = 1e-5
-        
-        # C++ implementation
-        # Signature: rms_norm_cuda(Input, eps, Weight, Output) -> Output
-        out_cuda = torch.empty_like(x)
-        fastvideo_kernel_ops.rms_norm_cuda(x, eps, w, out_cuda)
-        
-        # Reference
-        out_ref = rms_norm_ref(x, w, eps)
-        
-        torch.testing.assert_close(out_cuda, out_ref, atol=1e-2, rtol=1e-2)
-
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    @pytest.mark.parametrize("shape", [(2, 16, 128), (4, 32, 256)])
-    def test_layer_norm_triton(self, dtype, shape):
-        x = torch.randn(shape, dtype=dtype, device="cuda")
-        dim = shape[-1]
-        eps = 1e-5
-        
-        # With affine
-        w = torch.randn(dim, dtype=dtype, device="cuda")
-        b = torch.randn(dim, dtype=dtype, device="cuda")
-        
-        # Triton implementation
-        # Note: layernorm returns tuple now
-        res = turbodiffusion_ops.layernorm(x, w, b, eps, elementwise_affine=True)
-        if isinstance(res, tuple):
-            out_triton = res[0]
-        else:
-            out_triton = res
-        out_triton = out_triton.to(dtype)
-        
-        # Reference
-        ln = torch.nn.LayerNorm(dim, eps=eps, elementwise_affine=True, dtype=dtype).cuda()
-        ln.weight.data.copy_(w)
-        ln.bias.data.copy_(b)
-        out_ref = ln(x)
-        
-        torch.testing.assert_close(out_triton, out_ref, atol=1e-2, rtol=1e-2)
-        
-        # Without affine
-        res_no_affine = turbodiffusion_ops.layernorm(x, None, None, eps, elementwise_affine=False)
-        if isinstance(res_no_affine, tuple):
-            out_triton_no_affine = res_no_affine[0]
-        else:
-            out_triton_no_affine = res_no_affine
-        out_triton_no_affine = out_triton_no_affine.to(dtype)
-        ln_no_affine = torch.nn.LayerNorm(dim, eps=eps, elementwise_affine=False, dtype=dtype).cuda()
-        out_ref_no_affine = ln_no_affine(x)
-        
-        torch.testing.assert_close(out_triton_no_affine, out_ref_no_affine, atol=1e-2, rtol=1e-2)
-
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    @pytest.mark.parametrize("shape", [(2, 16, 128)])
-    def test_layer_norm_backward(self, dtype, shape):
-        x = torch.randn(shape, dtype=dtype, device="cuda", requires_grad=True)
-        dim = shape[-1]
-        eps = 1e-5
-        w = torch.randn(dim, dtype=dtype, device="cuda", requires_grad=True)
-        b = torch.randn(dim, dtype=dtype, device="cuda", requires_grad=True)
-        
-        # Triton implementation via Function
-        y = turbodiffusion_ops.LayerNormFunction.apply(x, w, b, eps, True)
-        loss = y.sum()
-        loss.backward()
-        
-        x_grad = x.grad
-        w_grad = w.grad
-        b_grad = b.grad
-        
-        # Reference
-        x_ref = x.detach().clone().requires_grad_()
-        ln = torch.nn.LayerNorm(dim, eps=eps, elementwise_affine=True, dtype=dtype).cuda()
-        ln.weight.data.copy_(w.detach())
-        ln.bias.data.copy_(b.detach())
-        
-        y_ref = ln(x_ref)
-        loss_ref = y_ref.sum()
-        loss_ref.backward()
-        
-        torch.testing.assert_close(x_grad, x_ref.grad, atol=1e-2, rtol=1e-2)
-        torch.testing.assert_close(w_grad, ln.weight.grad, atol=1e-2, rtol=1e-2)
-        torch.testing.assert_close(b_grad, ln.bias.grad, atol=1e-2, rtol=1e-2)
-
-
-    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
-    def test_layer_norm_cuda(self, dtype):
-        if fastvideo_kernel_ops is None or not hasattr(fastvideo_kernel_ops, "layer_norm_cuda"):
-            pytest.skip("layer_norm_cuda not available")
-            
-        shape = (16, 128)
-        x = torch.randn(shape, dtype=dtype, device="cuda")
-        dim = shape[-1]
-        eps = 1e-5
-        w = torch.randn(dim, dtype=dtype, device="cuda")
-        b = torch.randn(dim, dtype=dtype, device="cuda")
-        
-        # C++ implementation
-        # Signature: layer_norm_cuda(Input, eps, W, B, Output) -> Output
-        out_cuda = torch.empty_like(x)
-        fastvideo_kernel_ops.layer_norm_cuda(x, eps, w, b, out_cuda)
-        
-        # Reference
-        ln = torch.nn.LayerNorm(dim, eps=eps, elementwise_affine=True, dtype=dtype).cuda()
-        ln.weight.data.copy_(w)
-        ln.bias.data.copy_(b)
-        out_ref = ln(x)
-        
-        torch.testing.assert_close(out_cuda, out_ref, atol=1e-2, rtol=1e-2)
@@ -1,257 +0,0 @@
-import torch
-import pytest
-import sys
-import os
-import numpy as np
-from tqdm import tqdm
-
-from .utils import generate_block_sparse_mask_for_function, create_full_mask_from_block_mask
-# Use installed package
-# from fastvideo_kernel import video_sparse_attn as block_sparse_attn
-
-BLOCK_M = 64
-BLOCK_N = 64
-
-def pytorch_test(Q, K, V, block_sparse_mask, dO):
-    q_ = Q.clone().float().requires_grad_()
-    k_ = K.clone().float().requires_grad_()
-    v_ = V.clone().float().requires_grad_()
-
-    QK = torch.matmul(q_, k_.transpose(-2, -1))
-    QK /= (q_.size(-1) ** 0.5)
-    QK = QK.masked_fill(~block_sparse_mask.unsqueeze(0), float('-inf'))
-
-    QK = torch.nn.functional.softmax(QK, dim=-1)
-    output = torch.matmul(QK, v_)
-
-    dO_ = dO
-    output.backward(dO_)
-    return (
-        output.to(torch.bfloat16),
-        q_.grad.to(torch.bfloat16),
-        k_.grad.to(torch.bfloat16),
-        v_.grad.to(torch.bfloat16),
-    )
-
-
-def block_sparse_kernel_test(Q, K, V, block_sparse_mask, variable_block_sizes, q_non_pad_index, kv_non_pad_index, q_num_blocks, kv_num_blocks, dO):
-    Q = Q.detach().requires_grad_()
-    K = K.detach().requires_grad_()
-    V = V.detach().requires_grad_()
-    
-    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 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
-    )
-
-    output = output_padded[:, :, q_non_pad_index, :]
-    output.backward(dO)
-    return output, Q.grad, K.grad, V.grad
-
-
-def get_non_pad_index(
-    vid_len: torch.LongTensor,
-    n_win: int,
-    win_size: int,
-):
-    device = vid_len.device
-    starts_pad = torch.arange(n_win, device=device) * win_size
-    index_pad = starts_pad[:, None] + torch.arange(win_size, device=device)[None, :]
-    index_mask = torch.arange(win_size, device=device)[None, :] < vid_len[:, None]
-
-    return index_pad[index_mask]
-
-def generate_tensor(shape, dtype, device):
-    tensor = torch.randn(shape, dtype=dtype, device=device)
-    return tensor
-
-def generate_variable_block_sizes(num_blocks, min_size=16, max_size=64, device="cuda"):
-    return torch.randint(min_size, max_size + 1, (num_blocks,), device=device, dtype=torch.int32)
-
-
-def vsa_pad(x, non_pad_index, num_blocks, block_size):
-    padded_x = torch.zeros((1, x.shape[1], num_blocks * BLOCK_M, x.shape[3]), device=x.device, dtype=x.dtype)
-    padded_x[:, :, non_pad_index, :] = x
-    return padded_x
-
-def check_correctness(h, d, num_blocks, k,  num_iterations=20, error_mode='all'):
-    results = {
-        'gO': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-        'gQ': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-        'gK': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-        'gV': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-    }
-
-    device = "cuda" if torch.cuda.is_available() else "cpu"
-    variable_block_sizes = generate_variable_block_sizes(num_blocks, device=device)
-    S = int(variable_block_sizes.sum().item())
-    padded_S = num_blocks * BLOCK_M
-    non_pad_index = get_non_pad_index(variable_block_sizes, num_blocks, BLOCK_M)
-    block_mask = generate_block_sparse_mask_for_function(h, num_blocks, num_blocks, k, device)
-    full_mask = create_full_mask_from_block_mask(block_mask, variable_block_sizes, variable_block_sizes, device)
-    for _ in range(num_iterations):
-        Q = generate_tensor((1, h, S, d),  torch.bfloat16, device)
-        K = generate_tensor((1, h, S, d), torch.bfloat16, device)
-        V = generate_tensor((1, h, S, d),  torch.bfloat16, device)
-        dO = generate_tensor((1, h, S, d), torch.bfloat16, device)
-        
-        # print(Q.shape, K.shape, V.shape, dO.shape)
-
-        # dO_padded = torch.zeros_like(dO_padded)
-        # dO_padded[:, :, non_pad_index, :] = dO
-
-        pt_o, pt_qg, pt_kg, pt_vg = pytorch_test(Q, K, V, full_mask, dO)
-        bs_o, bs_qg, bs_kg, bs_vg = block_sparse_kernel_test(Q, K, V, block_mask.unsqueeze(0), variable_block_sizes, non_pad_index, non_pad_index, num_blocks, num_blocks, dO)
-        for name, (pt, bs) in zip(['gQ', 'gK', 'gV', 'gO'], [(pt_qg, bs_qg), (pt_kg, bs_kg), (pt_vg, bs_vg), (pt_o, bs_o)]):
-            if bs is not None:
-                diff = pt - bs
-                abs_diff = torch.abs(diff)
-                results[name]['sum_diff'] += torch.sum(abs_diff).item()
-                results[name]['sum_abs'] += torch.sum(torch.abs(pt)).item()
-                rel_max_diff = torch.max(abs_diff) / torch.mean(torch.abs(pt))
-                results[name]['max_diff'] = max(results[name]['max_diff'], rel_max_diff.item())
-        if torch.cuda.is_available():
-            torch.cuda.empty_cache()
-
-    total_elements = h * S * d * num_iterations
-    for name, data in results.items():
-        avg_diff = data['sum_diff'] / total_elements
-        max_diff = data['max_diff']
-        results[name] = {'avg_diff': avg_diff, 'max_diff': max_diff}
-
-    return results
-
-def check_correctness_qkdiff(h, d, num_q_blocks, num_kv_blocks, k, num_iterations=20, error_mode='all'):
-    results = {
-        'gO': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-        'gQ': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-        'gK': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-        'gV': {'sum_diff': 0.0, 'sum_abs': 0.0, 'max_diff': 0.0},
-    }
-
-    device = "cuda" if torch.cuda.is_available() else "cpu"
-    q_variable_block_sizes = generate_variable_block_sizes(num_q_blocks, device=device)
-    kv_variable_block_sizes = generate_variable_block_sizes(num_kv_blocks, device=device)
-    
-    S_q = int(q_variable_block_sizes.sum().item())
-    S_kv = int(kv_variable_block_sizes.sum().item())
-    
-    q_non_pad_index = get_non_pad_index(q_variable_block_sizes, num_q_blocks, BLOCK_M)
-    kv_non_pad_index = get_non_pad_index(kv_variable_block_sizes, num_kv_blocks, BLOCK_M)
-    
-    block_mask = generate_block_sparse_mask_for_function(h, num_q_blocks, num_kv_blocks, k, device)
-    full_mask = create_full_mask_from_block_mask(block_mask, q_variable_block_sizes, kv_variable_block_sizes, device)
-    
-    for _ in range(num_iterations):
-        Q = generate_tensor((1, h, S_q, d),  torch.bfloat16, device)
-        K = generate_tensor((1, h, S_kv, d), torch.bfloat16, device)
-        V = generate_tensor((1, h, S_kv, d),  torch.bfloat16, device)
-        dO = generate_tensor((1, h, S_q, d), torch.bfloat16, device)
-        
-        # print(Q.shape, K.shape, V.shape, dO.shape)
-
-        pt_o, pt_qg, pt_kg, pt_vg = pytorch_test(Q, K, V, full_mask, dO)
-        bs_o, bs_qg, bs_kg, bs_vg = block_sparse_kernel_test(Q, K, V, block_mask.unsqueeze(0), kv_variable_block_sizes, q_non_pad_index, kv_non_pad_index, num_q_blocks, num_kv_blocks, dO)
-        
-        for name, (pt, bs) in zip(['gQ', 'gK', 'gV', 'gO'], [(pt_qg, bs_qg), (pt_kg, bs_kg), (pt_vg, bs_vg), (pt_o, bs_o)]):
-            if bs is not None:
-                diff = pt - bs
-                abs_diff = torch.abs(diff)
-                results[name]['sum_diff'] += torch.sum(abs_diff).item()
-                results[name]['sum_abs'] += torch.sum(torch.abs(pt)).item()
-                rel_max_diff = torch.max(abs_diff) / torch.mean(torch.abs(pt))
-                results[name]['max_diff'] = max(results[name]['max_diff'], rel_max_diff.item())
-        if torch.cuda.is_available():
-            torch.cuda.empty_cache()
-
-    total_elements_q = h * S_q * d * num_iterations
-    total_elements_kv = h * S_kv * d * num_iterations
-    
-    for name, data in results.items():
-        total_elements = total_elements_q if name in ['gQ', 'gO'] else total_elements_kv
-        avg_diff = data['sum_diff'] / total_elements
-        max_diff = data['max_diff']
-        results[name] = {'avg_diff': avg_diff, 'max_diff': max_diff}
-
-    return results
-
-def generate_error_graphs(h, d, error_mode='all'):
-    test_configs = [
-        {"num_blocks": 16, "k": 2, "description": "Small sequence"},
-        {"num_blocks": 32, "k": 4, "description": "Medium sequence"},
-        {"num_blocks": 53, "k": 6, "description": "Large sequence"},
-    ]
-    
-    print(f"\nError Analysis for h={h}, d={d}, mode={error_mode}")
-    print("=" * 150)
-    print(f"{'Config':<20} {'Blocks':<8} {'K':<4} "
-          f"{'gQ Avg':<12} {'Rel gQ Max':<12} "
-          f"{'gK Avg':<12} {'Rel gK Max':<12} "
-          f"{'gV Avg':<12} {'Rel gV Max':<12} "
-          f"{'gO Avg':<12} {'Rel gO Max':<12}")
-    print("-" * 150)
-
-    for config in test_configs:
-        num_blocks = config["num_blocks"]
-        k = config["k"]
-        description = config["description"]
-        results = check_correctness(h, d, num_blocks, k, error_mode=error_mode)
-        print(f"{description:<20} {num_blocks:<8} {k:<4} "
-              f"{results['gQ']['avg_diff']:<12.6e} {results['gQ']['max_diff']:<12.6e} "
-              f"{results['gK']['avg_diff']:<12.6e} {results['gK']['max_diff']:<12.6e} "
-              f"{results['gV']['avg_diff']:<12.6e} {results['gV']['max_diff']:<12.6e} "
-              f"{results['gO']['avg_diff']:<12.6e} {results['gO']['max_diff']:<12.6e}")
-
-    print("-" * 150)
-
-def generate_error_graphs_qkdiff(h, d, error_mode='all'):
-    test_configs = [
-        {"num_q_blocks": 16, "num_kv_blocks": 32, "k": 2, "description": "Small Q, Med KV"},
-        {"num_q_blocks": 32, "num_kv_blocks": 16, "k": 4, "description": "Med Q, Small KV"},
-        {"num_q_blocks": 53, "num_kv_blocks": 32, "k": 6, "description": "Large Q, Med KV"},
-        {"num_q_blocks": 16, "num_kv_blocks": 48, "k": 2, "description": "Small Q, Large KV"},
-        {"num_q_blocks": 48, "num_kv_blocks": 16, "k": 2, "description": "Large Q, Small KV"},
-    ]
-    
-    print(f"\nError Analysis (QK Diff) for h={h}, d={d}, mode={error_mode}")
-    print("=" * 150)
-    print(f"{'Config':<20} {'Q Blks':<8} {'KV Blks':<8} {'K':<4} "
-          f"{'gQ Avg':<12} {'Rel gQ Max':<12} "
-          f"{'gK Avg':<12} {'Rel gK Max':<12} "
-          f"{'gV Avg':<12} {'Rel gV Max':<12} "
-          f"{'gO Avg':<12} {'Rel gO Max':<12}")
-    print("-" * 150)
-
-    for config in test_configs:
-        num_q_blocks = config["num_q_blocks"]
-        num_kv_blocks = config["num_kv_blocks"]
-        k = config["k"]
-        description = config["description"]
-        results = check_correctness_qkdiff(h, d, num_q_blocks, num_kv_blocks, k, error_mode=error_mode)
-        print(f"{description:<20} {num_q_blocks:<8} {num_kv_blocks:<8} {k:<4} "
-              f"{results['gQ']['avg_diff']:<12.6e} {results['gQ']['max_diff']:<12.6e} "
-              f"{results['gK']['avg_diff']:<12.6e} {results['gK']['max_diff']:<12.6e} "
-              f"{results['gV']['avg_diff']:<12.6e} {results['gV']['max_diff']:<12.6e} "
-              f"{results['gO']['avg_diff']:<12.6e} {results['gO']['max_diff']:<12.6e}")
-
-    print("-" * 150)
-
-def test_video_sparse_attention_backward():
-    if not torch.cuda.is_available():
-        return
-
-    h, d = 16, 128
-    print("Block Sparse Attention with Variable Block Sizes Analysis")
-    print("=" * 60)
-    for mode in ['backward']:
-        generate_error_graphs(h, d, error_mode=mode)
-        generate_error_graphs_qkdiff(h, d, error_mode=mode)
-    print("\nAnalysis completed for all modes.")
-
-if __name__ == "__main__":
-    test_video_sparse_attention_backward()
-
@@ -1,234 +0,0 @@
-import os
-import sys
-from typing import Tuple
-
-import torch
-import pytest
-
-from .utils import (
-    generate_block_sparse_mask_for_function,
-    create_full_mask_from_block_mask,
-)
-from .test_vsa import BLOCK_M  # Import from local test_vsa
-from . import test_vsa as ref
-
-def pytorch_forward(
-    Q: torch.Tensor,
-    K: torch.Tensor,
-    V: torch.Tensor,
-    block_sparse_mask: torch.Tensor,
-) -> torch.Tensor:
-    """
-    Dense PyTorch reference forward:
-    - Q: [1, h, S_q, d]
-    - K,V: [1, h, S_kv, d]
-    - block_sparse_mask: [h, S_q, S_kv] bool
-    """
-    q = Q.clone().float()
-    k = K.clone().float()
-    v = V.clone().float()
-
-    attn = torch.matmul(q, k.transpose(-2, -1))  # [1, h, S_q, S_kv]
-    attn = attn / (q.size(-1) ** 0.5)
-    attn = attn.masked_fill(~block_sparse_mask.unsqueeze(0), float("-inf"))
-    attn = torch.nn.functional.softmax(attn, dim=-1)
-    out = torch.matmul(attn, v)  # [1, h, S_q, d]
-    return out.to(torch.bfloat16)
-
-
-def block_sparse_forward_test(
-    Q: torch.Tensor,
-    K: torch.Tensor,
-    V: torch.Tensor,
-    block_sparse_mask: torch.Tensor,
-    variable_block_sizes: torch.Tensor,
-    q_non_pad_index: torch.Tensor,
-    kv_non_pad_index: torch.Tensor,
-    q_num_blocks: int,
-    kv_num_blocks: int,
-) -> torch.Tensor:
-    """
-    Forward-only wrapper
-    """
-    Q = Q.detach()
-    K = K.detach()
-    V = V.detach()
-
-    q_padded = ref.vsa_pad(Q, q_non_pad_index, q_num_blocks, BLOCK_M)
-    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 autograd-enabled wrapper (internally dispatches SM90 C++ vs Triton)
-    from fastvideo_kernel.block_sparse_attn import block_sparse_attn
-    try:
-        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, :]
-    return out
-
-
-def run_forward_equal_qk(
-    h: int = 16,
-    d: int = 128,
-    num_blocks: int = 16,
-    k: int = 2,
-    num_iterations: int = 5,
-) -> Tuple[float, float]:
-    """
-    Forward-only correctness test for the case S_q == S_kv.
-    Mirrors `check_correctness` but only compares forward outputs.
-    """
-    assert torch.cuda.is_available(), "VSA kernels require CUDA"
-    device = "cuda"
-
-    variable_block_sizes = ref.generate_variable_block_sizes(
-        num_blocks, device=device
-    )
-    S = int(variable_block_sizes.sum().item())
-    non_pad_index = ref.get_non_pad_index(
-        variable_block_sizes, num_blocks, BLOCK_M
-    )
-
-    block_mask = generate_block_sparse_mask_for_function(
-        h, num_blocks, num_blocks, k, device
-    )
-    full_mask = create_full_mask_from_block_mask(
-        block_mask, variable_block_sizes, variable_block_sizes, device
-    )
-    print(f"[qkequal] h: {h}, d: {d}, num_blocks: {num_blocks}, k: {k}")
-    print(f"[qkequal] variable_block_sizes: {variable_block_sizes}, non_pad_index: {non_pad_index.shape}, block_mask: {block_mask.shape}, full_mask: {full_mask.shape}")
-    sum_diff = 0.0
-    sum_abs = 0.0
-    max_rel_diff = 0.0
-
-    for i in range(num_iterations):
-        Q = ref.generate_tensor((1, h, S, d), torch.bfloat16, device)
-        K = ref.generate_tensor((1, h, S, d), torch.bfloat16, device)
-        V = ref.generate_tensor((1, h, S, d), torch.bfloat16, device)
-
-        if i == 0: print(f"[qkequal] Q: {Q.shape}, K: {K.shape}, V: {V.shape}, full_mask: {full_mask.shape}")
-        if i == 0: print(f"[qkequal] block_mask: {block_mask.shape}")
-
-        pt_o = pytorch_forward(Q, K, V, full_mask)
-        bs_o = block_sparse_forward_test(
-            Q,
-            K,
-            V,
-            block_mask.unsqueeze(0),
-            variable_block_sizes,
-            non_pad_index,
-            non_pad_index,
-            num_blocks,
-            num_blocks,
-        )
-
-        diff = (pt_o - bs_o).abs()
-        sum_diff += diff.sum().item()
-        sum_abs += pt_o.abs().sum().item()
-        rel_max = diff.max() / (pt_o.abs().mean() + 1e-6)
-        max_rel_diff = max(max_rel_diff, rel_max.item())
-
-    total_elems = h * S * d * num_iterations
-    avg_abs_err = sum_diff / total_elems
-    return avg_abs_err, max_rel_diff
-
-
-def run_forward_qk_diff(
-    h: int = 16,
-    d: int = 128,
-    num_q_blocks: int = 16,
-    num_kv_blocks: int = 32,
-    k: int = 2,
-    num_iterations: int = 5,
-) -> Tuple[float, float]:
-    """
-    Forward-only correctness test for the case S_q != S_kv.
-    """
-    assert torch.cuda.is_available(), "VSA kernels require CUDA"
-
-    device = "cuda"
-
-    q_variable_block_sizes = ref.generate_variable_block_sizes(
-        num_q_blocks, device=device
-    )
-    kv_variable_block_sizes = ref.generate_variable_block_sizes(
-        num_kv_blocks, device=device
-    )
-
-    S_q = int(q_variable_block_sizes.sum().item())
-    S_kv = int(kv_variable_block_sizes.sum().item())
-
-    q_non_pad_index = ref.get_non_pad_index(
-        q_variable_block_sizes, num_q_blocks, BLOCK_M
-    )
-    kv_non_pad_index = ref.get_non_pad_index(
-        kv_variable_block_sizes, num_kv_blocks, BLOCK_M
-    )
-
-    block_mask = generate_block_sparse_mask_for_function(
-        h, num_q_blocks, num_kv_blocks, k, device
-    )
-    full_mask = create_full_mask_from_block_mask(
-        block_mask, q_variable_block_sizes, kv_variable_block_sizes, device
-    )
-
-    sum_diff = 0.0
-    sum_abs = 0.0
-    max_rel_diff = 0.0
-
-    for _ in range(num_iterations):
-        Q = ref.generate_tensor((1, h, S_q, d), torch.bfloat16, device)
-        K = ref.generate_tensor((1, h, S_kv, d), torch.bfloat16, device)
-        V = ref.generate_tensor((1, h, S_kv, d), torch.bfloat16, device)
-
-        pt_o = pytorch_forward(Q, K, V, full_mask)
-        bs_o = block_sparse_forward_test(
-            Q,
-            K,
-            V,
-            block_mask.unsqueeze(0),
-            kv_variable_block_sizes,
-            q_non_pad_index,
-            kv_non_pad_index,
-            num_q_blocks,
-            num_kv_blocks,
-        )
-
-        diff = (pt_o - bs_o).abs()
-        sum_diff += diff.sum().item()
-        sum_abs += pt_o.abs().sum().item()
-        rel_max = diff.max() / (pt_o.abs().mean() + 1e-6)
-        max_rel_diff = max(max_rel_diff, rel_max.item())
-
-    total_elems = h * S_q * d * num_iterations
-    avg_abs_err = sum_diff / total_elems
-    return avg_abs_err, max_rel_diff
-
-
-def test_video_sparse_attention_forward():
-    if not torch.cuda.is_available():
-        return
-
-    h, d = 16, 128
-    print("Forward Block Sparse Attention Check (QK Equal)")
-    print("=" * 80)
-    avg_err_eq, max_rel_eq = run_forward_equal_qk(h, d, num_blocks=32, k=2)
-    print(f"QK equal: avg |ΔO| = {avg_err_eq:.6e}, max rel ΔO = {max_rel_eq:.6e}")
-
-    print("\nForward Block Sparse Attention Check (QK Different)")
-    print("=" * 80)
-    avg_err_diff, max_rel_diff = run_forward_qk_diff(
-        h, d, num_q_blocks=32, num_kv_blocks=48, k=2
-    )
-    print(
-        f"QK diff:  avg |ΔO| = {avg_err_diff:.6e}, max rel ΔO = {max_rel_diff:.6e}"
-    )
-
-if __name__ == "__main__":
-    test_video_sparse_attention_forward()
-
--- a/Show More
+++ b/Show More