fix cm

fix preprocessing before vae deocde
init lpips
2025-08-21 05:18:30 +00:00 · 2025-08-20 01:10:53 +00:00 · 2025-08-19 23:48:06 +00:00 · 2025-08-19 16:40:42 -07:00 · 2025-08-19 15:40:29 -07:00 · 2025-08-19 07:19:40 +00:00
331 changed files with 3933 additions and 15846 deletions
@@ -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,11 +117,11 @@ steps:
                  queue: "default"
            - path:
                - "fastvideo/**"
-                - "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"
+                - "csrc/attn/vsa/**"
+                - "csrc/attn/tk/**"
+                - "csrc/attn/setup_vsa.py"
+                - "csrc/attn/config_vsa.py"
+                - "csrc/attn/vsa.cpp"
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
@@ -145,10 +133,10 @@ steps:
                  queue: "default"
            - path:
                - "fastvideo/**"
-                - "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"
+                - "csrc/attn/st_attn/**"
+                - "csrc/attn/setup_sta.py"
+                - "csrc/attn/config_sta.py"
+                - "csrc/attn/st_attn.cpp"
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
@@ -159,10 +147,10 @@ steps:
                agents:
                  queue: "default"
            - path:
-                - "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"
+                - "csrc/attn/st_attn/**"
+                - "csrc/attn/setup_sta.py"
+                - "csrc/attn/config_sta.py"
+                - "csrc/attn/st_attn.cpp"
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
@@ -173,12 +161,12 @@ steps:
                agents:
                  queue: "default"
            - path:
-                - "csrc/attn/video_sparse_attn/**"
-                - "csrc/attn/video_sparse_attn/tk/**"
+                - "csrc/attn/vsa/**"
+                - "csrc/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"
+                - "csrc/attn/setup_vsa.py"
+                - "csrc/attn/config_vsa.py"
+                - "csrc/attn/vsa.cpp"
                - "pyproject.toml"
                - "docker/Dockerfile.python3.12"
              config:
@@ -188,37 +176,3 @@ steps:
                  - 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"
-              config:
-                command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Inference Tests VMoBA"
-                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"
@@ -109,23 +109,6 @@ case "$TEST_TYPE" in
        log "Running distillation DMD tests..."
        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"
-        ;;
-    "precision_vmoba")
-        log "Running V-MoBA precision tests..."
-        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_precision_tests_vmoba"
-        ;;
-    "unit_test")
-        log "Running unit tests..."
-        MODAL_COMMAND="$MODAL_ENV python3 -m modal run $MODAL_TEST_FILE::run_unit_test"
-        ;;
    *)
        log "Error: Unknown test type: $TEST_TYPE"
        exit 1
@@ -18,12 +18,6 @@ on:
        required: false
        default: false
        type: boolean
-      python_3_12_cuda_12_9:
-        description: 'Build Python 3.12 image Cuda 12.9'
-        required: false
-        default: false
-        type: boolean
-

 permissions:
  contents: read
@@ -55,13 +49,4 @@ jobs:
      python_version: '3.12'
      dockerfile_path: docker/Dockerfile.python3.12
      tag_suffix: py3.12
-    secrets: inherit
-
-  build-python-3-12-cuda-12-9:
-    if: ${{ github.event.inputs.python_3_12_cuda_12_9 == 'true' }}
-    uses: ./.github/workflows/build-image-template.yml
-    with:
-      python_version: '3.12'
-      dockerfile_path: docker/Dockerfile.python3.12.cuda12.9.1
-      tag_suffix: py3.12-cuda12.9.1
    secrets: inherit
@@ -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,21 +102,18 @@ 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/**'
-              - 'csrc/attn/sliding_tile_attn/tk/**'
-              - 'csrc/attn/sliding_tile_attn/setup.py'
-              - 'csrc/attn/sliding_tile_attn/config_sta.py'
-              - 'csrc/attn/sliding_tile_attn/st_attn.cpp'
+              - 'csrc/attn/st_attn/**'
+              - 'csrc/attn/setup_sta.py'
+              - 'csrc/attn/config_sta.py'
+              - 'csrc/attn/st_attn.cpp'
            vsa-kernel-paths: &vsa-kernel-paths
-              - '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'
+              - 'csrc/attn/vsa/**'
+              - 'csrc/attn/tk/**'
+              - 'csrc/attn/setup_vsa.py'
+              - 'csrc/attn/config_vsa.py'
+              - 'csrc/attn/vsa.cpp'
            vsa-paths: &vsa-paths
              - 'fastvideo/**'
              - *common-paths
@@ -158,9 +154,6 @@ jobs:
            precision-test-VSA:
              - *common-paths
              - *vsa-kernel-paths
-            unit-test:
-              - 'fastvideo/**'
-              - *common-paths

  encoder-test:
    needs: change-filter
@@ -241,7 +234,7 @@ jobs:
    secrets:
      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-
+  
  training-test:
    needs: change-filter
    if: >-
@@ -339,42 +332,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"
+      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
@@ -398,4 +372,4 @@ jobs:
          JOB_IDS: '["encoder-test", "vae-test", "transformer-test", "ssim-test-py3.10", "ssim-test-py3.11", "ssim-test-py3.12", "training-test", "training-test-VSA", "inference-test-STA", "precision-test-STA", "precision-test-VSA"]'
          RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
          GITHUB_RUN_ID: ${{ github.run_id }}
-        run: python .github/scripts/runpod_cleanup.py
+        run: python .github/scripts/runpod_cleanup.py
@@ -5,7 +5,7 @@ on:
    branches:
      - main
    paths:
-      - "csrc/attn/sliding_tile_attn/setup.py"
+      - "csrc/attn/setup_sta.py"
  workflow_dispatch:

 jobs:
@@ -23,13 +23,13 @@ jobs:
      - name: Check if version changed
        id: check-version
        run: |
-          cd csrc/attn/sliding_tile_attn
+          cd csrc/attn
          # Get current commit's version
-          NEW_VERSION=$(grep -oP 'VERSION\s*=\s*"\K[^"]+' setup.py)
+          NEW_VERSION=$(grep -oP 'VERSION\s*=\s*"\K[^"]+' setup_sta.py)
          echo "New version: $NEW_VERSION"
          
          # Get previous version from git history
-          OLD_VERSION=$(git show HEAD~1:./setup.py | grep -oP 'VERSION\s*=\s*"\K[^"]+' || echo "0.0.0")
+          OLD_VERSION=$(git show HEAD~1:./setup_sta.py | grep -oP 'VERSION\s*=\s*"\K[^"]+' || echo "0.0.0")
          echo "Old version: $OLD_VERSION"
          
          if [ "$NEW_VERSION" != "$OLD_VERSION" ]; then
@@ -144,13 +144,13 @@ jobs:
          pip install setuptools
          pip install ninja packaging wheel

-          cd csrc/attn/sliding_tile_attn  # Move into the correct folder
+          cd csrc/attn  # Move into the correct folder
          git submodule update --init --recursive # Ensure ThunderKittens submodule is initialized
-          python setup.py bdist_wheel --dist-dir=dist
+          python setup_sta.py bdist_wheel --dist-dir=dist

      - name: Rename wheel file
        run: |
-          cd csrc/attn/sliding_tile_attn
+          cd csrc/attn

          CUDA_SHORT_VERSION=$(echo ${{ matrix.cuda-version }} | cut -d. -f1,2 | sed 's/\.//g')
          TORCH_SHORT_VERSION=$(echo ${{ matrix.torch-version }} | cut -d. -f1,2)
@@ -165,7 +165,7 @@ jobs:
        uses: actions/upload-artifact@v4
        with:
          name: ${{ env.wheel_name }}
-          path: csrc/attn/sliding_tile_attn/dist/*.whl
+          path: csrc/attn/dist/*.whl
          retention-days: 90

  publish_package:
@@ -239,11 +239,11 @@ jobs:
          pip install setuptools
          pip install ninja packaging wheel

-          cd csrc/attn/sliding_tile_attn  # Move into the correct folder
+          cd csrc/attn  # Move into the correct folder
          git submodule update --init --recursive # Ensure ThunderKittens submodule is initialized
-          python setup.py sdist --dist-dir=dist
+          python setup_sta.py sdist --dist-dir=dist

      - name: Publish release distributions to PyPI
        uses: pypa/gh-action-pypi-publish@release/v1
        with:
-          packages-dir: csrc/attn/sliding_tile_attn/dist/
+          packages-dir: csrc/attn/dist/
@@ -5,7 +5,7 @@ on:
    branches:
      - main
    paths:
-      - "csrc/attn/video_sparse_attn/setup.py"
+      - "csrc/attn/setup_vsa.py"
  workflow_dispatch:

 jobs:
@@ -23,13 +23,13 @@ jobs:
      - name: Check if version changed
        id: check-version
        run: |
-          cd csrc/attn/video_sparse_attn
+          cd csrc/attn
          # Get current commit's version
-          NEW_VERSION=$(grep -oP 'VERSION\s*=\s*"\K[^"]+' setup.py)
+          NEW_VERSION=$(grep -oP 'VERSION\s*=\s*"\K[^"]+' setup_vsa.py)
          echo "New version: $NEW_VERSION"
          
          # Get previous version from git history
-          OLD_VERSION=$(git show HEAD~1:./setup.py | grep -oP 'VERSION\s*=\s*"\K[^"]+' || echo "0.0.0")
+          OLD_VERSION=$(git show HEAD~1:./setup_vsa.py | grep -oP 'VERSION\s*=\s*"\K[^"]+' || echo "0.0.0")
          echo "Old version: $OLD_VERSION"
          
          if [ "$NEW_VERSION" != "$OLD_VERSION" ]; then
@@ -152,13 +152,13 @@ jobs:
          pip install setuptools
          pip install ninja packaging wheel

-          cd csrc/attn/video_sparse_attn  # Move into the correct folder
+          cd csrc/attn  # Move into the correct folder
          git submodule update --init --recursive # Ensure ThunderKittens submodule is initialized
-          python setup.py bdist_wheel --dist-dir=dist
+          python setup_vsa.py bdist_wheel --dist-dir=dist

      - name: Rename wheel file
        run: |
-          cd csrc/attn/video_sparse_attn
+          cd csrc/attn

          CUDA_SHORT_VERSION=$(echo ${{ matrix.torch-cuda.cuda-version }} | cut -d. -f1,2 | sed 's/\.//g')
          TORCH_SHORT_VERSION=$(echo ${{ matrix.torch-cuda.torch-version }} | cut -d. -f1,2)
@@ -173,7 +173,7 @@ jobs:
        uses: actions/upload-artifact@v4
        with:
          name: ${{ env.wheel_name }}
-          path: csrc/attn/video_sparse_attn/dist/*.whl
+          path: csrc/attn/dist/*.whl
          retention-days: 90

  publish_package:
@@ -247,11 +247,11 @@ jobs:
          pip install setuptools
          pip install ninja packaging wheel

-          cd csrc/attn/video_sparse_attn  # Move into the correct folder
+          cd csrc/attn  # Move into the correct folder
          git submodule update --init --recursive # Ensure ThunderKittens submodule is initialized
-          python setup.py sdist --dist-dir=dist
+          python setup_vsa.py sdist --dist-dir=dist

      - name: Publish release distributions to PyPI
        uses: pypa/gh-action-pypi-publish@release/v1
        with:
-          packages-dir: csrc/attn/video_sparse_attn/dist/
+          packages-dir: csrc/attn/dist/
@@ -64,6 +64,3 @@ docs/source/distillation/examples/
 !docs/source/_static/images/**/*.png
 !comfyui/assets/**/*.png
 !comfyui/assets/**/*.gif
-
-dmd_t2v_output/
-preprocess_output_text/
@@ -1,7 +1,3 @@
-[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
+[submodule "csrc/attn/tk"]
+	path = csrc/attn/tk
 	url = https://github.com/HazyResearch/ThunderKittens.git
@@ -12,6 +12,9 @@ exclude: |
    scripts/.*|
    fastvideo/data_preprocess/.*|
    fastvideo/dataset/.*|
+    fastvideo/distill/.*|
+    fastvideo/distill\.py|
+    fastvideo/distill_adv\.py|
    fastvideo/models/.*|
    fastvideo/sample/.*|
    fastvideo/train\.py|
@@ -19,7 +22,6 @@ exclude: |
    examples/.*|
    .github/workflows/fastvideo-publish.yml|
    .github/workflows/sta-publish.yml|
-    .github/workflows/vsa-publish.yml|
    .github/workflows/build-image-template.yml|
    docs/source/inference/support_matrix.md
  )
@@ -41,10 +43,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:
@@ -1,5 +1,5 @@
 <div align="center">
-<img src=assets/logos/logo.svg width="30%"/>
+<img src=assets/logo.png width="30%"/>
 </div>

 **FastVideo is a unified post-training and inference framework for accelerated video generation.**
@@ -7,7 +7,7 @@
 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/q46BbX6" target="_blank"> <b> WeChat </b> </a> |
+    | 🕹️ <a href="https://fastwan.fastvideo.org/"<b>Online Demo</b></a> | <a href="https://hao-ai-lab.github.io/FastVideo"><b>Documentation</b></a> | <a href="https://hao-ai-lab.github.io/FastVideo/inference/inference_quick_start.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-38u6p1jqe-yDI1QJOCEnbtkLoaI5bjZQ" target="_blank"> <b>Slack</b> </a> |  🟣💬 <a href="https://ibb.co/qqPzbrw" target="_blank"> <b> WeChat </b> </a> |
 </p>

 <div align="center">
@@ -155,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,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,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,2 +1,2 @@
 recursive-include tk *
-include config_sta.py
+include config.py
@@ -25,9 +25,9 @@ sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave
 sudo apt update
 sudo apt install clang-11
 ```
-(If you use CUDA12.8)
+(If you use CUDA12.4)
 ```bash
-export CUDA_HOME=/usr/local/cuda-12.8
+export CUDA_HOME=/usr/local/cuda-12.4
 export PATH=${CUDA_HOME}/bin:${PATH} 
 export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
 ```
@@ -0,0 +1,4 @@
+    off_hz = tl.program_id(2)
+    b = off_hz // H
+    h = off_hz % H
+    meta_base = ((b * H + h) * q_tiles + q_blk)
@@ -1,7 +1,7 @@
 import os
 import subprocess

-from config_sta import kernels, sources, target
+from csrc.attn.config_sta import kernels, sources, target
 from setuptools import find_packages, setup
 from torch.utils.cpp_extension import BuildExtension, CUDAExtension

@@ -9,7 +9,7 @@ target = target.lower()

 # Package metadata
 PACKAGE_NAME = "st_attn"
-VERSION = "0.0.6"
+VERSION = "0.0.4"
 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"
@@ -9,10 +9,10 @@ target = target.lower()

 # Package metadata
 PACKAGE_NAME = "vsa"
-VERSION = "0.0.3"
+VERSION = "0.0.1"
 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"
+URL = "https://github.com/hao-ai-lab/FastVideo/tree/main/csrc/attn"

 # Set environment variables
 tk_root = os.getenv('THUNDERKITTENS_ROOT', os.path.abspath(os.path.join(os.getcwd(), 'tk/')))
@@ -1,87 +0,0 @@
-
-# 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.
@@ -13,9 +13,9 @@ 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_()
+    q_ = Q.clone().requires_grad_()
+    k_ = K.clone().requires_grad_()
+    v_ = V.clone().requires_grad_()

    QK = torch.matmul(q_, k_.transpose(-2, -1))
    QK /= (q_.size(-1) ** 0.5)
@@ -35,9 +35,9 @@ def pytorch_test(Q, K, V, block_sparse_mask, dO):


 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 = Q.clone().requires_grad_()
+    K = K.clone().requires_grad_()
+    V = V.clone().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)
@@ -60,9 +60,11 @@ def get_non_pad_index(

    return index_pad[index_mask]

-def generate_tensor(shape, dtype, device):
+def generate_tensor(shape, mean, std, dtype, device):
    tensor = torch.randn(shape, dtype=dtype, device=device)
-    return tensor
+    magnitude = torch.norm(tensor, dim=-1, keepdim=True)
+    scaled_tensor = tensor * (torch.randn(magnitude.shape, dtype=dtype, device=device) * std + mean) / magnitude
+    return scaled_tensor.contiguous()

 def 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)
@@ -73,7 +75,7 @@ def vsa_pad(x, non_pad_index, num_blocks, block_size):
    padded_x[:, :, non_pad_index, :] = x
    return padded_x

-def check_correctness(h, d, num_blocks, k,  num_iterations=20, error_mode='all'):
+def check_correctness(h, d, num_blocks, k, mean, std, 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},
@@ -89,10 +91,10 @@ def check_correctness(h, d, num_blocks, k,  num_iterations=20, error_mode='all')
    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)
+        Q = generate_tensor((1, h, S, d), mean, std, torch.bfloat16, device)
+        K = generate_tensor((1, h, S, d), mean, std, torch.bfloat16, device)
+        V = generate_tensor((1, h, S, d), mean, std, torch.bfloat16, device)
+        dO = generate_tensor((1, h, S, d), mean, std, torch.bfloat16, device)

        # dO_padded = torch.zeros_like(dO_padded)
        # dO_padded[:, :, non_pad_index, :] = dO
@@ -105,8 +107,7 @@ def check_correctness(h, d, num_blocks, k,  num_iterations=20, error_mode='all')
                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())
+                results[name]['max_diff'] = max(results[name]['max_diff'], torch.max(abs_diff).item())
        if torch.cuda.is_available():
            torch.cuda.empty_cache()

@@ -118,27 +119,27 @@ def check_correctness(h, d, num_blocks, k,  num_iterations=20, error_mode='all')

    return results

-def generate_error_graphs(h, d, error_mode='all'):
+def generate_error_graphs(h, d, mean, std, 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(f"\nError Analysis for h={h}, d={d}, mean={mean}, std={std}, 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}")
+          f"{'gQ Avg':<12} {'gQ Max':<12} "
+          f"{'gK Avg':<12} {'gK Max':<12} "
+          f"{'gV Avg':<12} {'gV Max':<12} "
+          f"{'gO Avg':<12} {'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)
+        results = check_correctness(h, d, num_blocks, k, mean, std, 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} "
@@ -149,8 +150,10 @@ def generate_error_graphs(h, d, error_mode='all'):

 if __name__ == "__main__":
    h, d = 16, 128
+    mean = 0.0
+    std = 1
    print("Block Sparse Attention with Variable Block Sizes Analysis")
    print("=" * 60)
    for mode in ['backward']:
-        generate_error_graphs(h, d, error_mode=mode)
+        generate_error_graphs(h, d, mean, std, error_mode=mode)
    print("\nAnalysis completed for all modes.")
@@ -1,2 +0,0 @@
-recursive-include tk *
-include config_vsa.py
@@ -1,61 +0,0 @@
-
-
-# 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.
@@ -1,32 +0,0 @@
-# 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
-```
-
@@ -1,26 +0,0 @@
-# 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=[
-        "flash-attn >= 2.7.1",
-    ]
-)
@@ -1,97 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-
-import torch
-import pytest
-import random
-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"):
-    """
-    Generates random data for testing the variable-length attention function.
-    """
-    torch.manual_seed(42)
-    random.seed(42)
-    torch.cuda.manual_seed_all(42)
-
-    # Generate sequence lengths for each item in the batch
-    if batch_size > 1:
-        # Ensure sequence lengths are reasonably distributed
-        avg_seqlen = total_seqlen // batch_size
-        seqlens = [random.randint(avg_seqlen // 2, avg_seqlen + avg_seqlen // 2) for _ in range(batch_size - 1)]
-        remaining_len = total_seqlen - sum(seqlens)
-        if remaining_len > 0:
-            seqlens.append(remaining_len)
-        else: # Adjust if sum exceeds total_seqlen
-            seqlens.append(avg_seqlen)
-            current_sum = sum(seqlens)
-            seqlens[-1] -= (current_sum - total_seqlen)
-        # Ensure all lengths are positive
-        seqlens = [max(1, s) for s in seqlens]
-        # Final adjustment to match total_seqlen
-        seqlens[-1] += total_seqlen - sum(seqlens)
-
-    else:
-        seqlens = [total_seqlen]
-
-    cu_seqlens = torch.tensor([0] + list(torch.cumsum(torch.tensor(seqlens), 0)), device=device, dtype=torch.int32)
-    max_seqlen = max(seqlens) if seqlens else 0
-
-    q = torch.randn((total_seqlen, num_heads, head_dim), dtype=dtype, device=device, requires_grad=False)
-    k = torch.randn((total_seqlen, num_heads, head_dim), dtype=dtype, device=device, requires_grad=False)
-    v = torch.randn((total_seqlen, num_heads, head_dim), dtype=dtype, device=device, requires_grad=False)
-
-    return q, k, v, cu_seqlens, max_seqlen
-
-
-@pytest.mark.parametrize("batch_size", [1, 2])
-@pytest.mark.parametrize("total_seqlen", [512, 1024])
-@pytest.mark.parametrize("num_heads", [8])
-@pytest.mark.parametrize("head_dim", [64])
-@pytest.mark.parametrize("moba_chunk_size", [64])
-@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.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
-):
-    """
-    Tests the forward pass of moba_attn_varlen for basic correctness.
-    It checks output shape, dtype, and for the presence of NaNs/Infs.
-    """
-    if dtype == torch.float32:
-        pytest.skip("float32 is not supported in flash attention")
-
-    q, k, v, cu_seqlens, max_seqlen = generate_test_data(
-        batch_size, total_seqlen, num_heads, head_dim, dtype
-    )
-
-    # Ensure chunk size is not larger than the smallest sequence length
-    min_seqlen = (cu_seqlens[1:] - cu_seqlens[:-1]).min().item()
-    if moba_chunk_size > min_seqlen:
-        pytest.skip("moba_chunk_size is larger than the minimum sequence length in the batch")
-
-    try:
-        output = moba_attn_varlen(
-            q=q,
-            k=k,
-            v=v,
-            cu_seqlens=cu_seqlens,
-            max_seqlen=max_seqlen,
-            moba_chunk_size=moba_chunk_size,
-            moba_topk=moba_topk,
-            select_mode=select_mode,
-            threshold_type=threshold_type,
-            simsum_threshold=0.5, # A reasonable default for threshold mode
-        )
-    except Exception as e:
-        pytest.fail(f"moba_attn_varlen forward pass failed with exception: {e}")
-
-    # 1. Check output shape
-    assert output.shape == q.shape, f"Expected output shape {q.shape}, but got {output.shape}"
-
-    # 2. Check output dtype
-    assert output.dtype == q.dtype, f"Expected output dtype {q.dtype}, but got {output.dtype}"
-
-    # 3. Check for NaNs or Infs in the output
-    assert torch.all(torch.isfinite(output)), "Output contains NaN or Inf values"
@@ -1,2 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-from .vmoba import moba_attn_varlen, process_moba_input, process_moba_output
@@ -1,868 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# Adapt from https://github.com/KwaiVGI/VMoBA/blob/main/src/vmoba.py
-
-import random
-import time
-import os
-import torch
-from typing import Tuple
-try:
-    from flash_attn import flash_attn_varlen_func  # Use the new flash attention function
-    from flash_attn.flash_attn_interface import _flash_attn_varlen_forward, _flash_attn_varlen_backward
-except ImportError:
-    def _unsupported(*args, **kwargs):
-        raise ImportError("flash-attn is not installed. Please install it, e.g., `pip install flash-attn`.")
-    _flash_attn_varlen_forward = _unsupported
-    _flash_attn_varlen_backward = _unsupported
-    flash_attn_varlen_func = _unsupported
-    
-from functools import lru_cache
-from einops import rearrange
-
-
-@lru_cache(maxsize=16)
-def calc_chunks(cu_seqlen, moba_chunk_size):
-    """
-    Calculate chunk boundaries.
-
-    For vision tasks we include all chunks (even the last one which might be shorter)
-    so that every chunk can be selected.
-    """
-    batch_sizes = cu_seqlen[1:] - cu_seqlen[:-1]
-    batch_num_chunk = (batch_sizes + (moba_chunk_size - 1)) // moba_chunk_size
-    cu_num_chunk = torch.ones(
-        batch_num_chunk.numel() + 1,
-        device=cu_seqlen.device,
-        dtype=batch_num_chunk.dtype,
-    )
-    cu_num_chunk[1:] = batch_num_chunk.cumsum(dim=0)
-    num_chunk = cu_num_chunk[-1]
-    chunk_sizes = torch.full(
-        (num_chunk + 1,), moba_chunk_size, dtype=torch.int32, device=cu_seqlen.device
-    )
-    chunk_sizes[0] = 0
-    batch_last_chunk_size = batch_sizes - (batch_num_chunk - 1) * moba_chunk_size
-    chunk_sizes[cu_num_chunk[1:]] = batch_last_chunk_size
-    cu_chunk = chunk_sizes.cumsum(dim=-1, dtype=torch.int32)
-    chunk_to_batch = torch.zeros(
-        (num_chunk,), dtype=torch.int32, device=cu_seqlen.device
-    )
-    chunk_to_batch[cu_num_chunk[1:-1]] = 1
-    chunk_to_batch = chunk_to_batch.cumsum(dim=0, dtype=torch.int32)
-
-    # Do not filter out any chunk
-    filtered_chunk_indices = torch.arange(
-        num_chunk, device=cu_seqlen.device, dtype=torch.int32
-    )
-    num_filtered_chunk = num_chunk
-
-    return cu_chunk, filtered_chunk_indices, num_filtered_chunk, chunk_to_batch
-
-
-# --- Threshold Selection Helper Functions ---
-
-def _select_threshold_query_head(
-    gate: torch.Tensor,
-    valid_gate_mask: torch.Tensor,
-    gate_self_chunk_mask: torch.Tensor,
-    simsum_threshold: float
-) -> torch.Tensor:
-    """
-    Selects chunks for each <query, head> pair based on threshold.
-    Normalization and sorting happen along the chunk dimension (dim=0).
-    """
-    C, H, S = gate.shape
-    eps = 1e-6
-
-    # LSE‐style normalization per <head, query> (across chunks)
-    gate_masked = torch.where(valid_gate_mask, gate, -torch.inf) # Use -inf for max
-    gate_min_val = torch.where(valid_gate_mask, gate, torch.inf) # Use +inf for min
-
-    row_min = gate_min_val.amin(dim=0)        # (H, S)
-    row_max = gate_masked.amax(dim=0)         # (H, S)
-    denom = row_max - row_min
-    denom = torch.where(denom <= eps, torch.ones_like(denom), denom) # avoid divide‑by‑zero
-
-    gate_norm = (gate - row_min.unsqueeze(0)) / denom.unsqueeze(0)
-    gate_norm = torch.where(valid_gate_mask, gate_norm, 0.0) # (C, H, S)
-
-    # 1) pull out the self‐chunk’s normalized weight for each <head,seq>
-    self_norm = (gate_norm * gate_self_chunk_mask).sum(dim=0)  # (H, S)
-
-    # 2) compute how much more normalized weight we need beyond self
-    total_norm_sum = gate_norm.sum(dim=0) # (H, S)
-    remain_ratio = simsum_threshold - self_norm / (total_norm_sum + eps)  # (H, S)
-    remain_ratio = torch.clamp(remain_ratio, min=0.0) # if already ≥ thresh, no extra needed
-
-    # 3) zero out the self‐chunk in a copy, so we only sort “others”
-    others_norm = gate_norm.clone()
-    others_norm[gate_self_chunk_mask] = 0.0
-
-    # 4) sort the other chunks by descending norm, per <head,seq>
-    sorted_norm, sorted_idx = torch.sort(others_norm, descending=True, dim=0) # (C, H, S)
-
-    # 5) cumulative‑sum the sorted norms per <head,seq>
-    cumsum_others = sorted_norm.cumsum(dim=0) # (C, H, S)
-
-    # 6) for each <head,seq>, find the smallest k where cumsum_ratio ≥ remain_ratio
-    ratio = cumsum_others / (total_norm_sum.unsqueeze(0) + eps) # (C, H, S)
-    cond = ratio >= remain_ratio.unsqueeze(0) # (C, H, S) boolean mask
-    any_cond = cond.any(dim=0) # (H, S)
-    # Find the index of the first True value along dim 0. If none, use C-1.
-    cutoff = torch.where(any_cond, cond.float().argmax(dim=0), torch.full_like(any_cond, fill_value=C - 1)) # (H, S)
-
-    # 7) build a mask in sorted order up to that cutoff
-    idx_range = torch.arange(C, device=gate.device).view(-1, 1, 1) # (C, 1, 1)
-    sorted_mask = idx_range <= cutoff.unsqueeze(0) # (C, H, S)
-
-    # 8) scatter it back to original chunk order
-    others_mask = torch.zeros_like(gate, dtype=torch.bool)
-    others_mask.scatter_(0, sorted_idx, sorted_mask)
-
-    # 9) finally, include every self‐chunk plus all selected others
-    final_gate_mask = valid_gate_mask & (others_mask | gate_self_chunk_mask)
-
-    return final_gate_mask
-
-
-def _select_threshold_block(
-    gate: torch.Tensor,
-    valid_gate_mask: torch.Tensor,
-    gate_self_chunk_mask: torch.Tensor,
-    simsum_threshold: float
-) -> torch.Tensor:
-    """
-    Selects <query, head> pairs for each block based on threshold.
-    Normalization and sorting happen across the head and sequence dimensions (dim=1, 2).
-    """
-    C, H, S = gate.shape
-    HS = H * S
-    eps = 1e-6
-
-    # LSE‐style normalization per block (across heads and queries)
-    gate_masked = torch.where(valid_gate_mask, gate, -torch.inf) # Use -inf for max
-    gate_min_val = torch.where(valid_gate_mask, gate, torch.inf) # Use +inf for min
-
-    block_max = gate_masked.amax(dim=(1, 2), keepdim=True) # (C, 1, 1)
-    block_min = gate_min_val.amin(dim=(1, 2), keepdim=True) # (C, 1, 1)
-    block_denom = block_max - block_min
-    block_denom = torch.where(block_denom <= eps, torch.ones_like(block_denom), block_denom) # (C, 1, 1)
-
-    gate_norm = (gate - block_min) / block_denom # (C, H, S)
-    gate_norm = torch.where(valid_gate_mask, gate_norm, 0.0) # (C, H, S)
-
-    # 1) identify normalized weights of entries that *are* self-chunks (from query perspective)
-    self_norm_entries = gate_norm * gate_self_chunk_mask # (C, H, S)
-    # Sum these weights *per block*
-    self_norm_sum_per_block = self_norm_entries.sum(dim=(1, 2)) # (C,)
-
-    # 2) compute how much more normalized weight each block needs beyond its self-chunk contributions
-    total_norm_sum_per_block = gate_norm.sum(dim=(1, 2)) # (C,)
-    remain_ratio = simsum_threshold - self_norm_sum_per_block / (total_norm_sum_per_block + eps) # (C,)
-    remain_ratio = torch.clamp(remain_ratio, min=0.0) # (C,)
-
-    # 3) zero out the self‐chunk entries in a copy, so we only sort “others”
-    others_norm = gate_norm.clone()
-    others_norm[gate_self_chunk_mask] = 0.0 # Zero out self entries
-
-    # 4) sort the other <head, seq> pairs by descending norm, per block
-    others_flat = others_norm.contiguous().view(C, HS) # (C, H*S)
-    sorted_others_flat, sorted_indices_flat = torch.sort(others_flat, dim=1, descending=True) # (C, H*S)
-
-    # 5) cumulative‑sum the sorted norms per block
-    cumsum_others_flat = sorted_others_flat.cumsum(dim=1) # (C, H*S)
-
-    # 6) for each block, find the smallest k where cumsum_ratio ≥ remain_ratio
-    ratio_flat = cumsum_others_flat / (total_norm_sum_per_block.unsqueeze(1) + eps) # (C, H*S)
-    cond_flat = ratio_flat >= remain_ratio.unsqueeze(1) # (C, H*S) boolean mask
-    any_cond = cond_flat.any(dim=1) # (C,)
-    # Find the index of the first True value along dim 1. If none, use HS-1.
-    cutoff_flat = torch.where(any_cond, cond_flat.float().argmax(dim=1), torch.full_like(any_cond, fill_value=HS - 1)) # (C,)
-
-    # 7) build a mask in sorted order up to that cutoff per block
-    idx_range_flat = torch.arange(HS, device=gate.device).unsqueeze(0) # (1, H*S)
-    sorted_mask_flat = idx_range_flat <= cutoff_flat.unsqueeze(1) # (C, H*S)
-
-    # 8) scatter it back to original <head, seq> order per block
-    others_mask_flat = torch.zeros_like(others_flat, dtype=torch.bool) # (C, H*S)
-    others_mask_flat.scatter_(1, sorted_indices_flat, sorted_mask_flat)
-    others_mask = others_mask_flat.view(C, H, S) # (C, H, S)
-
-    # 9) finally, include every self‐chunk entry plus all selected others
-    final_gate_mask = valid_gate_mask & (others_mask | gate_self_chunk_mask)
-
-    return final_gate_mask
-
-
-def _select_threshold_overall(
-    gate: torch.Tensor,
-    valid_gate_mask: torch.Tensor,
-    gate_self_chunk_mask: torch.Tensor,
-    simsum_threshold: float
-) -> torch.Tensor:
-    """
-    Selects <chunk, query, head> triplets globally based on threshold.
-    Normalization and sorting happen across all valid entries.
-    """
-    C, H, S = gate.shape
-    CHS = C * H * S
-    eps = 1e-6
-
-    # LSE‐style normalization globally across all valid entries
-    gate_masked = torch.where(valid_gate_mask, gate, -torch.inf) # Use -inf for max
-    gate_min_val = torch.where(valid_gate_mask, gate, torch.inf) # Use +inf for min
-
-    overall_max = gate_masked.max() # scalar
-    overall_min = gate_min_val.min() # scalar
-    overall_denom = overall_max - overall_min
-    overall_denom = torch.where(overall_denom <= eps, torch.tensor(1.0, device=gate.device, dtype=gate.dtype), overall_denom)
-
-    gate_norm = (gate - overall_min) / overall_denom # (C, H, S)
-    gate_norm = torch.where(valid_gate_mask, gate_norm, 0.0) # (C, H, S)
-
-    # 1) identify normalized weights of entries that *are* self-chunks
-    self_norm_entries = gate_norm * gate_self_chunk_mask # (C, H, S)
-    # Sum these weights globally
-    self_norm_sum_overall = self_norm_entries.sum() # scalar
-
-    # 2) compute how much more normalized weight is needed globally beyond self-chunk contributions
-    total_norm_sum_overall = gate_norm.sum() # scalar
-    remain_ratio = simsum_threshold - self_norm_sum_overall / (total_norm_sum_overall + eps) # scalar
-    remain_ratio = torch.clamp(remain_ratio, min=0.0) # scalar
-
-    # 3) zero out the self‐chunk entries in a copy, so we only sort “others”
-    others_norm = gate_norm.clone()
-    others_norm[gate_self_chunk_mask] = 0.0 # Zero out self entries
-
-    # 4) sort all other entries by descending norm, globally
-    others_flat = others_norm.flatten() # (C*H*S,)
-    valid_others_mask_flat = valid_gate_mask.flatten() & ~gate_self_chunk_mask.flatten() # Mask for valid, non-self entries
-    
-    # Only sort the valid 'other' entries
-    valid_others_indices = torch.where(valid_others_mask_flat)[0]
-    valid_others_values = others_flat[valid_others_indices]
-
-    sorted_others_values, sort_perm = torch.sort(valid_others_values, descending=True) # (N_valid_others,)
-    sorted_original_indices = valid_others_indices[sort_perm] # Original indices in C*H*S space, sorted by value
-
-    # 5) cumulative‑sum the sorted valid 'other' norms globally
-    cumsum_others_values = sorted_others_values.cumsum(dim=0) # (N_valid_others,)
-
-    # 6) find the smallest k where cumsum_ratio ≥ remain_ratio globally
-    ratio_values = cumsum_others_values / (total_norm_sum_overall + eps) # (N_valid_others,)
-    cond_values = ratio_values >= remain_ratio # (N_valid_others,) boolean mask
-    any_cond = cond_values.any() # scalar
-    
-    # Find the index of the first True value in the *sorted* list. If none, use all valid others.
-    cutoff_idx_in_sorted = torch.where(
-        any_cond, 
-        cond_values.float().argmax(dim=0), 
-        torch.tensor(len(sorted_others_values) - 1, device=gate.device, dtype=torch.long)
-    )
-
-    # 7) build a mask selecting the top-k others based on the cutoff
-    # Select the original indices corresponding to the top entries in the sorted list
-    selected_other_indices = sorted_original_indices[:cutoff_idx_in_sorted + 1]
-
-    # 8) create the mask in the original flat shape
-    others_mask_flat = torch.zeros_like(others_flat, dtype=torch.bool) # (C*H*S,)
-    if selected_other_indices.numel() > 0: # Check if any 'other' indices were selected
-        others_mask_flat[selected_other_indices] = True
-    others_mask = others_mask_flat.view(C, H, S) # (C, H, S)
-
-    # 9) finally, include every self‐chunk entry plus all selected others
-    final_gate_mask = valid_gate_mask & (others_mask | gate_self_chunk_mask)
-
-    return final_gate_mask
-
-
-def _select_threshold_head_global(
-    gate: torch.Tensor,
-    valid_gate_mask: torch.Tensor,
-    gate_self_chunk_mask: torch.Tensor,
-    simsum_threshold: float
-) -> torch.Tensor:
-    """
-    Selects <chunk, query> globally for each head based on threshold.
-    """
-    C, H, S = gate.shape
-    eps = 1e-6
-
-    # 1) LSE‐style normalization per head (across chunks and sequence dims)
-    gate_masked = torch.where(valid_gate_mask, gate, -torch.inf)
-    gate_min_val = torch.where(valid_gate_mask, gate, torch.inf)
-
-    max_per_head = gate_masked.amax(dim=(0, 2), keepdim=True)  # (1, H, 1)
-    min_per_head = gate_min_val.amin(dim=(0, 2), keepdim=True)  # (1, H, 1)
-    denom = max_per_head - min_per_head
-    denom = torch.where(denom <= eps, torch.ones_like(denom), denom)
-
-    gate_norm = (gate - min_per_head) / denom
-    gate_norm = torch.where(valid_gate_mask, gate_norm, 0.0)  # (C, H, S)
-
-    # 2) sum normalized self‐chunk contributions per head
-    self_norm_sum = (gate_norm * gate_self_chunk_mask).sum(dim=(0, 2))  # (H,)
-
-    # 3) total normalized sum per head
-    total_norm_sum = gate_norm.sum(dim=(0, 2))  # (H,)
-
-    # 4) how much more normalized weight needed per head
-    remain_ratio = simsum_threshold - self_norm_sum / (total_norm_sum + eps)  # (H,)
-    remain_ratio = torch.clamp(remain_ratio, min=0.0)
-
-    # 5) zero out self‐chunk entries to focus on "others"
-    others_norm = gate_norm.clone()
-    others_norm[gate_self_chunk_mask] = 0.0  # (C, H, S)
-
-    # 6) flatten chunk and sequence dims, per head
-    CS = C * S
-    others_flat = others_norm.permute(1, 0, 2).reshape(H, CS)  # (H, C*S)
-    valid_flat  = (valid_gate_mask & ~gate_self_chunk_mask) \
-                   .permute(1, 0, 2).reshape(H, CS)           # (H, C*S)
-
-    # 7) vectorized selection of “others” per head
-    masked_flat = torch.where(valid_flat, others_flat, torch.zeros_like(others_flat))
-    sorted_vals, sorted_idx = torch.sort(masked_flat, dim=1, descending=True)  # (H, C*S)
-
-    cumsum_vals = sorted_vals.cumsum(dim=1)                                    # (H, C*S)
-    ratio_vals  = cumsum_vals / (total_norm_sum.unsqueeze(1) + eps)            # (H, C*S)
-    cond        = ratio_vals >= remain_ratio.unsqueeze(1)                      # (H, C*S)
-
-    has_cutoff = cond.any(dim=1)                                               # (H,)
-    default    = torch.full((H,), CS - 1, device=gate.device, dtype=torch.long)
-    cutoff     = torch.where(has_cutoff, cond.float().argmax(dim=1), default)  # (H,)
-
-    idx_range   = torch.arange(CS, device=gate.device).unsqueeze(0)            # (1, C*S)
-    sorted_mask = idx_range <= cutoff.unsqueeze(1)                             # (H, C*S)
-
-    selected_flat = torch.zeros_like(valid_flat)                               # (H, C*S)
-    selected_flat.scatter_(1, sorted_idx, sorted_mask)                         # (H, C*S)
-
-    # 8) reshape selection mask back to (C, H, S)
-    others_mask = selected_flat.reshape(H, C, S).permute(1, 0, 2)               # (C, H, S)
-
-    # 9) include self‐chunks plus selected others, and obey valid mask
-    final_gate_mask = valid_gate_mask & (gate_self_chunk_mask | others_mask)
-
-    return final_gate_mask
-
-
-class MixedAttention(torch.autograd.Function):
-    @staticmethod
-    def forward(
-        ctx,
-        q,
-        k,
-        v,
-        self_attn_cu_seqlen,
-        moba_q,
-        moba_kv,
-        moba_cu_seqlen_q,
-        moba_cu_seqlen_kv,
-        max_seqlen,
-        moba_chunk_size,
-        moba_q_sh_indices,
-    ):
-        ctx.max_seqlen = max_seqlen
-        ctx.moba_chunk_size = moba_chunk_size
-        ctx.softmax_scale = softmax_scale = q.shape[-1] ** (-0.5)
-
-        # Non-causal self-attention branch
-        # return out, softmax_lse, S_dmask, rng_state
-        self_attn_out_sh, self_attn_lse_hs, _, _ = _flash_attn_varlen_forward(
-            q=q,
-            k=k,
-            v=v,
-            cu_seqlens_q=self_attn_cu_seqlen,
-            cu_seqlens_k=self_attn_cu_seqlen,
-            max_seqlen_q=max_seqlen,
-            max_seqlen_k=max_seqlen,
-            softmax_scale=softmax_scale,
-            causal=False,
-            dropout_p=0.0,
-        )
-        # MOBA attention branch (non-causal)
-        moba_attn_out, moba_attn_lse_hs, _, _ = _flash_attn_varlen_forward(
-            q=moba_q,
-            k=moba_kv[:, 0],
-            v=moba_kv[:, 1],
-            cu_seqlens_q=moba_cu_seqlen_q,
-            cu_seqlens_k=moba_cu_seqlen_kv,
-            max_seqlen_q=max_seqlen,
-            max_seqlen_k=moba_chunk_size,
-            softmax_scale=softmax_scale,
-            causal=False,
-            dropout_p=0.0,
-        )
-
-        self_attn_lse_sh = self_attn_lse_hs.t().contiguous()
-        moba_attn_lse = moba_attn_lse_hs.t().contiguous()
-
-        output = torch.zeros((q.shape[0], q.shape[1], q.shape[2]), device=q.device, dtype=torch.float32)
-        output_2d = output.view(-1, q.shape[2])
-
-        max_lse_1d = self_attn_lse_sh.view(-1)
-        max_lse_1d = max_lse_1d.index_reduce(
-            0, moba_q_sh_indices, moba_attn_lse.view(-1), "amax"
-        )
-        self_attn_lse_sh = self_attn_lse_sh - max_lse_1d.view_as(self_attn_lse_sh)
-        moba_attn_lse = (
-            moba_attn_lse.view(-1)
-            .sub(max_lse_1d.index_select(0, moba_q_sh_indices))
-            .reshape_as(moba_attn_lse)
-        )
-
-        mixed_attn_se_sh = self_attn_lse_sh.exp()
-        moba_attn_se = moba_attn_lse.exp()
-
-        mixed_attn_se_sh.view(-1).index_add_(
-            0, moba_q_sh_indices, moba_attn_se.view(-1)
-        )
-        mixed_attn_lse_sh = mixed_attn_se_sh.log()
-
-        # Combine self-attention output
-        factor = (self_attn_lse_sh - mixed_attn_lse_sh).exp()  # [S, H]
-        self_attn_out_sh = self_attn_out_sh * factor.unsqueeze(-1)
-        output_2d += self_attn_out_sh.reshape_as(output_2d)
-
-        # Combine MOBA attention output
-        mixed_attn_lse = (
-            mixed_attn_lse_sh.view(-1)
-            .index_select(0, moba_q_sh_indices)
-            .view_as(moba_attn_lse)
-        )
-        factor = (moba_attn_lse - mixed_attn_lse).exp()  # [S, H]
-        moba_attn_out = moba_attn_out * factor.unsqueeze(-1)
-        raw_attn_out = moba_attn_out.view(-1, moba_attn_out.shape[-1])
-        output_2d.index_add_(0, moba_q_sh_indices, raw_attn_out)
-        output = output.to(q.dtype)
-        mixed_attn_lse_sh = mixed_attn_lse_sh + max_lse_1d.view_as(mixed_attn_se_sh)
-        ctx.save_for_backward(
-            output,
-            mixed_attn_lse_sh,
-            q,
-            k,
-            v,
-            self_attn_cu_seqlen,
-            moba_q,
-            moba_kv,
-            moba_cu_seqlen_q,
-            moba_cu_seqlen_kv,
-            moba_q_sh_indices,
-        )
-
-        return output
-
-    @staticmethod
-    def backward(ctx, d_output):
-
-        max_seqlen = ctx.max_seqlen
-        moba_chunk_size = ctx.moba_chunk_size
-        softmax_scale = ctx.softmax_scale
-
-        (
-            output,
-            mixed_attn_vlse_sh,
-            q,
-            k,
-            v,
-            self_attn_cu_seqlen,
-            moba_q,
-            moba_kv,
-            moba_cu_seqlen_q,
-            moba_cu_seqlen_kv,
-            moba_q_sh_indices,
-        ) = ctx.saved_tensors
-
-        d_output = d_output.contiguous()
-
-        dq = torch.empty_like(q)
-        dk = torch.empty_like(k)
-        dv = torch.empty_like(v)
-        _ = _flash_attn_varlen_backward(
-            dout=d_output,
-            q=q,
-            k=k,
-            v=v,
-            out=output,
-            softmax_lse=mixed_attn_vlse_sh.t().contiguous(),
-            dq=dq,
-            dk=dk,
-            dv=dv,
-            cu_seqlens_q=self_attn_cu_seqlen,
-            cu_seqlens_k=self_attn_cu_seqlen,
-            max_seqlen_q=max_seqlen,
-            max_seqlen_k=max_seqlen,
-            softmax_scale=softmax_scale,
-            causal=False,
-            dropout_p=0.0,
-            softcap=0.0,
-            alibi_slopes=None,
-            deterministic=True,
-            window_size_left=-1,
-            window_size_right=-1
-        )
-
-        headdim = q.shape[-1]
-        d_moba_output = (
-            d_output.view(-1, headdim).index_select(0, moba_q_sh_indices).unsqueeze(1)
-        )
-        moba_output = (
-            output.view(-1, headdim).index_select(0, moba_q_sh_indices).unsqueeze(1)
-        )
-
-        mixed_attn_vlse = (
-            mixed_attn_vlse_sh.view(-1).index_select(0, moba_q_sh_indices).view(1, -1)
-        )
-
-        dmq = torch.empty_like(moba_q)
-        dmkv = torch.empty_like(moba_kv)
-        _ = _flash_attn_varlen_backward(
-            dout=d_moba_output,
-            q=moba_q,
-            k=moba_kv[:, 0],
-            v=moba_kv[:, 1],
-            out=moba_output,
-            softmax_lse=mixed_attn_vlse,
-            dq=dmq,
-            dk=dmkv[:,0],
-            dv=dmkv[:,1],
-            cu_seqlens_q=moba_cu_seqlen_q,
-            cu_seqlens_k=moba_cu_seqlen_kv,
-            max_seqlen_q=max_seqlen,
-            max_seqlen_k=moba_chunk_size,
-            softmax_scale=softmax_scale,
-            causal=False,
-            dropout_p=0.0,
-            softcap=0.0,
-            alibi_slopes=None,
-            deterministic=True,
-            window_size_left=-1,
-            window_size_right=-1
-        )
-
-        return dq, dk, dv, None, dmq, dmkv, None, None, None, None, None
-
-
-def moba_attn_varlen(
-    q: torch.Tensor,
-    k: torch.Tensor,
-    v: torch.Tensor,
-    cu_seqlens: torch.Tensor,
-    max_seqlen: int,
-    moba_chunk_size: int,
-    moba_topk: int,
-    select_mode: str = 'threshold', # "topk" or "threshold"
-    simsum_threshold: float = 0.25,
-    threshold_type: str = 'query_head',
-) -> torch.Tensor:
-    """
-    Accelerated MOBA attention for vision tasks with proper LSE normalization.
-
-    This version:
-      - Splits KV into chunks.
-      - For each query head, selects the top-k relevant KV chunks (including the self chunk)
-        by amplifying the diagonal (self-chunk) logits.
-      - Aggregates the attention outputs from the selected chunks using a log-sum-exp
-        reduction so that attending to each query over the selected chunks is equivalent
-        to the original algorithm.
-    """
-    # Stack keys and values.
-    kv = torch.stack((k, v), dim=1)
-    seqlen, num_head, head_dim = q.shape
-
-    # Compute chunk boundaries.
-    cu_chunk, filtered_chunk_indices, num_filtered_chunk, chunk_to_batch = calc_chunks(
-        cu_seqlens, moba_chunk_size
-    )
-
-    self_attn_cu_seqlen = cu_chunk
-
-    # Update top-k selection to include the self chunk.
-    moba_topk = min(moba_topk, num_filtered_chunk)
-
-    # --- Build filtered KV from chunks ---
-    chunk_starts = cu_chunk[filtered_chunk_indices]          # [num_filtered_chunk]
-    chunk_ends = cu_chunk[filtered_chunk_indices + 1]          # [num_filtered_chunk]
-    chunk_lengths = chunk_ends - chunk_starts                  # [num_filtered_chunk]
-    max_chunk_len = int(chunk_lengths.max().item())
-
-    range_tensor = torch.arange(max_chunk_len, device=kv.device, dtype=chunk_starts.dtype).unsqueeze(0)
-    indices = chunk_starts.unsqueeze(1) + range_tensor
-    indices = torch.clamp(indices, max=kv.shape[0] - 1)
-    valid_mask = range_tensor < chunk_lengths.unsqueeze(1)
-    gathered = kv[indices.view(-1)].view(num_filtered_chunk, max_chunk_len, *kv.shape[1:])
-    gathered = gathered * valid_mask.unsqueeze(-1).unsqueeze(-1).unsqueeze(-1).type_as(gathered)
-
-    # Compute key_gate_weight over valid tokens.
-    key_values = gathered[:, :, 0].float()  # [num_filtered_chunk, max_chunk_len, num_head, head_dim]
-    valid_mask_exp = valid_mask.unsqueeze(-1).unsqueeze(-1)
-    key_sum = (key_values * valid_mask_exp).sum(dim=1)
-    divisor = valid_mask.sum(dim=1).unsqueeze(-1).unsqueeze(-1)
-    key_gate_weight = key_sum / divisor  # [num_filtered_chunk, num_head, head_dim]
-
-    # Compute gate logits between key_gate_weight and queries.
-    q_float = q.float()
-    # gate = torch.einsum("nhd,shd->nhs", key_gate_weight, q_float)  # [num_filtered_chunk, num_head, seqlen]
-    gate = torch.bmm(key_gate_weight.permute(1, 0, 2), q_float.permute(1, 0, 2).transpose(1, 2)).permute(1, 0, 2)
-
-    # Amplify the diagonal (self chunk) contributions.
-    gate_seq_idx = torch.arange(seqlen, device=q.device, dtype=torch.int32).unsqueeze(0).expand(num_filtered_chunk, seqlen)
-    chunk_start = cu_chunk[filtered_chunk_indices]  # [num_filtered_chunk]
-    chunk_end = cu_chunk[filtered_chunk_indices + 1]  # [num_filtered_chunk]
-    gate_self_chunk_mask = ((gate_seq_idx >= chunk_start.unsqueeze(1)) &
-                            (gate_seq_idx < chunk_end.unsqueeze(1))).unsqueeze(1).expand(-1, num_head, -1)
-    amplification_factor = 1e9  # Example factor; adjust as needed.
-    origin_gate = gate.clone()
-    gate = gate.clone()
-    if select_mode == "topk":
-        gate[gate_self_chunk_mask] += amplification_factor
-
-    # Exclude positions that are outside the valid batch boundaries.
-    batch_starts = cu_seqlens[chunk_to_batch[filtered_chunk_indices]]
-    batch_ends = cu_seqlens[chunk_to_batch[filtered_chunk_indices] + 1]
-    gate_batch_start_mask = gate_seq_idx < batch_starts.unsqueeze(1)
-    gate_batch_end_mask = gate_seq_idx >= batch_ends.unsqueeze(1)
-    gate_inf_mask = gate_batch_start_mask | gate_batch_end_mask
-    gate.masked_fill_(gate_inf_mask.unsqueeze(1), -float("inf"))
-
-    if select_mode == 'topk':
-        # We amplify self‐chunk in gate already, so self entries will rank highest.
-        valid_gate_mask = gate != -float("inf")
-        if threshold_type == 'query_head':
-            # === per‐<head,seq> top-k across chunks (original behavior) ===
-            # gate: (C, H, S)
-            _, gate_topk_idx = torch.topk(gate, k=moba_topk, dim=0, largest=True, sorted=False)
-            gate_idx_mask = torch.zeros_like(gate, dtype=torch.bool)
-            gate_idx_mask.scatter_(0, gate_topk_idx, True)
-            gate_mask = valid_gate_mask & gate_idx_mask
-        elif threshold_type == 'overall':
-            # === global top-k across all (chunk, head, seq) entries ===
-            C, H, S = gate.shape
-            flat_gate = gate.flatten()
-            flat_mask = valid_gate_mask.flatten()
-            flat_gate_masked = torch.where(flat_mask, flat_gate, -float("inf"))
-            # pick topk global entries
-            vals, idx = torch.topk(flat_gate_masked, k=moba_topk * H * S, largest=True, sorted=False)
-            others_mask_flat = torch.zeros_like(flat_mask, dtype=torch.bool)
-            others_mask_flat[idx] = True
-            gate_mask = (valid_gate_mask.flatten() & others_mask_flat).view(gate.shape)
-        elif threshold_type == 'head_global':
-            # per-head top-k across all chunks and sequence positions
-            C, H, S = gate.shape
-            CS = C * S
-            flat_gate  = gate.permute(1, 0, 2).reshape(H, CS)
-            flat_valid = valid_gate_mask.permute(1, 0, 2).reshape(H, CS)
-            flat_gate_masked = torch.where(flat_valid, flat_gate, torch.full_like(flat_gate, -float('inf')))
-            # pick top-k indices per head
-            _, topk_idx = torch.topk(flat_gate_masked, k=moba_topk * S, dim=1, largest=True, sorted=False)
-            gate_idx_flat = torch.zeros_like(flat_valid, dtype=torch.bool)
-            gate_idx_flat.scatter_(1, topk_idx, True)
-            gate_mask = gate_idx_flat.reshape(H, C, S).permute(1, 0, 2)
-        else:
-            raise ValueError(
-                f"Invalid threshold_type for topk: {threshold_type}. "
-                "Choose 'query_head', 'block', or 'overall'."
-            )
-    elif select_mode == 'threshold':
-        # Delegate to the specific thresholding function
-        valid_gate_mask = gate != -float("inf")  # (num_chunk, num_head, seqlen)
-        if threshold_type == 'query_head':
-            gate_mask = _select_threshold_query_head(gate, valid_gate_mask, gate_self_chunk_mask, simsum_threshold)
-        elif threshold_type == 'block':
-            gate_mask = _select_threshold_block(gate, valid_gate_mask, gate_self_chunk_mask, simsum_threshold)
-        elif threshold_type == 'overall':
-            gate_mask = _select_threshold_overall(gate, valid_gate_mask, gate_self_chunk_mask, simsum_threshold)
-        elif threshold_type == 'head_global':
-            gate_mask = _select_threshold_head_global(gate, valid_gate_mask, gate_self_chunk_mask, simsum_threshold)
-        else:
-            raise ValueError(f"Invalid threshold_type: {threshold_type}. Choose 'query_head', 'block', or 'overall'.")
-    else:
-        raise ValueError(f"Invalid select_mode: {select_mode}. Choose 'topk' or 'threshold'.")
-    
-    # eliminate self_chunk in MoBA branch
-    gate_mask = gate_mask & ~gate_self_chunk_mask
-    # if gate_mask is all false, perform flash_attn instead
-    if gate_mask.sum() == 0:
-        return flash_attn_varlen_func(
-            q, k, v, cu_seqlens, cu_seqlens, max_seqlen, max_seqlen, causal=False
-        )
-
-    # Determine which query positions are selected.
-    # nonzero_indices has shape [N, 3] where each row is [chunk_index, head_index, seq_index].
-    moba_q_indices = gate_mask.reshape(gate_mask.shape[0], -1).nonzero(as_tuple=True)[-1] # [(h s k)]
-    moba_q_sh_indices = (moba_q_indices % seqlen) * num_head + (moba_q_indices // seqlen)
-    moba_q = rearrange(q, "s h d -> (h s) d").index_select(0, moba_q_indices).unsqueeze(1)
-
-    # Build cumulative sequence lengths for the selected queries.
-    moba_seqlen_q = gate_mask.sum(dim=-1).flatten()
-    q_zero_mask = moba_seqlen_q == 0
-    valid_expert_mask = ~q_zero_mask
-    if q_zero_mask.sum() > 0:
-        moba_seqlen_q = moba_seqlen_q[valid_expert_mask]
-    moba_cu_seqlen_q = torch.cat(
-        (
-            torch.tensor([0], device=q.device, dtype=moba_seqlen_q.dtype),
-            moba_seqlen_q.cumsum(dim=0),
-        ),
-        dim=0,
-    ).to(torch.int32)
-
-    # Rearrange gathered KV for the MOBA branch.
-    experts_tensor = rearrange(gathered, "nc cl two h d -> (nc h) cl two d")
-    valid_expert_lengths = chunk_lengths.unsqueeze(1).expand(num_filtered_chunk, num_head).reshape(-1).to(torch.int32)
-    if q_zero_mask.sum() > 0:
-        experts_tensor = experts_tensor[valid_expert_mask]
-        valid_expert_lengths = valid_expert_lengths[valid_expert_mask]
-
-    seq_range = torch.arange(experts_tensor.shape[1], device=experts_tensor.device).unsqueeze(0)
-    mask = seq_range < valid_expert_lengths.unsqueeze(1)
-    moba_kv = experts_tensor[mask]  # Shape: ((nc h cl_valid) two d)
-    moba_kv = moba_kv.unsqueeze(2)   # Shape: ((nc h cl_valid) two 1 d)
-
-    moba_cu_seqlen_kv = torch.cat(
-        [torch.zeros(1, device=experts_tensor.device, dtype=torch.int32),
-         valid_expert_lengths.cumsum(dim=0)],
-        dim=0,
-    ).to(torch.int32)
-
-    assert (
-        moba_cu_seqlen_kv.shape == moba_cu_seqlen_q.shape
-    ), f"Mismatch between moba_cu_seqlen_kv.shape and moba_cu_seqlen_q.shape: {moba_cu_seqlen_kv.shape} vs {moba_cu_seqlen_q.shape}"
-
-    return MixedAttention.apply(
-        q,
-        k,
-        v,
-        self_attn_cu_seqlen,
-        moba_q,
-        moba_kv,
-        moba_cu_seqlen_q,
-        moba_cu_seqlen_kv,
-        max_seqlen,
-        moba_chunk_size,
-        moba_q_sh_indices,
-    )
-
-
-def process_moba_input(
-    x,
-    patch_resolution,
-    chunk_size,
-):
-    """
-    Process inputs for the attention function.
-
-    Args:
-        x (torch.Tensor): Input tensor with shape [batch_size, num_patches, num_heads, head_dim].
-        patch_resolution (tuple): Tuple containing the patch resolution (t, h, w).
-        chunk_size (int): Size of the chunk. (maybe tuple or int, according to chunk type)
-
-    Returns:
-        torch.Tensor: Processed input tensor.
-    """
-    if isinstance(chunk_size, float) or isinstance(chunk_size, int):
-        moba_chunk_size = int(chunk_size * patch_resolution[1] * patch_resolution[2])
-    else:
-        assert isinstance(chunk_size, (Tuple, list)), f"chunk_size should be a tuple, list, or int, now it is: {type(chunk_size)}"
-        if len(chunk_size) == 2:
-            assert patch_resolution[1] % chunk_size[0] == 0 and patch_resolution[2] % chunk_size[1] == 0, f"spatial patch_resolution {patch_resolution[1:]} should be divisible by 2d chunk_size {chunk_size}"
-            nch, ncw = patch_resolution[1] // chunk_size[0], patch_resolution[2] // chunk_size[1]
-            x = rearrange(x, "b (t nch ch ncw cw) n d -> b (nch ncw t ch cw) n d", t=patch_resolution[0], nch=nch, ncw=ncw, ch=chunk_size[0], cw=chunk_size[1])
-            moba_chunk_size = patch_resolution[0] * chunk_size[0] * chunk_size[1]
-        elif len(chunk_size) == 3:
-            assert patch_resolution[0] % chunk_size[0] == 0 and patch_resolution[1] % chunk_size[1] == 0 and patch_resolution[2] % chunk_size[2] == 0, f"patch_resolution {patch_resolution} should be divisible by 3d chunk_size {chunk_size}"
-            nct, nch, ncw = patch_resolution[0] // chunk_size[0], patch_resolution[1] // chunk_size[1], patch_resolution[2] // chunk_size[2]
-            x = rearrange(x, "b (nct ct nch ch ncw cw) n d -> b (nct nch ncw ct ch cw) n d", nct=nct, nch=nch, ncw=ncw, ct=chunk_size[0], ch=chunk_size[1], cw=chunk_size[2])
-            moba_chunk_size = chunk_size[0] * chunk_size[1] * chunk_size[2]
-        else:
-            raise ValueError(f"chunk_size should be a int, or a tuple of length 2 or 3, now it is: {len(chunk_size)}")
-        
-    return x, moba_chunk_size
-
-
-def process_moba_output(
-    x,
-    patch_resolution,
-    chunk_size,
-):
-    if isinstance(chunk_size, float) or isinstance(chunk_size, int):
-        pass
-    elif len(chunk_size) == 2:
-        x = rearrange(x, "b (nch ncw t ch cw) n d -> b (t nch ch ncw cw) n d", nch=patch_resolution[1] // chunk_size[0], ncw=patch_resolution[2] // chunk_size[1], t=patch_resolution[0], ch=chunk_size[0], cw=chunk_size[1])
-    elif len(chunk_size) == 3:
-        x = rearrange(x, "b (nct nch ncw ct ch cw) n d -> b (nct ct nch ch ncw cw) n d", nct=patch_resolution[0] // chunk_size[0], nch=patch_resolution[1] // chunk_size[1], ncw=patch_resolution[2] // chunk_size[2], ct=chunk_size[0], ch=chunk_size[1], cw=chunk_size[2])
-
-    return x
-
-
-# TEST
-def generate_data(batch_size, seqlen, num_head, head_dim, dtype):
-    random.seed(0)
-    torch.manual_seed(0)
-    torch.cuda.manual_seed(0)
-    device = torch.cuda.current_device()
-
-    q = torch.randn((batch_size, seqlen, num_head, head_dim), requires_grad=True).to(dtype=dtype, device='cuda')
-    k = torch.randn((batch_size, seqlen, num_head, head_dim), requires_grad=True).to(dtype=dtype, device='cuda')
-    v = torch.randn((batch_size, seqlen, num_head, head_dim), requires_grad=True).to(dtype=dtype, device='cuda')
-    print(f"q.shape: {q.shape}, k.shape: {k.shape}, v.shape: {v.shape}")
-    cu_seqlens = torch.arange(0, q.shape[0] * q.shape[1] + 1, q.shape[1], dtype=torch.int32, device='cuda')
-    max_seqlen = q.shape[1]
-    q = rearrange(q, "b s ... -> (b s) ...")
-    k = rearrange(k, "b s ... -> (b s) ...")
-    v = rearrange(v, "b s ... -> (b s) ...")
-
-    return q, k, v, cu_seqlens, max_seqlen
-
-
-def test_attn_varlen_moba_speed(batch, head, seqlen, head_dim, moba_chunk_size, moba_topk, dtype=torch.bfloat16, select_mode='threshold', simsum_threshold=0.25, threshold_type='query_head'):
-    """Speed test comparing flash_attn vs moba_attention"""
-    # Get data
-    q, k, v, cu_seqlen, max_seqlen = generate_data(batch, seqlen, head, head_dim, dtype)
-    print(f"batch:{batch} head:{head} seqlen:{seqlen} chunk:{moba_chunk_size} topk:{moba_topk} select_mode: {select_mode} simsum_threshold:{simsum_threshold}")
-    vo_grad = torch.randn_like(q)
-    
-    # Warmup
-    warmup_iters = 3
-    perf_test_iters = 10
-
-    # Warmup
-    for _ in range(warmup_iters):
-        o = flash_attn_varlen_func(q, k, v, cu_seqlen, cu_seqlen, max_seqlen, max_seqlen, causal=False)
-        torch.autograd.backward(o, vo_grad)
-    
-    torch.cuda.synchronize()
-    start_flash = time.perf_counter()
-    for _ in range(perf_test_iters):
-        o = flash_attn_varlen_func(q, k, v, cu_seqlen, cu_seqlen, max_seqlen, max_seqlen, causal=False)
-        torch.autograd.backward(o, vo_grad)
-        
-    torch.cuda.synchronize()
-    time_flash = (time.perf_counter() - start_flash) / perf_test_iters * 1000
-
-    # Warmup
-    for _ in range(warmup_iters):
-        om = moba_attn_varlen(q, k, v, cu_seqlen, max_seqlen, moba_chunk_size=moba_chunk_size, moba_topk=moba_topk, select_mode=select_mode, simsum_threshold=simsum_threshold, threshold_type=threshold_type)
-        torch.autograd.backward(om, vo_grad)
-
-    torch.cuda.synchronize()
-    start_moba = time.perf_counter()
-    for _ in range(perf_test_iters):
-        om = moba_attn_varlen(q, k, v, cu_seqlen, max_seqlen, moba_chunk_size=moba_chunk_size, moba_topk=moba_topk, select_mode=select_mode, simsum_threshold=simsum_threshold, threshold_type=threshold_type)
-        torch.autograd.backward(om, vo_grad)
-    
-    torch.cuda.synchronize()
-    time_moba = (time.perf_counter() - start_moba) / perf_test_iters * 1000
-    
-    print(f"Flash: {time_flash:.2f}ms, MoBA: {time_moba:.2f}ms")
-    print(f"Speedup:  {time_flash / time_moba:.2f}x")
-
-
-
-if __name__ == "__main__":
-    """
-    CUDA_VISIBLE_DEVICES=1 \
-    python -u csrc/attn/vmoba_attn/vmoba/vmoba.py
-    """
-    test_attn_varlen_moba_speed(batch=1, head=12, seqlen=32760, head_dim=128, moba_chunk_size=32760 // 3 // 6 // 4, moba_topk=3, select_mode='threshold', simsum_threshold=0.3, threshold_type='query_head')
@@ -568,6 +568,7 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {
        __syncthreads(); // wait for sd_smem shared memory write
        warpgroup::mm_AtB(qg_reg, ds_smem_t[0], k_smem[0]); //delat dQ = dSK
        warpgroup::mma_commit_group();
+        tma::store_async_wait();
        warpgroup::mma_async_wait();
        // store qg to shared memory
        warpgroup::store(qg_smem, qg_reg);
@@ -577,7 +578,6 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {
        if (threadIdx.x / 32 == 0) {
            coord<qg_tile> tile_idx = {blockIdx.z, blockIdx.y, store_qg_block_index, 0};
            tma::store_add_async(g.qg, qg_smem, tile_idx);
-            tma::store_async_wait();
        }
    }

@@ -624,6 +624,7 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {
        __syncthreads(); // wait for sd_smem shared memory write
        warpgroup::mm_AtB(qg_reg, ds_smem_t[0], k_smem[0]); //delat dQ = dSK
        warpgroup::mma_commit_group();
+        tma::store_async_wait();
        warpgroup::mma_async_wait();
        // store qg to shared memory
        warpgroup::store(qg_smem, qg_reg);
@@ -633,14 +634,13 @@ void bwd_attend_ker(const __grid_constant__ bwd_globals<D> g) {
        if (threadIdx.x / 32 == 0) {
            coord<qg_tile> tile_idx = {blockIdx.z, blockIdx.y, store_qg_block_index, 0};
            tma::store_add_async(g.qg, qg_smem, tile_idx);
-            tma::store_async_wait();
        }
    }

    // store kq and vq

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

    warpgroup::store(kg_smem[0], kg_reg);
@@ -1174,4 +1174,4 @@ block_sparse_attention_backward(torch::Tensor q,

    return {qg, kg, vg};
    //cudadevicesynchronize();
-}
+}
@@ -1,452 +0,0 @@
-{
-  "data": [
-    {
-      "prompt": "Young man skating with a skateboard on the ramps with graffiti of a park with trees, on a sunny day.", 
-      "image_path": "images/mixkit-boy-skating-with-a-skateboard-in-a-park-with-ramps-34389.png"
-    },
-    {
-      "prompt": "In the midst of the joyous New Year's Eve celebration, the cheerful group of friends, their spirits lifted by the festivities, decides to immortalize the moment with a vibrant snapshot", 
-      "image_path": "images/mixkit-a-cheerful-group-of-friends-celebrate-new-years-eve-and-51525.png"
-    },
-    {
-      "prompt": "A man and a woman playing in a field with grass, during a bright afternoon, while cars pass by in the distance.", 
-      "image_path": "images/mixkit-a-cute-couple-playing-on-the-grass-4688.png"
-    },
-    {
-      "prompt": "Aerial view of a rocky mountain in the forest at a sunny day drone flight footage", 
-      "image_path": "images/mixkit-aerial-view-of-a-rocky-mountain-in-the-forest-50589.png"
-    },
-    {
-      "prompt": "A little girl wearing a pink security helmet and denim overall discovers the art of cycling amidst the serene park, as the camera captures her graceful progress.", 
-      "image_path": "images/mixkit-a-little-girl-cruises-through-the-forest-path-on-her-50088.png"
-    },
-    {
-      "prompt": "Aerial shot of a beach shore with sea waves. Big rocks on the sand at an alone beach.", 
-      "image_path": "images/mixkit-aerial-shot-of-a-beach-with-sea-waves-1087.png"
-    },
-    {
-      "prompt": "Young woman cleaning her house decorated with plants and decorations, while dancing happily to music in her headphones.", 
-      "image_path": "images/mixkit-woman-cleaning-her-house-dancing-happy-43379.png"
-    },
-    {
-      "prompt": "Aerial tour in a meadow surrounded by hills on the horizon, while some birds fly low over a lake.", 
-      "image_path": "images/mixkit-birds-flying-low-over-a-lake-in-a-meadow-41417.png"
-    },
-    {
-      "prompt": "In the video, a lone rider guides a majestic horse across an expansive, open field as the sun sets in the background. The rider, dressed in a classic blue shirt and wide-brimmed hat, sits confidently in the saddle, silhouetted against the warm glow of the evening sky. The horse moves gracefully, its mane and tail flowing with each step, creating a sense of harmony between horse and rider. Surrounding the pair, towering trees form a natural border, their leaves gently rustling in the breeze. The shadows lengthen on the ground, accentuating the serene and timeless feel of the scene. The distant hills and wooden fences frame the horizon, adding depth to the tranquil landscape. A few horses graze peacefully in the background, blending into the pastoral setting. The overall ambiance evokes a sense of calmness and quietude, capturing a perfect moment in the golden light of dusk.", 
-      "image_path": "images/mixkit-a-rancher-riding-a-horse-at-sunset-1143.png"
-    },
-    {
-      "prompt": "In the video, a martial artist dressed in a traditional white uniform with a black belt demonstrates a series of precise movements against a stark black background. The individual gracefully transitions between stances, embodying a sense of focused discipline and control. Each motion is executed with a deliberate pace, showcasing the fluidity of martial arts techniques. The soft lighting creates subtle highlights on the uniform, adding depth to the figure as it moves. The practitioner begins with an open-hand pose, feet firmly grounded, gradually shifting to a powerful forward punch. The fluidity of the sequence displays a mastery of balance and poise. Every trajectory of the limbs is precise and deliberate, capturing the elegance and strength of martial arts. The serene, isolated setting enhances the intensity and concentration of the practitioner. This visual presentation is an elegant interplay of motion and stillness, displaying the art form's discipline and grace.", 
-      "image_path": "images/mixkit-a-young-man-practicing-his-karate-moves-49635.png"
-    },
-    {
-      "prompt": "In a serene and softly lit yoga studio, three individuals engage in a yoga session, each performing an upward-facing stretch. The central figure is a woman with shoulder-length brown hair, dressed in a light cropped top and green leggings, her posture reflecting grace and concentration. To her right, another participant, a woman in a purple outfit, mirrors the pose with equal poise. On her left, a person with a bun focuses intently, supported slightly by yoga blocks beneath their hands. The warm-colored wooden floor contrasts soothingly with the soft pastel mural on the back wall, featuring an abstract design and partial visage of a serene face. Natural light floods the space from a large window on the right, where lush greens peek through, adding an element of tranquility. In the corner of the room, a collection of meditation instruments, including a gong and a Buddha statue, subtly frame the peaceful setting. The mood is calm yet focused, as all three participants are deeply engaged in their practice. The scene combines elements of balance, harmony, and a shared journey towards mindfulness. This depiction captures the essence of a yoga session that blends personal growth with collective experience.", 
-      "image_path": "images/mixkit-small-group-of-people-doing-yoga-together-43730.png"
-    },
-    {
-      "prompt": "In the deep blue expanse of the ocean, two dolphins glide effortlessly, their sleek bodies reflecting the sunlight filtering through the water. The prominent shadows and caustics create a shimmering effect on their skin, capturing the beauty of their natural habitat. Each dolphin moves with a fluid grace, occasionally interacting with gentle nudges, showcasing their playful and social nature. The scene is vibrant and dynamic, with the clear blue background accentuating the dolphins' movements, making it an ideal subject for AI recreation.", 
-      "image_path": "images/mixkit-dolphins-underwater-4133.png"
-    },
-    {
-      "prompt": "A bustling ski slope comes alive with skiers descending a pristine, snow-covered hill, surrounded by towering, snow-draped evergreens. Several figures stand atop the slope, silhouetted against a clear blue sky, preparing to embark on their ski run. The chair lift on the right continuously drops off eager adventurers, adding to the excitement at the hilltop. Each skier, clad in colorful winter gear, carves distinct paths into the textured snow as they weave their way down. The interplay of sunlight and shadows accentuates the myriad tracks etched into the slope, creating a dynamic visual rhythm. The scene captures a vibrant winter wonderland, full of action and the thrill of a perfect ski day.", 
-      "image_path": "images/mixkit-skiers-on-a-snowy-slope-3327.png"
-    },
-    {
-      "prompt": "A determined climber is scaling a massive rock face, showcasing exceptional strength and skill. The person, clad in a teal shirt and dark pants, climbs with precision, their movements measured and deliberate. They are secured by climbing gear, which includes ropes and a harness, emphasizing their commitment to safety. The rugged texture of the sandy-colored rock provides an imposing backdrop, adding drama and scale to the climb. In the distance, other large rock formations and sparse vegetation can be seen under a bright, overcast sky, contributing to the natural and adventurous atmosphere. The scene captures a moment of focus and challenge, highlighting the climber's tenacity and the breathtaking environment.",
-      "image_path": "images/mixkit-alpinist-climbing-a-huge-rock-in-a-desert-43306.png"
-    },
-    {
-      "prompt": "A silver SUV drives along a winding, snow-covered mountain road, with dense pine trees blanketed in snow lining both sides. The scene is serene, with the vehicle moving smoothly, possibly on a winter journey or vacation. As the SUV disappears around the bend, another, darker SUV follows, creating a sense of motion and perspective on the snow-dusted asphalt. The towering, snow-laden rock formation to the right contrasts with the dark green of the pines, highlighting the peacefulness of the wintry landscape.",
-      "image_path": "images/mixkit-curve-on-a-snowy-forest-road-3317.png"
-    },
-    {
-      "prompt": "A solitary boat glides across the expansive, tranquil expanse of a serene lake. The vessel leaves a gentle wake behind, creating delicate ripples across the mirror-like surface. The water appears a rich shade of teal, seamlessly blending with the sky at the horizon. Silhouettes of distant trees are faintly visible, creating a picturesque backdrop that enhances the solitary journey of the boat. The sky is a calm gradient, shifting from soft oranges near the shore to the pale blues above. In the distance, a few slender poles emerge from the water, remnants of an old structure or natural formation. The mood of the scene is one of peace and solitude, with the boat journeying steadily through the quiet landscape. There is a sense of endless possibilities as the boat moves toward the unseen beyond the frame. The simplicity and stillness of the scene invite contemplation and reflection, encapsulating a perfect moment of quietude on the water.",
-      "image_path": "images/mixkit-motorboat-on-a-large-lake-with-turquoise-blue-waters-4996.png"
-    },
-    {
-      "prompt": "Far shot from high behind a beach on a sunny day, with crowds of people everywhere, buildings and hotels behind, and the sea on the horizon as ships cross it.",
-      "image_path": "480P_i2v_mixkit_images/Far shot from high behind a beach on a sunny day, with crowds of people everywhere, buildings and ho.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Skillful cyclist doing a wheelie on a bike while riding through a forest, on a dirt road, surrounded by many trees, in the morning.",
-      "image_path": "480P_i2v_mixkit_images/Skillful cyclist doing a wheelie on a bike while riding through a forest, on a dirt road, surrounded.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Portrait of a man with a mask and fluorescent paint dancing in front of the camera, under a blue party light.",
-      "image_path": "480P_i2v_mixkit_images/Portrait of a man with a mask and fluorescent paint dancing in front of the camera, under a blue par.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Timelapse of a luxury swimming pool in a boutique hotel.",
-      "image_path": "480P_i2v_mixkit_images/Timelapse of a luxury swimming pool in a boutique hotel..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Man dressed in 80's style dances very happily in his kitchen while listening to music on his radio and drinking wine.",
-      "image_path": "480P_i2v_mixkit_images/Man dressed in 80's style dances very happily in his kitchen while listening to music on his radio a.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Romantic couple embracing and looking at each other in the middle of a forest, during a break on a road trip through nature.",
-      "image_path": "480P_i2v_mixkit_images/Romantic couple embracing and looking at each other in the middle of a forest, during a break on a r.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Big city with many big buildings and a big capitol, in a view from a slow drive through the heights in the afternoon.",
-      "image_path": "480P_i2v_mixkit_images/Big city with many big buildings and a big capitol, in a view from a slow drive through the heights .jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "A man appears in different parts of an abandoned building, standing in the dark, swinging a baseball bat, outside.",
-      "image_path": "480P_i2v_mixkit_images/A man appears in different parts of an abandoned building, standing in the dark, swinging a baseball.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Many small lights out of focus yellow and blue and of different intensities, at night on a blue background.",
-      "image_path": "480P_i2v_mixkit_images/Many small lights out of focus yellow and blue and of different intensities, at night on a blue back.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Pair of jazz musicians performing a song with their saxophone and trombone on an abandoned train.",
-      "image_path": "480P_i2v_mixkit_images/Pair of jazz musicians performing a song with their saxophone and trombone on an abandoned train..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Pair of mountain bikers climbing a slope in the forest covered with rocks, with abundant trees and vegetation around.",
-      "image_path": "480P_i2v_mixkit_images/Pair of mountain bikers climbing a slope in the forest covered with rocks, with abundant trees and v.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "A model with cyberpunk makeup poses in front of a neon light circle background.",
-      "image_path": "480P_i2v_mixkit_images/A model with cyberpunk makeup poses in front of a neon light circle background..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Couple of boy and girl walking hand in hand through a meadow among many cows that feed on grass and grass.",
-      "image_path": "480P_i2v_mixkit_images/Couple of boy and girl walking hand in hand through a meadow among many cows that feed on grass and .jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Natural coastal relief on a sunny day, with large hills covered with dry grass in an arid climate, in a panoramic aerial view.",
-      "image_path": "480P_i2v_mixkit_images/Natural coastal relief on a sunny day, with large hills covered with dry grass in an arid climate, i.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Young woman lying on her bed, while chatting on her cell phone, in a close-up shot of her hands typing on the phone.",
-      "image_path": "480P_i2v_mixkit_images/Young woman lying on her bed, while chatting on her cell phone, in a close-up shot of her hands typi.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Meadow with a frozen lake in a forest of pines covered in snow and fog.",
-      "image_path": "480P_i2v_mixkit_images/Meadow with a frozen lake in a forest of pines covered in snow and fog..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Aerial shot of buses and cars on the freeway with a cloudy sky in the background.",
-      "image_path": "480P_i2v_mixkit_images/Aerial shot of buses and cars on the freeway with a cloudy sky in the background..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Beautiful countryside meadow during summer with trees, grass and little yellow flowers in the ground on a sunny day.",
-      "image_path": "480P_i2v_mixkit_images/Beautiful countryside meadow during summer with trees, grass and little yellow flowers in the ground.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "A man wearing grey shorts jumps rope in a gym, weights and gym equipment in the background.",
-      "image_path": "480P_i2v_mixkit_images/A man wearing grey shorts jumps rope in a gym, weights and gym equipment in the background..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "A young woman wearing a black mini skirt, a red sleeveless shirt and a red Santa hat stands on a stage and sings into a microphone. In the background, a guitar player.",
-      "image_path": "480P_i2v_mixkit_images/A young woman wearing a black mini skirt, a red sleeveless shirt and a red Santa hat stands on a sta.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Little baby with a pacifier, playing with a teddy bear, accompanied by his mom, both sitting on a bed, in a close view.",
-      "image_path": "480P_i2v_mixkit_images/Little baby with a pacifier, playing with a teddy bear, accompanied by his mom, both sitting on a be.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "A saxophonist wearing a blazer dances while playing a song in a park.",
-      "image_path": "480P_i2v_mixkit_images/A saxophonist wearing a blazer dances while playing a song in a park..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "A young woman with short hair wearing pink sunglasses chews gum and makes a bubble gum with the city in the background.",
-      "image_path": "480P_i2v_mixkit_images/A young woman with short hair wearing pink sunglasses chews gum and makes a bubble gum with the city.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "White flowers from a bridal bouquet, with the background out of focus.",
-      "image_path": "480P_i2v_mixkit_images/White flowers from a bridal bouquet, with the background out of focus..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Silhouette of a person looking at an image projected in front, with spectra and light patterns rotating around a center in the background.",
-      "image_path": "480P_i2v_mixkit_images/Silhouette of a person looking at an image projected in front, with spectra and light patterns rotat.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Natural aerial landscape with a relief covered with abundant trees and vegetation and a thick layer of mist.",
-      "image_path": "480P_i2v_mixkit_images/Natural aerial landscape with a relief covered with abundant trees and vegetation and a thick layer .jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Boy dancing happily and cheerfully in a room set in the 80s.",
-      "image_path": "480P_i2v_mixkit_images/Boy dancing happily and cheerfully in a room set in the 80s..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Close up video of the legs of a young gymnast woman doing some acrobatics on a balance beam in a dark and empty gym.",
-      "image_path": "480P_i2v_mixkit_images/Close up video of the legs of a young gymnast woman doing some acrobatics on a balance beam in a dar.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Female scientist in a laboratory, with test tube, and more containers, looking through the lens of a microscope.",
-      "image_path": "480P_i2v_mixkit_images/Female scientist in a laboratory, with test tube, and more containers, looking through the lens of a.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Young woman getting makeup done by a professional makeup artist for a winter concept fashion photo shoot.",
-      "image_path": "480P_i2v_mixkit_images/Young woman getting makeup done by a professional makeup artist for a winter concept fashion photo s.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Loving couple sitting on a log on the shore of a lake outside, sharing an affectionate hug.",
-      "image_path": "480P_i2v_mixkit_images/Loving couple sitting on a log on the shore of a lake outside, sharing an affectionate hug..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Small family of father, mother and daughter, walking together through nature in a forest.",
-      "image_path": "480P_i2v_mixkit_images/Small family of father, mother and daughter, walking together through nature in a forest..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Flying over a peninsula covered in bushy trees, while discovering the sea around it, painted a beautiful turquoise blue, on a sunny day.",
-      "image_path": "480P_i2v_mixkit_images/Flying over a peninsula covered in bushy trees, while discovering the sea around it, painted a beaut.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "A woman with long blonde hair wearing a white t-shirt and black bikini bottoms stands on a wooden deck smiling, with the beach and palm trees in the background.",
-      "image_path": "480P_i2v_mixkit_images/A woman with long blonde hair wearing a white t-shirt and black bikini bottoms stands on a wooden de.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "View towards the horizon in a flat desert at night, as the full moon glows brightly red.",
-      "image_path": "480P_i2v_mixkit_images/View towards the horizon in a flat desert at night, as the full moon glows brightly red..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "100 dollar bill thrown in the street seen up close with the background out of focus, while a person arrives, picks it up and walks away.",
-      "image_path": "480P_i2v_mixkit_images/100 dollar bill thrown in the street seen up close with the background out of focus, while a person .jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Flying over a forest of abundant trees and vegetation, with a river and a dam, with houses and hills in the surroundings.",
-      "image_path": "480P_i2v_mixkit_images/Flying over a forest of abundant trees and vegetation, with a river and a dam, with houses and hills.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Artistic portrait of a man with his body and face stained with phosphorescent paint that glows with black light in a dark place.",
-      "image_path": "480P_i2v_mixkit_images/Artistic portrait of a man with his body and face stained with phosphorescent paint that glows with .jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Aerial view of the calm blue sea made with a drone showing the horizon and the sea.",
-      "image_path": "480P_i2v_mixkit_images/Aerial view of the calm blue sea made with a drone showing the horizon and the sea..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Cars driving on a two way street underneath a tree, with traffic signs, buildings and the sun in the background.",
-      "image_path": "480P_i2v_mixkit_images/Cars driving on a two way street underneath a tree, with traffic signs, buildings and the sun in the.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Face of a young woman, close up portrait",
-      "image_path": "480P_i2v_mixkit_images/Face of a young woman, close up portrait.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Man with his body splattered with fluorescent paint while doing a slow conceptual dance, under a black light in a dark place with colored spotlights around.",
-      "image_path": "480P_i2v_mixkit_images/Man with his body splattered with fluorescent paint while doing a slow conceptual dance, under a bla.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Two women dressed in colorful retro styles, having fun taking photos in a room with lots of pastel colors.",
-      "image_path": "480P_i2v_mixkit_images/Two women dressed in colorful retro styles, having fun taking photos in a room with lots of pastel c.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "The silhouette of a person who moves his head reflects inside a city with tall buildings next to a lake.",
-      "image_path": "480P_i2v_mixkit_images/The silhouette of a person who moves his head reflects inside a city with tall buildings next to a l.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "View from a car window of the trees lining the road speeding by, with the dazzling sunlight in the sky, in a retro looking video.",
-      "image_path": "480P_i2v_mixkit_images/View from a car window of the trees lining the road speeding by, with the dazzling sunlight in the s.jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Aerial view of two people on a tennis court playing tennis.",
-      "image_path": "480P_i2v_mixkit_images/Aerial view of two people on a tennis court playing tennis..jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "View from the floor upwards of a young woman with long dark hair dancing slowly in green and purple neon lights",
-      "image_path": "480P_i2v_mixkit_images/View from the floor upwards of a young woman with long dark hair dancing slowly in green and purple .jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    },
-    {
-      "prompt": "Some friends dancing and having fun together in circles, at a party surrounded by colored lights at a party, in a fancy old place, in a view from below them.",
-      "image_path": "480P_i2v_mixkit_images/Some friends dancing and having fun together in circles, at a party surrounded by colored lights at .jpg",
-      "num_inference_steps": 3,
-      "height": 448,
-      "width": 832,
-      "num_frames": 61
-    }
-  ]
-}
@@ -1,9 +1,7 @@
-FROM nvidia/cuda:12.8.0-devel-ubuntu22.04
+FROM nvidia/cuda:12.4.1-devel-ubuntu20.04

 ENV DEBIAN_FRONTEND=noninteractive

-SHELL ["/bin/bash", "-c"]
-
 WORKDIR /FastVideo

 RUN apt-get update && apt-get install -y --no-install-recommends \
@@ -11,25 +9,17 @@ RUN apt-get update && apt-get install -y --no-install-recommends \
    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
+RUN wget https://repo.anaconda.com/miniconda/Miniconda3-latest-Linux-x86_64.sh && \
+    bash Miniconda3-latest-Linux-x86_64.sh -b -p /opt/conda && \
+    rm Miniconda3-latest-Linux-x86_64.sh

-# Set CUDA environment variables
-ENV CUDA_HOME=/usr/local/cuda-12.8
-ENV PATH=${CUDA_HOME}/bin:${PATH}
-ENV LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
+ENV PATH=/opt/conda/bin:$PATH

-# Install uv and source its environment
-RUN curl -LsSf https://astral.sh/uv/install.sh | sh && \
-    echo 'source $HOME/.local/bin/env' >> /root/.bashrc
+RUN conda create --name fastvideo-dev python=3.10.0 -y
+
+SHELL ["/bin/bash", "-c"]

 # Copy just the pyproject.toml first to leverage Docker cache
 COPY pyproject.toml ./
@@ -37,36 +27,22 @@ COPY 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 && \
-    uv pip install --no-cache-dir .[dev] && \
-    uv pip install --no-cache-dir flash-attn==2.8.3 --no-build-isolation
+RUN conda run -n fastvideo-dev pip install --no-cache-dir --upgrade pip && \
+    conda run -n fastvideo-dev pip install --no-cache-dir .[dev] && \
+    conda run -n fastvideo-dev pip install --no-cache-dir flash-attn==2.7.4.post1 --no-build-isolation && \
+    conda clean -afy

 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 .[dev] && \
-    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
+RUN conda run -n fastvideo-dev pip install --no-cache-dir -e .[dev]

-# Install STA (Sliding Tile Attention)
-RUN source $HOME/.local/bin/env && \
-    source /opt/venv/bin/activate && \
-    cd csrc/attn/sliding_tile_attn && \
-    git submodule update --init --recursive && \
-    python setup.py install
+# Remove authentication headers
+RUN git config --unset-all http.https://github.com/.extraheader || true

-# 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
+# Set up automatic conda environment activation for all shells
+RUN echo 'source /opt/conda/etc/profile.d/conda.sh' >> /root/.bashrc && \
+    echo 'conda activate fastvideo-dev' >> /root/.bashrc && \
+    # Ensure .bashrc is sourced for SSH login shells
+    echo 'if [ -f ~/.bashrc ]; then . ~/.bashrc; fi' > /root/.profile

-EXPOSE 22
+EXPOSE 22
@@ -1,9 +1,7 @@
-FROM nvidia/cuda:12.8.0-devel-ubuntu22.04
+FROM nvidia/cuda:12.4.1-devel-ubuntu20.04

 ENV DEBIAN_FRONTEND=noninteractive

-SHELL ["/bin/bash", "-c"]
-
 WORKDIR /FastVideo

 RUN apt-get update && apt-get install -y --no-install-recommends \
@@ -11,25 +9,17 @@ RUN apt-get update && apt-get install -y --no-install-recommends \
    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
+RUN wget https://repo.anaconda.com/miniconda/Miniconda3-latest-Linux-x86_64.sh && \
+    bash Miniconda3-latest-Linux-x86_64.sh -b -p /opt/conda && \
+    rm Miniconda3-latest-Linux-x86_64.sh

-# Set CUDA environment variables
-ENV CUDA_HOME=/usr/local/cuda-12.8
-ENV PATH=${CUDA_HOME}/bin:${PATH}
-ENV LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
+ENV PATH=/opt/conda/bin:$PATH

-# Install uv and source its environment
-RUN curl -LsSf https://astral.sh/uv/install.sh | sh && \
-    echo 'source $HOME/.local/bin/env' >> /root/.bashrc
+RUN conda create --name fastvideo-dev python=3.11.11 -y
+
+SHELL ["/bin/bash", "-c"]

 # Copy just the pyproject.toml first to leverage Docker cache
 COPY pyproject.toml ./
@@ -37,36 +27,22 @@ COPY 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.11 --seed /opt/venv && \
-    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 flash-attn==2.8.3 --no-build-isolation
+RUN conda run -n fastvideo-dev pip install --no-cache-dir --upgrade pip && \
+    conda run -n fastvideo-dev pip install --no-cache-dir .[dev] && \
+    conda run -n fastvideo-dev pip install --no-cache-dir flash-attn==2.7.4.post1 --no-build-isolation && \
+    conda clean -afy

 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 .[dev] && \
-    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
+RUN conda run -n fastvideo-dev pip install --no-cache-dir -e .[dev]

-# Install STA (Sliding Tile Attention)
-RUN source $HOME/.local/bin/env && \
-    source /opt/venv/bin/activate && \
-    cd csrc/attn/sliding_tile_attn && \
-    git submodule update --init --recursive && \
-    python setup.py install
+# Remove authentication headers
+RUN git config --unset-all http.https://github.com/.extraheader || true

-# 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
+# Set up automatic conda environment activation for all shells
+RUN echo 'source /opt/conda/etc/profile.d/conda.sh' >> /root/.bashrc && \
+    echo 'conda activate fastvideo-dev' >> /root/.bashrc && \
+    # Ensure .bashrc is sourced for SSH login shells
+    echo 'if [ -f ~/.bashrc ]; then . ~/.bashrc; fi' > /root/.profile

-EXPOSE 22
+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 flash-attn==2.8.3 --no-build-isolation
+    uv pip install --no-cache-dir flash-attn==2.8.0.post2 --no-build-isolation

 COPY . .

@@ -58,15 +58,15 @@ RUN source $HOME/.local/bin/env && \
 # Install STA (Sliding Tile Attention)
 RUN source $HOME/.local/bin/env && \
    source /opt/venv/bin/activate && \
-    cd csrc/attn/sliding_tile_attn && \
+    cd csrc/attn && \
    git submodule update --init --recursive && \
-    python setup.py install
+    python setup_sta.py install

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

-EXPOSE 22
+EXPOSE 22
@@ -1,72 +0,0 @@
-FROM nvidia/cuda:12.9.1-cudnn-devel-ubuntu22.04
-
-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
-
-# Set CUDA environment variables
-ENV CUDA_HOME=/usr/local/cuda-12.9
-ENV PATH=${CUDA_HOME}/bin:${PATH}
-ENV LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
-
-# 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.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.12 --seed /opt/venv && \
-    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 flash-attn==2.8.3 --no-build-isolation
-
-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 .[dev] && \
-    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 STA (Sliding Tile Attention)
-RUN source $HOME/.local/bin/env && \
-    source /opt/venv/bin/activate && \
-    cd csrc/attn/sliding_tile_attn && \
-    git submodule update --init --recursive && \
-    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
@@ -96,7 +96,8 @@ copybutton_prompt_is_regexp = True
 #
 html_title = project
 html_theme = 'sphinx_book_theme'
-html_logo = '../../assets/logos/icon_simple.svg'
+html_logo = '../../assets/logo.jpg'
+#html_favicon = 'assets/logos/vllm-logo-only-light.ico'
 html_theme_options = {
    'path_to_docs': 'docs/source',
    'repository_url': 'https://github.com/hao-ai-lab/FastVideo/',
@@ -3,12 +3,12 @@

 If you prefer a containerized development environment or want to avoid managing dependencies manually, you can use our prebuilt Docker image:

-**Images:** [`ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev:py3.12-latest`](https://ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev)
+**Image:** [`ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev:latest`](https://ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev)

 ## Starting the container

 ```bash
-docker run --gpus all -it ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev:py3.12-latest
+docker run --gpus all -it ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev:latest
 ```

 This will:
@@ -6,7 +6,7 @@ You can easily use the FastVideo Docker image as a custom container on [RunPod](

 ## Creating a new pod

-Choose a GPU that supports CUDA 12.8
+Choose a GPU that supports CUDA 12.4

 Pick 1 or 2 L40S GPU(s)

@@ -22,20 +22,10 @@ source ~/.bashrc
 Create and activate a Conda environment for FastVideo:

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

-Install `uv` (optional, but recommended):
-
-From instructions on [uv](https://astral.sh/uv/):
-
-```
-curl -LsSf https://astral.sh/uv/install.sh | sh
-# or 
-wget -qO- https://astral.sh/uv/install.sh | sh
-```
-
 Clone the FastVideo repository and go to the FastVideo directory:

 ```
@@ -46,10 +36,10 @@ git clone https://github.com/hao-ai-lab/FastVideo.git && cd FastVideo
 Now you can install FastVideo and setup git hooks for running linting. By using `pre-commit`, the linters will run and have to pass before you'll be able to make a commit.

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

 # Can also install flash-attn (optional)
-uv pip install flash-attn --no-build-isolation 
+pip install flash-attn==2.7.4.post1 --no-build-isolation 

 # Linting, formatting and static type checking
 pre-commit install --hook-type pre-commit --hook-type commit-msg
@@ -60,14 +50,3 @@ pre-commit run --all-files
 # Unit tests
 pytest tests/
 ```
-
-If you are on a Hopper GPU, you should also install [FA3](https://github.com/Dao-AILab/flash-attention) for much better performance:
-
-```
-git clone https://github.com/Dao-AILab/flash-attention.git && cd flash-attention/hopper
-
-# make sure you have ninja installed
-uv pip install ninja
-
-python setup.py install
-```
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
SolitaryThinker	083c2258ea	fix cm	2025-08-21 05:18:30 +00:00
SolitaryThinker	60260a6be9	fix preprocessing before vae deocde	2025-08-20 01:10:53 +00:00
SolitaryThinker	06cc80ee93	init lpips	2025-08-19 23:48:06 +00:00
Matthew Noto	c6b3475f77	LPIPs fix	2025-08-19 16:40:42 -07:00
Matthew Noto	1d34380c58	LPIPS	2025-08-19 15:40:29 -07:00
SolitaryThinker	ddf8198e65	remove validation back sim; reg loss for back sim	2025-08-19 07:19:40 +00:00
SolitaryThinker	c38ee88294	cm	2025-08-18 23:22:37 +00:00
SolitaryThinker	ff1b77bdd6	update	2025-08-18 21:44:47 +00:00
SolitaryThinker	22ba3759ee	cm	2025-08-18 03:14:18 +00:00
SolitaryThinker	2c3a3893c0	cm	2025-08-18 02:46:24 +00:00
SolitaryThinker	6e9ba7ba9d	sim interval arg	2025-08-17 02:49:03 +00:00
SolitaryThinker	64f1295878	fix	2025-08-15 03:45:32 +00:00
SolitaryThinker	d9cc3b7e42	reg loss	2025-08-15 02:08:15 +00:00
SolitaryThinker	cb3230f5e9	reg loss	2025-08-15 01:30:46 +00:00
SolitaryThinker	1e91cec93d	fix	2025-08-13 07:05:51 +00:00
SolitaryThinker	5f0688084c	update	2025-08-13 07:05:24 +00:00
SolitaryThinker	f3023db253	fix	2025-08-13 07:01:18 +00:00
SolitaryThinker	c405941f21	fix	2025-08-13 04:05:01 +00:00
SolitaryThinker	f4e818305b	fix	2025-08-13 03:11:22 +00:00
SolitaryThinker	4c4ab65020	fix	2025-08-13 03:10:12 +00:00
SolitaryThinker	6b8bf495ee	fix	2025-08-13 02:19:52 +00:00
SolitaryThinker	ebd23ba8e7	add orig val	2025-08-12 20:00:33 +00:00
SolitaryThinker	0c1cb0c19d	runnable	2025-08-12 19:58:36 +00:00
SolitaryThinker	d87ebe051b	fix	2025-08-11 03:42:08 +00:00
SolitaryThinker	0ae669cd4c	fix	2025-08-11 03:31:02 +00:00
SolitaryThinker	63c124fec8	update	2025-08-10 23:37:01 +00:00
SolitaryThinker	ececbe1839	wip	2025-08-10 08:33:03 +00:00