wan2.2 training doesn't crash

Change Wan DiT to have 0 numerical diff with SF's Wan
fix rotary embedding
2025-09-20 05:08:58 +00:00 · 2025-09-18 22:50:07 +00:00 · 2025-09-13 08:31:53 +00:00 · 2025-09-10 11:12:48 +00:00 · 2025-09-10 08:57:54 +00:00 · 2025-09-09 23:30:55 +00:00
531 changed files with 8520 additions and 48252 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 60m .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
@@ -104,18 +104,6 @@ steps:
                  - TEST_TYPE=distillation_dmd
                agents:
                  queue: "default"
-            - path:
-                - "fastvideo/training/*self_forcing_distillation_pipeline.py"
-                - "fastvideo/tests/training/self-forcing/**"
-                - "pyproject.toml"
-                - "docker/Dockerfile.python3.12"
-              config:
-                command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Self-Forcing Tests"
-                env:
-                  - TEST_TYPE=self_forcing
-                agents:
-                  queue: "default"
            - path:
                - "fastvideo/**"
                - "pyproject.toml"
@@ -129,7 +117,11 @@ steps:
                  queue: "default"
            - path:
                - "fastvideo/**"
-                - "fastvideo-kernel/**"
+                - "csrc/attn/video_sparse_attn/**"
+                - "csrc/attn/video_sparse_attn/tk/**"
+                - "csrc/attn/video_sparse_attn/setup.py"
+                - "csrc/attn/video_sparse_attn/config_vsa.py"
+                - "csrc/attn/video_sparse_attn/vsa.cpp"
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
@@ -141,7 +133,10 @@ steps:
                  queue: "default"
            - path:
                - "fastvideo/**"
-                - "fastvideo-kernel/**"
+                - "csrc/attn/sliding_tile_attn/**"
+                - "csrc/attn/sliding_tile_attn/setup.py"
+                - "csrc/attn/sliding_tile_attn/config_sta.py"
+                - "csrc/attn/sliding_tile_attn/st_attn.cpp"
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
@@ -152,16 +147,48 @@ steps:
                agents:
                  queue: "default"
            - path:
-                - "fastvideo-kernel/**"
+                - "csrc/attn/sliding_tile_attn/**"
+                - "csrc/attn/sliding_tile_attn/setup.py"
+                - "csrc/attn/sliding_tile_attn/config_sta.py"
+                - "csrc/attn/sliding_tile_attn/st_attn.cpp"
                - "pyproject.toml"
                - "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/**"
+                - "csrc/attn/video_sparse_attn/**"
+                - "csrc/attn/video_sparse_attn/tk/**"
+                - "csrc/attn/tests/test_vsa.py"
+                - "csrc/attn/video_sparse_attn/setup.py"
+                - "csrc/attn/video_sparse_attn/config_vsa.py"
+                - "csrc/attn/video_sparse_attn/vsa.cpp"
+                - "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:
+                - "csrc/attn/vmoba_attn/**"
+                - "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:
+                - "csrc/attn/vmoba_attn/vmoba/**"
                - "fastvideo/attention/backends/vmoba.py"
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
@@ -171,26 +198,4 @@ steps:
                env:  
                  - TEST_TYPE=inference_vmoba
                agents:
-                  queue: "default"
-            - path:
-                - "fastvideo/**"
-                - "pyproject.toml"
-                - "docker/Dockerfile.python3.12"
-              config:
-                command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Unit Tests"
-                env:
-                  - 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"
+                  queue: "default"
@@ -31,9 +31,9 @@ log "Setting up Modal authentication from Buildkite secrets..."
 MODAL_TOKEN_ID=$(buildkite-agent secret get modal_token_id)
 MODAL_TOKEN_SECRET=$(buildkite-agent secret get modal_token_secret)

-# Retrieve other secrets
 WANDB_API_KEY=$(buildkite-agent secret get wandb_api_key)
-HF_API_KEY=$(buildkite-agent secret get hf_api_key)
+
+WANDB_API_KEY=$(buildkite-agent secret get wandb_api_key)

 if [ -n "$MODAL_TOKEN_ID" ] && [ -n "$MODAL_TOKEN_SECRET" ]; then
    log "Retrieved Modal credentials from Buildkite secrets"
@@ -63,19 +63,19 @@ MODAL_ENV="BUILDKITE_REPO=$BUILDKITE_REPO BUILDKITE_COMMIT=$BUILDKITE_COMMIT BUI
 case "$TEST_TYPE" in
    "encoder")
        log "Running encoder tests..."
-        MODAL_COMMAND="$MODAL_ENV HF_API_KEY=$HF_API_KEY python3 -m modal run $MODAL_TEST_FILE::run_encoder_tests"
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_encoder_tests"
        ;;
    "vae")
        log "Running VAE tests..."
-        MODAL_COMMAND="$MODAL_ENV HF_API_KEY=$HF_API_KEY python3 -m modal run $MODAL_TEST_FILE::run_vae_tests"
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_vae_tests"
        ;;
    "transformer")
        log "Running transformer tests..."
-        MODAL_COMMAND="$MODAL_ENV HF_API_KEY=$HF_API_KEY python3 -m modal run $MODAL_TEST_FILE::run_transformer_tests"
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_transformer_tests"
        ;;
    "ssim")
        log "Running SSIM tests..."
-        MODAL_COMMAND="$MODAL_ENV HF_API_KEY=$HF_API_KEY python3 -m modal run $MODAL_TEST_FILE::run_ssim_tests"
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_ssim_tests"
        ;;
    "training")
        log "Running training tests..."
@@ -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..."
@@ -106,21 +110,13 @@ case "$TEST_TYPE" in
        MODAL_COMMAND="$MODAL_ENV WANDB_API_KEY=$WANDB_API_KEY python3 -m modal run $MODAL_TEST_FILE::run_distill_dmd_tests"
        ;;
        # run_inference_tests_vmoba
-    "self_forcing")
-        log "Running self-forcing tests..."
-        MODAL_COMMAND="$MODAL_ENV WANDB_API_KEY=$WANDB_API_KEY python3 -m modal run $MODAL_TEST_FILE::run_self_forcing_tests"
-        ;;
    "inference_vmoba")
        log "Running V-MoBA inference tests..."
        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_inference_tests_vmoba"
        ;;
-    "unit_test")
-        log "Running unit tests..."
-        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_unit_test"
-        ;;
-    "lora_extraction")
-        log "Running LoRA extraction tests..."
-        MODAL_COMMAND="$MODAL_ENV HF_API_KEY=$HF_API_KEY python3 -m modal run $MODAL_TEST_FILE::run_lora_extraction_tests"
+    "precision_vmoba")
+        log "Running V-MoBA precision tests..."
+        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_precision_tests_vmoba"
        ;;
    *)
        log "Error: Unknown test type: $TEST_TYPE"
@@ -1,65 +1,82 @@
-name: Deploy Documentation
+# Sample workflow for building and deploying a Hugo site to GitHub Pages
+name: Deploy FastVideo Docs to Pages

 on:
+  # Runs on pushes targeting the default branch
  push:
-    branches: [ main ]
+    branches:
+      - main
    paths:
-      - 'docs/**'
-      - 'mkdocs.yml'
-      - 'requirements-mkdocs.txt'
-      - '.github/workflows/docs.yml'
+      - "docs/**/*.md"
+      - "fastvideo/examples/**/*.py"
  pull_request:
-    branches: [ main ]
+    branches: 
+      - main
+    types: [opened, ready_for_review, synchronize, reopened]
    paths:
-      - 'docs/**'
-      - 'mkdocs.yml'
-      - 'requirements-mkdocs.txt'
-      - '.github/workflows/docs.yml'
+      - "docs/**/*.md"
+      - "fastvideo/examples/**/*.py"

+  # Allows you to run this workflow manually from the Actions tab
+  workflow_dispatch:
+
+# Sets permissions of the GITHUB_TOKEN to allow deployment to GitHub Pages
 permissions:
  contents: read
  pages: write
  id-token: write

+# Allow only one concurrent deployment, skipping runs queued between the run in-progress and latest queued.
+# However, do NOT cancel in-progress runs as we want to allow these production deployments to complete.
 concurrency:
  group: "pages"
  cancel-in-progress: false

+# Default to bash
+defaults:
+  run:
+    shell: bash
+
 jobs:
+  pre-commit:
+    uses: ./.github/workflows/pre-commit.yml
+
+  # Build job
  build:
    runs-on: ubuntu-latest
+    needs: pre-commit
    steps:
      - name: Checkout
        uses: actions/checkout@v4
-
-      - name: Setup Python
+      - name: Setup Pages
+        id: pages
+        uses: actions/configure-pages@v5
+      - name: Set up Python
        uses: actions/setup-python@v5
        with:
-          python-version: '3.12'
-
+          python-version: "3.10"
      - name: Install dependencies
        run: |
-          python -m pip install --upgrade pip
-          pip install -r requirements-mkdocs.txt
-
-      - name: Setup Pages
-        uses: actions/configure-pages@v4
-
-      - name: Build documentation
-        run: mkdocs build
-
+          cd docs
+          pip install -r requirements-docs.txt
+      - name: Build docs
+        run: |
+          cd docs
+          make clean
+          make html
      - name: Upload artifact
        uses: actions/upload-pages-artifact@v3
        with:
-          path: ./site
+          path: ./docs/build/html

+  # Deployment job
  deploy:
    environment:
      name: github-pages
      url: ${{ steps.deployment.outputs.page_url }}
+    if: ${{ github.event_name == 'push' }}
    runs-on: ubuntu-latest
    needs: build
-    if: github.ref == 'refs/heads/main'
    steps:
      - name: Deploy to GitHub Pages
        id: deployment
@@ -1,207 +0,0 @@
-name: Publish FastVideo Kernel to PyPI on Version Change
-
-on:
-  push:
-    branches:
-      - main
-    paths:
-      - "fastvideo-kernel/pyproject.toml"
-  workflow_dispatch:
-
-jobs:
-  check-version-change:
-    runs-on: ubuntu-latest
-    outputs:
-      version-changed: ${{ steps.check-version.outputs.changed }}
-      new-version: ${{ steps.check-version.outputs.new-version }}
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-        with:
-          fetch-depth: 2
-
-      - name: Check if version changed
-        id: check-version
-        run: |
-          cd fastvideo-kernel
-          # Get current commit's version from pyproject.toml
-          # Use ^ to match start of line to avoid matching minimum-version
-          NEW_VERSION=$(grep -oP '^version\s*=\s*"\K[^"]+' pyproject.toml)
-          echo "New version: $NEW_VERSION"
-          
-          # Get previous version from git history
-          # Note: git show expects path relative to repo root
-          OLD_VERSION=$(git show HEAD~1:fastvideo-kernel/pyproject.toml | grep -oP '^version\s*=\s*"\K[^"]+' || echo "0.0.0")
-          echo "Old version: $OLD_VERSION"
-          
-          if [ "$NEW_VERSION" != "$OLD_VERSION" ]; then
-            echo "Version changed from $OLD_VERSION to $NEW_VERSION"
-            echo "changed=true" >> $GITHUB_OUTPUT
-            echo "new-version=$NEW_VERSION" >> $GITHUB_OUTPUT
-          else
-            echo "Version did not change"
-            echo "changed=false" >> $GITHUB_OUTPUT
-          fi
-
-  build_wheels:
-    name: Build Wheel
-    needs: check-version-change
-    if: ${{ needs.check-version-change.outputs.version-changed == 'true' || github.event_name == 'workflow_dispatch' }}
-    runs-on: ${{ matrix.os }}
-
-    strategy:
-      fail-fast: false
-      matrix:
-          os: [ubuntu-22.04]
-          python-version: ['3.10', '3.11', '3.12']
-          torch-cuda:
-            # - torch-version: '2.5.1'
-            #   cuda-version: '12.4.1'
-            #   torch-cuda-short: 'cu124'
-            # - torch-version: '2.6.0'
-            #   cuda-version: '12.6.3'
-            #   torch-cuda-short: 'cu126'
-            # - torch-version: '2.7.1'
-            #   cuda-version: '12.8.0'
-            #   torch-cuda-short: 'cu128'
-            - torch-version: '2.9.1'
-              cuda-version: '12.8.0'
-              torch-cuda-short: 'cu128'
-
-    steps:
-      - name: Free up disk space
-        run: |
-          echo "Initial disk space:"
-          df -h
-          
-          # Remove large directories
-          sudo rm -rf /usr/share/dotnet
-          sudo rm -rf /usr/local/lib/android
-          sudo rm -rf /opt/ghc
-          sudo rm -rf /usr/local/share/boost
-          sudo rm -rf /usr/share/swift
-          sudo rm -rf /usr/local/lib/node_modules
-          sudo rm -rf /usr/local/share/powershell
-          sudo rm -rf /usr/share/rust
-          sudo rm -rf /usr/local/.ghcup
-          
-          # Remove cached files
-          sudo rm -rf /var/lib/apt/lists/*
-          sudo rm -rf /var/cache/apt/archives/*
-          
-          echo "Disk space after cleanup:"
-          df -h
-
-      - name: Checkout
-        uses: actions/checkout@v4
-
-      - name: Set up Python
-        uses: actions/setup-python@v5
-        with:
-          python-version: ${{ matrix.python-version }}
-
-      - name: Install CUDA ${{ matrix.torch-cuda.cuda-version }}
-        uses: Jimver/cuda-toolkit@v0.2.21
-        id: cuda-toolkit
-        with:
-          cuda: ${{ matrix.torch-cuda.cuda-version }}
-          linux-local-args: '["--toolkit"]'
-          method: 'network'
-
-      - name: Install dependencies (GCC, Clang, CUDA Paths, Git)
-        run: |
-          sudo apt update
-          sudo apt install -y git patchelf gcc-11 g++-11 clang-11
-          sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
-
-          # Allow Git to Access Safe Directory
-          git config --global --add safe.directory /__w/FastVideo/FastVideo
-
-          # Set CUDA environment variables
-          export CUDA_HOME=/usr/local/cuda-${{ matrix.torch-cuda.cuda-version }}
-          export PATH=${CUDA_HOME}/bin:${PATH}
-          export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
-          
-          # Verify installation
-          gcc --version
-          g++ --version
-          clang-11 --version
-          nvcc --version
-
-      - name: Install PyTorch ${{ matrix.torch-cuda.torch-version }}+cu${{ matrix.torch-cuda.cuda-version }}
-        run: |
-          pip install --upgrade pip
-          pip install typing-extensions==4.12.2
-          pip install --no-cache-dir torch==${{ matrix.torch-cuda.torch-version }} --index-url https://download.pytorch.org/whl/${{matrix.torch-cuda.torch-cuda-short}}
-          nvcc --version
-          python --version
-          python -c "import torch; print('PyTorch:', torch.__version__)"
-          python -c "import torch; print('CUDA:', torch.version.cuda)"
-          python -c "from torch.utils import cpp_extension; print (cpp_extension.CUDA_HOME)"
-
-      - name: Build wheel
-        run: |
-          export PYTHONPATH=$GITHUB_WORKSPACE:$PYTHONPATH
-          
-          pip install setuptools ninja packaging wheel triton scikit-build-core cmake build
-          
-          cd fastvideo-kernel
-          git submodule update --init --recursive  # Ensure ThunderKittens submodule is initialized
-          # Release builds are produced on GPU-less runners, so force-enable TK and target Hopper.
-          export TORCH_CUDA_ARCH_LIST="9.0a"
-          export CMAKE_ARGS="${CMAKE_ARGS:-} -DFASTVIDEO_KERNEL_BUILD_TK=ON -DCMAKE_CUDA_ARCHITECTURES=90a"
-          
-          # Build standard wheel (no local version suffix) for PyPI
-          python -m build --wheel --outdir dist
-          
-          # Fix the wheel to be manylinux compliant
-          pip install auditwheel
-          # Target manylinux_2_35 (Ubuntu 22.04 native)
-          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/
-
-      - name: Upload wheel artifact
-        # Only upload if it's the "main" CUDA version we want on PyPI
-        # We upload all to artifacts for inspection/GH releases, but give them distinct artifact names
-        uses: actions/upload-artifact@v4
-        with:
-          name: fastvideo_kernel-py${{ matrix.python-version }}-${{ matrix.torch-cuda.torch-cuda-short }}-torch${{ matrix.torch-cuda.torch-version }}
-          path: fastvideo-kernel/dist/*.whl
-          retention-days: 90
-
-  publish_package:
-    name: Publish package
-    needs: [build_wheels, check-version-change]
-    if: ${{ needs.check-version-change.outputs.version-changed == 'true' || github.event_name == 'workflow_dispatch' }}
-    runs-on: ubuntu-22.04
-    permissions:
-      id-token: write  # Needed for OIDC Trusted Publishing
-
-    steps:
-      - uses: actions/checkout@v4
-
-      - uses: actions/setup-python@v5
-        with:
-          python-version: '3.10'
-
-      - name: Download PyPI wheels
-        uses: actions/download-artifact@v4
-        with:
-          path: fastvideo-kernel/dist/
-          pattern: 'fastvideo_kernel-py*'
-          merge-multiple: true
-
-      - name: Build source distribution
-        run: |
-          pip install build scikit-build-core cmake ninja
-          
-          cd fastvideo-kernel
-          # We don't need full CUDA/Torch to just package the source (sdist)
-          python -m build --sdist --outdir dist
-
-      - name: Publish release distributions to PyPI
-        uses: pypa/gh-action-pypi-publish@release/v1
-        with:
-          packages-dir: fastvideo-kernel/dist/
@@ -62,8 +62,8 @@ on:
        required: false
        default: false
        type: boolean
-      run_unit_test:
-        description: "Run unit-test"
+      run_nightly_test:
+        description: "Run nightly-test"
        required: false
        default: false
        type: boolean
@@ -93,7 +93,6 @@ jobs:
      inference-test-STA: ${{ steps.filter.outputs.inference-test-STA }}
      precision-test-STA: ${{ steps.filter.outputs.precision-test-STA }}
      precision-test-VSA: ${{ steps.filter.outputs.precision-test-VSA }}
-      unit-test: ${{ steps.filter.outputs.unit-test }}
    steps:
      - uses: actions/checkout@v4
      - uses: dorny/paths-filter@v3
@@ -103,8 +102,6 @@ jobs:
            # Define reusable path patterns
            common-paths: &common-paths
              - 'pyproject.toml'
-              - 'docker/Dockerfile.python3.10'
-              - 'docker/Dockerfile.python3.11'
              - 'docker/Dockerfile.python3.12'
            sta-kernel-paths: &sta-kernel-paths
              - 'csrc/attn/sliding_tile_attn/**'
@@ -158,9 +155,6 @@ jobs:
            precision-test-VSA:
              - *common-paths
              - *vsa-kernel-paths
-            unit-test:
-              - 'fastvideo/**'
-              - *common-paths

  encoder-test:
    needs: change-filter
@@ -339,42 +333,23 @@ jobs:
      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}

-  unit-test:
-    needs: change-filter
+  nightly-test:
    if: >-
-      (github.event_name != 'workflow_dispatch' && needs.change-filter.outputs.unit-test == 'true') || 
-      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_unit_test == 'true')
+      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_nightly_test == 'true')
    uses: ./.github/workflows/runpod-test.yml
    with:
-      job_id: "unit-test"
-      gpu_type: "NVIDIA L40S"
-      gpu_count: 1
+      job_id: "nightly-test"
+      gpu_type: "NVIDIA A40"
+      gpu_count: 4
      volume_size: 100
      disk_size: 100
      image: "ghcr.io/${{ github.repository }}/fastvideo-dev:py3.12-latest"
-      test_command: "uv pip install -e .[test] && pytest ./fastvideo/dataset/ -vs && pytest ./fastvideo/workflow/ -vs && pytest ./fastvideo/entrypoints/ -vs"
+      test_command: "wandb login $WANDB_API_KEY && uv pip install -e .[test] && pytest ./fastvideo/tests/nightly/test_e2e_overfit_single_sample.py -vs"
      timeout_minutes: 30
    secrets:
      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-
-  # nightly-test:
-  #   if: >-
-  #     (github.event_name == 'workflow_dispatch' && github.event.inputs.run_nightly_test == 'true')
-  #   uses: ./.github/workflows/runpod-test.yml
-  #   with:
-  #     job_id: "nightly-test"
-  #     gpu_type: "NVIDIA A40"
-  #     gpu_count: 4
-  #     volume_size: 100
-  #     disk_size: 100
-  #     image: "ghcr.io/${{ github.repository }}/fastvideo-dev:py3.12-latest"
-  #     test_command: "wandb login $WANDB_API_KEY && uv pip install -e .[test] && pytest ./fastvideo/tests/nightly/test_e2e_overfit_single_sample.py -vs"
-  #     timeout_minutes: 30
-  #   secrets:
-  #     RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-  #     RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-  #     WANDB_API_KEY: ${{ secrets.WANDB_API_KEY }}
+      WANDB_API_KEY: ${{ secrets.WANDB_API_KEY }}

  runpod-cleanup:
    # Add other jobs to this list as you create them
@@ -14,8 +14,6 @@ wandb/
 *.pt
 cache_dir/
 wandb/
-venv/
-.venv/
 runs/
 samples/
 *validation/
@@ -30,8 +28,6 @@ env
 **/build/
 **.pyc
 **.txt
-*.log
-weights/

 # Distribution / packaging
 build/
@@ -41,13 +37,12 @@ dist/
 eggs/
 .eggs/

-# MkDocs documentation
-site/
-docs/getting_started/examples/
-docs/inference/examples/
-docs/training/examples/
-docs/distillation/examples/
-!requirements-mkdocs.txt
+# Sphinx documentation
+docs/_build/
+docs/source/getting_started/examples/
+docs/source/inference/examples/
+docs/source/training/examples/
+docs/source/distillation/examples/

 # VSCode
 .vscode/
@@ -66,9 +61,7 @@ docs/distillation/examples/
 !fastvideo/tests/ssim/reference_videos/**/*.mp4

 # Static images
-!docs/assets/images/**/*.png
+!docs/source/_static/images/**/*.png
 !comfyui/assets/**/*.png
 !comfyui/assets/**/*.gif
-
-dmd_t2v_output/
-preprocess_output_text/
+dmd_t2v_output/
@@ -1,6 +1,7 @@
-[submodule "fastvideo-kernel/include/tk"]
-	path = fastvideo-kernel/include/tk
+[submodule "csrc/attn/video_sparse_attn/tk"]
+	path = csrc/attn/video_sparse_attn/tk
+	url = https://github.com/HazyResearch/ThunderKittens.git
+
+[submodule "csrc/attn/sliding_tile_attn/tk"]
+	path = csrc/attn/sliding_tile_attn/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
@@ -4,15 +4,17 @@ default_stages:
 exclude: |
  (?x)(
    fastvideo/third_party/.*|
-    fastvideo-kernel/.*|
+    csrc/.*|
    assets/.*|
    tests/.*|
    demo/.*|
    predict\.py|
    scripts/.*|
-    prompts/.*|
    fastvideo/data_preprocess/.*|
    fastvideo/dataset/.*|
+    fastvideo/distill/.*|
+    fastvideo/distill\.py|
+    fastvideo/distill_adv\.py|
    fastvideo/models/.*|
    fastvideo/sample/.*|
    fastvideo/train\.py|
@@ -42,10 +44,10 @@ repos:
  - id: codespell
    additional_dependencies: ['tomli']
    args: ['--toml', 'pyproject.toml']
-# - repo: https://github.com/PyCQA/isort
-#   rev: 6.0.1
-#   hooks:
-#   - id: isort
+- repo: https://github.com/PyCQA/isort
+  rev: 6.0.1
+  hooks:
+  - id: isort
 - repo: https://github.com/jackdewinter/pymarkdown
  rev: v0.9.30
  hooks:
@@ -2,46 +2,40 @@
 <img src=assets/logos/logo.svg width="30%"/>
 </div>

-<p align="center">
-     | <a href="https://hao-ai-lab.github.io/FastVideo"><b>Documentation</b></a> | <a href="https://hao-ai-lab.github.io/FastVideo/inference/inference_quick_start/"><b> Quick Start</b></a> | <a href="https://github.com/hao-ai-lab/FastVideo/discussions/982"  target="_blank"><b>Weekly Dev Meeting</b></a>  | 🟣💬 <a href="https://join.slack.com/t/fastvideo/shared_invite/zt-3f4lao1uq-u~Ipx6Lt4J27AlD2y~IdLQ" target="_blank"> <b>Slack</b> </a> |  🟣💬 <a href="https://ibb.co/sv3MMKyv" target="_blank"> <b> WeChat </b> </a> |
-</p>
-
 **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://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.html"><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-3csdw1isz-Euq8_Q8~baewG8hxjXs2gQ" target="_blank"> <b>Slack</b> </a> |  🟣💬 <a href="https://ibb.co/S7HLCSTh" target="_blank"> <b> WeChat </b> </a> |
+</p>
+
+<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/08/04```: Release [FastWan](https://hao-ai-lab.github.io/FastVideo/distillation/dmd.html) models and [Sparse-Distillation](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/).
 - ```2025/06/14```: Release finetuning and inference code for [VSA](https://arxiv.org/pdf/2505.13389)
 - ```2025/04/24```: [FastVideo V1](https://hao-ai-lab.github.io/blogs/fastvideo/) is released!
 - ```2025/02/18```: Release the inference code for [Sliding Tile Attention](https://hao-ai-lab.github.io/blogs/sta/).

-</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:
@@ -55,10 +49,10 @@ conda activate fastvideo
 pip install fastvideo
 ```

-Please see our [docs](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/) for more detailed installation instructions.
+Please see our [docs](https://hao-ai-lab.github.io/FastVideo/getting_started/installation.html) for more detailed installation instructions.

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

 See below for recipes and datasets:

@@ -70,7 +64,7 @@ See below for recipes and datasets:

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

 ```python
 import os
@@ -106,32 +100,35 @@ Run the script with:
 python example.py
 ```

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

 ### Other docs:

- [Design Overview](https://hao-ai-lab.github.io/FastVideo/design/overview/)
- [Contribution Guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/)
+- [Design Overview](https://hao-ai-lab.github.io/FastVideo/design/overview.html)
+- [Contribution Guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation.html)

 ## Distillation and Finetuning
- [Distillation Guide](https://hao-ai-lab.github.io/FastVideo/distillation/dmd/)
+- [Distillation Guide](https://hao-ai-lab.github.io/FastVideo/distillation/dmd.html)
 <!-- - [Finetuning Guide](https://hao-ai-lab.github.io/FastVideo/training/finetune.html) -->

-## Awesome work using FastVideo or our research projects
+## 📑 Development Plan
+<!-- - More distillation methods -->
+  <!-- - [ ] Add Distribution Matching Distillation -->
+More FastWan Models Coming Soon!
+- [ ] Add FastWan2.1-T2V-14B
+- [ ] Add FastWan2.2-T2V-14B
+- [ ] Add FastWan2.2-I2V-14B
+<!-- - Optimization features
+- Code updates -->
+  <!-- - [ ] fp8 support -->
+  <!-- - [ ] faster load model and save model support -->

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

 ## 🤝 Contributing

-We welcome all contributions. Please check out our guide [here](https://hao-ai-lab.github.io/FastVideo/contributing/overview/).
-See details in [development roadmap](https://github.com/hao-ai-lab/FastVideo/issues/899).
+We welcome all contributions. Please check out our guide [here](https://hao-ai-lab.github.io/FastVideo/contributing/overview.html)
+
 ## Acknowledgement
 We learned and reused code from the following projects:
 - [Wan-Video](https://github.com/Wan-Video)
@@ -158,8 +155,8 @@ If you find FastVideo useful, please considering citing our work:
 }

@article{zhang2025vsa,
-  title={Vsa: Faster video diffusion with trainable sparse attention},
-  author={Zhang, Peiyuan and Chen, Yongqi and Huang, Haofeng and Lin, Will and Liu, Zhengzhong and Stoica, Ion and Xing, Eric and Zhang, Hao},
+  title={VSA: Faster Video Diffusion with Trainable Sparse Attention},
+  author={Zhang, Peiyuan and Huang, Haofeng and Chen, Yongqi and Lin, Will and Liu, Zhengzhong and Stoica, Ion and Xing, Eric and Zhang, Hao},
  journal={arXiv preprint arXiv:2505.13389},
  year={2025}
 }
@@ -1,18 +0,0 @@
-<svg width="252" height="105" viewBox="0 0 252 105" fill="none" xmlns="http://www.w3.org/2000/svg">
-<path d="M89.4843 55.5457H101.361L87.7028 101H74.638L89.4843 55.5457Z" fill="#356CFF"/>
-<path fill-rule="evenodd" clip-rule="evenodd" d="M96.0167 1.00057H112.645L118.583 48.273H104.924L103.737 39.7882H79.9827L67.5117 55.5457H85.3273L43.1638 101H28.3174L22.3789 55.5457H33.6621L38.4129 91.3031L58.604 68.2729H44.3515L96.0167 1.00057ZM100.768 13.1217L87.7028 29.4852H103.143L100.768 13.1217Z" fill="#356CFF"/>
-<path d="M37.2252 1.00057L22.3789 48.273H36.0375L40.7884 30.6974L62.6727 30.6974L69.6727 21.0005L43.7576 21.0004L46.7269 11.9096L77.6727 11.9096L86 1.00057L37.2252 1.00057Z" fill="#356CFF"/>
-<path fill-rule="evenodd" clip-rule="evenodd" d="M108.488 55.5457L94.2351 101C94.2351 101 105.518 101 120.959 101C136.399 101 144.078 93.0133 148.276 79.788C152.432 68.0157 153.027 55.5457 136.399 55.5457C119.771 55.5457 108.488 55.5457 108.488 55.5457ZM109.081 90.697L116.802 65.8487C116.802 65.8487 120.959 65.8487 132.242 65.8487C143.525 65.8487 137.586 78.5759 135.211 84.0304C133.307 88.4021 127.491 90.697 122.74 90.697C117.989 90.697 109.081 90.697 109.081 90.697Z" fill="#356CFF"/>
-<path d="M173.188 1.00056L168.625 11.9096C168.625 11.9096 149.386 11.9092 142.525 11.9095C135.664 11.9098 136.586 20.3944 141.337 20.3944H159.747C168.654 20.3944 166.961 33.6899 163.904 38.5761C160.467 44.0675 157.371 48.273 148.463 48.273L125.188 48.273L124 37.97L147.87 37.97C153.808 37.97 156.184 29.4852 151.433 29.4852H131.836C120.142 29.4852 125.897 1.00043 141.337 1.00043L173.188 1.00056Z" fill="#356CFF"/>
-<path d="M179.938 1.00056L175.688 11.9096L191.221 11.9096L179.938 48.273H192.409L203.692 11.9096L219.132 11.9095L223.289 1.00043L179.938 1.00056Z" fill="#356CFF"/>
-<path d="M161.341 55.5457H202.845L198.5 65.8487H169.654L167.279 73.7268H188.5L184.749 82.8177H164.31L161.934 90.697H190.251L186.624 101H146.494L161.341 55.5457Z" fill="#356CFF"/>
-<path fill-rule="evenodd" clip-rule="evenodd" d="M230.821 54.9391C255.169 54.9391 251.776 67.0602 249.231 77.9692C246.686 88.8783 240.917 101 217.757 101C194.596 101 195.606 88.8783 199.347 77.9692C203.089 67.0602 206.473 54.9391 230.821 54.9391ZM237.948 77.9692C239.984 70.6965 240.917 65.242 228.446 65.242C215.975 65.242 211.818 71.9087 210.037 77.9692C208.255 84.0298 208.255 91.3025 219.538 91.3025C230.821 91.3025 235.911 85.2419 237.948 77.9692Z" fill="#356CFF"/>
-<path d="M173.188 1.00056L168.625 11.9096C168.625 11.9096 149.386 11.9092 142.525 11.9095C135.664 11.9098 136.586 20.3944 141.337 20.3944M173.188 1.00056C173.188 1.00056 156.777 1.00043 141.337 1.00043M173.188 1.00056L141.337 1.00043M141.337 20.3944C146.088 20.3944 150.839 20.3944 159.747 20.3944M141.337 20.3944H159.747M159.747 20.3944C168.654 20.3944 166.961 33.6899 163.904 38.5761C160.467 44.0675 157.371 48.273 148.463 48.273M148.463 48.273C139.556 48.273 125.188 48.273 125.188 48.273M148.463 48.273L125.188 48.273M125.188 48.273L124 37.97M124 37.97C124 37.97 141.931 37.97 147.87 37.97M124 37.97L147.87 37.97M147.87 37.97C153.808 37.97 156.184 29.4852 151.433 29.4852M151.433 29.4852C146.682 29.4852 138.962 29.4852 131.836 29.4852M151.433 29.4852H131.836M131.836 29.4852C120.142 29.4852 125.897 1.00043 141.337 1.00043M37.2252 1.00057L22.3789 48.273H36.0375L40.7884 30.6974L62.6727 30.6974L69.6727 21.0005L43.7576 21.0004L46.7269 11.9096L77.6727 11.9096L86 1.00057L37.2252 1.00057ZM96.0167 1.00057H112.645L118.583 48.273H104.924L103.737 39.7882H79.9827L67.5117 55.5457H85.3273L43.1638 101H28.3174L22.3789 55.5457H33.6621L38.4129 91.3031L58.604 68.2729H44.3515L96.0167 1.00057ZM87.7028 29.4852L100.768 13.1217L103.143 29.4852H87.7028ZM89.4843 55.5457H101.361L87.7028 101H74.638L89.4843 55.5457ZM108.488 55.5457L94.2351 101C94.2351 101 105.518 101 120.959 101C136.399 101 144.078 93.0133 148.276 79.788C152.432 68.0157 153.027 55.5457 136.399 55.5457C119.771 55.5457 108.488 55.5457 108.488 55.5457ZM116.802 65.8487L109.081 90.697C109.081 90.697 117.989 90.697 122.74 90.697C127.491 90.697 133.307 88.4021 135.211 84.0304C137.586 78.5759 143.525 65.8487 132.242 65.8487C120.959 65.8487 116.802 65.8487 116.802 65.8487ZM179.938 1.00056L175.688 11.9096L191.221 11.9096L179.938 48.273H192.409L203.692 11.9096L219.132 11.9095L223.289 1.00043L179.938 1.00056ZM161.341 55.5457H202.845L198.5 65.8487H169.654L167.279 73.7268H188.5L184.749 82.8177H164.31L161.934 90.697H190.251L186.624 101H146.494L161.341 55.5457ZM230.821 54.9391C255.169 54.9391 251.776 67.0602 249.231 77.9692C246.686 88.8783 240.917 101 217.757 101C194.596 101 195.606 88.8783 199.347 77.9692C203.089 67.0602 206.473 54.9391 230.821 54.9391ZM228.446 65.242C240.917 65.242 239.984 70.6965 237.948 77.9692C235.911 85.2419 230.821 91.3025 219.538 91.3025C208.255 91.3025 208.255 84.0298 210.037 77.9692C211.818 71.9087 215.975 65.242 228.446 65.242Z" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M15.2524 55.5451L21.191 100.999L24.7541 100.999L18.8156 55.5451L15.2524 55.5451Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M8.12646 55.5451L14.065 100.999L15.2527 100.999L9.31417 55.5451L8.12646 55.5451Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M1 55.5451L6.93853 100.999L7.53239 100.999L1.59385 55.5451L1 55.5451Z" fill="#356CFF" stroke="#356CFF" stroke-width="0.593853"/>
-<path d="M15.2524 48.2724L30.0988 1H33.6619L18.8156 48.2724H15.2524Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M8.12646 48.2724L22.9728 1H24.1605L9.31417 48.2724H8.12646Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M1 48.2724L15.8463 1H16.4402L1.59385 48.2724H1Z" fill="#356CFF" stroke="#356CFF" stroke-width="0.593853"/>
-<path d="M85.3271 55.5457H67.5116L87 12.7363L44.3513 68.2729H58.6038L43.1636 101L85.3271 55.5457Z" fill="#FDC717" stroke="#FDC717" stroke-width="1.18771" stroke-miterlimit="16"/>
-</svg>
@@ -1,6 +0,0 @@
-<svg width="160" height="93" viewBox="0 0 160 93" fill="none" xmlns="http://www.w3.org/2000/svg">
-<path d="M28.8511 91.66L57.6319 1.86368H64.5394L35.7585 91.66H28.8511Z" fill="#356CFF" stroke="#356CFF" stroke-width="2.30244"/>
-<path d="M15.0376 91.66L43.8185 1.86368H46.1209L17.3401 91.66H15.0376Z" fill="#356CFF" stroke="#356CFF" stroke-width="2.30244"/>
-<path d="M1.22217 91.66L30.003 1.86366H31.1543L2.3734 91.66H1.22217Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.15122"/>
-<path d="M71.4465 1.86483L42.666 91.6599H69.144L78.3538 58.2746H123.251L129.007 39.855H84.1099L89.866 22.5868H152.032L157.788 1.86483H71.4465Z" fill="#356CFF" stroke="#356CFF" stroke-width="2.30244"/>
-</svg>
@@ -1,106 +0,0 @@
-# FVD (Fréchet Video Distance) Benchmark
-
-Evaluate generated video quality using FVD with the I3D feature extractor.
-
-## Quick Start
-
-**Run the benchmark:**
-
-```bash
-bash benchmarks/scripts/run.sh
-```
-
-That's it! The script auto-installs dependencies and runs the benchmark.
-
-**To customize:** Edit `benchmarks/fvd/run_fvd.py` to change:
- Video paths (`real_dir`, `gen_dir`)
- Number of videos, frames, sampling strategy
- Device, batch size, caching, etc.
-
-## Advanced Usage (CLI)
-
-For more control without editing Python files, use the CLI.
-
-**First-time setup** (one-time per pod/environment):
-
-```bash
-bash benchmarks/scripts/setup_fvd.sh
-```
-
-Then run any configuration you want:
-
-```bash
-# Custom configuration
-python -m benchmarks.fvd.cli \
-    --real-path data/real/ \
-    --gen-path outputs/gen/ \
-    --num-videos 1024 \
-    --num-frames 32 \
-    --clip-strategy random \
-    --batch-size 32 \
-    --seed 42 \
-    --extractor clip
-```
-
-**Standard protocols:**
-
-```bash
-# Use predefined protocols
-python -m benchmarks.fvd.cli \
-    --real-path data/real/ \
-    --gen-path outputs/gen/ \
-    --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
-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)
-```
-
-Run `python -m benchmarks.fvd.cli --help` for all options.
-
-## Available Protocols
-
- `fvd2048_16f` - Standard (2048 videos, 16 frames)
- `fvd2048_128f` - Long videos (128 frames)
- `fvd2048_128f_subsample8` - Subsampled long videos
- `quick_test` - Fast testing (10 videos)
-
-## Configuration Options
-
-Key options in `FVDConfig`:
-
-```python
-num_videos=2048,              # Videos to evaluate
-num_frames_per_clip=16,       # Frames per clip
-clip_strategy='beginning',    # beginning|random|uniform|middle|sliding
-frame_stride=1,               # Frame subsampling
-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
-
-```python
-from benchmarks.fvd import compute_fvd_with_config, FVDConfig
-
-config = FVDConfig.fvd2048_16f()  # or custom config
-results = compute_fvd_with_config('data/real/', 'outputs/gen/', config)
-print(f"FVD: {results['fvd']:.2f}")
-```
-
-## Notes
-
- 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
@@ -1,38 +0,0 @@
-"""
-FastVideo Frechet Video Distance (FVD) Benchmark Module.
-    >>> from fastvideo.benchmarks.fvd import compute_fvd_with_config, FVDConfig
-    >>> config = FVDConfig.fvd2048_16f()  # Standard protocol
-    >>> results = compute_fvd_with_config('data/real/', 'outputs/gen/', config)
-    >>> print(f"FVD: {results['fvd']:.2f}")
-"""
-
-from .fvd import (
-    compute_fvd,
-    compute_fvd_with_config,
-    compute_frechet_distance,
-    compute_statistics,
-    FVDConfig,
-)
-from .feature_extractors import (BaseFeatureExtractor, I3DFeatureExtractor,
-                                 load_extractor)
-from .video_utils import (
-    load_video_auto,
-    sample_clips_from_video,
-    load_video_clips_streaming,
-    ClipSamplingStrategy,
-)
-
-__all__ = [
-    'compute_fvd',
-    'compute_fvd_with_config',
-    'compute_frechet_distance',
-    'compute_statistics',
-    'FVDConfig',
-    'BaseFeatureExtractor',
-    'I3DFeatureExtractor',
-    'load_extractor',
-    'load_video_auto',
-    'sample_clips_from_video',
-    'load_video_clips_streaming',
-    'ClipSamplingStrategy',
-]
@@ -1,107 +0,0 @@
-import argparse
-import sys
-import traceback
-from .fvd import compute_fvd_with_config, FVDConfig
-
-
-def main() -> int:
-    parser = argparse.ArgumentParser(
-        description='Compute Fréchet Video Distance (FVD)')
-
-    # Required arguments
-    parser.add_argument('--real-path',
-                        type=str,
-                        required=True,
-                        help='Path to real videos')
-    parser.add_argument('--gen-path',
-                        type=str,
-                        required=True,
-                        help='Path to generated videos')
-
-    # Extractor selection
-    parser.add_argument('--extractor',
-                        type=str,
-                        default='i3d',
-                        choices=['i3d', 'clip', 'videomae'],
-                        help='Feature extractor model to use (default: i3d)')
-
-    # Standard args
-    parser.add_argument('--seed',
-                        type=int,
-                        default=None,
-                        help='Random seed for reproducibility')
-    parser.add_argument('--protocol',
-                        type=str,
-                        default=None,
-                        choices=['fvd2048_16f', 'fvd2048_128f', 'quick_test'],
-                        help='Use standard protocol (overrides other settings)')
-    parser.add_argument('--num-videos',
-                        type=int,
-                        default=2048,
-                        help='Number of videos to use')
-    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')
-    parser.add_argument('--batch-size',
-                        type=int,
-                        default=32,
-                        help='Batch size for feature extraction')
-    parser.add_argument('--device',
-                        type=str,
-                        default='cuda',
-                        help='Device to use (cuda or cpu)')
-    parser.add_argument('--cache-real-features',
-                        type=str,
-                        default=None,
-                        help='Path to cache real video features')
-    parser.add_argument('--quiet',
-                        action='store_true',
-                        help='Suppress progress output')
-
-    args = parser.parse_args()
-
-    # Create config
-    if args.protocol:
-        protocol_map = {
-            'fvd2048_16f': FVDConfig.fvd2048_16f,
-            'fvd2048_128f': FVDConfig.fvd2048_128f,
-            'quick_test': FVDConfig.quick_test,
-        }
-        config = protocol_map[args.protocol]()
-        # Apply overrides
-        config.device = args.device
-        config.cache_real_features = args.cache_real_features
-        config.extractor_model = args.extractor  # Apply extractor arg
-    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)
-
-    try:
-        _ = compute_fvd_with_config(
-            args.real_path,  # noqa: F841
-            args.gen_path,
-            config,
-            verbose=not args.quiet)
-
-        return 0
-
-    except Exception as e:
-        print(f"Error: {e}", file=sys.stderr)
-        traceback.print_exc(file=sys.stderr)
-        return 1
-
-
-if __name__ == '__main__':
-    sys.exit(main())
@@ -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")
@@ -1,405 +0,0 @@
-import numpy as np
-import scipy.linalg
-import torch
-from pathlib import Path
-from collections.abc import Iterator
-import pickle
-from dataclasses import dataclass, field
-from .feature_extractors import BaseFeatureExtractor, load_extractor
-from .video_utils import ClipSamplingStrategy, load_video_clips_streaming
-
-
-def compute_statistics(features: np.ndarray) -> tuple[np.ndarray, np.ndarray]:
-    """Compute mean and covariance."""
-    mu = np.mean(features, axis=0)
-    sigma = np.cov(features, rowvar=False)
-    return mu, sigma
-
-
-def compute_frechet_distance(mu1: np.ndarray,
-                             sigma1: np.ndarray,
-                             mu2: np.ndarray,
-                             sigma2: np.ndarray,
-                             eps: float = 1e-6) -> float:
-    """
-    Compute Fréchet distance between two Gaussians.
-    """
-    sigma1 = sigma1 + eps * np.eye(sigma1.shape[0])
-    sigma2 = sigma2 + eps * np.eye(sigma2.shape[0])
-
-    diff = mu1 - mu2
-    mean_distance = np.sum(diff**2)
-
-    trace_sum = np.trace(sigma1 + sigma2)
-
-    covmean = scipy.linalg.sqrtm(sigma1 @ sigma2)
-
-    if np.iscomplexobj(covmean):
-        if not np.allclose(np.diagonal(covmean).imag, 0, atol=1e-3):
-            print(
-                f"Warning: Imaginary component: {np.max(np.abs(covmean.imag))}")
-        covmean = covmean.real
-
-    trace_product = np.trace(covmean)
-
-    fvd = mean_distance + trace_sum - 2 * trace_product
-
-    return float(fvd)
-
-
-@dataclass
-class FVDConfig:
-    # default configuration for FVD computation:
-
-    # 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
-    clip_strategy: str | ClipSamplingStrategy = 'beginning'
-
-    # Temporal subsampling
-    frame_stride: int = 1  # 1=no subsampling, 2=every 2nd, 8=every 8th
-    temporal_stride: int = 1  # For sliding window clips
-
-    # Data processing
-    video_extensions: list[str] = field(
-        default_factory=lambda: ['.mp4', '.avi', '.mov', '.mkv'])
-    support_frame_dirs: bool = True
-
-    # Computation
-    batch_size: int = 32
-    device: str = 'cuda'
-
-    use_streaming: bool = True
-    resize_before_extraction: bool = True
-
-    # Caching
-    cache_real_features: str | None = None
-    i3d_model_path: str | None = None
-
-    # Reproducibility
-    seed: int | None = None
-
-    @classmethod
-    def fvd2048_16f(cls) -> 'FVDConfig':
-        """Standard FVD protocol: 2048 videos, 16 frames, beginning clip."""
-        return cls(num_videos=2048,
-                   num_frames_per_clip=16,
-                   clip_strategy='beginning',
-                   use_streaming=True)
-
-    @classmethod
-    def fvd2048_128f(cls) -> 'FVDConfig':
-        """Long video protocol: 2048 videos, 128 frames."""
-        return cls(num_videos=2048,
-                   num_frames_per_clip=128,
-                   clip_strategy='beginning',
-                   use_streaming=True)
-
-    @classmethod
-    def quick_test(cls) -> 'FVDConfig':
-        """Quick test config: 100 videos, 16 frames."""
-        return cls(num_videos=100,
-                   num_frames_per_clip=16,
-                   clip_strategy='beginning')
-
-    def to_dict(self) -> dict:
-        """Export config to dict for logging"""
-        d = self.__dict__.copy()
-        d['clip_strategy'] = str(self.clip_strategy)
-        return d
-
-    def __str__(self) -> str:
-        """Human-readable protocol name"""
-        desc = f"FVD_{self.extractor_model.upper()}_{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:
-            desc += f"_{self.num_clips_per_video}clips"
-        if self.clip_strategy != 'beginning':
-            desc += f"_{self.clip_strategy}"
-        return desc
-
-
-def extract_features_streaming(video_generator: Iterator[torch.Tensor],
-                               extractor: BaseFeatureExtractor,
-                               batch_size: int = 32,
-                               max_clips: int | None = None,
-                               verbose: bool = True) -> np.ndarray:
-    """
-    Extract features from a video clip generator using streaming.
-    """
-    all_features = []
-    batch = []
-
-    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)
-
-            # 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:
-            batch_tensor = torch.stack(batch).to(extractor.device)
-            features = extractor.extract_features_batch(batch_tensor)
-            all_features.append(features.detach().cpu().numpy())
-
-    if len(all_features) == 0:
-        raise RuntimeError("No features extracted - check video loading")
-
-    features = np.concatenate(all_features, axis=0)
-
-    if verbose:
-        print(f"Extracted {len(features)} feature vectors")
-
-    return features
-
-
-def load_or_compute_features(videos: str | Path | torch.Tensor,
-                             extractor: BaseFeatureExtractor,
-                             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"
-
-        if cache_file.exists():
-            print(f"Loading cached features from {cache_file}")
-            with open(cache_file, 'rb') as f:
-                features = pickle.load(f)
-
-            # 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...")
-            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)):
-        target_size = (224, 224) if config.resize_before_extraction else None
-
-        video_generator = load_video_clips_streaming(
-            videos,
-            num_frames=config.num_frames_per_clip,
-            max_videos=config.num_videos,
-            clip_strategy=config.clip_strategy,
-            frame_stride=config.frame_stride,
-            num_clips_per_video=config.num_clips_per_video,
-            video_extensions=config.video_extensions,
-            support_frame_dirs=config.support_frame_dirs,
-            target_size=target_size,
-            verbose=True)
-
-        max_clips = config.num_videos * config.num_clips_per_video
-        features = extract_features_streaming(video_generator,
-                                              extractor,
-                                              batch_size=config.batch_size,
-                                              max_clips=max_clips,
-                                              verbose=True)
-    else:
-        print(f"Extracting features from {len(videos)} video tensors...")
-        features = extractor.extract_features(videos,
-                                              batch_size=config.batch_size,
-                                              verbose=True)
-        features = features.numpy()
-
-    # Validate feature count
-    expected_count = config.num_videos * config.num_clips_per_video
-    if len(features) < expected_count:
-        raise ValueError(
-            f"ERROR: Only extracted {len(features)} features, but need {expected_count}!\n"
-            f"Found fewer videos than expected. Check your video directory.")
-    elif len(features) > expected_count:
-        print(f"Truncating {len(features)} features to {expected_count}")
-        features = features[:expected_count]
-
-    # Cache features if requested
-    if cache_path is not None:
-        script_dir = Path(__file__).parent
-        cache_dir = script_dir / cache_path
-        cache_dir.mkdir(parents=True, exist_ok=True)
-        cache_file = cache_dir / f"{config.extractor_model}_{cache_name}.pkl"
-        print(f"Caching features to {cache_file}")
-        with open(cache_file, 'wb') as f:
-            pickle.dump(features, f)
-
-    return features
-
-
-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}")
-    """
-
-    # Seed for reproducibility
-    if config.seed is not None:
-        import random as _rnd
-        _rnd.seed(config.seed)
-        np.random.seed(config.seed)
-        torch.manual_seed(config.seed)
-        if torch.cuda.is_available():
-            torch.cuda.manual_seed_all(config.seed)
-
-    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
-    if verbose:
-        print(
-            f"\nInitializing {config.extractor_model.upper()} model on {config.device}..."
-        )
-
-    extractor = load_extractor(config.extractor_model, device=config.device)
-
-    # Extract features
-    if verbose:
-        print(f"\n{'='*70}")
-        print("Extracting REAL video features...")
-        print(f"{'='*70}")
-
-    real_features = load_or_compute_features(
-        videos=real_videos,
-        extractor=extractor,
-        config=config,
-        cache_path=config.cache_real_features,
-        cache_name="real_features")
-
-    if verbose:
-        print(f"\n{'='*70}")
-        print("Extracting GENERATED video features...")
-        print(f"{'='*70}")
-
-    gen_features = load_or_compute_features(videos=gen_videos,
-                                            extractor=extractor,
-                                            config=config,
-                                            cache_path=None,
-                                            cache_name="gen_features")
-
-    if verbose:
-        print(f"\nReal videos/clips: {len(real_features)}")
-        print(f"Generated videos/clips: {len(gen_features)}")
-        print(f"\n{'='*70}")
-        print("Computing statistics...")
-        print(f"{'='*70}")
-
-    mu_real, sigma_real = compute_statistics(real_features)
-    mu_gen, sigma_gen = compute_statistics(gen_features)
-
-    if verbose:
-        print(f"\n{'='*70}")
-        print("Computing Fréchet distance...")
-        print(f"{'='*70}")
-
-    fvd = compute_frechet_distance(mu_real, sigma_real, mu_gen, sigma_gen)
-
-    if verbose:
-        print(f"\n{'='*70}")
-        print(f"FVD Score ({config.extractor_model.upper()}): {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,142 +0,0 @@
-"""I3D Feature Extractor for FVD Computation"""
-
-import torch
-import torch.nn as nn
-import torch.nn.functional as F
-from pathlib import Path
-from huggingface_hub import hf_hub_download
-from tqdm import tqdm
-from contextlib import suppress
-
-
-class I3DFeatureExtractor(nn.Module):
-    """
-    I3D feature extractor for FVD computation.
-    Extracts 400-dimensional features from videos using I3D model
-    trained on Kinetics-400.
-    """
-
-    REPO_ID = 'flateon/FVD-I3D-torchscript'
-    MODEL_FILENAME = 'i3d_torchscript.pt'
-
-    def __init__(self,
-                 device: str = 'cuda',
-                 cache_dir: str | Path | None = None):
-        super().__init__()
-
-        self.device_str = device
-        if device == 'cuda' and not torch.cuda.is_available():
-            print(
-                "Warning: CUDA requested but not available – falling back to CPU"
-            )
-            self.device = torch.device('cpu')
-        else:
-            self.device = torch.device(device)
-
-        self.cache_dir: str | None
-        if cache_dir is not None:
-            self.cache_dir = str(Path(cache_dir).resolve())
-        else:
-            self.cache_dir = None  # Use HF default cache
-
-        self.model = self._load_model()
-        self.model.eval()
-
-        with suppress(Exception):
-            self.model.to(self.device)
-
-    def _load_model(self) -> torch.nn.Module:
-        """Download and load I3D TorchScript model from Hugging Face Hub."""
-        print(f"Loading I3D model from Hugging Face Hub ({self.REPO_ID})...")
-
-        try:
-            # Download model from Hugging Face Hub
-            model_path = hf_hub_download(repo_id=self.REPO_ID,
-                                         filename=self.MODEL_FILENAME,
-                                         cache_dir=self.cache_dir)
-
-            # Load directly to chosen device
-            model = torch.jit.load(model_path, map_location=self.device)
-            print("I3D model loaded successfully")
-            return model
-
-        except Exception as e:
-            raise RuntimeError(
-                f"Failed to load I3D model from Hugging Face Hub. Error: {e}\n"
-                f"Ensure you have internet connection and huggingface_hub installed:\n"
-                f"pip install huggingface_hub") from e
-
-    def preprocess(self, videos: torch.Tensor) -> torch.Tensor:
-        """
-        Preprocess videos for I3D.
-        
-        Args:
-            videos: [B, T, C, H, W], values in [0, 255]
-        
-        Returns:
-            Preprocessed videos [B, C, T, 224, 224] (normalized and resized)
-        """
-        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 needed
-        if videos.max() > 1.0:
-            videos = videos / 255.0
-
-        # Resize to 224x224 if needed
-        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)
-
-        # Convert to [B, C, T, H, W] format
-        videos = videos.permute(0, 2, 1, 3, 4).contiguous()
-
-        return videos
-
-    @torch.no_grad()
-    def extract_features(self,
-                         videos: torch.Tensor,
-                         batch_size: int = 32,
-                         verbose: bool = True) -> torch.Tensor:
-        """
-        Extract I3D features
-        
-        Args:
-            videos: [N, T, C, H, W], values in [0, 255]
-            batch_size: Batch size for processing
-            verbose: Show progress bar
-        
-        Returns:
-            Features [N, 400]
-        """
-        N = len(videos)
-        all_features = []
-
-        iterator = range(0, N, batch_size)
-        if verbose:
-            iterator = tqdm(iterator, desc="Extracting I3D features")
-
-        for i in iterator:
-            batch = videos[i:i + batch_size].to(self.device)
-            batch = self.preprocess(batch)  # Now returns [B, C, T, H, W]
-
-            # Use the HF model without rescale/resize (we handle it in preprocess)
-            features = self.model(batch,
-                                  rescale=False,
-                                  resize=False,
-                                  return_features=True)
-
-            all_features.append(features.cpu())
-
-        return torch.cat(all_features, dim=0)
-
-    def __call__(self,
-                 videos: torch.Tensor,
-                 batch_size: int = 32) -> torch.Tensor:
-        return self.extract_features(videos, batch_size=batch_size)
@@ -1,54 +0,0 @@
-import sys
-from pathlib import Path
-
-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:
-    script_dir = Path(__file__).parent.resolve()
-
-    # Define directories
-    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}")
-
-
-if __name__ == '__main__':
-    main()
@@ -1,97 +0,0 @@
-#!/usr/bin/env python3
-import sys
-from pathlib import Path
-import shutil
-import random
-from fvd import compute_fvd_with_config, FVDConfig
-
-script_path = Path(__file__).resolve()
-fastvideo_root = script_path.parent.parent.parent
-sys.path.insert(0, str(fastvideo_root))
-
-
-def split_videos(video_dir: Path, n_per_subset: int = 128, seed: int = 42):
-    subset_a = video_dir.parent / 'bair_full_subset_A'
-    subset_b = video_dir.parent / 'bair_full_subset_B'
-
-    if subset_a.exists():
-        shutil.rmtree(subset_a)
-    if subset_b.exists():
-        shutil.rmtree(subset_b)
-
-    subset_a.mkdir(parents=True)
-    subset_b.mkdir(parents=True)
-
-    videos = sorted(video_dir.glob('*.mp4'))
-
-    random.seed(seed)
-    shuffled = list(videos)
-    random.shuffle(shuffled)
-
-    needed = n_per_subset * 2
-    if len(shuffled) > needed:
-        shuffled = shuffled[:needed]
-
-    mid = len(shuffled) // 2
-
-    print(f"\nSplitting {len(shuffled)} BAIR FULL videos:")
-    print(f"  Subset A: {mid} videos")
-    print(f"  Subset B: {len(shuffled) - mid} videos")
-
-    for v in shuffled[:mid]:
-        shutil.copy2(v, subset_a / v.name)
-
-    for v in shuffled[mid:]:
-        shutil.copy2(v, subset_b / v.name)
-
-    return subset_a, subset_b, mid
-
-
-def validate_fvd(subset_a: Path, subset_b: Path, num_videos: int):
-    config = FVDConfig(num_videos=num_videos,
-                       num_frames_per_clip=16,
-                       clip_strategy='beginning',
-                       batch_size=8,
-                       device='cuda',
-                       seed=42)
-
-    print("\n" + "=" * 70)
-    print("TEST 1: Identity Test")
-    print("=" * 70)
-
-    result1 = compute_fvd_with_config(real_videos=str(subset_a),
-                                      gen_videos=str(subset_a),
-                                      config=config,
-                                      verbose=False)
-    fvd_identity = result1['fvd']
-    print(f"\nIdentity FVD: {fvd_identity:.2f}")
-
-    print("\n" + "=" * 70)
-    print("TEST 2: Real vs Real")
-    print("=" * 70)
-
-    result2 = compute_fvd_with_config(real_videos=str(subset_a),
-                                      gen_videos=str(subset_b),
-                                      config=config,
-                                      verbose=False)
-    fvd_real = result2['fvd']
-    print(f"\nReal vs Real FVD: {fvd_real:.2f}")
-
-    print("\n" + "=" * 70)
-    print("RESULTS")
-    print("=" * 70)
-    print(f"Identity:        {fvd_identity:.2f}")
-    print(f"Real vs Real:    {fvd_real:.2f}")
-
-
-def main() -> None:
-    bair_dir = Path('benchmarks/data/bair_full_videos')
-
-    subset_a, subset_b, count = split_videos(bair_dir,
-                                             n_per_subset=128,
-                                             seed=42)
-    validate_fvd(subset_a, subset_b, count)
-
-
-if __name__ == '__main__':
-    main()
@@ -1,490 +0,0 @@
-import torch
-import cv2
-import numpy as np
-from pathlib import Path
-from collections.abc import Iterator
-from tqdm import tqdm
-from enum import Enum
-
-
-class ClipSamplingStrategy(Enum):
-    """Clip sampling strategies for FVD evaluation."""
-    BEGINNING = 'beginning'  # Take first N frames (most common)
-    RANDOM = 'random'  # Random N consecutive frames
-    UNIFORM = 'uniform'  # Uniformly spaced frames across video
-    MIDDLE = 'middle'  # Middle N frames
-    SLIDING = 'sliding'  # Multiple sliding windows
-    ALL = 'all'  # All possible clips
-
-
-def _load_video_cv2(video_path: str | Path,
-                    num_frames: int | None = 16,
-                    sample_strategy: str = 'uniform') -> torch.Tensor:
-    """
-    Load video from video file using OpenCV.
-    
-    Args:
-        video_path: Path to video file (MP4, AVI, MOV, MKV)
-        num_frames: Number of frames to extract
-        sample_strategy: 'uniform' or 'random'
-    
-    Returns:
-        video: [T, C, H, W]
-    """
-    video_path = str(video_path)
-    cap = cv2.VideoCapture(video_path)
-
-    if not cap.isOpened():
-        raise RuntimeError(f"Cannot open video: {video_path}")
-
-    frames = []
-    total_frames = int(cap.get(cv2.CAP_PROP_FRAME_COUNT))
-
-    if num_frames is None:
-        # Read all available frames
-        while True:
-            ret, frame = cap.read()
-            if not ret:
-                break
-            frame = cv2.cvtColor(frame, cv2.COLOR_BGR2RGB)
-            frames.append(frame)
-
-        cap.release()
-        if len(frames) == 0:
-            raise RuntimeError(f"Video has 0 frames: {video_path}")
-
-        frames = np.stack(frames)  # [T, H, W, C]
-        frames = torch.from_numpy(frames).permute(0, 3, 1,
-                                                  2).float()  # [T, C, H, W]
-        return frames
-
-    if total_frames == 0:
-        raise RuntimeError(f"Video has 0 frames: {video_path}")
-
-    # Determine frame indices for sampling
-    if total_frames < num_frames:
-        frame_indices = list(range(
-            total_frames)) + [total_frames - 1] * (num_frames - total_frames)
-    elif sample_strategy == 'uniform':
-        frame_indices = np.linspace(0, total_frames - 1, num_frames,
-                                    dtype=int).tolist()
-    elif sample_strategy == 'random':
-        frame_indices = sorted(
-            np.random.choice(total_frames, num_frames, replace=False))
-    else:
-        raise ValueError(f"Unknown sample_strategy: {sample_strategy}")
-
-    # Extract frames
-    for idx in frame_indices:
-        cap.set(cv2.CAP_PROP_POS_FRAMES, idx)
-        ret, frame = cap.read()
-
-        if not ret:
-            if len(frames) > 0:
-                frames.append(frames[-1].copy())
-            else:
-                h = int(cap.get(cv2.CAP_PROP_FRAME_HEIGHT))
-                w = int(cap.get(cv2.CAP_PROP_FRAME_WIDTH))
-                frames.append(np.zeros((h, w, 3), dtype=np.uint8))
-            continue
-
-        frame = cv2.cvtColor(frame, cv2.COLOR_BGR2RGB)
-        frames.append(frame)
-
-    cap.release()
-
-    frames = np.stack(frames)  # [T, H, W, C]
-    frames = torch.from_numpy(frames).permute(0, 3, 1,
-                                              2).float()  # [T, C, H, W]
-
-    return frames
-
-
-def _load_video_from_frames(
-        frame_dir: str | Path,
-        num_frames: int | None = 16,
-        sample_strategy: str = 'uniform',
-        frame_extensions: list[str] | None = None) -> torch.Tensor:
-    """
-    Load video from directory of frame images.
-    
-    Args:
-        frame_dir: Directory containing frames
-        num_frames: Number of frames to sample
-        sample_strategy: 'uniform' or 'random'
-        frame_extensions: Image file extensions to look for
-    
-    Returns:
-        video: [T, C, H, W]
-    """
-    if frame_extensions is None:
-        frame_extensions = ['.jpg', '.png', '.jpeg', '.bmp']
-
-    frame_dir = Path(frame_dir)
-
-    if not frame_dir.exists():
-        raise FileNotFoundError(f"Frame directory not found: {frame_dir}")
-
-    # Find all frames
-    frame_files: list[Path] = []
-    for ext in frame_extensions:
-        frame_files.extend(frame_dir.glob(f"*{ext}"))
-
-    if len(frame_files) == 0:
-        raise ValueError(
-            f"No frames found in {frame_dir} with extensions {frame_extensions}"
-        )
-
-    frame_files = sorted(frame_files, key=lambda x: x.name)
-    total_frames = len(frame_files)
-
-    # Determine frame indices
-    if num_frames is None:
-        frame_indices = list(range(total_frames))
-    else:
-        if total_frames < num_frames:
-            frame_indices = list(range(total_frames)) + [total_frames - 1] * (
-                num_frames - total_frames)
-        elif sample_strategy == 'uniform':
-            frame_indices = np.linspace(0,
-                                        total_frames - 1,
-                                        num_frames,
-                                        dtype=int).tolist()
-        elif sample_strategy == 'random':
-            frame_indices = sorted(
-                np.random.choice(total_frames, num_frames, replace=False))
-        else:
-            raise ValueError(f"Unknown sample_strategy: {sample_strategy}")
-
-    # Load frames
-    frames = []
-    for idx in frame_indices:
-        frame_path = frame_files[idx]
-        frame = cv2.imread(str(frame_path))
-
-        if frame is None:
-            raise RuntimeError(f"Failed to load frame: {frame_path}")
-
-        frame = cv2.cvtColor(frame, cv2.COLOR_BGR2RGB)
-        frames.append(frame)
-
-    # Stack and convert to tensor
-    frames = np.stack(frames)  # [T, H, W, C]
-    frames = torch.from_numpy(frames).permute(0, 3, 1,
-                                              2).float()  # [T, C, H, W]
-
-    return frames
-
-
-def _detect_video_format(path: str | Path) -> str:
-    """
-    Detect if path is a video file or frame directory.
-    
-    Returns:
-        'video_file', 'frame_directory', or 'unknown'
-    """
-    path = Path(path)
-
-    if path.is_file():
-        return 'video_file'
-    elif path.is_dir():
-        # Check if contains image files
-        image_extensions = ['.jpg', '.jpeg', '.png', '.bmp']
-        for ext in image_extensions:
-            if list(path.glob(f"*{ext}")):
-                return 'frame_directory'
-        return 'unknown'
-    else:
-        raise ValueError(f"Path does not exist: {path}")
-
-
-def load_video_auto(video_path: str | Path,
-                    num_frames: int | None = 16,
-                    sample_strategy: str = 'uniform') -> torch.Tensor:
-    """
-    Automatically detect format and load video.
-    
-    Supports:
-    - Video files (MP4, AVI, MOV, MKV)
-    - Frame directories (JPG, PNG)
-    
-    Args:
-        video_path: Path to video file or frame directory
-        num_frames: Number of frames to extract
-        sample_strategy: 'uniform' or 'random'
-    
-    Returns:
-        video: [T, C, H, W]
-    """
-    format_type = _detect_video_format(video_path)
-
-    if format_type == 'video_file':
-        return _load_video_cv2(video_path, num_frames, sample_strategy)
-    elif format_type == 'frame_directory':
-        return _load_video_from_frames(video_path, num_frames, sample_strategy)
-    else:
-        raise ValueError(f"Unknown video format at {video_path}")
-
-
-def sample_clips_from_video(
-        video: torch.Tensor,
-        num_frames_per_clip: int = 16,
-        num_clips: int = 1,
-        strategy: str | ClipSamplingStrategy = ClipSamplingStrategy.BEGINNING,
-        frame_stride: int = 1,
-        temporal_stride: int = 1) -> list[torch.Tensor]:
-    """
-    Sample clips from a video with various strategies.
-        
-    Args:
-        video: [T, C, H, W] full video
-        num_frames_per_clip: Frames per clip
-        num_clips: Number of clips to extract
-        strategy: ClipSamplingStrategy or string ('beginning', 'random', etc.)
-        frame_stride: Skip frames (FPS control: 1=all, 2=every 2nd, 8=every 8th)
-        temporal_stride: Stride between clips for sliding window
-    
-    Returns:
-        List of clips, each [num_frames_per_clip, C, H, W]
-    
-    Examples:
-        >>> # Beginning clip (most common for FVD)
-        >>> clips = sample_clips_from_video(video, 16, strategy='beginning')
-        
-        >>> # Multiple random clips
-        >>> clips = sample_clips_from_video(video, 16, num_clips=4, strategy='random')
-        
-        >>> # Subsample FPS by 2x (every 2nd frame)
-        >>> clips = sample_clips_from_video(video, 16, frame_stride=2)
-        
-        >>> # Sliding window with overlap
-        >>> clips = sample_clips_from_video(video, 16, strategy='sliding', temporal_stride=8)
-    """
-    # Convert string to enum if needed
-    if isinstance(strategy, str):
-        strategy = ClipSamplingStrategy(strategy)
-
-    T, C, H, W = video.shape
-
-    # Apply frame stride (FPS subsampling)
-    if frame_stride > 1:
-        video = video[::frame_stride]
-        T = len(video)
-
-    effective_clip_length = num_frames_per_clip
-
-    # Handle videos shorter than clip length
-    if effective_clip_length > T:
-        pad_length = effective_clip_length - T
-        last_frame = video[-1:].repeat(pad_length, 1, 1, 1)
-        video = torch.cat([video, last_frame], dim=0)
-        T = len(video)
-
-    clips = []
-
-    if strategy == ClipSamplingStrategy.BEGINNING:
-        # Take first clip (most common for FVD evaluation)
-        clip = video[:effective_clip_length]
-        clips.append(clip)
-
-    elif strategy == ClipSamplingStrategy.MIDDLE:
-        # Take middle clip
-        start = (T - effective_clip_length) // 2
-        clip = video[start:start + effective_clip_length]
-        clips.append(clip)
-
-    elif strategy == ClipSamplingStrategy.RANDOM:
-        # Sample N random clips
-        for _ in range(num_clips):
-            if effective_clip_length == T:
-                start = 0
-            else:
-                start = np.random.randint(0, T - effective_clip_length + 1)
-            clip = video[start:start + effective_clip_length]
-            clips.append(clip)
-
-    elif strategy == ClipSamplingStrategy.UNIFORM:
-        # Uniformly spaced clips
-        if num_clips == 1:
-            # Single clip from middle
-            start = (T - effective_clip_length) // 2
-            clip = video[start:start + effective_clip_length]
-            clips.append(clip)
-        else:
-            # Multiple uniformly spaced clips
-            step = (T - effective_clip_length) / (num_clips -
-                                                  1) if num_clips > 1 else 0
-            for i in range(num_clips):
-                start = int(i * step)
-                start = min(start, T - effective_clip_length)
-                clip = video[start:start + effective_clip_length]
-                clips.append(clip)
-
-    elif strategy == ClipSamplingStrategy.SLIDING:
-        # Sliding window with stride
-        for start in range(0, T - effective_clip_length + 1, temporal_stride):
-            clip = video[start:start + effective_clip_length]
-            clips.append(clip)
-            if len(clips) >= num_clips:
-                break
-
-    elif strategy == ClipSamplingStrategy.ALL:
-        # All possible clips (overlapping)
-        for start in range(T - effective_clip_length + 1):
-            clip = video[start:start + effective_clip_length]
-            clips.append(clip)
-
-    else:
-        raise ValueError(f"Unknown strategy: {strategy}")
-
-    return clips
-
-
-def load_video_clips_streaming(directory: str | Path,
-                               num_frames: int = 16,
-                               max_videos: int | None = None,
-                               clip_strategy: str
-                               | ClipSamplingStrategy = 'beginning',
-                               frame_stride: int = 1,
-                               num_clips_per_video: int = 1,
-                               video_extensions: list[str] | None = None,
-                               support_frame_dirs: bool = True,
-                               target_size: tuple[int, int] | None = (224, 224),
-                               verbose: bool = True) -> Iterator[torch.Tensor]:
-    """
-    This generator yields clips one-by-one instead of loading all videos into RAM.
-    Perfect for large datasets where memory is limited.
-    
-    Args:
-        directory: Path to directory with videos
-        num_frames: Frames per clip
-        max_videos: Max videos to load
-        clip_strategy: 'beginning', 'random', 'uniform', etc.
-        frame_stride: Frame skip (1=all, 2=every 2nd, 8=every 8th)
-        num_clips_per_video: Number of clips per video
-        video_extensions: Video file extensions
-        support_frame_dirs: Also load frame directories
-        target_size: Resize clips to (H, W). If None, keep original size.
-        verbose: Show progress
-    
-    Yields:
-        clip: [T, C, H, W] individual clips
-    
-    Example:
-        >>> for clip in load_video_clips_streaming('data/videos/', num_frames=16):
-        >>>     features = model.extract_features(clip.unsqueeze(0))
-        >>>     # Process one clip at a time - low memory usage!
-    """
-    if video_extensions is None:
-        video_extensions = ['.mp4', '.avi', '.mov', '.mkv']
-
-    directory = Path(directory)
-
-    if not directory.exists():
-        raise FileNotFoundError(f"Directory not found: {directory}")
-
-    # Find video paths
-    video_paths: list[Path] = []
-
-    # Find video files
-    for ext in video_extensions:
-        video_paths.extend(directory.glob(f"**/*{ext}"))
-
-    # Find frame directories if enabled
-    if support_frame_dirs:
-        for subdir in directory.iterdir():
-            if subdir.is_dir():
-                # Check if it contains frames
-                image_extensions = ['.jpg', '.jpeg', '.png', '.bmp']
-                for ext in image_extensions:
-                    if list(subdir.glob(f"*{ext}")):
-                        video_paths.append(subdir)
-                        break
-
-    if len(video_paths) == 0:
-        raise ValueError(f"No videos found in {directory}")
-
-    video_paths = sorted(video_paths)
-
-    if max_videos is not None:
-        video_paths = video_paths[:max_videos]
-
-    if verbose:
-        print(f"Found {len(video_paths)} videos in {directory}")
-        if num_clips_per_video > 1:
-            print(f"Extracting {num_clips_per_video} clips per video...")
-        if frame_stride > 1:
-            print(f"Subsampling frames with stride {frame_stride}...")
-        if target_size:
-            print(f"Resizing clips to {target_size}...")
-
-    # Track statistics
-    failed_count = 0
-    total_clips = 0
-
-    iterator = tqdm(video_paths,
-                    desc="Loading videos") if verbose else video_paths
-
-    for video_path in iterator:
-        try:
-            # Load full video
-            video = load_video_auto(video_path,
-                                    num_frames=None,
-                                    sample_strategy='uniform')
-
-            # Sample clips from video
-            clips = sample_clips_from_video(video,
-                                            num_frames_per_clip=num_frames,
-                                            num_clips=num_clips_per_video,
-                                            strategy=clip_strategy,
-                                            frame_stride=frame_stride)
-
-            if target_size is not None:
-                resized_clips = []
-                for clip in clips:
-                    T, C, H, W = clip.shape
-                    if target_size != (H, W):
-                        # Resize to target size
-                        clip = clip.contiguous(
-                        )  # Fix non-contiguous tensors first
-                        clip_flat = clip.view(T * C, H,
-                                              W).unsqueeze(0)  # [1, T*C, H, W]
-                        clip_resized = torch.nn.functional.interpolate(
-                            clip_flat,
-                            size=target_size,
-                            mode='bilinear',
-                            align_corners=False)
-                        clip = clip_resized.squeeze(0).view(
-                            T, C, target_size[0],
-                            target_size[1])  # Back to [T, C, H, W]
-                    resized_clips.append(clip)
-                clips = resized_clips
-
-            # Yield clips one by one
-            for clip in clips:
-                yield clip
-                total_clips += 1
-
-            # Free memory
-            del video, clips
-
-        except Exception as e:
-            failed_count += 1
-            if verbose:
-                print(f"\nWarning: Failed to load {video_path}: {e}")
-            continue
-
-    # Validate
-    if total_clips == 0:
-        raise RuntimeError(f"Failed to load any videos from {directory}")
-
-    failure_rate = failed_count / len(video_paths)
-    if failure_rate > 0.1:  # More than 10% failed
-        print(
-            f"\nWARNING: {failure_rate:.1%} of videos failed to load ({failed_count}/{len(video_paths)})"
-        )
-
-    if verbose:
-        print(
-            f"\nSuccessfully loaded {total_clips} clips from {len(video_paths) - failed_count} videos"
-        )
@@ -1,7 +0,0 @@
-#!/bin/bash
-
-# 1. Install missing dependency
-pip install -q opencv-python-headless transformers huggingface_hub
-
-# 2. Run FVD script
-python benchmarks/fvd/run_fvd.py
@@ -1,4 +0,0 @@
-#!/bin/bash
-
-# 1. Install missing dependency
-pip install -q opencv-python-headless
@@ -0,0 +1,113 @@
+
+
+# Attention Kernel Used in FastVideo
+
+
+
+## Video Sparse Attention (VSA)
+
+### Installation
+We support H100 (via TK) and any other GPU (via triton) for VSA.
+```bash
+git submodule update --init --recursive
+python setup_vsa.py install
+```
+
+
+If you encounter error during installation, try below:
+Install C++20 for ThunderKittens:
+```bash
+sudo apt update
+sudo apt install gcc-11 g++-11
+
+sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
+
+sudo apt update
+sudo apt install clang-11
+```
+(If you use CUDA12.8)
+```bash
+export CUDA_HOME=/usr/local/cuda-12.8
+export PATH=${CUDA_HOME}/bin:${PATH} 
+export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
+```
+
+### Verify if you have successfully installed
+
+```bash
+# test numerical
+python tests/test_vsa.py
+# (For H100) test speed
+python benchmarks/bench_vsa_hopper.py
+```
+bench_vsa_hopper.py should print something like this:
+```bash
+Using topk=76 kv blocks per q block (out of 768 total kv blocks)
+
+=== BLOCK SPARSE ATTENTION BENCHMARK ===
+Block Sparse Forward  - TFLOPS: 5622.26
+Block Sparse Backward - TFLOPS: 3865.68
+```
+
+
+## Sliding Tile Attention (STA)
+We only support H100 for STA.
+```bash
+git submodule update --init --recursive
+python setup_sta.py install
+```
+
+
+
+
+###  Usage
+End-2-end inference with FastVideo:
+```bash
+bash scripts/inference/v1_inference_wan_STA.sh
+```
+
+If you want to use sliding tile attention in your custom model:
+```python
+from st_attn import sliding_tile_attention
+# assuming video size (T, H, W) = (30, 48, 80), text tokens = 256 with padding. 
+# q, k, v: [batch_size, num_heads, seq_length, head_dim], seq_length = T*H*W + 256
+# a tile is a cube of size (6, 8, 8)
+# window_size in tiles: [(window_t, window_h, window_w), (..)...]. For example, window size (3, 3, 3) means a query can attend to (3x6, 3x8, 3x8) = (18, 24, 24) tokens out of the total 30x48x80 video.
+# text_length: int ranging from 0 to 256
+# If your attention contains text token (Hunyuan)
+out = sliding_tile_attention(q, k, v, window_size, text_length)
+# If your attention does not contain text token (StepVideo)
+out = sliding_tile_attention(q, k, v, window_size, 0, False)
+```
+
+
+### Test
+```bash
+python tests/test_sta.py # test STA
+python tests/test_vsa.py # test VSA
+```
+### Benchmark
+```bash
+python benchmarks/bench_sta.py
+```
+
+
+### How Does STA Work?
+We give a demo for 2D STA with window size (6,6) operating on a (10, 10) image. 
+
+
+https://github.com/user-attachments/assets/f3b6dd79-7b43-4b60-a0fa-3d6495ec5747
+
+## Why is STA Fast?
+2D/3D Sliding Window Attention (SWA) creates many mixed blocks in the attention map. Even though mixed blocks have less output value,a mixed block is significantly slower than a dense block due to the GPU-unfriendly masking operation. 
+
+STA removes mixed blocks.
+
+
+<div align="center">
+<img src=../../assets/sliding_tile_attn_map.png width="80%"/>
+</div>
+
+## Acknowledgement
+
+We learned or reuse code from FlexAtteniton, NATEN, and ThunderKittens.
@@ -0,0 +1,145 @@
+import os
+from collections import defaultdict
+
+import matplotlib.pyplot as plt
+import numpy as np
+import torch
+from st_attn import sliding_tile_attention
+from triton.testing import do_bench
+
+
+def flops(batch, seqlen, nheads, headdim, causal, mode="fwd"):
+    assert mode in ["fwd", "bwd", "fwd_bwd"]
+    f = 4 * batch * seqlen**2 * nheads * headdim // (2 if causal else 1)
+    return f if mode == "fwd" else (2.5 * f if mode == "bwd" else 3.5 * f)
+
+
+def compute_TFLOPS(flops, ms):
+    flops = flops / 1e12
+    ms = ms / 1e3
+    return flops / ms
+
+
+def benchmark_attention(configurations):
+    results = {'fwd': defaultdict(list), 'bwd': defaultdict(list)}
+
+    for B, H, N, D, causal, dit_seq_shape, window_size in configurations:
+        print("=" * 60)
+        print(f"Timing forward and backward pass for B={B}, H={H}, N={N}, D={D}, causal={causal}")
+
+        q = torch.randn(B, H, N, D, dtype=torch.bfloat16, device='cuda', requires_grad=False).contiguous()
+        k = torch.randn(B, H, N, D, dtype=torch.bfloat16, device='cuda', requires_grad=False).contiguous()
+        v = torch.randn(B, H, N, D, dtype=torch.bfloat16, device='cuda', requires_grad=False).contiguous()
+
+        # grad_output = torch.randn_like(q, requires_grad=False).contiguous()
+        # qg = torch.zeros_like(q, requires_grad=False, dtype=torch.float).contiguous()
+        # kg = torch.zeros_like(k, requires_grad=False, dtype=torch.float).contiguous()
+        # vg = torch.zeros_like(v, requires_grad=False, dtype=torch.float).contiguous()
+
+
+        # # Warmup for forward pass
+        # for _ in range(10):
+        #     o = sliding_tile_attention(q, k, v, [[3, 6, 10]] * 24, 0, False, dit_seq_shape)
+
+        # # Time the forward pass
+        # for i in range(10):
+        #     start_events_fwd[i].record()
+        #     o = sliding_tile_attention(q, k, v, [[3, 6, 10]] * 24, 0, False, dit_seq_shape)
+        #     end_events_fwd[i].record()
+        ms = do_bench(lambda: sliding_tile_attention(q, k, v, [window_size] * 24, 0, False, dit_seq_shape))
+
+        # times_fwd = [s.elapsed_time(e) for s, e in zip(start_events_fwd, end_events_fwd)]
+        # time_us_fwd = np.mean(times_fwd) * 1000
+
+        tflops_fwd = compute_TFLOPS(flops(B, N, H, D, causal, 'fwd'), ms)
+        results['fwd'][(D, causal)].append((N, tflops_fwd))
+
+        print(f"Average time for forward pass (ms): {ms:.2f}")
+        print(f"Average TFLOPS: {tflops_fwd}")
+        print("-" * 60)
+
+        # torch.cuda.empty_cache()
+        # torch.cuda.synchronize()
+
+        # # Prepare for timing backward pass
+        # start_events_bwd = [torch.cuda.Event(enable_timing=True) for _ in range(10)]
+        # end_events_bwd = [torch.cuda.Event(enable_timing=True) for _ in range(10)]
+
+        # # Warmup for backward pass
+        # for _ in range(10):
+        #     qg, kg, vg = tk.mha_backward(q, k, v, o, l_vec, grad_output, causal)
+
+        # # Time the backward pass
+        # for i in range(10):
+        #     start_events_bwd[i].record()
+        #     qg, kg, vg = tk.mha_backward(q, k, v, o, l_vec, grad_output, causal)
+        #     end_events_bwd[i].record()
+
+        # torch.cuda.synchronize()
+        # times_bwd = [s.elapsed_time(e) for s, e in zip(start_events_bwd, end_events_bwd)]
+        # time_us_bwd = np.mean(times_bwd) * 1000
+
+        # tflops_bwd = compute_TFLOPS(flops(B, N, H, D, causal, 'bwd'), ms)
+        # results['bwd'][(D, causal)].append((N, tflops_bwd))
+
+        # print(f"Average time for backward pass(ms): {ms:.2f}")
+        # print(f"Average TFLOPS: {tflops_bwd}")
+        # print("=" * 60)
+
+    return results
+
+
+def plot_results(results):
+    os.makedirs('benchmark_results', exist_ok=True)
+    for mode in ['fwd', 'bwd']:
+        for (D, causal), values in results[mode].items():
+            seq_lens = [x[0] for x in values]
+            tflops = [x[1] for x in values]
+
+            plt.figure(figsize=(10, 6))
+            bars = plt.bar(range(len(seq_lens)), tflops, tick_label=seq_lens)
+            plt.xlabel('Sequence Length')
+            plt.ylabel('TFLOPS')
+            plt.title(f'{mode.upper()} Pass - Head Dim: {D}, Causal: {causal}')
+            plt.grid(True)
+
+            # Adding the numerical y value on top of each bar
+            for bar in bars:
+                yval = bar.get_height()
+                plt.text(bar.get_x() + bar.get_width() / 2, yval, round(yval, 2), ha='center', va='bottom')
+
+            filename = f'benchmark_results/{mode}_D{D}_causal{causal}.png'
+            plt.savefig(filename)
+            plt.close()
+
+
+# Example list of configurations to test
+configurations = [
+    (2, 24, 69120, 128, False, '18x48x80', [3, 6, 10]),
+    (2, 24, 69120, 128, True, '18x48x80', [3, 6, 10]),
+    (2, 24, 82944, 128, False, '36x48x48', [3, 3, 6]), # Stepvideo
+    (2, 24, 82944, 128, True, '36x48x48', [3, 3, 6]),
+    # (16, 16, 768*16,  128, False),
+    # (16, 16, 768*2,  128, False),
+    # (16, 16, 768*4,  128, False),
+    # (16, 16, 768*8,  128, False),
+    # (16, 16, 768*16, 128, False),
+    # (16, 16, 768,    128, True),
+    # (16, 16, 768*2,  128, True),
+    # (16, 16, 768*4,  128, True),
+    # (16, 16, 768*8,  128, True),
+    # (16, 16, 768*16, 128, True),
+    # (16, 32, 768,    64,  False),
+    # (16, 32, 768*2,  64,  False),
+    # (16, 32, 768*4,  64,  False),
+    # (16, 32, 768*8,  64,  False),
+    # (16, 32, 768*16, 64,  False),
+    # (16, 32, 768,    64,  True),
+    # (16, 32, 768*2,  64,  True),
+    # (16, 32, 768*4,  64,  True),
+    # (16, 32, 768*8,  64,  True),
+    # (16, 32, 768*16, 64,  True),
+]
+
+results = benchmark_attention(configurations)
+# plot_results(results)
@@ -0,0 +1,224 @@
+import torch
+import argparse
+from triton.testing import do_bench
+from vsa import block_sparse_fwd, block_sparse_bwd
+from vsa import BLOCK_M, BLOCK_N
+import  triton
+import numpy as np
+import random
+
+def set_seed(seed: int = 42):
+    # Python random module
+    random.seed(seed)
+
+    # NumPy
+    np.random.seed(seed)
+
+    # PyTorch
+    torch.manual_seed(seed)
+    torch.cuda.manual_seed(seed)
+    torch.cuda.manual_seed_all(seed)  # if using multi-GPU
+
+def parse_arguments():
+    parser = argparse.ArgumentParser(description='Benchmark Block Sparse Attention')
+    parser.add_argument('--batch_size', type=int, default=1, help='Batch size')
+    parser.add_argument('--num_heads', type=int, default=12, help='Number of heads')
+    parser.add_argument('--head_dim', type=int, default=128, help='Head dimension')
+    parser.add_argument('--topk', type=int, default=None, help='Number of kv blocks each q block attends to')
+    parser.add_argument('--seq_lengths', type=int, nargs='+', default=[49152], help='Sequence lengths to benchmark')
+    return parser.parse_args()
+
+def create_input_tensors(batch, head, seq_len, headdim):
+    """Create random input tensors for attention."""
+    q = torch.randn(batch, head, seq_len, headdim, dtype=torch.bfloat16, device="cuda")
+    k = torch.randn(batch, head, seq_len, headdim, dtype=torch.bfloat16, device="cuda")
+    v = torch.randn(batch, head, seq_len, headdim, dtype=torch.bfloat16, device="cuda")
+    return q, k, v
+
+def generate_block_sparse_pattern(bs, h, num_q_blocks, num_kv_blocks, k, device="cuda"):
+    """
+    Generate a block sparse pattern where each q block attends to exactly k kv blocks.
+    
+    Args:
+        bs: batch size
+        h: number of heads
+        num_q_blocks: number of query blocks
+        num_kv_blocks: number of key-value blocks
+        k: number of kv blocks each q block attends to
+        device: device to create tensors on
+        
+    Returns:
+        q2k_block_sparse_index: [bs, h, num_q_blocks, k]
+            Contains the indices of kv blocks that each q block attends to.
+        q2k_block_sparse_num: [bs, h, num_q_blocks]
+            Contains the number of kv blocks that each q block attends to (all equal to k).
+        k2q_block_sparse_index: [bs, h, num_kv_blocks, num_q_blocks]
+            Contains the indices of q blocks that attend to each kv block.
+        k2q_block_sparse_num: [bs, h, num_kv_blocks]
+            Contains the number of q blocks that attend to each kv block.
+        block_sparse_mask: [bs, h, num_q_blocks, num_kv_blocks]
+            Binary mask where 1 indicates attention connection.
+    """
+    # Ensure k is not larger than num_kv_blocks
+    k = min(k, num_kv_blocks)
+    
+    # Create random scores for sampling
+    scores = torch.rand(bs, h, num_q_blocks, num_kv_blocks, device=device)
+    
+    # Get top-k indices for each q block
+    _, q2k_block_sparse_index = torch.topk(scores, k, dim=-1)
+    q2k_block_sparse_index = q2k_block_sparse_index.to(torch.int32)
+    
+    # sort q2k_block_sparse_index
+    q2k_block_sparse_index, _ = torch.sort(q2k_block_sparse_index, dim=-1)
+
+    # All q blocks attend to exactly k kv blocks
+    q2k_block_sparse_num = torch.full((bs, h, num_q_blocks), k, dtype=torch.int32, device=device)
+    
+    # Create the corresponding mask
+    block_sparse_mask = torch.zeros(bs, h, num_q_blocks, num_kv_blocks, dtype=torch.bool, device=device)
+    
+    # Fill in the mask based on the indices
+    for b in range(bs):
+        for head in range(h):
+            for q_idx in range(num_q_blocks):
+                kv_indices = q2k_block_sparse_index[b, head, q_idx]
+                block_sparse_mask[b, head, q_idx, kv_indices] = True
+    
+    # Create the reverse mapping (k2q)
+    # First, initialize lists to collect q indices for each kv block
+    k2q_indices_list = [[[] for _ in range(num_kv_blocks)] for _ in range(bs * h)]
+    
+    # Populate the lists based on q2k mapping
+    for b in range(bs):
+        for head in range(h):
+            flat_idx = b * h + head
+            for q_idx in range(num_q_blocks):
+                kv_indices = q2k_block_sparse_index[b, head, q_idx].tolist()
+                for kv_idx in kv_indices:
+                    k2q_indices_list[flat_idx][kv_idx].append(q_idx)
+    
+    # Find the maximum number of q blocks that attend to any kv block
+    max_q_per_kv = 0
+    for flat_idx in range(bs * h):
+        for kv_idx in range(num_kv_blocks):
+            max_q_per_kv = max(max_q_per_kv, len(k2q_indices_list[flat_idx][kv_idx]))
+    
+    # Create tensors for k2q mapping
+    k2q_block_sparse_index = torch.full((bs, h, num_kv_blocks, max_q_per_kv), -1, 
+                                        dtype=torch.int32, device=device)
+    k2q_block_sparse_num = torch.zeros((bs, h, num_kv_blocks), 
+                                       dtype=torch.int32, device=device)
+    
+    # Fill the tensors
+    for b in range(bs):
+        for head in range(h):
+            flat_idx = b * h + head
+            for kv_idx in range(num_kv_blocks):
+                q_indices = k2q_indices_list[flat_idx][kv_idx]
+                num_q = len(q_indices)
+                k2q_block_sparse_num[b, head, kv_idx] = num_q
+                if num_q > 0:
+                    k2q_block_sparse_index[b, head, kv_idx, :num_q] = torch.tensor(
+                        q_indices, dtype=torch.int32, device=device)
+                
+    return q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, block_sparse_mask
+
+def benchmark_block_sparse_attention(q, k, v, q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, flops):
+    """Benchmark block sparse attention forward and backward passes."""
+    print("\n=== BLOCK SPARSE ATTENTION BENCHMARK ===")
+    
+    # Forward pass
+    # Warm-up run
+    variable_block_sizes = torch.ones(q2k_block_sparse_index.shape[2], device=q.device).int() * BLOCK_M
+    o, l_vec = block_sparse_fwd(q, k, v, q2k_block_sparse_index, q2k_block_sparse_num, variable_block_sizes)
+    torch.cuda.synchronize()
+    
+    # Benchmark forward
+    fwd_time = do_bench(
+        lambda: block_sparse_fwd(q, k, v, q2k_block_sparse_index, q2k_block_sparse_num, variable_block_sizes),
+        warmup=5,
+        rep=20,
+        quantiles=None
+    )
+    
+    sparse_tflops = flops / fwd_time * 1e-12 * 1e3
+    print(f"Block Sparse Forward - TFLOPS: {sparse_tflops:.2f}")
+    
+    # Backward pass
+    grad_output = torch.randn_like(o)
+    
+    # Warm-up runs
+    for _ in range(5):
+        block_sparse_bwd(q, k, v, o, l_vec, grad_output, k2q_block_sparse_index, k2q_block_sparse_num, variable_block_sizes)
+    torch.cuda.synchronize()
+    
+    # Benchmark backward
+    bwd_time = do_bench(
+        lambda: block_sparse_bwd(q, k, v, o, l_vec, grad_output, k2q_block_sparse_index, k2q_block_sparse_num, variable_block_sizes),
+        warmup=5,
+        rep=20,
+        quantiles=None
+    )
+    bwd_flops = 2.5 * flops  # Approximation
+
+    sparse_bwd_tflops = bwd_flops / bwd_time * 1e-12 * 1e3
+    print(f"Block Sparse Backward - TFLOPS: {sparse_bwd_tflops:.2f}")
+    
+    return sparse_tflops, sparse_bwd_tflops
+
+def main():
+    args = parse_arguments()
+
+    set_seed(42)
+    
+    # Extract parameters
+    batch = args.batch_size
+    head = args.num_heads
+    headdim = args.head_dim
+    
+    print(f"Block Sparse Attention Benchmark")
+    print(f"batch: {batch}, head: {head}, headdim: {headdim}")
+    
+    # Test with different sequence lengths
+    for seq_len in args.seq_lengths:
+        # Skip very long sequences if they might cause OOM
+        if seq_len > 16384 and batch > 1:
+            continue
+            
+        print("="*100)
+        print(f"\nSequence length: {seq_len}")
+        
+        # Calculate theoretical FLOPs for attention
+        flops = 4 * batch * head * headdim * seq_len * seq_len
+        
+        # Create input tensors
+        q, k, v = create_input_tensors(batch, head, seq_len, headdim)
+        
+        # Setup block sparse parameters
+        num_q_blocks = seq_len // BLOCK_M
+        num_kv_blocks = seq_len // BLOCK_N
+        
+        # Determine k value (number of kv blocks per q block)
+        topk = args.topk
+        if topk is None:
+            topk = num_kv_blocks // 10  # Default to ~90% sparsity if k is not specified
+        topk = max(1, topk)       
+        print(f"Using topk={topk} kv blocks per q block (out of {num_kv_blocks} total kv blocks)")
+        
+        # Generate block sparse pattern
+        q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, _ = generate_block_sparse_pattern(
+            batch, head, num_q_blocks, num_kv_blocks, topk, device="cuda")
+        
+        # Benchmark block sparse attention
+        sparse_fwd, sparse_bwd = benchmark_block_sparse_attention(
+            q, k, v, q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, flops
+        )
+        
+        # Print results
+        print("\n=== PERFORMANCE RESULTS ===")
+        print(f"Block Sparse Forward - TFLOPS: {sparse_fwd:.2f}")
+        print(f"Block Sparse Backward - TFLOPS: {sparse_bwd:.2f}")
+
+if __name__ == "__main__":
+    main()
@@ -0,0 +1,217 @@
+import torch
+import argparse
+import triton.testing
+from vsa import block_sparse_attn
+from vsa import BLOCK_M, BLOCK_N
+
+import numpy as np
+import random
+
+def set_seed(seed: int = 42):
+    # Python random module
+    random.seed(seed)
+
+    # NumPy
+    np.random.seed(seed)
+
+    # PyTorch
+    torch.manual_seed(seed)
+    torch.cuda.manual_seed(seed)
+    torch.cuda.manual_seed_all(seed)  # if using multi-GPU
+
+def parse_arguments():
+    parser = argparse.ArgumentParser(description='Benchmark Block Sparse Attention')
+    parser.add_argument('--batch_size', type=int, default=1, help='Batch size')
+    parser.add_argument('--num_heads', type=int, default=12, help='Number of heads')
+    parser.add_argument('--head_dim', type=int, default=64, help='Head dimension')
+    parser.add_argument('--topk', type=int, default=None, help='Number of kv blocks each q block attends to')
+    parser.add_argument('--seq_lengths', type=int, nargs='+', default=[49152], help='Sequence lengths to benchmark')
+    return parser.parse_args()
+
+def create_input_tensors(batch, head, seq_len, headdim):
+    """Create random input tensors for attention."""
+    q = torch.randn(batch, head, seq_len, headdim, dtype=torch.bfloat16, device="cuda")
+    k = torch.randn(batch, head, seq_len, headdim, dtype=torch.bfloat16, device="cuda")
+    v = torch.randn(batch, head, seq_len, headdim, dtype=torch.bfloat16, device="cuda")
+    return q, k, v
+
+def generate_block_sparse_pattern(bs, h, num_q_blocks, num_kv_blocks, k, device="cuda"):
+    """
+    Generate a block sparse pattern where each q block attends to exactly k kv blocks.
+    
+    Args:
+        bs: batch size
+        h: number of heads
+        num_q_blocks: number of query blocks
+        num_kv_blocks: number of key-value blocks
+        k: number of kv blocks each q block attends to
+        device: device to create tensors on
+        
+    Returns:
+        q2k_block_sparse_index: [bs, h, num_q_blocks, k]
+            Contains the indices of kv blocks that each q block attends to.
+        q2k_block_sparse_num: [bs, h, num_q_blocks]
+            Contains the number of kv blocks that each q block attends to (all equal to k).
+        k2q_block_sparse_index: [bs, h, num_kv_blocks, num_q_blocks]
+            Contains the indices of q blocks that attend to each kv block.
+        k2q_block_sparse_num: [bs, h, num_kv_blocks]
+            Contains the number of q blocks that attend to each kv block.
+        block_sparse_mask: [bs, h, num_q_blocks, num_kv_blocks]
+            Binary mask where 1 indicates attention connection.
+    """
+    # Ensure k is not larger than num_kv_blocks
+    k = min(k, num_kv_blocks)
+    
+    # Create random scores for sampling
+    scores = torch.rand(bs, h, num_q_blocks, num_kv_blocks, device=device)
+    
+    # Get top-k indices for each q block
+    _, q2k_block_sparse_index = torch.topk(scores, k, dim=-1)
+    q2k_block_sparse_index = q2k_block_sparse_index.to(torch.int32)
+    
+    # sort q2k_block_sparse_index
+    q2k_block_sparse_index, _ = torch.sort(q2k_block_sparse_index, dim=-1)
+
+    # All q blocks attend to exactly k kv blocks
+    q2k_block_sparse_num = torch.full((bs, h, num_q_blocks), k, dtype=torch.int32, device=device)
+    
+    # Create the corresponding mask
+    block_sparse_mask = torch.zeros(bs, h, num_q_blocks, num_kv_blocks, dtype=torch.bool, device=device)
+    
+    # Fill in the mask based on the indices
+    for b in range(bs):
+        for head in range(h):
+            for q_idx in range(num_q_blocks):
+                kv_indices = q2k_block_sparse_index[b, head, q_idx]
+                block_sparse_mask[b, head, q_idx, kv_indices] = True
+    
+    # Create the reverse mapping (k2q)
+    # First, initialize lists to collect q indices for each kv block
+    k2q_indices_list = [[[] for _ in range(num_kv_blocks)] for _ in range(bs * h)]
+    
+    # Populate the lists based on q2k mapping
+    for b in range(bs):
+        for head in range(h):
+            flat_idx = b * h + head
+            for q_idx in range(num_q_blocks):
+                kv_indices = q2k_block_sparse_index[b, head, q_idx].tolist()
+                for kv_idx in kv_indices:
+                    k2q_indices_list[flat_idx][kv_idx].append(q_idx)
+    
+    # Find the maximum number of q blocks that attend to any kv block
+    max_q_per_kv = 0
+    for flat_idx in range(bs * h):
+        for kv_idx in range(num_kv_blocks):
+            max_q_per_kv = max(max_q_per_kv, len(k2q_indices_list[flat_idx][kv_idx]))
+    
+    # Create tensors for k2q mapping
+    k2q_block_sparse_index = torch.full((bs, h, num_kv_blocks, max_q_per_kv), -1, 
+                                        dtype=torch.int32, device=device)
+    k2q_block_sparse_num = torch.zeros((bs, h, num_kv_blocks), 
+                                       dtype=torch.int32, device=device)
+    
+    # Fill the tensors
+    for b in range(bs):
+        for head in range(h):
+            flat_idx = b * h + head
+            for kv_idx in range(num_kv_blocks):
+                q_indices = k2q_indices_list[flat_idx][kv_idx]
+                num_q = len(q_indices)
+                k2q_block_sparse_num[b, head, kv_idx] = num_q
+                if num_q > 0:
+                    k2q_block_sparse_index[b, head, kv_idx, :num_q] = torch.tensor(
+                        q_indices, dtype=torch.int32, device=device)
+                
+    return q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, block_sparse_mask
+
+def benchmark_block_sparse_attention(q, k, v, q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, flops):
+    """Benchmark block sparse attention forward+backward pass."""
+    print("\n=== BLOCK SPARSE ATTENTION FORWARD+BACKWARD BENCHMARK ===")
+    
+    # Combined forward+backward pass
+    # Warm-up run
+    q_fwd = q.clone().requires_grad_(True)
+    k_fwd = k.clone().requires_grad_(True)
+    v_fwd = v.clone().requires_grad_(True)
+    o = block_sparse_attn(q_fwd, k_fwd, v_fwd, q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num)
+    grad_output = torch.randn_like(o)
+    o.backward(grad_output)
+    torch.cuda.synchronize()
+    
+    # Benchmark forward+backward
+    def forward_backward_fn():
+        q_fwd = q.clone().requires_grad_(True)
+        k_fwd = k.clone().requires_grad_(True)
+        v_fwd = v.clone().requires_grad_(True)
+        o = block_sparse_attn(q_fwd, k_fwd, v_fwd, q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num)
+        grad_output = torch.randn_like(o)
+        o.backward(grad_output)
+    
+    total_time = triton.testing.do_bench(
+        forward_backward_fn,
+        warmup=25,
+        rep=100,
+        return_mode='mean'
+    )
+    
+    # Total flops for forward + backward (forward + 2.5x backward approximation)
+    total_flops = flops + 2.5 * flops  # 3.5x the forward flops
+    sparse_tflops = total_flops / total_time * 1e-12 * 1e3
+    print(f"Block Sparse Forward+Backward - TFLOPS: {sparse_tflops:.2f}")
+    
+    return sparse_tflops
+
+def main():
+    args = parse_arguments()
+
+    set_seed(42)
+    
+    # Extract parameters
+    batch = args.batch_size
+    head = args.num_heads
+    headdim = args.head_dim
+    
+    print(f"Block Sparse Attention Benchmark")
+    print(f"batch: {batch}, head: {head}, headdim: {headdim}")
+    
+    # Test with different sequence lengths
+    for seq_len in args.seq_lengths:
+        # Skip very long sequences if they might cause OOM
+        if seq_len > 16384 and batch > 1:
+            continue
+            
+        print("="*100)
+        print(f"\nSequence length: {seq_len}")
+        
+        # Calculate theoretical FLOPs for attention
+        flops = 4 * batch * head * headdim * seq_len * seq_len
+        
+        # Create input tensors
+        q, k, v = create_input_tensors(batch, head, seq_len, headdim)
+        
+        # Setup block sparse parameters
+        num_q_blocks = seq_len // BLOCK_M
+        num_kv_blocks = seq_len // BLOCK_N
+        
+        # Determine k value (number of kv blocks per q block)
+        topk = args.topk
+        if topk is None:
+            topk = num_kv_blocks // 10  # Default to ~90% sparsity if k is not specified
+        topk = max(1, topk)       
+        print(f"Using topk={topk} kv blocks per q block (out of {num_kv_blocks} total kv blocks)")
+        
+        # Generate block sparse pattern
+        q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, _ = generate_block_sparse_pattern(
+            batch, head, num_q_blocks, num_kv_blocks, topk, device="cuda")
+        
+        # Benchmark block sparse attention
+        sparse_fwd = benchmark_block_sparse_attention(
+            q, k, v, q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, flops
+        )
+        
+        # Print results
+        print("\n=== PERFORMANCE RESULTS ===")
+        print(f"Block Sparse Forward+Backward - TFLOPS: {sparse_fwd:.2f}")
+
+if __name__ == "__main__":
+    main()
@@ -0,0 +1,2 @@
+recursive-include tk *
+include config_sta.py
@@ -0,0 +1,87 @@
+
+# Attention Kernel Used in FastVideo
+
+## Sliding Tile Attention (STA)
+We only support H100 for STA.
+
+### Installation
+```bash
+pip install st_attn
+``` 
+
+Install from source:
+
+```bash
+git submodule update --init --recursive
+python setup.py install
+```
+
+If you encounter error during installation, try below:
+Install C++20 for ThunderKittens:
+```bash
+sudo apt update
+sudo apt install gcc-11 g++-11
+
+sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
+
+sudo apt update
+sudo apt install clang-11
+```
+(If you use CUDA12.8)
+```bash
+export CUDA_HOME=/usr/local/cuda-12.8
+export PATH=${CUDA_HOME}/bin:${PATH} 
+export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
+```
+
+###  Usage
+End-2-end inference with FastVideo:
+```bash
+bash scripts/inference/v1_inference_wan_STA.sh
+```
+
+If you want to use sliding tile attention in your custom model:
+```python
+from st_attn import sliding_tile_attention
+# assuming video size (T, H, W) = (30, 48, 80), text tokens = 256 with padding. 
+# q, k, v: [batch_size, num_heads, seq_length, head_dim], seq_length = T*H*W + 256
+# a tile is a cube of size (6, 8, 8)
+# window_size in tiles: [(window_t, window_h, window_w), (..)...]. For example, window size (3, 3, 3) means a query can attend to (3x6, 3x8, 3x8) = (18, 24, 24) tokens out of the total 30x48x80 video.
+# text_length: int ranging from 0 to 256
+# If your attention contains text token (Hunyuan)
+out = sliding_tile_attention(q, k, v, window_size, text_length)
+# If your attention does not contain text token (StepVideo)
+out = sliding_tile_attention(q, k, v, window_size, 0, False)
+```
+
+
+### Test
+```bash
+python ../tests/test_sta.py # test STA
+python ../tests/test_vsa.py # test VSA
+```
+### Benchmark
+```bash
+python ../benchmarks/bench_sta.py
+```
+
+
+### How Does STA Work?
+We give a demo for 2D STA with window size (6,6) operating on a (10, 10) image. 
+
+
+https://github.com/user-attachments/assets/f3b6dd79-7b43-4b60-a0fa-3d6495ec5747
+
+## Why is STA Fast?
+2D/3D Sliding Window Attention (SWA) creates many mixed blocks in the attention map. Even though mixed blocks have less output value,a mixed block is significantly slower than a dense block due to the GPU-unfriendly masking operation. 
+
+STA removes mixed blocks.
+
+
+<div align="center">
+<img src=../../../assets/sliding_tile_attn_map.png width="80%"/>
+</div>
+
+## Acknowledgement
+
+We learned or reuse code from FlexAtteniton, NATEN, and ThunderKittens.
@@ -0,0 +1,15 @@
+### ADD TO THIS TO REGISTER NEW KERNELS
+sources = {
+    'st_attn': {
+        'source_files': {
+            'h100': 'st_attn/st_attn_h100.cu'  # define these source files for each GPU target desired.
+        }
+    }
+}
+
+### WHICH KERNELS DO WE WANT TO BUILD?
+# (oftentimes during development work you don't need to redefine them all.)
+kernels = ['st_attn']
+
+### WHICH GPU TARGET DO WE WANT TO BUILD FOR?
+target = 'h100'
@@ -0,0 +1,76 @@
+import os
+import subprocess
+
+from config_sta import kernels, sources, target
+from setuptools import find_packages, setup
+from torch.utils.cpp_extension import BuildExtension, CUDAExtension
+
+target = target.lower()
+
+# Package metadata
+PACKAGE_NAME = "st_attn"
+VERSION = "0.0.6"
+AUTHOR = "Hao AI Lab"
+DESCRIPTION = "Sliding Tile Atteniton Kernel Used in FastVideo"
+URL = "https://github.com/hao-ai-lab/FastVideo/tree/main/csrc/sliding_tile_attention"
+
+# Set environment variables
+tk_root = os.getenv('THUNDERKITTENS_ROOT', os.path.abspath(os.path.join(os.getcwd(), 'tk/')))
+python_include = subprocess.check_output(['python', '-c',
+                                          "import sysconfig; print(sysconfig.get_path('include'))"]).decode().strip()
+torch_include = subprocess.check_output([
+    'python', '-c',
+    "import torch; from torch.utils.cpp_extension import include_paths; print(' '.join(['-I' + p for p in include_paths()]))"
+]).decode().strip()
+print('st_attn root:', tk_root)
+print('Python include:', python_include)
+print('Torch include directories:', torch_include)
+
+# CUDA flags
+cuda_flags = [
+    '-DNDEBUG', '-Xcompiler=-Wno-psabi', '-Xcompiler=-fno-strict-aliasing', '--expt-extended-lambda',
+    '--expt-relaxed-constexpr', '-forward-unknown-to-host-compiler', '--use_fast_math', '-std=c++20', '-O3',
+    '-Xnvlink=--verbose', '-Xptxas=--verbose', '-Xptxas=--warn-on-spills', f'-I{tk_root}/include',
+    f'-I{tk_root}/prototype', f'-I{python_include}', '-DTORCH_COMPILE'
+] + torch_include.split()
+cpp_flags = ['-std=c++20', '-O3']
+
+if target == 'h100':
+    cuda_flags.append('-DKITTENS_HOPPER')
+    cuda_flags.append('-arch=sm_90a')
+else:
+    raise ValueError(f'Target {target} not supported')
+
+source_files = ['st_attn.cpp']
+for k in kernels:
+    if target not in sources[k]['source_files']:
+        raise KeyError(f'Target {target} not found in source files for kernel {k}')
+    if isinstance(sources[k]['source_files'][target], list):
+        source_files.extend(sources[k]['source_files'][target])
+    else:
+        source_files.append(sources[k]['source_files'][target])
+    cpp_flags.append(f'-DTK_COMPILE_{k.replace(" ", "_").upper()}')
+
+setup(name=PACKAGE_NAME,
+      version=VERSION,
+      author=AUTHOR,
+      description=DESCRIPTION,
+      url=URL,
+      packages=find_packages(),
+      ext_modules=[
+          CUDAExtension('st_attn_cuda',
+                        sources=source_files,
+                        extra_compile_args={
+                            'cxx': cpp_flags,
+                            'nvcc': cuda_flags
+                        },
+                        libraries=['cuda'])
+      ],
+      cmdclass={'build_ext': BuildExtension},
+      classifiers=[
+          "Programming Language :: Python :: 3",
+          "Environment :: GPU :: NVIDIA CUDA :: 12",
+          "License :: OSI Approved :: Apache Software License",
+      ],
+      python_requires='>=3.10',
+      install_requires=["torch>=2.5.0"])
@@ -0,0 +1,23 @@
+#include <torch/extension.h>
+#include <ATen/ATen.h>
+
+#include <vector>
+#include <cuda_fp16.h>
+#include <cuda_bf16.h>
+
+#include <cuda_runtime.h>
+
+#ifdef TK_COMPILE_ST_ATTN
+extern torch::Tensor sta_forward(
+    torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o, int kernel_t_size, int kernel_w_size, int kernel_h_size, int text_length, bool process_text, bool has_text, int kernel_aspect_ratio_flag
+); 
+#endif
+
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    m.doc() = "Sliding Block Attention Kernels"; // optional module docstring
+    
+#ifdef TK_COMPILE_ST_ATTN
+    m.def("sta_fwd",  torch::wrap_pybind_function(sta_forward), "sliding tile attention, assuming tile size is (6,8,8)");
+#endif
+}
@@ -0,0 +1,49 @@
+import math
+
+import torch
+from torch.utils.checkpoint import detach_variable
+try:
+    from st_attn_cuda import sta_fwd
+except ImportError:
+    sta_fwd = None
+
+def sliding_tile_attention(q_all, k_all, v_all, window_size, text_length, has_text=True, dit_seq_shape='30x48x80'):
+    seq_length = q_all.shape[2]
+    dit_seq_shape_mapping = {
+        '30x48x80':1,
+        '36x48x48':2,
+        '18x48x80':3, 
+    }
+    if has_text:
+        assert q_all.shape[
+            2] >= 115200 and q_all.shape[2] <= 115456, f"Unsupported {dit_seq_shape}, current shape is {q_all.shape}, only support '30x48x80' for HunyuanVideo"
+        assert q_all.shape[1] == len(window_size), "Number of heads must match the number of window sizes"
+        target_size = math.ceil(seq_length / 384) * 384
+        pad_size = target_size - seq_length
+        if pad_size > 0:
+            q_all = torch.cat([q_all, q_all[:, :, -pad_size:]], dim=2)
+            k_all = torch.cat([k_all, k_all[:, :, -pad_size:]], dim=2)
+            v_all = torch.cat([v_all, v_all[:, :, -pad_size:]], dim=2)
+    else:
+        if dit_seq_shape == '36x48x48': # Stepvideo 204x768x68
+            assert q_all.shape[2] == 82944
+        elif dit_seq_shape == '18x48x80': # Wan 69x768x1280
+            assert q_all.shape[2] == 69120
+        else:
+            raise ValueError(f"Unsupported {dit_seq_shape}, current shape is {q_all.shape}, only support '36x48x48' for Stepvideo and '18x48x80' for Wan")
+
+    kernel_aspect_ratio_flag = dit_seq_shape_mapping[dit_seq_shape]
+    hidden_states = torch.empty_like(q_all)
+    # This for loop is ugly. but it is actually quite efficient. The sequence dimension alone can already oversubscribe SMs
+    for head_index, (t_kernel, h_kernel, w_kernel) in enumerate(window_size):
+        for batch in range(q_all.shape[0]):
+            q_head, k_head, v_head, o_head = (q_all[batch:batch + 1, head_index:head_index + 1],
+                                              k_all[batch:batch + 1,
+                                                    head_index:head_index + 1], v_all[batch:batch + 1,
+                                                                                      head_index:head_index + 1],
+                                              hidden_states[batch:batch + 1, head_index:head_index + 1])
+
+            _ = sta_fwd(q_head, k_head, v_head, o_head, t_kernel, h_kernel, w_kernel, text_length, False, has_text, kernel_aspect_ratio_flag)
+    if has_text:
+        _ = sta_fwd(q_all, k_all, v_all, hidden_states, 3, 3, 3, text_length, True, True, kernel_aspect_ratio_flag)
+    return hidden_states[:, :, :seq_length]
@@ -451,10 +451,6 @@ sta_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o,

        globals g{qg_arg, kg_arg, vg_arg, lg_arg, og_arg, static_cast<int>(seq_len),  static_cast<int>(text_length), static_cast<int>(hr)};

-        // Shared memory size for the kernel.
-        // We use the maximum available shared memory (kittens::MAX_SHARED_MEMORY) 
-        // which is approximately 227KB on H100, necessary for the high-performance 
-        // TMA-based attention tiles with multiple stages.
        constexpr int mem_size = kittens::MAX_SHARED_MEMORY;
        int threads = NUM_WORKERS * kittens::WARP_THREADS;
        if (has_text) {
@@ -462,31 +458,104 @@ sta_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o,
            dim3 grid_image(seq_len/(CONSUMER_WARPGROUPS*kittens::TILE_ROW_DIM<bf16>*4)-2, qo_heads, batch);
            dim3 grid_text(2, qo_heads, batch);
            if (!process_text) {
-#define LAUNCH_IMAGE_KER(DT_VAL, DH_VAL, DW_VAL) \
-                    cudaFuncSetAttribute( \
-                        fwd_attend_ker<128, false, false, true, DT_VAL, DH_VAL, DW_VAL, 5, 6, 10>, \
-                        cudaFuncAttributeMaxDynamicSharedMemorySize, \
-                        mem_size \
-                    ); \
-                    fwd_attend_ker<128, false, false, true, DT_VAL, DH_VAL, DW_VAL, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 3) {

-                if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 3)      { LAUNCH_IMAGE_KER(1, 1, 1); }
-                else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 5) { LAUNCH_IMAGE_KER(1, 1, 2); }
-                else if (kernel_t_size == 5 && kernel_h_size == 3 && kernel_w_size == 3) { LAUNCH_IMAGE_KER(2, 1, 1); }
-                else if (kernel_t_size == 3 && kernel_h_size == 5 && kernel_w_size == 5) { LAUNCH_IMAGE_KER(1, 2, 2); }
-                else if (kernel_t_size == 5 && kernel_h_size == 6 && kernel_w_size == 1) { LAUNCH_IMAGE_KER(2, 3, 0); }
-                else if (kernel_t_size == 5 && kernel_h_size == 3 && kernel_w_size == 5) { LAUNCH_IMAGE_KER(2, 1, 2); }
-                else if (kernel_t_size == 5 && kernel_h_size == 5 && kernel_w_size == 5) { LAUNCH_IMAGE_KER(2, 2, 2); }
-                else if (kernel_t_size == 5 && kernel_h_size == 5 && kernel_w_size == 7) { LAUNCH_IMAGE_KER(2, 2, 3); }
-                else if (kernel_t_size == 5 && kernel_h_size == 6 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(2, 3, 5); }
-                else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(1, 3, 5); }
-                else if (kernel_t_size == 5 && kernel_h_size == 1 && kernel_w_size == 1) { LAUNCH_IMAGE_KER(2, 0, 0); }
-                else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(0, 3, 5); }
-                else if (kernel_t_size == 5 && kernel_h_size == 1 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(2, 0, 5); }
-                else {
-                    TORCH_CHECK(false, "Invalid kernel size: ", kernel_t_size, "x", kernel_h_size, "x", kernel_w_size);
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 1, 1, 1, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 1, 1, 1, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                }  else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 5) {
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 1, 1, 2, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true,1, 1, 2, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 5 && kernel_h_size == 3 && kernel_w_size == 3) {
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 1, 1, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 2, 1, 1, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                }else if (kernel_t_size ==3 && kernel_h_size == 5 && kernel_w_size == 5){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 1, 2, 2, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 1, 2, 2, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size ==5 && kernel_h_size == 6 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 3, 0, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 2, 3, 0, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size ==5 && kernel_h_size == 3 && kernel_w_size == 5){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 1, 2, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 2, 1, 2, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 5 && kernel_h_size == 5 && kernel_w_size == 5){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 2, 2, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 2, 2, 2, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 5 && kernel_h_size == 5 && kernel_w_size == 7){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 2, 3, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 2, 2, 3, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 5 && kernel_h_size == 6 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 3, 5, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 2, 3, 5, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 5 && kernel_h_size == 1 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 0, 0, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 2, 0, 0, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 0, 3, 5, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true, 0, 3, 5, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 5 && kernel_h_size == 1 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, true, 2, 0, 5, 5, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, true,2, 0, 5, 5, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else {
+                    // print error
+                    std::cout << "Invalid kernel size" << std::endl;
+                    //print kernel size
+                    std::cout << "Kernel size: " << kernel_t_size << " " << kernel_h_size << " " << kernel_w_size << std::endl;
                }
-#undef LAUNCH_IMAGE_KER
            } else {
                cudaFuncSetAttribute(
                    fwd_attend_ker<128, false, true, true, 1, 1, 1, 5, 6, 10>,
@@ -499,67 +568,266 @@ sta_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o,
        } else {
            dim3 grid_image(seq_len/(CONSUMER_WARPGROUPS*kittens::TILE_ROW_DIM<bf16>*4), qo_heads, batch);
            if (kernel_aspect_ratio_flag == 2){
-#define LAUNCH_IMAGE_KER(DT_VAL, DH_VAL, DW_VAL) \
-                    cudaFuncSetAttribute( \
-                        fwd_attend_ker<128, false, false, false, DT_VAL, DH_VAL, DW_VAL, 6, 6, 6>, \
-                        cudaFuncAttributeMaxDynamicSharedMemorySize, \
-                        mem_size \
-                    ); \
-                    fwd_attend_ker<128, false, false, false, DT_VAL, DH_VAL, DW_VAL, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 3) {
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 1, 1, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 1, 1, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);

-                if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 3)      { LAUNCH_IMAGE_KER(1, 1, 1); }
-                else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 6) { LAUNCH_IMAGE_KER(1, 1, 3); }
-                else if (kernel_t_size == 6 && kernel_h_size == 3 && kernel_w_size == 3) { LAUNCH_IMAGE_KER(3, 1, 1); }
-                else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 6) { LAUNCH_IMAGE_KER(1, 3, 3); }
-                else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 3) { LAUNCH_IMAGE_KER(1, 3, 1); }
-                else if (kernel_t_size == 6 && kernel_h_size == 3 && kernel_w_size == 6) { LAUNCH_IMAGE_KER(3, 1, 3); }
-                else if (kernel_t_size == 6 && kernel_h_size == 6 && kernel_w_size == 6) { LAUNCH_IMAGE_KER(3, 3, 3); }
-                else if (kernel_t_size == 6 && kernel_h_size == 1 && kernel_w_size == 1) { LAUNCH_IMAGE_KER(3, 0, 0); }
-                else if (kernel_t_size == 6 && kernel_h_size == 1 && kernel_w_size == 6) { LAUNCH_IMAGE_KER(3, 0, 3); }
-                else if (kernel_t_size == 6 && kernel_h_size == 6 && kernel_w_size == 1) { LAUNCH_IMAGE_KER(3, 3, 0); }
-                else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 6) { LAUNCH_IMAGE_KER(0, 3, 3); }
-                else if (kernel_t_size == 1 && kernel_h_size == 1 && kernel_w_size == 6) { LAUNCH_IMAGE_KER(0, 0, 3); }
-                else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 1) { LAUNCH_IMAGE_KER(0, 3, 0); }
-                else {
-                    TORCH_CHECK(false, "Invalid kernel size: ", kernel_t_size, "x", kernel_h_size, "x", kernel_w_size);
+                }  else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 6) {
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 1, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false,1, 1, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 6 && kernel_h_size == 3 && kernel_w_size == 3) {
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 1, 1, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 1, 1, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size ==3 && kernel_h_size == 6 && kernel_w_size == 6){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 3, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 3, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                }else if (kernel_t_size ==3 && kernel_h_size == 6 && kernel_w_size == 3){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 3, 1, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 3, 1, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size ==6 && kernel_h_size == 3 && kernel_w_size == 6){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 1, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 1, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 6 && kernel_h_size == 6 && kernel_w_size == 6){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 3, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 3, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 6 && kernel_h_size == 1 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 0, 0, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 0, 0, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 6 && kernel_h_size == 1 && kernel_w_size == 6){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 0, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 0, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                }  else if (kernel_t_size == 6 && kernel_h_size == 6 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 3, 0, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 3, 0, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 6){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 3, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 3, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                }  else if (kernel_t_size == 1 && kernel_h_size == 1 && kernel_w_size == 6){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 0, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 0, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 3, 0, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 3, 0, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 6 && kernel_h_size == 6 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 3, 0, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 3, 0, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 6 && kernel_h_size == 1 && kernel_w_size == 6){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 3, 0, 3, 6, 6, 6>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 3, 0, 3, 6, 6, 6><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else {
+                    // print error
+                    std::cout << "Invalid kernel size" << std::endl;
+                    //print kernel size
+                    std::cout << "Kernel size: " << kernel_t_size << " " << kernel_h_size << " " << kernel_w_size << std::endl;
                }
-#undef LAUNCH_IMAGE_KER
            }
            else if (kernel_aspect_ratio_flag == 3) {
-#define LAUNCH_IMAGE_KER(DT_VAL, DH_VAL, DW_VAL) \
-                    cudaFuncSetAttribute( \
-                        fwd_attend_ker<128, false, false, false, DT_VAL, DH_VAL, DW_VAL, 3, 6, 10>, \
-                        cudaFuncAttributeMaxDynamicSharedMemorySize, \
-                        mem_size \
-                    ); \
-                    fwd_attend_ker<128, false, false, false, DT_VAL, DH_VAL, DW_VAL, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 3) {

-                if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 3)      { LAUNCH_IMAGE_KER(1, 1, 1); }
-                else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 5) { LAUNCH_IMAGE_KER(1, 1, 2); }
-                else if (kernel_t_size == 3 && kernel_h_size == 5 && kernel_w_size == 5) { LAUNCH_IMAGE_KER(1, 2, 2); }
-                else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 1) { LAUNCH_IMAGE_KER(1, 3, 0); }
-                else if (kernel_t_size == 3 && kernel_h_size == 5 && kernel_w_size == 7) { LAUNCH_IMAGE_KER(1, 2, 3); }
-                else if (kernel_t_size == 3 && kernel_h_size == 5 && kernel_w_size == 9) { LAUNCH_IMAGE_KER(1, 2, 4); }
-                else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(1, 3, 5); }
-                else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 3) { LAUNCH_IMAGE_KER(1, 3, 1); }
-                else if (kernel_t_size == 3 && kernel_h_size == 1 && kernel_w_size == 1) { LAUNCH_IMAGE_KER(1, 0, 0); }
-                else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(0, 3, 5); }
-                else if (kernel_t_size == 1 && kernel_h_size == 5 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(0, 2, 5); }
-                else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 7) { LAUNCH_IMAGE_KER(0, 3, 3); }
-                else if (kernel_t_size == 1 && kernel_h_size == 5 && kernel_w_size == 7) { LAUNCH_IMAGE_KER(0, 2, 3); }
-                else if (kernel_t_size == 1 && kernel_h_size == 5 && kernel_w_size == 9) { LAUNCH_IMAGE_KER(0, 2, 4); }
-                else if (kernel_t_size == 3 && kernel_h_size == 1 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(1, 0, 5); }
-                else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(1, 1, 5); }
-                else if (kernel_t_size == 1 && kernel_h_size == 3 && kernel_w_size == 10){ LAUNCH_IMAGE_KER(0, 1, 5); }
-                else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 5) { LAUNCH_IMAGE_KER(0, 3, 2); }
-                else {
-                    TORCH_CHECK(false, "Invalid kernel size: ", kernel_t_size, "x", kernel_h_size, "x", kernel_w_size);
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 1, 1, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 1, 1, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 5) {
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 1, 2, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false,1, 1, 2, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 3 && kernel_h_size == 5 && kernel_w_size == 5){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 2, 2, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 2, 2, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 3, 0, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 3, 0, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+
+                } else if (kernel_t_size == 3 && kernel_h_size == 5 && kernel_w_size == 7){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 2, 3, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 2, 3, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 3 && kernel_h_size == 5 && kernel_w_size == 9){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 2, 4, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 2, 4, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 3, 5, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 3, 5, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 3 && kernel_h_size == 6 && kernel_w_size == 3){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 3, 1, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 3, 1, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 3 && kernel_h_size == 1 && kernel_w_size == 1){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 0, 0, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 1, 0, 0, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 3, 5, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 3, 5, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 5 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 2, 5, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 2, 5, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 7){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 3, 3, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 3, 3, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 5 && kernel_w_size == 7){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 2, 3, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 2, 3, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 5 && kernel_w_size == 9){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 2, 4, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false, 0, 2, 4, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 3 && kernel_h_size == 1 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 0, 5, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false,1, 0, 5, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 3 && kernel_h_size == 3 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 1, 1, 5, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false,1, 1, 5, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 3 && kernel_w_size == 10){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 1, 5, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false,0, 1, 5, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else if (kernel_t_size == 1 && kernel_h_size == 6 && kernel_w_size == 5){
+                    cudaFuncSetAttribute(
+                        fwd_attend_ker<128, false, false, false, 0, 3, 2, 3, 6, 10>,
+                        cudaFuncAttributeMaxDynamicSharedMemorySize,
+                        mem_size
+                    );
+                    fwd_attend_ker<128, false, false, false,0, 3, 2, 3, 6, 10><<<grid_image, (32*NUM_WORKERS), mem_size, stream>>>(g);
+                } else {
+                    // print error
+                    std::cout << "Invalid kernel size" << std::endl;
+                    //print kernel size
+                    std::cout << "Kernel size: " << kernel_t_size << " " << kernel_h_size << " " << kernel_w_size << std::endl;
                }
-#undef LAUNCH_IMAGE_KER
            }

            else {
-                TORCH_CHECK(false, "Unsupported kernel_aspect_ratio_flag: ", kernel_aspect_ratio_flag);
+                std::cout << "Unsupported kernel_aspect_ratio_flag: " << kernel_aspect_ratio_flag << std::endl;
            }

        }
@@ -1,6 +1,6 @@
 import torch
-from .support_flex_sta import get_sliding_tile_attention_mask
-from fastvideo_kernel import sliding_tile_attention
+from flex_sta_ref import get_sliding_tile_attention_mask
+from st_attn 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
@@ -73,23 +73,15 @@ def check_correctness(b, h, n, d, causal, mean, std, num_iterations=50, error_mo


 # 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()
+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']}")
@@ -0,0 +1,156 @@
+import torch
+import sys
+import os
+import numpy as np
+from tqdm import tqdm
+
+# Add the parent directory to the path to import block_sparse_attn
+sys.path.append(os.path.dirname(os.path.dirname(os.path.abspath(__file__))))
+from tests.utils import generate_block_sparse_mask_for_function, create_full_mask_from_block_mask
+from vsa import 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, non_pad_index, dO):
+    Q = Q.detach().requires_grad_()
+    K = K.detach().requires_grad_()
+    V = V.detach().requires_grad_()
+    
+    q_padded = vsa_pad(Q, non_pad_index, variable_block_sizes.shape[0], BLOCK_M)
+    k_padded = vsa_pad(K, non_pad_index, variable_block_sizes.shape[0], BLOCK_M)
+    v_padded = vsa_pad(V, non_pad_index, variable_block_sizes.shape[0], BLOCK_M)
+    output, _= block_sparse_attn(q_padded, k_padded, v_padded, block_sparse_mask, variable_block_sizes)
+    output = output[:, :, 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=32, 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, k, device)
+    full_mask = create_full_mask_from_block_mask(block_mask, 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)
+
+        # 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, 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 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)
+
+if __name__ == "__main__":
+    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)
+    print("\nAnalysis completed for all modes.")
@@ -0,0 +1,54 @@
+import torch
+
+def generate_block_sparse_mask_for_function(h, num_blocks, k, device="cuda"):
+    """
+    Generate block sparse mask of shape [h, num_blocks, num_blocks].
+    
+    Args:
+        h: number of heads
+        num_blocks: number of blocks
+        k: number of kv blocks each q block attends to
+        device: device to create tensors on
+
+    Returns:
+        block_sparse_mask: [h, num_blocks, num_blocks] bool tensor
+    """
+    k = min(k, num_blocks)
+    scores = torch.rand(h, num_blocks, num_blocks, device=device)
+    _, indices = torch.topk(scores, k, dim=-1)
+    block_sparse_mask = torch.zeros(h, num_blocks, num_blocks, dtype=torch.bool, device=device)
+
+    block_sparse_mask = block_sparse_mask.scatter_(2, indices, 1).bool()
+    return block_sparse_mask
+
+
+def create_full_mask_from_block_mask(block_sparse_mask, variable_block_sizes, device="cuda"):
+    """
+    Convert block-level sparse mask to full attention mask.
+    
+    Args:
+        block_sparse_mask: [h, num_blocks, num_blocks] bool tensor
+        variable_block_sizes: [num_blocks] tensor
+        device: device to create tensors on
+
+    Returns:
+        full_mask: [h, S, S] bool tensor where S = total sequence length
+    """
+    h, num_blocks, _ = block_sparse_mask.shape
+    total_seq_len = variable_block_sizes.sum().item()
+    cumsum = torch.cat([torch.tensor([0], device=device), variable_block_sizes.cumsum(dim=0)[:-1]])
+
+    full_mask = torch.zeros(h, total_seq_len, total_seq_len, dtype=torch.bool, device=device)
+
+    for head in range(h):
+        for q_block in range(num_blocks):
+            q_start = cumsum[q_block]
+            q_end = q_start + variable_block_sizes[q_block]
+
+            for kv_block in range(num_blocks):
+                if block_sparse_mask[head, q_block, kv_block]:
+                    kv_start = cumsum[kv_block]
+                    kv_end = kv_start + variable_block_sizes[kv_block]
+                    full_mask[head, q_start:q_end, kv_start:kv_end] = True
+
+    return full_mask
@@ -0,0 +1,2 @@
+recursive-include tk *
+include config_vsa.py
@@ -0,0 +1,61 @@
+
+
+# Attention Kernel Used in FastVideo
+
+## Video Sparse Attention (VSA)
+
+### Installation
+We support H100 (via TK) and any other GPU (via triton) for VSA.
+
+```bash
+pip install vsa
+``` 
+
+Install from source:
+
+```bash
+git submodule update --init --recursive
+python setup.py install
+```
+
+
+If you encounter error during installation, try below:
+Install C++20 for ThunderKittens:
+```bash
+sudo apt update
+sudo apt install gcc-11 g++-11
+
+sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
+
+sudo apt update
+sudo apt install clang-11
+```
+(If you use CUDA12.8)
+```bash
+export CUDA_HOME=/usr/local/cuda-12.8
+export PATH=${CUDA_HOME}/bin:${PATH} 
+export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
+```
+
+### Verify if you have successfully installed
+
+```bash
+# test numerical
+python ../tests/test_vsa.py
+# (For H100) test speed
+python ../benchmarks/bench_vsa_hopper.py
+```
+
+bench_vsa_hopper.py should print something like this:
+
+```bash
+Using topk=76 kv blocks per q block (out of 768 total kv blocks)
+
+=== BLOCK SPARSE ATTENTION BENCHMARK ===
+Block Sparse Forward  - TFLOPS: 5622.26
+Block Sparse Backward - TFLOPS: 3865.68
+```
+
+## Acknowledgement
+
+We learned or reuse code from FlexAtteniton, NATEN, and ThunderKittens.
@@ -0,0 +1,15 @@
+### ADD TO THIS TO REGISTER NEW KERNELS
+sources = {
+    'block_sparse': {
+        'source_files': {
+            'h100': 'vsa/block_sparse_h100.cu'
+        }
+    }
+}
+
+### WHICH KERNELS DO WE WANT TO BUILD?
+# (oftentimes during development work you don't need to redefine them all.)
+kernels = ['block_sparse']
+
+### WHICH GPU TARGET DO WE WANT TO BUILD FOR?
+target = 'h100'
@@ -0,0 +1,81 @@
+import os
+import subprocess
+
+from config_vsa import kernels, sources, target
+from setuptools import find_packages, setup
+from torch.utils.cpp_extension import BuildExtension, CUDAExtension
+
+target = target.lower()
+
+# Package metadata
+PACKAGE_NAME = "vsa"
+VERSION = "0.0.3"
+AUTHOR = "Hao AI Lab"
+DESCRIPTION = "Video Sparse Attention Kernel Used in FastVideo"
+URL = "https://github.com/hao-ai-lab/FastVideo/tree/main/csrc/attn/video_sparse_attn"
+
+# Set environment variables
+tk_root = os.getenv('THUNDERKITTENS_ROOT', os.path.abspath(os.path.join(os.getcwd(), 'tk/')))
+python_include = subprocess.check_output(['python', '-c',
+                                          "import sysconfig; print(sysconfig.get_path('include'))"]).decode().strip()
+torch_include = subprocess.check_output([
+    'python', '-c',
+    "import torch; from torch.utils.cpp_extension import include_paths; print(' '.join(['-I' + p for p in include_paths()]))"
+]).decode().strip()
+print('vsa root:', tk_root)
+print('Python include:', python_include)
+print('Torch include directories:', torch_include)
+
+# CUDA flags
+cuda_flags = [
+    '-DNDEBUG', '-Xcompiler=-Wno-psabi', '-Xcompiler=-fno-strict-aliasing', '--expt-extended-lambda',
+    '--expt-relaxed-constexpr', '-forward-unknown-to-host-compiler', '--use_fast_math', '-std=c++20', '-O3',
+    '-Xnvlink=--verbose', '-Xptxas=--verbose', '-Xptxas=--warn-on-spills', f'-I{tk_root}/include',
+    f'-I{tk_root}/prototype', f'-I{python_include}', '-DTORCH_COMPILE'
+] + torch_include.split()
+cpp_flags = ['-std=c++20', '-O3']
+
+if target == 'h100':
+    cuda_flags.append('-DKITTENS_HOPPER')
+    cuda_flags.append('-arch=sm_90a')
+else:
+    raise ValueError(f'Target {target} not supported')
+
+source_files = ['vsa.cpp']
+for k in kernels:
+    if target not in sources[k]['source_files']:
+        raise KeyError(f'Target {target} not found in source files for kernel {k}')
+    if isinstance(sources[k]['source_files'][target], list):
+        source_files.extend(sources[k]['source_files'][target])
+    else:
+        source_files.append(sources[k]['source_files'][target])
+    cpp_flags.append(f'-DTK_COMPILE_{k.replace(" ", "_").upper()}')
+
+
+ext_modules = [
+    CUDAExtension('vsa_cuda',
+        sources=source_files,
+        extra_compile_args={
+            'cxx': cpp_flags,
+            'nvcc': cuda_flags
+        },
+        libraries=['cuda'])
+]
+
+
+
+setup(name=PACKAGE_NAME,
+      version=VERSION,
+      author=AUTHOR,
+      description=DESCRIPTION,
+      url=URL,
+      packages=find_packages(),
+      ext_modules=ext_modules,
+      cmdclass={'build_ext': BuildExtension},
+      classifiers=[
+          "Programming Language :: Python :: 3",
+          "Environment :: GPU :: NVIDIA CUDA :: 12",
+          "License :: OSI Approved :: Apache Software License",
+      ],
+      python_requires='>=3.10',
+      install_requires=["torch>=2.5.0"])
@@ -0,0 +1,27 @@
+#include <torch/extension.h>
+#include <ATen/ATen.h>
+
+#include <vector>
+#include <cuda_fp16.h>
+#include <cuda_bf16.h>
+
+#include <cuda_runtime.h>
+
+#ifdef TK_COMPILE_BLOCK_SPARSE
+extern std::vector<torch::Tensor> block_sparse_attention_forward(
+    torch::Tensor q, torch::Tensor k, torch::Tensor v,  torch::Tensor q2k_block_sparse_index, torch::Tensor q2k_block_sparse_num, torch::Tensor block_size
+); 
+extern std::vector<torch::Tensor> block_sparse_attention_backward(
+    torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o, torch::Tensor l_vec, torch::Tensor og, torch::Tensor k2q_block_sparse_index, torch::Tensor k2q_block_sparse_num, torch::Tensor block_size
+);
+#endif
+
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    m.doc() = "Video Sparse Attention Kernels"; // optional module docstring
+
+#ifdef TK_COMPILE_BLOCK_SPARSE
+    m.def("block_sparse_fwd",  torch::wrap_pybind_function(block_sparse_attention_forward), "block sparse attention");
+    m.def("block_sparse_bwd",  torch::wrap_pybind_function(block_sparse_attention_backward), "block sparse attention backward");
+#endif
+}
@@ -0,0 +1,80 @@
+import torch
+from typing import Tuple
+block_sparse_attn=None
+import torch
+major, minor = torch.cuda.get_device_capability(0)
+if major == 9 and minor == 0:# check if H100
+    from vsa_cuda import block_sparse_fwd, block_sparse_bwd
+    from vsa.block_sparse_wrapper import block_sparse_attn_SM90
+    block_sparse_attn = block_sparse_attn_SM90
+else:
+    from vsa.block_sparse_wrapper import block_sparse_attn_triton
+    block_sparse_fwd = None
+    block_sparse_bwd = None
+    block_sparse_attn = block_sparse_attn_triton
+
+BLOCK_M = 64
+BLOCK_N = 64
+
+
+def torch_attention(q, k, v) -> Tuple[torch.Tensor, torch.Tensor]:
+    QK = torch.matmul(q, k.transpose(-2, -1))
+    QK /= (q.size(-1)**0.5)
+
+    # Causal mask removed since causal is always false
+
+    QK = torch.nn.functional.softmax(QK, dim=-1)
+    output = torch.matmul(QK, v)
+    return output, QK
+
+
+def video_sparse_attn(q, k, v, variable_block_sizes, topk, block_size, compress_attn_weight=None):
+    """
+    q: [batch_size, num_heads, seq_len, head_dim]
+    k: [batch_size, num_heads, seq_len, head_dim]
+    v: [batch_size, num_heads, seq_len, head_dim]
+    topk: int
+    block_size: int or tuple of 3 ints
+    video_shape: tuple of (T, H, W)
+    compress_attn_weight: [batch_size, num_heads, seq_len, head_dim]
+    select_attn_weight: [batch_size, num_heads, seq_len, head_dim]
+    NOTE: We assume q, k, v is zero padded!!
+    V1 of sparse attention. Include compress attn and sparse attn branch, use average pooling to compress. 
+    Assume q, k, v is flattened in this way: [batch_size, num_heads, T//block_size[0], H//block_size[1], W//block_size[2], block_size[0], block_size[1], block_size[2]]
+    """
+
+    if isinstance(block_size, int):
+        block_size = (block_size, block_size, block_size)
+
+    block_elements = block_size[0] * block_size[1] * block_size[2]
+    assert block_elements == 64
+    assert q.shape[2] % block_elements == 0
+    batch_size, num_heads, seq_len, head_dim = q.shape
+    # compress attn
+    q_compress = (q.view(batch_size, num_heads, seq_len // block_elements,
+                        block_elements, head_dim).float().sum(dim=3) / variable_block_sizes.view(1, 1, -1, 1)).to(q.dtype)
+    k_compress = (k.view(batch_size, num_heads, seq_len // block_elements,
+                        block_elements, head_dim).float().sum(dim=3) / variable_block_sizes.view(1, 1, -1, 1)).to(k.dtype)
+    v_compress = (v.view(batch_size, num_heads, seq_len // block_elements,
+                        block_elements, head_dim).float().sum(dim=3) / variable_block_sizes.view(1, 1, -1, 1)).to(v.dtype)
+
+    output_compress, block_attn_score = torch_attention(q_compress, k_compress,
+                                                        v_compress)
+
+    output_compress = output_compress.view(batch_size, num_heads,
+                                           seq_len // block_elements, 1,
+                                           head_dim)
+    output_compress = output_compress.repeat(1, 1, 1, block_elements,
+                                             1).view(batch_size, num_heads,
+                                                     seq_len, head_dim)
+
+    topK_indices = torch.topk(block_attn_score, topk, dim=-1).indices
+    block_mask = torch.zeros_like(block_attn_score, dtype=torch.bool).scatter_(-1, topK_indices, True)
+    output_select, _ = block_sparse_attn(q, k, v, block_mask, variable_block_sizes)
+
+    if compress_attn_weight is not None:
+        final_output = output_compress * compress_attn_weight + output_select
+    else:
+        final_output = output_compress + output_select
+    return final_output
+
@@ -8,6 +8,7 @@ This is a Triton implementation of the Flash Attention v2 algorithm from Tri Dao
 Credits: OpenAI kernel team
 """

+import pytest
 import torch
 import triton
 import triton.language as tl
@@ -16,6 +17,7 @@ import triton.language as tl
 import math  # small utility needed by the sparse wrapper
 # ──────────────────────────── SPARSE ADDITION END ─────────────────────────────

+
 # We don't run auto-tuning every time to keep the tutorial fast. Keeping
 # the code below and commenting out the equivalent parameters is convenient for
 # re-tuning.
@@ -27,92 +29,65 @@ configs = [
    for w in [4, 8]\
 ]

-
 # ──────────────────────────── SPARSE ADDITION BEGIN ───────────────────────────
@triton.autotune(configs, key=["N_CTX", "HEAD_DIM"])
@triton.jit
-def _attn_fwd_sparse(
-        Q,
-        K,
-        V,
-        sm_scale,  #
-        q2k_index,
-        q2k_num,
-        max_kv_blks,  #
-        variable_block_sizes,
-        M,
-        Out,  #
-        stride_qz,
-        stride_qh,
-        stride_qm,
-        stride_qk,
-        stride_kz,
-        stride_kh,
-        stride_kn,
-        stride_kk,
-        stride_vz,
-        stride_vh,
-        stride_vk,
-        stride_vn,
-        stride_oz,
-        stride_oh,
-        stride_om,
-        stride_on,
-        Z,
-        H,
-        N_CTX,  #
-        HEAD_DIM: tl.constexpr,  #
-        BLOCK_M: tl.constexpr,
-        BLOCK_N: tl.constexpr,
-        STAGE: tl.constexpr):
+def _attn_fwd_sparse(Q, K, V, sm_scale,                         #
+                     q2k_index, q2k_num, max_kv_blks,           #
+                     variable_block_sizes,
+                     M, Out,                                    #
+                     stride_qz, stride_qh, stride_qm, stride_qk,
+                     stride_kz, stride_kh, stride_kn, stride_kk,
+                     stride_vz, stride_vh, stride_vk, stride_vn,
+                     stride_oz, stride_oh, stride_om, stride_on,
+                     Z, H, N_CTX,                               #
+                     HEAD_DIM: tl.constexpr,                    #
+                     BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr,
+                     STAGE: tl.constexpr):
    """
    64×64 **block-sparse** forward kernel. Back-prop kernels remain dense
    (32×64 and 64×32) – memory footprint unchanged.
    """

    # ----- program-id mapping -----
-    q_blk = tl.program_id(0)  # Q-tile index
-    off_hz = tl.program_id(1)  # fused (batch, head)
-    b = off_hz // H
-    h = off_hz % H
+    q_blk   = tl.program_id(0)          # Q-tile index
+    off_hz  = tl.program_id(1)          # fused (batch, head)
+    b       = off_hz // H
+    h       = off_hz %  H
    q_tiles = N_CTX // BLOCK_M
    meta_base = ((b * H + h) * q_tiles + q_blk)

-    kv_blocks = tl.load(q2k_num + meta_base)  # int32
-    kv_ptr = q2k_index + meta_base * max_kv_blks  # ptr to list
+    kv_blocks = tl.load(q2k_num  + meta_base)                 # int32
+    kv_ptr    = q2k_index + meta_base * max_kv_blks           # ptr to list

    # ----- base pointers -----
-    qvk_off = (b.to(tl.int64) * stride_qz + h.to(tl.int64) * stride_qh)
+    qvk_off = (b.to(tl.int64) * stride_qz +
+               h.to(tl.int64) * stride_qh)

-    Q_ptr = tl.make_block_ptr(base=Q + qvk_off,
-                              shape=(N_CTX, HEAD_DIM),
-                              strides=(stride_qm, stride_qk),
-                              offsets=(q_blk * BLOCK_M, 0),
-                              block_shape=(BLOCK_M, HEAD_DIM),
-                              order=(1, 0))
+    Q_ptr = tl.make_block_ptr(
+        base=Q + qvk_off, shape=(N_CTX, HEAD_DIM),
+        strides=(stride_qm, stride_qk),
+        offsets=(q_blk * BLOCK_M, 0),
+        block_shape=(BLOCK_M, HEAD_DIM), order=(1, 0))

-    K_base = tl.make_block_ptr(base=K + qvk_off,
-                               shape=(HEAD_DIM, N_CTX),
-                               strides=(stride_kk, stride_kn),
-                               offsets=(0, 0),
-                               block_shape=(HEAD_DIM, BLOCK_N),
-                               order=(0, 1))
+    K_base = tl.make_block_ptr(
+        base=K + qvk_off, shape=(HEAD_DIM, N_CTX),
+        strides=(stride_kk, stride_kn),
+        offsets=(0, 0),
+        block_shape=(HEAD_DIM, BLOCK_N), order=(0, 1))

-    v_order: tl.constexpr = (0, 1) if V.dtype.element_ty == tl.float8e5 else (1,
-                                                                              0)
-    V_base = tl.make_block_ptr(base=V + qvk_off,
-                               shape=(N_CTX, HEAD_DIM),
-                               strides=(stride_vk, stride_vn),
-                               offsets=(0, 0),
-                               block_shape=(BLOCK_N, HEAD_DIM),
-                               order=v_order)
+    v_order: tl.constexpr = (0, 1) if V.dtype.element_ty == tl.float8e5 else (1, 0)
+    V_base = tl.make_block_ptr(
+        base=V + qvk_off, shape=(N_CTX, HEAD_DIM),
+        strides=(stride_vk, stride_vn),
+        offsets=(0, 0),
+        block_shape=(BLOCK_N, HEAD_DIM), order=v_order)

-    O_ptr = tl.make_block_ptr(base=Out + qvk_off,
-                              shape=(N_CTX, HEAD_DIM),
-                              strides=(stride_om, stride_on),
-                              offsets=(q_blk * BLOCK_M, 0),
-                              block_shape=(BLOCK_M, HEAD_DIM),
-                              order=(1, 0))
+    O_ptr = tl.make_block_ptr(
+        base=Out + qvk_off, shape=(N_CTX, HEAD_DIM),
+        strides=(stride_om, stride_on),
+        offsets=(q_blk * BLOCK_M, 0),
+        block_shape=(BLOCK_M, HEAD_DIM), order=(1, 0))

    # ----- accumulators -----
    offs_m = q_blk * BLOCK_M + tl.arange(0, BLOCK_M)
@@ -152,30 +127,23 @@ def _attn_fwd_sparse(
    acc = acc / l_i[:, None]
    tl.store(M + off_hz * N_CTX + offs_m, m_i)
    tl.store(O_ptr, acc.to(Out.type.element_ty))
-
-
 # ──────────────────────────── SPARSE ADDITION END ─────────────────────────────


+
+
@triton.jit
-def _attn_bwd_preprocess(
-        O,
-        DO,  #
-        Delta,  #
-        Z,
-        H,
-        N_CTX,  #
-        BLOCK_M: tl.constexpr,
-        HEAD_DIM: tl.constexpr  #
-):
+def _attn_bwd_preprocess(O, DO,  #
+                         Delta,  #
+                         Z, H, N_CTX,  #
+                         BLOCK_M: tl.constexpr, HEAD_DIM: tl.constexpr  #
+                         ):
    off_m = tl.program_id(0) * BLOCK_M + tl.arange(0, BLOCK_M)
    off_hz = tl.program_id(1)
    off_n = tl.arange(0, HEAD_DIM)
    # load
-    o = tl.load(O + off_hz * HEAD_DIM * N_CTX + off_m[:, None] * HEAD_DIM +
-                off_n[None, :])
-    do = tl.load(DO + off_hz * HEAD_DIM * N_CTX + off_m[:, None] * HEAD_DIM +
-                 off_n[None, :]).to(tl.float32)
+    o = tl.load(O + off_hz * HEAD_DIM * N_CTX + off_m[:, None] * HEAD_DIM + off_n[None, :])
+    do = tl.load(DO + off_hz * HEAD_DIM * N_CTX + off_m[:, None] * HEAD_DIM + off_n[None, :]).to(tl.float32)
    delta = tl.sum(o * do, axis=1)
    # write-back
    tl.store(Delta + off_hz * N_CTX + off_m, delta)
@@ -183,32 +151,19 @@ def _attn_bwd_preprocess(

 # The main inner-loop logic for computing dK and dV.
@triton.jit
-def _attn_bwd_dkdv(
-        dk,
-        dv,  #
-        Q,
-        k,
-        v,
-        sm_scale,  #
-        DO,  #
-        M,
-        D,  #
-        k2q_index,
-        k2q_num,
-        max_q_blks,
-        variable_block_sizes,
-        # shared by Q/K/V/DO.
-        stride_tok,
-        stride_d,  #
-        H,
-        N_CTX,
-        BLOCK_M1: tl.constexpr,  #
-        BLOCK_N1: tl.constexpr,  #
-        HEAD_DIM: tl.constexpr,  #
-        # Filled in by the wrapper.
-    start_n,
-        start_m,
-        num_steps):
+def _attn_bwd_dkdv(dk, dv,  #
+                   Q, k, v, sm_scale,  #
+                   DO,  #
+                   M, D,  #
+                   k2q_index, k2q_num, max_q_blks,
+                   variable_block_sizes,
+                   # shared by Q/K/V/DO.
+                   stride_tok, stride_d,  #
+                   H, N_CTX, BLOCK_M1: tl.constexpr,  #
+                   BLOCK_N1: tl.constexpr,  #
+                   HEAD_DIM: tl.constexpr,  #
+                   # Filled in by the wrapper.
+                   start_n, start_m, num_steps):
    offs_m = start_m + tl.arange(0, BLOCK_M1)
    offs_n = start_n + tl.arange(0, BLOCK_N1)
    offs_k = tl.arange(0, HEAD_DIM)
@@ -217,20 +172,21 @@ def _attn_bwd_dkdv(
    # BLOCK_N1 must be a multiple of BLOCK_M1, otherwise the code wouldn't work.
    tl.static_assert(BLOCK_N1 % BLOCK_M1 == 0)
    step_m = BLOCK_M1
-    kv_blk = tl.program_id(0)  # Q-tile index
-    off_hz = tl.program_id(2)  # fused (batch, head)
-    b = off_hz // H
-    h = off_hz % H
+    kv_blk   = tl.program_id(0)          # Q-tile index
+    off_hz  = tl.program_id(2)          # fused (batch, head)
+    b       = off_hz // H
+    h       = off_hz %  H
    q_tiles = N_CTX // BLOCK_N1
    meta_base = ((b * H + h) * q_tiles + kv_blk)

-    q_blocks = tl.load(k2q_num + meta_base)  # int32
-    q_ptr = k2q_index + meta_base * max_q_blks  # ptr to list
+    q_blocks = tl.load(k2q_num  + meta_base)                 # int32
+    q_ptr    = k2q_index + meta_base * max_q_blks           # ptr to list
    block_size = tl.load(variable_block_sizes + kv_blk)
-
-    for blk_idx in range(q_blocks * 2):
-        block_sparse_offset = (tl.load(q_ptr + blk_idx // 2).to(tl.int32) * 2 +
-                               blk_idx % 2) * step_m
+    
+    
+        
+    for blk_idx in range(q_blocks*2):
+        block_sparse_offset = (tl.load(q_ptr + blk_idx//2).to(tl.int32)*2 + blk_idx%2) *step_m
        qT = tl.load(qT_ptrs + block_sparse_offset * stride_tok)
        # Load m before computing qk to reduce pipeline stall.
        offs_m = start_m + block_sparse_offset + tl.arange(0, BLOCK_M1)
@@ -256,32 +212,21 @@ def _attn_bwd_dkdv(
    return dk, dv


+
 # the main inner-loop logic for computing dQ
@triton.jit
-def _attn_bwd_dq(
-        dq,
-        q,
-        K,
-        V,  #
-        do,
-        m,
-        D,
-        # shared by Q/K/V/DO.
-        q2k_index,
-        q2k_num,
-        max_kv_blks,
-        variable_block_sizes,
-        stride_tok,
-        stride_d,  #
-        H,
-        N_CTX,  #
-        BLOCK_M2: tl.constexpr,  #
-        BLOCK_N2: tl.constexpr,  #
-        HEAD_DIM: tl.constexpr,
-        # Filled in by the wrapper.
-        start_m,
-        start_n,
-        num_steps):
+def _attn_bwd_dq(dq, q, K, V,  #
+                 do, m, D,
+                 # shared by Q/K/V/DO.
+                 q2k_index, q2k_num, max_kv_blks,
+                 variable_block_sizes,
+                 stride_tok, stride_d,  #
+                 H, N_CTX,  #
+                 BLOCK_M2: tl.constexpr,  #
+                 BLOCK_N2: tl.constexpr,  #
+                 HEAD_DIM: tl.constexpr,
+                 # Filled in by the wrapper.
+                 start_m, start_n, num_steps):
    offs_m = start_m + tl.arange(0, BLOCK_M2)
    offs_n = start_n + tl.arange(0, BLOCK_N2)
    offs_k = tl.arange(0, HEAD_DIM)
@@ -292,27 +237,27 @@ def _attn_bwd_dq(
    # BLOCK_M2 must be a multiple of BLOCK_N2, otherwise the code wouldn't work.
    tl.static_assert(BLOCK_M2 % BLOCK_N2 == 0)
    step_n = BLOCK_N2
-
-    q_blk = tl.program_id(0)  # Q-tile index
-    off_hz = tl.program_id(2)  # fused (batch, head)
-    b = off_hz // H
-    h = off_hz % H
+    
+    q_blk   = tl.program_id(0)          # Q-tile index
+    off_hz  = tl.program_id(2)          # fused (batch, head)
+    b       = off_hz // H
+    h       = off_hz %  H
    q_tiles = N_CTX // BLOCK_M2
    meta_base = ((b * H + h) * q_tiles + q_blk)

-    kv_blocks = tl.load(q2k_num + meta_base)  # int32
-    kv_ptr = q2k_index + meta_base * max_kv_blks  # ptr to list
-    block_size = tl.load(variable_block_sizes + q_blk)
-
-    for blk_idx in range(kv_blocks * 2):
-        block_sparse_offset = (tl.load(kv_ptr + blk_idx // 2).to(tl.int32) * 2 +
-                               blk_idx % 2) * step_n * stride_tok
+    kv_blocks = tl.load(q2k_num  + meta_base)                 # int32
+    kv_ptr    = q2k_index + meta_base * max_kv_blks           # ptr to list
+    
+    
+    for blk_idx in range(kv_blocks*2):
+        block_sparse_offset = (tl.load(kv_ptr + blk_idx//2).to(tl.int32)*2 + blk_idx%2) *step_n * stride_tok
+        block_size = tl.load(variable_block_sizes + blk_idx//2) - (blk_idx%2) * step_n
        kT = tl.load(kT_ptrs + block_sparse_offset)
        vT = tl.load(vT_ptrs + block_sparse_offset)
        qk = tl.dot(q, kT)
        p = tl.math.exp2(qk - m)
        mask = tl.arange(0, BLOCK_N2) < block_size.to(tl.int32)
-        p = tl.where(mask[None, :], p, 0.0)
+        p = tl.where(mask[None, :], p , 0.0)
        # Compute dP and dS.
        dp = tl.dot(do, vT).to(tl.float32)
        ds = p * (dp - Di[:, None])
@@ -324,37 +269,23 @@ def _attn_bwd_dq(
    return dq


+
@triton.jit
-def _attn_bwd(
-        Q,
-        K,
-        V,
-        sm_scale,  #
-        DO,  #
-        DQ,
-        DK,
-        DV,  #
-        M,
-        D,
-        q2k_index,
-        q2k_num,
-        max_kv_blks,
-        k2q_index,
-        k2q_num,
-        max_q_blks,
-        variable_block_sizes,
-        # shared by Q/K/V/DO.
-        stride_z,
-        stride_h,
-        stride_tok,
-        stride_d,  #
-        H,
-        N_CTX,  #
-        BLOCK_M1: tl.constexpr,  #
-        BLOCK_N1: tl.constexpr,  #
-        BLOCK_M2: tl.constexpr,  #
-        BLOCK_N2: tl.constexpr,  #
-        HEAD_DIM: tl.constexpr):
+def _attn_bwd(Q, K, V, sm_scale,  #
+              DO,  #
+              DQ, DK, DV,  #
+              M, D,
+              q2k_index, q2k_num, max_kv_blks,
+              k2q_index, k2q_num, max_q_blks,
+              variable_block_sizes,
+              # shared by Q/K/V/DO.
+              stride_z, stride_h, stride_tok, stride_d,  #
+              H, N_CTX,  #
+              BLOCK_M1: tl.constexpr,  #
+              BLOCK_N1: tl.constexpr,  #
+              BLOCK_M2: tl.constexpr,  #
+              BLOCK_N2: tl.constexpr,  #
+              HEAD_DIM: tl.constexpr):
    LN2 = 0.6931471824645996  # = ln(2)

    bhid = tl.program_id(2)
@@ -388,32 +319,20 @@ def _attn_bwd(
    k = tl.load(K + offs_n[:, None] * stride_tok + offs_k[None, :] * stride_d)
    v = tl.load(V + offs_n[:, None] * stride_tok + offs_k[None, :] * stride_d)

+
    num_steps = N_CTX // BLOCK_M1

    dk, dv = _attn_bwd_dkdv(  #
-        dk,
-        dv,  #
-        Q,
-        k,
-        v,
-        sm_scale,  #
+        dk, dv,  #
+        Q, k, v, sm_scale,  #
        DO,  #
-        M,
-        D,  #
-        k2q_index,
-        k2q_num,
-        max_q_blks,
+        M, D,  #
+        k2q_index, k2q_num, max_q_blks,
        variable_block_sizes,
-        stride_tok,
-        stride_d,  #
-        H,
-        N_CTX,  #
-        BLOCK_M1,
-        BLOCK_N1,
-        HEAD_DIM,  #
-        start_n,
-        start_m,
-        num_steps  #
+        stride_tok, stride_d,  #
+        H, N_CTX,  #
+        BLOCK_M1, BLOCK_N1, HEAD_DIM,  #
+        start_n, start_m, num_steps  #
    )

    dv_ptrs = DV + offs_n[:, None] * stride_tok + offs_k[None, :] * stride_d
@@ -438,88 +357,54 @@ def _attn_bwd(
    m = m[:, None]

    num_steps = N_CTX // BLOCK_N2
-    dq = _attn_bwd_dq(
-        dq,
-        q,
-        K,
-        V,  #
-        do,
-        m,
-        D,  #
-        q2k_index,
-        q2k_num,
-        max_kv_blks,
-        variable_block_sizes,
-        stride_tok,
-        stride_d,  #
-        H,
-        N_CTX,  #
-        BLOCK_M2,
-        BLOCK_N2,
-        HEAD_DIM,  #
-        start_m,
-        end_n,
-        num_steps  #
-    )
+    dq = _attn_bwd_dq(dq, q, K, V,  #
+                      do, m, D,  #
+                      q2k_index, q2k_num, max_kv_blks,
+                      variable_block_sizes,
+                      stride_tok, stride_d,  #
+                      H, N_CTX,  #
+                      BLOCK_M2, BLOCK_N2, HEAD_DIM,  #
+                      start_m, end_n, num_steps  #
+                      )
    # Write back dQ.
    dq_ptrs = DQ + offs_m[:, None] * stride_tok + offs_k[None, :] * stride_d
    dq *= LN2
    tl.store(dq_ptrs, dq)
+    
+    
+


 # ──────────────────────────── SPARSE ADDITION BEGIN ───────────────────────────
-def triton_block_sparse_attn_forward(q, k, v, q2k_index, q2k_num,
-                                     variable_block_sizes):
+def triton_block_sparse_attn_forward(q, k, v, q2k_index, q2k_num, variable_block_sizes):
    B, H, T, D = q.shape
    sm_scale = 1.0 / math.sqrt(D)
    max_kv_blks = q2k_index.shape[-1]
    assert T % 64 == 0, f"T must be a multiple of 64, but got {T}"
-    assert q2k_num.shape[
-        -1] == T // 64, f"shape mismatch, T // 64 = {T // 64}, q2k_num.shape[-2] = {q2k_num.shape[-2]}"
+    assert T // 64 == q2k_num.shape[-1], f"shape mismatch, T // 64 = {T // 64}, q2k_num.shape[-2] = {q2k_num.shape[-2]}"
    o = torch.empty_like(q)
    M = torch.empty((B, H, T), dtype=torch.float32, device=q.device)

    grid = lambda _: (triton.cdiv(T, 64), B * H, 1)
-    _attn_fwd_sparse[grid](q,
-                           k,
-                           v,
-                           sm_scale,
-                           q2k_index,
-                           q2k_num,
-                           max_kv_blks,
-                           variable_block_sizes,
-                           M,
-                           o,
-                           q.stride(0),
-                           q.stride(1),
-                           q.stride(2),
-                           q.stride(3),
-                           k.stride(0),
-                           k.stride(1),
-                           k.stride(2),
-                           k.stride(3),
-                           v.stride(0),
-                           v.stride(1),
-                           v.stride(2),
-                           v.stride(3),
-                           o.stride(0),
-                           o.stride(1),
-                           o.stride(2),
-                           o.stride(3),
-                           B,
-                           H,
-                           T,
-                           HEAD_DIM=D,
-                           STAGE=3)
+    _attn_fwd_sparse[grid](
+        q, k, v, sm_scale,
+        q2k_index, q2k_num, max_kv_blks,
+        variable_block_sizes,
+        M, o,
+        q.stride(0), q.stride(1), q.stride(2), q.stride(3),
+        k.stride(0), k.stride(1), k.stride(2), k.stride(3),
+        v.stride(0), v.stride(1), v.stride(2), v.stride(3),
+        o.stride(0), o.stride(1), o.stride(2), o.stride(3),
+        B, H, T,
+        HEAD_DIM=D, STAGE=3
+    )

    return o, M

-
-def triton_block_sparse_attn_backward(do, q, k, v, o, M, q2k_index, q2k_num,
-                                      k2q_index, k2q_num, variable_block_sizes):
+def triton_block_sparse_attn_backward(do, q, k, v, o, M, q2k_index, q2k_num, k2q_index, k2q_num, variable_block_sizes):
    assert do.is_contiguous()
    assert q.stride() == k.stride() == v.stride() == o.stride() == do.stride()
-
+    
    B, H, T, D = q.shape
    sm_scale = 1.0 / math.sqrt(D)
    dq = torch.empty_like(q)
@@ -535,49 +420,30 @@ def triton_block_sparse_attn_backward(do, q, k, v, o, M, q2k_index, q2k_num,
    pre_grid = (N_CTX // PRE_BLOCK, BATCH * N_HEAD)
    delta = torch.empty_like(M)
    _attn_bwd_preprocess[pre_grid](
-        o,
-        do,  #
+        o, do,  #
        delta,  #
-        BATCH,
-        N_HEAD,
-        N_CTX,  #
-        BLOCK_M=PRE_BLOCK,
-        HEAD_DIM=D  #
+        BATCH, N_HEAD, N_CTX,  #
+        BLOCK_M=PRE_BLOCK, HEAD_DIM=D  #
    )
-
+    
+    
    max_q_blks = k2q_index.shape[-1]
    max_kv_blks = q2k_index.shape[-1]
-
+    
    grid = (N_CTX // BLOCK_N1, 1, BATCH * N_HEAD)
    _attn_bwd[grid](
-        q,
-        arg_k,
-        v,
-        sm_scale,
-        do,
-        dq,
-        dk,
-        dv,  #
-        M,
-        delta,  #
-        q2k_index,
-        q2k_num,
-        max_kv_blks,
-        k2q_index,
-        k2q_num,
-        max_q_blks,
+        q, arg_k, v, sm_scale, do, dq, dk, dv,  #
+        M, delta,  #
+        q2k_index, q2k_num, max_kv_blks,
+        k2q_index, k2q_num, max_q_blks,
        variable_block_sizes,
-        q.stride(0),
-        q.stride(1),
-        q.stride(2),
-        q.stride(3),  #
-        N_HEAD,
-        N_CTX,  #
-        BLOCK_M1=BLOCK_M1,
-        BLOCK_N1=BLOCK_N1,  #
-        BLOCK_M2=BLOCK_M2,
-        BLOCK_N2=BLOCK_N2,  #
+        q.stride(0), q.stride(1), q.stride(2), q.stride(3),  #
+        N_HEAD, N_CTX,  #
+        BLOCK_M1=BLOCK_M1, BLOCK_N1=BLOCK_N1,  #
+        BLOCK_M2=BLOCK_M2, BLOCK_N2=BLOCK_N2,  #
        HEAD_DIM=D  #
    )

    return dq, dk, dv
+
+
@@ -639,7 +639,8 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {

    // store kq and vq

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

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

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

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

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

    return {o, l_vec};
@@ -0,0 +1,185 @@
+import torch
+try:
+    from vsa_cuda import block_sparse_fwd, block_sparse_bwd
+except ImportError:
+    block_sparse_fwd = None
+    block_sparse_bwd = None
+from vsa.block_sparse_attn_triton import triton_block_sparse_attn_forward, triton_block_sparse_attn_backward
+assert torch.__version__ >= "2.4.0", "VSA requires PyTorch 2.4.0 or higher"
+from vsa.index import map_to_index
+from typing import Tuple, Optional
+
+
+
+@torch.library.custom_op("vsa::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.int()
+    q2k_block_sparse_index, q2k_block_sparse_num = map_to_index(block_map)
+    o, M = triton_block_sparse_attn_forward(q, k, v, q2k_block_sparse_index, q2k_block_sparse_num, variable_block_sizes)
+    return o, M
+
+
+@torch.library.register_fake("vsa::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]:
+    q = q.contiguous()
+    k = k.contiguous()
+    v = v.contiguous()
+    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("vsa::block_sparse_attn_backward_triton", mutates_args=(), device_types="cuda")
+def block_sparse_attn_backward_triton(
+    grad_output_padded: torch.Tensor, 
+    q_padded: torch.Tensor, 
+    k_padded: torch.Tensor, 
+    v_padded: torch.Tensor, 
+    o_padded: torch.Tensor, 
+    M: torch.Tensor, 
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    grad_output_padded = grad_output_padded.contiguous()
+    q2k_block_sparse_index, q2k_block_sparse_num = map_to_index(block_map)
+    k2q_block_sparse_index, k2q_block_sparse_num = map_to_index(block_map.transpose(-1, -2))
+    dq, dk, dv = triton_block_sparse_attn_backward(grad_output_padded, q_padded, k_padded, v_padded, o_padded, M, q2k_block_sparse_index, q2k_block_sparse_num, k2q_block_sparse_index, k2q_block_sparse_num, variable_block_sizes)
+    return dq, dk, dv
+
+@torch.library.register_fake("vsa::block_sparse_attn_backward_triton")
+def _block_sparse_attn_backward_triton_fake(
+    grad_output_padded: torch.Tensor, 
+    q_padded: torch.Tensor, 
+    k_padded: torch.Tensor, 
+    v_padded: torch.Tensor, 
+    o_padded: torch.Tensor, 
+    M: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    grad_output_padded = grad_output_padded.contiguous()
+    dq = torch.empty_like(grad_output_padded)
+    dk = torch.empty_like(grad_output_padded)
+    dv = torch.empty_like(grad_output_padded)
+    return dq, dk, dv
+
+
+def backward_triton(ctx, grad_output1, grad_output2):
+    q_padded, k_padded, v_padded, o_padded, M, block_map, variable_block_sizes = ctx.saved_tensors
+    dq, dk, dv = block_sparse_attn_backward_triton(grad_output1, q_padded, k_padded, v_padded, o_padded, M, block_map, variable_block_sizes)
+    return dq, dk, dv, None, None
+
+
+def setup_context_triton(ctx, inputs, output):
+    q_padded, k_padded, v_padded, block_map, variable_block_sizes = inputs
+    o_padded, M = output
+    ctx.save_for_backward(q_padded, k_padded, v_padded, o_padded, M, block_map, variable_block_sizes)
+    
+block_sparse_attn_triton.register_autograd(backward_triton, setup_context=setup_context_triton)
+
+
+major, minor = torch.cuda.get_device_capability(0)
+
+if major == 9 and minor == 0:# check if H100
+    @torch.library.custom_op("vsa::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]:
+        q_padded = q_padded.contiguous()
+        k_padded = k_padded.contiguous()
+        v_padded = v_padded.contiguous()
+        q2k_block_sparse_index, q2k_block_sparse_num = map_to_index(block_map)
+        variable_block_sizes = variable_block_sizes.int()
+        o_padded, lse_padded = block_sparse_fwd(q_padded, k_padded, v_padded, q2k_block_sparse_index, q2k_block_sparse_num, variable_block_sizes)
+        return o_padded, lse_padded
+
+
+
+
+    @torch.library.register_fake("vsa::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]:
+        q_padded, k_padded, v_padded = [x.contiguous() for x in (q_padded, k_padded, v_padded)]
+        B, H, S, D = q_padded.shape
+        o_padded = torch.empty_like(q_padded)
+        lse_padded = torch.empty((B, H, S, 1), device=q_padded.device, dtype=torch.float32)
+        return o_padded, lse_padded
+
+
+    @torch.library.custom_op("vsa::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]:
+        grad_output_padded = grad_output_padded.contiguous()
+        k2q_block_sparse_index, k2q_block_sparse_num = map_to_index(block_map.transpose(-1, -2))
+        grad_q_padded, grad_k_padded, grad_v_padded = block_sparse_bwd(
+            q_padded, k_padded, v_padded, o_padded, lse_padded, grad_output_padded, k2q_block_sparse_index, k2q_block_sparse_num, variable_block_sizes
+        )
+        grad_q_padded = grad_q_padded.to(grad_output_padded.dtype)
+        grad_k_padded = grad_k_padded.to(grad_output_padded.dtype)
+        grad_v_padded = grad_v_padded.to(grad_output_padded.dtype)
+        return grad_q_padded, grad_k_padded, grad_v_padded
+
+    @torch.library.register_fake("vsa::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, torch.Tensor]:
+        torch._check(grad_output_padded.dtype == torch.bfloat16)
+        torch._check(lse_padded.dtype == torch.float32)
+        grad_output_padded = grad_output_padded.contiguous()
+        dq = torch.empty_like(grad_output_padded)
+        dk = torch.empty_like(grad_output_padded)
+        dv = torch.empty_like(grad_output_padded)
+        return dq, dk, dv
+
+
+    def backward_SM90(ctx, grad_output1, grad_output2):
+        q_padded, k_padded, v_padded, o_padded, lse_padded, block_map, variable_block_sizes= ctx.saved_tensors
+        dq, dk, dv = block_sparse_attn_backward_SM90(grad_output1, q_padded, k_padded, v_padded, o_padded, lse_padded, block_map, variable_block_sizes)
+        return dq, dk, dv, None, None
+
+    def setup_context_SM90(ctx, inputs, output):
+        q_padded, k_padded, v_padded, block_map, variable_block_sizes = inputs
+        o_padded, lse_padded = output
+        ctx.save_for_backward(q_padded, k_padded, v_padded, o_padded, lse_padded, block_map, variable_block_sizes)
+
+
+    block_sparse_attn_SM90.register_autograd(backward_SM90, setup_context=setup_context_SM90)
@@ -1,9 +1,9 @@
+
 ## pytorch sdpa version of block sparse ##
 import triton
 import triton.language as tl
 import torch

-
@triton.jit
 def topk_index_to_map_kernel(
    map_ptr,
@@ -26,7 +26,6 @@ def topk_index_to_map_kernel(
        index = tl.load(index_ptr_base + i * index_kv_stride)
        tl.store(map_ptr_base + index * map_kv_stride, 1.0)

-
@triton.jit
 def map_to_index_kernel(
    map_ptr,
@@ -60,7 +59,6 @@ def map_to_index_kernel(
        index_num_ptr + b * index_num_bs_stride + h * index_num_h_stride +
        q * index_num_q_stride, num)

-
 def topk_index_to_map(index: torch.Tensor,
                      num_kv_blocks: int,
                      transpose_map: bool = False):
@@ -108,7 +106,6 @@ def topk_index_to_map(index: torch.Tensor,

    return block_map

-
 def map_to_index(block_map: torch.Tensor):
    """
    Convert a block map to indices and counts.
@@ -0,0 +1,32 @@
+# Attention Kernel Used in FastVideo
+
+## VMoBA: Mixture-of-Block Attention for Video Diffusion Models (VMoBA)
+
+### Installation
+Please ensure that you have installed FlashAttention version **2.7.1 or higher**, as some interfaces have changed in recent releases.
+
+### Usage
+
+You can use `moba_attn_varlen` in the following ways:
+
+**Install from source:**
+```bash
+python setup.py install
+```
+
+**Import after installation:**
+```python
+from vmoba import moba_attn_varlen
+```
+
+**Or import directly from the project root:**
+```python
+from csrc.attn.vmoba_attn.vmoba import moba_attn_varlen
+```
+
+### Verify if you have successfully installed
+
+```bash
+python csrc/attn/vmoba_attn/vmoba/vmoba.py
+```
+
@@ -0,0 +1,24 @@
+# SPDX-License-Identifier: Apache-2.0
+
+from setuptools import find_packages, setup
+
+PACKAGE_NAME = "vmoba"
+VERSION = "0.0.0"
+AUTHOR = "JianzongWu"
+DESCRIPTION = "VMoBA: Mixture-of-Block Attention for Video Diffusion Models"
+URL = "https://github.com/KwaiVGI/VMoBA"
+
+setup(
+    name=PACKAGE_NAME,
+    version=VERSION,
+    author=AUTHOR,
+    description=DESCRIPTION,
+    url=URL,
+    packages=find_packages(),
+    classifiers=[
+        "Programming Language :: Python :: 3",
+        "License :: OSI Approved :: Apache Software License",
+    ],
+    python_requires='>=3.12',
+    install_requires=[]
+)
@@ -3,7 +3,7 @@
 import torch
 import pytest
 import random
-from fastvideo_kernel import moba_attn_varlen
+from csrc.attn.vmoba_attn.vmoba  import moba_attn_varlen

 def generate_test_data(batch_size, total_seqlen, num_heads, head_dim, dtype, device="cuda"):
    """
@@ -51,7 +51,7 @@ def generate_test_data(batch_size, total_seqlen, num_heads, head_dim, dtype, dev
@pytest.mark.parametrize("moba_topk", [2, 4])
@pytest.mark.parametrize("select_mode", ["topk", "threshold"])
@pytest.mark.parametrize("threshold_type", ["query_head", "head_global", "overall"])
-@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
+@pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16])
 def test_moba_attn_varlen_forward(
    batch_size, total_seqlen, num_heads, head_dim, moba_chunk_size, moba_topk, select_mode, threshold_type, dtype
 ):
@@ -0,0 +1,2 @@
+# SPDX-License-Identifier: Apache-2.0
+from .vmoba import moba_attn_varlen, process_moba_input, process_moba_output
@@ -1,4 +1,4 @@
-FROM nvidia/cuda:12.8.0-cudnn-devel-ubuntu22.04
+FROM nvidia/cuda:12.8.0-devel-ubuntu22.04

 ENV DEBIAN_FRONTEND=noninteractive

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

 COPY . .

@@ -55,12 +55,18 @@ RUN source $HOME/.local/bin/env && \
    echo 'source /opt/venv/bin/activate' >> /root/.bashrc && \
    echo 'if [ -n "$ZSH_VERSION" ] && [ -f ~/.zshrc ]; then . ~/.zshrc; elif [ -f ~/.bashrc ]; then . ~/.bashrc; fi' > /root/.profile

-# Install FastVideo Unified Kernel
+# Install STA (Sliding Tile Attention)
 RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
-    cd fastvideo-kernel && \
+    cd csrc/attn/sliding_tile_attn && \
    git submodule update --init --recursive && \
-    ./build.sh
+    python setup.py install

+# Install VSA 
+RUN source $HOME/.local/bin/env && \
+    source /opt/venv/bin/activate && \
+    cd csrc/attn/video_sparse_attn && \
+    git submodule update --init --recursive && \
+    python setup.py install

 EXPOSE 22
@@ -1,4 +1,4 @@
-FROM nvidia/cuda:12.8.0-cudnn-devel-ubuntu22.04
+FROM nvidia/cuda:12.8.0-devel-ubuntu22.04

 ENV DEBIAN_FRONTEND=noninteractive

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

 COPY . .

@@ -55,12 +55,18 @@ RUN source $HOME/.local/bin/env && \
    echo 'source /opt/venv/bin/activate' >> /root/.bashrc && \
    echo 'if [ -n "$ZSH_VERSION" ] && [ -f ~/.zshrc ]; then . ~/.zshrc; elif [ -f ~/.bashrc ]; then . ~/.bashrc; fi' > /root/.profile

-# Install FastVideo Unified Kernel
+# Install STA (Sliding Tile Attention)
 RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
-    cd fastvideo-kernel && \
+    cd csrc/attn/sliding_tile_attn && \
    git submodule update --init --recursive && \
-    ./build.sh
+    python setup.py install

+# Install VSA 
+RUN source $HOME/.local/bin/env && \
+    source /opt/venv/bin/activate && \
+    cd csrc/attn/video_sparse_attn && \
+    git submodule update --init --recursive && \
+    python setup.py install

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

 COPY . .

@@ -55,11 +55,18 @@ RUN source $HOME/.local/bin/env && \
    echo 'source /opt/venv/bin/activate' >> /root/.bashrc && \
    echo 'if [ -n "$ZSH_VERSION" ] && [ -f ~/.zshrc ]; then . ~/.zshrc; elif [ -f ~/.bashrc ]; then . ~/.bashrc; fi' > /root/.profile

-# Install FastVideo Unified Kernel
+# Install STA (Sliding Tile Attention)
 RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
-    cd fastvideo-kernel && \
+    cd csrc/attn/sliding_tile_attn && \
    git submodule update --init --recursive && \
-    ./build.sh
+    python setup.py install
+
+# Install VSA 
+RUN source $HOME/.local/bin/env && \
+    source /opt/venv/bin/activate && \
+    cd csrc/attn/video_sparse_attn && \
+    git submodule update --init --recursive && \
+    python setup.py install

 EXPOSE 22
@@ -55,12 +55,18 @@ RUN source $HOME/.local/bin/env && \
    echo 'source /opt/venv/bin/activate' >> /root/.bashrc && \
    echo 'if [ -n "$ZSH_VERSION" ] && [ -f ~/.zshrc ]; then . ~/.zshrc; elif [ -f ~/.bashrc ]; then . ~/.bashrc; fi' > /root/.profile

-# Install FastVideo Unified Kernel
+# Install STA (Sliding Tile Attention)
 RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
-    cd fastvideo-kernel && \
+    cd csrc/attn/sliding_tile_attn && \
    git submodule update --init --recursive && \
-    ./build.sh
+    python setup.py install

+# Install VSA 
+RUN source $HOME/.local/bin/env && \
+    source /opt/venv/bin/activate && \
+    cd csrc/attn/video_sparse_attn && \
+    git submodule update --init --recursive && \
+    python setup.py install

 EXPOSE 22
@@ -1,58 +0,0 @@
-FROM rocm/pytorch:rocm7.1_ubuntu22.04_py3.10_pytorch_release_2.9.1
-
-ENV DEBIAN_FRONTEND=noninteractive
-
-SHELL ["/bin/bash", "-c"]
-
-WORKDIR /FastVideo
-
-RUN apt-get update && apt-get install -y --no-install-recommends \
-    wget \
-    git \
-    ca-certificates \
-    openssh-server \
-    zsh \
-    vim \
-    curl \
-    gcc-11 \
-    g++-11 \
-    clang-11 \
-    && rm -rf /var/lib/apt/lists/*
-
-# Set up C++20 compilers for ThunderKittens
-RUN update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
-
-# Install uv and source its environment
-RUN curl -LsSf https://astral.sh/uv/install.sh | sh && \
-    echo 'source $HOME/.local/bin/env' >> /root/.bashrc
-
-# Copy just the pyproject.toml first to leverage Docker cache
-COPY pyproject_other.toml ./pyproject.toml
-
-# Create a dummy README to satisfy the installation
-RUN echo "# Placeholder" > README.md
-
-# Create and activate virtual environment with specific Python version and seed
-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
-
-COPY . .
-
-# Install dependencies using uv and set up shell configuration
-RUN source $HOME/.local/bin/env && \
-    source /opt/venv/bin/activate && \
-    uv pip install --no-cache-dir -e .[rocm] && \
-    git config --unset-all http.https://github.com/.extraheader || true && \
-    echo 'source /opt/venv/bin/activate' >> /root/.bashrc && \
-    echo 'if [ -n "$ZSH_VERSION" ] && [ -f ~/.zshrc ]; then . ~/.zshrc; elif [ -f ~/.bashrc ]; then . ~/.bashrc; fi' > /root/.profile
-
-# Install FastVideo Unified Kernel
-RUN source $HOME/.local/bin/env && \
-    source /opt/venv/bin/activate && \
-    cd fastvideo-kernel && \
-    git submodule update --init --recursive && \
-    ./build.sh --rocm
-
-EXPOSE 22
@@ -0,0 +1,26 @@
+# Minimal makefile for Sphinx documentation
+#
+
+# You can set these variables from the command line, and also
+# from the environment for the first two.
+SPHINXOPTS    ?=
+SPHINXBUILD   ?= sphinx-build
+SOURCEDIR     = source
+BUILDDIR      = build
+
+# Put it first so that "make" without argument is like "make help".
+help:
+	@$(SPHINXBUILD) -M help "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O)
+
+.PHONY: help Makefile
+
+# Catch-all target: route all unknown targets to Sphinx using the new
+# "make mode" option.  $(O) is meant as a shortcut for $(SPHINXOPTS).
+%: Makefile
+	@$(SPHINXBUILD) -M $@ "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O)
+
+clean:
+	@$(SPHINXBUILD) -M clean "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O)
+	rm -rf "$(SOURCEDIR)/getting_started/examples"
+	rm -rf "$(SOURCEDIR)/inference/examples"
+	rm -rf "$(SOURCEDIR)/training/examples"
@@ -1,39 +1,20 @@
-# FastVideo Documentation
+# FastVideo documents

-This directory contains the FastVideo documentation built with MkDocs.
-
-## Build the docs locally
+## Build the docs

 ```bash
-# Install dependencies
-pip install -r requirements-mkdocs.txt
+# Install dependencies.
+pip install -r requirements-docs.txt

-# Serve docs with live reload (recommended for development)
-mkdocs serve
-
-# Or build static site
-mkdocs build
+# Build the docs.
+make clean
+make html
 ```

-## View the docs
-
-### Development server (with live reload)
+## Open the docs with your browser

 ```bash
-mkdocs serve
+python -m http.server -d build/html/
 ```

-Then open your browser to: http://127.0.0.1:8000
-
-### Static build
-
-```bash
-mkdocs build
-python -m http.server -d site/
-```
-
-Then open your browser to: http://localhost:8000
-
-## Automatic Deployment
-
-Documentation is automatically built and deployed to GitHub Pages when changes are pushed to the `main` branch via the `.github/workflows/docs.yml` workflow.
+Launch your browser and open localhost:8000.
@@ -1,248 +0,0 @@
-# FastVideo API Reference
-
-This page contains the complete API reference for the FastVideo library.
-
-## fastvideo
-
-### Modules
-
-| Name | Description |
-|------|-------------|
-| [attention](#fastvideoattention) | Attention mechanisms and backends for video generation |
-| [configs](#fastvideoconfigs) | Configuration classes for pipelines, models, and sampling |
-| [distributed](#fastvideodistributed) | Distributed execution and communication utilities |
-| [entrypoints](#fastvideoentrypoints) | Main API entry points for video generation |
-| [models](#fastvideomodels) | Model implementations (transformers, VAEs, schedulers) |
-| [pipelines](#fastvideopipelines) | Core pipeline classes for video diffusion |
-| [training](#fastvideotraining) | Training utilities and helpers |
-| [workflow](#fastvideoworkflow) | Workflow management and orchestration |
-| [dataset](#fastvideodataset) | Dataset handling and preprocessing |
-| [layers](#fastvideolayers) | Custom neural network layers |
-| [platforms](#fastvideoplatforms) | Platform-specific implementations |
-| [utils](#fastvideoutils) | Utility functions and helpers |
-| [worker](#fastvideoworker) | Execution workers for video generation |
-
-## fastvideo.attention
-
-::: fastvideo.attention
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.configs
-
-::: fastvideo.configs
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-### Submodules
-
-#### fastvideo.configs.pipelines
-
-::: fastvideo.configs.pipelines
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-#### fastvideo.configs.models
-
-::: fastvideo.configs.models
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-#### fastvideo.configs.sample
-
-::: fastvideo.configs.sample
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-## fastvideo.distributed
-
-::: fastvideo.distributed
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.entrypoints
-
-::: fastvideo.entrypoints
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.models
-
-::: fastvideo.models
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-### Submodules
-
-#### fastvideo.models.registry
-
-::: fastvideo.models.registry
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-#### fastvideo.models.loader
-
-::: fastvideo.models.loader
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-## fastvideo.pipelines
-
-::: fastvideo.pipelines
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-### Submodules
-
-#### fastvideo.pipelines.composed_pipeline_base
-
-::: fastvideo.pipelines.composed_pipeline_base
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-#### fastvideo.pipelines.lora_pipeline
-
-::: fastvideo.pipelines.lora_pipeline
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-#### fastvideo.pipelines.pipeline_batch_info
-
-::: fastvideo.pipelines.pipeline_batch_info
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-#### fastvideo.pipelines.pipeline_registry
-
-::: fastvideo.pipelines.pipeline_registry
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-#### fastvideo.pipelines.stages
-
-::: fastvideo.pipelines.stages
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 4
-
-## fastvideo.training
-
-::: fastvideo.training
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.workflow
-
-::: fastvideo.workflow
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.dataset
-
-::: fastvideo.dataset
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.layers
-
-::: fastvideo.layers
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.platforms
-
-::: fastvideo.platforms
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
-
-## fastvideo.utils
-
-::: fastvideo.utils
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      heading_level: 3
-
-## fastvideo.worker
-
-::: fastvideo.worker
-    options:
-      show_source: true
-      show_root_heading: true
-      show_root_toc_entry: true
-      show_submodules: true
-      heading_level: 3
@@ -1,27 +0,0 @@
-# API Summary
-
-This page provides a quick overview of the main FastVideo API components.
-
-## Video Generator
-
-::: fastvideo.VideoGenerator
-    options:
-      show_root_heading: false
-      show_source: false
-      heading_level: 3
-
-## Initialization Configuration
-
-::: fastvideo.PipelineConfig
-    options:
-      show_root_heading: false
-      show_source: false
-      heading_level: 3
-
-## Sampling Configuration
-
-::: fastvideo.SamplingParam
-    options:
-      show_root_heading: false
-      show_source: false
-      heading_level: 3
@@ -1,43 +0,0 @@
-.vertical-table-header th.head:not(.stub) {
-    writing-mode: sideways-lr;
-    white-space: nowrap;
-    max-width: 0;
-}
-
-/* Keep header cell paragraph content tight (avoid CSS nesting for compatibility) */
-.vertical-table-header th.head:not(.stub) p {
-    margin: 0;
-}
-
-/* Image sizing classes */
-.image-small {
-    max-width: 200px;
-    height: auto;
-}
-
-.image-medium {
-    max-width: 400px;
-    height: auto;
-}
-
-.image-large {
-    max-width: 600px;
-    height: auto;
-}
-
-.image-full {
-    max-width: 100%;
-    height: auto;
-}
-
-/* Responsive images */
-img {
-    max-width: 100%;
-    height: auto;
-}
-
-/* Center images */
-.image-center {
-    display: block;
-    margin: 0 auto;
-}
@@ -1,6 +0,0 @@
-<svg width="160" height="93" viewBox="0 0 160 93" fill="none" xmlns="http://www.w3.org/2000/svg">
-<path d="M28.8511 91.66L57.6319 1.86368H64.5394L35.7585 91.66H28.8511Z" fill="#356CFF" stroke="#356CFF" stroke-width="2.30244"/>
-<path d="M15.0376 91.66L43.8185 1.86368H46.1209L17.3401 91.66H15.0376Z" fill="#356CFF" stroke="#356CFF" stroke-width="2.30244"/>
-<path d="M1.22217 91.66L30.003 1.86366H31.1543L2.3734 91.66H1.22217Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.15122"/>
-<path d="M71.4465 1.86483L42.666 91.6599H69.144L78.3538 58.2746H123.251L129.007 39.855H84.1099L89.866 22.5868H152.032L157.788 1.86483H71.4465Z" fill="#356CFF" stroke="#356CFF" stroke-width="2.30244"/>
-</svg>
@@ -1,18 +0,0 @@
-<svg width="252" height="105" viewBox="0 0 252 105" fill="none" xmlns="http://www.w3.org/2000/svg">
-<path d="M89.4843 55.5457H101.361L87.7028 101H74.638L89.4843 55.5457Z" fill="#356CFF"/>
-<path fill-rule="evenodd" clip-rule="evenodd" d="M96.0167 1.00057H112.645L118.583 48.273H104.924L103.737 39.7882H79.9827L67.5117 55.5457H85.3273L43.1638 101H28.3174L22.3789 55.5457H33.6621L38.4129 91.3031L58.604 68.2729H44.3515L96.0167 1.00057ZM100.768 13.1217L87.7028 29.4852H103.143L100.768 13.1217Z" fill="#356CFF"/>
-<path d="M37.2252 1.00057L22.3789 48.273H36.0375L40.7884 30.6974L62.6727 30.6974L69.6727 21.0005L43.7576 21.0004L46.7269 11.9096L77.6727 11.9096L86 1.00057L37.2252 1.00057Z" fill="#356CFF"/>
-<path fill-rule="evenodd" clip-rule="evenodd" d="M108.488 55.5457L94.2351 101C94.2351 101 105.518 101 120.959 101C136.399 101 144.078 93.0133 148.276 79.788C152.432 68.0157 153.027 55.5457 136.399 55.5457C119.771 55.5457 108.488 55.5457 108.488 55.5457ZM109.081 90.697L116.802 65.8487C116.802 65.8487 120.959 65.8487 132.242 65.8487C143.525 65.8487 137.586 78.5759 135.211 84.0304C133.307 88.4021 127.491 90.697 122.74 90.697C117.989 90.697 109.081 90.697 109.081 90.697Z" fill="#356CFF"/>
-<path d="M173.188 1.00056L168.625 11.9096C168.625 11.9096 149.386 11.9092 142.525 11.9095C135.664 11.9098 136.586 20.3944 141.337 20.3944H159.747C168.654 20.3944 166.961 33.6899 163.904 38.5761C160.467 44.0675 157.371 48.273 148.463 48.273L125.188 48.273L124 37.97L147.87 37.97C153.808 37.97 156.184 29.4852 151.433 29.4852H131.836C120.142 29.4852 125.897 1.00043 141.337 1.00043L173.188 1.00056Z" fill="#356CFF"/>
-<path d="M179.938 1.00056L175.688 11.9096L191.221 11.9096L179.938 48.273H192.409L203.692 11.9096L219.132 11.9095L223.289 1.00043L179.938 1.00056Z" fill="#356CFF"/>
-<path d="M161.341 55.5457H202.845L198.5 65.8487H169.654L167.279 73.7268H188.5L184.749 82.8177H164.31L161.934 90.697H190.251L186.624 101H146.494L161.341 55.5457Z" fill="#356CFF"/>
-<path fill-rule="evenodd" clip-rule="evenodd" d="M230.821 54.9391C255.169 54.9391 251.776 67.0602 249.231 77.9692C246.686 88.8783 240.917 101 217.757 101C194.596 101 195.606 88.8783 199.347 77.9692C203.089 67.0602 206.473 54.9391 230.821 54.9391ZM237.948 77.9692C239.984 70.6965 240.917 65.242 228.446 65.242C215.975 65.242 211.818 71.9087 210.037 77.9692C208.255 84.0298 208.255 91.3025 219.538 91.3025C230.821 91.3025 235.911 85.2419 237.948 77.9692Z" fill="#356CFF"/>
-<path d="M173.188 1.00056L168.625 11.9096C168.625 11.9096 149.386 11.9092 142.525 11.9095C135.664 11.9098 136.586 20.3944 141.337 20.3944M173.188 1.00056C173.188 1.00056 156.777 1.00043 141.337 1.00043M173.188 1.00056L141.337 1.00043M141.337 20.3944C146.088 20.3944 150.839 20.3944 159.747 20.3944M141.337 20.3944H159.747M159.747 20.3944C168.654 20.3944 166.961 33.6899 163.904 38.5761C160.467 44.0675 157.371 48.273 148.463 48.273M148.463 48.273C139.556 48.273 125.188 48.273 125.188 48.273M148.463 48.273L125.188 48.273M125.188 48.273L124 37.97M124 37.97C124 37.97 141.931 37.97 147.87 37.97M124 37.97L147.87 37.97M147.87 37.97C153.808 37.97 156.184 29.4852 151.433 29.4852M151.433 29.4852C146.682 29.4852 138.962 29.4852 131.836 29.4852M151.433 29.4852H131.836M131.836 29.4852C120.142 29.4852 125.897 1.00043 141.337 1.00043M37.2252 1.00057L22.3789 48.273H36.0375L40.7884 30.6974L62.6727 30.6974L69.6727 21.0005L43.7576 21.0004L46.7269 11.9096L77.6727 11.9096L86 1.00057L37.2252 1.00057ZM96.0167 1.00057H112.645L118.583 48.273H104.924L103.737 39.7882H79.9827L67.5117 55.5457H85.3273L43.1638 101H28.3174L22.3789 55.5457H33.6621L38.4129 91.3031L58.604 68.2729H44.3515L96.0167 1.00057ZM87.7028 29.4852L100.768 13.1217L103.143 29.4852H87.7028ZM89.4843 55.5457H101.361L87.7028 101H74.638L89.4843 55.5457ZM108.488 55.5457L94.2351 101C94.2351 101 105.518 101 120.959 101C136.399 101 144.078 93.0133 148.276 79.788C152.432 68.0157 153.027 55.5457 136.399 55.5457C119.771 55.5457 108.488 55.5457 108.488 55.5457ZM116.802 65.8487L109.081 90.697C109.081 90.697 117.989 90.697 122.74 90.697C127.491 90.697 133.307 88.4021 135.211 84.0304C137.586 78.5759 143.525 65.8487 132.242 65.8487C120.959 65.8487 116.802 65.8487 116.802 65.8487ZM179.938 1.00056L175.688 11.9096L191.221 11.9096L179.938 48.273H192.409L203.692 11.9096L219.132 11.9095L223.289 1.00043L179.938 1.00056ZM161.341 55.5457H202.845L198.5 65.8487H169.654L167.279 73.7268H188.5L184.749 82.8177H164.31L161.934 90.697H190.251L186.624 101H146.494L161.341 55.5457ZM230.821 54.9391C255.169 54.9391 251.776 67.0602 249.231 77.9692C246.686 88.8783 240.917 101 217.757 101C194.596 101 195.606 88.8783 199.347 77.9692C203.089 67.0602 206.473 54.9391 230.821 54.9391ZM228.446 65.242C240.917 65.242 239.984 70.6965 237.948 77.9692C235.911 85.2419 230.821 91.3025 219.538 91.3025C208.255 91.3025 208.255 84.0298 210.037 77.9692C211.818 71.9087 215.975 65.242 228.446 65.242Z" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M15.2524 55.5451L21.191 100.999L24.7541 100.999L18.8156 55.5451L15.2524 55.5451Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M8.12646 55.5451L14.065 100.999L15.2527 100.999L9.31417 55.5451L8.12646 55.5451Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M1 55.5451L6.93853 100.999L7.53239 100.999L1.59385 55.5451L1 55.5451Z" fill="#356CFF" stroke="#356CFF" stroke-width="0.593853"/>
-<path d="M15.2524 48.2724L30.0988 1H33.6619L18.8156 48.2724H15.2524Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M8.12646 48.2724L22.9728 1H24.1605L9.31417 48.2724H8.12646Z" fill="#356CFF" stroke="#356CFF" stroke-width="1.18771"/>
-<path d="M1 48.2724L15.8463 1H16.4402L1.59385 48.2724H1Z" fill="#356CFF" stroke="#356CFF" stroke-width="0.593853"/>
-<path d="M85.3271 55.5457H67.5116L87 12.7363L44.3513 68.2729H58.6038L43.1636 101L85.3271 55.5457Z" fill="#FDC717" stroke="#FDC717" stroke-width="1.18771" stroke-miterlimit="16"/>
-</svg>
@@ -1,186 +0,0 @@
-# Adding a New Attention Backend
-
-FastVideo allows integrating new attention mechanisms easily. This guide walks you through adding a new backend (e.g., `MyNewAttn`).
-
-## 1. Implement the Backend (Python)
-
-Create a new file in `fastvideo/attention/backends/` (e.g., `mynew_attn.py`).
-
-Your implementation should inherit from `AttentionBackend` defined in `abstract.py`.
-
-```python
-# fastvideo/attention/backends/mynew_attn.py
-import torch
-from .abstract import AttentionBackend
-# Import the context manager to access metadata (optional)
-from fastvideo.forward_context import get_forward_context
-
-# Import compiled kernel if applicable (see Section 2)
-try:
-    # Import from the top-level package
-    from fastvideo_kernel import my_compiled_attn_func
-except ImportError:
-    my_compiled_attn_func = None
-
-class MyNewAttnBackend(AttentionBackend):
-    def process_inputs(self, q, k, v, **kwargs):
-        # Pre-process inputs if necessary
-        return q, k, v
-
-    def forward(self, q, k, v, **kwargs):
-        # Optional: Access extra metadata passed via ForwardContext
-        # Only needed if your backend requires global state (e.g. window_size)
-        try:
-            context = get_forward_context()
-            metadata = context.attn_metadata
-            # Example: window_size = metadata.window_size
-        except (AssertionError, AttributeError):
-            # Handle case where context is not set (e.g. standard inference)
-            pass
-
-        if my_compiled_attn_func is not None:
-            return my_compiled_attn_func(q, k, v)
-        else:
-            # Fallback implementation (e.g., Triton or pure PyTorch)
-            return self.fallback_impl(q, k, v)
-```
-
-## 2. Passing Extra Information via ForwardContext (Optional)
-
-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.
-
-## 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
-// 1. Declare external function
-#ifdef COMPILE_MYNEW_ATTN
-extern torch::Tensor mynew_attn_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v);
-#endif
-
-// 2. Register in module
-PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
-    // ... other kernels ...
-
-#ifdef COMPILE_MYNEW_ATTN
-    m.def("mynew_attn_fwd", torch::wrap_pybind_function(mynew_attn_forward), "My New Attention Forward");
-#endif
-}
-```
-
-### C. Update CMakeLists.txt
-
-Update `fastvideo-kernel/CMakeLists.txt` to compile your new files.
-
-**Case 1: General CUDA Kernel (Runs on all GPUs)**
-Add your source file directly to `EXTENSION_SOURCES` and define the compilation flag.
-
-```cmake
-# Add to EXTENSION_SOURCES
-list(APPEND EXTENSION_SOURCES csrc/attention/mynew_attn.cu)
-
-# Add compilation definition for common_extension.cpp
-list(APPEND COMPILE_DEFS COMPILE_MYNEW_ATTN)
-```
-
-**Case 2: ThunderKittens Kernel (Hopper H100 Only)**
-If your kernel uses ThunderKittens (TK), it requires specific architecture flags (`sm_90a`). Add it inside the `ENABLE_TK_KERNELS` block.
-
-```cmake
-if(ENABLE_TK_KERNELS)
-    # Add source only if TK is enabled
-    list(APPEND EXTENSION_SOURCES csrc/attention/mynew_attn_tk.cu)
-    
-    # Add definition to guard registration
-    list(APPEND COMPILE_DEFS TK_COMPILE_MYNEW_ATTN)
-endif()
-```
-
-### D. Expose in Python Ops
-
-Update `fastvideo-kernel/python/fastvideo_kernel/ops.py` to make the function importable and handle fallbacks gracefully.
-
-```python
-# fastvideo-kernel/python/fastvideo_kernel/ops.py
-
-# Try to load C++ extension symbols
-try:
-    from fastvideo_kernel._C import fastvideo_kernel_ops
-    mynew_attn_fwd = getattr(fastvideo_kernel_ops, "mynew_attn_fwd", None)
-except ImportError:
-    mynew_attn_fwd = None
-
-def my_compiled_attn_func(q, k, v):
-    # Runtime check: use C++ kernel if available, else fallback
-    if mynew_attn_fwd is not None:
-        return mynew_attn_fwd(q, k, v)
-    else:
-        # Call Triton/Python fallback
-        return mynew_attn_triton(q, k, v)
-```
-
-### E. Expose in Package Init
-
-Update `fastvideo-kernel/python/fastvideo_kernel/__init__.py` to export the function.
-
-```python
-from fastvideo_kernel.ops import (
-    my_compiled_attn_func,
-    # ...
-)
-
-__all__ = [
-    "my_compiled_attn_func",
-    # ...
-]
-```
-
-## 4. Register the Backend
-
-Update `fastvideo/attention/backends/__init__.py` to export your new class.
-
-```python
-from .mynew_attn import MyNewAttnBackend
-```
-
-## 5. Platform Integration
-
-If your backend requires specific platform checks (e.g., checking for H100 support), handle that in `fastvideo/platforms/cuda.py` or within your backend's `__init__`.
-
-## 6. Add Documentation
-
-Create a new documentation page for your backend to explain its usage, installation (if custom kernels are needed), and features.
-
-1. **Create Directory**: `docs/attention/mynew_attn/`
-2. **Create Index**: `docs/attention/mynew_attn/index.md`
-3. **Update Navigation**: Add an entry to `mkdocs.yml` under the "Attention" tab.
-
-## Checklist
-
-* [ ] Created `fastvideo/attention/backends/mynew_attn.py`.
-* [ ] (Optional) Added CUDA kernels in `fastvideo-kernel/csrc/attention/`.
-* [ ] (Optional) Updated `common_extension.cpp` and `CMakeLists.txt`.
-* [ ] (Optional) Exposed kernel in `fastvideo-kernel/python/fastvideo_kernel/ops.py`.
-* [ ] (Optional) Exported kernel in `fastvideo-kernel/python/fastvideo_kernel/__init__.py`.
-* [ ] Implemented `forward` method respecting the standard signature.
-* [ ] Added unit tests in `tests/`.
-* [ ] Added documentation in `docs/attention/` and updated `mkdocs.yml`.
@@ -1,53 +0,0 @@
-# FastVideo Attention Kernels
-
-FastVideo provides highly optimized custom attention kernels to accelerate video generation.
-
-## Supported Kernels
-
-* **[Video Sparse Attention (VSA)](vsa/index.md)**: Sparse attention mechanism selecting top-k blocks.
-* **[Sliding Tile Attention (STA)](sta/index.md)**: Optimized attention for window-based video generation.
-
-## General Build Instructions
-
-These instructions apply to building the `fastvideo-kernel` package from source, which includes both STA and VSA kernels.
-
-### Prerequisites
-
-* **PyTorch**: 2.5.0+
-* **CUDA**: 12.4+ (12.8 recommended for best performance)
-* **C++ Compiler**: GCC 11+ (C++20 support required for ThunderKittens)
-
-Install system dependencies:
-
-```bash
-sudo apt update
-sudo apt install -y gcc-11 g++-11 clang-11 ninja-build
-
-# Set gcc-11 as default
-sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
-```
-
-Set up your CUDA environment variables (adjust version as needed):
-
-```bash
-export CUDA_HOME=/usr/local/cuda-12.8
-export PATH=${CUDA_HOME}/bin:${PATH}
-export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
-```
-
-### Compile and Install
-
-Clone the repository and build the kernel:
-
-```bash
-# Clone recursively to get ThunderKittens submodule
-git clone https://github.com/hao-ai-lab/FastVideo.git
-cd FastVideo/fastvideo-kernel
-
-# Build and install
-./build.sh
-```
-
-The build script automatically detects your GPU architecture:
-* **H100 (sm_90a)**: Compiles optimized C++ ThunderKittens kernels.
-* **Other (A100, etc.)**: Skips C++ compilation; installs Python package with Triton kernels.
@@ -1,36 +0,0 @@
-# Sliding Tile Attention (STA)
-
-Optimized attention for window-based video generation (e.g., HunyuanVideo).
-
-## Installation
-
-STA is included in the `fastvideo-kernel` package. See the [main Attention page](../index.md) for build instructions.
-
-## Usage
-
-```python
-from fastvideo_kernel import sliding_tile_attention
-
-# q, k, v: [batch_size, num_heads, seq_length, head_dim]
-# window_size: List of (t, h, w) tiles. Tile size is (6, 8, 8).
-# text_length: Number of text tokens (0-256)
-
-out = sliding_tile_attention(
-    q, k, v, 
-    window_size=[(3, 3, 3)],  # Example window
-    text_length=256
-)
-```
-
-## Citation
-
-If you use Sliding Tile Attention in your research, please cite:
-
-```bibtex
-@article{zhang2025fast,
-  title={Fast video generation with sliding tile attention},
-  author={Zhang, Peiyuan and Chen, Yongqi and Su, Runlong and Ding, Hangliang and Stoica, Ion and Liu, Zhengzhong and Zhang, Hao},
-  journal={arXiv preprint arXiv:2502.04507},
-  year={2025}
-}
-```
@@ -1,36 +0,0 @@
-# Video Sparse Attention (VSA)
-
-Sparse attention mechanism selecting top-k blocks.
-
-## Installation
-
-VSA is included in the `fastvideo-kernel` package. See the [main Attention page](../index.md) for build instructions.
-
-## Usage
-
-```python
-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
-# topk: Number of blocks to attend
-
-output = video_sparse_attn(
-    q, k, v, 
-    variable_block_sizes=block_sizes, 
-    topk=32
-)
-```
-
-## Citation
-
-If you use Video Sparse Attention in your research, please cite:
-
-```bibtex
-@article{zhang2025vsa,
-  title={Vsa: Faster video diffusion with trainable sparse attention},
-  author={Zhang, Peiyuan and Chen, Yongqi and Huang, Haofeng and Lin, Will and Liu, Zhengzhong and Stoica, Ion and Xing, Eric and Zhang, Hao},
-  journal={arXiv preprint arXiv:2505.13389},
-  year={2025}
-}
-```
@@ -1,53 +0,0 @@
-# Profiling FastVideo
-
-!!! warning
-    Profiling is only intended for FastVideo developers and maintainers to understand the proportion of time spent in different parts of the codebase. **FastVideo end-users should never turn on profiling** as it will significantly slow down inference.
-
-## Profiling with PyTorch
-
-FastVideo exposes a process-wide torch profiler that you can enable via environment variables. Set `FASTVIDEO_TORCH_PROFILER_DIR` to an absolute directory path to start collecting traces, and specify the regions you want recorded with `FASTVIDEO_TORCH_PROFILE_REGIONS`:
-
-```bash
-FASTVIDEO_TORCH_PROFILER_DIR=/mnt/traces/fastvideo \
-FASTVIDEO_TORCH_PROFILE_REGIONS="profiler_region_model_loading,profiler_region_training_step"
-```
-
-All profiled regions must be registered in `fastvideo.profiler`; the current list includes:
-
- `profiler_region_model_loading` — pipeline/module loading
- `profiler_region_inference_pre_denoising`
- `profiler_region_inference_denoising`
- `profiler_region_inference_post_denoising`
- `profiler_region_training_checkpoint_saving`
- `profiler_region_training_dit`
- `profiler_region_training_validation`
- `profiler_region_training_epoch`
- `profiler_region_training_step`
- `profiler_region_training_backward`
- `profiler_region_training_optimizer`
- `profiler_region_distillation_teacher_forward`
- `profiler_region_distillation_student_forward`
- `profiler_region_distillation_loss`
- `profiler_region_distillation_update`
-
-While profiling is enabled, FastVideo records additional annotations:
-
- `fastvideo.region::<name>` spans are emitted when entering a region.
- `fastvideo.profiler.enable_collection` / `fastvideo.profiler.disable_collection` events mark when torch profiler collection is toggled on or off.
-
-Only one profiler instance is created per process; subsequent pipelines reuse the same controller. If you set `FASTVIDEO_TORCH_PROFILE_REGIONS` incorrectly (e.g. misspelled name), FastVideo logs a warning and ignores that entry.
-
-Additional knobs:
-
- `FASTVIDEO_TORCH_PROFILER_RECORD_SHAPES`
- `FASTVIDEO_TORCH_PROFILER_WITH_PROFILE_MEMORY`
- `FASTVIDEO_TORCH_PROFILER_WITH_STACK`
- `FASTVIDEO_TORCH_PROFILER_WITH_FLOPS`
-
-Traces can be visualized using <https://ui.perfetto.dev/>.
-
-### Best Practices
-
- Keep the profiled step count small; traces can be large and slow down job shutdown while the profiler flushes data.
- After profiling, clean up trace directories to avoid filling disk storage.
- When adding new regions, register them in `fastvideo.profiler` and wrap the corresponding code block with `with self.profiler_controller.region("your_region"):` or the `@profile_region` decorator.
@@ -1,131 +0,0 @@
-# Testing in FastVideo
-
-This guide explains how to add and run tests in FastVideo. The testing suite is divided into several categories to ensure correctness across components, training workflows, and inference quality.
-
-## Test Types
-
-* **Unit Tests**: Located in `fastvideo/tests/dataset`, `fastvideo/tests/entrypoints`, and `fastvideo/tests/workflow`. These test individual functions and classes.
-* **Component Tests**: Located in `fastvideo/tests/encoders`, `fastvideo/tests/transformers`, and `fastvideo/tests/vaes`. These verify the loading and basic functionality of model components.
-* **SSIM Tests**: Located in `fastvideo/tests/ssim`. These are regression tests that compare generated videos against reference videos using the Structural Similarity Index Measure (SSIM) to detect quality degradation.
-* **Training Tests**: Located in `fastvideo/tests/training`. These validate training loops, loss calculations, and specific training techniques like LoRA, Distillation, and VSA.
-* **Inference Tests**: Located in `fastvideo/tests/inference`. These test specialized inference pipelines and optimizations (e.g., STA, V-MoBA).
-
-For now, we will focus on **SSIM Tests**.
-
-## SSIM Tests
-
-SSIM tests are located in `fastvideo/tests/ssim`. These tests generate videos using specific models and parameters, and compare them against reference videos to ensure that changes in the codebase do not degrade generation quality or alter the output unexpectedly.
-
-!!! note
-    If you are adding an SSIM test, this serves as a safeguard. Any future code changes that break or cause errors with the specific arguments and configurations you defined will trigger a failure. Therefore, it is important to include multiple settings and arguments that cover the core features of your new pipeline to ensure robust regression testing.
-
-### Directory Structure
-
-```
-fastvideo/tests/ssim/
-├── <GPU>_reference_videos/   # Reference videos organized by GPU type (e.g., L40S_reference_videos)
-│   ├── <Model_Name>/
-│   │   ├── <Backend>/        # e.g., FLASH_ATTN, TORCH_SDPA
-│   │   │   └── <Video_File>
-├── test_causal_similarity.py
-├── test_inference_similarity.py
-├── update_reference_videos.sh
-└── ...
-```
-
-### Adding a New SSIM Test
-
-To add a new SSIM test, follow these steps:
-
-1. **Create or Update a Test File**: You can add a new test function to an existing file (like `test_inference_similarity.py`) or create a new one if testing a distinct category of models.
-
-2. **Define Model Parameters**: Define the configuration for the model you want to test. This includes model path, dimensions, inference steps, and other generation parameters. **Note:** Consider using lower `num_inference_steps` or reduced resolution (e.g., 480p instead of 720p) to keep test execution time reasonable, provided it doesn't compromise the test's ability to detect regression.
-
-   ```python
-   MY_MODEL_PARAMS = {
-       "num_gpus": 1,
-       "model_path": "organization/model-name",
-       "height": 480,
-       "width": 832,
-       "num_frames": 45,
-       "num_inference_steps": 20,
-       # ... other parameters
-   }
-   ```
-
-3. **Implement the Test Function**:
-   * Use `pytest.mark.parametrize` to run the test with different prompts, backends, and models.
-   * Set the attention backend environment variable.
-   * Initialize the `VideoGenerator`.
-   * Generate the video.
-   * Compare the generated video with the reference video using `compute_video_ssim_torchvision`.
-
-   Example structure:
-
-   ```python
-   @pytest.mark.parametrize("prompt", TEST_PROMPTS)
-   @pytest.mark.parametrize("ATTENTION_BACKEND", ["FLASH_ATTN"])
-   def test_my_model_similarity(prompt, ATTENTION_BACKEND):
-       # Setup output directories
-       # ...
-
-       # Initialize Generator
-       generator = VideoGenerator.from_pretrained(...)
-       generator.generate_video(prompt, ...)
-
-       # Compare with Reference
-       ssim_values = compute_video_ssim_torchvision(
-           reference_path, generated_path, use_ms_ssim=True
-       )
-       assert ssim_values[0] >= 0.98  # Threshold
-   ```
-
-4. **Reference Videos**:
-   * When running the test for the first time (or when updating the reference), the test will fail because the reference video is missing. The generated video will be saved in `fastvideo/tests/ssim/generated_videos`.
-   * Inspect the generated video to ensure it meets quality expectations.
-   * Move the generated video to the appropriate reference folder: `fastvideo/tests/ssim/<GPU>_reference_videos/<Model>/<Backend>/`.
-   * You can use the helper script `update_reference_videos.sh` to automate copying videos from `generated_videos` to `L40S_reference_videos`. Note: Check the script to ensure paths match your environment (it defaults to `L40S_reference_videos`).
-
-### Running Tests Locally
-
-To run the SSIM tests locally:
-
-```bash
-pytest fastvideo/tests/ssim/ -vs
-```
-
-Ensure you have the necessary GPUs available as defined in your test parameters.
-
-## Modal Workflow
-
-FastVideo uses [Modal](https://modal.com/) for running tests in a CI environment. The workflow scripts are located in `fastvideo/tests/modal/`.
-
-### `pr_test.py`
-
-The main entry point for CI tests is `fastvideo/tests/modal/pr_test.py`. This script defines Modal functions that execute the pytest suites on specific hardware (e.g., L40S, H100).
-
-### Updating Modal Configuration
-
-If you add a new test that requires:
-* **Different GPU Hardware**: You may need to change the `@app.function(gpu=...)` decorator.
-* **Longer Execution Time**: Increase the `timeout` parameter.
-* **New Environment Variables/Secrets**: Add them to `secrets=[...]` or the image environment. For example, if your model is gated on Hugging Face, ensure `HF_API_KEY` is passed.
-
-For SSIM tests, the `run_ssim_tests` function in `pr_test.py` currently runs:
-
-```python
-@app.function(gpu="L40S:2", image=image, timeout=2700, secrets=[modal.Secret.from_dict({"HF_API_KEY": os.environ.get("HF_API_KEY", "")})])
-def run_ssim_tests():
-    run_test("hf auth login --token $HF_API_KEY && pytest ./fastvideo/tests/ssim -vs")
-```
-
-If your new test file is inside `fastvideo/tests/ssim`, it will automatically be picked up by this command. However, ensure that the `gpu="L40S:2"` configuration is sufficient for your model. If your model requires more GPUs (e.g., 4 or 8), you might need to create a separate Modal function or update the existing one.
-
-### Workflow Scripts
-
-The shell script that triggers these tests in the CI pipeline is located at `.buildkite/scripts/pr_test.sh`. If you add a new test category (e.g., a new folder outside of `ssim`), you will need to:
-1. Add a new function in `fastvideo/tests/modal/pr_test.py`.
-2. Add a new case in `.buildkite/scripts/pr_test.sh` to handle the new test type.
-
-!!! note
-    If you are a maintainer, you'll need to finally manually update the workflow script in Buildkite. Otherwise, a maintainer will help you update.
@@ -1,12 +0,0 @@
-# 💡 Examples
-
-A collection of examples demonstrating usage of FastVideo.
-
-All documented examples are autogenerated using [generate_examples.py](https://github.com/hao-ai-lab/FastVideo/blob/main/docs/generate_examples.py) from examples found in the [examples](https://github.com/hao-ai-lab/FastVideo/tree/main/examples) directory.
-
-## Examples
-
- [Examples Distillation Index](distillation/examples/examples_distillation_index.md)
- [Examples Training Index](training/examples/examples_training_index.md)
- [Examples Inference Index](inference/examples/examples_inference_index.md)
-
@@ -1,45 +0,0 @@
-
-# 🔧 Installation
-
-FastVideo supports the following hardware platforms:
-
- [NVIDIA CUDA](installation/gpu.md)
- [Apple silicon](installation/mps.md)
-
-## Quick Installation
-
-### Using pip
-
-```bash
-# Create and activate a new conda environment
-conda create -n fastvideo python=3.12
-conda activate fastvideo
-
-pip install fastvideo
-```
-
-### From source
-
-```bash
-git clone https://github.com/hao-ai-lab/FastVideo.git
-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+
- **Apple Silicon**: macOS 12.0+ with M1/M2/M3 chips
- **CPU**: x86_64 architecture (for CPU-only inference)
-
-## Next Steps
-
- [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
@@ -1,83 +0,0 @@
-# 🚀 Quick Start
-
-Get up and running with FastVideo in minutes!
-
-## Installation
-
-First, install FastVideo:
-
-```bash
-# Create and activate a new conda environment
-conda create -n fastvideo python=3.12
-conda activate fastvideo
-
-# Install FastVideo
-pip install fastvideo
-```
-
-Also optionally install flash-attn:
-
-```bash
-pip install flash-attn --no-build-isolation
-```
-
-## Basic Usage
-
-### Text-to-Video Generation
-
-```python
-from fastvideo import VideoGenerator
-
-def main():
-    # Create a video generator with a pre-trained model
-    generator = VideoGenerator.from_pretrained(
-        "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
-        num_gpus=1,  # Adjust based on your hardware
-    )
-
-    # Define a prompt for your video
-    prompt = "A curious raccoon peers through a vibrant field of yellow sunflowers, its eyes wide with interest."
-
-    # Generate the video
-    video = generator.generate_video(
-        prompt,
-        return_frames=True,  # Also return frames from this call (defaults to False)
-        output_path="my_videos/",  # Controls where videos are saved
-        save_video=True
-    )
-
-if __name__ == '__main__':
-    main()
-```
-
-### Image-to-Video Generation
-
-```python
-from fastvideo import VideoGenerator, SamplingParam
-
-def main():
-    # Create the generator
-    model_name = "Wan-AI/Wan2.1-I2V-14B-480P-Diffusers"
-    generator = VideoGenerator.from_pretrained(model_name, num_gpus=1)
-
-    # Set up parameters with an initial image
-    sampling_param = SamplingParam.from_pretrained(model_name)
-    sampling_param.image_path = "https://huggingface.co/datasets/huggingface/documentation-images/resolve/main/diffusers/astronaut.jpg"
-    sampling_param.num_frames = 107
-
-    # Generate video based on the image
-    prompt = "A photograph coming to life with gentle movement"
-    generator.generate_video(prompt, sampling_param=sampling_param,
-                             output_path="my_videos/",
-                             save_video=True)
-
-if __name__ == '__main__':
-    main()
-```
-
-## Next Steps
-
- [Installation Guide](installation.md) - Detailed installation instructions
- [Configuration](../inference/configuration.md) - Learn about configuration options
- [Examples](../inference/examples/) - Explore more examples
- [Optimizations](../inference/optimizations.md) - Performance optimization tips
@@ -1,42 +0,0 @@
-# V1 API
-
-FastVideo's V1 API provides a streamlined interface for video generation tasks with powerful customization options. This page documents the primary components of the API.
-
-## Video Generator
-
-This class will be the primary Python API for generating videos and images.
-
-::: fastvideo.entrypoints.video_generator.VideoGenerator
-    options:
-      show_root_heading: true
-      show_source: false
-      members:
-        - from_pretrained
-      heading_level: 3
-
-`VideoGenerator.from_pretrained()` should be the primary way of creating a new video generator.
-
-## Configuring FastVideo
-
-The following two classes `PipelineConfig` and `SamplingParam` are used to configure initialization and sampling parameters, respectively.
-
-### PipelineConfig
-
-::: fastvideo.configs.pipelines.base.PipelineConfig
-    options:
-      show_root_heading: true
-      show_source: false
-      members:
-        - from_pretrained
-        - dump_to_json
-      heading_level: 4
-
-### SamplingParam
-
-::: fastvideo.configs.sample.base.SamplingParam
-    options:
-      show_root_heading: true
-      show_source: false
-      members:
-        - from_pretrained
-      heading_level: 4
@@ -1,82 +0,0 @@
-# Compatibility Matrix
-
-The table below shows every supported model and optimizations supported for them.
-
-The symbols used have the following meanings:
-
- ✅ = Full compatibility
- ❌ = No compatibility
- ⭕ = Does not apply to this model
-
-## Models x Optimization
-
-The `HuggingFace Model ID` can be directly pass to `from_pretrained()` methods and FastVideo will use the optimal default parameters when initializing and generating videos.
-
-<style>
-  /* Target tables in this section */
-  #models-x-optimization + p + table {
-    display: block;
-    overflow-x: auto;
-    width: 100%;
-    font-size: 0.85rem;
-  }
-  
-  #models-x-optimization + p + table td,
-  #models-x-optimization + p + table th {
-    text-align: center;
-    white-space: nowrap;
-    padding: 0.5em;
-  }
-  
-  /* First two columns can wrap */
-  #models-x-optimization + p + table td:nth-child(1),
-  #models-x-optimization + p + table td:nth-child(2) {
-    white-space: normal;
-    min-width: 120px;
-  }
-  
-  #models-x-optimization + p + table td:nth-child(2) code {
-    font-size: 0.75rem;
-  }
-</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 | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
-
-**Note**: Wan2.2 TI2V 5B has some quality issues when performing I2V generation. We are working on fixing this issue.
-
-## Special requirements
-
-### StepVideo T2V
- The self-attention in text-encoder (step_llm) only supports CUDA capabilities sm_80 sm_86 and sm_90
-
-### 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
@@ -0,0 +1,35 @@
+@ECHO OFF
+
+pushd %~dp0
+
+REM Command file for Sphinx documentation
+
+if "%SPHINXBUILD%" == "" (
+	set SPHINXBUILD=sphinx-build
+)
+set SOURCEDIR=source
+set BUILDDIR=build
+
+%SPHINXBUILD% >NUL 2>NUL
+if errorlevel 9009 (
+	echo.
+	echo.The 'sphinx-build' command was not found. Make sure you have Sphinx
+	echo.installed, then set the SPHINXBUILD environment variable to point
+	echo.to the full path of the 'sphinx-build' executable. Alternatively you
+	echo.may add the Sphinx directory to PATH.
+	echo.
+	echo.If you don't have Sphinx installed, grab it from
+	echo.https://www.sphinx-doc.org/
+	exit /b 1
+)
+
+if "%1" == "" goto help
+
+%SPHINXBUILD% -M %1 %SOURCEDIR% %BUILDDIR% %SPHINXOPTS% %O%
+goto end
+
+:help
+%SPHINXBUILD% -M help %SOURCEDIR% %BUILDDIR% %SPHINXOPTS% %O%
+
+:end
+popd
@@ -0,0 +1,15 @@
+sphinx==7.4.7
+sphinx-argparse==0.5.2
+sphinx-autodoc2==0.5.0
+sphinx-book-theme==1.1.4
+sphinx-copybutton==0.5.2
+sphinx-design==0.6.1
+sphinx-togglebutton==0.3.2
+myst-parser==3.0.1
+msgspec
+commonmark # Required by sphinx-argparse when using :markdownhelp:
+
+# packages to install to build the documentation
+cachetools
+# -f https://download.pytorch.org/whl/cpu
+torch
@@ -0,0 +1,51 @@
+# Seed Parameter Behavior in vLLM
+
+## Overview
+
+The `seed` parameter in vLLM is used to control the random states for various random number generators. This parameter can affect the behavior of random operations in user code, especially when working with models in vLLM.
+
+## Default Behavior
+
+By default, the `seed` parameter is set to `None`. When the `seed` parameter is `None`, the global random states for `random`, `np.random`, and `torch.manual_seed` are not set. This means that the random operations will behave as expected, without any fixed random states.
+
+## Specifying a Seed
+
+If a specific seed value is provided, the global random states for `random`, `np.random`, and `torch.manual_seed` will be set accordingly. This can be useful for reproducibility, as it ensures that the random operations produce the same results across multiple runs.
+
+## Example Usage
+
+### Without Specifying a Seed
+
+```python
+import random
+from vllm import LLM
+
+# Initialize a vLLM model without specifying a seed
+model = LLM(model="Qwen/Qwen2.5-0.5B-Instruct")
+
+# Try generating random numbers
+print(random.randint(0, 100))  # Outputs different numbers across runs
+```
+
+### Specifying a Seed
+
+```python
+import random
+from vllm import LLM
+
+# Initialize a vLLM model with a specific seed
+model = LLM(model="Qwen/Qwen2.5-0.5B-Instruct", seed=42)
+
+# Try generating random numbers
+print(random.randint(0, 100))  # Outputs the same number across runs
+```
+
+## Important Notes
+
+- If the `seed` parameter is not specified, the behavior of global random states remains unaffected.
+- If a specific seed value is provided, the global random states for `random`, `np.random`, and `torch.manual_seed` will be set to that value.
+- This behavior can be useful for reproducibility but may lead to non-intuitive behavior if the user is not explicitly aware of it.
+
+## Conclusion
+
+Understanding the behavior of the `seed` parameter in vLLM is crucial for ensuring the expected behavior of random operations in your code. By default, the `seed` parameter is set to `None`, which means that the global random states are not affected. However, specifying a seed value can help achieve reproducibility in your experiments.
@@ -0,0 +1,8 @@
+.vertical-table-header th.head:not(.stub) {
+    writing-mode: sideways-lr;
+    white-space: nowrap;
+    max-width: 0;
+    p {
+       margin: 0;
+    }
+}
@@ -0,0 +1,39 @@
+<style>
+  .notification-bar {
+    width: 100vw;
+    display: flex;
+    justify-content: center;
+    align-items: center;
+    font-size: 16px;
+    padding: 0 6px 0 6px;
+  }
+  .notification-bar p {
+    margin: 0;
+  }
+  .notification-bar a {
+    font-weight: bold;
+    text-decoration: none;
+  }
+
+  /* Light mode styles (default) */
+  .notification-bar {
+    background-color: #fff3cd;
+    color: #856404;
+  }
+  .notification-bar a {
+    color: #d97706;
+  }
+
+  /* Dark mode styles */
+  html[data-theme=dark] .notification-bar {
+    background-color: #333;
+    color: #ddd;
+  }
+  html[data-theme=dark] .notification-bar a {
+    color: #ffa500; /* Brighter color for visibility */
+  }
+</style>
+
+<!-- <div class="notification-bar">
+  <p>You are viewing the latest developer preview docs. <a href="https://docs.vllm.ai/en/stable/">Click here</a> to view docs for the latest stable release.</p>
+</div> -->
@@ -0,0 +1,19 @@
+# Summary
+
+## Video Generator
+
+```{autodoc2-summary}
+    fastvideo.VideoGenerator
+```
+
+## Initialization Configuration
+
+```{autodoc2-summary}
+    fastvideo.configs.pipelines.PipelineConfig
+```
+
+## Sampling Configuration
+
+```{autodoc2-summary}
+    fastvideo.configs.sample.SamplingParam
+```
@@ -0,0 +1,22 @@
+# type: ignore
+# SPDX-License-Identifier: Apache-2.0
+from docutils import nodes
+from myst_parser.parsers.sphinx_ import MystParser
+from sphinx.ext.napoleon import docstring
+
+
+class NapoleonParser(MystParser):
+
+    def parse(self, input_string: str, document: nodes.document) -> None:
+        # Get the Sphinx configuration
+        config = document.settings.env.config
+
+        parsed_content = str(
+            docstring.GoogleDocstring(
+                str(docstring.NumpyDocstring(input_string, config)),
+                config,
+            ))
+        return super().parse(parsed_content, document)
+
+
+Parser = NapoleonParser
@@ -0,0 +1,275 @@
+# SPDX-License-Identifier: Apache-2.0
+
+# Configuration file for the Sphinx documentation builder.
+#
+# This file only contains a selection of the most common options. For a full
+# list see the documentation:
+# https://www.sphinx-doc.org/en/master/usage/configuration.html
+
+# -- Path setup --------------------------------------------------------------
+
+# If extensions (or modules to document with autodoc) are in another directory,
+# add these directories to sys.path here. If the directory is relative to the
+# documentation root, use os.path.abspath to make it absolute, like shown here.
+
+import datetime
+import logging
+import os
+import re
+import sys
+from pathlib import Path
+
+import requests
+
+logger = logging.getLogger(__name__)
+REPO_ROOT = Path(__file__).resolve().parent.parent.parent
+print(os.path.abspath(REPO_ROOT))
+sys.path.append(os.path.abspath(REPO_ROOT))
+
+# -- Project information -----------------------------------------------------
+
+project = 'FastVideo'
+copyright = f'{datetime.datetime.now().year}, FastVideo Team'
+author = 'the FastVideo Team'
+
+# -- General configuration ---------------------------------------------------
+
+# Add any Sphinx extension module names here, as strings. They can be
+# extensions coming with Sphinx (named 'sphinx.ext.*') or your custom
+# ones.
+extensions = [
+    "sphinx.ext.napoleon",
+    "sphinx.ext.linkcode",
+    "sphinx.ext.intersphinx",
+    "sphinx_copybutton",
+    "autodoc2",
+    "myst_parser",
+    "sphinxarg.ext",
+    "sphinx_design",
+    "sphinx_togglebutton",
+]
+myst_enable_extensions = [
+    "colon_fence",
+    "fieldlist",
+]
+autodoc2_packages = [
+    {
+        "path": "../../fastvideo",
+        "exclude_dirs": ["__pycache__", "third_party"],
+    },
+]
+autodoc2_output_dir = "api"
+autodoc2_render_plugin = "myst"
+autodoc2_hidden_objects = ["dunder", "private", "inherited"]
+autodoc2_docstring_parser_regexes = [
+    (".*", "docs.source.autodoc2_docstring_parser"),
+]
+autodoc2_sort_names = True
+autodoc2_index_template = None
+autodoc2_skip_module_regexes = [
+    "fastvideo.dataset",
+    "fastvideo.distill",
+    "fastvideo.data_preprocess",
+    "fastvideo.models",
+    "fastvideo.sample",
+    "fastvideo.utils",
+    "fastvideo.distill_adv",
+    "fastvideo.train",
+]
+
+# Add any paths that contain templates here, relative to this directory.
+templates_path = ['_templates']
+
+# List of patterns, relative to source directory, that match files and
+# directories to ignore when looking for source files.
+# This pattern also affects html_static_path and html_extra_path.
+exclude_patterns: list[str] = ["**/*.template.md", "**/*.inc.md"]
+
+# Exclude the prompt "$" when copying code
+copybutton_prompt_text = r"\$ "
+copybutton_prompt_is_regexp = True
+
+# -- Options for HTML output -------------------------------------------------
+
+# The theme to use for HTML and HTML Help pages.  See the documentation for
+# a list of builtin themes.
+#
+html_title = project
+html_theme = 'sphinx_book_theme'
+html_logo = '../../assets/logos/icon_simple.svg'
+html_theme_options = {
+    'path_to_docs': 'docs/source',
+    'repository_url': 'https://github.com/hao-ai-lab/FastVideo/',
+    'use_repository_button': True,
+    'use_edit_page_button': True,
+    # Prevents the full API being added to the left sidebar of every page.
+    # Reduces build time by 2.5x and reduces build size from ~225MB to ~95MB.
+    'collapse_navbar': True,
+    # Makes API visible in the right sidebar on API reference pages.
+    'show_toc_level': 3,
+}
+# Add any paths that contain custom static files (such as style sheets) here,
+# relative to this directory. They are copied after the builtin static files,
+# so a file named "default.css" will overwrite the builtin "default.css".
+html_static_path = ["_static"]
+html_js_files = ["custom.js"]
+html_css_files = ["custom.css"]
+
+myst_url_schemes = {
+    'http': None,
+    'https': None,
+    'mailto': None,
+    'ftp': None,
+    "gh-issue": {
+        "url":
+        "https://github.com/hao-ai-lab/FastVideo/issues/{{path}}#{{fragment}}",
+        "title": "Issue #{{path}}",
+        "classes": ["github"],
+    },
+    "gh-pr": {
+        "url":
+        "https://github.com/hao-ai-lab/FastVideo/pull/{{path}}#{{fragment}}",
+        "title": "Pull Request #{{path}}",
+        "classes": ["github"],
+    },
+    "gh-dir": {
+        "url": "https://github.com/hao-ai-lab/FastVideo/tree/main/{{path}}",
+        "title": "{{path}}",
+        "classes": ["github"],
+    },
+    "gh-file": {
+        "url": "https://github.com/hao-ai-lab/FastVideo/blob/main/{{path}}",
+        "title": "{{path}}",
+        "classes": ["github"],
+    },
+}
+
+# see https://docs.readthedocs.io/en/stable/reference/environment-variables.html # noqa
+READTHEDOCS_VERSION_TYPE = os.environ.get('READTHEDOCS_VERSION_TYPE')
+if READTHEDOCS_VERSION_TYPE == "tag":
+    # remove the warning banner if the version is a tagged release
+    header_file = os.path.join(os.path.dirname(__file__),
+                               "_templates/sections/header.html")
+    # The file might be removed already if the build is triggered multiple times
+    # (readthedocs build both HTML and PDF versions separately)
+    if os.path.exists(header_file):
+        os.remove(header_file)
+
+
+# Generate additional rst documentation here.
+def setup(app):
+    from docs.source.generate_examples import generate_examples
+    generate_examples()
+
+
+_cached_base: str = ""
+_cached_branch: str = ""
+
+
+def get_repo_base_and_branch(pr_number: str) -> tuple[str | None, str | None]:
+    global _cached_base, _cached_branch
+    if _cached_base and _cached_branch:
+        return _cached_base, _cached_branch
+
+    url = f"https://api.github.com/repos/hao-ai-lab/FastVideo/pulls/{pr_number}"
+    response = requests.get(url)
+    if response.status_code == 200:
+        data = response.json()
+        _cached_base = data['head']['repo']['full_name']
+        _cached_branch = data['head']['ref']
+        return _cached_base, _cached_branch
+    else:
+        logger.error("Failed to fetch PR details: %s", response)
+        return None, None
+
+
+def linkcode_resolve(domain, info):
+    if domain != 'py':
+        return None
+    if not info['module']:
+        return None
+
+    # Get path from module name
+    file = Path(f"{info['module'].replace('.', '/')}.py")
+    path = REPO_ROOT / file
+    if not path.exists():
+        path = REPO_ROOT / file.with_suffix("") / "__init__.py"
+    if not path.exists():
+        return None
+
+    # Get the line number of the object
+    with open(path) as f:
+        lines = f.readlines()
+    name = info['fullname'].split(".")[-1]
+    pattern = fr"^( {{4}})*((def|class) )?{name}\b.*"
+    for lineno, line in enumerate(lines, 1):
+        if not line or line.startswith("#"):
+            continue
+        if re.match(pattern, line):
+            break
+
+    # If the line number is not found, return None
+    if lineno == len(lines):
+        return None
+
+    # If the line number is found, create the URL
+    filename = path.relative_to(REPO_ROOT)
+    if "checkouts" in path.parts:
+        # a PR build on readthedocs
+        pr_number = REPO_ROOT.name
+        base, branch = get_repo_base_and_branch(pr_number)
+        if base and branch:
+            return f"https://github.com/{base}/blob/{branch}/{filename}#L{lineno}"
+    # Otherwise, link to the source file on the main branch
+    return f"https://github.com/hao-ai-lab/FastVideo/blob/main/{filename}#L{lineno}"
+
+
+# Mock out external dependencies here, otherwise the autodoc pages may be blank.
+autodoc_mock_imports = [
+    "blake3",
+    "compressed_tensors",
+    "cpuinfo",
+    "cv2",
+    "torch",
+    "huggingface_hub",
+    "torchvision",
+    "transformers",
+    "psutil",
+    "prometheus_client",
+    "sentencepiece",
+    "vllm._C",
+    "PIL",
+    "numpy",
+    'triton',
+    "tqdm",
+    "tensorizer",
+    "pynvml",
+    "outlines",
+    "xgrammar",
+    "librosa",
+    "soundfile",
+    "gguf",
+    "lark",
+    "decord",
+]
+
+for mock_target in autodoc_mock_imports:
+    if mock_target in sys.modules:
+        logger.info(
+            "Potentially problematic mock target (%s) found; "
+            "autodoc_mock_imports cannot mock modules that have already "
+            "been loaded into sys.modules when the sphinx build starts.",
+            mock_target)
+
+intersphinx_mapping = {
+    "python": ("https://docs.python.org/3", None),
+    "typing_extensions":
+    ("https://typing-extensions.readthedocs.io/en/latest", None),
+    "aiohttp": ("https://docs.aiohttp.org/en/stable", None),
+    "pillow": ("https://pillow.readthedocs.io/en/stable", None),
+    "numpy": ("https://numpy.org/doc/stable", None),
+    "torch": ("https://pytorch.org/docs/stable", None),
+    "psutil": ("https://psutil.readthedocs.io/en/stable", None),
+}
+
+navigation_with_keys = False
@@ -1,4 +1,4 @@
-
+(docker)=
 # 🐳 Using the FastVideo Docker Image

 If you prefer a containerized development environment or want to avoid managing dependencies manually, you can use our prebuilt Docker image:
@@ -3,3 +3,11 @@
 # 🧰 Developer Environment

 Accelerate your FastVideo development workflow by leveraging Docker images and cloud GPUs for efficient experimentation and reproducible environments.
+
+:::{toctree}
+:caption: Contents
+:maxdepth: 1
+
+docker
+runpod
+:::
@@ -1,3 +1,4 @@
+(runpod)=

 # 📦 Developing FastVideo on RunPod

@@ -9,7 +10,7 @@ Choose a GPU that supports CUDA 12.8

 Pick 1 or 2 L40S GPU(s)

-![RunPod CUDA selection](../../assets/images/runpod_cuda.png)
+![RunPod CUDA selection](../../_static/images/runpod_cuda.png)

 When creating your pod template, use this image:

@@ -23,11 +24,11 @@ Paste Container Start Command to support SSH ([RunPod Docs](https://docs.runpod.
 bash -c "apt update;DEBIAN_FRONTEND=noninteractive apt-get install openssh-server -y;mkdir -p ~/.ssh;cd $_;chmod 700 ~/.ssh;echo \"$PUBLIC_KEY\" >> authorized_keys;chmod 700 authorized_keys;service ssh start;sleep infinity"
 ```

-![RunPod template configuration](../../assets/images/runpod_template.png)
+![RunPod template configuration](../../_static/images/runpod_template.png)

 After deploying, the pod will take a few minutes to pull the image and start the SSH service.

-![RunPod ssh](../../assets/images/runpod_ssh.png)
+![RunPod ssh](../../_static/images/runpod_ssh.png)

 ## Working with the pod

--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
RandNMR73	44840ac49d	wan2.2 training doesn't crash	2025-09-20 05:08:58 +00:00
JerryZhou54	c99b1d4d97	Change Wan DiT to have 0 numerical diff with SF's Wan	2025-09-18 22:50:07 +00:00
SolitaryThinker	edfe4dd1bf	fix rotary embedding	2025-09-13 08:31:53 +00:00
RandNMR73	93ebd15a0d	text preprocessing ready	2025-09-10 11:12:48 +00:00
JerryZhou54	1110474065	checkpoint	2025-09-10 08:57:54 +00:00
JerryZhou54	80baffd540	Enable timestep warping & using SelfForcing scheduler	2025-09-09 23:30:55 +00:00
JerryZhou54	918180048e	Stop backprop through kv_cache	2025-09-09 10:02:03 +00:00
RandNMR73	b7dbd7cb9e	new branch	2025-09-09 10:02:00 +00:00
RandNMR73	71159b6416	inference works after changes added	2025-09-09 10:01:31 +00:00