Remove explicit_package_bases from mypy settings

Remove explicit_package_bases setting from mypy configuration

.buildkite/pipeline.yml

@@ -2,20 +2,218 @@ env:
   IMAGE_VERSION: "py3.12-latest"
   BUILDKITE_CLEAN_CHECKOUT: true
 
+notify:
+  - github_commit_status:
+      context: "fastcheck-passed"
+    if: build.env("TEST_SCOPE") == "fastcheck" || build.env("TEST_SCOPE") == null
+  - github_commit_status:
+      context: "full-suite-passed"
+    if: build.env("TEST_SCOPE") == "full"
+  - github_commit_status:
+      context: "direct-test-completed"
+    if: build.env("TEST_SCOPE") == "direct"
+
 steps:
-  - label: "pre-commit"
-    command: ".buildkite/scripts/pre_commit.sh"
-    agents:
-      queue: "default"
+  # ============================================================
+  # Direct test: triggered by /test <name> slash command.
+  # Labels match fastcheck/full-suite counterparts so the GitHub
+  # check status overwrites the original failed check.
+  # Only ONE step executes per build (gated by TEST_TYPE).
+  # ============================================================
 
-  - wait
+  # --- Fastcheck-scope direct tests ---
+    - label: ":microscope: Encoder Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "encoder"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":microscope: VAE Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "vae"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":microscope: Transformer Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "transformer"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":microscope: Kernel Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "kernel_tests"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":microscope: Unit Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "unit_test"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
 
-  - label: "Trigger Tests"
-    plugins:
-      - monorepo-diff#v1.4.0:
-          diff: 'git fetch origin "$BUILDKITE_PULL_REQUEST_BASE_BRANCH" && git diff --name-only origin/"$BUILDKITE_PULL_REQUEST_BASE_BRANCH"...HEAD'
-          watch:
-            - path: 
+  # --- Full-suite-scope direct tests ---
+    - label: ":bar_chart: SSIM Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "ssim"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+          - exit_status: 1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: LoRA Inference Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "inference_lora"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: Training Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "training"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: Distillation DMD Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "distillation_dmd"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: Self-Forcing Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "self_forcing"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: LoRA Training Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "training_lora"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+          - exit_status: 1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: Training Tests VSA"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "training_vsa"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+          - exit_status: 1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: Inference Tests VMoBA"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "inference_vmoba"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: Performance Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "performance"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+    - label: ":test_tube: API Server Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "api_server"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
+
+  # ============================================================
+  # Fastcheck: Runs on every PR (~10-15 min parallel)
+  # Core component validation: encoders, VAEs, transformers,
+  # CUDA kernels, and unit tests.
+  # ============================================================
+    - label: "Trigger Fastcheck"
+      if: build.env("TEST_SCOPE") == "fastcheck" || build.env("TEST_SCOPE") == null
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      plugins:
+        - monorepo-diff#v1.4.0:
+            diff: 'git fetch origin "${BUILDKITE_PULL_REQUEST_BASE_BRANCH:-main}" && git diff --name-only "origin/${BUILDKITE_PULL_REQUEST_BASE_BRANCH:-main}...HEAD"'
+            watch:
+            - path:
                 - "fastvideo/models/encoders/**"
                 - "fastvideo/models/loader/**"
                 - "fastvideo/tests/encoders/**"
@@ -23,7 +221,7 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 20m .buildkite/scripts/pr_test.sh"
-                label: "Encoder Tests"
+                label: ":microscope: Encoder Tests"
                 env:
                   - TEST_TYPE=encoder
                 agents:
@@ -36,7 +234,7 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 20m .buildkite/scripts/pr_test.sh"
-                label: "VAE Tests"
+                label: ":microscope: VAE Tests"
                 env:
                   - TEST_TYPE=vae
                 agents:
@@ -51,20 +249,68 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Transformer Tests"
+                label: ":microscope: Transformer Tests"
                 env:
                   - TEST_TYPE=transformer
                 agents:
                   queue: "default"
-            - path: 
+            - path:
+                - "fastvideo-kernel/**"
+                - "pyproject.toml"
+                - "docker/Dockerfile.python3.12"
+              config:
+                command: "timeout 15m .buildkite/scripts/pr_test.sh"
+                label: ":microscope: Kernel Tests"
+                env:
+                  - TEST_TYPE=kernel_tests
+                agents:
+                  queue: "default"
+            - path:
+                - "fastvideo/**"
+                - ".buildkite/**"
+                - ".github/**"
+                - "pyproject.toml"
+                - "docker/Dockerfile.python3.12"
+              config:
+                command: "timeout 15m .buildkite/scripts/pr_test.sh"
+                label: ":microscope: Unit Tests"
+                env:
+                  - TEST_TYPE=unit_test
+                agents:
+                  queue: "default"
+
+  # ============================================================
+  # Full Suite: Runs when TEST_SCOPE=full
+  # Triggered by adding the 'ready' label (via ci-trigger-full-suite.yml)
+  # or on-demand via /test full slash command.
+  # Includes integration tests, SSIM regression, training pipelines,
+  # and performance benchmarks.
+  # ============================================================
+    - label: "Trigger Full Suite"
+      if: build.env("TEST_SCOPE") == "full"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      plugins:
+        - monorepo-diff#v1.4.0:
+            diff: 'git fetch origin "${BUILDKITE_PULL_REQUEST_BASE_BRANCH:-main}" && git diff --name-only "origin/${BUILDKITE_PULL_REQUEST_BASE_BRANCH:-main}...HEAD"'
+            watch:
+            - path:
                 - "fastvideo/**/*.py"
                 - "pyproject.toml"
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 90m .buildkite/scripts/pr_test.sh"
-                label: "SSIM Tests"
+                label: ":bar_chart: SSIM Tests"
                 env:
                   - TEST_TYPE=ssim
+                retry:
+                  automatic:
+                    - exit_status: 1
+                      limit: 2
                 agents:
                   queue: "default"
             - path:
@@ -77,7 +323,7 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 20m .buildkite/scripts/pr_test.sh"
-                label: "LoRA Inference Tests"
+                label: ":test_tube: LoRA Inference Tests"
                 env:
                   - TEST_TYPE=inference_lora
                 agents:
@@ -88,7 +334,7 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Training Tests"
+                label: ":test_tube: Training Tests"
                 env:
                   - TEST_TYPE=training
                 agents:
@@ -99,7 +345,7 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Distillation DMDTests"
+                label: ":test_tube: Distillation DMD Tests"
                 env:
                   - TEST_TYPE=distillation_dmd
                 agents:
@@ -111,7 +357,7 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Self-Forcing Tests"
+                label: ":test_tube: Self-Forcing Tests"
                 env:
                   - TEST_TYPE=self_forcing
                 agents:
@@ -122,9 +368,13 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "LoRA Training Tests"
+                label: ":test_tube: LoRA Training Tests"
                 env:
                   - TEST_TYPE=training_lora
+                retry:
+                  automatic:
+                    - exit_status: 1
+                      limit: 2
                 agents:
                   queue: "default"
             - path:
@@ -134,20 +384,15 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Training Tests VSA"
+                label: ":test_tube: Training Tests VSA"
                 env:
                   - TEST_TYPE=training_vsa
+                retry:
+                  automatic:
+                    - exit_status: 1
+                      limit: 2
                 agents:
                   queue: "default"
-            - path:
-                - "fastvideo-kernel/**"
-                - "pyproject.toml"
-                - "docker/Dockerfile.python3.12"
-              config:
-                command: "timeout 15m .buildkite/scripts/pr_test.sh"
-                label: "Kernel Tests"
-                env:
-                  - TEST_TYPE=kernel_tests
             - path:
                 - "fastvideo-kernel/**"
                 - "fastvideo/attention/backends/vmoba.py"
@@ -155,20 +400,9 @@ steps:
                 - "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"
+                label: ":test_tube: Inference Tests VMoBA"
                 env:
-                  - TEST_TYPE=unit_test
+                  - TEST_TYPE=inference_vmoba
                 agents:
                   queue: "default"
             - path:
@@ -184,7 +418,7 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 30m .buildkite/scripts/pr_test.sh"
-                label: "Performance Tests"
+                label: ":test_tube: Performance Tests"
                 env:
                   - TEST_TYPE=performance
                 agents:
@@ -197,19 +431,8 @@ steps:
                 - "docker/Dockerfile.python3.12"
               config:
                 command: "timeout 30m .buildkite/scripts/pr_test.sh"
-                label: "API Server Tests"
+                label: ":test_tube: API Server Tests"
                 env:
                   - TEST_TYPE=api_server
                 agents:
                   queue: "default"
-            # - path:
-            #     - "scripts/lora_extraction/**"
-            #     - "pyproject.toml"
-            #     - "docker/Dockerfile.python3.12"
-            #   config:
-            #     command: "timeout 90m .buildkite/scripts/pr_test.sh"
-            #     label: "LoRA Extraction Tests"
-            #     env:
-            #       - TEST_TYPE=lora_extraction
-            #     agents:
-            #       queue: "default"

.buildkite/scripts/pr_test.sh

@@ -59,7 +59,11 @@ if [ -z "${TEST_TYPE:-}" ]; then
 fi
 log "Test type: $TEST_TYPE"
 
-MODAL_ENV="BUILDKITE_REPO=$BUILDKITE_REPO BUILDKITE_COMMIT=$BUILDKITE_COMMIT BUILDKITE_PULL_REQUEST=$BUILDKITE_PULL_REQUEST IMAGE_VERSION=$IMAGE_VERSION"
+EFFECTIVE_PR=${BUILDKITE_PULL_REQUEST:-false}
+if [ "$EFFECTIVE_PR" = "false" ] && [ -n "${PR_NUMBER:-}" ]; then
+    EFFECTIVE_PR=$PR_NUMBER
+fi
+MODAL_ENV="BUILDKITE_REPO=$BUILDKITE_REPO BUILDKITE_COMMIT=$BUILDKITE_COMMIT BUILDKITE_PULL_REQUEST=$EFFECTIVE_PR IMAGE_VERSION=$IMAGE_VERSION"
 
 case "$TEST_TYPE" in
     "encoder")

.github/PULL_REQUEST_TEMPLATE.md

@@ -1,3 +1,22 @@
+<!--
+PR TITLE: Must start with a type tag, e.g.:
+  [feat] Add new model       [bugfix] Fix VAE tiling      [refactor] Restructure pipeline
+  [perf] Optimize kernel     [ci] Update tests             [docs] Add guide
+  [misc] Cleanup configs     [new-model] Port Flux2
+
+MERGE WORKFLOW:
+  1. Ensure pre-commit passes and you have at least 1 approval
+  2. Comment /merge (or add the "ready" label) to enter the Merge Queue
+  3. Full Test Suite runs automatically on a staging branch → auto-merge on success
+
+ON-DEMAND TESTING (write access required):
+  /test full       — Full Test Suite        /test ssim        — SSIM regression
+  /test training   — Training pipeline      /test encoder     — Encoder tests
+  /test transformer — Transformer tests     /test vae         — VAE tests
+  /test kernel     — CUDA kernel tests      /test unit        — Unit tests
+  See docs/contributing/pull_requests.md for all 17 test commands
+-->
+
 ## Purpose
 
 <!-- What does this PR do? Link the related issue if applicable. -->

.github/mergify.yml

@@ -0,0 +1,316 @@
+merge_protections:
+  - name: PR merge requirements
+    if:
+      - base = main
+    success_conditions:
+      - "title~=(?i)^\\[(feat|feature|bugfix|fix|refactor|perf|ci|doc|docs|misc|chore|kernel|new.?model)\\]"
+      - "#approved-reviews-by>=1"
+      - check-success~=pre-commit
+      - check-success=fastcheck-passed
+      - check-success=full-suite-passed
+
+pull_request_rules:
+
+  # ============================================================
+  # Type labels (from PR title prefix)
+  # ============================================================
+
+  - name: "label type: feat"
+    conditions:
+      - "title~=(?i)^\\[(feat|feature)\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: feat"]
+
+  - name: "label type: bugfix"
+    conditions:
+      - "title~=(?i)^\\[(bug)?fix\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: bugfix"]
+
+  - name: "label type: refactor"
+    conditions:
+      - "title~=(?i)^\\[refactor\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: refactor"]
+
+  - name: "label type: perf"
+    conditions:
+      - "title~=(?i)^\\[perf\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: perf"]
+
+  - name: "label type: ci"
+    conditions:
+      - "title~=(?i)^\\[ci\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: ci"]
+
+  - name: "label type: docs"
+    conditions:
+      - "title~=(?i)^\\[(doc|docs)\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: docs"]
+
+  - name: "label type: misc"
+    conditions:
+      - "title~=(?i)^\\[(misc|chore)\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: misc"]
+
+  - name: "label type: new-model"
+    conditions:
+      - "title~=(?i)^\\[new.?model\\]"
+      - -closed
+    actions:
+      label:
+        add: ["type: new-model"]
+
+  # ============================================================
+  # Scope labels (from changed files)
+  # ============================================================
+
+  - name: "label scope: training"
+    conditions:
+      - or:
+          - files~=^fastvideo/train/
+          - files~=^fastvideo/training/
+          - files~=^fastvideo/distillation/
+          - files~=^examples/train/
+          - files~=^examples/training/
+          - files~=^examples/distill/
+      - -closed
+    actions:
+      label:
+        add: ["scope: training"]
+
+  - name: "label scope: inference"
+    conditions:
+      - or:
+          - files~=^fastvideo/pipelines/basic/
+          - files~=^fastvideo/pipelines/stages/
+          - files~=^fastvideo/pipelines/samplers/
+          - files~=^fastvideo/entrypoints/
+          - files~=^fastvideo/worker/
+          - files~=^fastvideo/configs/sample/
+          - files~=^fastvideo/configs/pipelines/
+          - files~=^examples/inference/
+      - -closed
+    actions:
+      label:
+        add: ["scope: inference"]
+
+  - name: "label scope: attention"
+    conditions:
+      - files~=^fastvideo/attention/
+      - -closed
+    actions:
+      label:
+        add: ["scope: attention"]
+
+  - name: "label scope: kernel"
+    conditions:
+      - or:
+          - files~=^fastvideo-kernel/
+          - files~=^csrc/
+      - -closed
+    actions:
+      label:
+        add: ["scope: kernel"]
+
+  - name: "label scope: data"
+    conditions:
+      - or:
+          - files~=^fastvideo/dataset/
+          - files~=^fastvideo/pipelines/preprocess/
+          - files~=^examples/preprocessing/
+      - -closed
+    actions:
+      label:
+        add: ["scope: data"]
+
+  - name: "label scope: infra"
+    conditions:
+      - or:
+          - files~=^\.github/
+          - files~=^\.buildkite/
+          - files~=^fastvideo/tests/
+          - files~=^docker/
+      - -closed
+    actions:
+      label:
+        add: ["scope: infra"]
+
+  - name: "label scope: distributed"
+    conditions:
+      - files~=^fastvideo/distributed/
+      - -closed
+    actions:
+      label:
+        add: ["scope: distributed"]
+
+  - name: "label scope: docs"
+    conditions:
+      - files~=^docs/
+      - -closed
+    actions:
+      label:
+        add: ["scope: docs"]
+
+  - name: "label scope: ui"
+    conditions:
+      - files~=^ui/
+      - -closed
+    actions:
+      label:
+        add: ["scope: ui"]
+
+  - name: "label scope: model"
+    conditions:
+      - or:
+          - files~=^fastvideo/models/
+          - files~=^fastvideo/layers/
+          - files~=^fastvideo/configs/models/
+      - -closed
+    actions:
+      label:
+        add: ["scope: model"]
+
+  # ============================================================
+  # Pre-commit failure help comment
+  # ============================================================
+
+  - name: comment on pre-commit failure
+    conditions:
+      - check-failure~=pre-commit
+      - -closed
+    actions:
+      comment:
+        message: |
+          ## Pre-commit checks failed
+
+          Hi @{{author}}, the pre-commit checks have failed. To fix them locally:
+
+          ```bash
+          # Install pre-commit if you haven't already
+          uv pip install pre-commit
+          pre-commit install
+
+          # Run all checks and auto-fix what's possible
+          pre-commit run --all-files
+          ```
+
+          Common fixes:
+          - **yapf**: `yapf -i <file>` (formatting)
+          - **ruff**: `ruff check --fix <file>` (linting)
+          - **codespell**: `codespell --write-changes <file>` (spelling)
+
+          After fixing, commit and push the changes. The checks will re-run automatically.
+
+          For future commits, `pre-commit` will run automatically on changed files before each commit.
+
+
+  # ============================================================
+  # Merge conflict detection
+  # ============================================================
+
+  - name: label conflicting PRs
+    conditions:
+      - conflict
+      - -closed
+      - label!=stale
+    actions:
+      label:
+        add: [needs-rebase]
+      comment:
+        message: |
+          This PR has merge conflicts with the base branch. Please rebase:
+
+          ```bash
+          git fetch origin main
+          git rebase origin/main
+          # Resolve any conflicts, then:
+          git push --force-with-lease
+          ```
+
+  - name: remove conflict label when resolved
+    conditions:
+      - -conflict
+      - -closed
+      - label=needs-rebase
+    actions:
+      label:
+        remove: [needs-rebase]
+
+  # ============================================================
+  # Auto-merge and auto-rebase
+  # ============================================================
+
+  - name: auto-merge when ready and all checks pass
+    conditions:
+      - label=ready
+      - "title~=(?i)^\\[(feat|feature|bugfix|fix|refactor|perf|ci|doc|docs|misc|chore|kernel|new.?model)\\]"
+      - "#approved-reviews-by>=1"
+      - check-success~=pre-commit
+      - check-success=fastcheck-passed
+      - check-success=full-suite-passed
+      - -conflict
+      - -closed
+      - -draft
+    actions:
+      merge:
+        method: squash
+
+  - name: auto-update when ready
+    conditions:
+      - label=ready
+      - "#approved-reviews-by>=1"
+      - -conflict
+      - -closed
+      - -draft
+    actions:
+      update: {}
+
+  # ============================================================
+  # PR title format help
+  # ============================================================
+
+  - name: comment on invalid PR title format
+    conditions:
+      - -closed
+      - -draft
+      - "-title~=(?i)^\\[(feat|feature|bugfix|fix|refactor|perf|ci|doc|docs|misc|chore|kernel|new.?model)\\]"
+    actions:
+      comment:
+        message: |
+          ## ⚠️ PR title format required
+
+          Your PR title must start with a type tag in brackets. Examples:
+          - `[feat] Add new model support`
+          - `[bugfix] Fix VAE tiling corruption`
+          - `[refactor] Restructure training pipeline`
+          - `[perf] Optimize attention kernel`
+          - `[ci] Update test infrastructure`
+          - `[docs] Add inference guide`
+          - `[misc] Clean up configs`
+          - `[new-model] Port Flux2 to FastVideo`
+
+          Valid tags: `feat`, `feature`, `bugfix`, `fix`, `refactor`, `perf`, `ci`, `doc`, `docs`, `misc`, `chore`, `kernel`, `new-model`
+
+          Please update your PR title and the merge protection check will pass automatically.
+
+merge_protections_settings:
+  reporting_method: check-runs

.github/scripts/runpod_api.py

@@ -1,249 +0,0 @@
-import argparse
-import json
-import os
-import subprocess
-import sys
-import time
-
-import requests
-
-
-def parse_arguments():
-    """Parse command line arguments"""
-    parser = argparse.ArgumentParser(description='Run tests on RunPod GPU')
-    parser.add_argument('--gpu-type', type=str, help='GPU type to use')
-    parser.add_argument('--gpu-count',
-                        type=int,
-                        help='Number of GPUs to use',
-                        default=1)
-    parser.add_argument('--test-command', type=str, help='Test command to run')
-    parser.add_argument('--disk-size',
-                        type=int,
-                        default=20,
-                        help='Container disk size in GB (default: 20)')
-    parser.add_argument('--volume-size',
-                        type=int,
-                        default=20,
-                        help='Persistent volume size in GB (default: 20)')
-    parser.add_argument(
-        '--image',
-        type=str,
-        required=True,
-        help='Docker image to use')
-    return parser.parse_args()
-
-
-args = parse_arguments()
-API_KEY = os.environ['RUNPOD_API_KEY']
-RUN_ID = os.environ['GITHUB_RUN_ID']
-JOB_ID = os.environ['JOB_ID']
-PODS_API = "https://rest.runpod.io/v1/pods"
-HEADERS = {
-    "Content-Type": "application/json",
-    "Authorization": f"Bearer {API_KEY}"
-}
-
-
-def create_pod():
-    """Create a RunPod instance"""
-    # Ensure image name is lowercase (Docker requirement)
-    image_name = args.image.lower()
-    print(f"Using specified image: {image_name}")
-
-    docker_start_cmd = [
-        "bash", 
-        "-c", 
-        "apt update;DEBIAN_FRONTEND=noninteractive apt-get install openssh-server -y;mkdir -p ~/.ssh;cd $_;chmod 700 ~/.ssh;echo \"$PUBLIC_KEY\" >> authorized_keys;chmod 700 authorized_keys;service ssh start;sleep infinity"
-    ]
-    
-    print(f"Creating RunPod instance with GPU: {args.gpu_type}...")
-    payload = {
-        "name": f"fastvideo-{JOB_ID}-{RUN_ID}",
-        "containerDiskInGb": args.disk_size,
-        "volumeInGb": args.volume_size,
-        "gpuTypeIds": [args.gpu_type],
-        "gpuCount": args.gpu_count,
-        "imageName": image_name,
-        "allowedCudaVersions": ["12.4"],
-        "dockerStartCmd": docker_start_cmd
-    }
-
-    response = requests.post(PODS_API, headers=HEADERS, json=payload)
-    response_data = response.json()
-    print(f"Response: {json.dumps(response_data, indent=2)}")
-
-    return response_data["id"]
-
-
-def wait_for_pod(pod_id):
-    """Wait for pod to be in RUNNING state and fully ready with SSH access"""
-    print("Waiting for RunPod to be ready...")
-
-    # First wait for RUNNING status
-    max_attempts = 10
-    attempts = 0
-    while attempts < max_attempts:
-        response = requests.get(f"{PODS_API}/{pod_id}", headers=HEADERS)
-        pod_data = response.json()
-        status = pod_data["desiredStatus"]
-
-        if status == "RUNNING":
-            print("RunPod is running! Now waiting for ports to be assigned...")
-            break
-
-        print(
-            f"Current status: {status}, waiting... (attempt {attempts+1}/{max_attempts})"
-        )
-        time.sleep(2)
-        attempts += 1
-
-    if attempts >= max_attempts:
-        raise TimeoutError(
-            "Timed out waiting for RunPod to reach RUNNING state")
-
-    # Wait for ports to be assigned
-    max_attempts = 50
-    attempts = 0
-    while attempts < max_attempts:
-        response = requests.get(f"{PODS_API}/{pod_id}", headers=HEADERS)
-        pod_data = response.json()
-        port_mappings = pod_data.get("portMappings")
-
-        if (port_mappings is not None and "22" in port_mappings
-                and pod_data.get("publicIp", "") != ""):
-            print("RunPod is ready with SSH access!")
-            print(f"SSH IP: {pod_data['publicIp']}")
-            print(f"SSH Port: {port_mappings['22']}")
-            break
-
-        print(
-            f"Waiting for SSH port and public IP to be available... (attempt {attempts+1}/{max_attempts})"
-        )
-        time.sleep(20)
-        attempts += 1
-
-    if attempts >= max_attempts:
-        raise TimeoutError("Timed out waiting for RunPod SSH access")
-
-
-def execute_command(pod_id):
-    """Execute command on the pod via SSH using system SSH client"""
-    print(f"Running command: {args.test_command}")
-
-    response = requests.get(f"{PODS_API}/{pod_id}", headers=HEADERS)
-    pod_data = response.json()
-    ssh_ip = pod_data["publicIp"]
-    ssh_port = pod_data["portMappings"]["22"]
-
-    # Copy the repository to the pod using scp
-    repo_dir = os.path.abspath(os.getcwd())
-    repo_name = os.path.basename(repo_dir)
-
-    print(f"Copying repository from {repo_dir} to RunPod...")
-
-    tar_command = [
-        "tar", "-czf", "/tmp/repo.tar.gz", "-C",
-        os.path.dirname(repo_dir), repo_name
-    ]
-    subprocess.run(tar_command, check=True)
-
-    # Copy the tarball to the pod
-    scp_command = [
-        "scp", "-o", "StrictHostKeyChecking=no", "-o",
-        "UserKnownHostsFile=/dev/null", "-o", "ServerAliveInterval=60", "-o",
-        "ServerAliveCountMax=10", "-P",
-        str(ssh_port), "/tmp/repo.tar.gz", f"root@{ssh_ip}:/tmp/"
-    ]
-    subprocess.run(scp_command, check=True)
-
-    # For custom image, we can use the pre-configured environment
-    setup_steps = [
-        "tar -xzf /tmp/repo.tar.gz --no-same-owner -C /workspace/",
-        f"cd /workspace/{repo_name}",
-        "source $HOME/.local/bin/env && source /opt/venv/bin/activate",
-        args.test_command
-    ]
-    
-    remote_command = " && ".join(setup_steps)
-
-    ssh_command = [
-        "ssh", "-o", "StrictHostKeyChecking=no", "-o",
-        "UserKnownHostsFile=/dev/null", "-o", "ServerAliveInterval=60", "-o",
-        "ServerAliveCountMax=10", "-p",
-        str(ssh_port), f"root@{ssh_ip}", remote_command
-    ]
-
-    print(f"Connecting to {ssh_ip}:{ssh_port}...")
-
-    try:
-        process = subprocess.Popen(ssh_command,
-                                   stdout=subprocess.PIPE,
-                                   stderr=subprocess.STDOUT,
-                                   universal_newlines=True,
-                                   bufsize=0)
-
-        stdout_lines = []
-
-        print("Command output:")
-
-        for line in iter(process.stdout.readline, ''):
-            print(line.strip())
-            stdout_lines.append(line)
-
-        process.wait()
-
-        return_code = process.returncode
-        success = return_code == 0
-
-        stdout_str = "".join(stdout_lines)
-
-        if success:
-            print("Command executed successfully")
-        else:
-            print(f"Command failed with exit code {return_code}")
-
-        result = {
-            "success": success,
-            "return_code": return_code,
-            "stdout": stdout_str,
-            "stderr": ""
-        }
-        return result
-
-    except Exception as e:
-        print(f"Error executing SSH command: {str(e)}")
-        result = {"success": False, "error": str(e), "stdout": "", "stderr": ""}
-        return result
-
-
-def terminate_pod(pod_id):
-    """Terminate the pod"""
-    print("Terminating RunPod...")
-    requests.delete(f"{PODS_API}/{pod_id}", headers=HEADERS)
-    print(f"Terminated pod {pod_id}")
-
-
-def main():
-    pod_id = None
-    try:
-        pod_id = create_pod()
-        wait_for_pod(pod_id)
-        result = execute_command(pod_id)
-
-        if result.get("error") is not None:
-            print(f"Error executing command: {result['error']}")
-            sys.exit(1)
-
-        if not result.get("success", False):
-            print(
-                "Tests failed - check the output above for details on which tests failed"
-            )
-            sys.exit(1)
-
-    finally:
-        if pod_id:
-            terminate_pod(pod_id)
-
-
-if __name__ == "__main__":
-    main()

.github/scripts/runpod_cleanup.py

@@ -1,90 +0,0 @@
-import json
-import os
-import sys
-import uuid
-
-import requests
-
-API_KEY = os.environ['RUNPOD_API_KEY']
-RUN_ID = os.environ.get('GITHUB_RUN_ID', str(uuid.uuid4()))
-PODS_API = "https://rest.runpod.io/v1/pods"
-HEADERS = {
-    "Content-Type": "application/json",
-    "Authorization": f"Bearer {API_KEY}"
-}
-
-
-def get_job_ids():
-    """Parse job IDs from environment variable"""
-    job_ids_str = os.environ.get('JOB_IDS')
-    try:
-        job_ids = json.loads(job_ids_str)
-        if not isinstance(job_ids, list):
-            print("Error: JOB_IDS is not a list.")
-            sys.exit(1)
-        return job_ids
-    except json.JSONDecodeError as e:
-        print(f"Error parsing JOB_IDS: {e}")
-        sys.exit(1)
-
-
-def cleanup_pods():
-    """Find and terminate RunPod instances"""
-    print(f"Run ID: {RUN_ID}")
-
-    single_job_id = os.environ.get('JOB_ID')
-    
-    if single_job_id:
-        job_ids = [single_job_id]
-        print(f"Job ID: {single_job_id}")
-    else:
-        job_ids = get_job_ids()
-        print(f"Job IDs: {job_ids}")
-
-    # Get all pods associated with RunPod API_KEY
-    try:
-        response = requests.get(PODS_API, headers=HEADERS)
-        response.raise_for_status()
-        pods = response.json()
-    except requests.exceptions.RequestException as e:
-        print(f"Error getting pods: {e}")
-        sys.exit(1)
-
-    # Find and terminate pods created by this workflow run
-    terminated_pods = []
-    for pod in pods:
-        pod_name = pod.get("name", "")
-        pod_id = pod.get("id")
-
-        # Check if this pod was created by one of our jobs
-        if any(f"{job_id}-{RUN_ID}" in pod_name for job_id in job_ids):
-            print(f"Found pod: {pod_id} ({pod_name})")
-            try:
-                print(f"Terminating pod {pod_id}...")
-                term_response = requests.delete(f"{PODS_API}/{pod_id}",
-                                                headers=HEADERS)
-                term_response.raise_for_status()
-                terminated_pods.append(pod_id)
-                print(f"Successfully terminated pod {pod_id}")
-            except requests.exceptions.RequestException as e:
-                print(f"Error terminating pod {pod_id}: {e}")
-                sys.exit(1)
-    
-    if terminated_pods:
-        if single_job_id:
-            print(f"Terminated pod: {terminated_pods[0]}")
-        else:
-            print(f"Terminated {len(terminated_pods)} pods: {terminated_pods}")
-    else:
-        if single_job_id:
-            print(f"No pod found matching pattern: {single_job_id}-{RUN_ID}")
-        else:
-            print("No pods found to terminate.")
-
-
-def main():
-    cleanup_pods()
-
-
-if __name__ == "__main__":
-    main()

.github/workflows/ci-aggregate-status.yml

@@ -0,0 +1,80 @@
+name: Aggregate Test Status
+
+on:
+  status:
+
+permissions:
+  statuses: write
+
+jobs:
+  aggregate:
+    if: >-
+      github.event.context == 'direct-test-completed'
+      && github.event.state == 'success'
+    runs-on: ubuntu-latest
+    steps:
+      - name: Check and update aggregate status
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const sha = context.payload.sha;
+
+            const { data } = await github.rest.repos.getCombinedStatusForRef({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              ref: sha,
+              per_page: 100,
+            });
+
+            const bkStatuses = data.statuses.filter(
+              s => s.context.startsWith('buildkite/ci/')
+            );
+
+            const FASTCHECK_PREFIX = 'buildkite/ci/microscope-';
+            const FULL_SUITE_PREFIXES = [
+              'buildkite/ci/test-tube-',
+              'buildkite/ci/bar-chart-',
+            ];
+
+            const fastcheck = bkStatuses.filter(
+              s => s.context.startsWith(FASTCHECK_PREFIX)
+            );
+            const fullSuite = bkStatuses.filter(
+              s => FULL_SUITE_PREFIXES.some(p => s.context.startsWith(p))
+            );
+
+            if (
+              fastcheck.length > 0
+              && fastcheck.every(s => s.state === 'success')
+            ) {
+              core.info(
+                `All ${fastcheck.length} fastcheck tests passed — updating fastcheck-passed`
+              );
+              await github.rest.repos.createCommitStatus({
+                owner: context.repo.owner,
+                repo: context.repo.repo,
+                sha,
+                state: 'success',
+                context: 'fastcheck-passed',
+                description:
+                  `All ${fastcheck.length} fastcheck tests passed`,
+              });
+            }
+
+            if (
+              fullSuite.length > 0
+              && fullSuite.every(s => s.state === 'success')
+            ) {
+              core.info(
+                `All ${fullSuite.length} full suite tests passed — updating full-suite-passed`
+              );
+              await github.rest.repos.createCommitStatus({
+                owner: context.repo.owner,
+                repo: context.repo.repo,
+                sha,
+                state: 'success',
+                context: 'full-suite-passed',
+                description:
+                  `All ${fullSuite.length} full suite tests passed`,
+              });
+            }

.github/workflows/ci-precommit.yml

@@ -0,0 +1,32 @@
+name: pre-commit
+
+on:
+  pull_request:
+    branches: [main]
+  workflow_call:
+    inputs:
+      ref:
+        description: 'Git ref to checkout (defaults to github.ref)'
+        required: false
+        type: string
+
+permissions:
+  contents: read
+
+jobs:
+  pre-commit:
+    if: github.event_name == 'workflow_call' || github.event.pull_request.draft != true
+    runs-on: ubuntu-latest
+    steps:
+    - uses: actions/checkout@v4
+      with:
+        ref: ${{ inputs.ref || '' }}
+    - uses: actions/setup-python@v5
+      with:
+        python-version: "3.12"
+    - run: echo "::add-matcher::.github/workflows/matchers/actionlint.json"
+    - run: echo "::add-matcher::.github/workflows/matchers/mypy.json"
+    - run: echo "::add-matcher::.github/workflows/matchers/ruff.json"
+    - uses: pre-commit/action@v3.0.1
+      with:
+        extra_args: --all-files --hook-stage manual

.github/workflows/ci-slash-commands.yml

@@ -0,0 +1,271 @@
+name: Slash Commands
+
+on:
+  issue_comment:
+    types: [created]
+
+permissions:
+  contents: read
+  pull-requests: write
+  statuses: write
+
+jobs:
+  handle-merge:
+    if: >-
+      github.event.issue.pull_request != null
+      && startsWith(github.event.comment.body, '/merge')
+    runs-on: ubuntu-latest
+    steps:
+      - name: Check write permission
+        id: perm
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const { data: perm } = await github.rest.repos.getCollaboratorPermissionLevel({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              username: context.payload.comment.user.login,
+            });
+            const hasWrite = ['admin', 'write'].includes(perm.permission);
+            if (!hasWrite) {
+              core.setFailed(`User ${context.payload.comment.user.login} lacks write permission (has: ${perm.permission}).`);
+            }
+            core.setOutput('has_write', String(hasWrite));
+
+      - name: Add ready label and react
+        id: label
+        if: steps.perm.outputs.has_write == 'true'
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const owner = context.repo.owner;
+            const repo = context.repo.repo;
+            const prNumber = context.payload.issue.number;
+            try { await github.rest.issues.removeLabel({ owner, repo, issue_number: prNumber, name: 'ready' }); } catch {}
+            await github.rest.issues.addLabels({ owner, repo, issue_number: prNumber, labels: ['ready'] });
+            await github.rest.reactions.createForIssueComment({
+              owner, repo,
+              comment_id: context.payload.comment.id,
+              content: 'rocket',
+            });
+            const { data: pr } = await github.rest.pulls.get({ owner, repo, pull_number: prNumber });
+            core.setOutput('pr_sha', pr.head.sha);
+            core.setOutput('pr_branch', pr.head.ref);
+            core.setOutput('pr_number', String(prNumber));
+
+      - name: Trigger Full Suite
+        if: steps.perm.outputs.has_write == 'true'
+        env:
+          BUILDKITE_API_TOKEN: ${{ secrets.BUILDKITE_API_TOKEN }}
+          PR_SHA: ${{ steps.label.outputs.pr_sha }}
+          PR_BRANCH: ${{ steps.label.outputs.pr_branch }}
+          PR_NUMBER: ${{ steps.label.outputs.pr_number }}
+          BK_ORG: ${{ vars.BUILDKITE_ORG_SLUG }}
+          BK_PIPELINE: ${{ vars.BUILDKITE_PIPELINE_SLUG }}
+        run: |
+          curl -sS --fail-with-body -X POST \
+            "https://api.buildkite.com/v2/organizations/${BK_ORG}/pipelines/${BK_PIPELINE}/builds" \
+            -H "Authorization: Bearer $BUILDKITE_API_TOKEN" \
+            -H "Content-Type: application/json" \
+            --data-raw "$(jq -n \
+              --arg commit "$PR_SHA" \
+              --arg branch "$PR_BRANCH" \
+              --arg message "Full Suite for PR #${PR_NUMBER} (via /merge)" \
+              --argjson pr_id "$PR_NUMBER" \
+              '{
+                commit: $commit,
+                branch: $branch,
+                message: $message,
+                ignore_pipeline_branch_filters: true,
+                pull_request_id: $pr_id,
+                pull_request_base_branch: "main",
+                env: {
+                  TEST_SCOPE: "full",
+                  FULL_SUITE: "true",
+                  PR_NUMBER: ($pr_id | tostring)
+                }
+              }')"
+
+  parse-command:
+    if: >-
+      github.event.issue.pull_request != null
+      && startsWith(github.event.comment.body, '/test')
+    runs-on: ubuntu-latest
+    outputs:
+      test_type: ${{ steps.parse.outputs.test_type }}
+      test_scope: ${{ steps.parse.outputs.test_scope }}
+      full_suite: ${{ steps.parse.outputs.full_suite }}
+      pr_sha: ${{ steps.pr.outputs.sha }}
+      pr_branch: ${{ steps.pr.outputs.branch }}
+      has_write: ${{ steps.perm.outputs.has_write }}
+    steps:
+      - name: Check write permission
+        id: perm
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const { data: perm } = await github.rest.repos.getCollaboratorPermissionLevel({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              username: context.payload.comment.user.login,
+            });
+            const hasWrite = ['admin', 'write'].includes(perm.permission);
+            core.setOutput('has_write', String(hasWrite));
+            if (!hasWrite) {
+              core.info(`User ${context.payload.comment.user.login} lacks write permission — ignoring.`);
+            }
+
+      - name: Parse /test command
+        id: parse
+        if: steps.perm.outputs.has_write == 'true'
+        shell: bash
+        env:
+          COMMENT: ${{ github.event.comment.body }}
+        run: |
+          set -euo pipefail
+          TEST_NAME=$(echo "$COMMENT" | grep -oP '(?<=/test\s)\S+' | head -1 || true)
+
+          VALID="encoder vae transformer kernel unit ssim training lora-inference lora-training distillation self-forcing vsa vmoba performance api full fastcheck pre-commit"
+          if [ -z "$TEST_NAME" ] || ! echo "$VALID" | grep -qw "$TEST_NAME"; then
+            echo "Unknown test: '$TEST_NAME'. Valid: $VALID"
+            exit 1
+          fi
+
+          declare -A MAP=(
+            [encoder]=encoder [vae]=vae [transformer]=transformer
+            [kernel]=kernel_tests [unit]=unit_test
+            [ssim]=ssim [training]=training
+            [lora-inference]=inference_lora [lora-training]=training_lora
+            [distillation]=distillation_dmd [self-forcing]=self_forcing
+            [vsa]=training_vsa [vmoba]=inference_vmoba
+            [performance]=performance [api]=api_server
+          )
+
+          if [ "$TEST_NAME" = "full" ]; then
+            {
+              echo "test_type=all"
+              echo "test_scope=full"
+              echo "full_suite=true"
+            } >> "$GITHUB_OUTPUT"
+          elif [ "$TEST_NAME" = "fastcheck" ]; then
+            {
+              echo "test_type=fastcheck"
+              echo "test_scope=fastcheck"
+              echo "full_suite=false"
+            } >> "$GITHUB_OUTPUT"
+          elif [ "$TEST_NAME" = "pre-commit" ]; then
+            {
+              echo "test_type="
+              echo "test_scope=precommit"
+              echo "full_suite=false"
+            } >> "$GITHUB_OUTPUT"
+          else
+            {
+              echo "test_type=${MAP[$TEST_NAME]}"
+              echo "test_scope=direct"
+              echo "full_suite=false"
+            } >> "$GITHUB_OUTPUT"
+          fi
+
+      - name: Get PR details
+        id: pr
+        if: steps.perm.outputs.has_write == 'true'
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const { data: pr } = await github.rest.pulls.get({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              pull_number: context.payload.issue.number,
+            });
+            core.setOutput('sha', pr.head.sha);
+            core.setOutput('branch', pr.head.ref);
+
+      - name: React to comment
+        if: steps.perm.outputs.has_write == 'true'
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            await github.rest.reactions.createForIssueComment({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              comment_id: context.payload.comment.id,
+              content: 'rocket',
+            });
+
+  pre-commit:
+    needs: parse-command
+    if: >-
+      needs.parse-command.outputs.has_write == 'true'
+      && needs.parse-command.outputs.test_scope == 'precommit'
+    uses: ./.github/workflows/ci-precommit.yml
+    with:
+      ref: refs/pull/${{ github.event.issue.number }}/merge
+
+  post-precommit-status:
+    needs: [parse-command, pre-commit]
+    if: always() && needs.parse-command.outputs.test_scope == 'precommit'
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        env:
+          PR_SHA: ${{ needs.parse-command.outputs.pr_sha }}
+          RESULT: ${{ needs.pre-commit.result }}
+        with:
+          script: |
+            const state = process.env.RESULT === 'success' ? 'success' : 'failure';
+            await github.rest.repos.createCommitStatus({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              sha: process.env.PR_SHA,
+              state,
+              context: 'pre-commit',
+              description: `Triggered via /test pre-commit (${state})`,
+            });
+
+  trigger-buildkite:
+    needs: parse-command
+    if: >-
+      needs.parse-command.outputs.has_write == 'true'
+      && needs.parse-command.outputs.test_type != ''
+    runs-on: ubuntu-latest
+    steps:
+      - name: Trigger Buildkite
+        env:
+          BUILDKITE_API_TOKEN: ${{ secrets.BUILDKITE_API_TOKEN }}
+          PR_SHA: ${{ needs.parse-command.outputs.pr_sha }}
+          PR_BRANCH: ${{ needs.parse-command.outputs.pr_branch }}
+          PR_NUMBER: ${{ github.event.issue.number }}
+          TEST_SCOPE: ${{ needs.parse-command.outputs.test_scope }}
+          FULL_SUITE: ${{ needs.parse-command.outputs.full_suite }}
+          TEST_TYPE: ${{ needs.parse-command.outputs.test_type }}
+          BK_ORG: ${{ vars.BUILDKITE_ORG_SLUG }}
+          BK_PIPELINE: ${{ vars.BUILDKITE_PIPELINE_SLUG }}
+        run: |
+          curl -sS --fail-with-body -X POST \
+            "https://api.buildkite.com/v2/organizations/${BK_ORG}/pipelines/${BK_PIPELINE}/builds" \
+            -H "Authorization: Bearer $BUILDKITE_API_TOKEN" \
+            -H "Content-Type: application/json" \
+            --data-raw "$(jq -n \
+              --arg commit "$PR_SHA" \
+              --arg branch "$PR_BRANCH" \
+              --arg message "/test ${TEST_TYPE} on PR #${PR_NUMBER}" \
+              --argjson pr_id "$PR_NUMBER" \
+              --arg test_scope "$TEST_SCOPE" \
+              --arg full_suite "$FULL_SUITE" \
+              --arg test_type "$TEST_TYPE" \
+              --arg pr_number "$PR_NUMBER" \
+              '{
+                commit: $commit,
+                branch: $branch,
+                message: $message,
+                ignore_pipeline_branch_filters: true,
+                pull_request_id: $pr_id,
+                pull_request_base_branch: "main",
+                env: {
+                  TEST_SCOPE: $test_scope,
+                  FULL_SUITE: $full_suite,
+                  TEST_TYPE: $test_type,
+                  PR_NUMBER: $pr_number
+                }
+              }')"

.github/workflows/ci-trigger-full-suite.yml

@@ -0,0 +1,83 @@
+name: Trigger Full Suite
+
+on:
+  pull_request_target:
+    types: [labeled, synchronize]
+
+permissions:
+  contents: read
+  pull-requests: read
+
+concurrency:
+  group: full-suite-${{ github.event.pull_request.number }}
+  cancel-in-progress: false
+
+jobs:
+  trigger:
+    if: >-
+      (github.event.action == 'labeled' && github.event.label.name == 'ready')
+      || github.event.action == 'synchronize'
+    runs-on: ubuntu-latest
+    steps:
+      - name: Check ready label
+        id: check
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const { data: pr } = await github.rest.pulls.get({
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              pull_number: context.payload.pull_request.number,
+            });
+            const hasReady = pr.labels.some(l => l.name === 'ready');
+            core.setOutput('has_ready', String(hasReady));
+            if (!hasReady) core.info('No ready label — skipping Full Suite trigger.');
+
+      - name: Cancel previous Buildkite builds
+        if: steps.check.outputs.has_ready == 'true'
+        env:
+          BUILDKITE_API_TOKEN: ${{ secrets.BUILDKITE_API_TOKEN }}
+          PR_BRANCH: ${{ github.event.pull_request.head.ref }}
+        run: |
+          # Find running builds for this branch with TEST_SCOPE=full and cancel them
+          builds=$(curl -sS -H "Authorization: Bearer $BUILDKITE_API_TOKEN" \
+            "https://api.buildkite.com/v2/organizations/${{ vars.BUILDKITE_ORG_SLUG }}/pipelines/${{ vars.BUILDKITE_PIPELINE_SLUG }}/builds?branch=${PR_BRANCH}&state=running,scheduled" \
+            | jq -r '.[] | select(try (.env.TEST_SCOPE == "full") catch false) | .number')
+          for build_num in $builds; do
+            echo "Cancelling Buildkite build #$build_num"
+            curl -sS -X PUT -H "Authorization: Bearer $BUILDKITE_API_TOKEN" \
+              "https://api.buildkite.com/v2/organizations/${{ vars.BUILDKITE_ORG_SLUG }}/pipelines/${{ vars.BUILDKITE_PIPELINE_SLUG }}/builds/${build_num}/cancel"
+          done
+
+      - name: Trigger Buildkite Full Suite
+        if: steps.check.outputs.has_ready == 'true'
+        env:
+          BUILDKITE_API_TOKEN: ${{ secrets.BUILDKITE_API_TOKEN }}
+          PR_SHA: ${{ github.event.pull_request.head.sha }}
+          PR_BRANCH: ${{ github.event.pull_request.head.ref }}
+          PR_NUMBER: ${{ github.event.pull_request.number }}
+          BK_ORG: ${{ vars.BUILDKITE_ORG_SLUG }}
+          BK_PIPELINE: ${{ vars.BUILDKITE_PIPELINE_SLUG }}
+        run: |
+          curl -sS --fail-with-body -X POST \
+            "https://api.buildkite.com/v2/organizations/${BK_ORG}/pipelines/${BK_PIPELINE}/builds" \
+            -H "Authorization: Bearer $BUILDKITE_API_TOKEN" \
+            -H "Content-Type: application/json" \
+            --data-raw "$(jq -n \
+              --arg commit "$PR_SHA" \
+              --arg branch "$PR_BRANCH" \
+              --arg message "Full Suite for PR #${PR_NUMBER}" \
+              --argjson pr_id "$PR_NUMBER" \
+              '{
+                commit: $commit,
+                branch: $branch,
+                message: $message,
+                ignore_pipeline_branch_filters: true,
+                pull_request_id: $pr_id,
+                pull_request_base_branch: "main",
+                env: {
+                  TEST_SCOPE: "full",
+                  FULL_SUITE: "true",
+                  PR_NUMBER: ($pr_id | tostring)
+                }
+              }')"

.github/workflows/community-issue-labeler.yml

@@ -0,0 +1,65 @@
+name: Auto-Label Issues
+
+on:
+  issues:
+    types: [opened, edited]
+
+permissions:
+  issues: write
+
+jobs:
+  label-issues:
+    if: github.repository == 'hao-ai-lab/FastVideo'
+    runs-on: ubuntu-latest
+    steps:
+      - name: Label by keywords
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const title = context.payload.issue.title.toLowerCase();
+            const body = (context.payload.issue.body || '').toLowerCase();
+            const text = title + ' ' + body;
+            const labels = [];
+
+            const rules = [
+              // scope labels (shared with PR labeling via Mergify)
+              // Mapping: label → repo directories
+              //   scope: training    → fastvideo/train/, fastvideo/training/, fastvideo/distillation/
+              //   scope: inference   → fastvideo/pipelines/, fastvideo/entrypoints/, fastvideo/worker/
+              //   scope: attention   → fastvideo/attention/
+              //   scope: kernel      → fastvideo-kernel/, csrc/
+              //   scope: model       → fastvideo/models/, fastvideo/layers/, fastvideo/configs/models/
+              //   scope: data        → fastvideo/dataset/, fastvideo/pipelines/preprocess/
+              //   scope: distributed → fastvideo/distributed/
+              //   scope: docs        → docs/
+              { keywords: ['training', 'finetune', 'fine-tune', 'lora', 'fsdp', 'distill'], label: 'scope: training' },
+              { keywords: ['inference', 'generate', 'pipeline', 'slow', 'latency'], label: 'scope: inference' },
+              { keywords: ['attention', 'vsa', 'flash', 'sta', 'vmoba', 'sparse attn'], label: 'scope: attention' },
+              { keywords: ['kernel', 'csrc', 'cuda kernel', 'thunderkittens'], label: 'scope: kernel' },
+              { keywords: ['wan', 'hunyuan', 'mochi', 'ltx', 'cogvideo', 'flux', 'sd3', 'cosmos'], label: 'scope: model' },
+              { keywords: ['dataset', 'dataloader', 'preprocessing', 'preprocess'], label: 'scope: data' },
+              { keywords: ['distributed', 'sequence parallel', 'fsdp', 'tensor parallel', 'multi-node', 'multi-gpu'], label: 'scope: distributed' },
+              { keywords: ['docs', 'documentation', 'tutorial', 'example'], label: 'scope: docs' },
+              // issue-only labels (cross-module, no single repo directory)
+              { keywords: ['install', 'setup', 'pip', 'cuda', 'uv ', 'import error', 'modulenotfound'], label: 'installation' },
+              { keywords: ['memory', 'oom', 'out of memory', 'gpu memory', 'vram'], label: 'performance' },
+              { keywords: ['windows', 'macos', 'mac os', 'apple', 'mps', 'rocm', 'amd', 'npu'], label: 'platform' },
+            ];
+
+            for (const rule of rules) {
+              if (rule.keywords.some(kw => text.includes(kw))) {
+                labels.push(rule.label);
+              }
+            }
+
+            if (labels.length > 0) {
+              await github.rest.issues.addLabels({
+                owner: context.repo.owner,
+                repo: context.repo.repo,
+                issue_number: context.payload.issue.number,
+                labels: labels,
+              });
+              console.log(`Added labels: ${labels.join(', ')}`);
+            } else {
+              console.log('No keyword matches found');
+            }

.github/workflows/community-stale.yml

@@ -0,0 +1,51 @@
+name: Close Stale Issues and PRs
+
+on:
+  schedule:
+    # Daily at 1:30 AM UTC
+    - cron: '30 1 * * *'
+
+jobs:
+  stale:
+    if: github.repository == 'hao-ai-lab/FastVideo'
+    permissions:
+      issues: write
+      pull-requests: write
+      actions: write
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/stale@997185467fa4f803885201cee163a9f38240193d # v10.1.1
+        with:
+          operations-per-run: 500
+
+          exempt-draft-pr: true
+          exempt-issue-labels: 'keep-open,pinned,security,Bug,RFC'
+          exempt-pr-labels: 'keep-open,pinned'
+
+          labels-to-add-when-unstale: 'unstale'
+          labels-to-remove-when-stale: 'unstale'
+
+          days-before-issue-stale: 90
+          days-before-issue-close: 30
+          stale-issue-label: 'stale'
+          stale-issue-message: >
+            This issue has been automatically marked as stale because it has not
+            had any activity within 90 days. It will be automatically closed if
+            no further activity occurs within 30 days. Leave a comment if you
+            feel this issue should remain open. Thank you!
+          close-issue-message: >
+            This issue has been automatically closed due to inactivity. Please
+            feel free to reopen if you feel it is still relevant. Thank you!
+
+          days-before-pr-stale: 60
+          days-before-pr-close: 14
+          stale-pr-label: 'stale'
+          stale-pr-message: >
+            This pull request has been automatically marked as stale because it
+            has not had any activity within 60 days. It will be automatically
+            closed if no further activity occurs within 14 days. Leave a comment
+            if you feel this pull request should remain open. Thank you!
+          close-pr-message: >
+            This pull request has been automatically closed due to inactivity.
+            Please feel free to reopen if you intend to continue working on it.
+            Thank you!

.github/workflows/community-welcome.yml

@@ -0,0 +1,56 @@
+name: Welcome First-Time Contributors
+
+on:
+  issues:
+    types: [opened]
+  pull_request_target:
+    types: [opened]
+
+permissions:
+  issues: write
+  pull-requests: write
+
+jobs:
+  welcome:
+    if: github.repository == 'hao-ai-lab/FastVideo'
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/first-interaction@34f15e814fe48ac9312ccf29db4e74fa767cbab7 # v1.3.0
+        with:
+          repo-token: ${{ secrets.GITHUB_TOKEN }}
+          issue-message: |
+            Welcome to FastVideo! Thanks for opening your first issue.
+
+            To help us investigate, please include:
+            - **FastVideo version**: `pip show fastvideo`
+            - **GPU**: `nvidia-smi` output (GPU model, driver, CUDA version)
+            - **Python version**: `python --version`
+            - **OS**: e.g., Ubuntu 22.04
+
+            If this is a bug, a minimal reproduction script helps us fix it faster.
+
+            Useful links:
+            - [Documentation](https://hao-ai-lab.github.io/FastVideo)
+            - [Contributing Guide](https://hao-ai-lab.github.io/FastVideo/contributing/overview/)
+            - [Slack](https://join.slack.com/t/fastvideo/shared_invite/zt-3f4lao1uq-u~Ipx6Lt4J27AlD2y~IdLQ)
+          pr-message: |
+            Welcome to FastVideo! Thanks for your first pull request.
+
+            **How our CI works:**
+
+            PRs run a two-tier CI system:
+            1. **Pre-commit** — formatting (yapf), linting (ruff), type checking (mypy). Runs immediately on every PR.
+            2. **Fastcheck** — core GPU tests (encoders, VAEs, transformers, kernels, unit tests). Runs automatically via Buildkite on relevant file changes (~10-15 min).
+            3. **Full Suite** — integration tests, training pipelines, SSIM regression. Runs only when a reviewer adds the `ready` label.
+
+            **Before your PR is reviewed:**
+            - [ ] `pre-commit run --all-files` passes locally
+            - [ ] You've added or updated tests for your changes
+            - [ ] The PR description explains what and why
+
+            If pre-commit fails, a bot comment will explain how to fix it. Fastcheck and Full Suite results appear in the Checks section below.
+
+            **Useful links:**
+            - [Contributing Guide](https://hao-ai-lab.github.io/FastVideo/contributing/overview/)
+            - [Development Roadmap](https://github.com/hao-ai-lab/FastVideo/issues/899)
+            - [Slack](https://join.slack.com/t/fastvideo/shared_invite/zt-3f4lao1uq-u~Ipx6Lt4J27AlD2y~IdLQ)

.github/workflows/infra-build-image.yml

@@ -32,7 +32,7 @@ permissions:
 jobs:
   build-python-3-10:
     if: ${{ github.event.inputs.python_3_10 == 'true' }}
-    uses: ./.github/workflows/build-image-template.yml
+    uses: ./.github/workflows/_template-build-image.yml
     with:
       python_version: '3.10'
       dockerfile_path: docker/Dockerfile.python3.10
@@ -41,7 +41,7 @@ jobs:
 
   build-python-3-11:
     if: ${{ github.event.inputs.python_3_11 == 'true' }}
-    uses: ./.github/workflows/build-image-template.yml
+    uses: ./.github/workflows/_template-build-image.yml
     with:
       python_version: '3.11'
       dockerfile_path: docker/Dockerfile.python3.11
@@ -50,7 +50,7 @@ jobs:
 
   build-python-3-12:
     if: ${{ github.event.inputs.python_3_12 == 'true' }}
-    uses: ./.github/workflows/build-image-template.yml
+    uses: ./.github/workflows/_template-build-image.yml
     with:
       python_version: '3.12'
       dockerfile_path: docker/Dockerfile.python3.12
@@ -59,9 +59,9 @@ jobs:
 
   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
+    uses: ./.github/workflows/_template-build-image.yml
     with:
       python_version: '3.12'
       dockerfile_path: docker/Dockerfile.python3.12.cuda12.9.1
       tag_suffix: py3.12-cuda12.9.1
-    secrets: inherit
+    secrets: inherit

.github/workflows/infra-docs.yml

@@ -7,14 +7,14 @@ on:
       - 'docs/**'
       - 'mkdocs.yml'
       - 'requirements-mkdocs.txt'
-      - '.github/workflows/docs.yml'
+      - '.github/workflows/infra-docs.yml'
   pull_request:
     branches: [ main ]
     paths:
       - 'docs/**'
       - 'mkdocs.yml'
       - 'requirements-mkdocs.txt'
-      - '.github/workflows/docs.yml'
+      - '.github/workflows/infra-docs.yml'
 
 permissions:
   contents: read

.github/workflows/matchers/ruff.json

@@ -0,0 +1,17 @@
+{
+  "problemMatcher": [
+    {
+      "owner": "ruff",
+      "pattern": [
+        {
+          "regexp": "^(.+):(\\d+):(\\d+): (\\w+) (.+)$",
+          "file": 1,
+          "line": 2,
+          "column": 3,
+          "code": 4,
+          "message": 5
+        }
+      ]
+    }
+  ]
+}

.github/workflows/pr-test.yml

@@ -1,338 +0,0 @@
-name: PR Test
-
-on:
-  push:
-    branches: [main]
-    paths:
-      - "fastvideo/**/*.py"
-      - ".github/workflows/pr-test.yml"
-  pull_request:
-    branches: [main]
-    types: [opened, ready_for_review, synchronize, reopened]
-    paths:
-      - "fastvideo/**/*.py"
-      - ".github/workflows/pr-test.yml"
-      - "pyproject.toml"
-      - "docker/Dockerfile.python3.12"
-      - "csrc/**"
-  workflow_dispatch:
-    inputs:
-      run_encoder_test:
-        description: "Run encoder-test"
-        required: false
-        default: false
-        type: boolean
-      run_vae_test:
-        description: "Run vae-test"
-        required: false
-        default: false
-        type: boolean
-      run_transformer_test:
-        description: "Run transformer-test"
-        required: false
-        default: false
-        type: boolean
-      run_ssim_test:
-        description: "Run ssim-test"
-        required: false
-        default: false
-        type: boolean
-      run_training_test:
-        description: "Run training-test"
-        required: false
-        default: false
-        type: boolean
-      run_training_test_VSA:
-        description: "Run training-test-VSA"
-        required: false
-        default: false
-        type: boolean
-      run_precision_test_VSA:
-        description: "Run precision-test-VSA"
-        required: false
-        default: false
-        type: boolean
-      run_unit_test:
-        description: "Run unit-test"
-        required: false
-        default: false
-        type: boolean
-
-env:
-  PYTHONUNBUFFERED: "1"
-
-
-concurrency:
-  group: pr-test-${{ github.ref }}
-  cancel-in-progress: true
-
-jobs:
-  pre-commit:
-    uses: ./.github/workflows/pre-commit.yml
-
-  change-filter:
-    runs-on: ubuntu-latest
-    needs: pre-commit
-    if: ${{ github.event.pull_request.draft == false || github.event_name == 'workflow_dispatch' }}
-    outputs:
-      encoder-test: ${{ steps.filter.outputs.encoder-test }}
-      vae-test: ${{ steps.filter.outputs.vae-test }}
-      transformer-test: ${{ steps.filter.outputs.transformer-test }}
-      training-test: ${{ steps.filter.outputs.training-test }}
-      training-test-VSA: ${{ steps.filter.outputs.training-test-VSA }}
-      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
-        id: filter
-        with:
-          filters: |
-            # Define reusable path patterns
-            common-paths: &common-paths
-              - 'pyproject.toml'
-              - 'docker/Dockerfile.python3.10'
-              - 'docker/Dockerfile.python3.11'
-              - 'docker/Dockerfile.python3.12'
-            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'
-            vsa-paths: &vsa-paths
-              - 'fastvideo/**'
-              - *common-paths
-              - *vsa-kernel-paths
-            
-            # Actual tests
-            encoder-test:
-              - 'fastvideo/models/encoders/**'
-              - 'fastvideo/models/loader/**'
-              - 'fastvideo/tests/encoders/**'
-              - *common-paths
-            vae-test:
-              - 'fastvideo/models/vaes/**'
-              - 'fastvideo/models/loader/**'
-              - 'fastvideo/tests/vaes/**'
-              - *common-paths
-            transformer-test:
-              - 'fastvideo/models/dits/**'
-              - 'fastvideo/models/loader/**'
-              - 'fastvideo/tests/transformers/**'
-              - 'fastvideo/layers/**'
-              - 'fastvideo/attention/**'
-              - *common-paths
-            training-test:
-              - 'fastvideo/**'
-              - *common-paths
-            training-test-VSA:
-              - 'fastvideo/**'
-              - *common-paths
-              - *vsa-kernel-paths
-            precision-test-VSA:
-              - *common-paths
-              - *vsa-kernel-paths
-            unit-test:
-              - 'fastvideo/**'
-              - *common-paths
-
-  encoder-test:
-    needs: change-filter
-    if: >-
-      (github.event_name != 'workflow_dispatch' && needs.change-filter.outputs.encoder-test == 'true') || 
-      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_encoder_test == 'true')
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "encoder-test"
-      gpu_type: "NVIDIA A40"
-      gpu_count: 1
-      volume_size: 100
-      image: "ghcr.io/${{ github.repository }}/fastvideo-dev:py3.12-latest"
-      test_command: "uv pip install -e .[test] && pytest ./fastvideo/tests/encoders -s"
-      timeout_minutes: 30
-    secrets:
-      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-
-  vae-test:
-    needs: change-filter
-    if: >-
-      (github.event_name != 'workflow_dispatch' && needs.change-filter.outputs.vae-test == 'true') || 
-      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_vae_test == 'true')
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "vae-test"
-      gpu_type: "NVIDIA A40"
-      gpu_count: 1
-      volume_size: 100
-      image: "ghcr.io/${{ github.repository }}/fastvideo-dev:py3.12-latest"
-      test_command: "uv pip install -e .[test] && pytest ./fastvideo/tests/vaes -s"
-      timeout_minutes: 30
-    secrets:
-      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-
-  transformer-test:
-    needs: change-filter
-    if: >-
-      (github.event_name != 'workflow_dispatch' && needs.change-filter.outputs.transformer-test == 'true') || 
-      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_transformer_test == 'true')
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "transformer-test"
-      gpu_type: "NVIDIA L40S"
-      gpu_count: 1
-      volume_size: 100
-      image: "ghcr.io/${{ github.repository }}/fastvideo-dev:py3.12-latest"
-      test_command: "uv pip install -e .[test] && pytest ./fastvideo/tests/transformers -s"
-      timeout_minutes: 30
-    secrets:
-      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-      
-  ssim-test:
-    needs: change-filter
-    if: >-
-      github.event_name != 'workflow_dispatch' || (github.event_name == 'workflow_dispatch' && github.event.inputs.run_ssim_test == 'true')
-    strategy:
-      fail-fast: false
-      matrix:
-        python-version: [
-          # {version: "3.10", tag: "latest"},
-          # {version: "3.11", tag: "py3.11-latest"},
-          {version: "3.12", tag: "py3.12-latest"}
-        ]
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "ssim-test-py${{ matrix.python-version.version }}"
-      gpu_type: "NVIDIA A40"
-      gpu_count: 2
-      volume_size: 200
-      disk_size: 200
-      image: "ghcr.io/${{ github.repository }}/fastvideo-dev:${{ matrix.python-version.tag }}"
-      test_command: "uv pip install -e .[test] && pytest ./fastvideo/tests/ssim -vs"
-      timeout_minutes: 60
-    secrets:
-      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-
-  training-test:
-    needs: change-filter
-    if: >-
-      (github.event_name != 'workflow_dispatch' && needs.change-filter.outputs.training-test == 'true') || 
-      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_training_test == 'true')
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "training-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/training/Vanilla -srP"
-      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 }}
-
-  training-test-VSA:
-    needs: change-filter
-    if: >-
-      (github.event_name != 'workflow_dispatch' && needs.change-filter.outputs.training-test-VSA == 'true') || 
-      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_training_test_VSA == 'true')
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "training-test-VSA"
-      gpu_type: "NVIDIA H100 NVL"
-      gpu_count: 2
-      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/training/VSA -srP"
-      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 }}
-
-  precision-test-VSA:
-    needs: change-filter
-    if: >-
-      (github.event_name != 'workflow_dispatch' && needs.change-filter.outputs.precision-test-VSA == 'true') || 
-      (github.event_name == 'workflow_dispatch' && github.event.inputs.run_precision_test_VSA == 'true')
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "precision-test-VSA"
-      gpu_type: "NVIDIA H100 NVL"
-      gpu_count: 1
-      volume_size: 100
-      disk_size: 100
-      image: "ghcr.io/${{ github.repository }}/fastvideo-dev:py3.12-latest"
-      test_command: "uv pip install -e .[test] && python csrc/attn/tests/test_vsa.py"
-      timeout_minutes: 30
-    secrets:
-      RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-      RUNPOD_PRIVATE_KEY: ${{ secrets.RUNPOD_PRIVATE_KEY }}
-
-  unit-test:
-    needs: change-filter
-    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')
-    uses: ./.github/workflows/runpod-test.yml
-    with:
-      job_id: "unit-test"
-      gpu_type: "NVIDIA L40S"
-      gpu_count: 1
-      volume_size: 100
-      disk_size: 100
-      image: "ghcr.io/${{ github.repository }}/fastvideo-dev:py3.12-latest"
-      test_command: "uv pip install -e .[test] && pytest ./fastvideo/dataset/ -vs && pytest ./fastvideo/workflow/ -vs && pytest ./fastvideo/entrypoints/ -vs"
-      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 }}
-
-  runpod-cleanup:
-    # Add other jobs to this list as you create them
-    needs: [encoder-test, vae-test, transformer-test, ssim-test, training-test, training-test-VSA, precision-test-VSA] 
-    if: ${{ always() && ((github.event_name != 'workflow_dispatch' && github.event.pull_request.draft == false) || github.event_name == 'workflow_dispatch') }}
-    runs-on: ubuntu-latest
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-
-      - name: Set up Python
-        uses: actions/setup-python@v5
-        with:
-          python-version: "3.10"
-
-      - name: Install dependencies
-        run: pip install requests
-
-      - name: Cleanup all RunPod instances
-        env:
-          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", "precision-test-VSA"]'
-          RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-          GITHUB_RUN_ID: ${{ github.run_id }}
-        run: python .github/scripts/runpod_cleanup.py

.github/workflows/pre-commit.yml

@@ -1,18 +0,0 @@
-name: pre-commit
-
-on:
-    workflow_call:
-
-jobs:
-  pre-commit:
-    runs-on: ubuntu-latest
-    steps:
-    - uses: actions/checkout@v4
-    - uses: actions/setup-python@v5
-      with:
-        python-version: "3.12"
-    - run: echo "::add-matcher::.github/workflows/matchers/actionlint.json"
-    - run: echo "::add-matcher::.github/workflows/matchers/mypy.json"
-    - uses: pre-commit/action@v3.0.1
-      with:
-        extra_args: --all-files --hook-stage manual

.github/workflows/publish-kernel.yml

@@ -36,11 +36,11 @@ jobs:
           
           if [ "$NEW_VERSION" != "$OLD_VERSION" ]; then
             echo "Version changed from $OLD_VERSION to $NEW_VERSION"
-            echo "changed=true" >> $GITHUB_OUTPUT
-            echo "new-version=$NEW_VERSION" >> $GITHUB_OUTPUT
+            echo "changed=true" >> "$GITHUB_OUTPUT"
+            echo "new-version=$NEW_VERSION" >> "$GITHUB_OUTPUT"
           else
             echo "Version did not change"
-            echo "changed=false" >> $GITHUB_OUTPUT
+            echo "changed=false" >> "$GITHUB_OUTPUT"
           fi
 
   build_wheels:

.github/workflows/runpod-test.yml

@@ -1,94 +0,0 @@
-name: RunPod Test
-
-on:
-  workflow_call:
-    inputs:
-      job_id:
-        required: true
-        type: string
-        description: "Unique identifier for this test job"
-      gpu_type:
-        required: true
-        type: string
-        description: "GPU type to use (e.g. NVIDIA A40, NVIDIA L40S)"
-      gpu_count:
-        required: true
-        type: number
-        description: "Number of GPUs to use"
-      volume_size:
-        required: false
-        type: number
-        default: 20
-        description: "Volume size in GB"
-      disk_size:
-        required: false
-        type: number
-        default: 20
-        description: "Disk size in GB"
-      image:
-        required: true
-        type: string
-        description: "Docker image to use"
-      test_command:
-        required: true
-        type: string
-        description: "Command to run tests"
-      timeout_minutes:
-        required: false
-        type: number
-        default: 30
-        description: "Timeout in minutes"
-    secrets:
-      RUNPOD_API_KEY:
-        required: true
-      RUNPOD_PRIVATE_KEY:
-        required: true
-      WANDB_API_KEY:
-        required: false
-
-jobs:
-  run-test:
-    runs-on: ubuntu-latest
-    environment: runpod-runners
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-
-      - name: Set up Python
-        uses: actions/setup-python@v5
-        with:
-          python-version: "3.12"
-
-      - name: Set up SSH key
-        run: |
-          mkdir -p ~/.ssh
-          echo "${{ secrets.RUNPOD_PRIVATE_KEY }}" > ~/.ssh/id_rsa
-          chmod 600 ~/.ssh/id_rsa
-          ssh-keygen -y -f ~/.ssh/id_rsa > ~/.ssh/id_rsa.pub
-
-      - name: Install dependencies
-        run: pip install requests
-
-      - name: Run tests on RunPod
-        env:
-          JOB_ID: ${{ inputs.job_id }}
-          RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-          GITHUB_RUN_ID: ${{ github.run_id }}
-          WANDB_API_KEY: ${{ secrets.WANDB_API_KEY }}
-        timeout-minutes: ${{ inputs.timeout_minutes }}
-        run: >-
-          python .github/scripts/runpod_api.py 
-          --gpu-type "${{ inputs.gpu_type }}"
-          --gpu-count ${{ inputs.gpu_count }}
-          --volume-size ${{ inputs.volume_size }}
-          --disk-size ${{ inputs.disk_size }}
-          --image "${{ inputs.image }}"
-          --test-command "${{ inputs.test_command }}"
-
-      - name: Terminate RunPod Instances
-        if: ${{ always() }}
-        env:
-          RUNPOD_API_KEY: ${{ secrets.RUNPOD_API_KEY }}
-          GITHUB_RUN_ID: ${{ github.run_id }}
-          JOB_ID: ${{ inputs.job_id }}
-        run: python .github/scripts/runpod_cleanup.py 

.github/workflows/test.yml

@@ -1,31 +0,0 @@
-name: Run Tests
-
-on:
-  push:
-    branches: [ main ]
-  pull_request:
-    branches: [ main ]
-
-jobs:
-  test:
-    runs-on: ubuntu-latest
-    steps:
-      - name: Check out repository
-        uses: actions/checkout@v4
-
-      - name: Set up Python
-        uses: actions/setup-python@v5
-        with:
-          python-version: '3.12'  # or any version you need
-
-      - name: Install dependencies
-        run: |
-          python -m pip install --upgrade pip setuptools wheel
-          pip install torch
-          pip install packaging ninja
-          pip install -e .
-          pip install pytest
-
-      - name: Run Pytest
-        run: |
-          pytest --ignore csrc/attn/test

.github/workflows/vsa-publish.yml

@@ -1,257 +0,0 @@
-name: Publish Video Sparse Attention Kernel to PyPI on Version Change
-
-on:
-  push:
-    branches:
-      - main
-    paths:
-      - "csrc/attn/video_sparse_attn/setup.py"
-  workflow_dispatch:
-
-jobs:
-  check-version-change:
-    runs-on: ubuntu-latest
-    outputs:
-      version-changed: ${{ steps.check-version.outputs.changed }}
-      new-version: ${{ steps.check-version.outputs.new-version }}
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-        with:
-          fetch-depth: 2
-
-      - name: Check if version changed
-        id: check-version
-        run: |
-          cd csrc/attn/video_sparse_attn
-          # Get current commit's version
-          NEW_VERSION=$(grep -oP 'VERSION\s*=\s*"\K[^"]+' setup.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")
-          echo "Old version: $OLD_VERSION"
-          
-          if [ "$NEW_VERSION" != "$OLD_VERSION" ]; then
-            echo "Version changed from $OLD_VERSION to $NEW_VERSION"
-            echo "changed=true" >> $GITHUB_OUTPUT
-            echo "new-version=$NEW_VERSION" >> $GITHUB_OUTPUT
-          else
-            echo "Version did not change"
-            echo "changed=false" >> $GITHUB_OUTPUT
-          fi
-
-  build_wheels:
-    name: Build Wheel
-    needs: check-version-change
-    if: ${{ needs.check-version-change.outputs.version-changed == 'true' || github.event_name == 'workflow_dispatch' }}
-    runs-on: ${{ matrix.os }}
-
-    strategy:
-      fail-fast: false
-      matrix:
-          # Using ubuntu-20.04 instead of 22.04 for more compatibility (glibc). Ideally we'd use the
-          # manylinux docker image, but I haven't figured out how to install CUDA on manylinux.
-          os: [ubuntu-22.04]
-          python-version: ['3.10', '3.11', '3.12', '3.13']
-          # For version reference https://pytorch.org/get-started/previous-versions/
-          torch-cuda:
-            - torch-version: '2.5.1'
-              cuda-version: '12.4.1'
-              torch-cuda-short: 'cu124'
-            - torch-version: '2.6.0'
-              cuda-version: '12.6.3'
-              torch-cuda-short: 'cu126'
-            - torch-version: '2.7.1'
-              cuda-version: '12.8.0'
-              torch-cuda-short: 'cu128'
-
-    steps:
-      - name: Free up disk space
-        run: |
-          echo "Initial disk space:"
-          df -h
-          
-          # Remove large directories
-          sudo rm -rf /usr/share/dotnet
-          sudo rm -rf /usr/local/lib/android
-          sudo rm -rf /opt/ghc
-          sudo rm -rf /usr/local/share/boost
-          sudo rm -rf /usr/share/swift
-          sudo rm -rf /usr/local/lib/node_modules
-          sudo rm -rf /usr/local/share/powershell
-          sudo rm -rf /usr/share/rust
-          sudo rm -rf /usr/local/.ghcup
-          
-          # Remove cached files
-          sudo rm -rf /var/lib/apt/lists/*
-          sudo rm -rf /var/cache/apt/archives/*
-          
-          echo "Disk space after cleanup:"
-          df -h
-
-      - name: Checkout
-        uses: actions/checkout@v4
-
-      - name: Set up Python
-        uses: actions/setup-python@v5
-        with:
-          python-version: ${{ matrix.python-version }}
-
-      - name: Install CUDA ${{ matrix.torch-cuda.cuda-version }}
-        uses: Jimver/cuda-toolkit@v0.2.21
-        id: cuda-toolkit
-        with:
-          cuda: ${{ matrix.torch-cuda.cuda-version }}
-          linux-local-args: '["--toolkit"]'
-          method: 'network'
-
-      - name: Install dependencies (GCC, Clang, CUDA Paths, Git)
-        run: |
-          sudo apt update
-          sudo apt install -y git patchelf gcc-11 g++-11 clang-11
-          sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
-
-          # Allow Git to Access Safe Directory
-          git config --global --add safe.directory /__w/FastVideo/FastVideo
-
-          # Set CUDA environment variables
-          export CUDA_HOME=/usr/local/cuda-${{ matrix.torch-cuda.cuda-version }}
-          export PATH=${CUDA_HOME}/bin:${PATH}
-          export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
-          
-          # Verify installation
-          gcc --version
-          g++ --version
-          clang-11 --version
-          nvcc --version
-
-      - name: Install PyTorch ${{ matrix.torch-cuda.torch-version }}+cu${{ matrix.torch-cuda.cuda-version }}
-        run: |
-          pip install --upgrade pip
-          # With python 3.13 and torch 2.5.1, unless we update typing-extensions, we get error
-          # AttributeError: attribute '__default__' of 'typing.ParamSpec' objects is not writable
-          pip install typing-extensions==4.12.2
-          # We want to figure out the CUDA version to download pytorch
-          # e.g. we can have system CUDA version being 11.7 but if torch==1.12 then we need to download the wheel from cu116
-          # see https://github.com/pytorch/pytorch/blob/main/RELEASE.md#release-compatibility-matrix
-          pip install --no-cache-dir torch==${{ matrix.torch-cuda.torch-version }} --index-url https://download.pytorch.org/whl/${{matrix.torch-cuda.torch-cuda-short}}
-          nvcc --version
-          python --version
-          python -c "import torch; print('PyTorch:', torch.__version__)"
-          python -c "import torch; print('CUDA:', torch.version.cuda)"
-          python -c "from torch.utils import cpp_extension; print (cpp_extension.CUDA_HOME)"
-
-      - name: Build wheel
-        run: |
-          export PYTHONPATH=$GITHUB_WORKSPACE:$PYTHONPATH
-          
-          # We want setuptools >= 49.6.0 otherwise we can't compile the extension if system CUDA version is 11.7 and pytorch cuda version is 11.6
-          # https://github.com/pytorch/pytorch/blob/664058fa83f1d8eede5d66418abff6e20bd76ca8/torch/utils/cpp_extension.py#L810
-          # However this still fails so I'm using a newer version of setuptools
-          pip install setuptools
-          pip install ninja packaging wheel
-
-          cd csrc/attn/video_sparse_attn  # Move into the correct folder
-          git submodule update --init --recursive # Ensure ThunderKittens submodule is initialized
-          python setup.py bdist_wheel --dist-dir=dist
-
-      - name: Rename wheel file
-        run: |
-          cd csrc/attn/video_sparse_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)
-          # Get the correct version format
-          tmpname=cu${CUDA_SHORT_VERSION}torch${TORCH_SHORT_VERSION}
-          wheel_name=$(ls dist/*whl | xargs -n 1 basename | sed "s/-/+$tmpname-/2")
-          # Rename with version information
-          ls dist/*whl |xargs -I {} mv {} dist/${wheel_name}
-          echo "wheel_name=${wheel_name}" >> $GITHUB_ENV
-
-      - name: Upload wheel artifact
-        uses: actions/upload-artifact@v4
-        with:
-          name: ${{ env.wheel_name }}
-          path: csrc/attn/video_sparse_attn/dist/*.whl
-          retention-days: 90
-
-  publish_package:
-    name: Publish package
-    needs: [build_wheels, check-version-change]
-    if: ${{ needs.check-version-change.outputs.version-changed == 'true' || github.event_name == 'workflow_dispatch' }}
-    runs-on: ubuntu-22.04
-    permissions:
-      id-token: write  # Needed for OIDC Trusted Publishing
-
-    steps:
-      - uses: actions/checkout@v4
-
-      - uses: actions/setup-python@v5
-        with:
-          python-version: '3.10'
-
-      - name: Install CUDA 12.4.1
-        uses: Jimver/cuda-toolkit@v0.2.21
-        id: cuda-toolkit
-        with:
-          cuda: 12.4.1
-          linux-local-args: '["--toolkit"]'
-          method: 'network'
-          sub-packages: '["nvcc"]'
-
-      - name: Install dependencies (GCC, Clang, CUDA Paths, Git)
-        run: |
-          sudo apt update
-          sudo apt install -y git patchelf gcc-11 g++-11 clang-11
-          sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
-
-          # Allow Git to Access Safe Directory
-          git config --global --add safe.directory /__w/FastVideo/FastVideo
-
-          # Set CUDA environment variables
-          export CUDA_HOME=/usr/local/cuda-12.4.1
-          export PATH=${CUDA_HOME}/bin:${PATH}
-          export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:$LD_LIBRARY_PATH
-          
-          # Verify installation
-          gcc --version
-          g++ --version
-          clang-11 --version
-          nvcc --version
-
-      - name: Install PyTorch 2.5.1+cu12.4.1
-        run: |
-          pip install --upgrade pip
-          # With python 3.13 and torch 2.5.1, unless we update typing-extensions, we get error
-          # AttributeError: attribute '__default__' of 'typing.ParamSpec' objects is not writable
-          pip install typing-extensions==4.12.2
-          # We want to figure out the CUDA version to download pytorch
-          # e.g. we can have system CUDA version being 11.7 but if torch==1.12 then we need to download the wheel from cu116
-          # see https://github.com/pytorch/pytorch/blob/main/RELEASE.md#release-compatibility-matrix
-          export TORCH_CUDA_VERSION=124
-          pip install --no-cache-dir torch==2.5.1 --index-url https://download.pytorch.org/whl/cu${TORCH_CUDA_VERSION}
-          nvcc --version
-          python --version
-          python -c "import torch; print('PyTorch:', torch.__version__)"
-          python -c "import torch; print('CUDA:', torch.version.cuda)"
-          python -c "from torch.utils import cpp_extension; print (cpp_extension.CUDA_HOME)"
-
-      - name: Build source distribution
-        run: |
-          export PYTHONPATH=$GITHUB_WORKSPACE:$PYTHONPATH
-          
-          # We want setuptools >= 49.6.0 otherwise we can't compile the extension if system CUDA version is 11.7 and pytorch cuda version is 11.6
-          # https://github.com/pytorch/pytorch/blob/664058fa83f1d8eede5d66418abff6e20bd76ca8/torch/utils/cpp_extension.py#L810
-          # However this still fails so I'm using a newer version of setuptools
-          pip install setuptools
-          pip install ninja packaging wheel
-
-          cd csrc/attn/video_sparse_attn  # Move into the correct folder
-          git submodule update --init --recursive # Ensure ThunderKittens submodule is initialized
-          python setup.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/

.gitignore

@@ -85,6 +85,10 @@ docs/distillation/examples/
 dmd_t2v_output/
 preprocess_output_text/
 
+# Next.js / Node artifacts under ui/: see ui/.gitignore
+
 .claude/
 .codex/
+.sisyphus/
 openspec/
+fastvideo/tests/ssim/reference_videos/**

.nvimlog

@@ -0,0 +1 @@
+WRN 2026-03-26T13:46:33.469 ?.19646    server_start:193: Failed to start server: operation not permitted: /var/folders/z_/h_6myyk14d1b7z87z3vy4mjh0000gn/T/nvim.dsynkd/iSe0el/nvim.19646.0

.pre-commit-config.yaml

@@ -18,10 +18,8 @@ exclude: |
     fastvideo/train\.py|
     fastvideo/utils/.*|
     examples/.*|
-    .github/workflows/fastvideo-publish.yml|
-    .github/workflows/sta-publish.yml|
-    .github/workflows/vsa-publish.yml|
-    .github/workflows/build-image-template.yml|
+    .github/workflows/publish-fastvideo.yml|
+    .github/workflows/_template-build-image.yml|
     docs/source/inference/support_matrix.md
   )
 repos:

.python-version

@@ -0,0 +1 @@
+3.12

README.md

@@ -9,7 +9,7 @@
 **FastVideo is a unified post-training and real-time inference framework for accelerated video generation.**
 
 ## NEWS
-
+- `2026/03/17`: Release Live demo: [Into the Dreamverse: Vibe Directing in FastVideo](https://dreamverse.fastvideo.org/), check out the [Blog](https://haoailab.com/blogs/dreamverse/).
 - `2026/03/13`: Release Live demo: [Create a 5s 1080p Video in 4.5s with FastVideo on a Single GPU](https://1080p.fastvideo.org/), check out the [Blog](https://haoailab.com/blogs/fastvideo_realtime_1080p/).
 - `2025/11/19`: Release [CausalWan2.2 I2V A14B Preview](https://huggingface.co/FastVideo/CausalWan2.2-I2V-A14B-Preview-Diffusers) models, [Blog](https://hao-ai-lab.github.io/blogs/fastvideo_causalwan_preview/) and [Inference Code!](https://github.com/hao-ai-lab/FastVideo/blob/main/examples/inference/basic/basic_self_forcing_causal_wan2_2_i2v.py).
 - `2025/08/04`: Release [FastWan](https://hao-ai-lab.github.io/FastVideo/distillation/dmd) models and [Sparse-Distillation](https://hao-ai-lab.github.io/blogs/fastvideo_post_training/).

docs/README.md

@@ -36,4 +36,4 @@ Then open your browser to: http://localhost:8000
 
 ## Automatic Deployment
 
-Documentation is automatically built and deployed to GitHub Pages when changes are pushed to the `main` branch via the `.github/workflows/docs.yml` workflow.
+Documentation is automatically built and deployed to GitHub Pages when changes are pushed to the `main` branch via the `.github/workflows/infra-docs.yml` workflow.

docs/attention/attn_qat/index.md

@@ -0,0 +1,318 @@
+# Attention QAT
+
+Attention QAT in FastVideo covers two related, but different, backends:
+
+- `ATTN_QAT_INFER`: the inference-oriented CUDA kernel path
+- `ATTN_QAT_TRAIN`: the training-oriented Triton attention path
+
+Both are selected with `FASTVIDEO_ATTENTION_BACKEND`, but they are not
+interchangeable. The main practical split is:
+
+- use `ATTN_QAT_INFER` for standalone inference with the dedicated inference
+  kernel
+- use `ATTN_QAT_TRAIN` for finetuning, validation during training, or when you
+  specifically want to reproduce the training-side attention path
+
+## Quick Start
+
+If your goal is "run Wan 2.1 14B with Attention QAT inference weights", this is
+the shortest path:
+
+1. Build the in-repo kernel package so FastVideo can import `attn_qat_infer`.
+2. Download the Wan 2.1 14B QAT checkpoint.
+3. Edit the provided inference example to point at the 14B base model and the
+   downloaded QAT safetensors.
+4. Run the example with `ATTN_QAT_INFER`.
+
+### Step 1. Build the kernel package
+
+Before using either Attention QAT backend, build the in-repo
+`fastvideo-kernel` package from source:
+
+```bash
+git submodule update --init --recursive
+cd fastvideo-kernel
+./build.sh
+```
+
+After a successful build:
+
+- `ATTN_QAT_TRAIN` should be able to import `fastvideo_kernel`
+- `ATTN_QAT_INFER` should be able to import `attn_qat_infer`
+
+`ATTN_QAT_INFER` currently targets the Blackwell CUDA path under
+`fastvideo-kernel/attn_qat_infer/` and requires CUDA 12.8+.
+
+### Step 2. Download the Wan 2.1 14B QAT checkpoint
+
+FastVideo includes a helper script:
+
+- `examples/inference/optimizations/download_14B_qat.sh`
+
+By default it downloads:
+
+- Hugging Face repo: `FastVideo/14B_qat_400`
+- local directory: `checkpoints/14B_qat_400`
+
+Prerequisites:
+
+- `huggingface_hub` installed, for example:
+  `uv pip install huggingface_hub`
+- access to the model repo if it is private or gated:
+  `huggingface-cli login`
+
+Run the downloader:
+
+```bash
+bash examples/inference/optimizations/download_14B_qat.sh
+```
+
+To download into a custom directory:
+
+```bash
+bash examples/inference/optimizations/download_14B_qat.sh /path/to/14B_qat_400
+```
+
+The script prints a ready-to-copy `init_weights_from_safetensors=...` value at
+the end.
+
+### Step 3. Edit the provided inference example
+
+The example to start from is:
+
+- `examples/inference/optimizations/attn_qat_inference_example.py`
+
+Open that file and update these two values:
+
+1. Change the base model from `Wan-AI/Wan2.1-T2V-1.3B-Diffusers` to
+   `Wan-AI/Wan2.1-T2V-14B-Diffusers`
+2. Replace
+   `init_weights_from_safetensors="safetensors_path"` with the directory that
+   contains the downloaded `.safetensors` files
+
+Example:
+
+```python
+import os
+
+from fastvideo import VideoGenerator
+
+os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "ATTN_QAT_INFER"
+
+generator = VideoGenerator.from_pretrained(
+    "Wan-AI/Wan2.1-T2V-14B-Diffusers",
+    num_gpus=1,
+    use_fsdp_inference=True,
+    dit_cpu_offload=False,
+    vae_cpu_offload=False,
+    text_encoder_cpu_offload=True,
+    pin_cpu_memory=False,
+    init_weights_from_safetensors="checkpoints/14B_qat_400",
+)
+```
+
+Important:
+
+- the checked-in example currently uses the `1.3B` base model until you edit it
+- do not load the 14B QAT weights on top of the `1.3B` base model; the weights
+  and model config will not match
+
+### Step 4. Run the inference example
+
+```bash
+python examples/inference/optimizations/attn_qat_inference_example.py
+```
+
+Generated videos are written to `video_samples/` by default.
+
+## Backend Overview
+
+| Backend | Best for | Package requirement | Primary kernel location |
+|---------|----------|---------------------|-------------------------|
+| `ATTN_QAT_TRAIN` | finetuning, training-time validation, reproducing the training path | `fastvideo_kernel` | `fastvideo-kernel/python/fastvideo_kernel/triton_kernels/attn_qat_train.py` |
+| `ATTN_QAT_INFER` | standalone inference with the dedicated CUDA kernel | `attn_qat_infer` from the in-repo `fastvideo-kernel` checkout | `fastvideo-kernel/attn_qat_infer/` |
+
+FastVideo routes backend selection through:
+
+- `fastvideo/envs.py`
+- `fastvideo/platforms/cuda.py`
+- `fastvideo/attention/backends/attn_qat_train.py`
+- `fastvideo/attention/backends/attn_qat_infer.py`
+
+The legacy training pipeline also contains explicit Attention QAT integration:
+
+- `fastvideo/training/training_pipeline.py`
+
+That pipeline forces generator loading through `ATTN_QAT_TRAIN` when
+`FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_TRAIN` or `--generator_4bit_attn` is
+enabled.
+
+## Inference Workflows
+
+For standalone inference, prefer `ATTN_QAT_INFER` when the CUDA kernel is
+available. Use `ATTN_QAT_TRAIN` for inference only if you intentionally want to
+exercise the training-side attention path for debugging or parity checks.
+
+### Minimal Python example
+
+```python
+import os
+
+from fastvideo import VideoGenerator
+
+os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "ATTN_QAT_INFER"
+
+generator = VideoGenerator.from_pretrained(
+    "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
+    num_gpus=1,
+)
+
+generator.generate_video(
+    "A cinematic close-up of rain on a neon street at night.",
+    output_path="video_samples",
+    save_video=True,
+)
+```
+
+### Loading custom safetensors during inference
+
+FastVideo supports loading custom transformer weights through
+`init_weights_from_safetensors`.
+
+This value can point to either:
+
+- a directory containing one or more `.safetensors` files
+- a single `.safetensors` file
+
+For Wan 2.1 14B QAT inference, the common pattern is:
+
+```python
+generator = VideoGenerator.from_pretrained(
+    "Wan-AI/Wan2.1-T2V-14B-Diffusers",
+    num_gpus=1,
+    use_fsdp_inference=True,
+    init_weights_from_safetensors="checkpoints/14B_qat_400",
+)
+```
+
+### CLI example
+
+You can also force the backend from the command line:
+
+```bash
+FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_INFER \
+fastvideo generate \
+  --model-path Wan-AI/Wan2.1-T2V-14B-Diffusers \
+  --num-gpus 1 \
+  --sp-size 1 \
+  --tp-size 1 \
+  --height 480 \
+  --width 832 \
+  --num-frames 77 \
+  --num-inference-steps 50 \
+  --guidance-scale 6.0 \
+  --prompt "A cinematic close-up of rain on a neon street at night." \
+  --output-path outputs_video/
+```
+
+If you want to use custom QAT transformer weights from the CLI, pass the same
+custom weight override that the Python API uses:
+
+```bash
+FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_INFER \
+fastvideo generate \
+  --model-path Wan-AI/Wan2.1-T2V-14B-Diffusers \
+  --init-weights-from-safetensors checkpoints/14B_qat_400 \
+  --num-gpus 1 \
+  --output-path outputs_video/ \
+  --prompt "A cinematic close-up of rain on a neon street at night."
+```
+
+## Training Workflows
+
+Today the checked-in Attention QAT training launchers use the legacy training
+pipeline in `fastvideo/training/wan_training_pipeline.py`.
+
+### Ready-made launchers
+
+Use the provided SLURM scripts directly:
+
+```bash
+sbatch examples/training/finetune/wan_t2v_1.3B/crush_smol/finetune_t2v_qat_attn.sh
+sbatch examples/training/finetune/wan_t2v_14B/finetune_t2v_qat_attn.sh
+```
+
+Both scripts already set:
+
+```bash
+export FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_TRAIN
+```
+
+Before launching, update the script-local values that depend on your
+environment:
+
+- `WANDB_API_KEY`
+- `MODEL_PATH`
+- `DATA_DIR`
+- `VALIDATION_DATASET_FILE`
+- output directory and SLURM resource requests
+
+### What the launchers run
+
+The training scripts eventually invoke:
+
+```bash
+torchrun fastvideo/training/wan_training_pipeline.py ...
+```
+
+If you are adapting the workflow to your own cluster or running outside SLURM,
+the main Attention QAT requirement is still:
+
+```bash
+export FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_TRAIN
+```
+
+Then launch the normal Wan training pipeline with your preferred `torchrun`
+arguments and training flags.
+
+## Where The Code Lives
+
+Use these paths when you want to trace or modify the Attention QAT flow:
+
+| Location | Purpose |
+|----------|---------|
+| `fastvideo/attention/backends/attn_qat_train.py` | FastVideo wrapper that imports and calls the Triton training kernel |
+| `fastvideo/attention/backends/attn_qat_infer.py` | FastVideo wrapper that imports and calls the inference kernel |
+| `fastvideo-kernel/CMakeLists.txt` | Kernel build definition that compiles the `attn_qat_infer` inference extensions |
+| `fastvideo/platforms/cuda.py` | Chooses the concrete attention backend at runtime |
+| `fastvideo/envs.py` | Documents supported `FASTVIDEO_ATTENTION_BACKEND` values |
+| `fastvideo/training/training_pipeline.py` | Training-time forcing logic for the generator attention backend |
+| `fastvideo-kernel/python/fastvideo_kernel/triton_kernels/attn_qat_train.py` | Triton implementation for `ATTN_QAT_TRAIN` |
+| `fastvideo-kernel/attn_qat_infer/api.py` | Python API entrypoint for the inference kernel |
+| `fastvideo-kernel/benchmarks/benchmark_*.py` | Kernel-side benchmark scripts for FlashAttn2, SageAttention3, FP4, and comparison plots |
+| `fastvideo-kernel/attn_qat_infer/blackwell/api.cu` | CUDA implementation behind `ATTN_QAT_INFER` |
+| `fastvideo-kernel/tests/test_attn_qat_train.py` | Kernel-level test coverage for the training path |
+| `examples/inference/optimizations/attn_qat_inference_example.py` | Ready-to-edit inference example for custom Attention QAT weights |
+| `examples/inference/optimizations/download_14B_qat.sh` | Helper script for downloading the Wan 2.1 14B QAT checkpoint |
+| `examples/training/finetune/wan_t2v_1.3B/crush_smol/finetune_t2v_qat_attn.sh` | Ready-to-run Wan 1.3B Attention QAT finetune launcher |
+| `examples/training/finetune/wan_t2v_14B/finetune_t2v_qat_attn.sh` | Ready-to-run Wan 14B Attention QAT finetune launcher |
+
+## Troubleshooting
+
+- If `ATTN_QAT_TRAIN` fails to import, verify that `fastvideo-kernel` built
+  successfully and exposes `fastvideo_kernel`.
+- If `ATTN_QAT_INFER` fails to import, verify that the local build exposes the
+  `attn_qat_infer` package.
+- If the Wan 2.1 14B example fails after you changed only the checkpoint path,
+  make sure you also changed the base model to
+  `Wan-AI/Wan2.1-T2V-14B-Diffusers`.
+- If you hit issues with CPU memory pressure or obscure CUDA argument errors in
+  the example script, try setting `pin_cpu_memory=False`.
+- If you want a known-safe fallback for debugging, use
+  `FASTVIDEO_ATTENTION_BACKEND=TORCH_SDPA`.
+
+## Related Pages
+
+- [Attention Overview](../index.md)
+- [Inference Optimizations](../../inference/optimizations.md)
+- [Debugging](../../utilities/debugging.md)

docs/attention/index.md

@@ -5,6 +5,8 @@ FastVideo provides highly optimized custom attention kernels to accelerate video
 ## Supported Kernels
 
 * **[Video Sparse Attention (VSA)](vsa/index.md)**: Sparse attention mechanism selecting top-k blocks.
+* **[Attention QAT](attn_qat/index.md)**: Dedicated guide for Attention QAT
+  inference, training, checkpoint loading, and troubleshooting.
 * **[Sliding Tile Attention (STA)](sta/index.md)**: STA kernel support is kept in
   `fastvideo-kernel`; full FastVideo STA pipeline workflow is archived in
   `sta_do_not_delete`.

docs/contributing/ci_architecture.md

@@ -0,0 +1,331 @@
+# CI Architecture
+
+## Overview
+
+FastVideo uses a three-tier CI pipeline designed to keep feedback fast on every push while
+protecting `main` through a full GPU regression suite before any merge.
+
+```
+PR push
+  │
+  ├─► Tier 1: Pre-commit (~2 min)
+  │     GitHub Actions / ubuntu-latest
+  │     yapf, ruff, mypy, codespell, pymarkdown, actionlint, check-filenames
+  │
+  └─► Tier 2: Fastcheck (~10-20 min, path-filtered)
+        Buildkite / Modal GPU instances
+        Only runs tests for paths you changed
+
+              │ (developer comments /merge or maintainer adds 'ready' label)
+              ▼
+        Tier 3: Full Suite (~60-90 min)
+          Buildkite / Modal GPU instances
+          All integration, SSIM, training, and performance tests
+          Runs on the PR branch directly
+              │
+          pass ──► Mergify auto-squash-merges to main, branch deleted
+          fail ──► fix the regression, push, and /merge again
+```
+
+---
+
+## CI Tiers
+
+### Tier 1: Pre-commit (every PR push)
+
+| Attribute | Value |
+|-----------|-------|
+| Triggered by | Every push to any PR branch, plus pushes to `main` |
+| Runs on | GitHub Actions, `ubuntu-latest` |
+| Duration | ~2 minutes |
+
+**Checks run** (from `.pre-commit-config.yaml`, stage: `manual`):
+
+| Hook | What it checks |
+|------|---------------|
+| `yapf` | Python code formatting |
+| `ruff` | Python linting and auto-fixable style issues |
+| `codespell` | Spelling errors in code and docs |
+| `pymarkdown` | Markdown formatting |
+| `actionlint` | GitHub Actions workflow syntax |
+| `mypy` | Static type checking (Python 3.10 target) |
+| `check-filenames` | No spaces in tracked filenames |
+
+A failure here means code style or type issues. Run `pre-commit run --all-files` locally to
+replicate CI results before pushing.
+
+---
+
+### Tier 2: Fastcheck (path-filtered, every PR push)
+
+| Attribute | Value |
+|-----------|-------|
+| Triggered by | Every push; the monorepo-diff plugin skips tests for unchanged paths |
+| Runs on | Buildkite, Modal GPU instances |
+| Duration | ~10-20 minutes per test (run in parallel) |
+
+**Tests and their path triggers:**
+
+| Buildkite label | `TEST_TYPE` | Triggers when you change |
+|-----------------|-------------|--------------------------|
+| Encoder Tests | `encoder` | `fastvideo/models/encoders/**`, `fastvideo/models/loader/**`, `fastvideo/tests/encoders/**`, `pyproject.toml`, `docker/Dockerfile.python3.12` |
+| VAE Tests | `vae` | `fastvideo/models/vaes/**`, `fastvideo/models/loader/**`, `fastvideo/tests/vaes/**`, `pyproject.toml`, `docker/Dockerfile.python3.12` |
+| Transformer Tests | `transformer` | `fastvideo/models/dits/**`, `fastvideo/models/loader/**`, `fastvideo/tests/transformers/**`, `fastvideo/layers/**`, `fastvideo/attention/**`, `pyproject.toml`, `docker/Dockerfile.python3.12` |
+| Kernel Tests | `kernel_tests` | `fastvideo-kernel/**`, `pyproject.toml`, `docker/Dockerfile.python3.12` |
+| Unit Tests | `unit_test` | `fastvideo/**`, `.buildkite/**`, `.github/**`, `pyproject.toml`, `docker/Dockerfile.python3.12` |
+
+A Fastcheck failure means a component-level regression. Check the Buildkite build log for the
+failing test's output.
+
+---
+
+### Tier 3: Full Test Suite (triggered by `ready` label)
+
+| Attribute | Value |
+|-----------|-------|
+| Triggered by | Adding the `ready` label to the PR (via `/merge` command), or a `/test full` command |
+| Runs on | Buildkite, Modal GPU instances |
+| Duration | 60-90 minutes total (tests run in parallel, path-filtered) |
+
+**Tests:**
+
+| Buildkite label | `TEST_TYPE` | Timeout |
+|-----------------|-------------|---------|
+| SSIM Tests | `ssim` | 90 min |
+| LoRA Inference Tests | `inference_lora` | 20 min |
+| Training Tests | `training` | 15 min |
+| Distillation DMD Tests | `distillation_dmd` | 15 min |
+| Self-Forcing Tests | `self_forcing` | 15 min |
+| LoRA Training Tests | `training_lora` | 15 min |
+| Training Tests VSA | `training_vsa` | 15 min |
+| Inference Tests VMoBA | `inference_vmoba` | 15 min |
+| Performance Tests | `performance` | 30 min |
+| API Server Tests | `api_server` | 30 min |
+
+If a Full Suite test fails, check the Buildkite build log for the failing step's output.
+Fix the regression, push, and comment `/merge` again to re-trigger.
+
+---
+
+## Auto-merge Flow
+
+Mergify prevents untested code from landing on `main` by gating squash-merge on the Full
+Suite passing directly on the PR branch.
+
+**How it works:**
+
+1. A developer comments `/merge` on an approved PR (or a maintainer adds the `ready` label).
+2. The `ready` label triggers `ci-trigger-full-suite.yml`, which calls the Buildkite API to
+   run the Full Suite on the PR branch itself.
+3. While the Full Suite runs, Mergify also auto-rebases the PR branch against `main` if it
+   is behind and has no conflicts.
+4. Once the Full Suite posts `full-suite-passed`, Mergify checks all **merge conditions**:
+   - `pre-commit` check is green
+   - `fastcheck-passed` check is green
+   - `full-suite-passed` check is green
+   - At least 1 approved review (`#approved-reviews-by>=1`)
+   - PR title starts with a valid `[type]` tag
+   - PR is not a draft
+   - No merge conflicts
+5. If all conditions pass, Mergify squash-merges to `main` automatically. The branch is
+   deleted after merge.
+6. If the Full Suite fails, the developer fixes the issue, pushes, and comments `/merge`
+   again to re-trigger.
+
+**Merge conditions summary:**
+
+| Condition | Meaning |
+|-----------|---------|
+| `check-success~=pre-commit` | Tier 1 pre-commit must be green |
+| `check-success=fastcheck-passed` | Tier 2 Fastcheck must be green |
+| `check-success=full-suite-passed` | Tier 3 Full Suite must be green |
+| `#approved-reviews-by>=1` | At least one approved review |
+| `title~=(?i)^\[(feat|bugfix|...)` | PR title has a valid type tag |
+| `-draft` | PR is not in draft state |
+| `-conflict` | No merge conflicts with base branch |
+| `-closed` | PR is still open |
+
+---
+
+## Label System
+
+Labels are applied automatically. You don't need to set them manually.
+
+### Type Labels (from PR title prefix)
+
+Applied by Mergify based on the `[tag]` at the start of the PR title.
+
+| Label | Matched title prefix | Meaning |
+|-------|---------------------|---------|
+| `type: feat` | `[feat]` or `[feature]` | New feature or capability |
+| `type: bugfix` | `[bugfix]` or `[fix]` | Bug fix |
+| `type: refactor` | `[refactor]` | Code restructuring, no behavior change |
+| `type: perf` | `[perf]` | Performance improvement |
+| `type: ci` | `[ci]` | CI/CD or tooling changes |
+| `type: docs` | `[doc]` or `[docs]` | Documentation only |
+| `type: misc` | `[misc]` or `[chore]` | Housekeeping, dependency bumps |
+| `type: new-model` | `[new-model]` | Adding a new model |
+
+### Scope Labels (from changed files)
+
+Applied by Mergify based on which paths you modified. Multiple scope labels can be added.
+
+| Label | File paths that trigger it |
+|-------|---------------------------|
+| `scope: training` | `fastvideo/train/`, `fastvideo/training/`, `fastvideo/distillation/`, `examples/train/`, `examples/training/`, `examples/distill/` |
+| `scope: inference` | `fastvideo/pipelines/basic/`, `fastvideo/pipelines/stages/`, `fastvideo/pipelines/samplers/`, `fastvideo/entrypoints/`, `fastvideo/worker/`, `fastvideo/configs/sample/`, `fastvideo/configs/pipelines/`, `examples/inference/` |
+| `scope: attention` | `fastvideo/attention/` |
+| `scope: kernel` | `fastvideo-kernel/`, `csrc/` |
+| `scope: data` | `fastvideo/dataset/`, `fastvideo/pipelines/preprocess/`, `examples/preprocessing/` |
+| `scope: infra` | `.github/`, `.buildkite/`, `fastvideo/tests/`, `docker/` |
+| `scope: distributed` | `fastvideo/distributed/` |
+| `scope: docs` | `docs/` |
+| `scope: ui` | `ui/` |
+| `scope: model` | `fastvideo/models/`, `fastvideo/layers/`, `fastvideo/configs/models/` |
+
+### Process Labels
+
+| Label | Who sets it | Meaning |
+|-------|-------------|---------|
+| `ready` | Developer (`/merge` command) or maintainer | Triggers Full Suite and enables auto-merge |
+| `needs-rebase` | Mergify (automatic) | PR has merge conflicts; rebase needed |
+| `do-not-merge` | Maintainer | Blocks queue entry regardless of other conditions |
+
+---
+
+## PR Title Format
+
+All PR titles targeting `main` must start with a bracketed type tag. This is enforced by a
+Mergify merge protection rule and is required before a PR can be squash-merged.
+
+**Format:**
+
+```
+[type] Short description
+```
+
+**Valid type tags:**
+
+`feat`, `feature`, `bugfix`, `fix`, `refactor`, `perf`, `ci`, `doc`, `docs`, `misc`, `chore`,
+`kernel`, `new-model`
+
+**Valid examples:**
+
+```
+[feat] Add causal Wan pipeline
+[bugfix] Fix VAE temporal tiling corruption
+[refactor] Restructure training framework
+[perf] Optimize FlashAttention kernel dispatch
+[docs] Add inference guide for LoRA
+[new-model] Port HunyuanVideo 1.5 to FastVideo
+```
+
+**Invalid examples (will block merge):**
+
+```
+Add causal Wan pipeline          ← missing type tag
+FEAT: Add pipeline               ← wrong format
+feat: Add pipeline               ← square brackets required
+```
+
+If your title is invalid, Mergify posts a comment explaining the required format and the
+merge protection check will remain failed until you update the title.
+
+---
+
+## Slash Commands
+
+Slash commands let contributors and maintainers trigger CI actions directly from PR comments.
+**Write permission to the repository is required.**
+
+The command is recognized within a few seconds. The workflow reacts with a 🚀 emoji to confirm.
+
+### `/merge`
+
+```
+/merge
+```
+
+Adds the `ready` label to the PR, which triggers the Full Suite on your PR branch and
+enables Mergify to auto-squash-merge once all conditions pass.
+
+The command first removes the `ready` label if it is already present, then re-adds it. This
+ensures the `labeled` event fires and a fresh Full Suite build is started even on a re-try.
+
+### `/test <name>`
+
+```
+/test <name>
+```
+
+Triggers a specific Buildkite test or suite on the current PR branch.
+
+| Command | Runs | Maps to `TEST_TYPE` |
+|---------|------|---------------------|
+| `/test encoder` | Encoder Tests (Fastcheck) | `encoder` |
+| `/test vae` | VAE Tests (Fastcheck) | `vae` |
+| `/test transformer` | Transformer Tests (Fastcheck) | `transformer` |
+| `/test kernel` | Kernel Tests (Fastcheck) | `kernel_tests` |
+| `/test unit` | Unit Tests (Fastcheck) | `unit_test` |
+| `/test ssim` | SSIM regression tests | `ssim` |
+| `/test training` | Training pipeline tests | `training` |
+| `/test lora-inference` | LoRA inference tests | `inference_lora` |
+| `/test lora-training` | LoRA training tests | `training_lora` |
+| `/test distillation` | DMD distillation tests | `distillation_dmd` |
+| `/test self-forcing` | Self-Forcing tests | `self_forcing` |
+| `/test vsa` | VSA training tests | `training_vsa` |
+| `/test vmoba` | VMoBA inference tests | `inference_vmoba` |
+| `/test performance` | Performance benchmarks | `performance` |
+| `/test api` | API server integration tests | `api_server` |
+| `/test full` | Entire Full Suite | all (with `TEST_SCOPE=full`) |
+| `/test fastcheck` | Entire Fastcheck suite | fastcheck (with `TEST_SCOPE=fastcheck`) |
+| `/test pre-commit` | Pre-commit checks on PR code | — (runs `ci-precommit.yml` via `workflow_call`) |
+
+**Re-running failed tests:** When you use `/test <name>` to re-run a specific failed test,
+the resulting Buildkite check uses the same name as the original (e.g., `/test encoder`
+creates `buildkite/ci/microscope-encoder-tests`). This overwrites the failed check status.
+Once all tests in a tier pass, the aggregate status (`fastcheck-passed` or
+`full-suite-passed`) is automatically updated to `success` by the `ci-aggregate-status.yml`
+workflow.
+
+**How aggregate status refresh works:**
+
+1. `/test <name>` triggers a Buildkite build with `TEST_SCOPE=direct`. The test step uses
+   the same label as its fastcheck/full-suite counterpart, so the resulting GitHub check
+   overwrites the original.
+2. When the build completes, Buildkite's `notify` posts a `direct-test-completed` commit
+   status. This is the only signal that triggers the aggregate workflow — intermediate step
+   status updates do not trigger it.
+3. `ci-aggregate-status.yml` fires, calls `getCombinedStatusForRef` to fetch the latest
+   status for every context on that commit (each context returns only its most recent
+   state), groups them by prefix (`microscope-*` → fastcheck, `test-tube-*`/`bar-chart-*`
+   → full suite), and posts `fastcheck-passed: success` or `full-suite-passed: success` if
+   all entries in the group are `success`.
+4. Tests that were never triggered (skipped by monorepo-diff) have no status entry and do
+   not block the aggregate.
+
+---
+
+## Auto Branch Cleanup
+
+After a PR is squash-merged to `main`, Mergify automatically deletes the head branch.
+Protected branches (`main`, `master`, `release/*`) are never deleted.
+
+---
+
+## Workflow File Reference
+
+| Filename | Trigger | What it does |
+|----------|---------|-------------|
+| `ci-precommit.yml` | Every push / PR against `main` | Runs pre-commit hooks (yapf, ruff, mypy, codespell, pymarkdown, actionlint, check-filenames) |
+| `ci-trigger-full-suite.yml` | `ready` label added to a PR | Calls Buildkite API to run Full Suite on the PR branch |
+| `ci-slash-commands.yml` | PR comment starting with `/merge` or `/test` | Handles slash commands; adds `ready` label or triggers Buildkite |
+| `ci-aggregate-status.yml` | Any Buildkite commit status update | Checks if all tests in a tier passed; updates `fastcheck-passed` or `full-suite-passed` |
+| `community-issue-labeler.yml` | Issue opened or edited | Auto-labels issues by keyword matching against title and body |
+| `community-welcome.yml` | First contribution | Posts a welcome comment for first-time contributors |
+| `community-stale.yml` | Scheduled | Marks and closes stale issues and PRs |
+| `infra-build-image.yml` | Manual (`workflow_dispatch`) | Builds Docker images for CI |
+| `infra-docs.yml` | Changes to `docs/` merged to `main` | Builds and deploys documentation to GitHub Pages |
+| `publish-fastvideo.yml` | Version bump | Publishes `fastvideo` package to PyPI |
+| `publish-kernel.yml` | Version bump | Publishes `fastvideo-kernel` package to PyPI |
+| `publish-comfyui.yml` | Version bump | Publishes ComfyUI node package to PyPI |

docs/contributing/developer_env/runpod.md

@@ -1,65 +1,195 @@
+# RunPod Development Environment
 
-# 📦 Developing FastVideo on RunPod
+RunPod gives you on-demand cloud GPUs for FastVideo development. It's useful when you need a beefy GPU to test training runs, benchmark inference, or reproduce results without waiting for shared cluster time.
 
-You can easily use the FastVideo Pod Template on [RunPod](https://www.runpod.io) for development or experimentation.
+## Prerequisites
 
-## Creating a new pod
+- A [RunPod](https://www.runpod.io) account with billing configured
+- An SSH key pair. If you don't have one, generate it with `ssh-keygen -t ed25519`
+- Your public key (`~/.ssh/id_ed25519.pub`) ready to paste into RunPod
+
+## Step 1: Create a Pod
+
+**1. Verify your account**
+
+Make sure you're logged into the right RunPod account before spending credits.
 
-- Make sure you are using the correct RunPod account.
 ![RunPod Account Selection](../../assets/images/runpod_account.png)
 
-- Use "Additional Filters" to select CUDA 12.8.
+**2. Filter by CUDA version**
+
+Use "Additional Filters" to select CUDA 12.8.
+
 ![RunPod CUDA selection](../../assets/images/runpod_cuda.png)
 
-- Click "Deploy" and Pick a single A40 or RTX 4090 GPU.
+**3. Select a GPU**
+
+Click "Deploy" and pick a GPU. See [GPU Recommendations](#gpu-recommendations) below for guidance on which GPU to choose.
+
 ![RunPod GPU Selection](../../assets/images/runpod_deploy.png)
 
-- Select the "FastVideo" or "fastvideo-dev" Pod Template.
+**4. Pick the FastVideo template**
+
+Select the "FastVideo" or "fastvideo-dev" Pod Template. This pulls the pre-built image that includes all dependencies, Flash Attention, and a ready-to-use `uv` environment.
+
 ![RunPod Pod Template Selection](../../assets/images/runpod_create.png)
 
-- Set the Pod name to "`<name>-<FastVideo>-<date>`".
+**5. Name your pod**
+
+Use a memorable name like `yourname-fastvideo-2026-03-28`. This helps if you have multiple pods running.
+
+**6. Add a persistent volume (recommended)**
+
+Attach a network volume to `/root/.cache` or `/models` for storing downloaded model weights. Models can be 10-50 GB each, and re-downloading them every session wastes time and bandwidth.
+
+**7. Deploy**
+
+Click Deploy. The pod takes a few minutes to start while the image pulls. You'll see it transition to "Running" in your dashboard.
+
+## Step 2: Connect via SSH
+
+Once the pod is running, find the "SSH over exposed TCP" connection string in the pod dashboard.
 
-- Finally, once the pod is deployed (will take a few minutes as the image is being pulled), you can SSH into it using "SSH exposed over TCP". You'll need to use the matching private ssh key you provided.
 ![RunPod SSH](../../assets/images/runpod_ssh.png)
 
-## Working with the pod
-
-After SSH'ing into your pod, you'll find the correct `uv` environment already activated and you should be in /FastVideo/ directory. Make sure to use /FastVideo/ for all your work.
-
-To pull in the latest changes from the GitHub repo:
+Connect with:
 
 ```bash
-cd /FastVideo
-git pull
+ssh root@<pod-ip> -p <port> -i ~/.ssh/id_ed25519
 ```
 
-Run your development workflows as usual:
+RunPod also supports VS Code Remote SSH if you prefer an IDE.
 
-```bash
-# Run linters
-pre-commit run --all-files
+### Custom template (advanced)
 
-# Run tests
-pytest tests/
-```
-
-Make sure to push your changes back to the GitHub repo as nothing will be saved to the pod when it is terminated.
-
-After you are done with your work, you can terminate the pod by clicking the "Terminate" and "Delete" buttons. Remember if the pod is not completely deleted, Runpod will keep charging you for it.
-
-## Extra Information:
-If you need to customize the pod template this section has some useful information. For the most part you can leave the defaults of the FastVideo Pod Template.
-
-When creating your pod template, use this image:
+If you're setting up a pod from scratch instead of the FastVideo template, use this image:
 
 ```
 ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev:py3.12-latest
 ```
 
-Paste Container Start Command to support SSH ([RunPod Docs](https://docs.runpod.io/pods/configuration/use-ssh)):
+And paste this as the Container Start Command to enable SSH ([RunPod docs](https://docs.runpod.io/pods/configuration/use-ssh)):
 
 ```bash
 bash -c "apt update;DEBIAN_FRONTEND=noninteractive apt-get install openssh-server -y;mkdir -p ~/.ssh;cd $_;chmod 700 ~/.ssh;echo \"$PUBLIC_KEY\" >> authorized_keys;chmod 700 authorized_keys;service ssh start;sleep infinity"
 ```
 
 ![RunPod template configuration](../../assets/images/runpod_template.png)
+
+## Step 3: Set Up FastVideo
+
+After SSH'ing in, the `uv` virtual environment at `/opt/venv` is already activated (configured in `.bashrc` and `.profile`). You land in the `/FastVideo` directory.
+
+**Clone or pull the repo**
+
+If the pod already has the FastVideo repo:
+
+```bash
+cd /FastVideo
+git pull
+```
+
+If starting from a blank pod:
+
+```bash
+git clone https://github.com/hao-ai-lab/FastVideo.git /FastVideo
+cd /FastVideo
+```
+
+**Install the package**
+
+```bash
+uv pip install -e .[dev]
+```
+
+The Docker image already includes Flash Attention and most heavy dependencies, so this is fast.
+
+**Build the custom kernels (optional)**
+
+VSA and STA attention kernels aren't in the Docker image by default. Build them if you're working on attention backends or need maximum inference performance:
+
+```bash
+cd /FastVideo/fastvideo-kernel
+./build.sh
+```
+
+The build script detects your GPU architecture automatically. An A100 or H100 takes about 5-10 minutes.
+
+**Verify the setup**
+
+```bash
+cd /FastVideo
+python -c "import fastvideo; print('OK')"
+pytest tests/ -q --no-header
+```
+
+## Development Workflow
+
+### Editing code on RunPod
+
+Two common approaches:
+
+**Option A: Edit on RunPod directly**
+
+Use VS Code Remote SSH or `vim`/`nano` on the pod. Commit and push when you're ready:
+
+```bash
+cd /FastVideo
+git add .
+git commit -m "your change"
+git push
+```
+
+**Option B: Edit locally, sync to RunPod**
+
+Work in your local repo, then pull on the pod:
+
+```bash
+# On RunPod:
+cd /FastVideo
+git pull
+```
+
+This keeps your local tools (editor, linters) intact while running GPU workloads on the pod.
+
+### Running linters and tests
+
+```bash
+# Lint
+pre-commit run --all-files
+
+# Full test suite
+pytest tests/
+
+# Just package tests
+pytest fastvideo/tests/ -v
+```
+
+### Storing models
+
+If you attached a persistent volume, point your model downloads there:
+
+```bash
+export HF_HOME=/models/huggingface
+export TRANSFORMERS_CACHE=/models/huggingface
+```
+
+Add these to `/root/.bashrc` so they persist across SSH sessions. The volume survives pod termination, so you only download models once.
+
+### Terminating the pod
+
+When you're done, push any commits you want to keep. RunPod does not save pod storage after termination.
+
+Go to your RunPod dashboard, click "Terminate", then "Delete". A pod that's stopped but not deleted still charges you for storage. Fully delete it to stop all charges.
+
+## GPU Recommendations
+
+| GPU | VRAM | Good for |
+|-----|------|----------|
+| RTX 4090 | 24 GB | Inference testing, small model fine-tuning, quick iteration |
+| A40 | 48 GB | Mid-size training runs, 480p video generation |
+| A100 (40 GB) | 40 GB | Multi-GPU inference, training with sequence parallelism |
+| A100 (80 GB) | 80 GB | Large model training, 720p+ video generation |
+| H100 | 80 GB | Heavy training, benchmarking, kernel development |
+
+For most development work, a single RTX 4090 or A40 is sufficient and cost-effective. Use A100/H100 when you need to reproduce training results at scale or test multi-GPU features.

docs/contributing/pull_requests.md

@@ -0,0 +1,223 @@
+# Contributing via Pull Requests
+
+This guide walks through the PR workflow: title format, labels, CI pipeline, and getting
+your changes merged.
+
+---
+
+## PR Title Format (Required)
+
+Every PR targeting `main` must start with a type tag in square brackets. This is checked by
+Mergify before any merge is allowed.
+
+**Format:**
+
+```
+[type] Short description of the change
+```
+
+**Valid type tags:**
+
+| Tag | When to use |
+|-----|-------------|
+| `[feat]` or `[feature]` | New feature or capability |
+| `[bugfix]` or `[fix]` | Bug fix |
+| `[refactor]` | Code restructuring with no behavior change |
+| `[perf]` | Performance improvement |
+| `[ci]` | CI/CD or build tooling changes |
+| `[doc]` or `[docs]` | Documentation only |
+| `[misc]` or `[chore]` | Housekeeping, dependency bumps, minor cleanup |
+| `[kernel]` | CUDA kernel changes in `fastvideo-kernel/` |
+| `[new-model]` | Adding a new model or pipeline |
+
+**Examples:**
+
+```
+[feat] Add causal Wan 2.2 I2V pipeline
+[bugfix] Fix VAE temporal tiling corruption on H100
+[refactor] Restructure distributed attention dispatch
+[docs] Add LoRA finetuning guide
+[new-model] Port HunyuanVideo 1.5 to FastVideo
+```
+
+If your title is missing the tag, Mergify will post a comment listing the valid formats.
+Update the title and the check will re-evaluate automatically.
+
+---
+
+## Labels
+
+Labels are applied automatically based on your PR title and the files you changed. You don't
+need to set them manually.
+
+**Type label** — set from the `[tag]` in your title:
+`type: feat`, `type: bugfix`, `type: refactor`, `type: perf`, `type: ci`, `type: docs`,
+`type: misc`, `type: new-model`
+
+**Scope labels** — set from which files you modified (multiple labels can apply):
+`scope: training`, `scope: inference`, `scope: attention`, `scope: kernel`, `scope: data`,
+`scope: infra`, `scope: distributed`, `scope: docs`, `scope: ui`, `scope: model`
+
+**Process labels** — set during review and merge:
+
+| Label | Who sets it | Meaning |
+|-------|-------------|---------|
+| `ready` | You (`/merge` comment) or a maintainer | Triggers Full Suite and enables auto-merge |
+| `needs-rebase` | Mergify (automatic) | Your PR has conflicts; rebase against `main` |
+| `do-not-merge` | Maintainer | Blocks merge regardless of CI status |
+
+---
+
+## CI Pipeline
+
+Three tiers run automatically on every PR.
+
+**Tier 1: Pre-commit (~2 min) — runs on every push**
+
+GitHub Actions checks formatting, linting, type correctness, and spelling using pre-commit
+hooks: yapf, ruff, mypy, codespell, pymarkdown, actionlint, and check-filenames.
+
+**Tier 2: Fastcheck (~10-20 min) — runs on every push, path-filtered**
+
+Buildkite runs GPU tests only for the components you changed. If you only modified
+`fastvideo/models/vaes/`, only VAE Tests run. Tests run in parallel.
+
+**Tier 3: Full Suite (~60-90 min) — triggered by the `ready` label**
+
+When you comment `/merge` (or a maintainer adds the `ready` label), Buildkite runs the
+complete test suite on your PR branch: SSIM regression, LoRA inference and training,
+distillation, self-forcing, VSA, VMoBA, performance benchmarks, and API server tests.
+
+---
+
+## Getting Your PR Merged
+
+**Step-by-step:**
+
+1. Open a PR with a title that starts with a valid `[type]` tag.
+2. Push your changes. Pre-commit and Fastcheck run automatically.
+3. Fix any pre-commit failures locally (`pre-commit run --all-files`) and push again.
+4. Wait for at least one approving review.
+5. Once approved and pre-commit is green, comment `/merge` on the PR.
+6. The `ready` label is added, which triggers the Full Suite on your PR branch.
+7. Mergify also auto-rebases your branch against `main` if it is behind and conflict-free.
+8. If all Full Suite tests pass and all merge conditions are met (approval, valid title,
+   pre-commit green, fastcheck green, no draft, no conflicts), Mergify squash-merges to
+   `main` automatically. Your branch is deleted.
+9. If a Full Suite test fails, check the Buildkite build log for the failing step. Fix the
+   issue, push, and comment `/merge` again. You can also re-run individual failed tests
+   with `/test <name>` — see below.
+
+!!! note
+    Only contributors with write permission to the repository can trigger slash commands.
+    If you're an external contributor, ask a maintainer to run `/merge` or add the `ready`
+    label for you.
+
+---
+
+## Running Tests On Demand
+
+Comment on your PR to trigger specific tests independently of the auto-merge flow.
+
+**Trigger the entire Full Suite:**
+
+```
+/test full
+```
+
+**Trigger the Fastcheck suite:**
+
+```
+/test fastcheck
+```
+
+**Trigger individual tests:**
+
+```
+/test encoder          # Encoder component tests
+/test vae              # VAE component tests
+/test transformer      # Transformer / DiT tests
+/test kernel           # CUDA kernel tests
+/test unit             # Unit tests
+
+/test ssim             # SSIM regression tests
+/test training         # Training pipeline tests
+/test lora-inference   # LoRA inference tests
+/test lora-training    # LoRA training tests
+/test distillation     # DMD distillation tests
+/test self-forcing     # Self-Forcing distillation tests
+/test vsa              # VSA training tests
+/test vmoba            # VMoBA inference tests
+/test performance      # Performance benchmarks
+/test api              # API server integration tests
+/test pre-commit       # Pre-commit checks on PR code
+```
+
+The workflow reacts with a 🚀 emoji to confirm the command was received.
+
+When you re-run an individual test with `/test <name>`, the new result overwrites the
+original failed check (same Buildkite check name). Once all tests in a tier pass, the
+`fastcheck-passed` or `full-suite-passed` status is automatically updated.
+
+---
+
+## Troubleshooting
+
+### Pre-commit fails
+
+Run locally to reproduce and auto-fix:
+
+```bash
+# Install pre-commit if needed
+uv pip install pre-commit
+pre-commit install
+
+# Run all checks on all files
+pre-commit run --all-files
+```
+
+Common quick fixes:
+
+- **yapf**: `yapf -i <file>` (Python formatting)
+- **ruff**: `ruff check --fix <file>` (linting)
+- **codespell**: `codespell --write-changes <file>` (spelling)
+
+### PR title format check fails
+
+Update your title to start with a valid type tag. The Mergify merge protection check
+re-evaluates automatically after you save the title.
+
+Valid tags: `feat`, `feature`, `bugfix`, `fix`, `refactor`, `perf`, `ci`, `doc`, `docs`,
+`misc`, `chore`, `kernel`, `new-model`
+
+### My PR has merge conflicts (`needs-rebase` label)
+
+Rebase against `main` and force-push:
+
+```bash
+git fetch origin main
+git rebase origin/main
+# Resolve any conflicts, then:
+git push --force-with-lease
+```
+
+Mergify removes the `needs-rebase` label automatically once conflicts are resolved.
+
+### Full Suite failed after `/merge`
+
+The Full Suite found a regression. Check the failing Buildkite step's output for assertion
+errors or tracebacks.
+
+Common causes:
+
+- Test failures caused by your code changes
+- Missing dependency in `pyproject.toml`
+- GPU memory issue (some tests require specific hardware like L40S or H100)
+- Kernel build failure (if you changed `fastvideo-kernel/`)
+
+After fixing, push and comment `/merge` again.
+
+### I'm an external contributor without write permission
+
+You can't use slash commands directly. After your PR is approved, ask a maintainer to
+comment `/merge` or add the `ready` label.

docs/contributing/testing.md

@@ -23,13 +23,18 @@ SSIM tests are located in `fastvideo/tests/ssim`. These tests generate videos us
 
 ```
 fastvideo/tests/ssim/
-├── <GPU>_reference_videos/   # Reference videos organized by GPU type (e.g., L40S_reference_videos)
-│   ├── <Model_Name>/
-│   │   ├── <Backend>/        # e.g., FLASH_ATTN, TORCH_SDPA
-│   │   │   └── <Video_File>
+├── reference_videos/
+│   ├── default/
+│   │   └── <GPU>_reference_videos/
+│   │       ├── <Model_Name>/
+│   │       │   ├── <Backend>/        # e.g., FLASH_ATTN, TORCH_SDPA
+│   │       │   │   └── <Video_File>
+│   └── full_quality/
+│       └── <GPU>_reference_videos/
 ├── test_causal_similarity.py
-├── test_inference_similarity.py
-├── update_reference_videos.sh
+├── test_wan_t2v_similarity.py
+├── test_wan_i2v_similarity.py
+├── reference_videos_cli.py
 └── ...
 ```
 
@@ -37,7 +42,7 @@ fastvideo/tests/ssim/
 
 To add a new SSIM test, follow these steps:
 
-1. **Create or Update a Test File**: You can add a new test function to an existing file (like `test_inference_similarity.py`) or create a new one if testing a distinct category of models.
+1. **Create or Update a Test File**: Prefer model-specific files (for example `test_wan_t2v_similarity.py`) and create a new one when testing a distinct model or pipeline.
 
 2. **Define Model Parameters**: Define the configuration for the model you want to test. This includes model path, dimensions, inference steps, and other generation parameters. **Note:** Consider using lower `num_inference_steps` or reduced resolution (e.g., 480p instead of 720p) to keep test execution time reasonable, provided it doesn't compromise the test's ability to detect regression.
 
@@ -81,10 +86,15 @@ To add a new SSIM test, follow these steps:
    ```
 
 4. **Reference Videos**:
-   * When running the test for the first time (or when updating the reference), the test will fail because the reference video is missing. The generated video will be saved in `fastvideo/tests/ssim/generated_videos`.
+   * When running the test for the first time (or when updating the reference), the test will fail because the reference video is missing. The generated video will be saved in `fastvideo/tests/ssim/generated_videos/<quality-tier>/<GPU>_reference_videos`.
    * Inspect the generated video to ensure it meets quality expectations.
-   * Move the generated video to the appropriate reference folder: `fastvideo/tests/ssim/<GPU>_reference_videos/<Model>/<Backend>/`.
-   * You can use the helper script `update_reference_videos.sh` to automate copying videos from `generated_videos` to `L40S_reference_videos`. Note: Check the script to ensure paths match your environment (it defaults to `L40S_reference_videos`).
+   * Move the generated video to the appropriate quality/GPU reference folder:
+     `fastvideo/tests/ssim/reference_videos/<quality-tier>/<GPU>_reference_videos/<Model>/<Backend>/`.
+   * You can use the helper CLI to copy generated videos into a reference folder:
+     `python fastvideo/tests/ssim/reference_videos_cli.py copy-local --quality-tier default --reference-dir fastvideo/tests/ssim/reference_videos/default/L40S_reference_videos`
+   * Upload/download can target both quality tiers and specific GPU folders:
+     `python fastvideo/tests/ssim/reference_videos_cli.py upload --quality-tier all`
+     `python fastvideo/tests/ssim/reference_videos_cli.py download --quality-tier full_quality --device-folder H200_reference_videos`
 
 ### Running Tests Locally
 
@@ -96,36 +106,113 @@ pytest fastvideo/tests/ssim/ -vs
 
 Ensure you have the necessary GPUs available as defined in your test parameters.
 
-## Modal Workflow
+## CI Integration
 
-FastVideo uses [Modal](https://modal.com/) for running tests in a CI environment. The workflow scripts are located in `fastvideo/tests/modal/`.
+FastVideo uses [Modal](https://modal.com/) for running tests in a CI environment. The
+workflow scripts are located in `fastvideo/tests/modal/`.
+
+### Buildkite Pipeline
+
+Tests are orchestrated by Buildkite (`.buildkite/pipeline.yml`) and executed on Modal GPU
+instances. The pipeline runs in two modes:
+
+**Fastcheck** — runs on every PR push, path-filtered. Only tests for the components you
+changed are triggered. Tests run in parallel.
+
+**Full Suite** — runs when a PR enters the Merge Queue (or when triggered manually via
+`/test full`). Covers SSIM regression, training, distillation, inference, and performance.
 
 ### `pr_test.py`
 
-The main entry point for CI tests is `fastvideo/tests/modal/pr_test.py`. This script defines Modal functions that execute the pytest suites on specific hardware (e.g., L40S, H100).
+The main entry point for CI tests is `fastvideo/tests/modal/pr_test.py`. This script defines
+Modal functions that execute the pytest suites on specific hardware (e.g., L40S, H100).
 
 ### Updating Modal Configuration
 
 If you add a new test that requires:
+
 * **Different GPU Hardware**: You may need to change the `@app.function(gpu=...)` decorator.
 * **Longer Execution Time**: Increase the `timeout` parameter.
-* **New Environment Variables/Secrets**: Add them to `secrets=[...]` or the image environment. For example, if your model is gated on Hugging Face, ensure `HF_API_KEY` is passed.
+* **New Environment Variables/Secrets**: Add them to `secrets=[...]` or the image
+  environment. For example, if your model is gated on Hugging Face, ensure `HF_API_KEY`
+  is passed.
 
-For SSIM tests, the `run_ssim_tests` function in `pr_test.py` currently runs:
+For SSIM tests, use `fastvideo/tests/modal/ssim_test.py`:
 
-```python
-@app.function(gpu="L40S:2", image=image, timeout=2700, secrets=[modal.Secret.from_dict({"HF_API_KEY": os.environ.get("HF_API_KEY", "")})])
-def run_ssim_tests():
-    run_test("hf auth login --token $HF_API_KEY && pytest ./fastvideo/tests/ssim -vs")
+```bash
+python -m modal run fastvideo/tests/modal/ssim_test.py::run_ssim_tests
 ```
 
-If your new test file is inside `fastvideo/tests/ssim`, it will automatically be picked up by this command. However, ensure that the `gpu="L40S:2"` configuration is sufficient for your model. If your model requires more GPUs (e.g., 4 or 8), you might need to create a separate Modal function or update the existing one.
+Target specific SSIM files/models:
+
+```bash
+python -m modal run fastvideo/tests/modal/ssim_test.py::run_ssim_tests \
+  --test-files test_wan_t2v_similarity.py \
+  --model-ids Wan2.1-T2V-1.3B-Diffusers
+```
+
+If HF token env vars are not set (`HF_API_KEY` / `HUGGINGFACE_HUB_TOKEN` /
+`HF_TOKEN`), the local entrypoint fails fast. To export raw `generated_videos`
+from Modal to the shared volume:
+
+```bash
+python -m modal run fastvideo/tests/modal/ssim_test.py::run_ssim_tests \
+  --sync-generated-to-volume
+```
+
+The raw export path is quality-tiered:
+
+* default params: `ssim_generated_videos/default/<subdir>/generated_videos`
+* full-quality params: `ssim_generated_videos/full_quality/<subdir>/generated_videos`
+
+The printed `modal volume get` command also downloads into a quality-specific
+local directory under `./generated_videos_modal/<quality-tier>`.
+
+To turn downloaded Modal outputs into local reference videos, use the matching
+quality tier with `copy-local`, for example:
+
+```bash
+python fastvideo/tests/ssim/reference_videos_cli.py copy-local \
+  --quality-tier full_quality \
+  --generated-dir ./generated_videos_modal/full_quality/L40S_reference_videos \
+  --device-folder L40S_reference_videos
+```
 
 ### Workflow Scripts
 
-The shell script that triggers these tests in the CI pipeline is located at `.buildkite/scripts/pr_test.sh`. If you add a new test category (e.g., a new folder outside of `ssim`), you will need to:
+The shell script that triggers tests in CI is `.buildkite/scripts/pr_test.sh`. If you add
+a new test category (e.g., a new folder outside of `ssim`), you will need to:
+
 1. Add a new function in `fastvideo/tests/modal/pr_test.py`.
 2. Add a new case in `.buildkite/scripts/pr_test.sh` to handle the new test type.
 
 !!! note
-    If you are a maintainer, you'll need to finally manually update the workflow script in Buildkite. Otherwise, a maintainer will help you update.
+    If you are a maintainer, update the workflow script in Buildkite after merging. Otherwise,
+    ask a maintainer for help.
+
+## Triggering Tests via Slash Commands
+
+Maintainers and contributors with write access can trigger individual test suites directly
+from a PR comment. The workflow reacts with a 🚀 emoji to confirm the command was received.
+
+```
+/test ssim              # SSIM regression tests
+/test training          # Training pipeline tests
+/test lora-training     # LoRA training tests
+/test lora-inference    # LoRA inference tests
+/test distillation      # DMD distillation tests
+/test self-forcing      # Self-Forcing tests
+/test vsa               # VSA training tests
+/test vmoba             # VMoBA inference tests
+/test performance       # Performance benchmarks
+/test api               # API server integration tests
+/test encoder           # Encoder component tests (Fastcheck)
+/test vae               # VAE component tests (Fastcheck)
+/test transformer       # Transformer / DiT tests (Fastcheck)
+/test kernel            # CUDA kernel tests (Fastcheck)
+/test unit              # Unit tests (Fastcheck)
+/test full              # Entire Full Suite
+/test fastcheck         # Entire Fastcheck suite
+```
+
+See [CI Architecture](ci_architecture.md) for the complete reference.

docs/design/inference_schema_parity_inventory.yaml

@@ -0,0 +1,720 @@
+status_definitions:
+  kept: "Public field remains on a public adapter surface with the same meaning."
+  moved: "Public field remains supported but normalizes into a different nested path."
+  profile_owned: "Public field remains supported only through a model/profile-specific surface."
+  compatibility_only: "Legacy public field remains adapter-only during migration and is not part of the canonical typed schema."
+  private_only: "Field should only be handled by private adapters and is not a public FastVideo compatibility promise."
+  internal_only: "Field is runtime/config plumbing and should not be part of the new public typed inference API."
+
+surfaces:
+  fastvideo_args:
+    moved:
+      model_path: generator.model_path
+      workload_type: generator.pipeline.workload_type
+      distributed_executor_backend: generator.engine.execution_backend
+      trust_remote_code: generator.trust_remote_code
+      revision: generator.revision
+      num_gpus: generator.engine.num_gpus
+      tp_size: generator.engine.parallelism.tp_size
+      sp_size: generator.engine.parallelism.sp_size
+      hsdp_replicate_dim: generator.engine.parallelism.hsdp_replicate_dim
+      hsdp_shard_dim: generator.engine.parallelism.hsdp_shard_dim
+      dist_timeout: generator.engine.parallelism.dist_timeout
+      lora_path: generator.pipeline.components.lora_path
+      dit_cpu_offload: generator.engine.offload.dit
+      use_fsdp_inference: generator.engine.use_fsdp_inference
+      dit_layerwise_offload: generator.engine.offload.dit_layerwise
+      text_encoder_cpu_offload: generator.engine.offload.text_encoder
+      image_encoder_cpu_offload: generator.engine.offload.image_encoder
+      vae_cpu_offload: generator.engine.offload.vae
+      pin_cpu_memory: generator.engine.offload.pin_cpu_memory
+      enable_torch_compile: generator.engine.compile.enabled
+      torch_compile_kwargs: generator.engine.compile.kwargs
+      disable_autocast: generator.engine.disable_autocast
+      enable_stage_verification: generator.engine.enable_stage_verification
+      prompt_txt: request.inputs.prompt_path
+      override_text_encoder_safetensors: generator.pipeline.components.text_encoder_weights
+      override_text_encoder_quant: generator.engine.quantization.text_encoder_quant
+      transformer_quant: generator.engine.quantization.transformer_quant
+      override_transformer_cls_name: generator.pipeline.components.override_transformer_cls_name
+      init_weights_from_safetensors: generator.pipeline.components.transformer_weights
+      init_weights_from_safetensors_2: generator.pipeline.components.transformer_2_weights
+      override_pipeline_cls_name: generator.pipeline.components.override_pipeline_cls_name
+      boundary_ratio: request.sampling.boundary_ratio
+    profile_owned:
+      ltx2_vae_tiling: generator.pipeline.profile_overrides.ltx2.vae_tiling
+      ltx2_vae_spatial_tile_size_in_pixels: generator.pipeline.profile_overrides.ltx2.vae.spatial_tile_size_in_pixels
+      ltx2_vae_spatial_tile_overlap_in_pixels: generator.pipeline.profile_overrides.ltx2.vae.spatial_tile_overlap_in_pixels
+      ltx2_vae_temporal_tile_size_in_frames: generator.pipeline.profile_overrides.ltx2.vae.temporal_tile_size_in_frames
+      ltx2_vae_temporal_tile_overlap_in_frames: generator.pipeline.profile_overrides.ltx2.vae.temporal_tile_overlap_in_frames
+      ltx2_initial_latent_path: request.extensions.ltx2.initial_latent_path
+    compatibility_only:
+      mode: "Legacy multi-mode FastVideoArgs switch; typed inference config should not expose execution mode."
+      inference_mode: "Legacy boolean mirror of mode; kept only through adapters while FastVideoArgs remains."
+      lora_nickname: "Legacy adapter-selection surface pending LoRA API cleanup."
+      lora_target_modules: "Legacy LoRA configuration surface pending dedicated component API."
+      output_type: "Legacy output formatting surface pending GenerationResult cleanup."
+      VSA_sparsity: "Model-specific inference optimization not yet represented in the typed public schema."
+      moba_config_path: "Model-specific MoBA optimization surface not yet represented in the typed public schema."
+      master_port: "Executor/bootstrap compatibility field; not part of the canonical inference schema."
+    private_only:
+      ray_placement_group: "Ray deployment-only field."
+      ray_runtime_env: "Ray deployment-only field."
+    internal_only:
+      pipeline_config: "Legacy internal carrier object."
+      preprocess_config: "Legacy preprocess carrier object."
+      moba_config: "Derived runtime config loaded from moba_config_path."
+      model_paths: "Runtime bookkeeping."
+      model_loaded: "Runtime bookkeeping."
+
+  pipeline_config_base:
+    moved:
+      pipeline_config_path: generator.pipeline.components.pipeline_config_path
+    profile_owned:
+      embedded_cfg_scale: generator.pipeline.profile_overrides.embedded_cfg_scale
+      flow_shift: generator.pipeline.profile_overrides.flow_shift
+      flow_shift_sr: generator.pipeline.profile_overrides.flow_shift_sr
+      is_causal: generator.pipeline.profile_overrides.is_causal
+      vae_tiling: generator.pipeline.profile_overrides.vae_tiling
+      vae_sp: generator.pipeline.profile_overrides.vae_sp
+      dmd_denoising_steps: generator.pipeline.profile_overrides.dmd_denoising_steps
+      ti2v_task: generator.pipeline.profile_overrides.ti2v_task
+      boundary_ratio: generator.pipeline.profile_overrides.boundary_ratio
+    compatibility_only:
+      model_path: "Redundant with generator.model_path."
+      disable_autocast: "Duplicated by generator.engine.disable_autocast during migration."
+      dit_precision: "Precision override pending dedicated typed component precision design."
+      upsampler_precision: "Precision override pending dedicated typed component precision design."
+      vae_precision: "Precision override pending dedicated typed component precision design."
+      image_encoder_precision: "Precision override pending dedicated typed component precision design."
+      text_encoder_precisions: "Precision override pending dedicated typed component precision design."
+    internal_only:
+      dit_config: "Legacy internal component config object."
+      upsampler_config: "Legacy internal component config object."
+      vae_config: "Legacy internal component config object."
+      image_encoder_config: "Legacy internal component config object."
+      text_encoder_configs: "Legacy internal component config object."
+      preprocess_text_funcs: "Internal text preprocessing hooks."
+      postprocess_text_funcs: "Internal text postprocessing hooks."
+
+  pipeline_config_extensions:
+    profile_owned:
+      conditioning_strategy:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      max_num_conditional_frames:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      min_num_conditional_frames:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      sigma_conditional:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      sigma_data:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      state_ch:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      state_t:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      text_encoder_class:
+        sources:
+          - fastvideo.configs.pipelines.cosmos.CosmosConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      autoregressive_chunk_frames:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      autoregressive_overlap_frames:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      cfg_behavior:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      default_camera_rotation:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      default_movement_distance:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      default_negative_prompt:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      default_trajectory_type:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      filter_points_threshold:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      fps:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      frame_buffer_max:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      moge_model_name:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      noise_aug_strength:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      num_frames:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      offload_moge_after_depth:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      use_moge_depth:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      video_resolution:
+        sources:
+          - fastvideo.configs.pipelines.gen3c.Gen3CConfig
+          - fastvideo.configs.pipelines.gen3c.Gen3CInferenceConfig
+      text_encoder_crop_start:
+        sources:
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15I2V480PStepDistilledConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15I2V720PConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15SR1080PConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15T2V480PConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15T2V720PConfig
+          - fastvideo.configs.pipelines.hyworld.HYWorldConfig
+          - fastvideo.configs.pipelines.hyworld.Hunyuan15T2V480PConfig
+      text_encoder_max_lengths:
+        sources:
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15I2V480PStepDistilledConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15I2V720PConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15SR1080PConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15T2V480PConfig
+          - fastvideo.configs.pipelines.hunyuan15.Hunyuan15T2V720PConfig
+          - fastvideo.configs.pipelines.hyworld.HYWorldConfig
+          - fastvideo.configs.pipelines.hyworld.Hunyuan15T2V480PConfig
+      precision:
+        sources:
+          - fastvideo.configs.pipelines.lingbotworld.LingBotWorldI2V480PConfig
+          - fastvideo.configs.pipelines.lingbotworld.Wan2_2_I2V_A14B_Config
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionI2VConfig
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionI2V_A14B_Config
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionT2VConfig
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionT2V_14B_Config
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionT2V_1_3B_Config
+          - fastvideo.configs.pipelines.wan.FastWan2_1_T2V_480P_Config
+          - fastvideo.configs.pipelines.wan.FastWan2_2_TI2V_5B_Config
+          - fastvideo.configs.pipelines.wan.MatrixGameBaseI2V480PConfig
+          - fastvideo.configs.pipelines.wan.MatrixGameI2V480PConfig
+          - fastvideo.configs.pipelines.wan.SelfForcingWan2_2_T2V480PConfig
+          - fastvideo.configs.pipelines.wan.SelfForcingWanT2V480PConfig
+          - fastvideo.configs.pipelines.wan.WANV2VConfig
+          - fastvideo.configs.pipelines.wan.Wan2_2_I2V_A14B_Config
+          - fastvideo.configs.pipelines.wan.Wan2_2_T2V_A14B_Config
+          - fastvideo.configs.pipelines.wan.Wan2_2_TI2V_5B_Config
+          - fastvideo.configs.pipelines.wan.WanI2V480PConfig
+          - fastvideo.configs.pipelines.wan.WanI2V720PConfig
+          - fastvideo.configs.pipelines.wan.WanT2V480PConfig
+          - fastvideo.configs.pipelines.wan.WanT2V720PConfig
+      warp_denoising_step:
+        sources:
+          - fastvideo.configs.pipelines.lingbotworld.LingBotWorldI2V480PConfig
+          - fastvideo.configs.pipelines.lingbotworld.Wan2_2_I2V_A14B_Config
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionI2VConfig
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionI2V_A14B_Config
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionT2VConfig
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionT2V_14B_Config
+          - fastvideo.configs.pipelines.turbodiffusion.TurboDiffusionT2V_1_3B_Config
+          - fastvideo.configs.pipelines.wan.FastWan2_1_T2V_480P_Config
+          - fastvideo.configs.pipelines.wan.FastWan2_2_TI2V_5B_Config
+          - fastvideo.configs.pipelines.wan.MatrixGameBaseI2V480PConfig
+          - fastvideo.configs.pipelines.wan.MatrixGameI2V480PConfig
+          - fastvideo.configs.pipelines.wan.SelfForcingWan2_2_T2V480PConfig
+          - fastvideo.configs.pipelines.wan.SelfForcingWanT2V480PConfig
+          - fastvideo.configs.pipelines.wan.WANV2VConfig
+          - fastvideo.configs.pipelines.wan.Wan2_2_I2V_A14B_Config
+          - fastvideo.configs.pipelines.wan.Wan2_2_T2V_A14B_Config
+          - fastvideo.configs.pipelines.wan.Wan2_2_TI2V_5B_Config
+          - fastvideo.configs.pipelines.wan.WanI2V480PConfig
+          - fastvideo.configs.pipelines.wan.WanI2V720PConfig
+          - fastvideo.configs.pipelines.wan.WanT2V480PConfig
+          - fastvideo.configs.pipelines.wan.WanT2V720PConfig
+      bsa_cdf_threshold:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      bsa_chunk_k:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      bsa_chunk_q:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      bsa_params:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      bsa_sparsity:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      enable_bsa:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      enable_kv_cache:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      enhance_hf:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      offload_kv_cache:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      t_thresh:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      use_distill:
+        sources: [fastvideo.configs.pipelines.longcat.LongCatT2V480PConfig, fastvideo.configs.pipelines.longcat.LongCatT2V704PConfig]
+      scheduler_arch:
+        sources: [fastvideo.configs.pipelines.sd35.SD35Config]
+      text_encoder_archs:
+        sources: [fastvideo.configs.pipelines.sd35.SD35Config]
+      tokenizer_archs:
+        sources: [fastvideo.configs.pipelines.sd35.SD35Config]
+      transformer_arch:
+        sources: [fastvideo.configs.pipelines.sd35.SD35Config]
+      vae_arch:
+        sources: [fastvideo.configs.pipelines.sd35.SD35Config]
+      expand_timesteps:
+        sources:
+          - fastvideo.configs.pipelines.wan.FastWan2_2_TI2V_5B_Config
+          - fastvideo.configs.pipelines.wan.Wan2_2_TI2V_5B_Config
+      context_noise:
+        sources: [fastvideo.configs.pipelines.wan.MatrixGameI2V480PConfig]
+      num_frames_per_block:
+        sources: [fastvideo.configs.pipelines.wan.MatrixGameI2V480PConfig]
+    compatibility_only:
+      batch_size: "Gen3C inference-only tuning field pending typed batching design."
+      gradient_checkpointing: "Gen3C inference-only compatibility field pending typed batching design."
+      guidance_scale: "Gen3C pipeline-level default pending profile/default-request cleanup."
+      num_inference_steps: "Gen3C pipeline-level default pending profile/default-request cleanup."
+    internal_only:
+      audio_decoder_config: "Legacy internal component config object."
+      audio_decoder_precision: "Precision override pending dedicated component precision design."
+      vocoder_config: "Legacy internal component config object."
+      vocoder_precision: "Precision override pending dedicated component precision design."
+
+  sampling_param_base:
+    moved:
+      image_path: request.inputs.image_path
+      pil_image: request.inputs.pil_image
+      video_path: request.inputs.video_path
+      mouse_cond: request.inputs.mouse_cond
+      keyboard_cond: request.inputs.keyboard_cond
+      grid_sizes: request.inputs.grid_sizes
+      pose: request.inputs.pose
+      c2ws_plucker_emb: request.inputs.c2ws_plucker_emb
+      refine_from: request.inputs.refine_from
+      stage1_video: request.inputs.stage1_video
+      prompt: request.prompt
+      negative_prompt: request.negative_prompt
+      prompt_path: request.inputs.prompt_path
+      output_path: request.output.output_path
+      output_video_name: request.output.output_video_name
+      num_videos_per_prompt: request.sampling.num_videos_per_prompt
+      seed: request.sampling.seed
+      num_frames: request.sampling.num_frames
+      height: request.sampling.height
+      width: request.sampling.width
+      height_sr: request.sampling.height_sr
+      width_sr: request.sampling.width_sr
+      fps: request.sampling.fps
+      num_inference_steps: request.sampling.num_inference_steps
+      num_inference_steps_sr: request.sampling.num_inference_steps_sr
+      guidance_scale: request.sampling.guidance_scale
+      guidance_scale_2: request.sampling.guidance_scale_2
+      guidance_rescale: request.sampling.guidance_rescale
+      boundary_ratio: request.sampling.boundary_ratio
+      sigmas: request.sampling.sigmas
+      enable_teacache: request.runtime.enable_teacache
+      save_video: request.output.save_video
+      return_frames: request.output.return_frames
+      return_trajectory_latents: request.runtime.return_trajectory_latents
+      return_trajectory_decoded: request.runtime.return_trajectory_decoded
+    profile_owned:
+      t_thresh: request.stage_overrides.refine.t_thresh
+      spatial_refine_only: request.stage_overrides.refine.spatial_refine_only
+      num_cond_frames: request.stage_overrides.refine.num_cond_frames
+      trajectory_type: request.extensions.gen3c.trajectory_type
+      movement_distance: request.extensions.gen3c.movement_distance
+      camera_rotation: request.extensions.gen3c.camera_rotation
+    internal_only:
+      data_type: "Derived from the request shape and not a public input."
+
+  sampling_param_extensions:
+    moved: {}
+    profile_owned:
+      action_list:
+        target: request.extensions.hunyuangamecraft.action_list
+        sources:
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraftSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft65FrameSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft129FrameSamplingParam
+      action_speed_list:
+        target: request.extensions.hunyuangamecraft.action_speed_list
+        sources:
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraftSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft65FrameSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft129FrameSamplingParam
+      camera_states:
+        target: request.extensions.hunyuangamecraft.camera_states
+        sources:
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraftSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft65FrameSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft129FrameSamplingParam
+      camera_trajectory:
+        target: request.extensions.hunyuangamecraft.camera_trajectory
+        sources:
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraftSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft65FrameSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft129FrameSamplingParam
+      conditioning_mask:
+        target: request.extensions.hunyuangamecraft.conditioning_mask
+        sources:
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraftSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft65FrameSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft129FrameSamplingParam
+      gt_latents:
+        target: request.extensions.hunyuangamecraft.gt_latents
+        sources:
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraftSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft65FrameSamplingParam
+          - fastvideo.configs.sample.hunyuangamecraft.HunyuanGameCraft129FrameSamplingParam
+      prompt_attention_mask:
+        target: request.extensions.hyworld.prompt_attention_mask
+        sources: [fastvideo.configs.sample.hyworld.HYWorld_SamplingParam]
+      negative_attention_mask:
+        target: request.extensions.hyworld.negative_attention_mask
+        sources: [fastvideo.configs.sample.hyworld.HYWorld_SamplingParam]
+      ltx2_cfg_scale_audio:
+        target: request.extensions.ltx2.cfg_scale_audio
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_cfg_scale_video:
+        target: request.extensions.ltx2.cfg_scale_video
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_modality_scale_audio:
+        target: request.extensions.ltx2.modality_scale_audio
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_modality_scale_video:
+        target: request.extensions.ltx2.modality_scale_video
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_rescale_scale:
+        target: request.extensions.ltx2.rescale_scale
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_stg_blocks_audio:
+        target: request.extensions.ltx2.stg_blocks_audio
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_stg_blocks_video:
+        target: request.extensions.ltx2.stg_blocks_video
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_stg_scale_audio:
+        target: request.extensions.ltx2.stg_scale_audio
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+      ltx2_stg_scale_video:
+        target: request.extensions.ltx2.stg_scale_video
+        sources: [fastvideo.configs.sample.ltx2.LTX2BaseSamplingParam]
+
+  openai_image_request:
+    kept:
+      model: "HTTP adapter model-routing field."
+      response_format: "HTTP adapter response formatting field."
+      output_format: "HTTP adapter output-format field."
+      background: "HTTP adapter output-format field."
+      quality: "Compatibility field currently accepted by the adapter."
+      style: "Compatibility field currently accepted by the adapter."
+      user: "Compatibility field currently accepted by the adapter."
+    moved:
+      prompt: request.prompt
+      n: request.sampling.num_videos_per_prompt
+      size:
+        target: request.sampling.width,height
+        note: "Adapter parses OpenAI size strings as WIDTHxHEIGHT and forwards width then height."
+      num_inference_steps: request.sampling.num_inference_steps
+      guidance_scale: request.sampling.guidance_scale
+      true_cfg_scale: request.sampling.true_cfg_scale
+      seed: request.sampling.seed
+      negative_prompt: request.negative_prompt
+      enable_teacache: request.runtime.enable_teacache
+
+  openai_video_request:
+    kept:
+      model: "HTTP adapter model-routing field."
+    moved:
+      prompt: request.prompt
+      input_reference: request.inputs.image_path
+      reference_url: request.inputs.image_path
+      size:
+        target: request.sampling.width,height
+        note: "Adapter parses OpenAI size strings as WIDTHxHEIGHT and forwards width then height."
+      fps: request.sampling.fps
+      num_frames: request.sampling.num_frames
+      seed: request.sampling.seed
+      num_inference_steps: request.sampling.num_inference_steps
+      guidance_scale: request.sampling.guidance_scale
+      guidance_scale_2: request.sampling.guidance_scale_2
+      true_cfg_scale: request.sampling.true_cfg_scale
+      negative_prompt: request.negative_prompt
+      enable_teacache: request.runtime.enable_teacache
+      output_path: request.output.output_path
+    compatibility_only:
+      seconds:
+        target: request.sampling.num_frames
+        note: "HTTP adapter duration convenience field. If num_frames is omitted, the adapter computes num_frames = fps * seconds."
+
+cli:
+  notes:
+    - "CLI parity is checked against the actual generate/serve parser dest sets."
+    - "The inventory tracks parser dest names, excluding argparse's implicit help action."
+  generate:
+    explicit_local_fields:
+      - config
+    expected_dests:
+      - VSA_sparsity
+      - boundary_ratio
+      - bsa_cdf_threshold
+      - bsa_chunk_k
+      - bsa_chunk_q
+      - bsa_sparsity
+      - config
+      - disable_autocast
+      - dist_timeout
+      - distributed_executor_backend
+      - dit_config.prefix
+      - dit_config.quant_config
+      - dit_cpu_offload
+      - dit_layerwise_offload
+      - dit_precision
+      - dmd_denoising_steps
+      - embedded_cfg_scale
+      - enable_bsa
+      - enable_stage_verification
+      - enable_torch_compile
+      - flow_shift
+      - fps
+      - guidance_rescale
+      - guidance_scale
+      - height
+      - hsdp_replicate_dim
+      - hsdp_shard_dim
+      - image_encoder_cpu_offload
+      - image_encoder_precision
+      - image_path
+      - inference_mode
+      - init_weights_from_safetensors
+      - init_weights_from_safetensors_2
+      - lora_nickname
+      - lora_path
+      - lora_target_modules
+      - ltx2_initial_latent_path
+      - ltx2_vae_spatial_tile_overlap_in_pixels
+      - ltx2_vae_spatial_tile_size_in_pixels
+      - ltx2_vae_temporal_tile_overlap_in_frames
+      - ltx2_vae_temporal_tile_size_in_frames
+      - ltx2_vae_tiling
+      - master_port
+      - moba_config_path
+      - mode
+      - model_path
+      - negative_prompt
+      - num_cond_frames
+      - num_frames
+      - num_gpus
+      - num_inference_steps
+      - num_videos_per_prompt
+      - output_path
+      - output_type
+      - output_video_name
+      - override_pipeline_cls_name
+      - override_text_encoder_quant
+      - override_text_encoder_safetensors
+      - override_transformer_cls_name
+      - pin_cpu_memory
+      - pipeline_config_path
+      - preprocess.dataloader_num_workers
+      - preprocess.dataset_output_dir
+      - preprocess.dataset_path
+      - preprocess.dataset_type
+      - preprocess.do_temporal_sample
+      - preprocess.drop_short_ratio
+      - preprocess.flush_frequency
+      - preprocess.max_height
+      - preprocess.max_width
+      - preprocess.model_path
+      - preprocess.num_frames
+      - preprocess.preprocess_video_batch_size
+      - preprocess.samples_per_file
+      - preprocess.seed
+      - preprocess.speed_factor
+      - preprocess.train_fps
+      - preprocess.training_cfg_rate
+      - preprocess.video_length_tolerance_range
+      - preprocess.video_loader_type
+      - preprocess.with_audio
+      - prompt
+      - prompt_path
+      - prompt_txt
+      - refine_from
+      - return_frames
+      - return_trajectory_decoded
+      - return_trajectory_latents
+      - revision
+      - save_video
+      - seed
+      - sp_size
+      - spatial_refine_only
+      - t_thresh
+      - text_encoder_configs
+      - text_encoder_cpu_offload
+      - text_encoder_precisions
+      - torch_compile_kwargs
+      - transformer_quant
+      - tp_size
+      - trust_remote_code
+      - use_fsdp_inference
+      - vae_config.blend_num_frames
+      - vae_config.load_decoder
+      - vae_config.load_encoder
+      - vae_config.tile_sample_min_height
+      - vae_config.tile_sample_min_num_frames
+      - vae_config.tile_sample_min_width
+      - vae_config.tile_sample_stride_height
+      - vae_config.tile_sample_stride_num_frames
+      - vae_config.tile_sample_stride_width
+      - vae_config.use_parallel_tiling
+      - vae_config.use_temporal_tiling
+      - vae_config.use_tiling
+      - vae_cpu_offload
+      - vae_precision
+      - vae_sp
+      - vae_tiling
+      - video_path
+      - width
+      - workload_type
+  serve:
+    explicit_local_fields:
+      - config
+      - host
+      - output_dir
+      - port
+    expected_dests:
+      - VSA_sparsity
+      - bsa_cdf_threshold
+      - bsa_chunk_k
+      - bsa_chunk_q
+      - bsa_sparsity
+      - config
+      - disable_autocast
+      - dist_timeout
+      - distributed_executor_backend
+      - dit_config.prefix
+      - dit_config.quant_config
+      - dit_cpu_offload
+      - dit_layerwise_offload
+      - dit_precision
+      - dmd_denoising_steps
+      - embedded_cfg_scale
+      - enable_bsa
+      - enable_stage_verification
+      - enable_torch_compile
+      - flow_shift
+      - host
+      - hsdp_replicate_dim
+      - hsdp_shard_dim
+      - image_encoder_cpu_offload
+      - image_encoder_precision
+      - inference_mode
+      - init_weights_from_safetensors
+      - init_weights_from_safetensors_2
+      - lora_nickname
+      - lora_path
+      - lora_target_modules
+      - ltx2_initial_latent_path
+      - ltx2_vae_spatial_tile_overlap_in_pixels
+      - ltx2_vae_spatial_tile_size_in_pixels
+      - ltx2_vae_temporal_tile_overlap_in_frames
+      - ltx2_vae_temporal_tile_size_in_frames
+      - ltx2_vae_tiling
+      - master_port
+      - mode
+      - model_path
+      - num_gpus
+      - output_dir
+      - output_type
+      - override_pipeline_cls_name
+      - override_text_encoder_quant
+      - override_text_encoder_safetensors
+      - override_transformer_cls_name
+      - pin_cpu_memory
+      - pipeline_config_path
+      - port
+      - preprocess.dataloader_num_workers
+      - preprocess.dataset_output_dir
+      - preprocess.dataset_path
+      - preprocess.dataset_type
+      - preprocess.do_temporal_sample
+      - preprocess.drop_short_ratio
+      - preprocess.flush_frequency
+      - preprocess.max_height
+      - preprocess.max_width
+      - preprocess.model_path
+      - preprocess.num_frames
+      - preprocess.preprocess_video_batch_size
+      - preprocess.samples_per_file
+      - preprocess.seed
+      - preprocess.speed_factor
+      - preprocess.train_fps
+      - preprocess.training_cfg_rate
+      - preprocess.video_length_tolerance_range
+      - preprocess.video_loader_type
+      - preprocess.with_audio
+      - prompt_txt
+      - revision
+      - sp_size
+      - text_encoder_cpu_offload
+      - text_encoder_precisions
+      - torch_compile_kwargs
+      - transformer_quant
+      - tp_size
+      - trust_remote_code
+      - use_fsdp_inference
+      - vae_config.blend_num_frames
+      - vae_config.load_decoder
+      - vae_config.load_encoder
+      - vae_config.tile_sample_min_height
+      - vae_config.tile_sample_min_num_frames
+      - vae_config.tile_sample_min_width
+      - vae_config.tile_sample_stride_height
+      - vae_config.tile_sample_stride_num_frames
+      - vae_config.tile_sample_stride_width
+      - vae_config.use_parallel_tiling
+      - vae_config.use_temporal_tiling
+      - vae_config.use_tiling
+      - vae_cpu_offload
+      - vae_precision
+      - vae_sp
+      - vae_tiling
+      - workload_type

docs/design/overview.md

@@ -167,6 +167,11 @@ How this maps to FastVideo:
 
 - Attention backends live in `fastvideo/attention/` and can be selected via
   `FASTVIDEO_ATTENTION_BACKEND`.
+- SageAttention3 is split into two selectable backends:
+  `SAGE_ATTN_THREE` for the regular upstream package and
+  `ATTN_QAT_INFER` for the FastVideoKernel-backed inference variant.
+- `ATTN_QAT_TRAIN` is a separate FastVideoKernel Triton backend for the QAT attention
+  path.
 - `LocalAttention` is used for cross-attention and most attention layers.
 - `DistributedAttention` is used for full-sequence self-attention in the DiT.
 - Tensor-parallel layers live in `fastvideo/layers/`.

docs/inference/gen3c.md

@@ -0,0 +1,129 @@
+# GEN3C: 3D-Informed Camera-Controlled Video Generation
+
+[GEN3C](https://arxiv.org/abs/2503.03751) is NVIDIA's Cosmos-7B-based video model for camera-controlled generation from a single image. The FastVideo integration supports the GEN3C I2V workflow, including 3D cache conditioning and tokenizer-based conditioning latents.
+
+## Key Features
+
+- **Camera trajectory control**: `left/right/up/down/zoom_in/zoom_out/clockwise/counterclockwise`
+- **3D cache conditioning**: depth prediction -> point cloud cache -> forward warping -> latent conditioning
+- **Single-image to video generation**: 121-frame generation with camera motion
+- **Official raw checkpoint conversion**: `model.pt` -> Diffusers/FastVideo layout
+
+## Model Sources
+
+- Official raw checkpoint (not Diffusers): `nvidia/GEN3C-Cosmos-7B`
+- Diffusers-format checkpoint: `FastVideo/GEN3C-Cosmos-7B-Diffusers`
+
+## Prerequisites
+
+- Install MoGe:
+
+```bash
+pip install git+https://github.com/microsoft/MoGe.git
+```
+
+- If you hit `ImportError: libGL.so.1` (common on Ubuntu/headless nodes), you can try installing OpenCV runtime libs:
+
+```bash
+sudo apt-get update
+sudo apt-get install -y libgl1 libglib2.0-0 libsm6 libxext6 libxrender1
+```
+
+## Quick Start
+
+### Option A: Use Diffusers-format weights directly
+
+```bash
+python examples/inference/basic/basic_gen3c.py \
+  --model_path FastVideo/GEN3C-Cosmos-7B-Diffusers \
+  --image_path /path/to/input.png \
+  --prompt "" \
+  --trajectory left \
+  --movement_distance 0.3 \
+  --camera_rotation center_facing \
+  --num_inference_steps 35 \
+  --guidance_scale 1.0 \
+  --output_path outputs_video/gen3c_output.mp4
+```
+
+### Option B: Convert official raw checkpoint locally
+
+1. Download:
+
+```bash
+huggingface-cli download nvidia/GEN3C-Cosmos-7B --local-dir official_weights/GEN3C-Cosmos-7B
+```
+
+1. Convert:
+
+```bash
+python scripts/checkpoint_conversion/convert_gen3c_to_fastvideo.py \
+  --source official_weights/GEN3C-Cosmos-7B/model.pt \
+  --output converted_weights/GEN3C-Cosmos-7B
+```
+
+1. Run:
+
+```bash
+python examples/inference/basic/basic_gen3c.py \
+  --model_path converted_weights/GEN3C-Cosmos-7B \
+  --image_path /path/to/input.png \
+  --prompt "" \
+  --trajectory left \
+  --movement_distance 0.3 \
+  --camera_rotation center_facing \
+  --num_inference_steps 35 \
+  --guidance_scale 1.0 \
+  --output_path outputs_video/gen3c_output.mp4
+```
+
+## FastVideo Defaults
+
+GEN3C defaults in FastVideo:
+
+- `height=704`, `width=1280`
+- `num_frames=121`
+- `num_inference_steps=35`
+- `guidance_scale=1.0`
+- `fps=24`
+
+These values are defined in:
+
+- `fastvideo/configs/sample/gen3c.py`
+- `fastvideo/configs/pipelines/gen3c.py`
+
+and align with the official GEN3C inference defaults in:
+
+- `tmp/GEN3C/cosmos_predict1/diffusion/inference/inference_utils.py`
+
+## Scheduler Note
+
+The converted GEN3C Diffusers layout may include a FlowMatch scheduler config, but GEN3C denoising uses EDM preconditioning behavior. FastVideo's GEN3C pipeline enforces an EDM scheduler at runtime for parity with official inference behavior.
+
+Implementation path:
+
+- `fastvideo/pipelines/basic/gen3c/gen3c_pipeline.py`
+
+## 3D Cache Conditioning Path
+
+FastVideo GEN3C conditioning stage performs:
+
+1. MoGe depth estimation from input image
+2. 3D cache initialization
+3. Camera trajectory generation
+4. Forward rendering of warped frames + masks
+5. VAE/tokenizer encoding of conditioning buffers
+6. Denoising with condition mask + condition pose channels
+
+Main implementation:
+
+- `fastvideo/pipelines/basic/gen3c/gen3c_pipeline.py`
+- `fastvideo/pipelines/basic/gen3c/cache_3d.py`
+- `fastvideo/pipelines/basic/gen3c/depth_estimation.py`
+- `fastvideo/models/vaes/gen3c_tokenizer_vae.py`
+
+## References
+
+- [GEN3C Paper](https://arxiv.org/abs/2503.03751)
+- [Official Repository](https://github.com/nv-tlabs/GEN3C)
+- [Official Checkpoint (raw)](https://huggingface.co/nvidia/GEN3C-Cosmos-7B)

docs/inference/inference_quick_start.md

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

docs/inference/optimizations.md

@@ -21,6 +21,8 @@ This page describes the various options for speeding up generation times in Fast
 - Video Sparse Attention: `FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_ATTN`
 - Sage Attention: `FASTVIDEO_ATTENTION_BACKEND=SAGE_ATTN`
 - Sage Attention 3: `FASTVIDEO_ATTENTION_BACKEND=SAGE_ATTN_THREE`
+- Attn QAT Infer: `FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_INFER`
+- Attn QAT Train: `FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_TRAIN`
 - Video MoBA Attention: `FASTVIDEO_ATTENTION_BACKEND=VMOBA_ATTN`
 - Sparse Linear Attention: `FASTVIDEO_ATTENTION_BACKEND=SLA_ATTN`
 - SageSLA Attention: `FASTVIDEO_ATTENTION_BACKEND=SAGE_SLA_ATTN`
@@ -103,6 +105,14 @@ python setup.py install  # or pip install -e .
 
 ### Sage Attention 3
 
+FastVideo now exposes two SageAttention3-compatible backends with distinct
+environment variable values:
+
+- `SAGE_ATTN_THREE`: the regular upstream SageAttention3 backend imported from
+  the `sageattn3` package.
+- `ATTN_QAT_INFER`: the inference CUDA-kernel backend imported from the
+  in-repo `attn_qat_infer` package.
+
 **`SAGE_ATTN_THREE`**
 
 [SageAttention 3](https://github.com/thu-ml/SageAttention/tree/main/sageattention3_blackwell) is an advanced attention mechanism that leverages FP4 quantization and Blackwell GPU Tensor Cores for significant performance improvements.
@@ -117,6 +127,53 @@ Note that Sage Attention 3 requires `python>=3.13`, `torch>=2.8.0`, `CUDA >=12.8
 
 To use Sage Attention 3 in FastVideo, follow the `README.md` in the linked repository to install the package from source.
 
+### Attn QAT Infer
+
+**`ATTN_QAT_INFER`**
+
+This backend uses the `attn_qat_infer` implementation that lives in the
+`fastvideo-kernel` repository alongside the `fastvideo_kernel` Triton kernels.
+Use this backend when you want to run the dedicated FP4 inference CUDA kernel
+directly during inference.
+
+For the full Attention QAT guide, including Wan 2.1 14B checkpoint download,
+example editing steps, training launchers, and troubleshooting, see
+[Attention QAT](../attention/attn_qat/index.md).
+
+This backend currently assumes access to the in-repo `fastvideo-kernel`
+checkout or an equivalent editable/source install that exposes:
+
+- `attn_qat_infer`
+
+Example:
+
+```python
+os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "ATTN_QAT_INFER"
+```
+
+### QAT Attention
+
+**`ATTN_QAT_TRAIN`**
+
+This backend uses the FastVideoKernel Triton attention implementation from
+`fastvideo_kernel.triton_kernels.attn_qat_train`. Use it when you specifically
+want the training-oriented Triton attention path rather than the
+`attn_qat_infer` CUDA kernel path.
+
+The dedicated [Attention QAT](../attention/attn_qat/index.md) page covers when
+to use `ATTN_QAT_TRAIN` versus `ATTN_QAT_INFER`, the ready-made training
+launchers, and the end-to-end Wan 2.1 14B inference workflow.
+
+This backend currently assumes access to an install that exposes:
+
+- `fastvideo_kernel`
+
+Example:
+
+```python
+os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "ATTN_QAT_TRAIN"
+```
+
 ### V-MoBA / SLA / SageSLA
 
 These backends are model-specific and require the corresponding kernels and

docs/inference/support_matrix.md

@@ -73,6 +73,7 @@ pipeline initialization and sampling.
 | Matrix Game 2.0 Base | `FastVideo/Matrix-Game-2.0-Base-Diffusers` | 352x640 | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
 | Matrix Game 2.0 GTA | `FastVideo/Matrix-Game-2.0-GTA-Diffusers` | 352x640 | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
 | Matrix Game 2.0 TempleRun | `FastVideo/Matrix-Game-2.0-TempleRun-Diffusers` | 352x640 | ⭕ | ⭕ | ⭕ | ⭕ | ⭕ |
+| GEN3C Cosmos 7B | `FastVideo/GEN3C-Cosmos-7B-Diffusers` | 704px1280p | ❌ | ❌ | ❌ | ⭕ | ⭕ |
 
 **Note**: Wan2.2 TI2V 5B has some quality issues when performing I2V generation. We are working on fixing this issue.
 
@@ -85,6 +86,11 @@ The authoritative source for model-ID recognition is
 `fastvideo/registry.py`. If a model ID is registered there, FastVideo can
 resolve default pipeline and sampling configuration for it.
 
+**Note (GEN3C)**: The official `nvidia/GEN3C-Cosmos-7B` repo provides a raw
+`model.pt` checkpoint. Use a Diffusers-format repo (for example,
+`FastVideo/GEN3C-Cosmos-7B-Diffusers`) or convert locally with
+`scripts/checkpoint_conversion/convert_gen3c_to_fastvideo.py`.
+
 ## Special requirements
 
 ### Sliding Tile Attention

docs/utilities/debugging.md

@@ -27,7 +27,8 @@ Useful variables:
 - `FASTVIDEO_LOGGING_LEVEL`: `DEBUG`, `INFO`, `WARNING`, `ERROR`
 - `FASTVIDEO_STAGE_LOGGING`: print per-stage timings during pipeline execution
 - `FASTVIDEO_ATTENTION_BACKEND`: force an attention backend (for example
-  `TORCH_SDPA` or `FLASH_ATTN`)
+  `TORCH_SDPA`, `FLASH_ATTN`, `SAGE_ATTN_THREE`, or
+  `ATTN_QAT_INFER`, or `ATTN_QAT_TRAIN`)
 
 ## Common Failure Modes
 
@@ -52,7 +53,11 @@ If forcing a backend fails, verify optional dependencies are installed:
 - `VIDEO_SPARSE_ATTN`: `fastvideo-kernel`
 - `SLIDING_TILE_ATTN`: STA legacy workflow in
   `sta_do_not_delete` + `fastvideo-kernel`
-- `SAGE_ATTN` / `SAGE_ATTN_THREE`: SageAttention packages
+- `SAGE_ATTN`: SageAttention package
+- `SAGE_ATTN_THREE`: upstream `sageattn3` package
+- `ATTN_QAT_INFER`: `fastvideo-kernel` checkout/source install that exposes
+  `attn_qat_infer`
+- `ATTN_QAT_TRAIN`: `fastvideo-kernel` install exposing `fastvideo_kernel`
 
 As a fallback, use:
 

examples/distill/SFWan2.2-A14B/distill_dmd.sh

@@ -22,7 +22,7 @@ export NODE_RANK=$SLURM_PROCID
 nodes=( $(scontrol show hostnames $SLURM_JOB_NODELIST) )
 export MASTER_ADDR=${nodes[0]}
 export TOKENIZERS_PARALLELISM=false
-export WANDB_API_KEY="2f25ad37933894dbf0966c838c0b8494987f9f2f"
+export WANDB_API_KEY=YOUR_WANDB_API_KEY
 # export WANDB_API_KEY='your_wandb_api_key_here'
 export WANDB_BASE_URL="https://api.wandb.ai"
 export WANDB_MODE=online

examples/distill/Wan2.1-T2V/Wan-Syn-Data-480P/README.md

@@ -9,7 +9,7 @@ pip install vsa
 
 ### 1. Download dataset:
 ```bash
-bash examples/distill/Wan-Syn-480P/download_dataset.sh
+bash examples/distill/Wan2.1-T2V/Wan-Syn-Data-480P/download_dataset.sh
 ```
 
 ### 2. Configure and run distillation:

examples/distill/Wan2.1-T2V/Wan-Syn-Data-480P/download_dataset.sh

@@ -1,3 +1,3 @@
 #!/bin/bash
-
-python scripts/huggingface/download_hf.py --repo_id "FastVideo/Wan-Syn_77x448x832_600k" --local_dir "FastVideo/Wan-Syn_77x448x832_600k" --repo_type "dataset"
+mkdir -p data
+python scripts/huggingface/download_hf.py --repo_id "FastVideo/Wan-Syn_77x448x832_600k" --local_dir "data/Wan-Syn_77x448x832_600k" --repo_type "dataset"

examples/inference/basic/README.md

@@ -28,6 +28,11 @@ For an example running DMD+VSA inference:
 python examples/inference/basic/basic_dmd.py
 ```
 
+For the typed config/request path added during the inference API refactor:
+```
+python examples/inference/basic/basic_dmd_new_api.py
+```
+
 ## Basic Walkthrough
 
 All you need to generate videos using multi-gpus from state-of-the-art diffusion pipelines is the following few lines!

examples/inference/basic/basic_dmd_new_api.py

@@ -0,0 +1,98 @@
+import os
+import time
+
+from fastvideo import VideoGenerator
+from fastvideo.api import (
+    EngineConfig,
+    GenerationRequest,
+    GeneratorConfig,
+    OffloadConfig,
+    OutputConfig,
+    PipelineSelection,
+)
+
+OUTPUT_PATH = "video_samples_dmd2_typed"
+
+
+def main():
+    os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "VIDEO_SPARSE_ATTN"
+
+    model_name = "FastVideo/FastWan2.1-T2V-1.3B-Diffusers"
+    generator_config = GeneratorConfig(
+        model_path=model_name,
+        engine=EngineConfig(
+            num_gpus=1,
+            use_fsdp_inference=False,
+            offload=OffloadConfig(
+                text_encoder=True,
+                pin_cpu_memory=True,
+                dit=False,
+                vae=False,
+            ),
+        ),
+        # PR 2 still routes a few advanced inference knobs through the
+        # compatibility bridge until they get first-class typed fields.
+        pipeline=PipelineSelection(
+            experimental={
+                "VSA_sparsity": 0.8,
+            },
+        ),
+    )
+
+    load_start_time = time.perf_counter()
+    generator = VideoGenerator.from_config(generator_config)
+    load_end_time = time.perf_counter()
+    load_time = load_end_time - load_start_time
+
+    prompt = (
+        "A neon-lit alley in futuristic Tokyo during a heavy rainstorm at night. "
+        "The puddles reflect glowing signs in kanji, advertising ramen, karaoke, "
+        "and VR arcades. A woman in a translucent raincoat walks briskly with an "
+        "LED umbrella. Steam rises from a street food cart, and a cat darts "
+        "across the screen. Raindrops are visible on the camera lens, creating "
+        "a cinematic bokeh effect."
+    )
+    request = GenerationRequest(
+        prompt=prompt,
+        output=OutputConfig(
+            output_path=OUTPUT_PATH,
+            save_video=True,
+            return_frames=False,
+        ),
+    )
+
+    start_time = time.perf_counter()
+    result = generator.generate(request)
+    end_time = time.perf_counter()
+    gen_time = end_time - start_time
+
+    prompt2 = (
+        "A majestic lion strides across the golden savanna, its powerful frame "
+        "glistening under the warm afternoon sun. The tall grass ripples gently "
+        "in the breeze, enhancing the lion's commanding presence. The tone is "
+        "vibrant, embodying the raw energy of the wild. Low angle, steady "
+        "tracking shot, cinematic."
+    )
+    request2 = GenerationRequest(
+        prompt=prompt2,
+        output=OutputConfig(
+            output_path=OUTPUT_PATH,
+            save_video=True,
+            return_frames=False,
+        ),
+    )
+
+    start_time = time.perf_counter()
+    result2 = generator.generate(request2)
+    end_time = time.perf_counter()
+    gen_time2 = end_time - start_time
+
+    print(f"Time taken to load model: {load_time} seconds")
+    print(f"Time taken to generate video: {gen_time} seconds")
+    print(f"First output written to: {result.video_path}")
+    print(f"Time taken to generate video2: {gen_time2} seconds")
+    print(f"Second output written to: {result2.video_path}")
+
+
+if __name__ == "__main__":
+    main()

examples/inference/basic/basic_gen3c.py

@@ -0,0 +1,109 @@
+"""
+GEN3C: 3D-aware camera-controlled video generation.
+
+This example generates a video from a single input image with camera control.
+The pipeline uses MoGe depth estimation, 3D point cloud forward warping,
+and the GEN3C diffusion model.
+
+Requirements:
+  1. Install MoGe:
+     pip install git+https://github.com/microsoft/MoGe.git
+     If you hit `ImportError: libGL.so.1`, install:
+     sudo apt-get update && sudo apt-get install -y libgl1 libglib2.0-0 libsm6 libxext6 libxrender1
+  2. Download and convert weights:
+     huggingface-cli download nvidia/GEN3C-Cosmos-7B --local-dir official_weights/GEN3C-Cosmos-7B
+     python scripts/checkpoint_conversion/convert_gen3c_to_fastvideo.py \
+       --source ./official_weights/GEN3C-Cosmos-7B/model.pt \
+       --output ./converted_weights/GEN3C-Cosmos-7B \
+       --components-source nvidia/Cosmos-Predict2-2B-Video2World
+  3. Provide an input image for 3D-conditioned generation.
+"""
+
+import argparse
+
+from fastvideo import VideoGenerator
+
+
+def main():
+    parser = argparse.ArgumentParser(description="GEN3C video generation")
+    parser.add_argument("--model_path",
+                        type=str,
+                        default="converted_weights/GEN3C-Cosmos-7B")
+    parser.add_argument("--image_path",
+                        type=str,
+                        default=None,
+                        help="Input image for 3D cache conditioning")
+    parser.add_argument("--prompt",
+                        type=str,
+                        default="A slow camera pan over a sunlit landscape.")
+    parser.add_argument(
+        "--negative_prompt",
+        type=str,
+        default=(
+            "The video captures a series of frames showing ugly scenes, static with no motion, motion blur, "
+            "over-saturation, shaky footage, low resolution, grainy texture, pixelated images, poorly lit areas, "
+            "underexposed and overexposed scenes, poor color balance, washed out colors, choppy sequences, "
+            "jerky movements, low frame rate, artifacting, color banding, unnatural transitions, outdated special "
+            "effects, fake elements, unconvincing visuals, poorly edited content, jump cuts, visual noise, and "
+            "flickering. Overall, the video is of poor quality."
+        ),
+    )
+    parser.add_argument("--trajectory",
+                        type=str,
+                        default="left",
+                        choices=[
+                            "left", "right", "up", "down", "zoom_in",
+                            "zoom_out", "clockwise", "counterclockwise", "none"
+                        ])
+    parser.add_argument("--movement_distance", type=float, default=0.3)
+    parser.add_argument("--camera_rotation",
+                        type=str,
+                        default="center_facing",
+                        choices=[
+                            "center_facing", "no_rotation",
+                            "trajectory_aligned"
+                        ])
+    parser.add_argument("--height", type=int, default=704)
+    parser.add_argument("--width", type=int, default=1280)
+    parser.add_argument("--num_frames", type=int, default=121)
+    parser.add_argument("--num_inference_steps", type=int, default=35)
+    parser.add_argument("--guidance_scale", type=float, default=1.0)
+    parser.add_argument("--output_path",
+                        type=str,
+                        default="outputs_video/gen3c.mp4")
+    parser.add_argument("--seed", type=int, default=42)
+    args = parser.parse_args()
+
+    generator = VideoGenerator.from_pretrained(
+        args.model_path,
+        num_gpus=1,
+        use_fsdp_inference=False,
+        dit_cpu_offload=False,
+        vae_cpu_offload=True,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=True,
+    )
+
+    video = generator.generate_video(
+        args.prompt,
+        negative_prompt=args.negative_prompt,
+        image_path=args.image_path,
+        trajectory_type=args.trajectory,
+        movement_distance=args.movement_distance,
+        camera_rotation=args.camera_rotation,
+        height=args.height,
+        width=args.width,
+        num_frames=args.num_frames,
+        num_inference_steps=args.num_inference_steps,
+        guidance_scale=args.guidance_scale,
+        fps=24,
+        seed=args.seed,
+        output_path=args.output_path,
+        save_video=True,
+    )
+
+    generator.shutdown()
+
+
+if __name__ == "__main__":
+    main()

examples/inference/optimizations/README.md

@@ -1,5 +1,79 @@
 # Optimization Examples
 
+## Wan 2.1 QAT Attention 14B Inference
+
+Use these files for Wan 2.1 14B inference with the `ATTN_QAT_INFER` backend:
+
+- `examples/inference/optimizations/download_14B_qat.sh`
+- `examples/inference/optimizations/attn_qat_inference_example.py`
+
+### 1. Download the 14B QAT checkpoint
+
+The helper script downloads the QAT safetensors from
+`FastVideo/14B_qat_400` into `checkpoints/14B_qat_400` by default.
+
+Prerequisites:
+
+- `huggingface_hub` installed, for example: `uv pip install huggingface_hub`
+- access to the model repo if it is private or gated: `huggingface-cli login`
+
+Run:
+
 ```bash
-python examples/inference/optimizations/attention_example.py
+bash examples/inference/optimizations/download_14B_qat.sh
 ```
+
+To download into a custom directory, pass it as the first argument:
+
+```bash
+bash examples/inference/optimizations/download_14B_qat.sh /path/to/14B_qat_400
+```
+
+### 2. Edit the inference example for Wan 2.1 14B
+
+Open `examples/inference/optimizations/attn_qat_inference_example.py` and
+update these two values:
+
+1. Change the base model from `Wan-AI/Wan2.1-T2V-1.3B-Diffusers` to
+   `Wan-AI/Wan2.1-T2V-14B-Diffusers`.
+2. Replace the placeholder
+   `init_weights_from_safetensors="safetensors_path"` with the directory that
+   contains the downloaded `.safetensors` files.
+
+Example:
+
+```python
+generator = VideoGenerator.from_pretrained(
+    "Wan-AI/Wan2.1-T2V-14B-Diffusers",
+    num_gpus=1,
+    use_fsdp_inference=True,
+    dit_cpu_offload=False,
+    vae_cpu_offload=False,
+    text_encoder_cpu_offload=True,
+    pin_cpu_memory=False,
+    init_weights_from_safetensors="checkpoints/14B_qat_400",
+)
+```
+
+The script already sets:
+
+```python
+os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "ATTN_QAT_INFER"
+```
+
+### 3. Run the example
+
+```bash
+python examples/inference/optimizations/attn_qat_inference_example.py
+```
+
+The generated videos are written to `video_samples/` by default.
+
+### Notes
+
+- `ATTN_QAT_INFER` requires the in-repo `fastvideo-kernel` build to expose the
+  `attn_qat_infer` package.
+- If you have not built the kernel yet, run `cd fastvideo-kernel && ./build.sh`
+  first.
+- If you keep the example on the `1.3B` base model while loading the 14B QAT
+  weights, the model/config will not match.

examples/inference/optimizations/attn_qat_inference_example.py

@@ -0,0 +1,54 @@
+from fastvideo import VideoGenerator
+import os
+from pathlib import Path
+# from fastvideo.configs.sample import SamplingParam
+
+OUTPUT_PATH = "video_samples"
+os.environ["FASTVIDEO_ATTENTION_BACKEND"] = "ATTN_QAT_INFER"
+
+CHECKPOINT_PATH = Path(__file__).parent.parent.parent 
+
+def main():
+    # FastVideo will automatically use the optimal default arguments for the
+    # model.
+    # If a local path is provided, FastVideo will make a best effort
+    # attempt to identify the optimal arguments.
+    generator = VideoGenerator.from_pretrained(
+        "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
+        # FastVideo will automatically handle distributed setup
+        num_gpus=1,
+        use_fsdp_inference=True,
+        dit_cpu_offload=False,
+        vae_cpu_offload=False,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False, # set to false if low CPU RAM or hit obscure "CUDA error: Invalid argument"
+        # image_encoder_cpu_offload=False,
+        # Load custom weights from checkpoint 
+        init_weights_from_safetensors="safetensors_path"
+    )
+
+    # sampling_param = SamplingParam.from_pretrained("Wan-AI/Wan2.1-T2V-1.3B-Diffusers")
+    # sampling_param.num_frames = 45
+    # sampling_param.image_path = "https://huggingface.co/datasets/huggingface/documentation-images/resolve/main/diffusers/astronaut.jpg"
+    # Generate videos with the same simple API, regardless of GPU count
+    prompt = (
+        "A curious raccoon peers through a vibrant field of yellow sunflowers, its eyes "
+        "wide with interest. The playful yet serene atmosphere is complemented by soft "
+        "natural light filtering through the petals. Mid-shot, warm and cheerful tones."
+    )
+    video = generator.generate_video(prompt, output_path=OUTPUT_PATH, save_video=True)
+    # video = generator.generate_video(prompt, sampling_param=sampling_param, output_path="wan_t2v_videos/")
+
+    # Generate another video with a different prompt, without reloading the
+    # model!
+    prompt2 = (
+        "A majestic lion strides across the golden savanna, its powerful frame "
+        "glistening under the warm afternoon sun. The tall grass ripples gently in "
+        "the breeze, enhancing the lion's commanding presence. The tone is vibrant, "
+        "embodying the raw energy of the wild. Low angle, steady tracking shot, "
+        "cinematic.")
+    video2 = generator.generate_video(prompt2, output_path=OUTPUT_PATH, save_video=True)
+
+
+if __name__ == "__main__":
+    main()

examples/inference/optimizations/download_14B_qat.sh

@@ -0,0 +1,58 @@
+#!/usr/bin/env bash
+
+set -euo pipefail
+
+SCRIPT_DIR="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)"
+REPO_ROOT="$(cd -- "${SCRIPT_DIR}/../../.." && pwd)"
+
+HF_REPO_ID="${HF_REPO_ID:-FastVideo/14B_qat_400}"
+HF_REVISION="${HF_REVISION:-main}"
+LOCAL_DIR="${1:-${REPO_ROOT}/checkpoints/14B_qat_400}"
+PYTHON_BIN="${PYTHON:-python}"
+
+if ! command -v "${PYTHON_BIN}" >/dev/null 2>&1; then
+    echo "Python executable not found: ${PYTHON_BIN}" >&2
+    exit 1
+fi
+
+if ! "${PYTHON_BIN}" -c "import huggingface_hub" >/dev/null 2>&1; then
+    echo "Missing dependency: huggingface_hub" >&2
+    echo "Install it with: uv pip install huggingface_hub" >&2
+    exit 1
+fi
+
+mkdir -p "${LOCAL_DIR}"
+
+echo "Downloading ${HF_REPO_ID}@${HF_REVISION}"
+echo "Local directory: ${LOCAL_DIR}"
+
+"${PYTHON_BIN}" -c '
+import argparse
+from huggingface_hub import snapshot_download
+
+parser = argparse.ArgumentParser()
+parser.add_argument("--repo-id", required=True)
+parser.add_argument("--revision", required=True)
+parser.add_argument("--local-dir", required=True)
+args = parser.parse_args()
+
+snapshot_download(
+    repo_id=args.repo_id,
+    revision=args.revision,
+    repo_type="model",
+    local_dir=args.local_dir,
+    local_dir_use_symlinks=False,
+    resume_download=True,
+)
+' \
+    --repo-id "${HF_REPO_ID}" \
+    --revision "${HF_REVISION}" \
+    --local-dir "${LOCAL_DIR}"
+
+echo
+echo "Download complete."
+echo "Use this in your inference script:"
+echo "init_weights_from_safetensors=\"${LOCAL_DIR}\""
+echo
+echo "If the repo is private or gated, make sure you are logged in with:"
+echo "huggingface-cli login"

examples/inference/optimizations/quantization_example.py

@@ -0,0 +1,88 @@
+import torch
+from fastvideo import VideoGenerator
+from fastvideo.configs.pipelines.base import PipelineConfig
+
+OUTPUT_PATH = "video_samples"
+
+
+def main():
+    print("=== FP4 Quantization Video Generation Example ===")
+
+    if not torch.cuda.is_available():
+        print("Warning: CUDA not available. FP4 quantization requires GPU.")
+        return
+    
+    gpu_capability = torch.cuda.get_device_capability()
+    if gpu_capability[0] < 9:  # H100 and newer
+        print(f"Warning: GPU capability {gpu_capability} may not support FP4. Recommended: 9.0+")
+    
+    print(f"GPU: {torch.cuda.get_device_name()}")
+    print(f"GPU Capability: {gpu_capability}")
+    
+    model_id = "Wan-AI/Wan2.1-T2V-1.3B-Diffusers"
+    # model_id = "Wan-AI/Wan2.1-T2V-14B-Diffusers"
+    pipeline_config = PipelineConfig.from_pretrained(model_id)
+    pipeline_config.dit_precision = "bf16"
+
+    print("\nLoading model with FP4 quantization...")
+    generator = VideoGenerator.from_pretrained(
+        model_id,
+        pipeline_config=pipeline_config,
+        num_gpus=1,
+        use_fsdp_inference=True,
+        transformer_quant="fp4",
+        dit_cpu_offload=False,
+        vae_cpu_offload=False,
+        text_encoder_cpu_offload=True,
+        pin_cpu_memory=False,
+    )
+    
+    print("FP4 configuration applied. Generating videos...")
+    
+    print("\n=== Generating Video with FP4 Quantization ===")
+    
+    prompt1 = (
+        "A curious raccoon peers through a vibrant field of yellow sunflowers, its eyes "
+        "wide with interest. The playful yet serene atmosphere is complemented by soft "
+        "natural light filtering through the petals. Mid-shot, warm and cheerful tones."
+    )
+    
+    print(f"Prompt: {prompt1}")
+    print("Generating video...")
+    
+    try:
+        video1 = generator.generate_video(
+            prompt1,
+            output_path=OUTPUT_PATH,
+            save_video=True,
+        )
+        print("✓ First video generated successfully with FP4 quantization!")
+        
+        # # Generate a second video to show the model can be reused
+        prompt2 = (
+            "A majestic lion strides across the golden savanna, its powerful frame "
+            "glistening under the warm afternoon sun. The tall grass ripples gently in "
+            "the breeze, enhancing the lion's commanding presence. The tone is vibrant, "
+            "embodying the raw energy of the wild. Low angle, steady tracking shot, "
+            "cinematic."
+        )
+        
+        print(f"\nGenerating second video...")
+        print(f"Prompt: {prompt2}")
+        
+        video2 = generator.generate_video(
+            prompt2,
+            output_path=OUTPUT_PATH,
+            save_video=True,
+        )
+        print("✓ Second video generated successfully with FP4 quantization!")
+        
+    except Exception as e:
+        print(f"Error during video generation: {e}")
+        return
+    
+    print(f"Videos saved to: {OUTPUT_PATH}")
+
+
+if __name__ == "__main__":
+    main() 

examples/train/README.md

@@ -0,0 +1,93 @@
+# Training
+
+## Quick Start
+
+### Single-node
+
+```bash
+bash examples/train/run.sh <config.yaml> [--dotted.key value ...]
+```
+
+The script auto-detects available GPUs. Override with `NUM_GPUS`:
+
+```bash
+NUM_GPUS=4 bash examples/train/run.sh \
+    examples/train/configs/fine_tuning/wan/t2v.yaml
+```
+
+### Multi-node (Slurm)
+
+```bash
+bash examples/train/run_slurm.sh <config.yaml> <num_nodes> [--dotted.key value ...]
+```
+
+```bash
+bash examples/train/run_slurm.sh \
+    examples/train/configs/fine_tuning/wan/t2v.yaml 4 \
+    --training.distributed.hsdp_shard_dim 8 \
+    --training.distributed.hsdp_replicate_dim 4
+```
+
+Slurm environment variables:
+
+| Variable | Default | Description |
+|---|---|---|
+| `PARTITION` | `main` | Slurm partition |
+| `NUM_GPUS` | `8` | GPUs per node |
+| `CPUS_PER_TASK` | `128` | CPUs per task |
+| `MEM` | `1440G` | Memory per node |
+| `EXCLUDE` | `""` | Nodes to exclude |
+
+### Overriding config values
+
+Any config field can be overridden from the command line using dotted keys:
+
+```bash
+bash examples/train/run.sh examples/train/configs/fine_tuning/wan/t2v.yaml \
+    --training.optimizer.learning_rate 1e-5 \
+    --training.loop.max_train_steps 1000 \
+    --training.checkpoint.output_dir outputs/my_experiment
+```
+
+### Resuming from a checkpoint
+
+```bash
+bash examples/train/run.sh examples/train/configs/fine_tuning/wan/t2v.yaml \
+    --training.checkpoint.resume_from_checkpoint outputs/my_experiment/checkpoint-500
+```
+
+## W&B Logging
+
+Training metrics and validation videos are logged to
+[Weights & Biases](https://wandb.ai). Set your API key before launching:
+
+```bash
+export WANDB_API_KEY=your_key_here
+```
+
+To disable W&B (e.g. for local debugging):
+
+```bash
+export WANDB_MODE=offline
+```
+
+The project name and run name are set in the config under `training.tracker`:
+
+```yaml
+training:
+  tracker:
+    project_name: my_project
+    run_name: my_run
+```
+
+## Directory Layout
+
+```
+examples/train/
+├── configs/      # Single-step training configs (by method and model)
+├── scenario/     # Multi-step end-to-end training pipelines
+├── run.sh        # Single-node launcher
+└── run_slurm.sh  # Multi-node Slurm launcher
+```
+
+See `configs/README.md` and `scenario/README.md` for details.

examples/train/configs/README.md

@@ -0,0 +1,21 @@
+# Training Configs
+
+Single-step training configurations organized by method and model.
+
+```
+configs/
+├── fine_tuning/          # Standard finetuning and DFSFT
+├── distribution_matching/  # DMD2 and Self-Forcing
+├── knowledge_distillation/ # KD from teacher to student
+└── example.yaml          # Annotated reference config with all fields
+```
+
+Each method directory contains per-model subdirectories (e.g. `wan/`, `hunyuan/`).
+
+Launch any config with:
+
+```bash
+bash examples/train/run.sh examples/train/configs/<method>/<model>/<config>.yaml
+```
+
+For multi-step training pipelines, see `examples/train/scenario/`.

examples/train/configs/fine_tuning/hunyuan/t2v.yaml

@@ -0,0 +1,72 @@
+# HunyuanVideo T2V finetune config.
+#
+# Data must be preprocessed with Hunyuan VAE + dual text encoders
+# (LLaMA + CLIP) into parquet format before training.
+# See scripts/preprocess/ for preprocessing pipeline.
+
+models:
+  student:
+    _target_: fastvideo.train.models.hunyuan.HunyuanModel
+    init_from: hunyuanvideo-community/HunyuanVideo
+    trainable: true
+
+method:
+  _target_: fastvideo.train.methods.fine_tuning.finetune.FineTuneMethod
+
+training:
+  distributed:
+    num_gpus: 8
+    sp_size: 1
+    tp_size: 1
+    hsdp_replicate_dim: 1
+    hsdp_shard_dim: 8
+
+  data:
+    data_path: data/hunyuan_training_data
+    dataloader_num_workers: 4
+    train_batch_size: 1
+    training_cfg_rate: 0.0
+    seed: 1000
+    num_latent_t: 32
+    num_height: 720
+    num_width: 1280
+    num_frames: 125
+
+  optimizer:
+    learning_rate: 1.0e-5
+    betas: [0.9, 0.999]
+    weight_decay: 0.01
+    lr_scheduler: constant
+    lr_warmup_steps: 0
+
+  loop:
+    max_train_steps: 5000
+    gradient_accumulation_steps: 1
+
+  checkpoint:
+    output_dir: outputs/hunyuan_finetune
+    training_state_checkpointing_steps: 500
+    checkpoints_total_limit: 3
+    resume_from_checkpoint: latest
+
+  tracker:
+    project_name: fastvideo_hunyuan
+    run_name: hunyuan_finetune
+
+  model:
+    enable_gradient_checkpointing_type: full
+
+callbacks:
+  grad_clip:
+    _target_: fastvideo.train.callbacks.grad_clip.GradNormClipCallback
+    max_grad_norm: 1.0
+  validation:
+    _target_: fastvideo.train.callbacks.validation.ValidationCallback
+    pipeline_target: fastvideo.pipelines.basic.hunyuan.hunyuan_pipeline.HunyuanVideoPipeline
+    dataset_file: data/hunyuan_training_data/validation_prompts.json
+    every_steps: 100
+    sampling_steps: [50]
+    guidance_scale: 1.0
+
+pipeline:
+  flow_shift: 7

examples/train/configs/fine_tuning/wan/t2v.yaml

@@ -0,0 +1,68 @@
+# Wan 2.1 T2V 1.3B finetune.
+
+models:
+  student:
+    _target_: fastvideo.train.models.wan.WanModel
+    init_from: Wan-AI/Wan2.1-T2V-1.3B-Diffusers
+    trainable: true
+
+method:
+  _target_: fastvideo.train.methods.fine_tuning.finetune.FineTuneMethod
+
+training:
+  distributed:
+    num_gpus: 8
+    sp_size: 1
+    tp_size: 1
+    hsdp_replicate_dim: 8
+    hsdp_shard_dim: 1
+
+  data:
+    data_path: data/preprocessed
+    dataloader_num_workers: 4
+    train_batch_size: 1
+    training_cfg_rate: 0.0
+    seed: 1000
+    num_latent_t: 20
+    num_height: 448
+    num_width: 832
+    num_frames: 77
+
+  optimizer:
+    learning_rate: 1.0e-6
+    betas: [0.9, 0.999]
+    weight_decay: 0.01
+    lr_scheduler: constant
+    lr_warmup_steps: 0
+
+  loop:
+    max_train_steps: 4000
+    gradient_accumulation_steps: 1
+
+  checkpoint:
+    output_dir: outputs/wan_finetune
+    training_state_checkpointing_steps: 100
+    checkpoints_total_limit: 3
+    resume_from_checkpoint: latest
+
+  tracker:
+    project_name: fastvideo
+    run_name: wan_finetune
+
+  model:
+    enable_gradient_checkpointing_type: full
+
+callbacks:
+  grad_clip:
+    _target_: fastvideo.train.callbacks.grad_clip.GradNormClipCallback
+    max_grad_norm: 1.0
+  validation:
+    _target_: fastvideo.train.callbacks.validation.ValidationCallback
+    pipeline_target: fastvideo.pipelines.basic.wan.wan_pipeline.WanPipeline
+    dataset_file: data/validation_prompts.json
+    every_steps: 100
+    sampling_steps: [50]
+    guidance_scale: 5.0
+
+pipeline:
+  flow_shift: 3

examples/train/configs/kdtest/usage.md

@@ -1,8 +0,0 @@
-# Step 1 — KD training                                                                                      
-  bash examples/train/run.sh examples/train/configs/kdtest/kd_causal.yaml                                     
-                                                                                                              
-  # Step 2 — Export ode_init                                                                                  
-python -m fastvideo.train.entrypoint.dcp_to_diffusers --role student --checkpoint outputs/kdtest/kd_causal/checkpoint-300 --output-dir outputs/kdtest/kd_ode_init
-                                                                                                              
-  # Step 3 — Self-Forcing                                                                                     
-  bash examples/train/run.sh examples/train/configs/kdtest/self_forcing_causal.yaml

examples/train/configs/overfit_hunyuan_t2v.yaml

@@ -0,0 +1,78 @@
+# HunyuanVideo T2V overfitting test config.
+#
+# Overfits on 3 short videos (480x832, 77 frames) to verify the
+# training pipeline works end-to-end.
+#
+# Preprocess data first:
+#   CUDA_VISIBLE_DEVICES=0 python scripts/preprocess_hunyuan_overfit.py
+#
+# Run:
+#   bash examples/train/run.sh examples/train/configs/overfit_hunyuan_t2v.yaml
+
+models:
+  student:
+    _target_: fastvideo.train.models.hunyuan.HunyuanModel
+    init_from: hunyuanvideo-community/HunyuanVideo
+    trainable: true
+    enable_gradient_checkpointing_type: full
+
+method:
+  _target_: fastvideo.train.methods.fine_tuning.finetune.FineTuneMethod
+
+training:
+  distributed:
+    num_gpus: 8
+    sp_size: 1
+    tp_size: 1
+    hsdp_replicate_dim: 1
+    hsdp_shard_dim: 8
+
+  data:
+    data_path: data/hunyuan_overfit_preprocessed
+    dataloader_num_workers: 0
+    train_batch_size: 1
+    training_cfg_rate: 0.0
+    seed: 42
+    num_latent_t: 20
+    num_height: 480
+    num_width: 832
+    num_frames: 77
+
+  optimizer:
+    learning_rate: 5.0e-5
+    betas: [0.9, 0.999]
+    weight_decay: 0.0
+    lr_scheduler: constant
+    lr_warmup_steps: 0
+
+  loop:
+    max_train_steps: 300
+    gradient_accumulation_steps: 1
+
+  checkpoint:
+    output_dir: outputs/hunyuan_overfit
+    training_state_checkpointing_steps: 50
+    checkpoints_total_limit: 2
+
+  tracker:
+    project_name: fastvideo_hunyuan
+    run_name: hunyuan_overfit
+
+  model:
+    precondition_outputs: true
+    enable_gradient_checkpointing_type: full
+
+callbacks:
+  grad_clip:
+    _target_: fastvideo.train.callbacks.grad_clip.GradNormClipCallback
+    max_grad_norm: 1.0
+  validation:
+    _target_: fastvideo.train.callbacks.validation.ValidationCallback
+    pipeline_target: fastvideo.pipelines.basic.hunyuan.hunyuan_pipeline.HunyuanVideoPipeline
+    dataset_file: data/hunyuan_overfit_preprocessed/validation_prompts.json
+    every_steps: 50
+    sampling_steps: [50]
+    guidance_scale: 1.0
+
+pipeline:
+  flow_shift: 7

examples/train/run.sh

@@ -5,12 +5,12 @@
 #   bash examples/train/run.sh <config.yaml> [--dotted.key value ...]
 #
 # Examples:
-#   bash examples/train/run.sh examples/train/finetune_wan2.1_t2v_1.3B_vsa_phase3.4_0.9sparsity.yaml
-#   bash examples/train/run.sh examples/train/distill_wan2.1_t2v_1.3B_dmd2.yaml --dry-run
-#   bash examples/train/run.sh examples/train/distill_wan2.1_t2v_1.3B_dmd2.yaml \
+#   bash examples/train/run.sh examples/train/configs/fine_tuning/wan/t2v.yaml
+#   bash examples/train/run.sh examples/train/configs/distribution_matching/wan/dmd2_t2v.yaml --dry-run
+#   bash examples/train/run.sh examples/train/configs/distribution_matching/wan/dmd2_t2v.yaml \
 #       --training.distributed.num_gpus 4 \
 #       --training.optimizer.learning_rate 1e-5
-#   bash examples/train/run.sh examples/train/distill_wan2.1_t2v_1.3B_dmd2.yaml \
+#   bash examples/train/run.sh examples/train/configs/distribution_matching/wan/dmd2_t2v.yaml \
 #       --training.checkpoint.resume_from_checkpoint outputs/my_run/checkpoint-1000
 #
 # Logs are written to logs/<config_name>_<timestamp>.log (and also printed to stdout).

examples/train/run_slurm.sh

@@ -6,7 +6,7 @@
 #
 # Examples:
 #   bash examples/train/run_slurm.sh examples/train/configs/example.yaml 2
-#   bash examples/train/run_slurm.sh examples/train/configs/distill_wan2.1_t2v_1.3B_dmd2.yaml 4 \
+#   bash examples/train/run_slurm.sh examples/train/configs/distribution_matching/wan/dmd2_t2v.yaml 4 \
 #       --training.optimizer.learning_rate 1e-5
 #   bash examples/train/run_slurm.sh examples/train/configs/example.yaml 8 \
 #       --training.checkpoint.resume_from_checkpoint outputs/my_run/checkpoint-1000

examples/train/scenario/README.md

@@ -0,0 +1,13 @@
+# Training Scenarios
+
+End-to-end multi-step training pipelines. Each scenario directory contains
+all configs, scripts, and data needed to run a complete workflow.
+
+```
+scenario/
+└── ode_init_self_forcing_wan_causal/   # KD → export → Self-Forcing
+```
+
+See the `usage.md` inside each scenario for step-by-step instructions.
+
+For single-step configs, see `examples/train/configs/`.

examples/train/scenario/ode_init_self_forcing_wan_causal/decode_kd_cache.py

@@ -4,7 +4,7 @@ Uses fastvideo's AutoencoderKLWan and its built-in denormalization
 (latent * std + mean, matching normalize_dit_input's inverse).
 
 Usage:
-    python examples/train/configs/kdtest/decode_kd_cache.py \
+    python examples/train/scenario/ode_init_self_forcing_wan_causal/decode_kd_cache.py \
         --cache_dir data/kd_test_cache_small \
         --out_dir   data/kd_test_cache_small/video \
         --fps 16

examples/train/scenario/ode_init_self_forcing_wan_causal/step1_kd.yaml

@@ -5,7 +5,7 @@
 #
 # Step 1 — Run KD training:
 #   torchrun --nproc_per_node=1 -m fastvideo.train.entrypoint.train \
-#       --config examples/train/configs/kdtest/kd_causal.yaml
+#       --config examples/train/scenario/ode_init_self_forcing_wan_causal/step1_kd.yaml
 #
 # Step 2 — Export ode_init checkpoint:
 #   python -m fastvideo.train.entrypoint.dcp_to_diffusers \
@@ -15,7 +15,7 @@
 #
 # Step 3 — Run Self-Forcing with the exported ode_init:
 #   torchrun --nproc_per_node=1 -m fastvideo.train.entrypoint.train \
-#       --config examples/train/configs/kdtest/self_forcing_causal.yaml
+#       --config examples/train/scenario/ode_init_self_forcing_wan_causal/step3_self_forcing.yaml
 #
 # Success criteria:
 #   - Cache generates 16 .pt files under data/kd_test_cache/samples/

examples/train/scenario/ode_init_self_forcing_wan_causal/step2_export.sh

@@ -0,0 +1,21 @@
+#!/usr/bin/env bash
+# Step 2 — Export ODE-init weights from KD checkpoint to diffusers format.
+#
+# Usage:
+#   bash examples/train/scenario/ode_init_self_forcing_wan_causal/step2_export.sh
+
+set -euo pipefail
+
+CHECKPOINT_DIR="${1:-outputs/kdtest/kd_causal/checkpoint-300}"
+OUTPUT_DIR="${2:-outputs/kdtest/kd_ode_init}"
+
+echo "Exporting ODE-init weights..."
+echo "  checkpoint: ${CHECKPOINT_DIR}"
+echo "  output:     ${OUTPUT_DIR}"
+
+python -m fastvideo.train.entrypoint.dcp_to_diffusers \
+    --role student \
+    --checkpoint "${CHECKPOINT_DIR}" \
+    --output-dir "${OUTPUT_DIR}"
+
+echo "Done. Use the exported weights in step3_self_forcing.yaml."

examples/train/scenario/ode_init_self_forcing_wan_causal/step3_self_forcing.yaml

@@ -1,7 +1,7 @@
 # Self-Forcing Test: Causal Wan 1.3B
 #
 # Smoke-test for SelfForcingMethod on a tiny 16-sample dataset.
-# Run this AFTER kd_causal.yaml + dcp_to_diffusers export.
+# Run this AFTER step1_kd.yaml + dcp_to_diffusers export.
 #
 # Prerequisite — export ode_init from the KD run:
 #   python -m fastvideo.train.entrypoint.dcp_to_diffusers \
@@ -11,7 +11,7 @@
 #
 # Then run this config:
 #   torchrun --nproc_per_node=1 -m fastvideo.train.entrypoint.train \
-#       --config examples/train/configs/kdtest/self_forcing_causal.yaml
+#       --config examples/train/scenario/ode_init_self_forcing_wan_causal/step3_self_forcing.yaml
 #
 # Success criteria:
 #   - generator_loss and critic_loss both finite at step 0
@@ -100,7 +100,7 @@ callbacks:
     max_grad_norm: 1.0
   validation:
     pipeline_target: fastvideo.pipelines.basic.wan.wan_causal_dmd_pipeline.WanCausalDMDPipeline
-    dataset_file: examples/train/configs/kdtest/validation.json
+    dataset_file: examples/train/scenario/ode_init_self_forcing_wan_causal/validation.json
     every_steps: 20
     sampling_steps: [4]
     sampling_timesteps: [1000, 750, 500, 250]

examples/train/scenario/ode_init_self_forcing_wan_causal/usage.md

@@ -0,0 +1,19 @@
+# ODE-Init Self-Forcing: Wan 2.1 Causal 1.3B
+
+End-to-end scenario that trains a causal Wan 1.3B model via
+ODE-init knowledge distillation followed by self-forcing.
+
+## Steps
+
+```bash
+# Step 1 — KD training (produces ODE-init checkpoint)
+bash examples/train/run.sh \
+    examples/train/scenario/ode_init_self_forcing_wan_causal/step1_kd.yaml
+
+# Step 2 — Export ODE-init weights to diffusers format
+bash examples/train/scenario/ode_init_self_forcing_wan_causal/step2_export.sh
+
+# Step 3 — Self-forcing with the exported ODE-init
+bash examples/train/run.sh \
+    examples/train/scenario/ode_init_self_forcing_wan_causal/step3_self_forcing.yaml
+```

examples/training/finetune/wan_t2v_1.3B/crush_smol/finetune_t2v.sh

@@ -1,27 +1,47 @@
 #!/bin/bash
+#SBATCH --job-name=wan_t2v_1.3B_finetune
+#SBATCH --partition=all
+#SBATCH --nodes=1
+#SBATCH --gres=gpu:4
+#SBATCH --ntasks-per-node=1
+#SBATCH --output=logs/wan_t2v_1.3B_finetune.out
+#SBATCH --error=logs/wan_t2v_1.3B_finetune.err
+
+source .venv/bin/activate
+
 
 export WANDB_BASE_URL="https://api.wandb.ai"
 export WANDB_MODE=online
 export TOKENIZERS_PARALLELISM=false
 # export FASTVIDEO_ATTENTION_BACKEND=TORCH_SDPA
 
+# export TRITON_PRINT_AUTOTUNING=1  # to print the best config
+export WANDB_API_KEY=YOUR_WANDB_API_KEY
 MODEL_PATH="Wan-AI/Wan2.1-T2V-1.3B-Diffusers"
-DATA_DIR="data/crush-smol_processed_t2v/combined_parquet_dataset/"
-VALIDATION_DATASET_FILE="$(dirname "$0")/validation.json"
-NUM_GPUS=4
+DATA_DIR=data/Wan-Syn_77x448x832_600k
+VALIDATION_DATASET_FILE="examples/training/finetune/wan_t2v_1.3B/crush_smol/validation.json"
+NUM_GPUS=1
 # export CUDA_VISIBLE_DEVICES=4,5
 
+set -euo pipefail
+
+# ---- torchrun rendezvous (multi-node) ----
+# Launch ONE torchrun per node (via srun) and let torchrun spawn 4 workers per node.
+MASTER_ADDR="$(scontrol show hostnames "$SLURM_JOB_NODELIST" | head -n 1)"
+MASTER_PORT="${MASTER_PORT:-29500}"
+export MASTER_ADDR MASTER_PORT
+
 
 # Training arguments
 training_args=(
-  --tracker_project_name "wan_t2v_finetune"
-  --output_dir "checkpoints/wan_t2v_finetune"
-  --max_train_steps 5000
+  --tracker_project_name "wan_t2v_finetune_qat"
+  --output_dir "checkpoints/wan_t2v_finetune_1.3B_77"
+  --max_train_steps 4000
   --train_batch_size 1
   --train_sp_batch_size 1
-  --gradient_accumulation_steps 8
+  --gradient_accumulation_steps 1
   --num_latent_t 20
-  --num_height 480
+  --num_height 448  
   --num_width 832
   --num_frames 77
   --enable_gradient_checkpointing_type "full"
@@ -30,7 +50,7 @@ training_args=(
 # Parallel arguments
 parallel_args=(
   --num_gpus $NUM_GPUS 
-  --sp_size $NUM_GPUS 
+  --sp_size 1
   --tp_size 1
   --hsdp_replicate_dim 1
   --hsdp_shard_dim $NUM_GPUS
@@ -45,7 +65,7 @@ model_args=(
 # Dataset arguments
 dataset_args=(
   --data_path $DATA_DIR
-  --dataloader_num_workers 1
+  --dataloader_num_workers 4
 )
 
 # Validation arguments
@@ -54,16 +74,16 @@ validation_args=(
   --validation_dataset_file $VALIDATION_DATASET_FILE
   --validation_steps 200
   --validation_sampling_steps "50" 
-  --validation_guidance_scale "3.0"
+  --validation_guidance_scale "5.0"
 )
 
 # Optimizer arguments
 optimizer_args=(
-  --learning_rate 5e-5
+  --learning_rate 1e-6
   --mixed_precision "bf16"
   --weight_only_checkpointing_steps 1000
   --training_state_checkpointing_steps 1000
-  --weight_decay 1e-4
+  --weight_decay 0.01
   --max_grad_norm 1.0
 )
 
@@ -72,23 +92,24 @@ miscellaneous_args=(
   --inference_mode False
   --checkpoints_total_limit 3
   --training_cfg_rate 0.1
-  --multi_phased_distill_schedule "4000-1"
-  --not_apply_cfg_solver
   --dit_precision "fp32"
-  --num_euler_timesteps 50
   --ema_start_step 0
-  --enable_gradient_checkpointing_type "full"
-  # --resume_from_checkpoint "checkpoints/wan_t2v_finetune/checkpoint-2500"
+  --flow_shift 5
+  --seed 1000
 )
 
-torchrun \
-  --nnodes 1 \
-  --nproc_per_node $NUM_GPUS \
-    fastvideo/training/wan_training_pipeline.py \
-    "${parallel_args[@]}" \
-    "${model_args[@]}" \
-    "${dataset_args[@]}" \
-    "${training_args[@]}" \
-    "${optimizer_args[@]}" \
-    "${validation_args[@]}" \
-    "${miscellaneous_args[@]}"
+srun --nodes="$SLURM_NNODES" --ntasks="$SLURM_NNODES" --ntasks-per-node=1 \
+  torchrun \
+    --nnodes "$SLURM_NNODES" \
+    --nproc_per_node 4 \
+    --rdzv_backend c10d \
+    --rdzv_endpoint "${MASTER_ADDR}:${MASTER_PORT}" \
+    --rdzv_id "$SLURM_JOB_ID" \
+      fastvideo/training/wan_training_pipeline.py \
+      "${parallel_args[@]}" \
+      "${model_args[@]}" \
+      "${dataset_args[@]}" \
+      "${training_args[@]}" \
+      "${optimizer_args[@]}" \
+      "${validation_args[@]}" \
+      "${miscellaneous_args[@]}"

examples/training/finetune/wan_t2v_1.3B/crush_smol/finetune_t2v_qat_attn.sh

@@ -0,0 +1,124 @@
+#!/bin/bash
+#SBATCH --job-name=wan_t2v_1.3B_finetune_qat_16
+#SBATCH --partition=all
+#SBATCH --nodes=4
+#SBATCH --gres=gpu:4
+#SBATCH --ntasks-per-node=1
+#SBATCH --output=logs/wan_t2v_1.3B_finetune_qat_16.out
+#SBATCH --error=logs/wan_t2v_1.3B_finetune_qat_16.err
+
+source .venv/bin/activate
+
+
+export WANDB_BASE_URL="https://api.wandb.ai"
+export WANDB_MODE=online
+export TOKENIZERS_PARALLELISM=false
+export FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_TRAIN
+
+# export TRITON_PRINT_AUTOTUNING=1  # to print the best config
+export WANDB_API_KEY=YOUR_WANDB_API_KEY
+# Use node-local Triton cache to avoid stale file handle errors on shared filesystems
+export TRITON_CACHE_DIR="/tmp/triton_cache_${SLURM_JOB_ID}_${SLURM_NODEID}"
+MODEL_PATH="Wan-AI/Wan2.1-T2V-1.3B-Diffusers"
+DATA_DIR=YOUR_DATA_DIR
+VALIDATION_DATASET_FILE="examples/training/finetune/wan_t2v_1.3B/crush_smol/validation.json"
+NUM_GPUS=16
+# export CUDA_VISIBLE_DEVICES=4,5
+
+set -euo pipefail
+
+# ---- torchrun rendezvous (multi-node) ----
+# 1. Get the hostname of the first node (Master)
+nodes=( $( scontrol show hostnames $SLURM_JOB_NODELIST ) )
+nodes_array=($nodes)
+head_node=${nodes_array[0]}
+MASTER_ADDR=$(srun --nodes=1 --ntasks=1 -w "$head_node" hostname --ip-address)
+MASTER_PORT=29500
+
+# 2. Get the node count automatically
+NNODES=$SLURM_NNODES
+GPUS_PER_NODE=$SLURM_GPUS_ON_NODE
+NUM_GPUS=$((NNODES * GPUS_PER_NODE))
+
+echo "MASTER_ADDR=$MASTER_ADDR MASTER_PORT=$MASTER_PORT NNODES=$NNODES"
+
+# Training arguments
+training_args=(
+  --tracker_project_name "wan_t2v_finetune_qat"
+  --output_dir "checkpoints/wan_1.3B_t2v_finetune_qat"
+  --max_train_steps 4000
+  --train_batch_size 1
+  --train_sp_batch_size 1
+  --gradient_accumulation_steps 1
+  --num_latent_t 20
+  --num_height 448  
+  --num_width 832
+  --num_frames 77
+  --enable_gradient_checkpointing_type "full" # if OOM enable this
+)
+
+# Parallel arguments
+parallel_args=(
+  --num_gpus $NUM_GPUS 
+  --sp_size 1
+  --tp_size 1
+  --hsdp_replicate_dim $NUM_GPUS
+  --hsdp_shard_dim 1
+)
+
+# Model arguments
+model_args=(
+  --model_path $MODEL_PATH
+  --pretrained_model_name_or_path $MODEL_PATH
+)
+
+# Dataset arguments
+dataset_args=(
+  --data_path "$DATA_DIR"
+  --dataloader_num_workers 4
+)
+
+# Validation arguments
+validation_args=(
+  --log_validation 
+  --validation_dataset_file $VALIDATION_DATASET_FILE
+  --validation_steps 200
+  --validation_sampling_steps "50" 
+  --validation_guidance_scale "5.0"
+)
+
+# Optimizer arguments
+optimizer_args=(
+  --learning_rate 1e-6
+  --mixed_precision "bf16"
+  --weight_only_checkpointing_steps 200
+  --training_state_checkpointing_steps 200
+  --weight_decay 0.01
+  --max_grad_norm 1.0
+)
+
+# Miscellaneous arguments
+miscellaneous_args=(
+  --inference_mode False
+  --checkpoints_total_limit 3
+  --training_cfg_rate 0.1
+  --dit_precision "fp32"
+  --ema_start_step 0
+  --flow_shift 1
+  --seed 1000
+)
+
+srun torchrun \
+    --nnodes $NNODES \
+    --nproc_per_node $GPUS_PER_NODE \
+    --node_rank $SLURM_PROCID \
+    --rdzv_backend c10d \
+    --rdzv_endpoint $MASTER_ADDR:$MASTER_PORT \
+      fastvideo/training/wan_training_pipeline.py \
+      "${parallel_args[@]}" \
+      "${model_args[@]}" \
+      "${dataset_args[@]}" \
+      "${training_args[@]}" \
+      "${optimizer_args[@]}" \
+      "${validation_args[@]}" \
+      "${miscellaneous_args[@]}"

examples/training/finetune/wan_t2v_1.3B/crush_smol/validation.json

@@ -1,31 +1,131 @@
 {
   "data": [
     {
-      "caption": "A large metal cylinder is seen pressing down on a pile of Oreo cookies, flattening them as if they were under a hydraulic press.",
-      "image_path": null,
-      "video_path": null,
-      "num_inference_steps": 50,
-      "height": 480,
+      "caption": "In the video, a woman is elegantly showcasing her earrings, bringing attention to their intricate design with a gentle touch of her fingers. She is bathed in ambient purple and pink lighting, which casts a soft glow on her delicate features and enhances the vivid tones of her lipstick and eye makeup. Her hair is styled to frame her face smoothly, emphasizing the contours of her jawline and cheekbones. The background features a blurred neon light, adding an artistic and modern touch to the overall aesthetic.",
+      "video_path": "Fashion/mixkit-face-of-an-elegant-and-captivating-woman-41914_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
       "width": 832,
       "num_frames": 77
     },
     {
-      "caption": "A large metal cylinder is seen compressing colorful clay into a compact shape, demonstrating the power of a hydraulic press.",
-      "image_path": null,
-      "video_path": null,
-      "num_inference_steps": 50,
-      "height": 480,
+      "caption": "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.",
+      "video_path": "Man/mixkit-a-rancher-riding-a-horse-at-sunset-1143_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
       "width": 832,
       "num_frames": 77
     },
     {
-      "caption": "A large metal cylinder is seen pressing down on a pile of colorful candies, flattening them as if they were under a hydraulic press. The candies are crushed and broken into small pieces, creating a mess on the table.",
-      "image_path": null,
-      "video_path": null,
-      "num_inference_steps": 50,
-      "height": 480,
+      "caption": "In a dimly lit, eerie setting, a mysterious pink bottle labeled \"Authentic 100% organic POISON\" sits prominently in the foreground, casting a menacing aura. The bottle is accentuated by green fog, which swirls lightly around it, enhancing its sinister allure. Behind it, a shadowy golden bottle adorned with a spider emblem subtly emerges, adding an extra layer of mystery to the scene. Dim candles provide faint, flickering light, which complements the dark atmosphere, making the setting ideal for an illusion of hidden dangers.",
+      "video_path": "smoke/mixkit-poison-in-halloween-ritual-33879_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
       "width": 832,
       "num_frames": 77
-    }
-  ]
+    },
+    {
+      "caption": "The video opens with a tranquil scene in the heart of a dense forest, emphasizing two large, textured tree trunks in the foreground framing the view. Sunlight filters through the canopy above, casting intricate patterns of light and shadow on the trees and the ground. Between the tree trunks, a clear view of a calm, muddy river unfolds, its surface shimmering under the gentle sunlight. The riverbank is decorated with a variety of small bushes and vibrant foliage, subtly transitioning into the deep greens of tall, leafy plants. In the background, the dense forest looms, filled with dark, towering trees, their branches intertwining to form an intricate canopy. The scene is bathed in the soft glow of the sun, creating a serene and picturesque setting. Occasional sunbeams pierce through the foliage, adding a magical aura to the landscape. The vibrant reds and oranges of the smaller plants add contrast, bringing warmth to the earthy tones of the scenery. Overall, this harmonious blend of natural elements creates a peaceful and idyllic forest setting.",
+      "video_path": "forest/mixkit-view-of-a-river-between-two-old-trees-560_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "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.",
+      "video_path": "Man/mixkit-a-young-man-practicing-his-karate-moves-49635_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "A tranquil coastal scene unfolds with a drone's aerial view capturing a serene beach landscape. The camera glides over a quiet stretch of sandy shoreline, where gentle waves kiss the shore under a clear blue sky. Nestled amidst lush palm trees are a series of traditional thatched-roof huts, their earthy tones blending harmoniously with the natural surroundings. The sandy beach stretches endlessly, bordered by the rhythmic dance of ocean waves on one side and verdant greenery on the other. A pair of white umbrellas is set up on the sand, suggesting a place to relax and enjoy the sun. In the distance, two small human figures can be seen walking leisurely along the water's edge, leaving faint footprints behind them. The scene exudes a calm and inviting atmosphere, with the soft rustle of palm leaves and the whisper of the ocean breeze almost audible. The overall composition is a captivating blend of nature's tranquility and architectural simplicity. This picturesque setting invites viewers to imagine themselves steps away from this idyllic coastal escape.",
+      "video_path": "beach/mixkit-sunny-beach-in-a-dynamic-shot-from-a-drone-44383_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "A lone figure stands on a large, moss-covered rock, surrounded by the soft rush of a nearby stream. The figure is wearing white sneakers and shorts, with a plaid shirt that hangs loosely in the breeze. The lighting creates dramatic shadows, enhancing the textures of the rock and the subtle movement of the water below. In the background, a waterfall cascades into the stream, completing this tranquil and serene nature scene.",
+      "video_path": "forest/mixkit-woman-standing-in-front-of-waterfall-559_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "In an industrial setting, a person leans casually against a railing, exuding a sense of confidence and composure. They are wearing a striking outfit, consisting of a vibrant, patterned jacket over a simple white crop top, creating a bold contrast. The atmosphere is infused with warm, ambient lighting that casts soft shadows on the concrete walls and metallic surfaces. Intricate wiring and pipes form an intricate backdrop, enhancing the urban aesthetic. Their relaxed posture and direct, engaging gaze suggest a sense of ease in this industrial environment. This scene encapsulates a blend of modern fashion and gritty, urban architecture, creating a visually compelling narrative.",
+      "video_path": "Fashion/mixkit-portrait-of-a-hipster-woman-walking-down-a-stairs-1297_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "A man is energetically stretching in an open-air setting, surrounded by rows of vibrant red seats that suggest an amphitheater or outdoor venue. He wears a sleeveless black shirt layered with a hooded vest, emphasizing his athletic build as he engages in a warm-up routine. Behind him, the striking modern architecture of the building features geometric panels, with large sections of glass and overlapping metallic beams creating a dynamic backdrop. The scene captures the contrast between his focused movements and the static, bold design of the structure, while the surrounding greenery adds a touch of nature to the environment. The overall atmosphere is one of preparation and anticipation, with the man appearing determined and ready for an upcoming event or performance.",
+      "video_path": "Sport/mixkit-man-doing-arm-stretches-595_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "A young woman is seated on the floor in front of a plush, beige tufted couch, fully engrossed in sorting through a stack of papers. Her dark hair falls loosely past her shoulders, and she wears a green plaid shirt, contributing to the casual yet focused atmosphere. She gently places the papers onto a small round white table, occasionally lifting individual sheets to examine them more closely. Her expression shifts subtly, reflecting concentration and contemplation as she processes the information on the pages. Two small, round nested tables hold her documents, along with a small plant in a gray pot, adding a touch of greenery to the scene. The background features a dark paneled wall, creating a contrasting backdrop for the light-colored furniture. The setting is tranquil and organized, the couch and tables arranged symmetrically, conveying a sense of harmony. A calculator rests on the smaller table, hinting at a task involving calculations or budgeting.",
+      "video_path": "Woman/mixkit-frustrated-woman-throws-paperwork-on-the-floor-4526_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "A heavily rusted metal gate stands firmly locked, with two vertical bars joined by a thick, old chain that loops elegantly around them. The chain's texture is coarse and rugged, its surface reflecting varying shades of orange and brown, indicative of years exposed to the elements. At the heart of the chain, a black iron padlock, slightly worn yet imposing, secures the gate, its curves and edges smooth against the aged links. The gate's metalwork is outlined by a backdrop of soft, blurred greenery, suggesting a serene and isolated location beyond the barrier. Tall trees rise in the distance, their trunks and leaves creating a lush, forest-like setting that contrasts with the gate's severe rust. A pathway leads away from the gate, its surface uneven with patches of moss and weathered stone visible in the soft focus, inviting yet inaccessible. The ambiance is quiet and mysterious, with a sense of abandonment hanging subtly in the air, evoking curiosity about what lies beyond. Shadows play across the gate, cast by branches swaying gently in the breeze, adding to the dynamic interaction of light and texture. This scene, rich in detail and atmosphere, captures the viewer's imagination, evoking both the allure of the forbidden and the beauty of decay.",
+      "video_path": "forest/mixkit-rusty-fence-with-a-chain-of-a-property-in-nature-5294_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "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.",
+      "video_path": "People/mixkit-small-group-of-people-doing-yoga-together-43730_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "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.",
+      "video_path": "sea/mixkit-dolphins-underwater-4133_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "In the video, a young woman stands against a vibrant graffiti-covered wall, deeply engrossed in her smartphone. Her expression reflects a mix of focus and subtle satisfaction as she interacts with the screen. She wears a black floral-patterned top, which contrasts with the bright, abstract shapes and bold colors of the mural behind her. As she continues to engage with her phone, a series of like count notifications appear on the screen, indicating a growing online appreciation. The wall behind her features a striking mix of geometric and organic shapes, including swirls of teal, orange, and black, with large humanoid figures in a pop-art style. Her long, light-brown hair frames her face, adding a calm, composed aura amidst the lively backdrop. The video captures a blend of contemporary digital interaction and expressive urban art, creating a dynamic yet harmonious scene.",
+      "video_path": "Girl/mixkit-girl-looking-at-the-likes-in-her-post-4914_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "A young mother and her baby sit comfortably on a bed, surrounded by an inviting, cozy atmosphere. The woman, wearing a sleeveless top and jeans, is gently engaging with the baby, who is dressed in an adorable animal-print onesie. The child is seated on the bed with colorful toys scattered around, including a plush toy and a board book. The warm glow from a hanging lamp casts a soft light on them, enhancing the serene environment. Pillows are propped up against the headboard, providing a cushioned backdrop as the mother leans slightly over to interact with the baby. A small bottle is visible beside her, suggesting a nurturing setting. Her hand gestures animatedly as she holds up a soft, white cushion with red and blue accents, likely stimulating the baby\u2019s curiosity. Their shared moment is filled with affection and joy, a perfect snapshot of familial bonding.",
+      "video_path": "Baby/mixkit-loving-mother-and-her-baby-playing-with-soft-toys-49966_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+    },
+    {
+      "caption": "A young girl with long brown hair sits at a round wooden table, engrossed in working on her laptop. The laptop screen is a vivid green, suggesting a green screen effect is in use. To her left, a doll dressed in a yellow and white outfit is casually laid on top of some books, adding a playful and innocent touch to the scene. The setting is cozy, with sheer curtains in the background allowing soft natural light to spill into the room. The girl's posture and focused attention on the laptop suggest she is either playing a game or learning something new. This serene and domestic atmosphere is complemented by the slight blur of a dark couch in the foreground, framing the focused activity of the child.",
+      "video_path": "Girl/mixkit-little-girl-doing-homework-on-a-laptop-4757_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 448,
+      "width": 832,
+      "num_frames": 77
+  }  ]
 }

examples/training/finetune/wan_t2v_14B/finetune_t2v_qat_attn.sh

@@ -0,0 +1,123 @@
+#!/bin/bash
+#SBATCH --job-name=wan_t2v_1.3B_finetune_qat_16
+#SBATCH --partition=main
+#SBATCH --nodes=4
+#SBATCH --ntasks-per-node=1
+#SBATCH --gres=gpu:8
+#SBATCH --cpus-per-task=128
+#SBATCH --mem=1440G
+#SBATCH --output=logs/wan_t2v_1.3B_finetune_qat_16.out
+#SBATCH --error=logs/wan_t2v_1.3B_finetune_qat_16.err
+#SBATCH --exclusive
+
+source ~/conda/miniconda/bin/activate
+conda activate matthew-fv
+
+# Basic Info
+export WANDB_MODE="online"
+export NCCL_P2P_DISABLE=1
+export TORCH_NCCL_ENABLE_MONITORING=0
+# different cache dir for different processes
+export TRITON_CACHE_DIR=/tmp/triton_cache_${SLURM_PROCID}
+export MASTER_PORT=29500
+export NODE_RANK=$SLURM_PROCID
+nodes=( $(scontrol show hostnames $SLURM_JOB_NODELIST) )
+export MASTER_ADDR=${nodes[0]}
+export CUDA_VISIBLE_DEVICES=$SLURM_LOCALID
+export TOKENIZERS_PARALLELISM=false
+export WANDB_BASE_URL="https://api.wandb.ai"
+export WANDB_MODE=online
+export FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_TRAIN
+
+echo "MASTER_ADDR: $MASTER_ADDR"
+echo "NODE_RANK: $NODE_RANK"
+
+# export TRITON_PRINT_AUTOTUNING=1  # to print the best config
+export WANDB_API_KEY=YOUR_WANDB_API_KEY
+MODEL_PATH="Wan-AI/Wan2.1-T2V-14B-Diffusers"
+DATA_DIR=YOUR_DATA_DIR
+VALIDATION_DATASET_FILE="examples/training/finetune/wan_t2v_1.3B/crush_smol/validation.json"
+NUM_GPUS_PER_NODE=8
+TOTAL_GPUS=$((NUM_GPUS_PER_NODE * SLURM_JOB_NUM_NODES))
+# export CUDA_VISIBLE_DEVICES=4,5
+
+# Training arguments
+training_args=(
+  --tracker_project_name "wan_t2v_finetune_qat"
+  --output_dir "checkpoints/wan_14B_t2v_finetune_qat"
+  --max_train_steps 4000
+  --train_batch_size 1
+  --train_sp_batch_size 1
+  --gradient_accumulation_steps 1
+  --num_latent_t 20
+  --num_height 768 
+  --num_width 1280
+  --num_frames 77
+  --enable_gradient_checkpointing_type "full" # if OOM enable this
+)
+
+# Parallel arguments
+parallel_args=(
+  --num_gpus $TOTAL_GPUS 
+  --sp_size 4
+  --tp_size 1
+  --hsdp_replicate_dim 4
+  --hsdp_shard_dim 8
+)
+
+# Model arguments
+model_args=(
+  --model_path $MODEL_PATH
+  --pretrained_model_name_or_path $MODEL_PATH
+)
+
+# Dataset arguments
+dataset_args=(
+  --data_path "$DATA_DIR"
+  --dataloader_num_workers 4
+)
+
+# Validation arguments
+validation_args=(
+  # --log_validation 
+  --validation_dataset_file $VALIDATION_DATASET_FILE
+  --validation_steps 200
+  --validation_sampling_steps "50" 
+  --validation_guidance_scale "5.0"
+)
+
+# Optimizer arguments
+optimizer_args=(
+  --learning_rate 1e-6
+  --mixed_precision "bf16"
+  --weight_only_checkpointing_steps 200
+  --training_state_checkpointing_steps 200
+  --weight_decay 0.01
+  --max_grad_norm 1.0
+)
+
+# Miscellaneous arguments
+miscellaneous_args=(
+  --inference_mode False
+  --checkpoints_total_limit 3
+  --training_cfg_rate 0.1
+  --dit_precision "fp32"
+  --ema_start_step 0
+  --flow_shift 5
+  --seed 1000
+)
+
+srun torchrun \
+--nnodes $SLURM_JOB_NUM_NODES \
+--nproc_per_node $NUM_GPUS_PER_NODE \
+--node_rank $SLURM_PROCID \
+--rdzv_backend=c10d \
+--rdzv_endpoint="$MASTER_ADDR:$MASTER_PORT" \
+    fastvideo/training/wan_training_pipeline.py \
+    "${parallel_args[@]}" \
+    "${model_args[@]}" \
+    "${dataset_args[@]}" \
+    "${training_args[@]}" \
+    "${optimizer_args[@]}" \
+    "${validation_args[@]}" \
+    "${miscellaneous_args[@]}"

examples/training/finetune/wan_t2v_14B/validation.json

@@ -0,0 +1,131 @@
+{
+  "data": [
+    {
+      "caption": "In the video, a woman is elegantly showcasing her earrings, bringing attention to their intricate design with a gentle touch of her fingers. She is bathed in ambient purple and pink lighting, which casts a soft glow on her delicate features and enhances the vivid tones of her lipstick and eye makeup. Her hair is styled to frame her face smoothly, emphasizing the contours of her jawline and cheekbones. The background features a blurred neon light, adding an artistic and modern touch to the overall aesthetic.",
+      "video_path": "Fashion/mixkit-face-of-an-elegant-and-captivating-woman-41914_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "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.",
+      "video_path": "Man/mixkit-a-rancher-riding-a-horse-at-sunset-1143_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "In a dimly lit, eerie setting, a mysterious pink bottle labeled \"Authentic 100% organic POISON\" sits prominently in the foreground, casting a menacing aura. The bottle is accentuated by green fog, which swirls lightly around it, enhancing its sinister allure. Behind it, a shadowy golden bottle adorned with a spider emblem subtly emerges, adding an extra layer of mystery to the scene. Dim candles provide faint, flickering light, which complements the dark atmosphere, making the setting ideal for an illusion of hidden dangers.",
+      "video_path": "smoke/mixkit-poison-in-halloween-ritual-33879_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "The video opens with a tranquil scene in the heart of a dense forest, emphasizing two large, textured tree trunks in the foreground framing the view. Sunlight filters through the canopy above, casting intricate patterns of light and shadow on the trees and the ground. Between the tree trunks, a clear view of a calm, muddy river unfolds, its surface shimmering under the gentle sunlight. The riverbank is decorated with a variety of small bushes and vibrant foliage, subtly transitioning into the deep greens of tall, leafy plants. In the background, the dense forest looms, filled with dark, towering trees, their branches intertwining to form an intricate canopy. The scene is bathed in the soft glow of the sun, creating a serene and picturesque setting. Occasional sunbeams pierce through the foliage, adding a magical aura to the landscape. The vibrant reds and oranges of the smaller plants add contrast, bringing warmth to the earthy tones of the scenery. Overall, this harmonious blend of natural elements creates a peaceful and idyllic forest setting.",
+      "video_path": "forest/mixkit-view-of-a-river-between-two-old-trees-560_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "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.",
+      "video_path": "Man/mixkit-a-young-man-practicing-his-karate-moves-49635_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "A tranquil coastal scene unfolds with a drone's aerial view capturing a serene beach landscape. The camera glides over a quiet stretch of sandy shoreline, where gentle waves kiss the shore under a clear blue sky. Nestled amidst lush palm trees are a series of traditional thatched-roof huts, their earthy tones blending harmoniously with the natural surroundings. The sandy beach stretches endlessly, bordered by the rhythmic dance of ocean waves on one side and verdant greenery on the other. A pair of white umbrellas is set up on the sand, suggesting a place to relax and enjoy the sun. In the distance, two small human figures can be seen walking leisurely along the water's edge, leaving faint footprints behind them. The scene exudes a calm and inviting atmosphere, with the soft rustle of palm leaves and the whisper of the ocean breeze almost audible. The overall composition is a captivating blend of nature's tranquility and architectural simplicity. This picturesque setting invites viewers to imagine themselves steps away from this idyllic coastal escape.",
+      "video_path": "beach/mixkit-sunny-beach-in-a-dynamic-shot-from-a-drone-44383_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "A lone figure stands on a large, moss-covered rock, surrounded by the soft rush of a nearby stream. The figure is wearing white sneakers and shorts, with a plaid shirt that hangs loosely in the breeze. The lighting creates dramatic shadows, enhancing the textures of the rock and the subtle movement of the water below. In the background, a waterfall cascades into the stream, completing this tranquil and serene nature scene.",
+      "video_path": "forest/mixkit-woman-standing-in-front-of-waterfall-559_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "In an industrial setting, a person leans casually against a railing, exuding a sense of confidence and composure. They are wearing a striking outfit, consisting of a vibrant, patterned jacket over a simple white crop top, creating a bold contrast. The atmosphere is infused with warm, ambient lighting that casts soft shadows on the concrete walls and metallic surfaces. Intricate wiring and pipes form an intricate backdrop, enhancing the urban aesthetic. Their relaxed posture and direct, engaging gaze suggest a sense of ease in this industrial environment. This scene encapsulates a blend of modern fashion and gritty, urban architecture, creating a visually compelling narrative.",
+      "video_path": "Fashion/mixkit-portrait-of-a-hipster-woman-walking-down-a-stairs-1297_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "A man is energetically stretching in an open-air setting, surrounded by rows of vibrant red seats that suggest an amphitheater or outdoor venue. He wears a sleeveless black shirt layered with a hooded vest, emphasizing his athletic build as he engages in a warm-up routine. Behind him, the striking modern architecture of the building features geometric panels, with large sections of glass and overlapping metallic beams creating a dynamic backdrop. The scene captures the contrast between his focused movements and the static, bold design of the structure, while the surrounding greenery adds a touch of nature to the environment. The overall atmosphere is one of preparation and anticipation, with the man appearing determined and ready for an upcoming event or performance.",
+      "video_path": "Sport/mixkit-man-doing-arm-stretches-595_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "A young woman is seated on the floor in front of a plush, beige tufted couch, fully engrossed in sorting through a stack of papers. Her dark hair falls loosely past her shoulders, and she wears a green plaid shirt, contributing to the casual yet focused atmosphere. She gently places the papers onto a small round white table, occasionally lifting individual sheets to examine them more closely. Her expression shifts subtly, reflecting concentration and contemplation as she processes the information on the pages. Two small, round nested tables hold her documents, along with a small plant in a gray pot, adding a touch of greenery to the scene. The background features a dark paneled wall, creating a contrasting backdrop for the light-colored furniture. The setting is tranquil and organized, the couch and tables arranged symmetrically, conveying a sense of harmony. A calculator rests on the smaller table, hinting at a task involving calculations or budgeting.",
+      "video_path": "Woman/mixkit-frustrated-woman-throws-paperwork-on-the-floor-4526_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "A heavily rusted metal gate stands firmly locked, with two vertical bars joined by a thick, old chain that loops elegantly around them. The chain's texture is coarse and rugged, its surface reflecting varying shades of orange and brown, indicative of years exposed to the elements. At the heart of the chain, a black iron padlock, slightly worn yet imposing, secures the gate, its curves and edges smooth against the aged links. The gate's metalwork is outlined by a backdrop of soft, blurred greenery, suggesting a serene and isolated location beyond the barrier. Tall trees rise in the distance, their trunks and leaves creating a lush, forest-like setting that contrasts with the gate's severe rust. A pathway leads away from the gate, its surface uneven with patches of moss and weathered stone visible in the soft focus, inviting yet inaccessible. The ambiance is quiet and mysterious, with a sense of abandonment hanging subtly in the air, evoking curiosity about what lies beyond. Shadows play across the gate, cast by branches swaying gently in the breeze, adding to the dynamic interaction of light and texture. This scene, rich in detail and atmosphere, captures the viewer's imagination, evoking both the allure of the forbidden and the beauty of decay.",
+      "video_path": "forest/mixkit-rusty-fence-with-a-chain-of-a-property-in-nature-5294_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "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.",
+      "video_path": "People/mixkit-small-group-of-people-doing-yoga-together-43730_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "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.",
+      "video_path": "sea/mixkit-dolphins-underwater-4133_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "In the video, a young woman stands against a vibrant graffiti-covered wall, deeply engrossed in her smartphone. Her expression reflects a mix of focus and subtle satisfaction as she interacts with the screen. She wears a black floral-patterned top, which contrasts with the bright, abstract shapes and bold colors of the mural behind her. As she continues to engage with her phone, a series of like count notifications appear on the screen, indicating a growing online appreciation. The wall behind her features a striking mix of geometric and organic shapes, including swirls of teal, orange, and black, with large humanoid figures in a pop-art style. Her long, light-brown hair frames her face, adding a calm, composed aura amidst the lively backdrop. The video captures a blend of contemporary digital interaction and expressive urban art, creating a dynamic yet harmonious scene.",
+      "video_path": "Girl/mixkit-girl-looking-at-the-likes-in-her-post-4914_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "A young mother and her baby sit comfortably on a bed, surrounded by an inviting, cozy atmosphere. The woman, wearing a sleeveless top and jeans, is gently engaging with the baby, who is dressed in an adorable animal-print onesie. The child is seated on the bed with colorful toys scattered around, including a plush toy and a board book. The warm glow from a hanging lamp casts a soft light on them, enhancing the serene environment. Pillows are propped up against the headboard, providing a cushioned backdrop as the mother leans slightly over to interact with the baby. A small bottle is visible beside her, suggesting a nurturing setting. Her hand gestures animatedly as she holds up a soft, white cushion with red and blue accents, likely stimulating the baby\u2019s curiosity. Their shared moment is filled with affection and joy, a perfect snapshot of familial bonding.",
+      "video_path": "Baby/mixkit-loving-mother-and-her-baby-playing-with-soft-toys-49966_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+    },
+    {
+      "caption": "A young girl with long brown hair sits at a round wooden table, engrossed in working on her laptop. The laptop screen is a vivid green, suggesting a green screen effect is in use. To her left, a doll dressed in a yellow and white outfit is casually laid on top of some books, adding a playful and innocent touch to the scene. The setting is cozy, with sheer curtains in the background allowing soft natural light to spill into the room. The girl's posture and focused attention on the laptop suggest she is either playing a game or learning something new. This serene and domestic atmosphere is complemented by the slight blur of a dark couch in the foreground, framing the focused activity of the child.",
+      "video_path": "Girl/mixkit-little-girl-doing-homework-on-a-laptop-4757_clip_1.mp4",
+      "num_inference_steps":  50,
+      "height": 768,
+      "width": 1280,
+      "num_frames": 77
+  }  ]
+}

fastvideo-kernel/CMakeLists.txt

@@ -12,6 +12,18 @@ else()
     enable_language(CUDA)
     # Ensure CUDA toolkit targets (CUDA::cudart, CUDA::cuda_driver, etc.) are available.
     find_package(CUDAToolkit REQUIRED)
+    if(NOT DEFINED CUDA_TOOLKIT_ROOT_DIR)
+        if(DEFINED CUDAToolkit_ROOT)
+            set(CUDA_TOOLKIT_ROOT_DIR "${CUDAToolkit_ROOT}" CACHE PATH
+                "CUDA toolkit root directory" FORCE)
+        elseif(DEFINED ENV{CUDAToolkit_ROOT})
+            set(CUDA_TOOLKIT_ROOT_DIR "$ENV{CUDAToolkit_ROOT}" CACHE PATH
+                "CUDA toolkit root directory" FORCE)
+        elseif(DEFINED ENV{CUDA_HOME})
+            set(CUDA_TOOLKIT_ROOT_DIR "$ENV{CUDA_HOME}" CACHE PATH
+                "CUDA toolkit root directory" FORCE)
+        endif()
+    endif()
 endif()
 
 # Import common utils if needed, but we keep it simple for now
@@ -19,13 +31,46 @@ endif()
 # Find Python and Torch
 find_package(Python COMPONENTS Interpreter Development.Module REQUIRED)
 
-# Robustly find Torch include paths using Python
+# Locate the installed torch package without importing it. This keeps CMake
+# configure working even on nodes where CUDA runtime libraries are not yet on
+# the dynamic loader path.
 execute_process(
-    COMMAND "${Python_EXECUTABLE}" -c "import torch; from torch.utils.cpp_extension import include_paths; print(';'.join(include_paths()))"
-    OUTPUT_VARIABLE TORCH_INCLUDE_PATHS
+    COMMAND "${Python_EXECUTABLE}" -c "import sysconfig; print(sysconfig.get_path('platlib'))"
+    OUTPUT_VARIABLE PYTHON_PLATLIB
     OUTPUT_STRIP_TRAILING_WHITESPACE
 )
-list(APPEND TORCH_INCLUDE_DIRS ${TORCH_INCLUDE_PATHS})
+execute_process(
+    COMMAND "${Python_EXECUTABLE}" -c "import sysconfig; print(sysconfig.get_path('purelib'))"
+    OUTPUT_VARIABLE PYTHON_PURELIB
+    OUTPUT_STRIP_TRAILING_WHITESPACE
+)
+
+set(TORCH_PYTHON_PACKAGE_DIR "")
+foreach(_candidate
+    "${PYTHON_PLATLIB}/torch"
+    "${PYTHON_PURELIB}/torch"
+)
+    if(EXISTS "${_candidate}")
+        set(TORCH_PYTHON_PACKAGE_DIR "${_candidate}")
+        break()
+    endif()
+endforeach()
+
+if(NOT TORCH_PYTHON_PACKAGE_DIR)
+    message(FATAL_ERROR "Could not locate the installed torch Python package.")
+endif()
+
+list(APPEND TORCH_INCLUDE_DIRS
+    "${TORCH_PYTHON_PACKAGE_DIR}/include"
+    "${TORCH_PYTHON_PACKAGE_DIR}/include/torch/csrc/api/include"
+)
+
+if(NOT Torch_DIR)
+    set(_TORCH_CONFIG_DIR "${TORCH_PYTHON_PACKAGE_DIR}/share/cmake/Torch")
+    if(EXISTS "${_TORCH_CONFIG_DIR}/TorchConfig.cmake")
+        set(Torch_DIR "${_TORCH_CONFIG_DIR}" CACHE PATH "Path to Torch CMake config" FORCE)
+    endif()
+endif()
 
 # Find Torch package (still useful for libraries)
 find_package(Torch REQUIRED)
@@ -50,6 +95,21 @@ include_directories(
 set(FASTVIDEO_KERNEL_BUILD_TK "AUTO" CACHE STRING "Build ThunderKittens kernels: AUTO/ON/OFF")
 set_property(CACHE FASTVIDEO_KERNEL_BUILD_TK PROPERTY STRINGS AUTO ON OFF)
 
+set(_FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER_DEFAULT "AUTO")
+if(DEFINED FASTVIDEO_KERNEL_BUILD_MODIFIED_SAGE3 AND NOT DEFINED CACHE{FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER})
+    set(_FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER_DEFAULT "${FASTVIDEO_KERNEL_BUILD_MODIFIED_SAGE3}")
+endif()
+
+set(FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER "${_FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER_DEFAULT}" CACHE STRING
+    "Build attn_qat_infer Blackwell inference kernels: AUTO/ON/OFF")
+set_property(CACHE FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER PROPERTY STRINGS AUTO ON OFF)
+
+if(DEFINED FASTVIDEO_KERNEL_BUILD_MODIFIED_SAGE3)
+    message(DEPRECATION
+        "FASTVIDEO_KERNEL_BUILD_MODIFIED_SAGE3 is deprecated. "
+        "Use FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER instead.")
+endif()
+
 # Prefer environment variable (used by CI) if CMake var is not explicitly set.
 if(NOT DEFINED TORCH_CUDA_ARCH_LIST AND DEFINED ENV{TORCH_CUDA_ARCH_LIST})
     set(TORCH_CUDA_ARCH_LIST "$ENV{TORCH_CUDA_ARCH_LIST}")
@@ -57,6 +117,7 @@ endif()
 
 message(STATUS "TORCH_CUDA_ARCH_LIST (cmake/env): ${TORCH_CUDA_ARCH_LIST}")
 message(STATUS "FASTVIDEO_KERNEL_BUILD_TK: ${FASTVIDEO_KERNEL_BUILD_TK}")
+message(STATUS "FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER: ${FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER}")
 
 set(ENABLE_TK_KERNELS OFF)
 if(FASTVIDEO_KERNEL_BUILD_TK STREQUAL "ON")
@@ -91,6 +152,54 @@ else()
     message(STATUS "ThunderKittens kernels: DISABLED (will use Triton fallbacks at runtime)")
 endif()
 
+set(ENABLE_ATTN_QAT_INFER OFF)
+if(GPU_BACKEND STREQUAL "ROCM")
+    message(STATUS "attn_qat_infer kernels: DISABLED (ROCm build)")
+else()
+    set(_WANTS_ATTN_QAT_INFER OFF)
+    if(FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER STREQUAL "ON")
+        set(_WANTS_ATTN_QAT_INFER ON)
+    elseif(FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER STREQUAL "AUTO")
+        if(TORCH_CUDA_ARCH_LIST)
+            string(REGEX MATCH
+                "(^|[; ,])((12\\.0a)|(120a)|(sm_120a))([; ,]|$)"
+                _HAS_120A "${TORCH_CUDA_ARCH_LIST}")
+            if(_HAS_120A)
+                set(_WANTS_ATTN_QAT_INFER ON)
+            endif()
+        else()
+            execute_process(
+                COMMAND "${Python_EXECUTABLE}" -c
+                "import torch; print('1' if (torch.cuda.is_available() and torch.version.cuda and torch.cuda.get_device_capability()[0] >= 12) else '0')"
+                OUTPUT_VARIABLE _LOCAL_HAS_BLACKWELL
+                OUTPUT_STRIP_TRAILING_WHITESPACE
+                ERROR_QUIET
+            )
+            if(_LOCAL_HAS_BLACKWELL STREQUAL "1")
+                set(_WANTS_ATTN_QAT_INFER ON)
+            endif()
+        endif()
+    endif()
+
+    if(_WANTS_ATTN_QAT_INFER)
+        if(CUDAToolkit_VERSION VERSION_LESS 12.8)
+            message(WARNING
+                "attn_qat_infer kernels require CUDA Toolkit 12.8+. "
+                "Skipping because CUDAToolkit_VERSION=${CUDAToolkit_VERSION}.")
+        else()
+            set(ENABLE_ATTN_QAT_INFER ON)
+        endif()
+    endif()
+
+    if(ENABLE_ATTN_QAT_INFER)
+        message(STATUS "attn_qat_infer kernels: ENABLED")
+    else()
+        message(STATUS
+            "attn_qat_infer kernels: DISABLED "
+            "(requires CUDA 12.8+ and Blackwell sm_120a)")
+    endif()
+endif()
+
 # Always try to build the extension if CUDA is available, but conditionally add sources/flags
 set(BUILD_CXX_KERNELS ON)
 
@@ -161,12 +270,15 @@ if(BUILD_CXX_KERNELS)
 
     # Also link against libtorch_python to satisfy Python-binding symbols
     # (e.g., torch::PyWarningHandler) required by torch/extension.h.
-    execute_process(
-        COMMAND "${Python_EXECUTABLE}" -c "import torch; from pathlib import Path; p=Path(torch.__file__).parent/'lib'; m=sorted(p.glob('libtorch_python*')); print(str(m[0]) if m else '')"
-        OUTPUT_VARIABLE TORCH_PYTHON_LIBRARY_PATH
-        OUTPUT_STRIP_TRAILING_WHITESPACE
-        ERROR_QUIET
+    file(GLOB TORCH_PYTHON_LIBRARY_CANDIDATES
+        "${TORCH_PYTHON_PACKAGE_DIR}/lib/libtorch_python*"
     )
+    list(LENGTH TORCH_PYTHON_LIBRARY_CANDIDATES _TORCH_PYTHON_LIBRARY_COUNT)
+    if(_TORCH_PYTHON_LIBRARY_COUNT GREATER 0)
+        list(GET TORCH_PYTHON_LIBRARY_CANDIDATES 0 TORCH_PYTHON_LIBRARY_PATH)
+    else()
+        set(TORCH_PYTHON_LIBRARY_PATH "")
+    endif()
     if(TORCH_PYTHON_LIBRARY_PATH)
         message(STATUS "TORCH_PYTHON_LIBRARY_PATH: ${TORCH_PYTHON_LIBRARY_PATH}")
         target_link_libraries(fastvideo_kernel_ops PRIVATE "${TORCH_PYTHON_LIBRARY_PATH}")
@@ -183,3 +295,73 @@ if(BUILD_CXX_KERNELS)
     install(TARGETS fastvideo_kernel_ops LIBRARY DESTINATION fastvideo_kernel/_C)
 endif()
 
+if(ENABLE_ATTN_QAT_INFER)
+    set(ATTN_QAT_INFER_DIR ${CMAKE_SOURCE_DIR}/attn_qat_infer)
+    set(ATTN_QAT_INFER_INCLUDE_DIRS
+        ${ATTN_QAT_INFER_DIR}
+        ${CMAKE_SOURCE_DIR}/include/cutlass/include
+        ${CMAKE_SOURCE_DIR}/include/cutlass/tools/util/include
+        ${TORCH_INCLUDE_DIRS}
+    )
+    set(ATTN_QAT_INFER_CUDA_FLAGS
+        "-O3"
+        "-std=c++17"
+        "-U__CUDA_NO_HALF_OPERATORS__"
+        "-U__CUDA_NO_HALF_CONVERSIONS__"
+        "-U__CUDA_NO_BFLOAT16_OPERATORS__"
+        "-U__CUDA_NO_BFLOAT16_CONVERSIONS__"
+        "-U__CUDA_NO_BFLOAT162_OPERATORS__"
+        "-U__CUDA_NO_BFLOAT162_CONVERSIONS__"
+        "--expt-relaxed-constexpr"
+        "--expt-extended-lambda"
+        "--use_fast_math"
+        "--ptxas-options=--verbose,--warn-on-local-memory-usage"
+        "-lineinfo"
+        "-DCUTLASS_DEBUG_TRACE_LEVEL=0"
+        "-DNDEBUG"
+        "-DQBLKSIZE=128"
+        "-DKBLKSIZE=128"
+        "-DCTA256"
+        "-DDQINRMEM"
+    )
+
+    Python_add_library(fp4attn_cuda MODULE WITH_SOABI
+        attn_qat_infer/blackwell/api.cu
+    )
+    target_include_directories(fp4attn_cuda PRIVATE ${ATTN_QAT_INFER_INCLUDE_DIRS})
+    target_compile_definitions(fp4attn_cuda PRIVATE TORCH_EXTENSION_NAME=fp4attn_cuda)
+    target_compile_options(fp4attn_cuda PRIVATE
+        $<$<COMPILE_LANGUAGE:CXX>:-O3 -std=c++17>
+        $<$<COMPILE_LANGUAGE:CUDA>:${ATTN_QAT_INFER_CUDA_FLAGS}>
+    )
+    set_target_properties(fp4attn_cuda PROPERTIES
+        CUDA_ARCHITECTURES "120a"
+        CXX_STANDARD 17
+        CUDA_STANDARD 17
+    )
+    target_link_libraries(fp4attn_cuda PRIVATE ${TORCH_LIBRARIES} CUDA::cudart CUDA::cuda_driver)
+
+    Python_add_library(fp4quant_cuda MODULE WITH_SOABI
+        attn_qat_infer/quantization/fp4_quantization_4d.cu
+    )
+    target_include_directories(fp4quant_cuda PRIVATE ${ATTN_QAT_INFER_INCLUDE_DIRS})
+    target_compile_definitions(fp4quant_cuda PRIVATE TORCH_EXTENSION_NAME=fp4quant_cuda)
+    target_compile_options(fp4quant_cuda PRIVATE
+        $<$<COMPILE_LANGUAGE:CXX>:-O3 -std=c++17>
+        $<$<COMPILE_LANGUAGE:CUDA>:${ATTN_QAT_INFER_CUDA_FLAGS}>
+    )
+    set_target_properties(fp4quant_cuda PROPERTIES
+        CUDA_ARCHITECTURES "120a"
+        CXX_STANDARD 17
+        CUDA_STANDARD 17
+    )
+    target_link_libraries(fp4quant_cuda PRIVATE ${TORCH_LIBRARIES} CUDA::cudart CUDA::cuda_driver)
+
+    if(TORCH_PYTHON_LIBRARY_PATH)
+        target_link_libraries(fp4attn_cuda PRIVATE "${TORCH_PYTHON_LIBRARY_PATH}")
+        target_link_libraries(fp4quant_cuda PRIVATE "${TORCH_PYTHON_LIBRARY_PATH}")
+    endif()
+
+    install(TARGETS fp4attn_cuda LIBRARY DESTINATION .)
+    install(TARGETS fp4quant_cuda LIBRARY DESTINATION .)
+endif()

fastvideo-kernel/MANIFEST.in

@@ -2,5 +2,6 @@ include LICENSE
 include README.md
 include pyproject.toml
 recursive-include python/fastvideo_kernel *.py
+recursive-include attn_qat_infer *.py *.cu *.cuh *.cpp *.h
 recursive-include csrc *.cu *.cuh *.cpp *.h
 recursive-include include/tk *.cu *.cuh *.cpp *.h *.src

fastvideo-kernel/README.md

@@ -20,6 +20,11 @@ cd fastvideo-kernel
 ./build.sh
 ```
 
+On supported Blackwell environments, the same install also packages
+`attn_qat_infer` and builds its `fp4attn_cuda` / `fp4quant_cuda`
+extensions directly from `fastvideo-kernel/attn_qat_infer/`. This path
+requires CUDA Toolkit 12.8+ and targets `sm_120a`.
+
 ### Rocm Build
 If you are in a rocm environment without the compilation toolchaine of CUDA.
 
@@ -58,6 +63,18 @@ cd fastvideo-kernel
 python benchmarks/bench_vsa.py --batch_size 1 --num_heads 16 --head_dim 128 --q_seq_lens 49152 --topk 64
 ```
 
+### Attn QAT Attention Benchmarks
+
+The Attn QAT microbenchmarks now live alongside the kernel package:
+
+```bash
+cd fastvideo-kernel
+python benchmarks/benchmark_flashattn2.py --batch-size 1 --num-heads 16 --seq-len 4096 --head-dim 128
+python benchmarks/benchmark_sageattn3.py --batch-size 1 --num-heads 16 --seq-len 4096 --head-dim 128
+python benchmarks/benchmark_blockscaled_fp4_attn.py --batch-size 1 --num-heads 16 --seq-len 4096 --head-dim 128
+python benchmarks/benchmark_combined.py --output benchmark_attention.png
+```
+
 ### TurboDiffusion Kernels
 
 This package also includes kernels from [TurboDiffusion](https://github.com/thu-ml/TurboDiffusion), including INT8 GEMM, Quantization, RMSNorm and LayerNorm.

fastvideo-kernel/attn_qat_infer/__init__.py

@@ -0,0 +1,16 @@
+"""
+Copyright (c) 2025 by SageAttention team.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+"""
+from .api import sageattn_blackwell

fastvideo-kernel/attn_qat_infer/api.py

@@ -0,0 +1,185 @@
+# Modified from the original SageATtention3 code
+"""
+Copyright (c) 2025 by SageAttention team.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+"""
+import torch
+import triton
+import triton.language as tl
+import torch.nn.functional as F
+from typing import Tuple
+from torch.nn.functional import scaled_dot_product_attention as sdpa
+import fp4attn_cuda
+import fp4quant_cuda
+
+# Centralized block size configuration for sageattn_blackwell kernels
+# These should match the values in fastvideo/attention/backends/sageattn/blackwell/block_config.h
+BLOCK_M = 128  # Block size for M dimension (query sequence length)
+BLOCK_N = 128  # Block size for N dimension (key/value sequence length)
+
+
+@triton.jit
+def group_mean_kernel(
+    q_ptr,          
+    q_out_ptr,      
+    qm_out_ptr,     
+    B, H, L, D: tl.constexpr,    
+    stride_qb, stride_qh, stride_ql, stride_qd,  
+    stride_qmb, stride_qmh, stride_qml, stride_qmd,  
+    GROUP_SIZE: tl.constexpr
+):
+    pid_b = tl.program_id(0)
+    pid_h = tl.program_id(1)
+    pid_group = tl.program_id(2)
+    
+    group_start = pid_group * GROUP_SIZE
+    offsets = group_start + tl.arange(0, GROUP_SIZE)
+    
+    q_offsets = pid_b * stride_qb + pid_h * stride_qh + offsets[:, None] * stride_ql + tl.arange(0, D)[None, :] * stride_qd
+    q_group = tl.load(q_ptr + q_offsets)
+    
+    qm_group = tl.sum(q_group, axis=0) / GROUP_SIZE
+    
+    q_group = q_group - qm_group
+    tl.store(q_out_ptr + q_offsets, q_group)
+
+    qm_offset = pid_b * stride_qmb + pid_h * stride_qmh + pid_group * stride_qml + tl.arange(0, D) * stride_qmd
+    tl.store(qm_out_ptr + qm_offset, qm_group)
+
+
+def triton_group_mean(q: torch.Tensor):
+    B, H, L, D = q.shape
+    GROUP_SIZE = BLOCK_M
+    num_groups = L // GROUP_SIZE
+    
+    q_out = torch.empty_like(q)  # [B, H, L, D]
+    qm = torch.empty(B, H, num_groups, D, device=q.device, dtype=q.dtype) 
+    
+    grid = (B, H, num_groups)
+    
+    group_mean_kernel[grid](
+        q, q_out, qm,
+        B, H, L, D,
+        q.stride(0), q.stride(1), q.stride(2), q.stride(3),
+        qm.stride(0), qm.stride(1), qm.stride(2), qm.stride(3),
+        GROUP_SIZE=GROUP_SIZE
+    )
+    return q_out, qm
+
+
+def preprocess_qkv(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, per_block_mean: bool = True, enable_smoothing_q: bool = False, enable_smoothing_k: bool = False):
+
+    def pad_to_block_size(x):
+        L = x.size(2)
+        pad_len = (BLOCK_M - L % BLOCK_M) % BLOCK_M
+        if pad_len == 0:
+            return x.contiguous()
+        return F.pad(x, (0, 0, 0, pad_len), value=0).contiguous()
+    
+    if enable_smoothing_k:
+        k -= k.mean(dim=-2, keepdim=True)  
+    q, k, v = map(lambda x: pad_to_block_size(x), [q, k, v])
+    if per_block_mean and enable_smoothing_q:
+        q, qm = triton_group_mean(q)
+    elif enable_smoothing_q:
+        qm = q.mean(dim=-2, keepdim=True)
+        q = q - qm
+    if enable_smoothing_q:
+        delta_s = torch.matmul(qm, k.transpose(-2, -1)).to(torch.float32).contiguous()
+    else:  # used to disable q smoothing
+        B, H, L, D = q.shape
+        delta_s = torch.zeros((B, H, L // BLOCK_M, k.shape[2]), device=q.device, dtype=torch.float32)
+        
+    return q, k, v, delta_s
+
+def scale_and_quant_fp4(x: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
+    assert x.ndim == 4
+    B, H, N, D = x.shape
+    packed_fp4 = torch.empty((B, H, N, D // 2), device=x.device, dtype=torch.uint8)
+    fp8_scale = torch.empty((B, H, N, D // 16), device=x.device, dtype=torch.float8_e4m3fn)
+    fp4quant_cuda.scaled_fp4_quant(x, packed_fp4, fp8_scale, 1)
+    return packed_fp4, fp8_scale
+
+def scale_and_quant_fp4_permute(x: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
+    assert x.ndim == 4
+    B, H, N, D = x.shape
+    packed_fp4 = torch.empty((B, H, N, D // 2), device=x.device, dtype=torch.uint8)
+    fp8_scale = torch.empty((B, H, N, D // 16), device=x.device, dtype=torch.float8_e4m3fn)
+    fp4quant_cuda.scaled_fp4_quant_permute(x, packed_fp4, fp8_scale, 1)
+    return packed_fp4, fp8_scale
+
+def scale_and_quant_fp4_transpose(x: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
+    assert x.ndim == 4
+    B, H, N, D = x.shape
+    packed_fp4 = torch.empty((B, H, D, N // 2), device=x.device, dtype=torch.uint8)
+    fp8_scale = torch.empty((B, H, D, N // 16), device=x.device, dtype=torch.float8_e4m3fn)
+    fp4quant_cuda.scaled_fp4_quant_trans(x, packed_fp4, fp8_scale, 1)
+    return packed_fp4, fp8_scale
+
+def blockscaled_fp4_attn(qlist: Tuple, 
+                         klist: Tuple,
+                         vlist: Tuple,
+                         delta_s: torch.Tensor,
+                         KL: int,
+                         is_causal: bool = False, 
+                         per_block_mean: bool = True,
+                         is_bf16: bool = True,
+                         single_level_p_quant: bool = False
+                        ):
+    softmax_scale = (qlist[0].shape[-1] * 2) ** (-0.5)
+    return fp4attn_cuda.fwd(qlist[0], klist[0], vlist[0], qlist[1], klist[1], vlist[1], delta_s, KL, None, softmax_scale, is_causal, per_block_mean, is_bf16, single_level_p_quant)
+
+
+def sageattn_blackwell(q, k, v, attn_mask = None, is_causal = False, per_block_mean = True, single_level_p_quant = True, **kwargs):
+    """
+    SageAttention3 Blackwell kernel for FP4 attention.
+    
+    Args:
+        q: Query tensor [B, H, L, D]
+        k: Key tensor [B, H, L, D]
+        v: Value tensor [B, H, L, D]
+        attn_mask: Attention mask (not used)
+        is_causal: Whether to use causal masking
+        per_block_mean: Whether to use per-block mean for Q smoothing
+        single_level_p_quant: If True, use single-level quantization: s_P2, P̂_2 = φ(P̃) directly
+                              (standard per-block FP4 quantization like V, no s_P1).
+                              If False (default), use two-level quantization:
+                              s_P1 = rowmax(P̃)/(448×6), then s_P2, P̂_2 = φ(P̃/s_P1).
+        **kwargs: Additional arguments (ignored)
+    
+    Returns:
+        Output tensor [B, H, L, D]
+    """
+    if q.size(-1) >= 256:
+        print(f"Unsupported Headdim {q.size(-1)}")
+        return sdpa(q, k, v, is_causal = is_causal)
+    QL = q.size(2)
+    KL = k.size(2)
+    is_bf16 = q.dtype == torch.bfloat16
+    q, k, v, delta_s = preprocess_qkv(q, k, v, per_block_mean)
+    qlist_from_cuda = scale_and_quant_fp4(q)
+    klist_from_cuda = scale_and_quant_fp4_permute(k)
+    vlist_from_cuda = scale_and_quant_fp4_transpose(v)
+    o_fp4 = blockscaled_fp4_attn(
+    qlist_from_cuda,
+    klist_from_cuda, 
+    vlist_from_cuda,
+    delta_s,
+    KL,
+    is_causal,
+    per_block_mean,
+    is_bf16,
+    single_level_p_quant
+    )[0][:, :, :QL, :].contiguous()
+    return o_fp4

fastvideo-kernel/attn_qat_infer/blackwell/__init__.py

@@ -0,0 +1 @@
+__version__ = "3.0.0.b1"

fastvideo-kernel/attn_qat_infer/blackwell/api.cu

@@ -0,0 +1,346 @@
+// Modified from the original SageAttention3 code
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+// Include these 2 headers instead of torch/extension.h since we don't need all of the torch headers.
+#include <torch/python.h>
+#include <torch/nn/functional.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
+
+#include <cutlass/numeric_types.h>
+
+#include "params.h"
+#include "launch.h"
+#include "static_switch.h"
+#include "block_config.h"
+
+#define CHECK_DEVICE(x) TORCH_CHECK(x.is_cuda(), #x " must be on CUDA")
+#define CHECK_SHAPE(x, ...) TORCH_CHECK(x.sizes() == torch::IntArrayRef({__VA_ARGS__}), #x " must have shape (" #__VA_ARGS__ ")")
+#define CHECK_CONTIGUOUS(x) TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
+
+
+void set_params_fprop(Flash_fwd_params &params,
+                      // sizes
+                      const size_t b,
+                      const size_t seqlen_q,
+                      const size_t seqlen_k,
+                      const size_t unpadded_seqlen_k,
+                      const size_t seqlen_q_rounded,
+                      const size_t seqlen_k_rounded,
+                      const size_t h,
+                      const size_t h_k,
+                      const size_t d,
+                      const size_t d_rounded,
+                      // device pointers
+                      const at::Tensor q,
+                      const at::Tensor k,
+                      const at::Tensor v,
+                      const at::Tensor delta_s,
+                      at::Tensor out,
+                      const at::Tensor sfq,
+                      const at::Tensor sfk,
+                      const at::Tensor sfv,
+                      void *cu_seqlens_q_d,
+                      void *cu_seqlens_k_d,
+                      void *seqused_k,
+                      void *p_d,
+                      void *softmax_lse_d,
+                      float p_dropout,
+                      float softmax_scale,
+                      int window_size_left,
+                      int window_size_right,
+                      bool per_block_mean,
+                      bool is_bf16,
+                      bool single_level_p_quant=false,
+                      bool seqlenq_ngroups_swapped=false) {
+
+    // Reset the parameters
+    params = {};
+    // Set the pointers and strides.
+    params.q_ptr = q.data_ptr();
+    params.k_ptr = k.data_ptr();
+    params.v_ptr = v.data_ptr();
+    params.delta_s_ptr = delta_s.data_ptr();
+    params.sfq_ptr = sfq.data_ptr();
+    params.sfk_ptr = sfk.data_ptr();
+    params.sfv_ptr = sfv.data_ptr();
+    
+    // All stride are in elements, not bytes.
+    params.q_row_stride = q.stride(-2) * 2;
+    params.k_row_stride = k.stride(-2) * 2;
+    params.v_row_stride = v.stride(-2) * 2;;
+    params.q_head_stride = q.stride(-3) * 2;
+    params.k_head_stride = k.stride(-3) * 2;
+    params.v_head_stride = v.stride(-3) * 2; // for packed q k v
+
+    params.ds_row_stride = delta_s.stride(-2);
+    params.ds_head_stride = delta_s.stride(-3);
+    
+    params.sfq_row_stride = sfq.stride(-2);
+    params.sfk_row_stride = sfk.stride(-2);
+    params.sfv_row_stride = sfv.stride(-2);
+    params.sfq_head_stride = sfq.stride(-3);
+    params.sfk_head_stride = sfk.stride(-3);
+    params.sfv_head_stride = sfv.stride(-3);
+    params.o_ptr = out.data_ptr();
+    params.o_row_stride = out.stride(-2);
+    params.o_head_stride = out.stride(-3);
+
+    if (cu_seqlens_q_d == nullptr) {
+        params.q_batch_stride = q.stride(0) * 2;
+        params.k_batch_stride = k.stride(0) * 2;
+        params.v_batch_stride = v.stride(0) * 2;
+        params.ds_batch_stride = delta_s.stride(0);
+        params.sfq_batch_stride = sfq.stride(0);
+        params.sfk_batch_stride = sfk.stride(0);
+        params.sfv_batch_stride = sfv.stride(0);
+        params.o_batch_stride = out.stride(0);
+        if (seqlenq_ngroups_swapped) {
+             params.q_batch_stride *= seqlen_q;
+             params.o_batch_stride *= seqlen_q;
+        }
+    }
+
+    params.cu_seqlens_q = static_cast<int *>(cu_seqlens_q_d);
+    params.cu_seqlens_k = static_cast<int *>(cu_seqlens_k_d);
+    params.seqused_k = static_cast<int *>(seqused_k);
+
+    // P = softmax(QK^T)
+    params.p_ptr = p_d;
+
+    // Softmax sum
+    params.softmax_lse_ptr = softmax_lse_d;
+
+    // Set the dimensions.
+    params.b = b;
+    params.h = h;
+    params.h_k = h_k;
+    params.h_h_k_ratio = h / h_k;
+    params.seqlen_q = seqlen_q;
+    params.seqlen_k = seqlen_k;
+    params.unpadded_seqlen_k = unpadded_seqlen_k;
+    params.seqlen_q_rounded = seqlen_q_rounded;
+    params.seqlen_k_rounded = seqlen_k_rounded;
+    params.d = d;
+    params.d_rounded = d_rounded;
+
+    params.head_divmod = cutlass::FastDivmod(int(h));
+
+    // Set the different scale values.
+    params.scale_softmax = softmax_scale;
+    params.scale_softmax_log2 = softmax_scale * M_LOG2E;
+    __half scale_softmax_log2_half = __float2half(params.scale_softmax_log2);
+    __half2 scale_softmax_log2_half2 = __half2(scale_softmax_log2_half, scale_softmax_log2_half);
+    params.scale_softmax_log2_half2 = reinterpret_cast<uint32_t&>(scale_softmax_log2_half2);
+
+    // Set this to probability of keeping an element to simplify things.
+    params.p_dropout = 1.f - p_dropout;
+    // Convert p from float to int so we don't have to convert the random uint to float to compare.
+    // [Minor] We want to round down since when we do the comparison we use <= instead of <
+    // params.p_dropout_in_uint = uint32_t(std::floor(params.p_dropout * 4294967295.0));
+    // params.p_dropout_in_uint16_t = uint16_t(std::floor(params.p_dropout * 65535.0));
+    params.p_dropout_in_uint8_t = uint8_t(std::floor(params.p_dropout * 255.0));
+    params.rp_dropout = 1.f / params.p_dropout;
+    params.scale_softmax_rp_dropout = params.rp_dropout * params.scale_softmax;
+    TORCH_CHECK(p_dropout < 1.f);
+    #ifdef FLASHATTENTION_DISABLE_DROPOUT
+        TORCH_CHECK(p_dropout == 0.0f, "This flash attention build does not support dropout.");
+    #endif
+
+    // Causal is the special case where window_size_right == 0 and window_size_left < 0.
+    // Local is the more general case where window_size_right >= 0 or window_size_left >= 0.
+    params.is_causal = window_size_left < 0 && window_size_right == 0;
+    params.per_block_mean = per_block_mean;
+    if (per_block_mean) {
+        params.seqlen_s = seqlen_q;
+    } else {
+        params.seqlen_s = flash::BLOCK_M; // size of BLOCK_M
+    }
+    if (window_size_left < 0 && window_size_right >= 0) { window_size_left = seqlen_k; }
+    if (window_size_left >= 0 && window_size_right < 0) { window_size_right = seqlen_k; }
+    params.window_size_left = window_size_left;
+    params.window_size_right = window_size_right;
+
+    #ifdef FLASHATTENTION_DISABLE_LOCAL
+        TORCH_CHECK(params.is_causal || (window_size_left < 0 && window_size_right < 0),
+            "This flash attention build does not support local attention.");
+    #endif
+
+    params.is_seqlens_k_cumulative = true;
+    params.is_bf16 = is_bf16;
+    params.single_level_p_quant = single_level_p_quant;
+    #ifdef FLASHATTENTION_DISABLE_UNEVEN_K
+        TORCH_CHECK(d == d_rounded, "This flash attention build does not support headdim not being a multiple of 32.");
+    #endif
+}
+
+template<bool IsBF16>
+void run_mha_fwd_dispatch_dtype(Flash_fwd_params &params, cudaStream_t stream) {
+    using OType = std::conditional_t<IsBF16, cutlass::bfloat16_t, cutlass::half_t>;
+    if (params.d == 64) {
+        run_mha_fwd_<cutlass::nv_float4_t<cutlass::float_e2m1_t>, 64, OType>(params, stream);
+    } else if (params.d == 128) {
+        run_mha_fwd_<cutlass::nv_float4_t<cutlass::float_e2m1_t>, 128, OType>(params, stream);
+    }
+}
+
+void run_mha_fwd(Flash_fwd_params &params, cudaStream_t stream, bool force_split_kernel = false) {
+    BOOL_SWITCH(params.is_bf16, IsBF16, ([&] {
+        run_mha_fwd_dispatch_dtype<IsBF16>(params, stream);
+    }));
+}
+
+std::vector<at::Tensor>
+mha_fwd(at::Tensor &q,         // batch_size x seqlen_q x num_heads x (head_size // 2)
+        const at::Tensor &k,         // batch_size x seqlen_k x num_heads_k x (head_size // 2)
+        const at::Tensor &v,         // batch_size x seqlen_k x num_heads_k x (head_size // 2)
+        const at::Tensor &sfq,
+        const at::Tensor &sfk,
+        const at::Tensor &sfv,
+        const at::Tensor &delta_s,
+        int unpadded_k,
+        c10::optional<at::Tensor> &out_,             // batch_size x seqlen_q x num_heads x head_size
+        const float softmax_scale,
+        bool is_causal, 
+        bool per_block_mean,
+        bool is_bf16,
+        bool single_level_p_quant=false  // If true, use only per-row scale s_P2 (no per-block s_P1)
+    ) {
+
+    auto dprops = at::cuda::getCurrentDeviceProperties();
+    bool is_sm120 = dprops->major == 12 && dprops->minor == 0;
+    TORCH_CHECK(is_sm120, "only supports Blackwell GPUs or newer.");
+
+    auto q_dtype = q.dtype();
+    auto sfq_dtype = sfq.dtype();
+    TORCH_CHECK(q_dtype == torch::kUInt8, "q dtype must be uint8");
+    TORCH_CHECK(k.dtype() == q_dtype, "query and key must have the same dtype");
+    TORCH_CHECK(v.dtype() == q_dtype, "query and value must have the same dtype");
+    CHECK_DEVICE(q); CHECK_DEVICE(k); CHECK_DEVICE(v);
+
+    TORCH_CHECK(sfq_dtype == torch::kFloat8_e4m3fn, "q dtype must be uint8");
+    TORCH_CHECK(sfk.dtype() == sfq_dtype, "query and key must have the same dtype");
+    TORCH_CHECK(sfv.dtype() == sfq_dtype, "query and value must have the same dtype");
+    CHECK_DEVICE(sfq); CHECK_DEVICE(sfk); CHECK_DEVICE(sfv);
+    
+    TORCH_CHECK(q.stride(-1) == 1, "Input tensor must have contiguous last dimension");
+    TORCH_CHECK(k.stride(-1) == 1, "Input tensor must have contiguous last dimension");
+    TORCH_CHECK(v.stride(-1) == 1, "Input tensor must have contiguous last dimension");
+    TORCH_CHECK(delta_s.stride(-1) == 1, "Input tensor must have contiguous last dimension");
+
+    TORCH_CHECK(q.is_contiguous(), "Input tensor must be contiguous");
+    TORCH_CHECK(k.is_contiguous(), "Input tensor must be contiguous");
+    TORCH_CHECK(v.is_contiguous(), "Input tensor must be contiguous");
+
+    const auto sizes = q.sizes();
+    auto opts = q.options();
+    const int batch_size = sizes[0];
+    int seqlen_q = sizes[2];
+    int num_heads = sizes[1];
+    const int head_size_og = sizes[3];
+    const int unpacked_head_size = head_size_og * 2;
+    const int seqlen_k = k.size(2);
+    const int num_heads_k = k.size(1);
+    
+    TORCH_CHECK(batch_size > 0, "batch size must be postive");
+    TORCH_CHECK(unpacked_head_size <= 256, "FlashAttention forward only supports head dimension at most 256");
+    TORCH_CHECK(num_heads % num_heads_k == 0, "Number of heads in key/value must divide number of heads in query");
+    TORCH_CHECK(num_heads == num_heads_k, "We do not support MQA/GQA yet");
+
+    TORCH_CHECK(unpacked_head_size == 64 || unpacked_head_size == 128 || unpacked_head_size == 256, "Only support head size 64, 128, and 256 for now");
+
+    CHECK_SHAPE(q, batch_size, num_heads, seqlen_q, head_size_og);
+    CHECK_SHAPE(k, batch_size, num_heads_k, seqlen_k, head_size_og);
+    CHECK_SHAPE(v, batch_size, num_heads_k, unpacked_head_size, seqlen_k/2);
+    // CHECK_SHAPE(delta_s, batch_size, num_heads, seqlen_q / 128, seqlen_k);
+    // CHECK_SHAPE(sfq, batch_size, seqlen_q, num_heads, unpacked_head_size);
+    // CHECK_SHAPE(sfk, batch_size, seqlen_k, num_heads_k, unpacked_head_size);
+    // CHECK_SHAPE(sfv, batch_size, unpacked_head_size, num_heads_k, seqlen_k);
+    TORCH_CHECK(unpacked_head_size % 8 == 0, "head_size must be a multiple of 8");
+    
+    auto dtype = is_bf16 ? at::ScalarType::BFloat16 : at::ScalarType::Half;
+    at::Tensor out = torch::empty({batch_size, num_heads, seqlen_q, unpacked_head_size}, opts.dtype(dtype));
+
+    auto round_multiple = [](int x, int m) { return (x + m - 1) / m * m; };
+    // const int head_size = round_multiple(head_size_og, 8);
+    // const int head_size_rounded = round_multiple(head_size, 32);
+    const int seqlen_q_rounded = round_multiple(seqlen_q, flash::BLOCK_M);
+    const int seqlen_k_rounded = round_multiple(seqlen_k, flash::BLOCK_N);
+
+    // Otherwise the kernel will be launched from cuda:0 device
+    // Cast to char to avoid compiler warning about narrowing
+    at::cuda::CUDAGuard device_guard{(char)q.get_device()};
+
+    
+
+    auto softmax_lse = torch::empty({batch_size, num_heads, seqlen_q}, opts.dtype(at::kFloat));
+    at::Tensor p;
+
+    Flash_fwd_params params;
+    set_params_fprop(params,
+                     batch_size,
+                     seqlen_q, seqlen_k, unpadded_k,
+                     seqlen_q_rounded, seqlen_k_rounded,
+                     num_heads, num_heads_k,
+                     unpacked_head_size, unpacked_head_size,
+                     q, k, v, delta_s, out, 
+                     sfq, sfk, sfv,
+                     /*cu_seqlens_q_d=*/nullptr,
+                     /*cu_seqlens_k_d=*/nullptr,
+                     /*seqused_k=*/nullptr,
+                     nullptr,
+                     softmax_lse.data_ptr(),
+                     /*p_dropout=*/0.f,
+                     softmax_scale,
+                     /*window_size_left=*/-1,
+                     /*window_size_right=*/is_causal ? 0 : -1,
+                     per_block_mean,
+                     is_bf16,
+                     single_level_p_quant
+                    );
+    // TODO: 132 sm count?
+    auto tile_count_semaphore = is_causal ? torch::full({1}, 132, opts.dtype(torch::kInt32)) : torch::empty({1}, opts.dtype(torch::kInt32));
+    params.tile_count_semaphore = tile_count_semaphore.data_ptr<int>();
+
+    if (seqlen_k > 0) {
+        auto stream = at::cuda::getCurrentCUDAStream().stream();
+        run_mha_fwd(params, stream);
+    } else {
+        // If seqlen_k == 0, then we have an empty tensor. We need to set the output to 0.
+        out.zero_();
+        softmax_lse.fill_(std::numeric_limits<float>::infinity());
+    }
+
+    // at::Tensor out_padded = out;
+    // if (head_size_og % 8 != 0) {
+    //     out = out.index({"...", torch::indexing::Slice(torch::indexing::None, head_size_og)});
+    //     if (out_.has_value()) { out_.value().copy_(out); }
+    // }
+
+    // return {out, q_padded, k_padded, v_padded, out_padded, softmax_lse, p};
+    // cudaDeviceSynchronize();
+    // auto err = cudaGetLastError();
+    // printf("%s\n", cudaGetErrorString(err));
+    return {out, softmax_lse};
+}
+
+
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    m.doc() = "FlashAttention";
+    m.def("fwd", &mha_fwd, "Forward pass");
+}

fastvideo-kernel/attn_qat_infer/blackwell/block_config.h

@@ -0,0 +1,28 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+// Centralized block size configuration for sageattn_blackwell kernels
+// Block sizes for M and N dimensions
+namespace flash {
+    // Block size for M dimension (query sequence length)
+    static constexpr int BLOCK_M = 128;
+    
+    // Block size for N dimension (key/value sequence length)
+    static constexpr int BLOCK_N = 128;
+}
+

fastvideo-kernel/attn_qat_infer/blackwell/block_info.h

@@ -0,0 +1,60 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * This code is based on code from FlashAttention3, https://github.com/Dao-AILab/flash-attention
+ * Copyright (c) 2024, Jay Shah, Ganesh Bikshandi, Ying Zhang, Vijay Thakkar, Pradeep Ramani, Tri Dao.
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+namespace flash {
+
+////////////////////////////////////////////////////////////////////////////////////////////////////
+
+template<bool Varlen=true>
+struct BlockInfo {
+
+    template<typename Params>
+    __device__ BlockInfo(const Params &params, const int bidb)
+        : sum_s_q(!Varlen || params.cu_seqlens_q == nullptr ? -1 : params.cu_seqlens_q[bidb])
+        , sum_s_k(!Varlen || params.cu_seqlens_k == nullptr || !params.is_seqlens_k_cumulative ? -1 : params.cu_seqlens_k[bidb])
+        , actual_seqlen_q(!Varlen || params.cu_seqlens_q == nullptr ? params.seqlen_q : params.cu_seqlens_q[bidb + 1] - sum_s_q)
+        // If is_seqlens_k_cumulative, then seqlen_k is cu_seqlens_k[bidb + 1] - cu_seqlens_k[bidb].
+        // Otherwise it's cu_seqlens_k[bidb], i.e., we use cu_seqlens_k to store the sequence lengths of K.
+        , seqlen_k_cache(!Varlen || params.cu_seqlens_k == nullptr ? params.seqlen_k : (params.is_seqlens_k_cumulative ? params.cu_seqlens_k[bidb + 1] - sum_s_k : params.cu_seqlens_k[bidb]))
+        , actual_seqlen_k(params.seqused_k ? params.seqused_k[bidb] : seqlen_k_cache + (params.knew_ptr == nullptr ? 0 : params.seqlen_knew))
+        {
+        }
+
+    template <typename index_t>
+    __forceinline__ __device__ index_t q_offset(const index_t batch_stride, const index_t row_stride, const int bidb) const {
+        return sum_s_q == -1 ? bidb * batch_stride : uint32_t(sum_s_q) * row_stride;
+    }
+
+    template <typename index_t>
+    __forceinline__ __device__ index_t k_offset(const index_t batch_stride, const index_t row_stride, const int bidb) const {
+        return sum_s_k == -1 ? bidb * batch_stride : uint32_t(sum_s_k) * row_stride;
+    }
+
+    const int sum_s_q;
+    const int sum_s_k;
+    const int actual_seqlen_q;
+    // We have to have seqlen_k_cache declared before actual_seqlen_k, otherwise actual_seqlen_k is set to 0.
+    const int seqlen_k_cache;
+    const int actual_seqlen_k;
+};
+
+////////////////////////////////////////////////////////////////////////////////////////////////////
+
+}  // namespace flash

fastvideo-kernel/attn_qat_infer/blackwell/blockscaled_layout.h

@@ -0,0 +1,149 @@
+/***************************************************************************************************
+ * Copyright (c) 2023 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+    \brief Blocked Scale configs specific for SM100 BlockScaled MMA
+*/
+
+#pragma once
+
+#include "cutlass/layout/matrix.h"
+
+#include "cute/int_tuple.hpp"
+#include "cute/atom/mma_traits_sm100.hpp"
+
+namespace flash {
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+using namespace cute;
+
+template<int SFVecSize, UMMA::Major major = UMMA::Major::K>
+struct BlockScaledBasicChunk {
+
+  using Blk_MN    = _64;
+  using Blk_SF    =   _4; 
+
+  using SfAtom  = Layout< Shape< Shape<_16,_4>, Shape<Int<SFVecSize>, _4>>, 
+                               Stride<Stride<_16,_4>, Stride<           _0, _1>>>;
+};
+
+template<int SFVecSize_>
+struct BlockScaledConfig {
+  // We are creating the SFA and SFB tensors' layouts in the collective since they always have the same layout.
+  // k-major order
+  static constexpr int SFVecSize = SFVecSize_;
+  static constexpr int MMA_NSF = 4; // SFVecSize, MMA_NSF
+  using BlkScaledChunk = BlockScaledBasicChunk<SFVecSize>;
+  using Blk_MN    = _64;
+  using Blk_SF    =   _4; 
+  using mnBasicBlockShape  =  Shape<_16,_4>;
+  using mnBasicBlockStride = Stride<_16,_4>;
+  using kBasicBlockShape  = Shape<Int<SFVecSize>, Int<MMA_NSF>>; // SFVecSize, MMA_NSF
+  using kBasicBlockStride = Stride<_0, _1>;
+  using SfAtom  = Layout< Shape< mnBasicBlockShape, kBasicBlockShape>, 
+                          Stride<mnBasicBlockStride, kBasicBlockStride>>;
+
+  using LayoutSF = decltype(blocked_product(SfAtom{}, 
+                                make_layout(
+                                    make_shape(int32_t(0), int32_t(0), int32_t(0), int32_t(0)),
+                                    make_stride(int32_t(0), _1{}, int32_t(0), int32_t(0)))));
+  // A single indivisible block will hold 4 scale factors of 64 rows/columns (A/B matrix).
+  // 4 is chosen to make consecutive 32bits of data to have scale factors for only a single row (col). 32bits corresponds to the TMEM word size 
+  using Blk_Elems = decltype(Blk_MN{} * Blk_SF{});
+  using sSF_strideMN = decltype(prepend(Blk_Elems{},  mnBasicBlockStride{}));
+    
+
+  // The following function is provided for user fill dynamic problem size to the layout_SFA.
+  template < class ProblemShape>
+  CUTE_HOST_DEVICE
+  static constexpr auto
+  tile_atom_to_shape_SFQKV(ProblemShape problem_shape) {
+    auto [Seqlen, Dim, HeadNum, Batch] = problem_shape;
+    return tile_to_shape(SfAtom{}, make_shape(Seqlen, Dim, HeadNum, Batch), Step<_2,_1,_3,_4>{});
+  }
+
+  // The following function is provided for user fill dynamic problem size to the layout_SFB.
+  template <class ProblemShape>
+  CUTE_HOST_DEVICE
+  static constexpr auto
+  tile_atom_to_shape_SFVt(ProblemShape problem_shape) {
+    auto [Dim, Seqlen, HeadNum, Batch] = problem_shape;
+    return tile_to_shape(SfAtom{}, make_shape(Dim, Seqlen, HeadNum, Batch), Step<_2,_1,_3,_4>{});
+  }
+
+  template<class TiledMma, class TileShape_MNK>
+  CUTE_HOST_DEVICE
+  static constexpr auto
+  deduce_smem_layoutSFQ(TiledMma tiled_mma, TileShape_MNK tileshape_mnk) {
+    
+    using sSFQ_shapeK = decltype(prepend(make_shape(Blk_SF{}/Int<MMA_NSF>{}, size<2>(TileShape_MNK{}) / Int<SFVecSize>{} / Blk_SF{}), kBasicBlockShape{}));
+    using sSFQ_shapeM = decltype(prepend(size<0>(TileShape_MNK{}) / Blk_MN{}, mnBasicBlockShape{}));
+    using sSFQ_strideM = sSF_strideMN;
+    using sSFQ_strideK = decltype(prepend(make_stride(Int<MMA_NSF>{}, size<0>(TileShape_MNK{}) / Blk_MN{} * Blk_Elems{}), kBasicBlockStride{}));
+    using sSFQ_shape = decltype(make_shape(sSFQ_shapeM{}, sSFQ_shapeK{}));
+    using sSFQ_stride = decltype(make_stride(sSFQ_strideM{}, sSFQ_strideK{}));
+    using SmemLayoutAtomSFQ = decltype(make_layout(sSFQ_shape{},  sSFQ_stride{}));
+    return SmemLayoutAtomSFQ{};
+  }
+
+  template<class TiledMma, class TileShape_MNK>
+  CUTE_HOST_DEVICE
+  static constexpr auto
+  deduce_smem_layoutSFKV(TiledMma tiled_mma, TileShape_MNK tileshape_mnk) {
+  
+    using sSFK_shapeK = decltype(prepend(make_shape(Blk_SF{}/Int<MMA_NSF>{}, size<2>(TileShape_MNK{}) / Int<SFVecSize>{} / Blk_SF{}), kBasicBlockShape{}));
+    using sSFK_shapeN = decltype(prepend(size<1>(TileShape_MNK{}) / Blk_MN{}, mnBasicBlockShape{}));
+    using sSFK_strideN = sSF_strideMN;
+    using sSFK_strideK = decltype(prepend(make_stride(Int<MMA_NSF>{}, size<1>(TileShape_MNK{}) / Blk_MN{} * Blk_Elems{}), kBasicBlockStride{}));
+    using sSFK_shape = decltype(make_shape(sSFK_shapeN{}, sSFK_shapeK{}));
+    using sSFK_stride = decltype(make_stride(sSFK_strideN{}, sSFK_strideK{}));
+    using SmemLayoutAtomSFK = decltype(make_layout(sSFK_shape{}, sSFK_stride{}));
+    return SmemLayoutAtomSFK{};
+  }
+
+  template<class TiledMma, class TileShape_MNK>
+  CUTE_HOST_DEVICE
+  static constexpr auto
+  deduce_smem_layoutSFVt(TiledMma tiled_mma, TileShape_MNK tileshape_mnk) {
+  
+    using sSFVt_shapeK = decltype(prepend(make_shape(Blk_SF{}/Int<MMA_NSF>{}, size<2>(TileShape_MNK{}) / Int<SFVecSize>{} / Blk_SF{}), kBasicBlockShape{}));
+    using sSFVt_shapeN = decltype(prepend(size<1>(TileShape_MNK{}) / Blk_MN{}, mnBasicBlockShape{}));
+    using sSFVt_strideN = sSF_strideMN;
+    using sSFVt_strideK = decltype(prepend(make_stride(Int<MMA_NSF>{}, size<1>(TileShape_MNK{}) / Blk_MN{} * Blk_Elems{}), kBasicBlockStride{}));
+    using sSFVt_shape = decltype(make_shape(sSFVt_shapeN{}, sSFVt_shapeK{}));
+    using sSFVt_stride = decltype(make_stride(sSFVt_strideN{}, sSFVt_strideK{}));
+    using SmemLayoutAtomSFVt = decltype(make_layout(sSFVt_shape{}, sSFVt_stride{}));
+    return SmemLayoutAtomSFVt{};
+  }
+};
+
+
+} // namespace flash

fastvideo-kernel/attn_qat_infer/blackwell/cute_extension.h

@@ -0,0 +1,327 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#include "cute/arch/mma_sm120.hpp"
+#include "cute/atom/mma_traits_sm120.hpp"
+#include "cute/atom/mma_atom.hpp"
+#include "cutlass/cutlass.h"
+#include "cutlass/float8.h"
+#include "cutlass/float_subbyte.h"
+
+namespace cute::SM120::BLOCKSCALED {
+
+using cutlass::float_e2m1_t;
+using cutlass::float_ue4m3_t;
+
+// MMA.SF 16x32x64 TN E2M1 x E2M1 with SF E4M3
+struct SM120_16x32x64_TN_VS_NVFP4 {
+  using DRegisters = float[16];
+  using ARegisters = uint32_t[4];
+  using BRegisters = uint32_t[8];
+  using CRegisters = float[16];
+
+  static constexpr int SFBits = 32;
+  using RegTypeSF = cute::uint_bit_t<SFBits>;
+
+  using SFARegisters = RegTypeSF[1];
+  using SFBRegisters = RegTypeSF[1];
+
+  CUTE_HOST_DEVICE static void
+  fma(float         & d0 , float         & d1 , float         & d2 , float         & d3 ,
+      float         & d4 , float         & d5 , float         & d6 , float         & d7 ,
+      float         & d8 , float         & d9 , float         & d10, float         & d11,
+      float         & d12, float         & d13, float         & d14, float         & d15,
+      uint32_t const& a0 , uint32_t const& a1 , uint32_t const& a2 , uint32_t const& a3 ,
+      uint32_t const& b0 , uint32_t const& b1 , uint32_t const& b2 , uint32_t const& b3 ,
+      uint32_t const& b4 , uint32_t const& b5 , uint32_t const& b6 , uint32_t const& b7 ,
+      float const   & c0 , float const   & c1 , float const   & c2 , float const   & c3 ,
+      float const   & c4 , float const   & c5 , float const   & c6 , float const   & c7 ,
+      float const   & c8 , float const   & c9 , float const   & c10 , float const   & c11,
+      float const   & c12, float const   & c13, float const   & c14, float const   & c15,
+      RegTypeSF const& sfa0,
+      RegTypeSF const& sfb0)
+  {
+    static constexpr uint16_t tidA = 0;
+    static constexpr uint16_t bidA = 0;
+    static constexpr uint16_t bidB = 0;
+    static constexpr uint16_t tidB0 = 0;
+    static constexpr uint16_t tidB1 = 1;
+    static constexpr uint16_t tidB2 = 2;
+    static constexpr uint16_t tidB3 = 3;
+
+#if defined(CUTE_ARCH_MXF4NVF4_4X_UE4M3_MMA_ENABLED)
+    asm volatile(
+      "mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X.m16n8k64.row.col.f32.e2m1.e2m1.f32.ue4m3 "
+      "{%0,  %1,  %2,  %3},"
+      "{%4,  %5,  %6,  %7},"
+      "{%8,  %9},"
+      "{%10, %11, %12, %13},"
+      "{%14},"
+      "{%15, %16},"
+      "{%17},"
+      "{%18, %19};\n"
+      :  "=f"(d0),  "=f"(d1),  "=f"(d8),  "=f"(d9)
+      :   "r"(a0),   "r"(a1),   "r"(a2),   "r"(a3),
+          "r"(b0),   "r"(b1),
+          "f"(c0),   "f"(c1),   "f"(c8),   "f"(c9),
+          "r"(uint32_t(sfa0)) , "h"(bidA), "h"(tidA),
+          "r"(uint32_t(sfb0)) , "h"(bidB), "h"(tidB0));
+
+    asm volatile(
+      "mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X.m16n8k64.row.col.f32.e2m1.e2m1.f32.ue4m3 "
+      "{%0,  %1,  %2,  %3},"
+      "{%4,  %5,  %6,  %7},"
+      "{%8,  %9},"
+      "{%10, %11, %12, %13},"
+      "{%14},"
+      "{%15, %16},"
+      "{%17},"
+      "{%18, %19};\n"
+      :  "=f"(d2),  "=f"(d3),  "=f"(d10),  "=f"(d11)
+      :   "r"(a0),   "r"(a1),   "r"(a2),   "r"(a3),
+          "r"(b2),   "r"(b3),
+          "f"(c2),   "f"(c3),   "f"(c10),   "f"(c11),
+          "r"(uint32_t(sfa0)) , "h"(bidA), "h"(tidA),
+          "r"(uint32_t(sfb0)) , "h"(bidB), "h"(tidB1));
+
+    asm volatile(
+      "mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X.m16n8k64.row.col.f32.e2m1.e2m1.f32.ue4m3 "
+      "{%0,  %1,  %2,  %3},"
+      "{%4,  %5,  %6,  %7},"
+      "{%8,  %9},"
+      "{%10, %11, %12, %13},"
+      "{%14},"
+      "{%15, %16},"
+      "{%17},"
+      "{%18, %19};\n"
+      :  "=f"(d4),  "=f"(d5),  "=f"(d12),  "=f"(d13)
+      :   "r"(a0),   "r"(a1),   "r"(a2),   "r"(a3),
+          "r"(b4),   "r"(b5),
+          "f"(c4),   "f"(c5),   "f"(c12),   "f"(c13),
+          "r"(uint32_t(sfa0)) , "h"(bidA), "h"(tidA),
+          "r"(uint32_t(sfb0)) , "h"(bidB), "h"(tidB2));
+
+    asm volatile(
+      "mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X.m16n8k64.row.col.f32.e2m1.e2m1.f32.ue4m3 "
+      "{%0,  %1,  %2,  %3},"
+      "{%4,  %5,  %6,  %7},"
+      "{%8,  %9},"
+      "{%10, %11, %12, %13},"
+      "{%14},"
+      "{%15, %16},"
+      "{%17},"
+      "{%18, %19};\n"
+      :  "=f"(d6),  "=f"(d7),  "=f"(d14),  "=f"(d15)
+      :   "r"(a0),   "r"(a1),   "r"(a2),   "r"(a3),
+          "r"(b6),   "r"(b7),
+          "f"(c6),   "f"(c7),   "f"(c14),   "f"(c15),
+          "r"(uint32_t(sfa0)) , "h"(bidA), "h"(tidA),
+          "r"(uint32_t(sfb0)) , "h"(bidB), "h"(tidB3));
+#else
+    CUTE_INVALID_CONTROL_PATH("Attempting to use SM120::BLOCKSCALED::SM120_16x8x64_TN_VS without CUTE_ARCH_MXF4NVF4_4X_UE4M3_MMA_ENABLED");
+#endif
+
+  }
+};
+
+} // namespace cute::SM120::BLOCKSCALED
+
+namespace cute {
+
+// MMA NVFP4 16x32x64 TN
+template <>
+struct MMA_Traits<SM120::BLOCKSCALED::SM120_16x32x64_TN_VS_NVFP4>
+{
+  // The MMA accepts 4-bit inputs regardless of the types for A and B
+  using ValTypeA = uint4_t;
+  using ValTypeB = uint4_t;
+
+  using ValTypeD = float;
+  using ValTypeC = float;
+
+  using ValTypeSF = cutlass::float_ue4m3_t;
+  constexpr static int SFVecSize = 16;
+
+  using Shape_MNK = Shape<_16,_32,_64>;
+  using ThrID     = Layout<_32>;
+
+  // (T32,V32) -> (M16,K64)
+  using ALayout   = Layout<Shape <Shape <  _4,_8>,Shape < _8,_2,  _2>>,
+                           Stride<Stride<_128,_1>,Stride<_16,_8,_512>>>;
+  // (T32,V64) -> (N32,K64)
+  using BLayout   = Layout<Shape <Shape < _4,_8>,Shape <_8,  _2, _4>>,
+                           Stride<Stride<_256,_1>,Stride<_32,_1024, _8>>>;
+  // (T32,V64) -> (M16,K64)
+  using SFALayout = Layout<Shape <Shape <_2,_2,_8>,_64>,
+                           Stride<Stride<_8,_0,_1>,_16>>;
+  // (T32,V64) -> (N32,K64)
+  using SFBLayout = Layout<Shape <Shape <_4,_8>,_64>,
+                           Stride<Stride<_8,_1>, _32>>;
+  // (T32,V16) -> (M16,N32)
+  using CLayout = Layout<Shape <Shape < _4,_8>,Shape < Shape<_2, _4>,_2>>,
+                              Stride<Stride<_32,_1>,Stride<Stride<_16, _128>,_8>>>;
+};
+
+
+template <class SFATensor, class Atom, class TiledThr, class TiledPerm>
+CUTE_HOST_DEVICE constexpr
+auto
+thrfrg_SFA(SFATensor&& sfatensor, TiledMMA<Atom, TiledThr, TiledPerm>& mma)
+{
+  CUTE_STATIC_ASSERT_V(rank(sfatensor) >= Int<2>{});
+
+  using AtomShape_MNK  = typename Atom::Shape_MNK;
+  using AtomLayoutSFA_TV = typename Atom::Traits::SFALayout;
+
+  auto permutation_mnk = TiledPerm{};
+  auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+
+  // Reorder the tensor for the TiledAtom
+  auto t_tile = make_tile(get<0>(permutation_mnk),
+                          get<2>(permutation_mnk));
+  auto t_tensor = logical_divide(sfatensor, t_tile);                 // (PermM,PermK)
+
+  // Tile the tensor for the Atom
+  auto a_tile = make_tile(make_layout(size<0>(AtomShape_MNK{})),
+                          make_layout(size<2>(AtomShape_MNK{})));
+  auto a_tensor = zipped_divide(t_tensor, a_tile);                 // ((AtomM,AtomK),(RestM,RestK))
+
+  // Transform the Atom mode from (M,K) to (Thr,Val)
+  auto tv_tensor = a_tensor.compose(AtomLayoutSFA_TV{},_);           // ((ThrV,FrgV),(RestM,RestK))
+
+  // Tile the tensor for the Thread
+  auto thr_tile = make_tile(_,
+                            make_tile(make_layout(size<1>(thr_layout_vmnk)),
+                                      make_layout(size<3>(thr_layout_vmnk))));
+  auto thr_tensor = zipped_divide(tv_tensor, thr_tile);            // ((ThrV,(ThrM,ThrK)),(FrgV,(RestM,RestK)))
+
+  return thr_tensor;
+}
+
+template <class SFBTensor, class Atom, class TiledThr, class TiledPerm>
+CUTE_HOST_DEVICE constexpr
+auto
+thrfrg_SFB(SFBTensor&& sfbtensor, TiledMMA<Atom, TiledThr, TiledPerm>& mma)
+{
+  CUTE_STATIC_ASSERT_V(rank(sfbtensor) >= Int<2>{});
+
+  using AtomShape_MNK  = typename Atom::Shape_MNK;
+  using AtomLayoutSFB_TV = typename Atom::Traits::SFBLayout;
+
+  auto permutation_mnk = TiledPerm{};
+  auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+
+  // Reorder the tensor for the TiledAtom
+  auto t_tile = make_tile(get<1>(permutation_mnk),
+                          get<2>(permutation_mnk));
+  auto t_tensor = logical_divide(sfbtensor, t_tile);                 // (PermN,PermK)
+
+  // Tile the tensor for the Atom
+  auto a_tile = make_tile(make_layout(size<1>(AtomShape_MNK{})),
+                          make_layout(size<2>(AtomShape_MNK{})));
+  auto a_tensor = zipped_divide(t_tensor, a_tile);                 // ((AtomN,AtomK),(RestN,RestK))
+
+  // Transform the Atom mode from (M,K) to (Thr,Val)
+  auto tv_tensor = a_tensor.compose(AtomLayoutSFB_TV{},_);           // ((ThrV,FrgV),(RestN,RestK))
+
+  // Tile the tensor for the Thread
+  auto thr_tile = make_tile(_,
+                            make_tile(make_layout(size<2>(thr_layout_vmnk)),
+                                      make_layout(size<3>(thr_layout_vmnk))));
+  auto thr_tensor = zipped_divide(tv_tensor, thr_tile);            // ((ThrV,(ThrN,ThrK)),(FrgV,(RestN,RestK)))
+  return thr_tensor;
+}
+
+template <class SFATensor, class ThrMma>
+CUTE_HOST_DEVICE constexpr
+auto
+partition_SFA(SFATensor&& sfatensor, ThrMma& thread_mma) {
+  auto thr_tensor = make_tensor(static_cast<SFATensor&&>(sfatensor).data(), thrfrg_SFA(sfatensor.layout(),thread_mma));
+  auto thr_vmnk = thread_mma.thr_vmnk_;
+  auto thr_vmk = make_coord(get<0>(thr_vmnk), make_coord(get<1>(thr_vmnk), get<3>(thr_vmnk)));
+  return thr_tensor(thr_vmk, make_coord(_, repeat<rank<1,1>(thr_tensor)>(_)));
+}
+
+template <class SFATensor, class ThrMma>
+CUTE_HOST_DEVICE constexpr
+auto
+partition_fragment_SFA(SFATensor&& sfatensor, ThrMma& thread_mma) {
+  using ValTypeSF = typename ThrMma::Atom::Traits::ValTypeSF;
+  return make_fragment_like<ValTypeSF>(partition_SFA(sfatensor, thread_mma));
+}
+
+template <class SFBTensor, class ThrMma>
+CUTE_HOST_DEVICE constexpr
+auto
+partition_SFB(SFBTensor&& sfbtensor, ThrMma& thread_mma) {
+  auto thr_tensor = make_tensor(static_cast<SFBTensor&&>(sfbtensor).data(), thrfrg_SFB(sfbtensor.layout(),thread_mma));
+  auto thr_vmnk = thread_mma.thr_vmnk_;
+  auto thr_vnk = make_coord(get<0>(thr_vmnk), make_coord(get<2>(thr_vmnk), get<3>(thr_vmnk)));
+  return thr_tensor(thr_vnk, make_coord(_, repeat<rank<1,1>(thr_tensor)>(_)));
+}
+
+template <class SFBTensor, class ThrMma>
+CUTE_HOST_DEVICE constexpr
+auto
+partition_fragment_SFB(SFBTensor&& sfbtensor, ThrMma& thread_mma) {
+  using ValTypeSF = typename ThrMma::Atom::Traits::ValTypeSF;
+  return make_fragment_like<ValTypeSF>(partition_SFB(sfbtensor, thread_mma));
+}
+
+template<class TiledMma>
+CUTE_HOST_DEVICE constexpr
+auto
+get_layoutSFA_TV(TiledMma& mma)
+{
+  // (M,K) -> (M,K)
+  auto tile_shape_mnk = tile_shape(mma);
+  auto ref_A = make_layout(make_shape(size<0>(tile_shape_mnk), size<2>(tile_shape_mnk)));
+  auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+
+  // (ThrV,(ThrM,ThrK)) -> (ThrV,(ThrM,ThrN,ThrK))
+  auto atile = make_tile(_,
+                        make_tile(make_layout(make_shape (size<1>(thr_layout_vmnk), size<2>(thr_layout_vmnk)),
+                                              make_stride(               Int<1>{} ,                Int<0>{} )),
+                                  _));
+
+  // thr_idx -> (ThrV,ThrM,ThrN,ThrK)
+  auto thridx_2_thrid = right_inverse(thr_layout_vmnk);
+  // (thr_idx,val) -> (M,K)
+  return thrfrg_SFA(ref_A, mma).compose(atile, _).compose(thridx_2_thrid, _);
+}
+
+template<class TiledMma>
+CUTE_HOST_DEVICE constexpr
+auto
+get_layoutSFB_TV(TiledMma& mma)
+{
+  // (N,K) -> (N,K)
+  auto tile_shape_mnk = tile_shape(mma);
+  auto ref_B = make_layout(make_shape(size<1>(tile_shape_mnk), size<2>(tile_shape_mnk)));
+  auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+
+  // (ThrV,(ThrM,ThrK)) -> (ThrV,(ThrM,ThrN,ThrK))
+  auto btile = make_tile(_,
+                        make_tile(make_layout(make_shape (size<1>(thr_layout_vmnk), size<2>(thr_layout_vmnk)),
+                                              make_stride(               Int<0>{} ,                Int<1>{} )),
+                                  _));
+
+  // thr_idx -> (ThrV,ThrM,ThrN,ThrK)
+  auto thridx_2_thrid = right_inverse(thr_layout_vmnk);
+  // (thr_idx,val) -> (M,K)
+  return thrfrg_SFB(ref_B, mma).compose(btile, _).compose(thridx_2_thrid, _);
+}
+
+} // namespace cute

fastvideo-kernel/attn_qat_infer/blackwell/epilogue_tma_ws.h

@@ -0,0 +1,222 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <cutlass/cutlass.h>
+#include "cute/tensor.hpp"
+
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "named_barrier.h"
+#include "utils.h"
+
+namespace flash {
+
+using namespace cute;
+
+template <typename Ktraits>
+struct CollectiveEpilogueFwd{
+
+    using Element = typename Ktraits::ElementOut;
+    static constexpr int kBlockM = Ktraits::kBlockM;
+    static constexpr int kBlockN = Ktraits::kBlockN;
+    static constexpr int kHeadDim = Ktraits::kHeadDim;
+    using TileShape_MNK = Shape<Int<kBlockM>, Int<kBlockN>, Int<kHeadDim>>;
+    static constexpr int kNWarps = Ktraits::kNWarps;
+    static constexpr int kNThreads = kNWarps * cutlass::NumThreadsPerWarp;
+    static constexpr int NumMmaThreads = kNThreads - cutlass::NumThreadsPerWarpGroup;
+
+    using GmemTiledCopyOTMA = cute::SM90_TMA_STORE;
+
+    // These are for storing the output tensor without TMA (e.g., for setting output to zero)
+    static constexpr int kGmemElemsPerLoad = sizeof(cute::uint128_t) / sizeof(Element);
+    static_assert(kHeadDim % kGmemElemsPerLoad == 0, "kHeadDim must be a multiple of kGmemElemsPerLoad");
+    static constexpr int kGmemThreadsPerRow = kHeadDim / kGmemElemsPerLoad;
+    static_assert(NumMmaThreads % kGmemThreadsPerRow == 0, "NumMmaThreads must be a multiple of kGmemThreadsPerRow");
+    using GmemLayoutAtom = Layout<Shape <Int<NumMmaThreads / kGmemThreadsPerRow>, Int<kGmemThreadsPerRow>>,
+                                  Stride<Int<kGmemThreadsPerRow>, _1>>;
+    using GmemTiledCopyO = decltype(
+        make_tiled_copy(Copy_Atom<DefaultCopy, Element>{}, 
+                        GmemLayoutAtom{},
+                        Layout<Shape<_1, Int<kGmemElemsPerLoad>>>{}));  // Val layout, 8 or 16 vals per store
+
+    using SmemLayoutO = typename Ktraits::SmemLayoutO;
+
+    using SmemCopyAtomO = Copy_Atom<SM90_U32x2_STSM_N, Element>;
+    using SharedStorage = cute::array_aligned<Element, cute::cosize_v<SmemLayoutO>>;
+
+    using ShapeO = cute::Shape<int32_t, int32_t, int32_t, int32_t>;  // (seqlen_q, d, head, batch)
+    using StrideO = cute::Stride<int64_t, _1, int64_t, int64_t>;
+    using StrideLSE = cute::Stride<_1, int64_t, int64_t>;            // (seqlen_q, head, batch)
+
+    using TMA_O = decltype(make_tma_copy(
+        GmemTiledCopyOTMA{},
+        make_tensor(make_gmem_ptr(static_cast<Element*>(nullptr)), repeat_like(StrideO{}, int32_t(0)), StrideO{}),
+        SmemLayoutO{},
+        select<0, 2>(TileShape_MNK{}),
+        _1{}));  // no mcast for O
+
+    // Host side kernel arguments
+    struct Arguments {
+        Element* ptr_O;
+        ShapeO const shape_O;
+        StrideO const stride_O;
+        float* ptr_LSE;
+        StrideLSE const stride_LSE;
+    };
+
+    // Device side kernel params
+    struct Params {
+        Element* ptr_O;
+        ShapeO const shape_O;
+        StrideO const stride_O;
+        float* ptr_LSE;
+        StrideLSE const stride_LSE;
+        TMA_O tma_store_O;
+    };
+
+    static Params
+    to_underlying_arguments(Arguments const& args) {
+        Tensor mO = make_tensor(make_gmem_ptr(args.ptr_O), args.shape_O, args.stride_O);
+        TMA_O tma_store_O = make_tma_copy(
+            GmemTiledCopyOTMA{},
+            mO,
+            SmemLayoutO{},
+            select<0, 2>(TileShape_MNK{}),
+            _1{}); // no mcast for O
+        return {args.ptr_O, args.shape_O, args.stride_O, args.ptr_LSE, args.stride_LSE, tma_store_O};
+    }
+
+    /// Issue Tma Descriptor Prefetch -- ideally from a single thread for best performance
+    CUTLASS_DEVICE
+    static void prefetch_tma_descriptors(Params const& epilogue_params) {
+        cute::prefetch_tma_descriptor(epilogue_params.tma_store_O.get_tma_descriptor());
+    }
+
+    template <typename SharedStorage, typename FrgTensorO, typename TiledMma>
+    CUTLASS_DEVICE void
+    mma_store(
+        SharedStorage& shared_storage,
+        TiledMma tiled_mma,
+        FrgTensorO const& tOrO,
+        int thread_idx
+    ){
+        Tensor sO = cute::as_position_independent_swizzle_tensor(make_tensor(make_smem_ptr(shared_storage.smem_o.begin()), SmemLayoutO{}));
+        auto smem_tiled_copy_O = make_tiled_copy_C(SmemCopyAtomO{}, tiled_mma);
+        auto smem_thr_copy_O = smem_tiled_copy_O.get_thread_slice(thread_idx);
+        constexpr int numel = decltype(size(tOrO))::value;
+        cutlass::NumericArrayConverter<Element, float, numel> convert_op;
+        // HACK: this requires tensor to be "contiguous"
+        auto frag = convert_op(*reinterpret_cast<const cutlass::Array<float, numel> *>(tOrO.data()));
+        auto tOrO_out = make_tensor(make_rmem_ptr<Element>(&frag), tOrO.layout());
+        Tensor taccOrO = smem_thr_copy_O.retile_S(tOrO_out);        // ((Atom,AtomNum), MMA_M, MMA_N)
+        Tensor taccOsO = smem_thr_copy_O.partition_D(sO);     // ((Atom,AtomNum),PIPE_M,PIPE_N)
+        cute::copy(smem_tiled_copy_O, taccOrO, taccOsO);
+        cutlass::arch::fence_view_async_shared(); // ensure smem writes are visible to TMA
+    }
+
+    template<typename SharedStorage, typename Params, typename WorkTileInfo, typename SchedulerParams>
+    CUTLASS_DEVICE void
+    tma_store(
+        SharedStorage& shared_storage,
+        Params const& epilogue_params,
+        WorkTileInfo work_tile_info,
+        SchedulerParams const& scheduler_params,
+        int thread_idx
+    ) {
+        auto [m_block, bidh, bidb] = work_tile_info.get_block_coord(scheduler_params);
+        Tensor sO = cute::as_position_independent_swizzle_tensor(make_tensor(make_smem_ptr(shared_storage.smem_o.begin()), SmemLayoutO{}));
+        Tensor mO = epilogue_params.tma_store_O.get_tma_tensor(epilogue_params.shape_O);
+        Tensor gO = local_tile(mO(_, _, bidh, bidb), select<0, 2>(TileShape_MNK{}), make_coord(m_block, _0{}));  // (M, K)
+        auto block_tma_O = epilogue_params.tma_store_O.get_slice(_0{});
+        Tensor tOgO = block_tma_O.partition_D(gO);  // (TMA, TMA_M, TMA_K)
+        Tensor tOsO = block_tma_O.partition_S(sO); // (TMA, TMA_M, TMA_K)
+
+        // auto shape_LSE = select<0, 2, 3>(epilogue_params.shape_O);
+        // Tensor mLSE = make_tensor(make_gmem_ptr(epilogue_params.ptr_LSE), shape_LSE, epilogue_params.stride_LSE);
+        // Tensor gLSE = local_tile(mLSE(_, bidh, bidb), Shape<Int<kBlockM>>{}, make_coord(m_block));
+
+        // Tensor caccO = cute::make_identity_tensor(select<0, 2>(TileShape_MNK{}));
+        // auto thread_mma = tiled_mma.get_thread_slice(thread_idx);
+        // Tensor taccOcO = thread_mma.partition_C(caccO);                           // (MMA,MMA_M,MMA_K)
+        // static_assert(decltype(size<0, 0>(taccOcO))::value == 2);
+        // static_assert(decltype(size<0, 1>(taccOcO))::value == 2);
+        // // // // taccOcO has shape ((2, 2, V), MMA_M, MMA_K), we only take only the row indices.
+        // Tensor taccOcO_row = taccOcO(make_coord(_0{}, _), _, _0{});
+        // CUTE_STATIC_ASSERT_V(size(lse) == size(taccOcO_row));                     // MMA_M
+        // if (get<1>(taccOcO_row(_0{})) == 0) {
+        //     #pragma unroll
+        //     for (int mi = 0; mi < size(lse); ++mi) {
+        //         const int row = get<0>(taccOcO_row(mi));
+        //         if (row < get<0>(shape_LSE) - m_block * kBlockM) { gLSE(row) = lse(mi); }
+        //     }
+        // }
+
+        // if (cutlass::canonical_warp_idx_sync() == kNWarps - 1) {
+        //     cutlass::arch::NamedBarrier::sync(NumMmaThreads + cutlass::NumThreadsPerWarp,
+        //                         static_cast<uint32_t>(FP4NamedBarriers::EpilogueBarrier));
+        //     int const lane_predicate = cute::elect_one_sync();
+        //     if (lane_predicate) {
+        //         cute::copy(epilogue_params.tma_store_O, tOsO, tOgO);
+        //         tma_store_arrive();
+        //     }
+        // }
+        cute::copy(epilogue_params.tma_store_O, tOsO, tOgO);
+        tma_store_arrive();
+    }
+
+    CUTLASS_DEVICE void
+    store_tail() {
+        tma_store_wait<0>();
+    }
+
+    // Write 0 to output and -inf to LSE
+    CUTLASS_DEVICE void
+    store_zero(
+         Params const& epilogue_params,
+         int thread_idx,
+         cute::tuple<int32_t, int32_t, int32_t> const& block_coord
+         ) {
+        auto [m_block, bidh, bidb] = block_coord;
+        Tensor mO = make_tensor(make_gmem_ptr(epilogue_params.ptr_O), epilogue_params.shape_O, epilogue_params.stride_O);
+        Tensor gO = local_tile(mO(_, _, bidh, bidb), select<0, 2>(TileShape_MNK{}), make_coord(m_block, _0{}));  // (M, K)
+        auto shape_LSE = select<0, 2, 3>(epilogue_params.shape_O);
+        Tensor mLSE = make_tensor(make_gmem_ptr(epilogue_params.ptr_LSE), shape_LSE, epilogue_params.stride_LSE);
+        Tensor gLSE = local_tile(mLSE(_, bidh, bidb), Shape<Int<kBlockM>>{}, make_coord(m_block));
+
+        GmemTiledCopyO gmem_tiled_copy_O;
+        auto gmem_thr_copy_O = gmem_tiled_copy_O.get_thread_slice(thread_idx);
+        Tensor tOgO = gmem_thr_copy_O.partition_D(gO);
+        Tensor tOrO = make_fragment_like(tOgO);
+        clear(tOrO);
+        // Construct identity layout for sO
+        Tensor cO = cute::make_identity_tensor(select<0, 2>(TileShape_MNK{}));  // (BLK_M,BLK_K) -> (blk_m,blk_k)
+        // Repeat the partitioning with identity layouts
+        Tensor tOcO = gmem_thr_copy_O.partition_D(cO);
+        Tensor tOpO = make_tensor<bool>(make_shape(size<2>(tOgO)));
+        #pragma unroll
+        for (int k = 0; k < size(tOpO); ++k) { tOpO(k) = get<1>(tOcO(_0{}, _0{}, k)) < get<1>(epilogue_params.shape_O); }
+        // Clear_OOB_K must be false since we don't want to write zeros to gmem
+        flash::copy</*Is_even_MN=*/false, /*Is_even_K=*/false, /*Clear_OOB_MN=*/false, /*Clear_OOB_K=*/false>(
+            gmem_tiled_copy_O, tOrO, tOgO, tOcO, tOpO, get<0>(epilogue_params.shape_O) - m_block * kBlockM
+        );
+        static_assert(kBlockM <= NumMmaThreads);
+        if (thread_idx < get<0>(shape_LSE) - m_block * kBlockM) { gLSE(thread_idx) = INFINITY; }
+    }
+
+};
+
+} // namespace flash

fastvideo-kernel/attn_qat_infer/blackwell/kernel_traits.h

@@ -0,0 +1,202 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include "cute/algorithm/copy.hpp"
+#include "cute/atom/mma_atom.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cute/tensor.hpp"
+#include "cutlass/cutlass.h"
+#include "cutlass/layout/layout.h"
+#include "cutlass/numeric_types.h"
+#include "cutlass/pipeline/pipeline.hpp"
+
+#include "blockscaled_layout.h"
+#include "cute_extension.h"
+#include "named_barrier.h"
+using namespace cute;
+
+template <
+    int kStages,
+    int EpiStages,
+    typename Element,
+    typename ElementSF,
+    typename OutputType,
+    typename SmemLayoutQ,
+    typename SmemLayoutK,
+    typename SmemLayoutV,
+    typename SmemLayoutDS,
+    typename SmemLayoutO,
+    typename SmemLayoutSFQ,
+    typename SmemLayoutSFK,
+    typename SmemLayoutSFV
+>
+struct SharedStorageQKVOwithSF : cute::aligned_struct<128, _0>{
+    
+    alignas(1024) cute::ArrayEngine<Element, cute::cosize_v<SmemLayoutQ>> smem_q;
+    alignas(1024) cute::ArrayEngine<Element, cute::cosize_v<SmemLayoutK>> smem_k;
+    cute::ArrayEngine<ElementSF, cute::cosize_v<SmemLayoutSFQ>> smem_SFQ;
+    cute::ArrayEngine<ElementSF, cute::cosize_v<SmemLayoutSFK>> smem_SFK;
+    cute::ArrayEngine<ElementSF, cute::cosize_v<SmemLayoutSFV>> smem_SFV;
+    alignas(1024) cute::ArrayEngine<float, cute::cosize_v<SmemLayoutDS>> smem_ds;
+    alignas(1024) cute::ArrayEngine<Element, cute::cosize_v<SmemLayoutV>> smem_v;
+    alignas(1024) cute::ArrayEngine<OutputType, cute::cosize_v<SmemLayoutO>> smem_o;
+    
+    struct {
+        alignas(16) typename cutlass::PipelineTmaAsync<1>::SharedStorage pipeline_q;
+        alignas(16) typename cutlass::PipelineTmaAsync<kStages>::SharedStorage pipeline_k;
+        alignas(16) typename cutlass::PipelineTmaAsync<kStages>::SharedStorage pipeline_v;
+        alignas(16) typename flash::OrderedSequenceBarrierVarGroupSize<EpiStages, 2>::SharedStorage barrier_o;
+        int tile_count_semaphore;
+    };
+  };
+
+template <
+    int kHeadDim_, 
+    int kBlockM_, 
+    int kBlockN_, 
+    int kStages_,  
+    int kClusterM_, 
+    bool BlockMean_,
+    typename ElementPairType_ = cutlass::nv_float4_t<cutlass::float_e2m1_t>, 
+    typename ElementOut_ = cutlass::bfloat16_t
+>
+struct Flash_fwd_kernel_traits {
+    static constexpr int kBlockM = kBlockM_;
+    static constexpr int kBlockN = kBlockN_;
+    static constexpr int kHeadDim = kHeadDim_;
+    static constexpr bool BlockMean = BlockMean_;
+    static constexpr bool SmoothQ = true;
+    static_assert(kHeadDim % 32 == 0);
+    static_assert(kBlockM == 64 || kBlockM == 128);
+    static constexpr int kNWarps = kBlockM == 128 ? 12 : 8;
+    static constexpr int kNThreads = kNWarps * cutlass::NumThreadsPerWarp;
+    static constexpr int kClusterM = kClusterM_;
+    static constexpr int kStages = kStages_;
+    static constexpr int EpiStages = 1;
+    static constexpr int NumSFQK = kHeadDim / 16;
+    static constexpr int NumSFPV = kBlockN / 16;
+    using ElementSF = cutlass::float_ue4m3_t;
+    using Element = cutlass::float_e2m1_t;
+    using ElementAccum = float;
+    using ElementOut = ElementOut_;
+    using index_t = int64_t;
+    static constexpr auto SFVectorSize = 16;
+    using TileShape_MNK = Shape<Int<kBlockM>, Int<kBlockN>, Int<kHeadDim>>;
+    using ClusterShape_MNK = Shape<_1, _1, _1>;
+    using PermTileM = decltype(cute::min(size<0>(TileShape_MNK{}), _128{}));
+    using PermTileN = _32;
+    using PermTileK = Int<kHeadDim>;
+    
+    using ElementQMma = decltype(cutlass::gemm::collective::detail::sm1xx_kernel_input_element_to_mma_input_element<Element>());
+    using ElementKMma = decltype(cutlass::gemm::collective::detail::sm1xx_kernel_input_element_to_mma_input_element<Element>());
+
+    using AtomLayoutMNK = std::conditional_t<kBlockM == 128,
+                                            Layout<Shape<_8, _1, _1>>,
+                                            Layout<Shape<_4, _1, _1>>
+                                            >;
+    using TiledMmaQK = decltype(cute::make_tiled_mma(
+        cute::SM120::BLOCKSCALED::SM120_16x32x64_TN_VS_NVFP4{},
+        AtomLayoutMNK{},
+        Tile<PermTileM, PermTileN, PermTileK>{}
+      ));
+    
+    using TiledMmaPV = decltype(cute::make_tiled_mma(
+        cute::SM120::BLOCKSCALED::SM120_16x32x64_TN_VS_NVFP4{},
+        AtomLayoutMNK{},
+        Tile<PermTileM, _32, PermTileK>{}
+      ));
+    
+    static constexpr int MMA_NSF = size<2>(typename TiledMmaQK::AtomShape_MNK{}) / SFVectorSize;
+
+    using GmemTiledCopy = SM90_TMA_LOAD;
+    using GmemTiledCopySF = SM90_TMA_LOAD;
+
+    using SmemLayoutAtomQ = decltype(cutlass::gemm::collective::detail::sm120_rr_smem_selector<Element, decltype(size<2>(TileShape_MNK{}))>());
+    using SmemLayoutAtomK = decltype(cutlass::gemm::collective::detail::sm120_rr_smem_selector<Element, decltype(size<2>(TileShape_MNK{}))>());
+    using SmemLayoutAtomV = decltype(cutlass::gemm::collective::detail::sm120_rr_smem_selector<Element, decltype(size<2>(TileShape_MNK{}))>());
+    using SmemLayoutAtomVt = decltype(cutlass::gemm::collective::detail::sm120_rr_smem_selector<Element, decltype(size<1>(TileShape_MNK{}))>());
+    using SmemLayoutQ = decltype(tile_to_shape(SmemLayoutAtomQ{}, select<0, 2>(TileShape_MNK{})));
+    using SmemLayoutK =
+        decltype(tile_to_shape(SmemLayoutAtomK{},
+                 make_shape(shape<1>(TileShape_MNK{}), shape<2>(TileShape_MNK{}), Int<kStages>{})));
+    using SmemLayoutV =
+        decltype(tile_to_shape(SmemLayoutAtomV{},
+                 make_shape(shape<1>(TileShape_MNK{}), shape<2>(TileShape_MNK{}), Int<kStages>{})));
+    using SmemLayoutVt =
+        decltype(tile_to_shape(SmemLayoutAtomVt{},
+                 make_shape(shape<2>(TileShape_MNK{}), shape<1>(TileShape_MNK{}), Int<kStages>{})));
+    using SmemLayoutAtomDS = Layout<Shape<Int<kBlockM>, Int<kBlockN>>, Stride<_0, _1>>;
+    using SmemLayoutDS = 
+        decltype(tile_to_shape(SmemLayoutAtomDS{},
+            make_shape(shape<0>(TileShape_MNK{}), shape<1>(TileShape_MNK{}), Int<kStages>{})));
+
+    using SmemCopyAtomQ = Copy_Atom<SM75_U32x4_LDSM_N, Element>;
+    using SmemCopyAtomKV = Copy_Atom<SM75_U32x4_LDSM_N, Element>;
+    using SmemCopyAtomSF = Copy_Atom<UniversalCopy<ElementSF>, ElementSF>;
+    using SmemCopyAtomDS = Copy_Atom<UniversalCopy<float>, float>;
+
+    using BlkScaledConfig = flash::BlockScaledConfig<SFVectorSize>;
+    using LayoutSF = typename BlkScaledConfig::LayoutSF;
+    using SfAtom = typename BlkScaledConfig::SfAtom;
+    using SmemLayoutAtomSFQ = decltype(BlkScaledConfig::deduce_smem_layoutSFQ(TiledMmaQK{}, TileShape_MNK{}));
+    using SmemLayoutAtomSFK = decltype(BlkScaledConfig::deduce_smem_layoutSFKV(TiledMmaQK{}, TileShape_MNK{}));
+    using SmemLayoutAtomSFV = decltype(BlkScaledConfig::deduce_smem_layoutSFKV(TiledMmaPV{}, TileShape_MNK{}));
+    using SmemLayoutAtomSFVt = decltype(BlkScaledConfig::deduce_smem_layoutSFVt(TiledMmaPV{}, Shape<Int<kBlockM>, Int<kHeadDim>, Int<kBlockN>>{}));
+    using LayoutSFP = decltype(
+      make_layout(
+          make_shape(make_shape(_16{}, _4{}), _1{}, Int<kBlockN / 64>{}),
+          make_stride(make_stride(_0{}, _1{}), _0{}, _4{})
+      )
+    );
+    using LayoutP = decltype(
+      make_layout(
+        make_shape(make_shape(_8{}, _2{}, _2{}), _1{}, Int<kBlockN / 64>{}),
+        make_stride(make_stride(_1{}, _8{}, _16{}), _0{}, _32{})
+      )
+    );
+    using SmemLayoutSFQ = decltype(make_layout(
+        shape(SmemLayoutAtomSFQ{}),
+        stride(SmemLayoutAtomSFQ{})
+      ));
+    using SmemLayoutSFK = decltype(make_layout(
+        append(shape(SmemLayoutAtomSFK{}), Int<kStages>{}),
+        append(stride(SmemLayoutAtomSFK{}), size(filter_zeros(SmemLayoutAtomSFK{})))
+      ));
+    using SmemLayoutSFV = decltype(make_layout(
+        append(shape(SmemLayoutAtomSFV{}), Int<kStages>{}),
+        append(stride(SmemLayoutAtomSFV{}), size(filter_zeros(SmemLayoutAtomSFV{})))
+      ));
+    using SmemLayoutSFVt = decltype(make_layout(
+        append(shape(SmemLayoutAtomSFVt{}), Int<kStages>{}),
+        append(stride(SmemLayoutAtomSFVt{}), size(filter_zeros(SmemLayoutAtomSFVt{})))
+      ));
+
+    using SmemLayoutAtomO = decltype(cutlass::gemm::collective::detail::ss_smem_selector<GMMA::Major::K, ElementOut,
+        decltype(cute::get<0>(TileShape_MNK{})), decltype(cute::get<2>(TileShape_MNK{}))>());
+    using SmemLayoutO = decltype(tile_to_shape(SmemLayoutAtomO{}, select<0, 2>(TileShape_MNK{}), Step<_1, _2>{}));
+    using SharedStorage = SharedStorageQKVOwithSF<kStages, EpiStages, Element, ElementSF, ElementOut,
+        SmemLayoutQ, SmemLayoutK, SmemLayoutV, SmemLayoutDS, 
+        SmemLayoutO, SmemLayoutSFQ, SmemLayoutSFK, SmemLayoutSFVt>;
+    using MainloopPipeline = typename cutlass::PipelineTmaAsync<kStages>;
+    using PipelineState = typename cutlass::PipelineState<kStages>;
+    using MainloopPipelineQ = cutlass::PipelineTmaAsync<1>;
+    using PipelineParamsQ = typename MainloopPipelineQ::Params;
+    using PipelineStateQ = typename cutlass::PipelineState<1>;
+    using EpilogueBarrier = typename flash::OrderedSequenceBarrierVarGroupSize<EpiStages, 2>;
+};
+

fastvideo-kernel/attn_qat_infer/blackwell/kernel_ws.h

@@ -0,0 +1,204 @@
+// Modified from the original SageAttention3 code
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include "cute/tensor.hpp"
+
+#include <cutlass/cutlass.h>
+#include <cutlass/arch/reg_reconfig.h>
+#include <cutlass/array.h>
+#include <cutlass/numeric_types.h>
+#include <cutlass/numeric_conversion.h>
+#include "cutlass/pipeline/pipeline.hpp"
+
+#include "params.h"
+#include "utils.h"
+#include "tile_scheduler.h"
+#include "mainloop_tma_ws.h"
+#include "epilogue_tma_ws.h"
+#include "named_barrier.h"
+#include "softmax_fused.h"
+
+namespace flash {
+
+using namespace cute;
+
+template <typename Ktraits, bool Is_causal, typename TileScheduler>
+__global__ void __launch_bounds__(Ktraits::kNWarps * cutlass::NumThreadsPerWarp, 1)
+    compute_attn_ws(CUTE_GRID_CONSTANT Flash_fwd_params const params,
+                    CUTE_GRID_CONSTANT typename CollectiveMainloopFwd<Ktraits, Is_causal>::Params const mainloop_params,
+                    CUTE_GRID_CONSTANT typename CollectiveEpilogueFwd<Ktraits>::Params const epilogue_params,
+                    CUTE_GRID_CONSTANT typename TileScheduler::Params const scheduler_params
+                    ) {
+
+    using Element = typename Ktraits::Element;
+    using ElementAccum = typename Ktraits::ElementAccum;
+    using SoftType = ElementAccum;
+    using TileShape_MNK = typename Ktraits::TileShape_MNK;
+    using ClusterShape = typename Ktraits::ClusterShape_MNK;
+
+    static constexpr int NumMmaThreads = size(typename Ktraits::TiledMmaQK{});
+    static constexpr int NumCopyThreads = cutlass::NumThreadsPerWarpGroup;
+    static constexpr int kBlockM = Ktraits::kBlockM;
+
+    using CollectiveMainloop = CollectiveMainloopFwd<Ktraits, Is_causal>;
+    using CollectiveEpilogue = CollectiveEpilogueFwd<Ktraits>;
+
+    using MainloopPipeline = typename Ktraits::MainloopPipeline;
+    using PipelineParams = typename MainloopPipeline::Params;
+    using PipelineState = typename MainloopPipeline::PipelineState;
+    using MainloopPipelineQ = typename Ktraits::MainloopPipelineQ;
+    using PipelineParamsQ = typename Ktraits::PipelineParamsQ;
+    using PipelineStateQ = typename Ktraits::PipelineStateQ;
+    using EpilogueBarrier = typename Ktraits::EpilogueBarrier;
+
+
+    enum class WarpGroupRole {
+        Producer = 0,
+        Consumer0 = 1,
+        Consumer1 = 2
+    };
+    enum class ProducerWarpRole {
+        Mainloop = 0,
+        Epilogue = 1,
+        Warp2 = 2,
+        Warp3 = 3
+    };
+
+    extern __shared__ char shared_memory[];
+    auto &shared_storage = *reinterpret_cast<typename Ktraits::SharedStorage*>(shared_memory);
+
+    int const lane_predicate = cute::elect_one_sync();
+    int const warp_idx = cutlass::canonical_warp_idx_sync();
+    int warp_group_idx = cutlass::canonical_warp_group_idx();
+    int const warp_group_thread_idx = threadIdx.x % cutlass::NumThreadsPerWarpGroup;
+    int warp_idx_in_warp_group = warp_idx % cutlass::NumWarpsPerWarpGroup;
+    auto warp_group_role = WarpGroupRole(warp_group_idx);
+    auto producer_warp_role = ProducerWarpRole(warp_idx_in_warp_group);
+
+    // Issue Tma Descriptor Prefetch from a single thread
+    if (warp_idx == 0 && lane_predicate) {
+        CollectiveMainloop::prefetch_tma_descriptors(mainloop_params);
+        CollectiveEpilogue::prefetch_tma_descriptors(epilogue_params);
+    }
+
+    // Obtain warp index
+
+    PipelineParams pipeline_params_v;
+    pipeline_params_v.transaction_bytes = CollectiveMainloop::TmaTransactionBytesV;
+    pipeline_params_v.role = warp_group_role == WarpGroupRole::Producer
+        ? MainloopPipeline::ThreadCategory::Producer
+        : MainloopPipeline::ThreadCategory::Consumer;
+    pipeline_params_v.is_leader = warp_group_thread_idx == 0;
+    pipeline_params_v.num_consumers = NumMmaThreads;
+
+    PipelineParams pipeline_params_k;
+    pipeline_params_k.transaction_bytes = CollectiveMainloop::TmaTransactionBytesK;
+    pipeline_params_k.role = warp_group_role == WarpGroupRole::Producer
+        ? MainloopPipeline::ThreadCategory::Producer
+        : MainloopPipeline::ThreadCategory::Consumer;
+    pipeline_params_k.is_leader = warp_group_thread_idx == 0;
+    pipeline_params_k.num_consumers = NumMmaThreads;
+
+    PipelineParamsQ pipeline_params_q;
+    pipeline_params_q.transaction_bytes = CollectiveMainloop::TmaTransactionBytesQ;
+    pipeline_params_q.role = warp_group_role == WarpGroupRole::Producer
+        ? MainloopPipelineQ::ThreadCategory::Producer
+        : MainloopPipelineQ::ThreadCategory::Consumer;
+    pipeline_params_q.is_leader = warp_group_thread_idx == 0;
+    pipeline_params_q.num_consumers = NumMmaThreads;
+
+    // We're counting on pipeline_k to call cutlass::arch::fence_barrier_init();
+    MainloopPipelineQ pipeline_q(shared_storage.pipeline_q, pipeline_params_q, ClusterShape{});
+    MainloopPipeline pipeline_k(shared_storage.pipeline_k, pipeline_params_k, ClusterShape{});
+    MainloopPipeline pipeline_v(shared_storage.pipeline_v, pipeline_params_v, ClusterShape{});
+
+    uint32_t epilogue_barrier_group_size_list[2] = {cutlass::NumThreadsPerWarp, NumMmaThreads};
+    typename EpilogueBarrier::Params params_epilogue_barrier;
+    params_epilogue_barrier.group_id = (warp_group_role == WarpGroupRole::Producer);
+    params_epilogue_barrier.group_size_list = epilogue_barrier_group_size_list;
+    EpilogueBarrier barrier_o(shared_storage.barrier_o, params_epilogue_barrier);
+
+    CollectiveMainloop collective_mainloop;
+    CollectiveEpilogue collective_epilogue;
+    __syncthreads();
+
+    if (warp_group_role == WarpGroupRole::Producer) {
+        cutlass::arch::warpgroup_reg_dealloc<24>();
+        TileScheduler scheduler;
+        
+        if (producer_warp_role == ProducerWarpRole::Mainloop) {  // Load Q, K, V
+            PipelineStateQ smem_pipe_write_q = cutlass::make_producer_start_state<MainloopPipelineQ>();
+            PipelineState smem_pipe_write_k = cutlass::make_producer_start_state<MainloopPipeline>();
+            PipelineState smem_pipe_write_v = cutlass::make_producer_start_state<MainloopPipeline>();
+            
+        int work_idx = 0;
+            for (auto work_tile_info = scheduler.get_initial_work(); work_tile_info.is_valid(scheduler_params); work_tile_info = scheduler.get_next_work(scheduler_params, work_tile_info)) {
+                int tile_count_semaphore = 0;
+                collective_mainloop.load(mainloop_params, scheduler_params, 
+                                         pipeline_q, pipeline_k, pipeline_v, 
+                                         smem_pipe_write_q, smem_pipe_write_k, smem_pipe_write_v,
+                                         shared_storage, work_tile_info, work_idx, tile_count_semaphore);
+            }
+            collective_mainloop.load_tail(pipeline_q, pipeline_k, pipeline_v, 
+                                          smem_pipe_write_q, smem_pipe_write_k, smem_pipe_write_v);
+        } else if (producer_warp_role == ProducerWarpRole::Epilogue) {
+            for (auto work_tile_info = scheduler.get_initial_work(); work_tile_info.is_valid(scheduler_params); work_tile_info = scheduler.get_next_work(scheduler_params, work_tile_info)) {
+                barrier_o.wait();
+                collective_epilogue.tma_store(shared_storage, epilogue_params, work_tile_info, scheduler_params, threadIdx.x);
+                collective_epilogue.store_tail();
+                barrier_o.arrive();
+            }
+            
+        }
+    } else if (warp_group_role == WarpGroupRole::Consumer0 || warp_group_role == WarpGroupRole::Consumer1) {
+        cutlass::arch::warpgroup_reg_alloc<232>();
+        typename Ktraits::TiledMmaPV tiled_mma_pv;
+        TileScheduler scheduler{};
+        PipelineState smem_pipe_read_k, smem_pipe_read_v;
+        PipelineStateQ smem_pipe_read_q;
+
+        int work_idx = 0;
+
+        CUTLASS_PRAGMA_NO_UNROLL
+        for (auto work_tile_info = scheduler.get_initial_work(); work_tile_info.is_valid(scheduler_params); work_tile_info = scheduler.get_next_work(scheduler_params, work_tile_info)) {
+            // Attention output (GEMM-II) accumulator.
+            Tensor tOrO = partition_fragment_C(tiled_mma_pv, select<0, 2>(TileShape_MNK{}));
+            // flash::Softmax<2 * (2 * kBlockM / NumMmaThreads)> softmax;
+            // Pass single_level_p_quant flag to control P quantization mode
+            flash::SoftmaxFused<2 * (2 * kBlockM / NumMmaThreads)> softmax_fused(params.single_level_p_quant);
+            auto block_coord = work_tile_info.get_block_coord(scheduler_params);
+            auto [m_block, bidh, bidb] = block_coord;
+
+            int n_block_max = collective_mainloop.get_n_block_max(mainloop_params, m_block);
+            if (Is_causal && n_block_max <= 0) {  // We exit early and write 0 to gO and -inf to gLSE.
+                collective_epilogue.store_zero(epilogue_params, threadIdx.x - NumCopyThreads, block_coord);
+                continue;
+            }
+
+            collective_mainloop.mma(mainloop_params, pipeline_q, pipeline_k, pipeline_v, smem_pipe_read_q, smem_pipe_read_k, smem_pipe_read_v,
+                                    tOrO, softmax_fused, n_block_max, threadIdx.x - NumCopyThreads, work_idx, m_block, shared_storage);
+            barrier_o.wait();
+            collective_epilogue.mma_store(shared_storage, tiled_mma_pv, tOrO, threadIdx.x - NumCopyThreads); 
+            barrier_o.arrive();
+            ++work_idx;
+        }
+    }
+}
+
+} // namespace flash

fastvideo-kernel/attn_qat_infer/blackwell/launch.h

@@ -0,0 +1,114 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <ATen/cuda/CUDAContext.h>
+
+#include "cute/tensor.hpp"
+
+#include "cutlass/cluster_launch.hpp"
+
+#include "static_switch.h"
+#include "params.h"
+#include "tile_scheduler.h"
+#include "kernel_ws.h"
+#include "kernel_traits.h"
+#include "block_config.h"
+
+
+template<typename Kernel_traits, bool Is_causal>
+void run_flash_fwd(Flash_fwd_params &params, cudaStream_t stream) {
+    using Element = typename Kernel_traits::Element;
+    using ElementSF = typename Kernel_traits::ElementSF;
+    using ElementOut = typename Kernel_traits::ElementOut;
+    using TileShape_MNK = typename Kernel_traits::TileShape_MNK;
+    using ClusterShape = typename Kernel_traits::ClusterShape_MNK;
+    using CollectiveMainloop = flash::CollectiveMainloopFwd<Kernel_traits, Is_causal>;
+    using CollectiveEpilogue = flash::CollectiveEpilogueFwd<Kernel_traits>;
+    // using Scheduler = flash::SingleTileScheduler;
+    using Scheduler = flash::StaticPersistentTileScheduler;
+    typename CollectiveMainloop::Params mainloop_params =
+        CollectiveMainloop::to_underlying_arguments({
+            static_cast<Element const*>(params.q_ptr),
+            {params.seqlen_q, params.d, params.h, params.b},  // shape_Q
+            {params.q_row_stride, _1{}, params.q_head_stride, params.q_batch_stride},  // stride_Q
+            static_cast<Element const*>(params.k_ptr),
+            {params.seqlen_k, params.d, params.h_k, params.b},  // shape_K
+            {params.k_row_stride, _1{}, params.k_head_stride, params.k_batch_stride},  // stride_K
+            {params.unpadded_seqlen_k, params.d, params.h_k, params.b},  // shape_K
+            static_cast<Element const*>(params.v_ptr),
+            {params.d, params.seqlen_k, params.h_k, params.b},  // shape_Vt
+            {params.v_row_stride, _1{}, params.v_head_stride, params.v_batch_stride},  // stride_Vt
+            static_cast<ElementSF const*>(params.sfq_ptr),
+            {params.seqlen_q, params.d, params.h, params.b},  // shape_SFQ
+            static_cast<ElementSF const*>(params.sfk_ptr),
+            {params.seqlen_k, params.d, params.h_k, params.b},  // shape_SFK
+            static_cast<ElementSF const*>(params.sfv_ptr),
+            {params.d, params.seqlen_k, params.h_k, params.b},  // shape_SFVt
+            static_cast<float const*>(params.delta_s_ptr),
+            {params.seqlen_s, params.seqlen_k, params.h_k, params.b},
+            {params.ds_row_stride, _1{}, params.ds_head_stride, params.ds_batch_stride},
+            params.scale_softmax_log2
+        });
+    typename CollectiveEpilogue::Params epilogue_params =
+        CollectiveEpilogue::to_underlying_arguments({
+            static_cast<ElementOut*>(params.o_ptr),
+            {params.seqlen_q, params.d, params.h, params.b},  // shape_O
+            {params.o_row_stride, _1{}, params.o_head_stride, params.o_batch_stride},  // stride_O
+            static_cast<float*>(params.softmax_lse_ptr),
+            {_1{}, params.seqlen_q, params.h * params.seqlen_q},  // stride_LSE
+        });
+        
+    int num_blocks_m = cutlass::ceil_div(params.seqlen_q, Kernel_traits::kBlockM);
+    num_blocks_m = cutlass::ceil_div(num_blocks_m, size<0>(ClusterShape{})) * size<0>(ClusterShape{});
+    typename Scheduler::Arguments scheduler_args = {num_blocks_m, params.h, params.b};
+    typename Scheduler::Params scheduler_params = Scheduler::to_underlying_arguments(scheduler_args);
+    // Get the ptr to kernel function.
+    void *kernel;
+    kernel = (void *)flash::compute_attn_ws<Kernel_traits, Is_causal, Scheduler>;
+    int smem_size = sizeof(typename Kernel_traits::SharedStorage);
+    if (smem_size >= 48 * 1024) {
+       C10_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size));
+    }
+    static constexpr int ctaSize = Kernel_traits::kNWarps * 32;
+    params.m_block_divmod = cutlass::FastDivmod(num_blocks_m);
+    params.total_blocks = num_blocks_m * params.h * params.b;
+    dim3 grid_dims = Scheduler::get_grid_dim(scheduler_args, 170);
+    dim3 block_dims(ctaSize);
+    dim3 cluster_dims(size<0>(ClusterShape{}), size<1>(ClusterShape{}), size<2>(ClusterShape{}));
+    cutlass::ClusterLaunchParams launch_params{grid_dims, block_dims, cluster_dims, smem_size, stream};
+    cutlass::launch_kernel_on_cluster(launch_params, kernel, params, mainloop_params, epilogue_params, scheduler_params);
+    
+    C10_CUDA_KERNEL_LAUNCH_CHECK();
+}
+
+
+template<typename T, int Headdim, typename O = cutlass::bfloat16_t>
+void run_mha_fwd_(Flash_fwd_params &params, cudaStream_t stream) {
+    BOOL_SWITCH(params.is_causal, Is_causal, [&] {
+        BOOL_SWITCH(params.per_block_mean, per_block, [&] {
+            if constexpr (Headdim == 64 || Headdim == 128) {
+                run_flash_fwd<
+                    Flash_fwd_kernel_traits<Headdim, flash::BLOCK_M, flash::BLOCK_N, 3, 1, per_block, T, O>,
+                    Is_causal
+                >(params, stream);
+            } else {
+                static_assert(Headdim == 64 || Headdim == 128, "Unsupported Headdim");
+            }
+        });
+    });
+}

fastvideo-kernel/attn_qat_infer/blackwell/mainloop_tma_ws.h

@@ -0,0 +1,908 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <cutlass/cutlass.h>
+#include <cutlass/array.h>
+#include <cutlass/numeric_types.h>
+#include <cutlass/numeric_conversion.h>
+#include "cutlass/pipeline/pipeline.hpp"
+
+#include "cute/tensor.hpp"
+
+#include "cutlass/gemm/collective/collective_builder.hpp"
+
+#include "utils.h"
+#include "named_barrier.h"
+namespace flash {
+
+using namespace cute;
+
+template <typename Ktraits, bool Is_causal>
+struct CollectiveMainloopFwd {
+
+    using Element = typename Ktraits::Element;
+    using ElementSF = typename Ktraits::ElementSF;
+    // using TMAElement = Element;
+    // using TMAElementSF = typename Ktraits::ElementSF;
+    using TileShape_MNK = typename Ktraits::TileShape_MNK;
+    using ClusterShape = typename Ktraits::ClusterShape_MNK;
+
+    static constexpr int kStages = Ktraits::kStages;
+    static constexpr int kHeadDim = Ktraits::kHeadDim;
+    static constexpr int BlockMean = Ktraits::BlockMean;
+    using GmemTiledCopy = typename Ktraits::GmemTiledCopy;
+    using SmemLayoutQ = typename Ktraits::SmemLayoutQ;
+    using SmemLayoutK = typename Ktraits::SmemLayoutK;
+    using SmemLayoutV = typename Ktraits::SmemLayoutV;
+    using SmemLayoutVt = typename Ktraits::SmemLayoutVt;
+    using SmemLayoutDS = typename Ktraits::SmemLayoutDS;
+    using SmemLayoutAtomDS = typename Ktraits::SmemLayoutAtomDS;
+    using LayoutDS = decltype(
+        blocked_product(
+            SmemLayoutAtomDS{}, 
+            make_layout(
+            make_shape(int32_t(0), int32_t(0), int32_t(0), int32_t(0)),
+            make_stride(int32_t(0), _1{}, int32_t(0), int32_t(0)))
+        )
+        );
+    using ShapeQKV = cute::Shape<int32_t, int32_t, int32_t, int32_t>;  // (seqlen, d, head, batch)
+    using StrideQKV = cute::Stride<int64_t, _1, int64_t, int64_t>;
+    using ShapeSF = cute::Shape<int32_t, int32_t, int32_t, int32_t>;  // (seqlen, d // 16, head, batch)
+    using LayoutSF = typename Ktraits::LayoutSF;
+    using LayoutP = typename Ktraits::LayoutP;
+    using LayoutSFP = typename Ktraits::LayoutSFP;
+    using SfAtom = typename Ktraits::SfAtom;
+    using TMA_Q = decltype(make_tma_copy(
+        GmemTiledCopy{},
+        make_tensor(make_gmem_ptr(static_cast<Element const*>(nullptr)), repeat_like(StrideQKV{}, int32_t(0)), StrideQKV{}),
+        SmemLayoutQ{},
+        select<0, 2>(TileShape_MNK{}),
+        _1{}));
+
+    using TMA_KV = decltype(make_tma_copy(
+        GmemTiledCopy{},
+        make_tensor(make_gmem_ptr(static_cast<Element const*>(nullptr)), repeat_like(StrideQKV{}, int32_t(0)), StrideQKV{}),
+        take<0, 2>(SmemLayoutK{}),
+        select<1, 2>(TileShape_MNK{}),
+        _1{})); 
+    
+    using TMA_Vt = decltype(make_tma_copy(
+        GmemTiledCopy{},
+        make_tensor(make_gmem_ptr(static_cast<Element const*>(nullptr)), repeat_like(StrideQKV{}, int32_t(0)), StrideQKV{}),
+        take<0, 2>(SmemLayoutVt{}),
+        make_shape(shape<2>(TileShape_MNK{}), shape<1>(TileShape_MNK{})),
+        _1{})); 
+    
+    using TMA_DS = decltype(make_tma_copy(
+        GmemTiledCopy{},
+        make_tensor(make_gmem_ptr(static_cast<float const*>(nullptr)), LayoutDS{}),
+        take<0, 2>(SmemLayoutDS{}),
+        make_shape(shape<0>(TileShape_MNK{}), shape<1>(TileShape_MNK{})),
+        _1{}));
+    
+    using BlkScaledConfig = typename Ktraits::BlkScaledConfig;
+    using GmemTiledCopySF = typename Ktraits::GmemTiledCopySF;
+    using SmemLayoutSFQ = typename Ktraits::SmemLayoutSFQ;
+    using SmemLayoutSFK = typename Ktraits::SmemLayoutSFK;
+    using SmemLayoutSFV = typename Ktraits::SmemLayoutSFV;
+    using SmemLayoutSFVt = typename Ktraits::SmemLayoutSFVt;
+
+    using TMA_SFQ = decltype(make_tma_copy<uint16_t>(
+        GmemTiledCopySF{},
+        make_tensor(static_cast<ElementSF const*>(nullptr), LayoutSF{}),
+        SmemLayoutSFQ{},
+        make_shape(shape<0>(TileShape_MNK{}), shape<2>(TileShape_MNK{})),
+        _1{}));  // No programmatic multicast
+
+
+    using TMA_SFKV = decltype(make_tma_copy<uint16_t>(
+        GmemTiledCopySF{},
+        make_tensor(static_cast<ElementSF const*>(nullptr), LayoutSF{}),
+        SmemLayoutSFK{}(_,_,cute::Int<0>{}),
+        make_shape(shape<1>(TileShape_MNK{}), shape<2>(TileShape_MNK{})),
+        _1{}));
+
+    using TMA_SFVt = decltype(make_tma_copy<uint16_t>(
+        GmemTiledCopySF{},
+        make_tensor(static_cast<ElementSF const*>(nullptr), LayoutSF{}),
+        SmemLayoutSFVt{}(_,_,cute::Int<0>{}),
+        make_shape(shape<2>(TileShape_MNK{}), shape<1>(TileShape_MNK{})),
+        _1{}));
+
+    using SmemCopyAtomQ = typename Ktraits::SmemCopyAtomQ;
+    using SmemCopyAtomKV = typename Ktraits::SmemCopyAtomKV;
+    using SmemCopyAtomSF = typename Ktraits::SmemCopyAtomSF;
+    using TiledMmaQK = typename Ktraits::TiledMmaQK;
+    using TiledMmaPV = typename Ktraits::TiledMmaPV;
+    static constexpr int NumMmaThreads = size(TiledMmaQK{});
+    using MainloopPipeline = typename Ktraits::MainloopPipeline;
+    using PipelineParams = typename MainloopPipeline::Params;
+    using PipelineState = typename MainloopPipeline::PipelineState;
+    using MainloopPipelineQ = typename Ktraits::MainloopPipelineQ;
+    using PipelineParamsQ = typename Ktraits::PipelineParamsQ;
+    using PipelineStateQ = typename Ktraits::PipelineStateQ;
+    using EpilogueBarrier = typename Ktraits::EpilogueBarrier;
+
+    // Set the bytes transferred in this TMA transaction (may involve multiple issues)
+    static constexpr uint32_t TmaTransactionBytesQ = static_cast<uint32_t>(
+        cutlass::bits_to_bytes(cosize((SmemLayoutSFQ{})) * cute::sizeof_bits_v<ElementSF>) +
+        cutlass::bits_to_bytes(size((SmemLayoutQ{})) * sizeof_bits<Element>::value));
+    
+    static constexpr uint32_t TmaTransactionBytesK = static_cast<uint32_t>(
+        cutlass::bits_to_bytes(cosize(take<0,2>(SmemLayoutSFK{})) * cute::sizeof_bits_v<ElementSF>) +
+        cutlass::bits_to_bytes(cosize(take<0,2>(SmemLayoutDS{})) * cute::sizeof_bits_v<float>) +
+        cutlass::bits_to_bytes(size(take<0,2>(SmemLayoutK{})) * sizeof_bits<Element>::value));
+    
+    static constexpr uint32_t TmaTransactionBytesV = static_cast<uint32_t>(
+        cutlass::bits_to_bytes(cosize(take<0,2>(SmemLayoutSFVt{})) * cute::sizeof_bits_v<ElementSF>) +
+        cutlass::bits_to_bytes(size(take<0,2>(SmemLayoutVt{})) * sizeof_bits<Element>::value));
+
+    // Host side kernel arguments
+    struct Arguments {
+        Element const* ptr_Q;
+        ShapeQKV const shape_Q;
+        StrideQKV const stride_Q;
+        Element const* ptr_K;
+        ShapeQKV const shape_K;
+        StrideQKV const stride_K;
+        ShapeQKV const unpadded_shape_K;
+        Element const* ptr_Vt;
+        ShapeQKV const shape_Vt;
+        StrideQKV const stride_Vt;
+        ElementSF const* ptr_SFQ{nullptr};
+        ShapeSF const shape_SFQ{};
+        ElementSF const* ptr_SFK{nullptr};
+        ShapeSF const shape_SFK{};
+        ElementSF const* ptr_SFVt{nullptr};
+        ShapeSF const shape_SFVt{};
+        float const* ptr_ds;
+        ShapeQKV const shape_ds;
+        StrideQKV const stride_ds;
+        float const softmax_scale_log2;
+    };
+
+    // Device side kernel params
+    struct Params {
+        ShapeQKV const shape_Q;
+        LayoutSF const layout_SFQ;
+        ShapeQKV const shape_K;
+        ShapeQKV const unpadded_shape_K;
+        LayoutSF const layout_SFK;
+        ShapeQKV const shape_Vt;
+        LayoutSF const layout_SFVt;
+        LayoutDS const layout_DS;
+        TMA_Q tma_load_Q;
+        TMA_SFQ tma_load_SFQ;
+        TMA_KV tma_load_K;
+        TMA_SFKV tma_load_SFK;
+        TMA_Vt tma_load_Vt;
+        TMA_SFVt tma_load_SFVt;
+        TMA_DS tma_load_DS;
+        float const softmax_scale_log2;
+    };
+
+
+    static Params
+    to_underlying_arguments(Arguments const& args) {
+        Tensor mQ = make_tensor(make_gmem_ptr(args.ptr_Q), args.shape_Q, args.stride_Q);
+        TMA_Q tma_load_Q = make_tma_copy(
+            GmemTiledCopy{},
+            mQ,
+            SmemLayoutQ{},
+            select<0, 2>(TileShape_MNK{}),
+            _1{}); // no mcast for Q
+        Tensor mK = make_tensor(make_gmem_ptr(args.ptr_K), args.shape_K, args.stride_K);
+        TMA_KV tma_load_K = make_tma_copy(
+            GmemTiledCopy{},
+            mK,
+            SmemLayoutK{}(_, _, _0{}),
+            select<1, 2>(TileShape_MNK{}),
+            _1{}); // mcast along M mode for this N load, if any
+        Tensor mVt = make_tensor(make_gmem_ptr(args.ptr_Vt), args.shape_Vt, args.stride_Vt);
+        TMA_Vt tma_load_Vt = make_tma_copy(
+            GmemTiledCopy{},
+            mVt,
+            SmemLayoutVt{}(_, _, _0{}),
+            make_shape(shape<2>(TileShape_MNK{}), shape<1>(TileShape_MNK{})),
+            _1{}); // mcast along M mode for this N load, if any
+        auto [Seqlen_Q, Seqlen_K, HeadNum, Batch] = args.shape_ds;
+        LayoutDS layout_ds = tile_to_shape(SmemLayoutAtomDS{}, make_shape(Seqlen_Q, Seqlen_K, HeadNum, Batch), Step<_2,_1,_3,_4>{});
+        Tensor mDS = make_tensor(make_gmem_ptr(args.ptr_ds), layout_ds);
+        TMA_DS tma_load_ds = make_tma_copy (
+            GmemTiledCopy{},
+            mDS,
+            SmemLayoutDS{}(_, _, _0{}),
+            make_shape(shape<0>(TileShape_MNK{}), shape<1>(TileShape_MNK{})),
+            _1{});
+        LayoutSF layout_sfq = BlkScaledConfig::tile_atom_to_shape_SFQKV(args.shape_SFQ);
+        Tensor mSFQ = make_tensor(make_gmem_ptr(args.ptr_SFQ), layout_sfq);
+        TMA_SFQ tma_load_sfq = make_tma_copy<uint16_t>(
+            GmemTiledCopySF{},
+            mSFQ,
+            SmemLayoutSFQ{},
+            make_shape(shape<0>(TileShape_MNK{}), shape<2>(TileShape_MNK{})),
+            _1{});
+        LayoutSF layout_sfk = BlkScaledConfig::tile_atom_to_shape_SFQKV(args.shape_SFK);
+        Tensor mSFK = make_tensor(make_gmem_ptr(args.ptr_SFK), layout_sfk);
+        TMA_SFKV tma_load_sfk = make_tma_copy<uint16_t>(
+            GmemTiledCopySF{},
+            mSFK,
+            SmemLayoutSFK{}(_, _, _0{}),
+            make_shape(shape<1>(TileShape_MNK{}), shape<2>(TileShape_MNK{})),
+            _1{});
+        LayoutSF layout_sfvt = BlkScaledConfig::tile_atom_to_shape_SFVt(args.shape_SFVt);
+        Tensor mSFVt = make_tensor(make_gmem_ptr(args.ptr_SFVt), layout_sfvt);
+        TMA_SFVt tma_load_sfvt = make_tma_copy<uint16_t>(
+            GmemTiledCopySF{},
+            mSFVt,
+            SmemLayoutSFVt{}(_, _, _0{}),
+            make_shape(shape<2>(TileShape_MNK{}), shape<1>(TileShape_MNK{})),
+            _1{});
+        return {args.shape_Q, layout_sfq,
+                args.shape_K, args.unpadded_shape_K, layout_sfk,
+                args.shape_Vt, layout_sfvt,
+                layout_ds,
+                tma_load_Q, tma_load_sfq,
+                tma_load_K, tma_load_sfk,
+                tma_load_Vt, tma_load_sfvt,
+                tma_load_ds, 
+                args.softmax_scale_log2};
+    }
+
+    /// Issue Tma Descriptor Prefetch -- ideally from a single thread for best performance
+    CUTLASS_DEVICE
+    static void prefetch_tma_descriptors(Params const& mainloop_params) {
+        cute::prefetch_tma_descriptor(mainloop_params.tma_load_Q.get_tma_descriptor());
+        cute::prefetch_tma_descriptor(mainloop_params.tma_load_K.get_tma_descriptor());
+        cute::prefetch_tma_descriptor(mainloop_params.tma_load_Vt.get_tma_descriptor());
+        cute::prefetch_tma_descriptor(mainloop_params.tma_load_SFQ.get_tma_descriptor());
+        cute::prefetch_tma_descriptor(mainloop_params.tma_load_SFK.get_tma_descriptor());
+        cute::prefetch_tma_descriptor(mainloop_params.tma_load_SFVt.get_tma_descriptor());
+        cute::prefetch_tma_descriptor(mainloop_params.tma_load_DS.get_tma_descriptor());
+    }
+
+    CUTLASS_DEVICE
+    int get_n_block_max(Params const& mainloop_params, int m_block) {
+        static constexpr int kBlockM = get<0>(TileShape_MNK{});
+        static constexpr int kBlockN = get<1>(TileShape_MNK{});
+        int const seqlen_q = get<0>(mainloop_params.shape_Q);
+        int const seqlen_k = get<0>(mainloop_params.shape_K);
+        int n_block_max = cute::ceil_div(seqlen_k, kBlockN);
+        if constexpr (Is_causal) {
+            n_block_max = std::min(n_block_max,
+                                   cute::ceil_div((m_block + 1) * kBlockM + seqlen_k - seqlen_q, kBlockN));
+        }
+        return n_block_max;
+    }
+
+    template <class SFATensor, class Atom, class TiledThr, class TiledPerm>
+    CUTE_HOST_DEVICE constexpr
+    auto
+    thrfrg_SFA(SFATensor&& sfatensor, TiledMMA<Atom, TiledThr, TiledPerm>& mma)
+    {
+      CUTE_STATIC_ASSERT_V(rank(sfatensor) >= Int<2>{});
+  
+      using AtomShape_MNK  = typename Atom::Shape_MNK;
+      using AtomLayoutSFA_TV = typename Atom::Traits::SFALayout;
+  
+      auto permutation_mnk = TiledPerm{};
+      auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+  
+      // Reorder the tensor for the TiledAtom
+      auto t_tile = make_tile(get<0>(permutation_mnk),
+                              get<2>(permutation_mnk));
+      auto t_tensor = logical_divide(sfatensor, t_tile);                 // (PermM,PermK)
+  
+      // Tile the tensor for the Atom
+      auto a_tile = make_tile(make_layout(size<0>(AtomShape_MNK{})),
+                              make_layout(size<2>(AtomShape_MNK{})));
+      auto a_tensor = zipped_divide(t_tensor, a_tile);                 // ((AtomM,AtomK),(RestM,RestK))
+  
+      // Transform the Atom mode from (M,K) to (Thr,Val)
+      auto tv_tensor = a_tensor.compose(AtomLayoutSFA_TV{},_);           // ((ThrV,FrgV),(RestM,RestK))
+  
+      // Tile the tensor for the Thread
+      auto thr_tile = make_tile(_,
+                                make_tile(make_layout(size<1>(thr_layout_vmnk)),
+                                          make_layout(size<3>(thr_layout_vmnk))));
+      auto thr_tensor = zipped_divide(tv_tensor, thr_tile);            // ((ThrV,(ThrM,ThrK)),(FrgV,(RestM,RestK)))
+  
+      return thr_tensor;
+    }
+  
+    template <class SFBTensor, class Atom, class TiledThr, class TiledPerm>
+    CUTE_HOST_DEVICE constexpr
+    auto
+    thrfrg_SFB(SFBTensor&& sfbtensor, TiledMMA<Atom, TiledThr, TiledPerm>& mma)
+    {
+      CUTE_STATIC_ASSERT_V(rank(sfbtensor) >= Int<2>{});
+  
+      using AtomShape_MNK  = typename Atom::Shape_MNK;
+      using AtomLayoutSFB_TV = typename Atom::Traits::SFBLayout;
+  
+      auto permutation_mnk = TiledPerm{};
+      auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+  
+      // Reorder the tensor for the TiledAtom
+      auto t_tile = make_tile(get<1>(permutation_mnk),
+                              get<2>(permutation_mnk));
+      auto t_tensor = logical_divide(sfbtensor, t_tile);                 // (PermN,PermK)
+  
+      // Tile the tensor for the Atom
+      auto a_tile = make_tile(make_layout(size<1>(AtomShape_MNK{})),
+                              make_layout(size<2>(AtomShape_MNK{})));
+      auto a_tensor = zipped_divide(t_tensor, a_tile);                 // ((AtomN,AtomK),(RestN,RestK))
+  
+      // Transform the Atom mode from (M,K) to (Thr,Val)
+      auto tv_tensor = a_tensor.compose(AtomLayoutSFB_TV{},_);           // ((ThrV,FrgV),(RestN,RestK))
+  
+      // Tile the tensor for the Thread
+      auto thr_tile = make_tile(_,
+                                make_tile(make_layout(size<2>(thr_layout_vmnk)),
+                                          make_layout(size<3>(thr_layout_vmnk))));
+      auto thr_tensor = zipped_divide(tv_tensor, thr_tile);            // ((ThrV,(ThrN,ThrK)),(FrgV,(RestN,RestK)))
+      return thr_tensor;
+    }
+
+    template <class SFATensor, class ThrMma>
+    CUTE_HOST_DEVICE constexpr
+    auto
+    partition_fragment_SFA(SFATensor&& sfatensor, ThrMma& thread_mma)
+    {
+      using ValTypeSF = typename ThrMma::Atom::Traits::ValTypeSF;
+      auto thr_tensor = make_tensor(static_cast<SFATensor&&>(sfatensor).data(), thrfrg_SFA(sfatensor.layout(),thread_mma));
+      auto thr_vmnk = thread_mma.thr_vmnk_;
+      auto thr_vmk = make_coord(get<0>(thr_vmnk), make_coord(get<1>(thr_vmnk), get<3>(thr_vmnk)));
+      auto partition_SFA =  thr_tensor(thr_vmk, make_coord(_, repeat<rank<1,1>(thr_tensor)>(_)));
+      return make_fragment_like<ValTypeSF>(partition_SFA);
+    }
+  
+    template <class SFBTensor, class ThrMma>
+    CUTE_HOST_DEVICE constexpr
+    auto
+    partition_fragment_SFB(SFBTensor&& sfbtensor, ThrMma& thread_mma)
+    {
+      using ValTypeSF = typename ThrMma::Atom::Traits::ValTypeSF;
+      auto thr_tensor = make_tensor(static_cast<SFBTensor&&>(sfbtensor).data(), thrfrg_SFB(sfbtensor.layout(),thread_mma));
+      auto thr_vmnk = thread_mma.thr_vmnk_;
+      auto thr_vnk = make_coord(get<0>(thr_vmnk), make_coord(get<2>(thr_vmnk), get<3>(thr_vmnk)));
+      auto partition_SFB =  thr_tensor(thr_vnk, make_coord(_, repeat<rank<1,1>(thr_tensor)>(_)));
+      return make_fragment_like<ValTypeSF>(partition_SFB);
+    }
+
+    template<class TiledMma>
+    CUTE_HOST_DEVICE constexpr
+    auto
+    get_layoutSFA_TV(TiledMma& mma)
+    {
+      // (M,K) -> (M,K)
+      auto tile_shape_mnk = tile_shape(mma);
+      auto ref_A = make_layout(make_shape(size<0>(tile_shape_mnk), size<2>(tile_shape_mnk)));
+      auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+  
+      // (ThrV,(ThrM,ThrK)) -> (ThrV,(ThrM,ThrN,ThrK))
+      auto atile = make_tile(_,
+                            make_tile(make_layout(make_shape (size<1>(thr_layout_vmnk), size<2>(thr_layout_vmnk)),
+                                                  make_stride(               Int<1>{} ,                Int<0>{} )),
+                                      _));
+  
+      // thr_idx -> (ThrV,ThrM,ThrN,ThrK)
+      auto thridx_2_thrid = right_inverse(thr_layout_vmnk);
+      // (thr_idx,val) -> (M,K)
+      return thrfrg_SFA(ref_A, mma).compose(atile, _).compose(thridx_2_thrid, _);
+    }
+  
+    template<class TiledMma>
+    CUTE_HOST_DEVICE constexpr
+    auto
+    get_layoutSFB_TV(TiledMma& mma)
+    {
+      // (N,K) -> (N,K)
+      auto tile_shape_mnk = tile_shape(mma);
+      auto ref_B = make_layout(make_shape(size<1>(tile_shape_mnk), size<2>(tile_shape_mnk)));
+      auto thr_layout_vmnk = mma.get_thr_layout_vmnk();
+  
+      // (ThrV,(ThrM,ThrK)) -> (ThrV,(ThrM,ThrN,ThrK))
+      auto btile = make_tile(_,
+                            make_tile(make_layout(make_shape (size<1>(thr_layout_vmnk), size<2>(thr_layout_vmnk)),
+                                                  make_stride(               Int<0>{} ,                Int<1>{} )),
+                                      _));
+  
+      // thr_idx -> (ThrV,ThrM,ThrN,ThrK)
+      auto thridx_2_thrid = right_inverse(thr_layout_vmnk);
+      // (thr_idx,val) -> (M,K)
+      return thrfrg_SFB(ref_B, mma).compose(btile, _).compose(thridx_2_thrid, _);
+    }
+
+    template <typename SchedulerParams, typename SharedStorage, typename WorkTileInfo>
+    CUTLASS_DEVICE void
+    load(Params const& mainloop_params,
+         SchedulerParams const& scheduler_params,
+         MainloopPipelineQ pipeline_q,
+         MainloopPipeline pipeline_k,
+         MainloopPipeline pipeline_v,
+         PipelineStateQ& smem_pipe_write_q,
+         PipelineState& smem_pipe_write_k,
+         PipelineState& smem_pipe_write_v,
+         SharedStorage &shared_storage,
+         WorkTileInfo work_tile_info,
+         int& work_idx,
+         int& tile_count_semaphore
+         ) {
+
+        static constexpr int kBlockM = get<0>(TileShape_MNK{});
+        static constexpr int kBlockN = get<1>(TileShape_MNK{});
+
+        auto [m_block, bidh, bidb] = work_tile_info.get_block_coord(scheduler_params);
+
+        int n_block_max = get_n_block_max(mainloop_params, m_block);
+
+        Tensor sQ = make_tensor(make_smem_ptr(shared_storage.smem_q.begin()), SmemLayoutQ{});
+        Tensor sK = make_tensor(make_smem_ptr(shared_storage.smem_k.begin()), SmemLayoutK{});
+        Tensor sVt = make_tensor(make_smem_ptr(shared_storage.smem_v.begin()), SmemLayoutVt{});
+        Tensor sSFQ = make_tensor(make_smem_ptr(shared_storage.smem_SFQ.begin()), SmemLayoutSFQ{});
+        Tensor sSFK = make_tensor(make_smem_ptr(shared_storage.smem_SFK.begin()), SmemLayoutSFK{});
+        Tensor sSFVt = make_tensor(make_smem_ptr(shared_storage.smem_SFV.begin()), SmemLayoutSFVt{});
+        Tensor sDS = make_tensor(make_smem_ptr(shared_storage.smem_ds.begin()), SmemLayoutDS{});
+
+        Tensor mQ = mainloop_params.tma_load_Q.get_tma_tensor(mainloop_params.shape_Q);
+        Tensor mK = mainloop_params.tma_load_K.get_tma_tensor(mainloop_params.shape_K);
+        Tensor mVt = mainloop_params.tma_load_Vt.get_tma_tensor(mainloop_params.shape_Vt);
+        Tensor mDS = mainloop_params.tma_load_DS.get_tma_tensor(shape(mainloop_params.layout_DS));
+        Tensor mSFQ = mainloop_params.tma_load_SFQ.get_tma_tensor(shape(mainloop_params.layout_SFQ));
+        Tensor mSFK = mainloop_params.tma_load_SFK.get_tma_tensor(shape(mainloop_params.layout_SFK));
+        Tensor mSFVt = mainloop_params.tma_load_SFVt.get_tma_tensor(shape(mainloop_params.layout_SFVt));
+        uint32_t block_rank_in_cluster = cute::block_rank_in_cluster();
+        constexpr uint32_t cluster_shape_x = get<0>(ClusterShape());
+        uint2 cluster_local_block_id = {block_rank_in_cluster % cluster_shape_x, block_rank_in_cluster / cluster_shape_x};
+        Tensor gQ = local_tile(mQ(_, _, bidh, bidb), select<0, 2>(TileShape_MNK{}), make_coord(m_block, _0{}));  // (M, K)
+        Tensor gK = local_tile(mK(_, _, bidh, bidb), select<1, 2>(TileShape_MNK{}), make_coord(_, _0{}));  // (N, K, _)
+        Tensor gVt = local_tile(mVt(_, _, bidh, bidb), make_shape(shape<2>(TileShape_MNK{}), shape<1>(TileShape_MNK{})), make_coord(_0{}, _));  // (N, K, _)
+        Tensor gDS = [&] {
+                        if constexpr (BlockMean) {
+                            return local_tile(mDS(_, _, bidh, bidb), select<0, 1>(TileShape_MNK{}), make_coord(m_block, _));
+                        } else {
+                            return local_tile(mDS(_, _, bidh, bidb), select<0, 1>(TileShape_MNK{}), make_coord(_0{}, _));
+                        }
+                    }();
+        Tensor gSFQ = local_tile(mSFQ(_, _, bidh, bidb), select<0, 2>(TileShape_MNK{}), make_coord(m_block, _0{}));
+        Tensor gSFK = local_tile(mSFK(_, _, bidh, bidb), select<1, 2>(TileShape_MNK{}), make_coord(_, _0{}));
+        Tensor gSFVt = local_tile(mSFVt(_, _, bidh, bidb), make_shape(shape<2>(TileShape_MNK{}), shape<1>(TileShape_MNK{})), make_coord(_0{}, _));
+        auto block_tma_q = mainloop_params.tma_load_Q.get_slice(_0{});
+        Tensor tQgQ = block_tma_q.partition_S(gQ);
+        Tensor tQsQ = block_tma_q.partition_D(sQ);
+        auto block_tma_sfq = mainloop_params.tma_load_SFQ.get_slice(_0{});
+        Tensor tQgSFQ = block_tma_sfq.partition_S(gSFQ);
+        Tensor tQsSFQ = block_tma_sfq.partition_D(sSFQ);
+        auto block_tma_k = mainloop_params.tma_load_K.get_slice(cluster_local_block_id.x);
+        Tensor tKgK = group_modes<0, 3>(block_tma_k.partition_S(gK));
+        Tensor tKsK = group_modes<0, 3>(block_tma_k.partition_D(sK));
+        auto block_tma_sfk = mainloop_params.tma_load_SFK.get_slice(cluster_local_block_id.x);
+        Tensor tKgSFK = group_modes<0, 3>(block_tma_sfk.partition_S(gSFK));
+        Tensor tKsSFK = group_modes<0, 3>(block_tma_sfk.partition_D(sSFK));
+        auto block_tma_vt = mainloop_params.tma_load_Vt.get_slice(cluster_local_block_id.x);
+        Tensor tVgVt = group_modes<0, 3>(block_tma_vt.partition_S(gVt));
+        Tensor tVsVt = group_modes<0, 3>(block_tma_vt.partition_D(sVt));
+        auto block_tma_sfvt = mainloop_params.tma_load_SFVt.get_slice(cluster_local_block_id.x);
+        Tensor tVgSFVt = group_modes<0, 3>(block_tma_sfvt.partition_S(gSFVt));
+        Tensor tVsSFVt = group_modes<0, 3>(block_tma_sfvt.partition_D(sSFVt));
+        auto block_tma_ds = mainloop_params.tma_load_DS.get_slice(cluster_local_block_id.x);
+        Tensor tDSgDS = group_modes<0, 3>(block_tma_ds.partition_S(gDS));
+        Tensor tDSsDS = group_modes<0, 3>(block_tma_ds.partition_D(sDS));
+        uint16_t mcast_mask_kv = 0;
+
+        int n_block = n_block_max - 1;
+        int lane_predicate = cute::elect_one_sync();
+        if (lane_predicate) {
+        pipeline_q.producer_acquire(smem_pipe_write_q);
+        copy(mainloop_params.tma_load_Q.with(*pipeline_q.producer_get_barrier(smem_pipe_write_q), 0), tQgQ, tQsQ);
+        copy(mainloop_params.tma_load_SFQ.with(*pipeline_q.producer_get_barrier(smem_pipe_write_q), 0), tQgSFQ, tQsSFQ);
+        ++smem_pipe_write_q;
+        pipeline_k.producer_acquire(smem_pipe_write_k);
+        copy(mainloop_params.tma_load_K.with(*pipeline_k.producer_get_barrier(smem_pipe_write_k), mcast_mask_kv),
+            tKgK(_, n_block), tKsK(_, smem_pipe_write_k.index()));
+        copy(mainloop_params.tma_load_SFK.with(*pipeline_k.producer_get_barrier(smem_pipe_write_k), mcast_mask_kv),
+            tKgSFK(_, n_block), tKsSFK(_, smem_pipe_write_k.index()));
+        copy(mainloop_params.tma_load_DS.with(*pipeline_k.producer_get_barrier(smem_pipe_write_k), mcast_mask_kv),
+            tDSgDS(_, n_block), tDSsDS(_, smem_pipe_write_k.index()));
+        ++smem_pipe_write_k;
+        pipeline_v.producer_acquire(smem_pipe_write_v);
+        copy(mainloop_params.tma_load_Vt.with(*pipeline_v.producer_get_barrier(smem_pipe_write_v), mcast_mask_kv),
+            tVgVt(_, n_block), tVsVt(_, smem_pipe_write_v.index()));
+        copy(mainloop_params.tma_load_SFVt.with(*pipeline_v.producer_get_barrier(smem_pipe_write_v), mcast_mask_kv),
+            tVgSFVt(_, n_block), tVsSFVt(_, smem_pipe_write_v.index()));
+        ++smem_pipe_write_v;
+        }
+
+        n_block--;
+        if (lane_predicate) {
+            // CUTLASS_PRAGMA_NO_UNROLL
+            #pragma unroll 2
+            for (; n_block >= 0; --n_block) {
+                pipeline_k.producer_acquire(smem_pipe_write_k);
+                copy(mainloop_params.tma_load_K.with(*pipeline_k.producer_get_barrier(smem_pipe_write_k), mcast_mask_kv),
+                    tKgK(_, n_block), tKsK(_, smem_pipe_write_k.index()));
+                copy(mainloop_params.tma_load_SFK.with(*pipeline_k.producer_get_barrier(smem_pipe_write_k), mcast_mask_kv),
+                    tKgSFK(_, n_block), tKsSFK(_, smem_pipe_write_k.index()));
+                copy(mainloop_params.tma_load_DS.with(*pipeline_k.producer_get_barrier(smem_pipe_write_k), mcast_mask_kv),
+                    tDSgDS(_, n_block), tDSsDS(_, smem_pipe_write_k.index()));
+                ++smem_pipe_write_k;
+                pipeline_v.producer_acquire(smem_pipe_write_v);
+                copy(mainloop_params.tma_load_Vt.with(*pipeline_v.producer_get_barrier(smem_pipe_write_v), mcast_mask_kv),
+                    tVgVt(_, n_block), tVsVt(_, smem_pipe_write_v.index()));
+                copy(mainloop_params.tma_load_SFVt.with(*pipeline_v.producer_get_barrier(smem_pipe_write_v), mcast_mask_kv),
+                    tVgSFVt(_, n_block), tVsSFVt(_, smem_pipe_write_v.index()));
+                ++smem_pipe_write_v;
+            }
+        }
+        ++work_idx;
+    }
+
+    /// Perform a Producer Epilogue to prevent early exit of blocks in a Cluster
+    CUTLASS_DEVICE void
+    load_tail(MainloopPipelineQ pipeline_q,
+              MainloopPipeline pipeline_k, 
+              MainloopPipeline pipeline_v,
+              PipelineStateQ& smem_pipe_write_q,
+              PipelineState& smem_pipe_write_k, 
+              PipelineState& smem_pipe_write_v) {
+        int lane_predicate = cute::elect_one_sync();
+        // Issue the epilogue waits
+        if (lane_predicate) {
+          pipeline_q.producer_tail(smem_pipe_write_q);
+          pipeline_k.producer_tail(smem_pipe_write_k);
+          pipeline_v.producer_tail(smem_pipe_write_v);
+        }
+    }
+
+    template <typename SharedStorage, typename FrgTensorO, typename SoftmaxFused>
+    CUTLASS_DEVICE void
+    mma(Params const& mainloop_params,
+        MainloopPipelineQ pipeline_q,
+        MainloopPipeline pipeline_k,
+        MainloopPipeline pipeline_v,
+        PipelineStateQ& smem_pipe_read_q,
+        PipelineState& smem_pipe_read_k,
+        PipelineState& smem_pipe_read_v,
+        FrgTensorO& tOrO_store,
+        SoftmaxFused& softmax_fused,
+        int n_block_count,
+        int thread_idx,
+        int work_idx,
+        int m_block,
+        SharedStorage& shared_storage
+        ) {
+
+        static_assert(is_rmem<FrgTensorO>::value, "O tensor must be rmem resident.");
+
+        static constexpr int kBlockM = get<0>(TileShape_MNK{});
+        static constexpr int kBlockN = get<1>(TileShape_MNK{});
+        static constexpr int kBlockK = get<2>(TileShape_MNK{});
+        Tensor sQ = make_tensor(make_smem_ptr(shared_storage.smem_q.begin()), SmemLayoutQ{});
+        Tensor sK = make_tensor(make_smem_ptr(shared_storage.smem_k.begin()), SmemLayoutK{});
+        Tensor sVt = make_tensor(make_smem_ptr(shared_storage.smem_v.begin()), SmemLayoutVt{});
+        Tensor sDS = make_tensor(make_smem_ptr(shared_storage.smem_ds.begin()), SmemLayoutDS{});
+        Tensor sSFQ = make_tensor(make_smem_ptr(shared_storage.smem_SFQ.begin()), SmemLayoutSFQ{});
+        Tensor sSFK = make_tensor(make_smem_ptr(shared_storage.smem_SFK.begin()), SmemLayoutSFK{});
+        Tensor sSFVt = make_tensor(make_smem_ptr(shared_storage.smem_SFV.begin()), SmemLayoutSFVt{});
+
+        Tensor cQ = make_identity_tensor(make_shape(size<0>(sQ), size<1>(sQ)));
+        Tensor cKV = make_identity_tensor(make_shape(size<0>(sK), size<1>(sK)));
+        TiledMmaQK tiled_mma_qk;
+        TiledMmaPV tiled_mma_pv;
+        auto thread_mma_qk = tiled_mma_qk.get_thread_slice(thread_idx);
+        auto thread_mma_pv = tiled_mma_pv.get_thread_slice(thread_idx);
+
+        Tensor tSrQ = thread_mma_qk.partition_fragment_A(sQ);
+        Tensor tSrK = thread_mma_qk.partition_fragment_B(sK(_,_,Int<0>{}));
+        Tensor tOrVt = thread_mma_pv.partition_fragment_B(sVt(_,_,Int<0>{}));
+        Tensor tOrP = make_tensor_like<Element>(LayoutP{});
+        Tensor tSrSFQ = partition_fragment_SFA(sSFQ, thread_mma_qk);
+        Tensor tSrSFK = partition_fragment_SFB(sSFK(_,_,Int<0>{}), thread_mma_qk);
+        Tensor tOrSFVt = partition_fragment_SFB(sSFVt(_,_,Int<0>{}), thread_mma_pv);
+        Tensor tOrSFP = make_tensor<ElementSF>(LayoutSFP{});
+        Tensor tOrSFP_flt = filter_zeros(tOrSFP);
+        Tensor tSrDS = make_tensor<float>(make_shape(_8{}, _4{}), make_stride(_1{}, _8{}));
+        // copy qk and sf from smem to rmem
+        auto smem_tiled_copy_Q = make_tiled_copy_A(SmemCopyAtomQ{}, tiled_mma_qk);
+        auto smem_thr_copy_Q = smem_tiled_copy_Q.get_thread_slice(thread_idx);
+        Tensor tSsQ = smem_thr_copy_Q.partition_S(as_position_independent_swizzle_tensor(sQ));
+        Tensor tSrQ_copy_view = smem_thr_copy_Q.retile_D(tSrQ);
+
+        auto smem_tiled_copy_K = make_tiled_copy_B(SmemCopyAtomKV{}, tiled_mma_qk);
+        auto smem_thr_copy_K = smem_tiled_copy_K.get_thread_slice(thread_idx);
+        Tensor tSsK = smem_thr_copy_K.partition_S(as_position_independent_swizzle_tensor(sK));
+        Tensor tSrK_copy_view = smem_thr_copy_K.retile_D(tSrK);
+
+        auto smem_tiled_copy_V = make_tiled_copy_B(SmemCopyAtomKV{}, tiled_mma_pv);
+        auto smem_thr_copy_V = smem_tiled_copy_V.get_thread_slice(thread_idx);
+        Tensor tOsVt = smem_thr_copy_V.partition_S(as_position_independent_swizzle_tensor(sVt));
+        Tensor tOrVt_copy_view = smem_thr_copy_V.retile_D(tOrVt); 
+
+        auto tile_shape_mnk = tile_shape(tiled_mma_qk);
+        auto smem_tiled_copy_SFQ = make_tiled_copy_impl(SmemCopyAtomSF{}, 
+                                                        get_layoutSFA_TV(tiled_mma_qk),
+                                                        make_shape(size<0>(tile_shape_mnk), size<2>(tile_shape_mnk))
+                                                        );
+        auto smem_thr_copy_SFQ = smem_tiled_copy_SFQ.get_thread_slice(thread_idx);
+        Tensor tSsSFQ = smem_thr_copy_SFQ.partition_S(as_position_independent_swizzle_tensor(sSFQ));
+        Tensor tSrSFQ_copy_view = smem_thr_copy_SFQ.retile_D(tSrSFQ);
+
+        auto smem_tiled_copy_SFK = make_tiled_copy_impl(SmemCopyAtomSF{}, 
+                                                        get_layoutSFB_TV(tiled_mma_qk),
+                                                        make_shape(size<1>(tile_shape_mnk), size<2>(tile_shape_mnk))
+                                                        );
+        auto smem_thr_copy_SFK = smem_tiled_copy_SFK.get_thread_slice(thread_idx);
+        Tensor tSsSFK = smem_thr_copy_SFK.partition_S(as_position_independent_swizzle_tensor(sSFK));
+        Tensor tSrSFK_copy_view = smem_thr_copy_SFK.retile_D(tSrSFK);
+
+        auto smem_tiled_copy_SFV = make_tiled_copy_impl(SmemCopyAtomSF{}, 
+                                                        get_layoutSFB_TV(tiled_mma_pv),
+                                                        make_shape(size<1>(tile_shape_mnk), size<2>(tile_shape_mnk))
+                                                        );
+        auto smem_thr_copy_SFV = smem_tiled_copy_SFV.get_thread_slice(thread_idx);
+        Tensor tOsSFVt = smem_thr_copy_SFV.partition_S(as_position_independent_swizzle_tensor(sSFVt));
+        Tensor tOrSFVt_copy_view = smem_thr_copy_SFV.retile_D(tOrSFVt);
+
+        auto consumer_wait = [](auto& pipeline, auto& smem_pipe_read) {
+            auto barrier_token = pipeline.consumer_try_wait(smem_pipe_read);
+            pipeline.consumer_wait(smem_pipe_read, barrier_token);
+        };
+
+        int const seqlen_q = get<0>(mainloop_params.shape_Q);
+        int const seqlen_k = get<0>(mainloop_params.shape_K);
+        int const unpadded_seqlen_k = get<0>(mainloop_params.unpadded_shape_K);
+        int n_block = n_block_count - 1;
+
+        auto copy_k_block = [&](auto block_id) {
+            auto tSsK_stage = tSsK(_, _, _, smem_pipe_read_k.index());
+            auto tSsSFK_stage = tSsSFK(_, _, _, smem_pipe_read_k.index());
+            copy(smem_tiled_copy_K, tSsK_stage(_, _, block_id), tSrK_copy_view(_, _, block_id));
+            copy(smem_tiled_copy_SFK, tSsSFK_stage(_, _, block_id), tSrSFK_copy_view(_, _, block_id));
+        };
+        
+        auto copy_v_block = [&](auto block_id) {
+            auto tOsVt_stage = tOsVt(_, _, _, smem_pipe_read_v.index());
+            auto tOsSFVt_stage = tOsSFVt(_, _, _, smem_pipe_read_v.index());
+            copy(smem_tiled_copy_V, tOsVt_stage(_, _, block_id), tOrVt_copy_view(_, _, block_id));
+            copy(smem_tiled_copy_SFV, tOsSFVt_stage(_, _, block_id), tOrSFVt_copy_view(_, _, block_id));
+        };
+        // auto gemm_qk = [&](auto block_id) {
+        //     cute::gemm(tiled_mma_qk, make_zip_tensor(tSrQ(_, _, block_id), tSrSFQ(_, _, block_id)), make_zip_tensor(tSrK(_, _, block_id), tSrSFK(_, _, block_id)), tSrS);
+        // };
+        // auto gemm_pv = [&](auto block_id) {
+        //     cute::gemm(tiled_mma_pv, make_zip_tensor(tOrP(_, _, block_id), tOrSFP(_, _, block_id)), make_zip_tensor(tOrVt(_, _, block_id), tOrSFVt(_, _, block_id)), tOrO);
+        // };
+        auto add_delta_s = [&](auto& acc) {
+            auto tSsDS_stage = recast<float4>(sDS(_, _, smem_pipe_read_k.index()));
+            auto acc_float4 = recast<float4>(acc);
+            int quad_id = (threadIdx.x % 4) * 2;
+            for (int i = 0; i < 4; i++) {
+                auto num = quad_id + i * 8;
+                float4 delta_s_0 = tSsDS_stage(make_coord(_0{}, _0{}), make_coord(num, _0{}));
+                float4 delta_s_1 = tSsDS_stage(make_coord(_0{}, _0{}), make_coord(num + 1, _0{}));
+                acc_float4(make_coord(make_coord(_0{}, _0{}), _0{}), _0{}, i) = delta_s_0;
+                acc_float4(make_coord(make_coord(_0{}, _0{}), _1{}), _0{}, i) = delta_s_0;
+                acc_float4(make_coord(make_coord(_0{}, _1{}), _0{}), _0{}, i) = delta_s_1;
+                acc_float4(make_coord(make_coord(_0{}, _1{}), _1{}), _0{}, i) = delta_s_1;
+            }
+        };
+        consumer_wait(pipeline_q, smem_pipe_read_q);
+        copy(smem_tiled_copy_Q, tSsQ, tSrQ_copy_view);
+        copy(smem_tiled_copy_SFQ, tSsSFQ, tSrSFQ_copy_view);
+        pipeline_q.consumer_release(smem_pipe_read_q);
+        ++smem_pipe_read_q;
+
+        Tensor tSrS = partition_fragment_C(tiled_mma_qk, select<0, 1>(TileShape_MNK{}));
+        Tensor tSrS_converion_view = make_tensor(tSrS.data(), flash::convert_to_conversion_layout(tSrS.layout()));
+        Tensor AbsMaxP = make_tensor_like<float>(
+            make_layout(shape(group<1, 4>(flatten(tSrS_converion_view.layout()(make_coord(_0{}, _), _, _)))))
+        );
+        consumer_wait(pipeline_k, smem_pipe_read_k);
+        copy_k_block(_0{});
+        add_delta_s(tSrS);
+        CUTLASS_PRAGMA_UNROLL
+        for (int k_block = 0; k_block < size<2>(tSrQ); ++k_block) {
+            cute::gemm(tiled_mma_qk, make_zip_tensor(tSrQ(_, _, k_block), tSrSFQ(_, _, k_block)), 
+                                    make_zip_tensor(tSrK(_, _, k_block), tSrSFK(_, _, k_block)), tSrS);
+            if (k_block < size<2>(tSrQ) - 1) {
+                copy_k_block(k_block + 1);
+            } else {
+                pipeline_k.consumer_release(smem_pipe_read_k);
+                ++smem_pipe_read_k;
+            }
+        }
+        
+         
+        auto col_limit_causal = [&](int row, int n_block) {
+            return row + 1 + seqlen_k - n_block * kBlockN - seqlen_q + m_block * kBlockM;
+        };
+        {
+            Tensor cS = cute::make_identity_tensor(select<0, 1>(TileShape_MNK{}));
+            Tensor tScS = thread_mma_qk.partition_C(cS);
+            CUTLASS_PRAGMA_UNROLL
+            for (int i = 0; i < size(tSrS); ++i) {
+                if constexpr (!Is_causal) {  // Just masking based on col
+                    if (int(get<1>(tScS(i))) >= int(unpadded_seqlen_k - n_block * kBlockN)) { tSrS(i) = -INFINITY; }
+                } else { 
+                    if (int(get<1>(tScS(i))) >= std::min(seqlen_k - n_block * kBlockN,
+                                                        col_limit_causal(int(get<0>(tScS(i))), n_block))) {
+                        tSrS(i) = -INFINITY;
+                    }
+                }
+            }
+        }
+        auto quantize = [&](auto mma_k, auto acc_conversion_view) {
+            Tensor AbsMaxP_stagek = AbsMaxP(_, make_coord(_, _, mma_k));
+            Tensor acc_conversion_stagek = acc_conversion_view(_, _, mma_k);
+            Tensor SFP = make_tensor_like<cutlass::float_ue4m3_t>(AbsMaxP_stagek.layout());
+            Tensor SFP_uint32_view = recast<uint32_t>(SFP);
+            CUTLASS_PRAGMA_UNROLL
+            for (int i = 0; i < size(AbsMaxP_stagek); i += 4) {
+                uint32_t& tmp = SFP_uint32_view(i / 4);
+                flash::packed_float_to_ue4m3(
+                    AbsMaxP_stagek(i), 
+                    AbsMaxP_stagek(i + 1), 
+                    AbsMaxP_stagek(i + 2), 
+                    AbsMaxP_stagek(i + 3), 
+                    tmp
+                );
+            }
+            int const quad_id = threadIdx.x & 3;
+            uint32_t MASK = (0xFF00FF) << ((quad_id & 1) * 8);
+            Tensor tOrSFP_uint32_view = recast<uint32_t>(tOrSFP(_, _, mma_k));
+            Tensor tOrP_uint32_view = recast<uint32_t>(tOrP(_, _, mma_k));
+        
+            CUTLASS_PRAGMA_UNROLL
+            for (int mma_m = 0; mma_m < size<1>(tOrP); ++mma_m) {
+                    CUTLASS_PRAGMA_UNROLL
+                    for (int i = 0; i < 4; ++i) {
+                        flash::packed_float_to_e2m1(
+                            acc_conversion_stagek(make_coord(_0{}, i), mma_m),
+                            acc_conversion_stagek(make_coord(_1{}, i), mma_m),
+                            acc_conversion_stagek(make_coord(_2{}, i), mma_m),
+                            acc_conversion_stagek(make_coord(_3{}, i), mma_m),
+                            acc_conversion_stagek(make_coord(_4{}, i), mma_m),
+                            acc_conversion_stagek(make_coord(_5{}, i), mma_m),
+                            acc_conversion_stagek(make_coord(_6{}, i), mma_m),
+                            acc_conversion_stagek(make_coord(_7{}, i), mma_m),
+                            tOrP_uint32_view(i, mma_m)
+                        );
+                    }
+                    uint32_t local_sfp = SFP_uint32_view(_0{}, _0{}, mma_m);
+                    uint32_t peer_sfp  = __shfl_xor_sync(int32_t(-1), local_sfp, 2);
+                    if ((quad_id & 1) == 0) {
+                        uint32_t sfp = (local_sfp & MASK) | ((peer_sfp & MASK) << 8);
+                        tOrSFP_uint32_view(_0{}, mma_m) = sfp;
+                    } else {
+                        uint32_t sfp = (peer_sfp & MASK) | ((local_sfp & MASK) >> 8);
+                        tOrSFP_uint32_view(_0{}, mma_m) = sfp;
+                    }
+            }
+        };
+
+        softmax_fused.template online_softmax_with_quant</*Is_first=*/true>(tSrS, AbsMaxP, mainloop_params.softmax_scale_log2);
+
+        consumer_wait(pipeline_v, smem_pipe_read_v);
+        copy_v_block(_0{});
+        quantize(_0{}, tSrS_converion_view);
+        CUTLASS_PRAGMA_UNROLL
+        for (int v_block = 0; v_block < size<2>(tOrP); ++v_block) {
+            cute::gemm(tiled_mma_pv, make_zip_tensor(tOrP(_, _, v_block), tOrSFP(_, _, v_block)), 
+                                    make_zip_tensor(tOrVt(_, _, v_block), tOrSFVt(_, _, v_block)), tOrO_store);
+            if (v_block < size<2>(tOrP) - 1) {
+                copy_v_block(v_block + 1);
+                quantize(v_block + 1, tSrS_converion_view);
+            } else {
+                pipeline_v.consumer_release(smem_pipe_read_v);
+                ++smem_pipe_read_v;
+            }
+        }
+        
+        n_block--;
+        constexpr int n_masking_steps = !Is_causal ? 1 : cute::ceil_div(kBlockM, kBlockN) + 1;
+        // // Only go through these if Is_causal, since n_masking_steps = 1 when !Is_causal
+        CUTLASS_PRAGMA_UNROLL
+        for (int masking_step = 0; masking_step < n_masking_steps - 1 && n_block >= 0; ++masking_step, --n_block) {
+            Tensor tSrS = partition_fragment_C(tiled_mma_qk, select<0, 1>(TileShape_MNK{}));
+            Tensor tSrS_converion_view = make_tensor(tSrS.data(), flash::convert_to_conversion_layout(tSrS.layout()));
+            consumer_wait(pipeline_k, smem_pipe_read_k);
+            copy_k_block(_0{});
+            add_delta_s(tSrS);
+            CUTLASS_PRAGMA_UNROLL
+            for (int k_block = 0; k_block < size<2>(tSrQ); ++k_block) {
+                cute::gemm(tiled_mma_qk, make_zip_tensor(tSrQ(_, _, k_block), tSrSFQ(_, _, k_block)), 
+                                    make_zip_tensor(tSrK(_, _, k_block), tSrSFK(_, _, k_block)), tSrS);
+                if (k_block < size<2>(tSrQ) - 1) {
+                    copy_k_block(k_block + 1);
+                }
+            }
+            pipeline_k.consumer_release(smem_pipe_read_k);  // release K
+            ++smem_pipe_read_k;
+            Tensor cS = cute::make_identity_tensor(select<0, 1>(TileShape_MNK{}));
+            Tensor tScS = thread_mma_qk.partition_C(cS);
+            #pragma unroll
+            for (int i = 0; i < size(tSrS); ++i) {
+                if (int(get<1>(tScS(i))) >= col_limit_causal(int(get<0>(tScS(i))), n_block - 1)) {
+                    tSrS(i) = -INFINITY;
+                }
+            }
+            softmax_fused.template online_softmax_with_quant</*Is_first=*/false>(tSrS, AbsMaxP, mainloop_params.softmax_scale_log2);
+            Tensor tOrO = make_fragment_like(tOrO_store);
+            consumer_wait(pipeline_v, smem_pipe_read_v);
+            copy_v_block(_0{});
+            quantize(_0{}, tSrS_converion_view);
+            CUTLASS_PRAGMA_UNROLL
+            for (int v_block = 0; v_block < size<2>(tOrP); ++v_block) {
+                cute::gemm(tiled_mma_pv, make_zip_tensor(tOrP(_, _, v_block), tOrSFP(_, _, v_block)), 
+                                    make_zip_tensor(tOrVt(_, _, v_block), tOrSFVt(_, _, v_block)), tOrO);
+                if (v_block < size<2>(tOrP) - 1) {
+                    copy_v_block(v_block + 1);
+                    quantize(v_block + 1, tSrS_converion_view);
+                }
+            }
+            pipeline_v.consumer_release(smem_pipe_read_v);
+            ++smem_pipe_read_v;
+            if (masking_step > 0) { softmax_fused.rescale_o(tOrO_store, tOrO); }
+        }
+
+        #pragma unroll 1
+        for (; n_block >= 0; --n_block) {
+            Tensor tSrS = partition_fragment_C(tiled_mma_qk, select<0, 1>(TileShape_MNK{}));
+            Tensor tSrS_converion_view = make_tensor(tSrS.data(), flash::convert_to_conversion_layout(tSrS.layout()));
+            consumer_wait(pipeline_k, smem_pipe_read_k);
+            copy_k_block(_0{});
+            add_delta_s(tSrS);
+            CUTLASS_PRAGMA_UNROLL
+            for (int k_block = 0; k_block < size<2>(tSrQ); ++k_block) {
+                cute::gemm(tiled_mma_qk, make_zip_tensor(tSrQ(_, _, k_block), tSrSFQ(_, _, k_block)), 
+                                    make_zip_tensor(tSrK(_, _, k_block), tSrSFK(_, _, k_block)), tSrS);
+                if (k_block < size<2>(tSrQ) - 1) {
+                    copy_k_block(k_block + 1);
+                } else {
+                    pipeline_k.consumer_release(smem_pipe_read_k);
+                    ++smem_pipe_read_k;
+                }
+            }
+            softmax_fused.template online_softmax_with_quant</*Is_first=*/false>(tSrS, AbsMaxP, mainloop_params.softmax_scale_log2);
+            Tensor tOrO = make_fragment_like(tOrO_store);
+            consumer_wait(pipeline_v, smem_pipe_read_v);
+            copy_v_block(_0{});
+            quantize(_0{}, tSrS_converion_view);
+            CUTLASS_PRAGMA_UNROLL
+            for (int v_block = 0; v_block < size<2>(tOrP); ++v_block) {
+                cute::gemm(tiled_mma_pv, make_zip_tensor(tOrP(_, _, v_block), tOrSFP(_, _, v_block)), 
+                                    make_zip_tensor(tOrVt(_, _, v_block), tOrSFVt(_, _, v_block)), tOrO);
+                if (v_block < size<2>(tOrP) - 1) {
+                    copy_v_block(v_block + 1);
+                    quantize(v_block + 1, tSrS_converion_view);
+                } else {
+                    pipeline_v.consumer_release(smem_pipe_read_v);
+                    ++smem_pipe_read_v;
+                }
+            }
+            softmax_fused.rescale_o(tOrO_store, tOrO);
+        }
+        softmax_fused.finalize(tOrO_store);
+        return;
+    }
+
+};
+
+} // namespace flash
+

fastvideo-kernel/attn_qat_infer/blackwell/named_barrier.h

@@ -0,0 +1,119 @@
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#pragma once
+
+#include "cutlass/arch/barrier.h"
+#include "cutlass/pipeline/sm90_pipeline.hpp"
+
+namespace flash {
+
+enum class FP4NamedBarriers {
+    QueryEmpty = 1,
+    WarpSpecializedConsumer = 2,
+    WarpSpecializedPingPongConsumer1 = 3,
+    WarpSpecializedPingPongConsumer2 = 4,
+    ProducerEnd = 5,
+    ConsumerEnd = 6,
+    EpilogueBarrier = 7
+};
+
+template<int SequenceDepth, int SequenceLength>
+struct OrderedSequenceBarrierVarGroupSizeSharedStorage {
+  using Barrier = cutlass::arch::ClusterBarrier;
+  Barrier barrier_[SequenceDepth][SequenceLength];
+};
+
+ template<int SequenceDepth_, int SequenceLength_>
+class OrderedSequenceBarrierVarGroupSize {
+public:
+  static constexpr int SequenceDepth = SequenceDepth_;
+  static constexpr int SequenceLength = SequenceLength_;
+  using Barrier = cutlass::arch::ClusterBarrier;
+  using SharedStorage = flash::OrderedSequenceBarrierVarGroupSizeSharedStorage<SequenceDepth, SequenceLength>;
+  
+
+  struct Params {
+    uint32_t group_id;
+    uint32_t* group_size_list;
+  };
+
+private :
+  // In future this Params object can be replaced easily with a CG object
+  Params params_;
+  Barrier *barrier_ptr_;
+  cutlass::PipelineState<SequenceDepth> stage_;
+
+  static constexpr int Depth = SequenceDepth;
+  static constexpr int Length = SequenceLength;
+
+public:
+  OrderedSequenceBarrierVarGroupSize() = delete;
+  OrderedSequenceBarrierVarGroupSize(const OrderedSequenceBarrierVarGroupSize&) = delete;
+  OrderedSequenceBarrierVarGroupSize(OrderedSequenceBarrierVarGroupSize&&) = delete;
+  OrderedSequenceBarrierVarGroupSize& operator=(const OrderedSequenceBarrierVarGroupSize&) = delete;
+  OrderedSequenceBarrierVarGroupSize& operator=(OrderedSequenceBarrierVarGroupSize&&) = delete;
+  ~OrderedSequenceBarrierVarGroupSize() = default;
+
+  CUTLASS_DEVICE
+  OrderedSequenceBarrierVarGroupSize(SharedStorage& storage, Params const& params) :
+      params_(params),
+      barrier_ptr_(&storage.barrier_[0][0]),
+      // Group 0 - starts with an opposite phase
+      stage_({0, params.group_id == 0, 0}) {
+    int warp_idx = cutlass::canonical_warp_idx_sync();
+    int lane_predicate = cute::elect_one_sync();
+
+    // Barrier FULL, EMPTY init
+    // Init is done only by the one elected thread of the block
+    if (warp_idx == 0 && lane_predicate) {
+      for (int d = 0; d < Depth; ++d) {
+        for (int l = 0; l < Length; ++l) {
+          barrier_ptr_[d * Length + l].init(*(params.group_size_list + l));
+        }
+      }
+    }
+    cutlass::arch::fence_barrier_init();
+  }
+
+  // Wait on a stage to be unlocked
+  CUTLASS_DEVICE
+  void wait() {
+    get_barrier_for_current_stage(params_.group_id).wait(stage_.phase());
+  }
+
+  // Signal completion of Stage and move to the next stage
+  // (group_id) signals to (group_id+1)
+  CUTLASS_DEVICE
+  void arrive() {
+    int signalling_id = (params_.group_id + 1) % Length;
+    get_barrier_for_current_stage(signalling_id).arrive();
+    ++stage_;
+  }
+
+  CUTLASS_DEVICE
+  void advance() {
+    ++stage_;
+  }
+
+private:
+
+  CUTLASS_DEVICE
+  Barrier& get_barrier_for_current_stage(int group_id) {
+    return barrier_ptr_[stage_.index() * Length + group_id];
+  }
+};
+ 
+ } // flash

fastvideo-kernel/attn_qat_infer/blackwell/params.h

@@ -0,0 +1,180 @@
+// Modified from the original SageAttention3 code
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+ #pragma once
+
+ #include <cuda.h>
+ #include <vector>
+ 
+ #ifdef OLD_GENERATOR_PATH
+ #include <ATen/CUDAGeneratorImpl.h>
+ #else
+ #include <ATen/cuda/CUDAGeneratorImpl.h>
+ #endif
+ 
+ #include <ATen/cuda/CUDAGraphsUtils.cuh> // For at::cuda::philox::unpack
+ 
+ #include "cutlass/fast_math.h"  // For cutlass::FastDivmod
+ 
+ ////////////////////////////////////////////////////////////////////////////////////////////////////
+ 
+ struct Qkv_params {
+     using index_t = int64_t;
+     // The QKV matrices.
+     void *__restrict__ q_ptr;
+     void *__restrict__ k_ptr;
+     void *__restrict__ v_ptr;
+     void *__restrict__ delta_s_ptr;
+     // The QKV scale factor matrices.
+     void *__restrict__ sfq_ptr;
+     void *__restrict__ sfk_ptr;
+     void *__restrict__ sfv_ptr;
+     // The stride between rows of the Q, K and V matrices.
+     index_t q_batch_stride;
+     index_t k_batch_stride;
+     index_t v_batch_stride;
+     index_t q_row_stride;
+     index_t k_row_stride;
+     index_t v_row_stride;
+     index_t q_head_stride;
+     index_t k_head_stride;
+     index_t v_head_stride;
+     index_t ds_batch_stride;
+     index_t ds_row_stride;
+     index_t ds_head_stride;
+     // The stride of the Q, K and V scale factor matrices.
+     index_t sfq_batch_stride;
+     index_t sfk_batch_stride;
+     index_t sfv_batch_stride;
+     index_t sfq_row_stride;
+     index_t sfk_row_stride;
+     index_t sfv_row_stride;
+     index_t sfq_head_stride;
+     index_t sfk_head_stride;
+     index_t sfv_head_stride;
+ 
+     // The number of heads.
+     int h, h_k;
+     // In the case of multi-query and grouped-query attention (MQA/GQA), nheads_k could be
+     // different from nheads (query).
+     int h_h_k_ratio; // precompute h / h_k,
+ };
+ 
+ ////////////////////////////////////////////////////////////////////////////////////////////////////
+ 
+ struct Flash_fwd_params : public Qkv_params {
+ 
+     // The O matrix (output).
+     void * __restrict__ o_ptr;
+     void * __restrict__ oaccum_ptr;
+     void * __restrict__ s_ptr;
+ 
+     // The stride between rows of O.
+     index_t o_batch_stride;
+     index_t o_row_stride;
+     index_t o_head_stride;
+ 
+     // The pointer to the P matrix.
+     void * __restrict__ p_ptr;
+ 
+     // The pointer to the softmax sum.
+     void * __restrict__ softmax_lse_ptr;
+     void * __restrict__ softmax_lseaccum_ptr;
+ 
+     // The dimensions.
+     int b, seqlen_q, seqlen_k, seqlen_knew, d, seqlen_q_rounded, seqlen_k_rounded, d_rounded, rotary_dim, unpadded_seqlen_k;
+     cutlass::FastDivmod head_divmod, m_block_divmod;
+     int total_blocks;
+     int seqlen_s;
+ 
+     // The scaling factors for the kernel.
+     float scale_softmax;
+     float scale_softmax_log2;
+     uint32_t scale_softmax_log2_half2;
+ 
+     // array of length b+1 holding starting offset of each sequence.
+     int * __restrict__ cu_seqlens_q;
+     int * __restrict__ cu_seqlens_k;
+ 
+     // If provided, the actual length of each k sequence.
+     int * __restrict__ seqused_k;
+ 
+     int *__restrict__ blockmask;
+ 
+     // The K_new and V_new matrices.
+     void * __restrict__ knew_ptr;
+     void * __restrict__ vnew_ptr;
+ 
+     // The stride between rows of the Q, K and V matrices.
+     index_t knew_batch_stride;
+     index_t vnew_batch_stride;
+     index_t knew_row_stride;
+     index_t vnew_row_stride;
+     index_t knew_head_stride;
+     index_t vnew_head_stride;
+ 
+     // The cos and sin matrices for rotary embedding.
+     void * __restrict__ rotary_cos_ptr;
+     void * __restrict__ rotary_sin_ptr;
+ 
+     // The indices to index into the KV cache.
+     int * __restrict__ cache_batch_idx;
+ 
+     // Paged KV cache
+     int * __restrict__ block_table;
+     index_t block_table_batch_stride;
+     int page_block_size;
+ 
+     // The dropout probability (probability of keeping an activation).
+     float p_dropout;
+     // uint32_t p_dropout_in_uint;
+     // uint16_t p_dropout_in_uint16_t;
+     uint8_t p_dropout_in_uint8_t;
+ 
+     // Scale factor of 1 / (1 - p_dropout).
+     float rp_dropout;
+     float scale_softmax_rp_dropout;
+ 
+     // Local window size
+     int window_size_left, window_size_right;
+ 
+     // Random state.
+     at::PhiloxCudaState philox_args;
+ 
+     // Pointer to the RNG seed (idx 0) and offset (idx 1).
+     uint64_t * rng_state;
+ 
+     bool is_bf16;
+     bool is_e4m3;
+     bool is_causal;
+     bool per_block_mean;
+     bool single_level_p_quant;  // If true, use single-level 1x16 block scale quantization for P (like V), instead of two-level quantization
+     // If is_seqlens_k_cumulative, then seqlen_k is cu_seqlens_k[bidb + 1] - cu_seqlens_k[bidb].
+     // Otherwise it's cu_seqlens_k[bidb], i.e., we use cu_seqlens_k to store the sequence lengths of K.
+     bool is_seqlens_k_cumulative;
+ 
+     bool is_rotary_interleaved;
+ 
+     int num_splits;  // For split-KV version
+ 
+     void * __restrict__ alibi_slopes_ptr;
+     index_t alibi_slopes_batch_stride;
+ 
+     int * __restrict__ tile_count_semaphore;
+ };
+ 
+ ////////////////////////////////////////////////////////////////////////////////////////////////////

fastvideo-kernel/attn_qat_infer/blackwell/softmax_fused.h

@@ -0,0 +1,190 @@
+// Modified from the original SageAttention3 code
+/*
+ * Copyright (c) 2025 by SageAttention team.
+ * 
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+ #pragma once
+
+ #include <cmath>
+ #include "cute/tensor.hpp"
+ #include "cutlass/numeric_types.h"
+ #include "utils.h"
+ 
+ namespace flash {
+ 
+ using namespace cute;
+ 
+ template <int Rows>
+ struct SoftmaxFused{
+ 
+     using TensorT = decltype(make_fragment_like<float>(Shape<Int<Rows>>{}));
+     TensorT row_sum, row_max, scores_scale;
+     static constexpr float fp8_scalexfp4_scale = 1.f / (448 * 6);
+     static constexpr float fp8_scalexfp4_scale_log2 = -11.392317422778762f; //log2f(fp8_scalexfp4_scale)
+     static constexpr float fp4_scale_log2 = -2.584962500721156f; // log2f(fp4_scale)
+     static constexpr int RowReductionThr = 4;
+     
+     // If true, use single-level quantization: s_P2, P̂_2 = φ(P̃) directly (standard per-block FP4 quantization like V)
+     // If false (default), use two-level quantization: s_P1 = rowmax(P̃)/(448×6), then s_P2, P̂_2 = φ(P̃/s_P1)
+     bool single_level_p_quant;
+ 
+     CUTLASS_DEVICE SoftmaxFused(bool single_level = false) : single_level_p_quant(single_level) {};
+ 
+     template<bool FirstTile, bool InfCheck = false, typename TensorAcc, typename TensorMax>
+     CUTLASS_DEVICE auto online_softmax_with_quant(
+         TensorAcc& acc, 
+         TensorMax& AbsMaxP,
+         const float softmax_scale_log2
+     ) {
+         Tensor acc_reduction_view = make_tensor(acc.data(), flash::convert_to_reduction_layout(acc.layout()));
+         Tensor acc_conversion_view = make_tensor(acc.data(), flash::convert_to_conversion_layout(acc.layout()));
+         Tensor acc_conversion_flatten = group_modes<1, 5>(group_modes<0, 2>(flatten(acc_conversion_view)));
+         
+         if constexpr (FirstTile) {
+             fill(row_max, -INFINITY);
+             clear(row_sum);
+             fill(scores_scale, 1.f);
+ 
+             CUTLASS_PRAGMA_UNROLL
+             for (int mi = 0; mi < size<0>(acc_reduction_view); mi++) {
+                 CUTLASS_PRAGMA_UNROLL
+                 for (int ni = 0; ni < size<1, 1>(acc_reduction_view); ni++) {
+                     CUTLASS_PRAGMA_UNROLL
+                     for (int ei = 0; ei < size<1, 0>(acc_reduction_view); ei++) {
+                         AbsMaxP(mi, ni) = fmaxf(AbsMaxP(mi, ni), acc_reduction_view(mi, make_coord(ei, ni)));
+                     }
+                     float max_recv = __shfl_xor_sync(int32_t(-1), AbsMaxP(mi, ni), 1); // exchange max with neighbour thread of 8 elements
+                     AbsMaxP(mi, ni) = fmaxf(AbsMaxP(mi, ni), max_recv);
+                     row_max(mi) = fmaxf(row_max(mi), AbsMaxP(mi, ni));
+                 }
+                 
+                 float max_recv = __shfl_xor_sync(int32_t(-1), row_max(mi), 2); // exchange max in a quad in a row
+                 row_max(mi) = fmaxf(row_max(mi), max_recv);
+ 
+                 // Two-level P quantization (default): s_P1 = rowmax(P̃)/(448×6), then s_P2,P̂_2 = φ(P̃/s_P1)
+                 //   - Pre-scales P to [0, 448×6] range before φ, output scaled by s_P1
+                 // Single-level P quantization: s_P2, P̂_2 = φ(P̃) directly (like V quantization)
+                 //   - No s_P1, just standard per-block FP4 quantization φ
+                 const float s_P1_offset = single_level_p_quant ? 0.f : fp8_scalexfp4_scale_log2;
+                 const float max_scaled = InfCheck
+                                         ? (row_max(mi) == -INFINITY ? 0.f : (row_max(mi) * softmax_scale_log2 + s_P1_offset))
+                                         : (row_max(mi) * softmax_scale_log2 + s_P1_offset);
+                 CUTLASS_PRAGMA_UNROLL
+                 for (int ni = 0; ni < size<1>(acc_reduction_view); ni++) {
+                     acc_reduction_view(mi, ni) = flash::ptx_exp2(acc_reduction_view(mi, ni) * softmax_scale_log2 - max_scaled);
+                 }
+                 // s_P2 = max(P_block)/6 — per-block scale factor from φ function (same formula for both modes)
+                 // The difference is in max_scaled: two-level includes 448×6 pre-scaling, single-level doesn't
+                 CUTLASS_PRAGMA_UNROLL
+                 for (int sfi = 0; sfi < size<1>(AbsMaxP); sfi++) {
+                     AbsMaxP(mi, sfi) = flash::ptx_exp2(AbsMaxP(mi, sfi) * softmax_scale_log2 - max_scaled + fp4_scale_log2);
+                 }
+             }
+             CUTLASS_PRAGMA_UNROLL
+             for (int mi = 0; mi < size<0>(acc_reduction_view); mi++) {
+                 CUTLASS_PRAGMA_UNROLL
+                 for (int ni = 0; ni < size<1>(acc_reduction_view); ni++) {
+                     row_sum(mi) += acc_reduction_view(mi, ni);
+                 }
+             }
+         }
+         else {
+             Tensor scores_max_prev = make_fragment_like(row_max);
+             cute::copy(row_max, scores_max_prev);
+             CUTLASS_PRAGMA_UNROLL
+             for (int mi = 0; mi < size<0>(acc_reduction_view); mi++) {
+                 CUTLASS_PRAGMA_UNROLL
+                 for (int ni = 0; ni < size<1, 1>(acc_reduction_view); ni++) {
+                     float local_max = -INFINITY;
+                     CUTLASS_PRAGMA_UNROLL
+                     for (int ei = 0; ei < size<1, 0>(acc_reduction_view); ei++) {
+                         local_max = fmaxf(local_max, acc_reduction_view(mi, make_coord(ei, ni)));
+                     }
+                     float max_recv = __shfl_xor_sync(int32_t(-1), local_max, 1); // exchange max with neighbour thread of 8 elements
+                     AbsMaxP(mi, ni) = fmaxf(local_max, max_recv);
+                     row_max(mi) = fmaxf(row_max(mi), AbsMaxP(mi, ni));
+                 }
+                 
+                 float max_recv = __shfl_xor_sync(int32_t(-1), row_max(mi), 2); // exchange max in a quad in a row
+                 row_max(mi) = fmaxf(row_max(mi), max_recv);
+ 
+                 float scores_max_cur = !InfCheck
+                                         ? row_max(mi)
+                                         : (row_max(mi) == -INFINITY ? 0.0f : row_max(mi));
+                 scores_scale(mi) = flash::ptx_exp2((scores_max_prev(mi) - scores_max_cur) * softmax_scale_log2);
+ 
+                 // Two-level P quantization (default): s_P1 = rowmax(P̃)/(448×6), then s_P2,P̂_2 = φ(P̃/s_P1)
+                 // Single-level P quantization: s_P2, P̂_2 = φ(P̃) directly (like V quantization)
+                 const float s_P1_offset = single_level_p_quant ? 0.f : fp8_scalexfp4_scale_log2;
+                 const float max_scaled = InfCheck
+                                         ? (row_max(mi) == -INFINITY ? 0.f : (row_max(mi) * softmax_scale_log2 + s_P1_offset))
+                                         : (row_max(mi) * softmax_scale_log2 + s_P1_offset);
+                 row_sum(mi) = row_sum(mi) * scores_scale(mi);
+                 CUTLASS_PRAGMA_UNROLL
+                 for (int ni = 0; ni < size<1>(acc_reduction_view); ni++) {
+                     acc_reduction_view(mi, ni) = flash::ptx_exp2(acc_reduction_view(mi, ni) * softmax_scale_log2 - max_scaled);
+                     row_sum(mi) += acc_reduction_view(mi, ni);
+                 }
+                 // s_P2 = max(P_block)/6 — per-block scale factor from φ function
+                 CUTLASS_PRAGMA_UNROLL
+                 for (int sfi = 0; sfi < size<1>(AbsMaxP); sfi++) {
+                     AbsMaxP(mi, sfi) = flash::ptx_exp2(AbsMaxP(mi, sfi) * softmax_scale_log2 - max_scaled + fp4_scale_log2);
+                 }
+                 // scores_scale(mi) = max_scaled;
+             }
+         }
+         CUTLASS_PRAGMA_UNROLL
+         for (int i = 0; i < size(AbsMaxP); ++i) {
+             CUTLASS_PRAGMA_UNROLL
+             for (int j = 0; j < size<0>(acc_conversion_flatten); ++j)
+                 acc_conversion_flatten(j, i) /= AbsMaxP(i);
+         }
+     }
+ 
+     template<typename TensorAcc>
+     CUTLASS_DEVICE void finalize(TensorAcc& o_store) {
+         Tensor o_store_reduction_view = make_tensor(o_store.data(), flash::convert_to_reduction_layout(o_store.layout()));
+         CUTLASS_PRAGMA_UNROLL
+         for (int mi = 0; mi < size(row_max); ++mi) {
+             CUTLASS_PRAGMA_UNROLL
+             for (int i = 1; i < RowReductionThr; i <<= 1) {
+                 float sum_recv = __shfl_xor_sync(int32_t(-1), row_sum(mi), i);
+                 row_sum(mi) += sum_recv;
+             }
+             float sum = row_sum(mi);
+             float inv_sum = (sum == 0.f || sum != sum) ? 0.f : 1 / sum;
+             CUTLASS_PRAGMA_UNROLL
+             for (int ni = 0; ni < size<1>(o_store_reduction_view); ++ni) { 
+                 o_store_reduction_view(mi, ni) *= inv_sum;
+              }
+         }
+     }
+ 
+     template<typename TensorAcc>
+     CUTLASS_DEVICE void rescale_o(TensorAcc& o_store, TensorAcc const& o_tmp) {
+         Tensor o_store_reduction_view = make_tensor(o_store.data(), flash::convert_to_reduction_layout(o_store.layout()));
+         Tensor o_tmp_reduction_view = make_tensor(o_tmp.data(), flash::convert_to_reduction_layout(o_tmp.layout()));
+         CUTLASS_PRAGMA_UNROLL
+         for (int mi = 0; mi < size(row_max); ++mi) {
+             CUTLASS_PRAGMA_UNROLL
+             for (int ni = 0; ni < size<1>(o_store_reduction_view); ++ni) { 
+                 o_store_reduction_view(mi, ni) = o_store_reduction_view(mi, ni) * scores_scale(mi) + o_tmp_reduction_view(mi, ni);
+              }
+         }
+ 
+     }
+ 
+ 
+ };
+ } // namespace flash