[perf] Probe H3 chunked all-to-all with direct output slices

Two defects measured on 4x GB200, MiniMax-H3 geometry, per attention layer
(NCCL baseline 2421us at sp=4, 1295us at sp=8 across two trays):

  as shipped   sp=4 2035us (1.19x)   sp=8 3102us (2.4x SLOWER than NCCL)
  with these   sp=4 1590us (1.52x)   sp=8 declines, 1300us

ncclTeamLsa answers "addressable", not "fast". NCCL 2.29 extends the LSA
team across a multi-node NVLink domain, so on a GB200 rack the gate passes
for ranks on different trays and the kernel arms. Its fine-grained 16B
remote stores are far slower there than NCCL's bulk transfers. Require a
single host, which is the regime the slab decomposition was tuned for, and
which matches flashinfer's own gate. torch 2.12 (the pin) bundles NCCL
2.29.7, so this is reachable today.

_can_attempt now computes its local verdict without collectives and then
runs one unconditional all_gather_object carrying (hostname, local_ok).
A collective behind a rank-local early return hangs the group whenever
ranks disagree -- exactly the case this gate exists to detect. Verified:
with one rank reporting the kernel unavailable, the earlier ordering timed
out at 180s while this completes with correct results on every rank.

The per-call agreement cost a flat ~227us regardless of operand size -- a
host-side gloo all_gather before every collective. The signature is
architectural: two distinct values (scatter and gather shapes) across 50
layers x 4 steps. Cache the verdict so the collective runs twice per
generation instead of ~400 times.

The cache trades one property: a rank whose signature diverges mid-run now
misses the cache and calls the collective alone, hanging rather than
falling back. Re-voting every N calls would bound that if wanted.

.agents/skills/ci-runner/SKILL.md

@@ -0,0 +1,99 @@
+---
+name: ci-runner
+description: Work on FastVideo's Slurm-only, change-aware GPU CI lanes, static Buildkite graph, trusted ci-runner policy, lane scripts, and GB200 validation.
+---
+
+# Slinky Slurm CI lanes
+
+FastVideo's `ci-runner` Buildkite queue is the control plane for all active
+GPU CI. A private host-owned dispatcher leases GPUs from the Slinky Slurm tray
+and runs the immutable PR SHA inside an isolated Enroot container. Buildkite
+pipeline upload and Slurm submission occur on the login plane; every test
+payload executes on Slurm compute.
+
+The files under `fastvideo/tests/modal/` and `.buildkite/scripts/pr_test.sh`
+are dormant rollback code. Never add an active Buildkite or slash-command
+route to them. `pr_test.sh` must continue to reject Buildkite invocations.
+
+The private operator bundle is deliberately outside this repository because
+it contains site paths and credentials. See
+`docs/contributing/ci_architecture.md`; this skill covers the repository half
+and the coordination contract with that bundle.
+
+## Invariants
+
+- `.buildkite/pipeline.yml` contains exactly one static step for every active
+  GPU lane. Each step pins a unique key and label, a 90-minute timeout, the
+  trusted `/opt/fastvideo-ci-runner/run-ci` command (`run-unit` is the one
+  compatibility wrapper), step-level internal `TEST_TYPE`, and
+  `queue: "ci-runner"`.
+- Active CI contains no `pr_test.sh` command, Modal invocation, default queue,
+  Buildkite plugin, `soft_fail`, or job-controlled artifact glob.
+- The six Fastcheck lanes use `:microscope:` labels. Full-Suite-only lanes use
+  `:test_tube:` or `:bar_chart:` so direct reruns update the right aggregate.
+- SSIM and vanilla training request all four GPUs. Keep both in the
+  `fastvideo/slinky/whole-tray` Buildkite concurrency group with a limit of one
+  so the second job does not consume an agent or command timeout while waiting
+  for the same tray.
+- `/test full` schedules all twenty lanes. `/merge`, `ready`, and new pushes to
+  ready PRs use the trusted base-branch planner in
+  `.github/scripts/plan_merge_ci.py`: automatic Fastcheck remains the universal
+  six-lane baseline, and the merge build adds only path-relevant integration
+  lanes. Unknown source/build paths fail closed to all fourteen additive lanes.
+  The trusted uploader still normalizes and validates the complete static graph
+  before Buildkite evaluates its plan conditions.
+- Focused merge builds may pass allowlisted golden-gate and SSIM test basenames.
+  The private host validates the lane plan and basenames before staging them,
+  and the in-container scripts validate them again. Direct `/test ssim`,
+  explicit `/test full`, and the weekly main-branch schedule run the complete
+  SSIM matrix.
+- The trusted uploader serves exactly three entry pipelines:
+  `pr-fastcheck` for automatic PR builds, `ci` for slash-command/ready-label
+  API builds, and `fastvideo-performance-lane` for the weekly schedule. Keep
+  incoming GitHub webhook processing disabled on `ci` so it cannot duplicate
+  `pr-fastcheck` on every PR update.
+- Test payloads live in `.buildkite/scripts/unit_test.sh` or executable
+  `.buildkite/scripts/lanes/<lane>.sh`. Backend policy (GPU count, extras,
+  secrets, kernel build, artifacts) stays in the agent-owned lane table.
+- Tests must preserve an inherited `MASTER_PORT`. Packed containers share the
+  tray network namespace, so the private runner assigns a distinct port range
+  per GPU lease and the SSIM scheduler assigns task offsets within its range.
+- The ARM64 runner image includes the pinned FA4 CuTe overlay validated on
+  GB200. Keep SSIM at `FASTVIDEO_FA4=1` because its references were seeded with
+  FA4; keep lanes with FA2 baselines at `FASTVIDEO_FA4=0`. A runner image change
+  must revalidate both the FA4 import and an actual GB200 forward kernel.
+- `fastvideo/tests/ssim/ci_runner.py` is the active four-GPU SSIM scheduler.
+  New SSIM files are discovered through `REQUIRED_GPUS` and
+  `*_MODEL_TO_PARAMS`; do not wire them through the dormant Modal scheduler.
+- The host policy fail-closes unknown tuples. A repository-side lane change is
+  inert until the operator updates the private lane table and uploader policy
+  in the same rollout.
+
+## Adding or changing a lane
+
+1. Read the closest `AGENTS.md` and the domain-specific testing guide.
+2. Add or update the executable lane payload under `.buildkite/scripts/`.
+   Keep it deterministic and free of host-specific paths or credential fetches.
+3. Add the static pipeline step and canonical `/test <name>` mapping. Keep the
+   `<name>-ci` alias only when compatibility requires it.
+4. Add its source/test path ownership to `.github/scripts/plan_merge_ci.py`.
+   Prefer the narrowest correctness-preserving lane set; leave unknown paths
+   fail-closed. Extend `fastvideo/tests/contract/test_ci_test_collection.py`,
+   `test_merge_ci_plan.py`, and focused CPU-only scheduler/policy tests.
+5. Coordinate the private lane row: GPU count (1-4), wall time, script, scope
+   pairs, step key, command, HF cache/token, tracking mode, extras, attention
+   backend policy, kernel policy, and artifact relay. Active training lanes
+   keep W&B offline and do not stage a W&B credential.
+6. Update the trusted pipeline-uploader schema. A mismatch must reject the
+   pipeline rather than silently skip a lane.
+7. Run `pre-commit run --files <changed paths>`, the planner's representative
+   diff matrix, contract tests, private driver tests, and a real GB200 canary.
+   Multi-GPU, hardware-reference, training, performance, and SSIM changes need
+   their own target-hardware evidence.
+
+## Rollback
+
+Rollback the Slurm routing/configuration change or pause the `ci-runner` queue.
+Do not silently reactivate Modal. A manual Modal experiment requires the
+explicit local opt-in documented in `ci_architecture.md`; returning it to
+production CI needs a separate reviewed decision.

.agents/skills/reseed-ssim-references/SKILL.md

@@ -1,6 +1,6 @@
 ---
 name: reseed-ssim-references
-description: Re-seed HF reference videos for a single existing SSIM test on Modal L40S. Always backs up current refs locally first, regenerates on Modal, pauses for the user to eyeball before-vs-after quality, then overwrites the targeted `<model_id>` subtree on `FastVideo/ssim-reference-videos` with `--force`. Use when an intentional code change (model port fix, attention backend swap, kernel upgrade, hyperparameter change) has invalidated existing refs and they need to be regenerated. Pairs with `seed-ssim-references`, which is for first-time seeding only.
+description: Re-seed HF reference videos for a single existing SSIM test on Modal L40S. Always backs up current refs locally first, regenerates on Modal, pauses for the user to eyeball before-vs-after quality, then overwrites the targeted model subtree on `FastVideo/ssim-reference-videos` with `--force`. Use when an intentional code change (model port fix, attention backend swap, kernel upgrade, hyperparameter change) has invalidated existing refs and they need to be regenerated. Pairs with `seed-ssim-references`, which is for first-time seeding only.
 ---
 
 # Re-seed SSIM Reference Videos
@@ -13,7 +13,7 @@ on HF — the old refs are overwritten — so the skill always:
 
 1. Confirms intent with a one-liner the user has to type.
 2. Downloads the existing refs as a local, timestamped backup.
-3. Regenerates on Modal L40S (same code path that CI uses).
+3. Regenerates through the manual legacy Modal L40S maintenance path.
 4. Pauses for a side-by-side eyeball of backup vs new mp4s.
 5. Uploads with `--force`, scoped to the single `--model-id`.
 6. Reminds the user to keep the backup until the PR lands.
@@ -51,8 +51,9 @@ harder to recover from than failing closed.
 
 Hardcoded:
 
-- Modal GPU: **L40S** (matches CI; re-seeding from another SKU produces refs
-  that L40S CI cannot match).
+- Modal GPU: **L40S**. This is a manual reference-maintenance target, not the
+  active Slurm CI compute path; changing the SKU also changes the historical
+  `L40S_reference_videos` contract.
 - Quality tier: **`default`**. `full_quality` is a separate, deliberate
   operation.
 - HF repo: `FastVideo/ssim-reference-videos` (override via

.agents/skills/seed-ssim-references/SKILL.md

@@ -35,7 +35,8 @@ The skill is run **manually**, once per new test. Before invoking it, the user
 has already sanity-tested the new test locally — it launches `VideoGenerator`
 and writes an artefact without crashing (the missing-reference assertion at
 the end is expected). The skill does not re-test locally; it goes straight
-to Modal L40S (which is what CI uses).
+to the manual legacy Modal L40S reference-maintenance target. Active CI runs
+on the Slinky Slurm cluster and only consumes the resulting references.
 
 ## When to use
 
@@ -61,7 +62,8 @@ Prompt the user for it if they didn't supply it.
 
 Everything else is fixed:
 
-- Modal runner GPU: **L40S** (hardcoded in `fastvideo/tests/modal/ssim_test.py`).
+- Modal maintenance GPU: **L40S** (hardcoded in
+  `fastvideo/tests/modal/ssim_test.py`; this is not the active CI compute path).
 - Device folder: `L40S_reference_videos`.
 - Quality tier: `default` (the tier CI runs). The `full_quality` tier is not
   seeded by this skill.

.buildkite/scripts/lanes/api_server.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the OpenAI-compatible API lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/entrypoints/test_openai_api_integration.py -vs

.buildkite/scripts/lanes/distillation_dmd.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the distillation-DMD lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/training/distill/test_distill_dmd.py -vs

.buildkite/scripts/lanes/dreamverse.sh

@@ -0,0 +1,87 @@
+#!/usr/bin/env bash
+# DreamVerse needs a GPU for import-time device resolution, but it does not
+# build or exercise fastvideo-kernel. A checksummed Node archive is installed
+# in the disposable Slurm container because the shared CI image is
+# Python/CUDA focused.
+set -euo pipefail
+
+node_version=v22.23.2
+case $(uname -m) in
+  aarch64 | arm64)
+    node_arch=arm64
+    node_archive_sha256=013b59cfd2819703a6f4a14ab891fc46fc2a4e3f5bcd92de3fb4929b43e35b30
+    ;;
+  x86_64 | amd64)
+    node_arch=x64
+    node_archive_sha256=b294a556e639d64338823920e5866c21c02741742d2e1529ee1a225c1ec9252a
+    ;;
+  *)
+    echo "Unsupported architecture for DreamVerse Node runtime: $(uname -m)" >&2
+    exit 2
+    ;;
+esac
+node_archive="node-${node_version}-linux-${node_arch}.tar.gz"
+node_runtime_root=$(mktemp -d -t fastvideo-node.XXXXXX)
+node_archive_path="${node_runtime_root}/${node_archive}"
+node_install_dir="${node_runtime_root}/${node_archive%.tar.gz}"
+curl --proto '=https' --tlsv1.2 --retry 5 --retry-all-errors \
+  --location --fail --silent --show-error \
+  "https://nodejs.org/dist/${node_version}/${node_archive}" \
+  --output "$node_archive_path"
+printf '%s  %s\n' "$node_archive_sha256" "$node_archive_path" | sha256sum --check --status
+tar -xzf "$node_archive_path" -C "$node_runtime_root"
+export PATH="${node_install_dir}/bin:${PATH}"
+node --version
+npm --version
+
+export PYTHONPATH="$(pwd)/apps/dreamverse${PYTHONPATH:+:$PYTHONPATH}"
+pytest apps/dreamverse/dreamverse/tests -q
+
+cd apps/dreamverse/web
+npm ci
+npm run typecheck
+npm test
+machine_arch=$(uname -m)
+if [[ $machine_arch =~ ^(aarch64|arm64)$ ]]; then
+  npx playwright install --with-deps chromium firefox
+else
+  npx playwright install --with-deps chromium webkit firefox
+fi
+
+master_port=${MASTER_PORT:-7959}
+BACKEND_PORT=${BACKEND_PORT:-$((master_port + 50))}
+python -m uvicorn dreamverse.mock_server:app --host 127.0.0.1 --port "$BACKEND_PORT" &
+mock_server_pid=$!
+cleanup() {
+  kill "$mock_server_pid" 2>/dev/null || true
+  wait "$mock_server_pid" 2>/dev/null || true
+}
+trap cleanup EXIT INT TERM
+
+for _ in {1..30}; do
+  curl -fsS "http://127.0.0.1:$BACKEND_PORT/healthz" && break
+  sleep 1
+done
+curl -fsS "http://127.0.0.1:$BACKEND_PORT/healthz"
+
+if [[ $machine_arch =~ ^(aarch64|arm64)$ ]]; then
+  # Playwright WebKit traps before opening a page on Linux ARM64, and its
+  # bundled Chromium lacks the H.264/AAC codecs used by the fMP4 assertions.
+  # Firefox covers every flow, including streaming. Chromium and its mobile
+  # profile still cover all codec-independent UI behavior on GB200.
+  BACKEND_HOST=127.0.0.1 BACKEND_PORT="$BACKEND_PORT" CI=1 \
+    npm run e2e -- --project=firefox
+  BACKEND_HOST=127.0.0.1 BACKEND_PORT="$BACKEND_PORT" CI=1 \
+    npm run e2e -- \
+      --project=chromium \
+      --project=mobile-chromium \
+      --grep-invert='streams, plays, and surfaces a downloadable clip|starts a new project and switches back to the prior session|saved projects persist across a page reload'
+else
+  BACKEND_HOST=127.0.0.1 BACKEND_PORT="$BACKEND_PORT" CI=1 \
+    npm run e2e -- \
+      --project=chromium \
+      --project=webkit \
+      --project=firefox \
+      --project=mobile-safari \
+      --project=mobile-chromium
+fi

.buildkite/scripts/lanes/encoder.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the encoder lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/encoders -vs

.buildkite/scripts/lanes/eval.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the evaluation lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/eval -vs

.buildkite/scripts/lanes/golden_gate.sh

@@ -0,0 +1,35 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the golden-gate lane. Environment (HF_HOME
+# and authentication) is the runner's responsibility.
+set -euo pipefail
+
+golden_root=./fastvideo/tests/golden_gate
+selected=${FASTVIDEO_GOLDEN_TEST_FILES-}
+if [ -z "$selected" ]; then
+  if [ "${TEST_SCOPE:-}" = merge ]; then
+    echo "Missing FASTVIDEO_GOLDEN_TEST_FILES for merge scope" >&2
+    exit 2
+  fi
+  selected=all
+fi
+if [ "$selected" = all ]; then
+  exec pytest "$golden_root" -vs
+fi
+
+[[ $selected =~ ^test_[a-z0-9_]+\.py(,test_[a-z0-9_]+\.py)*$ ]] || {
+  echo "Invalid FASTVIDEO_GOLDEN_TEST_FILES selection" >&2
+  exit 2
+}
+
+IFS=, read -r -a golden_files <<< "$selected"
+golden_paths=()
+for golden_file in "${golden_files[@]}"; do
+  golden_path="$golden_root/$golden_file"
+  [ -f "$golden_path" ] || {
+    echo "Selected golden test does not exist: $golden_file" >&2
+    exit 2
+  }
+  golden_paths+=("$golden_path")
+done
+
+exec pytest "${golden_paths[@]}" -vs

.buildkite/scripts/lanes/inference_lora.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the LoRA-inference lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/inference/lora/test_lora_inference_similarity.py -vs

.buildkite/scripts/lanes/inference_vmoba.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the VMoBA-inference lane.
+set -euo pipefail
+
+exec python fastvideo/tests/inference/vmoba/test_vmoba_inference.py

.buildkite/scripts/lanes/kernel_tests.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the custom-kernel lane.
+set -euo pipefail
+
+exec pytest fastvideo-kernel/tests/ -vs

.buildkite/scripts/lanes/lora_extraction.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the LoRA-extraction lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/lora_extraction/test_lora_extraction.py -vs

.buildkite/scripts/lanes/performance.sh

@@ -0,0 +1,52 @@
+#!/usr/bin/env bash
+# Canonical Slurm performance lane. Reports are written outside the checkout
+# so the trusted host driver can upload them after untrusted code exits.
+set -uo pipefail
+
+export PERFORMANCE_TRACKING_ROOT=/tmp/perf-tracking
+export PERF_REPORTS_DIR=/workspace/artifacts/performance
+mkdir -p "$PERF_REPORTS_DIR"
+
+if [[ ${BUILDKITE_PULL_REQUEST:-false} =~ ^[1-9][0-9]*$ ]]; then
+  export PERF_RUN_SOURCE=pr
+  export PERF_UPLOAD_POLICY=pass
+elif [ "${BUILDKITE_BRANCH:-}" = main ] \
+  && { [ "${BUILDKITE_SOURCE:-}" = schedule ] || [ "${TEST_SCOPE:-}" = full ]; }; then
+  export PERF_RUN_SOURCE=scheduled_main
+  export PERF_UPLOAD_POLICY=always
+elif [ "${TEST_SCOPE:-}" = direct ]; then
+  export PERF_RUN_SOURCE=unknown
+  export PERF_UPLOAD_POLICY=pass
+else
+  export PERF_RUN_SOURCE=unknown
+  export PERF_UPLOAD_POLICY=never
+fi
+
+nvidia-smi \
+  --query-gpu=index,timestamp,clocks.sm,clocks.max.sm,power.draw,power.limit,temperature.gpu \
+  --format=csv -l 10 > "$PERF_REPORTS_DIR/gpu_telemetry.csv" 2>/dev/null &
+telemetry_pid=$!
+cleanup() {
+  kill "$telemetry_pid" 2>/dev/null || true
+  wait "$telemetry_pid" 2>/dev/null || true
+}
+trap cleanup EXIT INT TERM
+
+pytest ./fastvideo/tests/performance -vs
+pytest_rc=$?
+compare_rc=0
+if [ "$pytest_rc" -eq 0 ] || [ "$PERF_UPLOAD_POLICY" = always ]; then
+  PERF_PYTEST_RC=$pytest_rc python ./fastvideo/tests/performance/compare_baseline.py
+  compare_rc=$?
+fi
+python ./fastvideo/tests/performance/dashboard.py || true
+cp -f fastvideo/tests/performance/results/*.json "$PERF_REPORTS_DIR/" 2>/dev/null || true
+
+echo "--- GPU telemetry (clocks.sm vs clocks.max.sm reveals capped hosts) ---"
+cat "$PERF_REPORTS_DIR/gpu_telemetry.csv" || true
+
+final_rc=$pytest_rc
+if [ "$final_rc" -eq 0 ]; then
+  final_rc=$compare_rc
+fi
+exit "$final_rc"

.buildkite/scripts/lanes/self_forcing.sh

@@ -0,0 +1,6 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the self-forcing lane.
+set -euo pipefail
+
+export WANDB_MODE=offline
+exec pytest ./fastvideo/tests/training/self-forcing/test_self_forcing.py -vs

.buildkite/scripts/lanes/ssim.sh

@@ -0,0 +1,40 @@
+#!/usr/bin/env bash
+# Canonical four-GPU SSIM lane for the Slinky Slurm worker.
+set -euo pipefail
+
+args=()
+if [ "${FASTVIDEO_SSIM_BOOTSTRAP_MODE:-0}" = 1 ]; then
+  args+=(--bootstrap-mode)
+fi
+selected=${FASTVIDEO_SSIM_TEST_FILES-}
+if [ -z "$selected" ]; then
+  if [ "${TEST_SCOPE:-}" = merge ]; then
+    echo "Missing FASTVIDEO_SSIM_TEST_FILES for merge scope" >&2
+    exit 2
+  fi
+  selected=all
+fi
+if [ "$selected" != all ]; then
+  [[ $selected =~ ^test_[a-z0-9_]+\.py(,test_[a-z0-9_]+\.py)*$ ]] || {
+    echo "Invalid FASTVIDEO_SSIM_TEST_FILES selection" >&2
+    exit 2
+  }
+  IFS=, read -r -a ssim_files <<< "$selected"
+  for ssim_file in "${ssim_files[@]}"; do
+    args+=(--test-file "$ssim_file")
+  done
+fi
+
+# MoGe's utils3d dependency builds glcontext from source on ARM64. The current
+# runner image predates the baked-in X11 headers below, so keep this guarded
+# bootstrap until every deployed image digest contains libx11-dev.
+if [ ! -f /usr/include/X11/Xlib.h ]; then
+  apt-get -o Acquire::Retries=5 update
+  apt-get -o Acquire::Retries=5 install -y --no-install-recommends libx11-dev
+  rm -rf /var/lib/apt/lists/*
+fi
+
+uv pip install git+https://github.com/microsoft/MoGe.git
+uv pip install k_diffusion einops_exts alias_free_torch torchsde
+
+exec python fastvideo/tests/ssim/ci_runner.py "${args[@]}"

.buildkite/scripts/lanes/train_framework.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the modular training-framework lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/train/models ./fastvideo/tests/train/methods -vs

.buildkite/scripts/lanes/training.sh

@@ -0,0 +1,6 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the legacy vanilla-training lane.
+set -euo pipefail
+
+export WANDB_MODE=offline
+exec pytest ./fastvideo/tests/training/Vanilla -srP

.buildkite/scripts/lanes/training_lora.sh

@@ -0,0 +1,6 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the legacy LoRA-training lane.
+set -euo pipefail
+
+export WANDB_MODE=offline
+exec pytest ./fastvideo/tests/training/lora/test_lora_training.py -srP

.buildkite/scripts/lanes/training_vsa.sh

@@ -0,0 +1,6 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the legacy VSA-training lane.
+set -euo pipefail
+
+export WANDB_MODE=offline
+exec pytest ./fastvideo/tests/training/VSA -srP

.buildkite/scripts/lanes/transformer.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the transformer lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/transformers -vs

.buildkite/scripts/lanes/vae.sh

@@ -0,0 +1,5 @@
+#!/usr/bin/env bash
+# Canonical Slurm CI selection for the VAE lane.
+set -euo pipefail
+
+exec pytest ./fastvideo/tests/vaes -vs

.buildkite/scripts/pr_test.sh

@@ -1,6 +1,19 @@
 #!/bin/bash
 set -uo pipefail
 
+# DORMANT ROLLBACK ONLY. Active CI is Slurm-only and pipeline.yml never calls
+# this launcher. Refuse every Buildkite invocation even if a stale step or
+# operator typo reaches this file; local rollback experiments require an
+# explicit opt-in.
+if [ -n "${BUILDKITE:-}" ]; then
+  echo "Legacy Modal CI is disabled; use the Slinky Slurm runner." >&2
+  exit 2
+fi
+if [ "${FASTVIDEO_ENABLE_LEGACY_MODAL_CI:-0}" != 1 ]; then
+  echo "Legacy Modal CI is dormant. Set FASTVIDEO_ENABLE_LEGACY_MODAL_CI=1 only for a manual rollback test." >&2
+  exit 2
+fi
+
 log() {
   echo "[$(date '+%Y-%m-%d %H:%M:%S')] $1"
 }

.buildkite/scripts/unit_test.sh

@@ -7,6 +7,9 @@ exec pytest \
   ./fastvideo/tests/dataset/ \
   ./fastvideo/tests/workflow/ \
   ./fastvideo/tests/entrypoints/ \
+  ./fastvideo/tests/loader/ \
+  ./fastvideo/tests/pipelines/ \
+  ./fastvideo/tests/platforms/ \
   ./fastvideo/tests/train/ \
   ./fastvideo/tests/stages/ \
   ./fastvideo/tests/ops/ \

.github/PULL_REQUEST_TEMPLATE.md

@@ -8,10 +8,10 @@ PR TITLE: Must start with a type tag, e.g.:
 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
+  3. A path-aware merge gate runs only relevant integration tests → auto-merge on success
 
 ON-DEMAND TESTING (write access required):
-  /test full       — Full Test Suite        /test ssim        — SSIM regression
+  /test full       — Explicit all-lane run  /test ssim        — Full 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

.github/scripts/gate_full_suite.sh

@@ -1,14 +1,14 @@
 #!/usr/bin/env bash
-# Gate the expensive Buildkite full suite on the cheap GitHub checks.
+# Gate the path-aware Buildkite merge plan on the cheap GitHub checks.
 #
 # Polls the workflow runs for the PR head commit and only exits 0 once the
 # watched cheap workflows (pre-commit, docs build) have succeeded, so the
-# 'ready' label cannot burn ~20 GPU lanes on a head that a cheap check has
-# already doomed.
+# 'ready' label cannot burn path-selected GPU lanes on a head that a cheap
+# check has already doomed.
 #
 # Semantics:
 #   - watched run completed with a bad conclusion  -> exit 1 (fail CLOSED:
-#     no full suite; the next push re-arms via the 'synchronize' trigger)
+#     no merge gate; the next push re-arms via the 'synchronize' trigger)
 #   - watched run cancelled                         -> still pending: the docs
 #     workflow's repo-global 'pages' concurrency group cancels runs superseded
 #     by unrelated pushes, so 'cancelled' is not a verdict on this PR
@@ -29,7 +29,7 @@ set -euo pipefail
 : "${PR_NUMBER:?PR_NUMBER (pull request number) is required}"
 : "${GITHUB_REPOSITORY:?GITHUB_REPOSITORY is required}"
 
-# Workflow-level `name:` values that must be green before the full suite
+# Workflow-level `name:` values that must be green before the merge gate
 # may start. "Deploy Documentation" is path-filtered on PRs, so its run may
 # legitimately never exist; pre-commit always runs, so it must appear.
 WATCHED_NAMES='["pre-commit", "Deploy Documentation"]'
@@ -56,7 +56,7 @@ recheck_ready_label() {
   if pr_json=$(gh_api "repos/${GITHUB_REPOSITORY}/pulls/${PR_NUMBER}" 2>/dev/null); then
     if ! jq -e '[.labels[]?.name] | index("ready")' <<<"$pr_json" >/dev/null 2>&1; then
       echo "::error::PR #${PR_NUMBER} no longer has the 'ready' label —" \
-        "NOT triggering the Buildkite full suite. Re-add the label to re-arm."
+        "NOT triggering the Buildkite merge gate. Re-add the label to re-arm."
       exit 1
     fi
   else
@@ -84,7 +84,7 @@ while true; do
         | map(.name) | join(", ")' <<<"$state")
     if [ -n "$failed" ]; then
       echo "::error::Cheap check(s) failed on ${PR_SHA}: ${failed}." \
-        "NOT triggering the Buildkite full suite. Push a fix (the 'ready'" \
+        "NOT triggering the Buildkite merge gate. Push a fix (the 'ready'" \
         "label re-arms on every push), or re-run the failed check and then" \
         "re-run this workflow."
       exit 1
@@ -97,7 +97,7 @@ while true; do
     if [ "$pending" -eq 0 ]; then
       if [ -z "$missing" ]; then
         recheck_ready_label
-        echo "All watched cheap checks are green — full suite may proceed."
+        echo "All watched cheap checks are green — merge gate may proceed."
         exit 0
       fi
       case "$missing" in
@@ -119,14 +119,14 @@ while true; do
     echo "::warning::GitHub API error querying workflow runs for ${PR_SHA} (attempt ${api_fails}/3)."
     if [ "$api_fails" -ge 3 ]; then
       recheck_ready_label
-      echo "::warning::FAILING OPEN: cannot query GitHub check status — triggering the full suite WITHOUT the cheap-check gate."
+      echo "::warning::FAILING OPEN: cannot query GitHub check status — triggering the merge gate WITHOUT the cheap-check gate."
       exit 0
     fi
   fi
 
   if [ "$elapsed" -ge "$MAX_WAIT_SECS" ]; then
     recheck_ready_label
-    echo "::warning::FAILING OPEN: watched checks still pending after $(( MAX_WAIT_SECS / 60 )) min${missing:+ (never appeared: ${missing})} — triggering the full suite anyway."
+    echo "::warning::FAILING OPEN: watched checks still pending after $(( MAX_WAIT_SECS / 60 )) min${missing:+ (never appeared: ${missing})} — triggering the merge gate anyway."
     exit 0
   fi
   sleep "$POLL_SECS"

.github/scripts/plan_merge_ci.py

@@ -0,0 +1,570 @@
+#!/usr/bin/env python3
+"""Select the additive GPU integration lanes needed by a PR diff.
+
+Fastcheck is the universal six-lane baseline and is intentionally not repeated
+here.  This planner selects only the more expensive merge-gate lanes.  Unknown
+source/build paths fail closed to the complete integration set, while explicit
+documentation and repository-metadata paths require no additional GPU work.
+"""
+
+from __future__ import annotations
+
+import argparse
+import fnmatch
+import re
+from dataclasses import dataclass, field
+from pathlib import Path
+from typing import TextIO
+
+MERGE_LANES = (
+    "golden-gate",
+    "ssim",
+    "lora-inference",
+    "lora-extraction",
+    "training",
+    "distillation",
+    "self-forcing",
+    "lora-training",
+    "training-vsa",
+    "inference-vmoba",
+    "performance",
+    "api-server",
+    "train-framework",
+    "eval",
+)
+
+LANE_SCRIPT_TO_KEY = {
+    "api_server.sh": "api-server",
+    "distillation_dmd.sh": "distillation",
+    "eval.sh": "eval",
+    "golden_gate.sh": "golden-gate",
+    "inference_lora.sh": "lora-inference",
+    "inference_vmoba.sh": "inference-vmoba",
+    "lora_extraction.sh": "lora-extraction",
+    "performance.sh": "performance",
+    "self_forcing.sh": "self-forcing",
+    "ssim.sh": "ssim",
+    "train_framework.sh": "train-framework",
+    "training.sh": "training",
+    "training_lora.sh": "lora-training",
+    "training_vsa.sh": "training-vsa",
+}
+
+FASTCHECK_LANE_SCRIPTS = {
+    "dreamverse.sh",
+    "encoder.sh",
+    "kernel_tests.sh",
+    "transformer.sh",
+    "vae.sh",
+}
+
+LEGACY_TRAINING_LANES = (
+    "training",
+    "distillation",
+    "self-forcing",
+    "lora-training",
+    "training-vsa",
+)
+
+ALL_TRAINING_LANES = (*LEGACY_TRAINING_LANES, "train-framework")
+
+SSIM_SMOKE_TESTS = (
+    "test_flux_t2i_similarity.py",
+    "test_wan_t2v_similarity.py",
+)
+
+SAFE_PATTERNS = (
+    "*.md",
+    "*.rst",
+    ".agents/**",
+    ".claude/**",
+    ".codex/**",
+    ".github/ISSUE_TEMPLATE/**",
+    ".github/PULL_REQUEST_TEMPLATE.md",
+    ".github/dependabot.yml",
+    ".github/mergify.yml",
+    ".github/scripts/**",
+    ".github/workflows/**",
+    ".buildkite/scripts/pre_commit.sh",
+    ".git-blame-ignore-revs",
+    ".gitattributes",
+    ".gitignore",
+    ".pre-commit-config.yaml",
+    "AGENTS.md",
+    "CITATION.cff",
+    "CODE_OF_CONDUCT.md",
+    "CONTRIBUTING.md",
+    "LICENSE",
+    "NOTICE",
+    "__init__.py",
+    "collect_env.py",
+    "SECURITY.md",
+    "assets/**",
+    "comfyui/**",
+    "docs/**",
+    "examples/**",
+    "mkdocs.yml",
+    "requirements-mkdocs.in",
+    "requirements-mkdocs.txt",
+    "scripts/**",
+    "tests/__init__.py",
+    "tests/local_tests/**",
+)
+
+ALL_IMPACT_PATTERNS = (
+    ".buildkite/pipeline.yml",
+    "docker/**",
+    "pyproject.toml",
+    "requirements*.txt",
+    "setup.cfg",
+    "setup.py",
+    "uv.lock",
+)
+
+
+@dataclass(frozen=True)
+class FamilyCoverage:
+    pattern: re.Pattern[str]
+    golden_tests: tuple[str, ...]
+    ssim_tests: tuple[str, ...]
+
+
+FAMILY_COVERAGE = (
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])dreamx(_world)?([/_.-]|$)"),
+        ("test_dreamx.py", ),
+        ("test_dreamx_world_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])flux[_-]?2([/_.-]|$)"),
+        ("test_flux2_klein.py", ),
+        ("test_flux2_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])flux(?![_-]?2)([/_.-]|$)"),
+        ("test_flux.py", ),
+        ("test_flux_t2i_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])(hunyuan)?gamecraft([/_.-]|$)"),
+        ("test_gamecraft.py", ),
+        ("test_gamecraft_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])gen3c([/_.-]|$)"),
+        ("test_gen3c.py", ),
+        ("test_gen3c_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])glm[_-]?image([/_.-]|$)"),
+        ("test_glm_image.py", ),
+        ("test_glm_image_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])kandinsky[_-]?5([/_.-]|$)"),
+        ("test_kandinsky5.py", ),
+        ("test_kandinsky5_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])lingbot([a-z0-9_-]*)([/_.-]|$)"),
+        ("test_lingbot.py", ),
+        ("test_lingbot_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])longcat([/_.-]|$)"),
+        ("test_longcat.py", ),
+        ("test_longcat_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])ltx[_-]?2([/_.-]|$)"),
+        ("test_ltx2.py", ),
+        ("test_ltx2_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])matrixgame[_-]?2([/_.-]|$)"),
+        ("test_matrixgame.py", ),
+        ("test_matrixgame2_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])matrixgame[_-]?3([/_.-]|$)"),
+        ("test_matrixgame.py", ),
+        ("test_matrixgame3_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])minimax[_-]?h3([/_.-]|$)"),
+        ("test_minimax_h3_t2v.py", ),
+        ("test_minimax_h3_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])sd[_-]?3([._-]?5)?([/_.-]|$)"),
+        ("test_sd35.py", ),
+        ("test_sd35_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])stable[_-]?audio([/_.-]|$)"),
+        ("test_stable_audio.py", ),
+        ("test_stable_audio_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])turbo(diffusion)?([/_.-]|$)"),
+        (),
+        ("test_turbodiffusion_similarity.py", ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])wan(video)?([/_.-]|$)"),
+        ("test_wan_t2v.py", ),
+        (
+            "test_causal_similarity.py",
+            "test_wan_i2v_similarity.py",
+            "test_wan_t2v_similarity.py",
+        ),
+    ),
+    FamilyCoverage(
+        re.compile(r"(^|[/_.-])z[_-]?image([/_.-]|$)"),
+        ("test_zimage.py", ),
+        ("test_zimage_similarity.py", ),
+    ),
+)
+
+
+@dataclass
+class MergePlan:
+    lanes: set[str] = field(default_factory=set)
+    golden_tests: set[str] = field(default_factory=set)
+    ssim_tests: set[str] = field(default_factory=set)
+    golden_all: bool = False
+    ssim_all: bool = False
+    reasons: list[str] = field(default_factory=list)
+
+    def add_lanes(self, *lanes: str, reason: str) -> None:
+        unknown = set(lanes) - set(MERGE_LANES)
+        if unknown:
+            raise ValueError(f"Unknown merge lanes: {sorted(unknown)}")
+        self.lanes.update(lanes)
+        self.reasons.append(reason)
+
+    def add_golden(self, tests: tuple[str, ...], reason: str) -> None:
+        self.add_lanes("golden-gate", reason=reason)
+        self.golden_tests.update(tests)
+
+    def add_ssim(self, tests: tuple[str, ...], reason: str) -> None:
+        self.add_lanes("ssim", reason=reason)
+        self.ssim_tests.update(tests)
+
+    def require_all(self, reason: str) -> None:
+        self.lanes.update(MERGE_LANES)
+        self.golden_all = True
+        self.ssim_all = True
+        self.reasons.append(reason)
+
+    def ordered_lanes(self) -> tuple[str, ...]:
+        return tuple(lane for lane in MERGE_LANES if lane in self.lanes)
+
+    def encoded_lanes(self) -> str:
+        lanes = self.ordered_lanes()
+        return "," + ",".join(lanes or ("none", )) + ","
+
+    def encoded_golden_tests(self) -> str:
+        if "golden-gate" not in self.lanes:
+            return "none"
+        if self.golden_all or not self.golden_tests:
+            return "all"
+        return ",".join(sorted(self.golden_tests))
+
+    def encoded_ssim_tests(self) -> str:
+        if "ssim" not in self.lanes:
+            return "none"
+        if self.ssim_all or not self.ssim_tests:
+            return "all"
+        return ",".join(sorted(self.ssim_tests))
+
+
+def _matches_any(path: str, patterns: tuple[str, ...]) -> bool:
+    return any(fnmatch.fnmatchcase(path, pattern) for pattern in patterns)
+
+
+def _family_coverage(path: str) -> tuple[set[str], set[str]]:
+    normalized = path.lower()
+    golden: set[str] = set()
+    ssim: set[str] = set()
+    for family in FAMILY_COVERAGE:
+        if family.pattern.search(normalized):
+            golden.update(family.golden_tests)
+            ssim.update(family.ssim_tests)
+    return golden, ssim
+
+
+def _select_output_coverage(plan: MergePlan, path: str) -> None:
+    golden, ssim = _family_coverage(path)
+    if golden:
+        plan.add_golden(tuple(sorted(golden)), reason=f"model-family golden coverage: {path}")
+    else:
+        plan.golden_all = True
+        plan.add_lanes("golden-gate", reason=f"shared output golden coverage: {path}")
+    if ssim:
+        plan.add_ssim(tuple(sorted(ssim)), reason=f"model-family SSIM coverage: {path}")
+    else:
+        plan.add_ssim(SSIM_SMOKE_TESTS, reason=f"shared output SSIM smoke coverage: {path}")
+
+
+def classify_paths(paths: list[str]) -> MergePlan:
+    plan = MergePlan()
+    normalized_paths: list[str] = []
+    for raw_path in paths:
+        path = raw_path.strip()
+        while path.startswith("./"):
+            path = path[2:]
+        if path:
+            normalized_paths.append(path)
+    normalized_paths = sorted(set(normalized_paths))
+    if not normalized_paths:
+        plan.require_all("changed-file list was empty; failing closed")
+        return plan
+
+    for path in normalized_paths:
+        if path == "__FASTVIDEO_CI_PLAN_ALL__":
+            plan.require_all("changed-file API failed; failing closed")
+            continue
+
+        if path in {"requirements-mkdocs.in", "requirements-mkdocs.txt"}:
+            plan.reasons.append(f"documentation dependencies need no GPU integration: {path}")
+            continue
+
+        if _matches_any(path, ALL_IMPACT_PATTERNS):
+            plan.require_all(f"cross-cutting build/runtime surface: {path}")
+            continue
+
+        lane_script_prefix = ".buildkite/scripts/lanes/"
+        if path.startswith(lane_script_prefix):
+            script_name = Path(path).name
+            lane = LANE_SCRIPT_TO_KEY.get(script_name)
+            if lane is None:
+                if script_name in FASTCHECK_LANE_SCRIPTS:
+                    plan.reasons.append(f"covered by automatic Fastcheck lane: {path}")
+                else:
+                    plan.require_all(f"unknown lane script: {path}")
+            elif lane == "golden-gate":
+                plan.golden_all = True
+                plan.add_lanes(lane, reason=f"golden lane implementation: {path}")
+            elif lane == "ssim":
+                plan.ssim_all = True
+                plan.add_lanes(lane, reason=f"SSIM lane implementation: {path}")
+            else:
+                plan.add_lanes(lane, reason=f"lane implementation: {path}")
+            continue
+
+        if path.startswith("fastvideo/tests/golden_gate/"):
+            name = Path(path).name
+            if name.startswith("test_") and name.endswith(".py"):
+                plan.add_golden((name, ), reason=f"changed golden test: {path}")
+            elif name in {"AGENTS.md", "README.md"}:
+                plan.reasons.append(f"golden documentation only: {path}")
+            else:
+                plan.golden_all = True
+                plan.add_lanes("golden-gate", reason=f"shared golden harness/reference: {path}")
+            continue
+
+        if path.startswith("fastvideo/tests/ssim/"):
+            name = Path(path).name
+            if name.startswith("test_") and name.endswith(".py"):
+                plan.add_ssim((name, ), reason=f"changed SSIM test: {path}")
+            elif path.endswith((".py", ".json", ".pt", ".png", ".mp4")):
+                plan.ssim_all = True
+                plan.add_lanes("ssim", reason=f"shared SSIM harness/reference: {path}")
+            continue
+
+        if path.startswith("fastvideo/tests/performance/") or path.startswith(".buildkite/performance-benchmarks/"):
+            plan.add_lanes("performance", reason=f"performance coverage: {path}")
+            continue
+        if path.startswith(("fastvideo/performance/", "fastvideo/performance_dashboard/",
+                            "apps/performance_dashboard/")):
+            plan.add_lanes("performance", reason=f"performance implementation: {path}")
+            continue
+        if path.startswith("fastvideo/benchmarks/"):
+            if "/mlx_" in path or Path(path).name.startswith("mlx_"):
+                plan.reasons.append(f"covered by the path-filtered macOS MLX workflow: {path}")
+            else:
+                plan.add_lanes("performance", reason=f"benchmark implementation: {path}")
+            continue
+        if path.startswith("fastvideo/tests/eval/") or path.startswith("fastvideo/eval/"):
+            plan.add_lanes("eval", reason=f"evaluation coverage: {path}")
+            continue
+        if path.startswith("fastvideo/third_party/eval/"):
+            plan.add_lanes("eval", reason=f"vendored evaluation implementation: {path}")
+            continue
+        if path.startswith("fastvideo/tests/lora_extraction/") or path.startswith("scripts/lora_extraction/"):
+            plan.add_lanes("lora-extraction", reason=f"LoRA extraction coverage: {path}")
+            continue
+        if path.startswith("fastvideo/tests/inference/lora/"):
+            plan.add_lanes("lora-inference", reason=f"LoRA inference coverage: {path}")
+            continue
+        if path.startswith("fastvideo/tests/inference/vmoba/"):
+            plan.add_lanes("inference-vmoba", reason=f"VMoBA inference coverage: {path}")
+            continue
+        if path.startswith(("fastvideo/dataset/", "fastvideo/workflow/", "fastvideo/pipelines/preprocess/",
+                            "fastvideo/pipelines/training/")):
+            plan.add_lanes(*ALL_TRAINING_LANES, reason=f"shared data/training input surface: {path}")
+            continue
+        if path.startswith("fastvideo/tests/train/") or path.startswith("fastvideo/train/"):
+            plan.add_lanes("train-framework", reason=f"modular training coverage: {path}")
+            continue
+
+        if path.startswith("fastvideo/tests/training/"):
+            lowered = path.lower()
+            if "/vanilla/" in lowered:
+                plan.add_lanes("training", reason=f"vanilla training coverage: {path}")
+            elif "/distill/" in lowered:
+                plan.add_lanes("distillation", reason=f"distillation coverage: {path}")
+            elif "/self-forcing/" in lowered:
+                plan.add_lanes("self-forcing", reason=f"self-forcing coverage: {path}")
+            elif "/lora/" in lowered:
+                plan.add_lanes("lora-training", reason=f"LoRA training coverage: {path}")
+            elif "/vsa/" in lowered:
+                plan.add_lanes("training-vsa", reason=f"VSA training coverage: {path}")
+            else:
+                plan.add_lanes(*LEGACY_TRAINING_LANES, reason=f"shared legacy training coverage: {path}")
+            continue
+
+        if path.startswith("fastvideo/training/"):
+            lowered = path.lower()
+            if "self_forcing" in lowered:
+                plan.add_lanes("self-forcing", reason=f"self-forcing implementation: {path}")
+            elif "distill" in lowered:
+                plan.add_lanes("distillation", reason=f"distillation implementation: {path}")
+            elif "lora" in lowered:
+                plan.add_lanes("lora-training", reason=f"LoRA training implementation: {path}")
+            else:
+                plan.add_lanes(*LEGACY_TRAINING_LANES, reason=f"shared legacy training implementation: {path}")
+            continue
+
+        lowered = path.lower()
+        if "vmoba" in lowered and path.startswith(("fastvideo/", ".buildkite/")):
+            plan.add_lanes("inference-vmoba", reason=f"VMoBA implementation: {path}")
+            plan.add_golden(("test_wan_t2v.py", ), reason=f"VMoBA end-to-end coverage: {path}")
+            continue
+        if "lora" in lowered and path.startswith("fastvideo/"):
+            plan.add_lanes(
+                "lora-inference",
+                "lora-extraction",
+                "lora-training",
+                reason=f"shared LoRA implementation: {path}",
+            )
+            _select_output_coverage(plan, path)
+            continue
+
+        if path.startswith("fastvideo/entrypoints/") or path.startswith("fastvideo/api/"):
+            plan.add_lanes("api-server", reason=f"API/entrypoint integration: {path}")
+            if "openai" not in lowered and "/cli/" not in lowered:
+                _select_output_coverage(plan, path)
+            continue
+        if path.startswith("fastvideo/worker/"):
+            plan.add_lanes("api-server", reason=f"worker/API integration: {path}")
+            _select_output_coverage(plan, path)
+            continue
+        if path.startswith("fastvideo/distributed/"):
+            plan.add_lanes(
+                "training",
+                "train-framework",
+                reason=f"distributed runtime integration: {path}",
+            )
+            _select_output_coverage(plan, path)
+            continue
+        if path.startswith(("fastvideo/hooks/", "fastvideo/platforms/", "fastvideo/third_party/")):
+            _select_output_coverage(plan, path)
+            continue
+        if path.startswith(("fastvideo/models/", "fastvideo/pipelines/", "fastvideo/configs/",
+                            "fastvideo/layers/", "fastvideo/attention/")):
+            _select_output_coverage(plan, path)
+            continue
+        if path in {
+                "fastvideo/fastvideo_args.py",
+                "fastvideo/forward_context.py",
+                "fastvideo/image_processor.py",
+                "fastvideo/registry.py",
+                "fastvideo/utils.py",
+        }:
+            _select_output_coverage(plan, path)
+            continue
+        if path.startswith("fastvideo/mlx_runtime/"):
+            plan.reasons.append(f"covered by the path-filtered macOS MLX workflow: {path}")
+            continue
+        if path.startswith("fastvideo/logging_utils/") or path in {
+                "fastvideo/__init__.py",
+                "fastvideo/envs.py",
+                "fastvideo/logger.py",
+                "fastvideo/profiler.py",
+                "fastvideo/version.py",
+        }:
+            plan.reasons.append(f"covered by automatic Fastcheck: {path}")
+            continue
+        if path.startswith(("fastvideo-kernel/", "csrc/")):
+            plan.add_golden(("test_wan_t2v.py", ), reason=f"kernel integration smoke: {path}")
+            plan.add_ssim(("test_wan_t2v_similarity.py", ), reason=f"kernel numerical smoke: {path}")
+            continue
+
+        if path.startswith("apps/dreamverse/"):
+            # DreamVerse is already one of the six automatic Fastcheck lanes.
+            plan.reasons.append(f"covered by automatic DreamVerse Fastcheck: {path}")
+            continue
+        if path.startswith("fastvideo/tests/"):
+            # The automatic unit/component Fastcheck lanes own the remaining
+            # package tests. Domain-specific expensive test roots were handled
+            # above.
+            plan.reasons.append(f"covered by automatic Fastcheck: {path}")
+            continue
+        if path in {".buildkite/scripts/unit_test.sh", ".buildkite/scripts/pr_test.sh"}:
+            plan.reasons.append(f"covered by automatic unit Fastcheck: {path}")
+            continue
+        if _matches_any(path, SAFE_PATTERNS):
+            plan.reasons.append(f"no additional GPU integration needed: {path}")
+            continue
+
+        plan.require_all(f"unclassified path; failing closed: {path}")
+
+    return plan
+
+
+def _write_github_output(output: TextIO, plan: MergePlan) -> None:
+    output.write(f"merge_test_plan={plan.encoded_lanes()}\n")
+    output.write(f"merge_golden_tests={plan.encoded_golden_tests()}\n")
+    output.write(f"merge_ssim_tests={plan.encoded_ssim_tests()}\n")
+    output.write(f"merge_plan_label={','.join(plan.ordered_lanes()) or 'none'}\n")
+
+
+def _write_summary(output: TextIO, plan: MergePlan) -> None:
+    output.write("## Change-aware merge test plan\n\n")
+    output.write("| Selection | Value |\n|---|---|\n")
+    output.write(f"| Additional Slurm lanes | `{','.join(plan.ordered_lanes()) or 'none'}` |\n")
+    output.write(f"| Golden tests | `{plan.encoded_golden_tests()}` |\n")
+    output.write(f"| SSIM tests | `{plan.encoded_ssim_tests()}` |\n\n")
+    output.write("Fastcheck remains the universal six-lane baseline.\n")
+
+
+def parse_args() -> argparse.Namespace:
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--paths-file", type=Path, required=True)
+    parser.add_argument("--github-output", type=Path)
+    parser.add_argument("--summary-file", type=Path)
+    return parser.parse_args()
+
+
+def main() -> int:
+    args = parse_args()
+    paths = args.paths_file.read_text(encoding="utf-8").splitlines()
+    plan = classify_paths(paths)
+    print(f"MERGE_TEST_PLAN={plan.encoded_lanes()}")
+    print(f"MERGE_GOLDEN_TESTS={plan.encoded_golden_tests()}")
+    print(f"MERGE_SSIM_TESTS={plan.encoded_ssim_tests()}")
+    for reason in plan.reasons:
+        print(f"- {reason}")
+    if args.github_output:
+        with args.github_output.open("a", encoding="utf-8") as output:
+            _write_github_output(output, plan)
+    if args.summary_file:
+        with args.summary_file.open("a", encoding="utf-8") as output:
+            _write_summary(output, plan)
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

.github/scripts/test_gate_full_suite.sh

@@ -53,7 +53,7 @@ PC_PENDING='{"name": "pre-commit", "id": 1, "status": "in_progress", "conclusion
 DOCS_OK='{"name": "Deploy Documentation", "id": 2, "status": "completed", "conclusion": "success"}'
 DOCS_BAD='{"name": "Deploy Documentation", "id": 2, "status": "completed", "conclusion": "failure"}'
 DOCS_CANCELLED='{"name": "Deploy Documentation", "id": 2, "status": "completed", "conclusion": "cancelled"}'
-OTHER='{"name": "Trigger Full Suite", "id": 3, "status": "in_progress", "conclusion": null}'
+OTHER='{"name": "Trigger Merge Gate", "id": 3, "status": "in_progress", "conclusion": null}'
 NULL_NAME='{"name": null, "id": 4, "status": "completed", "conclusion": "failure"}'
 PC_OK_RERUN='{"name": "pre-commit", "id": 5, "status": "completed", "conclusion": "success"}'
 

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

@@ -26,29 +26,48 @@ jobs:
               per_page: 100,
             });
 
-            const bkStatuses = data.statuses.filter(
-              s => s.context.startsWith('buildkite/ci/')
-            );
-
-            const FASTCHECK_PREFIX = 'buildkite/ci/microscope-';
+            // Buildkite derives the GitHub context prefix from the label emoji.
+            // Keep hard Full Suite lanes in test-tube/bar-chart namespaces and
+            // Fastcheck lanes in microscope so targeted reruns cannot clear the
+            // wrong aggregate status. Automatic PR jobs use pr-fastcheck while
+            // slash-command and Full Suite jobs use ci; normalize the suffix
+            // and keep the newest status for each logical lane.
+            const FASTCHECK_PREFIXES = [
+              'buildkite/pr-fastcheck/microscope-',
+              '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))
-            );
+            function newestByLane(prefixes) {
+              const statuses = new Map();
+              for (const status of data.statuses) {
+                const prefix = prefixes.find(p => status.context.startsWith(p));
+                if (!prefix) continue;
+                const lane = status.context.slice(prefix.length);
+                const previous = statuses.get(lane);
+                if (!previous || Date.parse(status.updated_at) > Date.parse(previous.updated_at)) {
+                  statuses.set(lane, status);
+                }
+              }
+              return statuses;
+            }
 
-            if (
-              fastcheck.length > 0
-              && fastcheck.every(s => s.state === 'success')
-            ) {
+            const fastcheck = newestByLane(FASTCHECK_PREFIXES);
+            const fullSuiteOnly = newestByLane(FULL_SUITE_PREFIXES);
+            const fastcheckPassed =
+              fastcheck.size === 6
+              && [...fastcheck.values()].every(s => s.state === 'success');
+            const fullSuitePassed =
+              fastcheckPassed
+              && fullSuiteOnly.size === 14
+              && [...fullSuiteOnly.values()].every(s => s.state === 'success');
+
+            if (fastcheckPassed) {
               core.info(
-                `All ${fastcheck.length} fastcheck tests passed — updating fastcheck-passed`
+                `All ${fastcheck.size} fastcheck tests passed — updating fastcheck-passed`
               );
               await github.rest.repos.createCommitStatus({
                 owner: context.repo.owner,
@@ -56,17 +75,13 @@ jobs:
                 sha,
                 state: 'success',
                 context: 'fastcheck-passed',
-                description:
-                  `All ${fastcheck.length} fastcheck tests passed`,
+                description: `All ${fastcheck.size} fastcheck tests passed`,
               });
             }
 
-            if (
-              fullSuite.length > 0
-              && fullSuite.every(s => s.state === 'success')
-            ) {
+            if (fullSuitePassed) {
               core.info(
-                `All ${fullSuite.length} full suite tests passed — updating full-suite-passed`
+                'All 20 full suite tests passed — updating full-suite-passed'
               );
               await github.rest.repos.createCommitStatus({
                 owner: context.repo.owner,
@@ -74,7 +89,6 @@ jobs:
                 sha,
                 state: 'success',
                 context: 'full-suite-passed',
-                description:
-                  `All ${fullSuite.length} full suite tests passed`,
+                description: 'All 20 full suite tests passed',
               });
             }

.github/workflows/ci-macos-mlx.yml

@@ -78,6 +78,7 @@ jobs:
             fastvideo/tests/mlx/test_mlx_dit_parity.py \
             fastvideo/tests/mlx/test_mlx_compile_parity.py \
             fastvideo/tests/mlx/test_mlx_checkpoint.py \
+            fastvideo/tests/mlx/test_mlx_checkpoint_compat.py \
             fastvideo/tests/mlx/test_mlx_fastwan_benchmark.py \
             fastvideo/tests/mlx/test_taehv_decode.py \
             fastvideo/tests/mlx/test_frame_upsample.py \
@@ -134,6 +135,7 @@ jobs:
             fastvideo/tests/mlx/test_mlx_dit_parity.py \
             fastvideo/tests/mlx/test_mlx_compile_parity.py \
             fastvideo/tests/mlx/test_mlx_checkpoint.py \
+            fastvideo/tests/mlx/test_mlx_checkpoint_compat.py \
             fastvideo/tests/mlx/test_mlx_fastwan_benchmark.py \
             fastvideo/tests/mlx/test_taehv_decode.py \
             fastvideo/tests/mlx/test_frame_upsample.py \

.github/workflows/ci-scheduled-ssim.yml

@@ -0,0 +1,44 @@
+name: Scheduled Full SSIM
+
+on:
+  schedule:
+    - cron: "0 5 * * 0"
+  workflow_dispatch:
+
+permissions:
+  contents: read
+
+jobs:
+  trigger:
+    if: github.repository == 'hao-ai-lab/FastVideo'
+    runs-on: ubuntu-latest
+    steps:
+      - name: Trigger weekly full SSIM on Slinky Slurm
+        env:
+          BUILDKITE_API_TOKEN: ${{ secrets.BUILDKITE_API_TOKEN }}
+          SOURCE_SHA: ${{ github.sha }}
+          SOURCE_BRANCH: ${{ github.event.repository.default_branch }}
+          BK_ORG: ${{ vars.BUILDKITE_ORG_SLUG }}
+          BK_PIPELINE: ${{ vars.BUILDKITE_PIPELINE_SLUG }}
+        run: |
+          set -euo pipefail
+          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 "$SOURCE_SHA" \
+              --arg branch "$SOURCE_BRANCH" \
+              '{
+                commit: $commit,
+                branch: $branch,
+                message: "Weekly full SSIM on Slinky Slurm",
+                ignore_pipeline_branch_filters: true,
+                env: {
+                  TEST_SCOPE: "scheduled",
+                  FULL_SUITE: "false",
+                  TEST_TYPE: "ssim",
+                  PR_NUMBER: "false",
+                  PR_TITLE: "Scheduled full SSIM"
+                }
+              }')"

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

@@ -33,7 +33,6 @@ jobs:
             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:
@@ -48,47 +47,6 @@ jobs:
               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));
-            core.setOutput('pr_title', pr.title);
-
-      - 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 }}
-          PR_TITLE: ${{ steps.label.outputs.pr_title }}
-          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)" \
-              --arg pr_title "$PR_TITLE" \
-              --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),
-                    PR_TITLE: $pr_title
-                  }
-                }')"
 
   parse-command:
     if: >-
@@ -129,7 +87,7 @@ jobs:
           set -euo pipefail
           TEST_NAME=$(echo "$COMMENT" | grep -oP '(?<=/test\s)\S+' | head -1 || true)
 
-          VALID="encoder vae transformer kernel unit unit-ci dreamverse ssim golden-gate training lora-inference lora-training lora-extraction distillation self-forcing vsa vmoba performance api train-framework eval full fastcheck pre-commit"
+          VALID="encoder vae transformer kernel unit dreamverse ssim golden-gate training lora-inference lora-training lora-extraction distillation self-forcing vsa vmoba performance api train-framework eval unit-ci kernel-ci dreamverse-ci ssim-ci golden-gate-ci encoder-ci vae-ci transformer-ci lora-inference-ci lora-training-ci lora-extraction-ci training-ci distillation-ci self-forcing-ci vsa-ci vmoba-ci performance-ci api-ci train-framework-ci eval-ci full fastcheck pre-commit"
           if [ -z "$TEST_NAME" ] || ! echo "$VALID" | grep -qw "$TEST_NAME"; then
             echo "Unknown test: '$TEST_NAME'. Valid: $VALID"
             exit 1
@@ -138,6 +96,15 @@ jobs:
           declare -A MAP=(
             [encoder]=encoder [vae]=vae [transformer]=transformer
             [kernel]=kernel_tests [unit]=unit_test [unit-ci]=unit_test_ci
+            [kernel-ci]=kernel_tests_ci [dreamverse-ci]=dreamverse_app_ci
+            [ssim-ci]=ssim_ci [vmoba-ci]=inference_vmoba_ci
+            [golden-gate-ci]=golden_gate_ci [training-ci]=training_ci
+            [encoder-ci]=encoder_ci [vae-ci]=vae_ci [transformer-ci]=transformer_ci
+            [lora-inference-ci]=inference_lora_ci [lora-training-ci]=training_lora_ci
+            [lora-extraction-ci]=lora_extraction_ci [distillation-ci]=distillation_dmd_ci
+            [self-forcing-ci]=self_forcing_ci [vsa-ci]=training_vsa_ci
+            [performance-ci]=performance_ci [api-ci]=api_server_ci
+            [train-framework-ci]=train_framework_ci [eval-ci]=eval_ci
             [dreamverse]=dreamverse_app
             [ssim]=ssim [golden-gate]=golden_gate [training]=training
             [lora-inference]=inference_lora [lora-training]=training_lora

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

@@ -1,4 +1,4 @@
-name: Trigger Full Suite
+name: Trigger Merge Gate
 
 on:
   pull_request_target:
@@ -10,7 +10,7 @@ permissions:
   actions: read
 
 concurrency:
-  group: full-suite-${{ github.event.pull_request.number }}
+  group: merge-gate-${{ github.event.pull_request.number }}
   cancel-in-progress: false
 
 jobs:
@@ -34,29 +34,72 @@ jobs:
             });
             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.');
+            core.setOutput('changed_files', String(pr.changed_files));
+            if (!hasReady) core.info('No ready label — skipping merge-gate 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 }}
+          PR_NUMBER: ${{ github.event.pull_request.number }}
         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')
+          # Match both branch and PR number: forks can reuse the same branch name.
+          builds=$(curl -sS --get -H "Authorization: Bearer $BUILDKITE_API_TOKEN" \
+            --data-urlencode "branch=$PR_BRANCH" \
+            --data-urlencode "state=running,scheduled" \
+            "https://api.buildkite.com/v2/organizations/${{ vars.BUILDKITE_ORG_SLUG }}/pipelines/${{ vars.BUILDKITE_PIPELINE_SLUG }}/builds" \
+            | jq -r --arg pr_number "$PR_NUMBER" \
+              '.[] | select((.env.TEST_SCOPE? == "merge") and (.env.PR_NUMBER? == $pr_number)) | .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
 
-      # Checks out the BASE branch (default for pull_request_target), so PR
-      # authors cannot tamper with the gate script.
-      - name: Checkout gate script
+      # Check out the immutable BASE SHA: pull_request_target must never run a
+      # planner or gate script from the untrusted PR head.
+      - name: Checkout trusted merge planner
         if: steps.check.outputs.has_ready == 'true'
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+        with:
+          ref: ${{ github.event.pull_request.base.sha }}
+          persist-credentials: false
+
+      - name: Collect changed paths
+        if: steps.check.outputs.has_ready == 'true'
+        env:
+          GH_TOKEN: ${{ github.token }}
+          PR_NUMBER: ${{ github.event.pull_request.number }}
+          EXPECTED_CHANGED_FILES: ${{ steps.check.outputs.changed_files }}
+        run: |
+          set -euo pipefail
+          changed_json="$RUNNER_TEMP/merge-changed-files.json"
+          changed_paths="$RUNNER_TEMP/merge-changed-paths.txt"
+          if gh api --paginate --slurp \
+              "repos/${GITHUB_REPOSITORY}/pulls/${PR_NUMBER}/files?per_page=100" \
+              > "$changed_json"; then
+            observed=$(jq '[.[][] | .filename] | unique | length' "$changed_json")
+            if [ "$observed" = "$EXPECTED_CHANGED_FILES" ]; then
+              jq -r '.[][] | .filename, (.previous_filename // empty)' "$changed_json" \
+                | sort -u > "$changed_paths"
+            else
+              echo "::warning::Changed-file API returned $observed of $EXPECTED_CHANGED_FILES paths; selecting all merge lanes."
+              echo '__FASTVIDEO_CI_PLAN_ALL__' > "$changed_paths"
+            fi
+          else
+            echo "::warning::Changed-file API failed; selecting all merge lanes."
+            echo '__FASTVIDEO_CI_PLAN_ALL__' > "$changed_paths"
+          fi
+
+      - name: Select minimal merge tests
+        id: plan
+        if: steps.check.outputs.has_ready == 'true'
+        run: |
+          python3 .github/scripts/plan_merge_ci.py \
+            --paths-file "$RUNNER_TEMP/merge-changed-paths.txt" \
+            --github-output "$GITHUB_OUTPUT" \
+            --summary-file "$GITHUB_STEP_SUMMARY"
 
       - name: Wait for pre-commit and docs build
         if: steps.check.outputs.has_ready == 'true'
@@ -66,7 +109,7 @@ jobs:
           PR_NUMBER: ${{ github.event.pull_request.number }}
         run: bash .github/scripts/gate_full_suite.sh
 
-      - name: Trigger Buildkite Full Suite
+      - name: Trigger Buildkite merge gate
         if: steps.check.outputs.has_ready == 'true'
         env:
           BUILDKITE_API_TOKEN: ${{ secrets.BUILDKITE_API_TOKEN }}
@@ -76,6 +119,10 @@ jobs:
           PR_TITLE: ${{ github.event.pull_request.title }}
           BK_ORG: ${{ vars.BUILDKITE_ORG_SLUG }}
           BK_PIPELINE: ${{ vars.BUILDKITE_PIPELINE_SLUG }}
+          MERGE_TEST_PLAN: ${{ steps.plan.outputs.merge_test_plan }}
+          MERGE_GOLDEN_TESTS: ${{ steps.plan.outputs.merge_golden_tests }}
+          MERGE_SSIM_TESTS: ${{ steps.plan.outputs.merge_ssim_tests }}
+          MERGE_PLAN_LABEL: ${{ steps.plan.outputs.merge_plan_label }}
         run: |
           curl -sS --fail-with-body -X POST \
             "https://api.buildkite.com/v2/organizations/${BK_ORG}/pipelines/${BK_PIPELINE}/builds" \
@@ -84,8 +131,11 @@ jobs:
             --data-raw "$(jq -n \
               --arg commit "$PR_SHA" \
               --arg branch "$PR_BRANCH" \
-              --arg message "Full Suite for PR #${PR_NUMBER}" \
+              --arg message "Merge gate [${MERGE_PLAN_LABEL}] for PR #${PR_NUMBER}" \
               --arg pr_title "$PR_TITLE" \
+              --arg merge_test_plan "$MERGE_TEST_PLAN" \
+              --arg merge_golden_tests "$MERGE_GOLDEN_TESTS" \
+              --arg merge_ssim_tests "$MERGE_SSIM_TESTS" \
               --argjson pr_id "$PR_NUMBER" \
               '{
                 commit: $commit,
@@ -95,8 +145,11 @@ jobs:
                 pull_request_id: $pr_id,
                 pull_request_base_branch: "main",
                 env: {
-                  TEST_SCOPE: "full",
+                  TEST_SCOPE: "merge",
                   FULL_SUITE: "true",
+                  MERGE_TEST_PLAN: $merge_test_plan,
+                  MERGE_GOLDEN_TESTS: $merge_golden_tests,
+                  MERGE_SSIM_TESTS: $merge_ssim_tests,
                   PR_NUMBER: ($pr_id | tostring),
                   PR_TITLE: $pr_title
                 }

.github/workflows/community-welcome.yml

@@ -38,17 +38,17 @@ jobs:
 
             **How our CI works:**
 
-            PRs run a two-tier CI system:
+            PRs run a three-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.
+            2. **Fastcheck** — six core GPU lanes run automatically via Buildkite (~10-15 min).
+            3. **Merge gate** — a reviewer adds `ready`; changed paths select only the relevant integration, training, golden, or SSIM coverage.
 
             **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.
+            If pre-commit fails, a bot comment will explain how to fix it. Fastcheck and merge-gate results appear in the Checks section below.
 
             **Useful links:**
             - [Contributing Guide](https://hao-ai-lab.github.io/FastVideo/contributing/overview/)

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

@@ -203,8 +203,8 @@ jobs:
           docker buildx imagetools create "${TAG_ARGS[@]}" "${IMAGE_REFS[@]}"
           docker buildx imagetools inspect "${TAGS[0]}"
 
-  # The CI runner is ARM64 like DGX Spark, but its GB200 GPUs are sm_100 rather
-  # than sm_121. Publish a single-architecture variant so Slurm CI can reuse
+  # The CI runner is ARM64 like DGX Spark, but targets sm_100 rather than sm_121.
+  # Publish a single-architecture variant so the self-hosted CI runner can reuse
   # the exact prebuilt kernel instead of compiling it in every job.
   build-ci-runner-image:
     if: ${{ (github.event_name == 'push' && github.repository == 'hao-ai-lab/FastVideo') || github.event.inputs.build_ci_runner_image == 'true' }}

.github/workflows/infra-docs.yml

@@ -6,6 +6,7 @@ on:
     paths:
       - 'docs/**'
       - 'examples/**'
+      - 'scripts/inference/**'
       - 'mkdocs.yml'
       - 'requirements-mkdocs.in'
       - 'requirements-mkdocs.txt'
@@ -16,6 +17,7 @@ on:
     paths:
       - 'docs/**'
       - 'examples/**'
+      - 'scripts/inference/**'
       - 'mkdocs.yml'
       - 'requirements-mkdocs.in'
       - 'requirements-mkdocs.txt'

.github/workflows/publish-kernel.yml

@@ -62,8 +62,9 @@ jobs:
               cuda-version: '13.0.0'
               torch-cuda-short: 'cu130'
           platform:
-            # x86_64 builds the full cu126 + cu130 set (cu130 ships the consumer
-            # Blackwell sm_120a FP4 kernels).
+            # x86_64 builds the full cu126 + cu130 set. cu130 ships the
+            # data-center Blackwell sm_100a VSA and consumer sm_120a FP4
+            # kernels.
             - os: ubuntu-22.04
               arch: x86_64
               wheel-plat: manylinux_2_35_x86_64
@@ -124,7 +125,7 @@ jobs:
       - 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 apt install -y git 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
@@ -168,7 +169,8 @@ jobs:
           #     covers sm_120a; turbodiffusion covers sm_100a+sm_120a. The sm_100 FP4
           #     forward is the FA4 CuTe DSL path in the fastvideo package (PR #1221),
           #     JIT-compiled at runtime — not built into this wheel.
-          #   * x86_64 cu130 = Hopper TK + consumer Blackwell sm_120a FP4.
+          #   * x86_64 cu130 = Hopper TK + data-center Blackwell sm_100a VSA
+          #     + consumer Blackwell sm_120a FP4.
           #   * x86_64 cu126 = Hopper TK only (older drivers; CUDA < 12.8 has no FP4).
           # The per-arch split in CMakeLists pins the FP4 targets to sm_120a and builds
           # the main extension for the full arch list. CMAKE_BUILD_PARALLEL_LEVEL caps
@@ -178,7 +180,7 @@ jobs:
             export CMAKE_ARGS="${CMAKE_ARGS:-} -DFASTVIDEO_KERNEL_BUILD_TK=OFF -DFASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER=ON"
             export CMAKE_BUILD_PARALLEL_LEVEL=1
           elif [ "${{ matrix.torch-cuda.torch-cuda-short }}" = "cu130" ]; then
-            export TORCH_CUDA_ARCH_LIST="9.0a;12.0a"
+            export TORCH_CUDA_ARCH_LIST="9.0a;10.0a;12.0a"
             export CMAKE_ARGS="${CMAKE_ARGS:-} -DFASTVIDEO_KERNEL_BUILD_TK=ON -DFASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER=ON -DCMAKE_CUDA_ARCHITECTURES=90a"
             # A single FP4 TU (attn_qat_infer) can use ~8-12 GB on its own, so serialize.
             export CMAKE_BUILD_PARALLEL_LEVEL=1
@@ -194,7 +196,11 @@ jobs:
           python -m build --wheel --outdir dist
 
           # Fix the wheel to be manylinux compliant
-          uv pip install --system auditwheel
+          # Ubuntu 22.04 ships patchelf 0.14.3, while current auditwheel
+          # requires at least 0.14.5. Use the stable PyPI binary on both
+          # x86_64 and aarch64 release runners.
+          uv pip install --system auditwheel patchelf==0.17.2.4
+          patchelf --version
           # Point auditwheel at torch libs, but do not vendor them into the wheel.
           TORCH_LIB_DIR=$(python - <<'PY'
           import os
@@ -211,7 +217,8 @@ jobs:
             --exclude libtorch.so \
             --exclude libc10.so \
             --exclude libc10_cuda.so \
-            --exclude libtorch_python.so
+            --exclude libtorch_python.so \
+            --exclude libnccl.so.2
           # Move fixed wheels back to dist for upload consistency
           rm dist/*.whl
           mv fixed_dist/*.whl dist/

README.md

@@ -9,6 +9,8 @@
 **FastVideo is a unified post-training and real-time inference framework for accelerated video generation.**
 
 ## NEWS
+- `2026/08/23`: [FastH3 Preview v0.2](https://huggingface.co/FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2) is a 4-step DMD2-distilled MiniMax-H3 checkpoint that generates synchronized video and audio. Run the verified [basic FastH3 example](examples/inference/basic/basic_fasth3.py), see the [inference guide](examples/inference/basic/README.md#fasth3-preview), or have a coding agent install FastVideo with the [agent setup prompt](#install-with-an-ai-coding-agent).
+- `2026/08/19`: FastVideo now supports MLX on Apple Silicon with [FastMetal-QAD](https://huggingface.co/collections/FastVideo/fastmetal), a family of 1.3B, 5B, and 14B models optimized for Mac—follow the [Apple Silicon guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mps/) and read the [Blog](https://haoailab.com/blogs/fastmetal/).
 - `2026/06/23`: Release FastWan-QAD: 5s of Video generated in 1.8s E2E. See the [FastWan-QAD models](https://huggingface.co/FastVideo/FastWan-QAD-FP8-1.3B), [Attn-QAT training guide](https://haoailab.com/FastVideo/training/attn_qat/), and [blog](https://haoailab.com/blogs/fastwan-qad/).
 - `2026/03/17`: Release demo: Into the Dreamverse: Vibe Directing in FastVideo, check out the [Blog](https://haoailab.com/blogs/dreamverse/).
 - `2026/03/13`: Release demo: Create a 5s 1080p Video in 4.5s with FastVideo on a Single GPU, check out the [Blog](https://haoailab.com/blogs/fastvideo_realtime_1080p/).
@@ -62,9 +64,10 @@ UV_TORCH_BACKEND=cu126 uv pip install fastvideo
 Use `UV_TORCH_BACKEND=cu130` on CUDA 13. Apple silicon users should follow the
 [MPS installation guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mps/).
 
-> **On an Apple Silicon Mac?** FastVideo runs FastWan text-to-video natively
-> through an MLX runtime — a 5-second 480p clip generated locally, no cloud,
-> no discrete GPU. Install with `uv pip install -e '.[mlx]'` and follow the
+> **On an Apple Silicon Mac?** FastVideo runs FastMetal-QAD through an MLX
+> runtime. Install with `uv pip install -e '.[mlx]'`, download
+> [`FastVideo/FastMetal-1.3B-QAD`](https://huggingface.co/FastVideo/FastMetal-1.3B-QAD),
+> and follow the
 > [Apple Silicon guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mps/).
 
 Please see our [docs](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/) for more detailed installation instructions.

docker/Dockerfile

@@ -92,11 +92,26 @@ RUN apt-get update && apt-get install -y --no-install-recommends \
     ffmpeg \
     libgl1 \
     libglib2.0-0 \
+    libx11-dev \
     gcc-11 \
     g++-11 \
     clang-11 \
+    cmake \
+    pkg-config \
+    build-essential \
+    libssl-dev \
     && rm -rf /var/lib/apt/lists/*
 
+# Rust toolchain: some dependencies only ship sdists on aarch64 and need cargo
+# to build. The dormant legacy Modal image layers the identical apt set +
+# rustup on top of this image (fastvideo/tests/modal/pr_test.py); baking both
+# here keeps its manual rollback path reproducible without changing the Slurm
+# runner's package surface.
+RUN set -o pipefail && \
+    curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --default-toolchain stable --profile minimal && \
+    /root/.cargo/bin/cargo --version && /root/.cargo/bin/rustc --version
+ENV PATH=/root/.cargo/bin:${PATH}
+
 # Set up C++20 compilers for ThunderKittens
 RUN update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 100 --slave /usr/bin/g++ g++ /usr/bin/g++-11
 
@@ -173,26 +188,21 @@ RUN --mount=type=cache,target=/opt/uv/cache \
 # flash_attn/__init__.py), so FA2/varlen/bert_padding stay from the install above;
 # rmtree clears the wheel's stale cute files first to avoid an install conflict.
 # Then verify both survive so a broken overlay fails the build instead of shipping
-# an FA2-less image. x86 only: the FA4 stack (quack-kernels etc.) is unvalidated on
-# arm64 / GB10 (sm_121), so there we skip the overlay; FA4 is opt-in
-# (FASTVIDEO_FA4=1) and errors if set without the overlay, so leave it unset on
-# arm64 and the image runs FA3/FA2 as usual.
+# an FA2-less image. The pinned stack is validated on ARM64 GB200 (sm_100) as well
+# as x86; FA4 remains opt-in through FASTVIDEO_FA4=1 so lanes with FA2 baselines
+# keep their existing numerics.
 RUN --mount=type=cache,target=/opt/uv/cache \
     source $HOME/.local/bin/env && \
     source /opt/venv/bin/activate && \
-    if [ "${TARGETARCH}" = "arm64" ]; then \
-        echo "Skipping FA4 cute overlay on arm64 (FA4 stack unvalidated there; do not set FASTVIDEO_FA4)"; \
-    else \
-        python -c "import glob, shutil; [shutil.rmtree(d, ignore_errors=True) for d in glob.glob('/opt/venv/lib/python*/site-packages/flash_attn/cute')]" && \
-        uv pip install "flash-attn-4 @ git+https://github.com/Dao-AILab/flash-attention.git@${FA4_CUTE_REF}#subdirectory=flash_attn/cute" && \
-        python -c "import flash_attn; assert hasattr(flash_attn, 'flash_attn_func'), 'FA2 was clobbered by the cute overlay'; import flash_attn.cute; print('FA2 + FA4 cute OK')"; \
-    fi
+    python -c "import glob, shutil; [shutil.rmtree(d, ignore_errors=True) for d in glob.glob('/opt/venv/lib/python*/site-packages/flash_attn/cute')]" && \
+    uv pip install "flash-attn-4 @ git+https://github.com/Dao-AILab/flash-attention.git@${FA4_CUTE_REF}#subdirectory=flash_attn/cute" && \
+    python -c "import flash_attn; assert hasattr(flash_attn, 'flash_attn_func'), 'FA2 was clobbered by the cute overlay'; import flash_attn.cute; print('FA2 + FA4 cute OK')"
 
 COPY . .
 
 # Build immutable FastVideo kernel wheels for the published image. The requested
 # architecture remains installed for normal image users; amd64 images also carry
-# an SM89 artifact so the predominant L40S Modal lanes can reuse it exactly.
+# an SM89 artifact for L40S users and the dormant legacy rollback path.
 ARG FASTVIDEO_KERNEL_PREBUILT_DIR=/opt/fastvideo-kernel-prebuilt
 RUN --mount=type=cache,target=/opt/uv/cache \
     source $HOME/.local/bin/env && \

docs/assets/cookbook-recipes.json

@@ -0,0 +1,52 @@
+{
+  "recipes": [
+    {
+      "id": "fastwan21-t2v",
+      "task": "Text to video",
+      "label": "FastWan2.1 1.3B (distilled + VSA)",
+      "model": "FastVideo/FastWan2.1-T2V-1.3B-Diffusers",
+      "source": "scripts/inference/inference_wan_VSA_DMD_1_3B.yaml",
+      "command": "FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_ATTN fastvideo generate --config scripts/inference/inference_wan_VSA_DMD_1_3B.yaml"
+    },
+    {
+      "id": "wan22-t2v",
+      "task": "Text to video",
+      "label": "Wan2.2 A14B",
+      "model": "Wan-AI/Wan2.2-T2V-A14B-Diffusers",
+      "source": "examples/inference/basic/basic_wan2_2.py",
+      "command": "python examples/inference/basic/basic_wan2_2.py"
+    },
+    {
+      "id": "wan21-i2v",
+      "task": "Image to video",
+      "label": "Wan2.1 14B 480P",
+      "model": "Wan-AI/Wan2.1-I2V-14B-480P-Diffusers",
+      "source": "scripts/inference/inference_wan_i2v.yaml",
+      "command": "fastvideo generate --config scripts/inference/inference_wan_i2v.yaml"
+    },
+    {
+      "id": "turbowan22-i2v",
+      "task": "Image to video",
+      "label": "TurboWan2.2 A14B",
+      "model": "loayrashid/TurboWan2.2-I2V-A14B-Diffusers",
+      "source": "examples/inference/basic/basic_turbodiffusion_i2v.py",
+      "command": "python examples/inference/basic/basic_turbodiffusion_i2v.py"
+    },
+    {
+      "id": "wan22-ti2v",
+      "task": "Text or image to video",
+      "label": "Wan2.2 TI2V 5B",
+      "model": "Wan-AI/Wan2.2-TI2V-5B-Diffusers",
+      "source": "examples/inference/basic/basic_wan2_2_ti2v.py",
+      "command": "python examples/inference/basic/basic_wan2_2_ti2v.py"
+    },
+    {
+      "id": "matrix-game-2",
+      "task": "Interactive world",
+      "label": "Matrix Game 2.0",
+      "model": "FastVideo/Matrix-Game-2.0-Base-Distilled-Diffusers",
+      "source": "examples/inference/basic/basic_matrixgame2.py",
+      "command": "python examples/inference/basic/basic_matrixgame2.py"
+    }
+  ]
+}

docs/assets/cookbook.js

@@ -0,0 +1,60 @@
+(() => {
+  let recipesPromise;
+
+  const loadRecipes = (url) => {
+    recipesPromise ||= fetch(url).then((response) => {
+      if (!response.ok) throw new Error(`HTTP ${response.status}`);
+      return response.json();
+    });
+    return recipesPromise;
+  };
+
+  const init = () => {
+    document.querySelectorAll("[data-cookbook]").forEach(async (root) => {
+      if (root.dataset.initialized) return;
+      root.dataset.initialized = "true";
+
+      const select = root.querySelector("[data-cookbook-recipe]");
+      const model = root.querySelector("[data-cookbook-model]");
+      const source = root.querySelector("[data-cookbook-source]");
+      const command = root.querySelector("[data-cookbook-command]");
+      const status = root.querySelector("[data-cookbook-status]");
+
+      try {
+        const { recipes } = await loadRecipes(root.dataset.recipes);
+        const byId = new Map(recipes.map((recipe) => [recipe.id, recipe]));
+        const groups = new Map();
+
+        select.replaceChildren();
+        recipes.forEach((recipe) => {
+          if (!groups.has(recipe.task)) {
+            const group = document.createElement("optgroup");
+            group.label = recipe.task;
+            groups.set(recipe.task, group);
+            select.append(group);
+          }
+          groups.get(recipe.task).append(new Option(recipe.label, recipe.id));
+        });
+
+        const render = () => {
+          const recipe = byId.get(select.value);
+          model.textContent = recipe.model;
+          source.textContent = recipe.source;
+          source.href = `https://github.com/hao-ai-lab/FastVideo/blob/main/${recipe.source}`;
+          command.textContent = recipe.command;
+          status.textContent = `${recipe.label} selected.`;
+        };
+
+        select.addEventListener("change", render);
+        select.disabled = false;
+        render();
+      } catch (error) {
+        status.textContent = "Recipes could not be loaded. Use the examples link below.";
+        console.error("Failed to load FastVideo cookbook recipes", error);
+      }
+    });
+  };
+
+  if (window.document$) window.document$.subscribe(init);
+  else document.addEventListener("DOMContentLoaded", init);
+})();

docs/assets/custom.css

@@ -42,6 +42,46 @@ img {
     margin: 0 auto;
 }
 
+.cookbook-picker {
+    padding: 1rem;
+    border: 0.05rem solid var(--md-default-fg-color--lightest);
+    border-radius: 0.2rem;
+}
+
+.cookbook-picker select {
+    width: 100%;
+    padding: 0.6rem;
+    color: var(--md-default-fg-color);
+    background: var(--md-default-bg-color);
+    border: 0.05rem solid var(--md-default-fg-color--lighter);
+    border-radius: 0.2rem;
+}
+
+.cookbook-picker dl {
+    display: grid;
+    grid-template-columns: max-content 1fr;
+    gap: 0.25rem 1rem;
+}
+
+.cookbook-picker dt {
+    font-weight: 700;
+}
+
+.cookbook-picker dd {
+    margin: 0;
+    min-width: 0;
+    overflow-wrap: anywhere;
+}
+
+.cookbook-picker__status {
+    position: absolute;
+    width: 1px;
+    height: 1px;
+    overflow: hidden;
+    clip: rect(0, 0, 0, 0);
+    white-space: nowrap;
+}
+
 .md-typeset .copy-page-button.md-button {
     float: right;
     margin: 0 0 1rem 1rem;

docs/contributing/ci_architecture.md

@@ -6,7 +6,8 @@ lives in [Testing](testing.md).
 
 ## Overview
 
-FastVideo splits validation across GitHub Actions, Buildkite, and Modal:
+FastVideo splits validation across GitHub Actions, Buildkite, Slinky Slurm,
+and Mergify:
 
 ```text
 PR opened or updated
@@ -16,17 +17,24 @@ PR opened or updated
   |     style, lint, type, spelling, Markdown, workflow syntax, filenames
   |
   |-- Tier 2: Fastcheck
-  |     Buildkite orchestrates Modal GPU jobs
-  |     path-filtered component and unit checks
+  |     Buildkite schedules six lanes on the Slinky Slurm cluster
+  |     encoder, VAE, transformer, kernel, unit, DreamVerse
   |
-  |-- /merge, /test full, or ready label
+  |-- /merge or ready label
         |
-        `-- Tier 3: Full Suite
-              Buildkite orchestrates Modal GPU jobs
-              path-filtered integration, SSIM, training, eval, and performance checks
+        `-- Tier 3: Change-aware merge gate
+              trusted base-branch planner classifies the complete PR diff
+              Buildkite adds only relevant integration, quality, training,
+              API, or performance lanes on Slinky Slurm
               |
               pass -> Mergify squash-merges when all merge conditions pass
               fail -> fix, push, and re-run
+
+  |-- /test full
+        `-- Explicit all-20-lane diagnostic run
+
+  `-- weekly schedule on main
+        `-- Complete four-GPU SSIM matrix
 ```
 
 CI is not one monolithic job:
@@ -34,10 +42,31 @@ CI is not one monolithic job:
 - GitHub Actions owns pre-commit, slash-command handling, aggregate status
   updates, docs deployment, image builds, package publishing, and community
   automations.
-- Buildkite owns the GPU test pipeline and path filtering.
-- Modal owns the actual GPU execution environment for test jobs.
+- Buildkite owns the GPU test graph, statuses, and trusted dispatch control
+  plane. Its agent runs on the Slurm login plane; it does not execute test
+  payloads.
+- Slinky Slurm is the only active CI compute backend. A host-owned dispatcher
+  leases GPUs from a persistent four-GPU allocation and runs each lane in an
+  isolated Enroot container at the immutable PR SHA.
 - Mergify owns merge protection, labeling, and the final squash merge.
 
+The old files under `fastvideo/tests/modal/` are retained as dormant manual
+rollback code. `.buildkite/scripts/pr_test.sh` rejects Buildkite invocations,
+and no pipeline or slash-command route calls Modal.
+
+Three Buildkite entry pipelines share the validated graph:
+
+| Pipeline | Trigger | Scope |
+|---|---|---|
+| `pr-fastcheck` | Automatic pull-request webhook | Six Fastcheck lanes |
+| `ci` | `/merge`, `ready`, schedules, and `/test` API builds | Change-aware merge gates, scheduled SSIM, explicit Full Suite, Fastcheck reruns, or one direct lane |
+| `fastvideo-performance-lane` | Weekly scheduler | Direct performance lane |
+
+Each entry pipeline starts with the same trusted `pipeline-upload` job on the
+`ci-runner` queue. The `ci` pipeline's incoming GitHub webhook is disabled;
+otherwise it would duplicate the automatic `pr-fastcheck` build. API and
+scheduled builds continue to work with webhook processing disabled.
+
 ## CI Tiers
 
 ### Tier 1: Pre-commit
@@ -70,34 +99,37 @@ debugging a hook implementation.
 | Attribute | Value |
 |---|---|
 | Triggered by | Buildkite PR builds with `TEST_SCOPE=fastcheck` or unset |
-| Runner | Buildkite agent that launches Modal GPU jobs |
+| Compute | Slinky Slurm (`ci-runner` queue) |
 | Definition | `.buildkite/pipeline.yml` |
-| Entrypoint | `.buildkite/scripts/pr_test.sh` -> `fastvideo/tests/modal/pr_test.py` |
+| Entrypoint | Trusted host driver -> `.buildkite/scripts/unit_test.sh` or `.buildkite/scripts/lanes/*.sh` |
 
-Fastcheck uses Buildkite's `monorepo-diff` plugin. Jobs whose watched paths did
-not change are skipped and do not block the aggregate `fastcheck-passed`
-status.
+Fastcheck always schedules these six lanes: encoder, VAE, transformer, custom
+kernels, unit tests, and DreamVerse. Static steps replace the former
+host-side path-filter plugin: the login plane never checks out or executes PR
+code.
 
-| Buildkite label | `TEST_TYPE` | Main watched paths |
-|---|---|---|
-| Encoder Tests | `encoder` | `fastvideo/models/encoders/**`, `fastvideo/models/loader/**`, `fastvideo/tests/encoders/**`, `pyproject.toml`, `docker/Dockerfile` |
-| VAE Tests | `vae` | `fastvideo/models/vaes/**`, `fastvideo/models/loader/**`, `fastvideo/tests/vaes/**`, `pyproject.toml`, `docker/Dockerfile` |
-| Transformer Tests | `transformer` | `fastvideo/models/dits/**`, `fastvideo/models/loader/**`, `fastvideo/tests/transformers/**`, `fastvideo/layers/**`, `fastvideo/attention/**`, `pyproject.toml`, `docker/Dockerfile` |
-| Kernel Tests | `kernel_tests` | `fastvideo-kernel/**`, `pyproject.toml`, `docker/Dockerfile` |
-| Unit Tests | `unit_test` | `fastvideo/**`, `.buildkite/**`, `.github/**`, `pyproject.toml`, `docker/Dockerfile` |
-| DreamVerse App Tests | `dreamverse_app` | `apps/dreamverse/**`, `pyproject.toml` |
-
-### Tier 3: Full Suite
+### Tier 3: Change-Aware Merge Gate
 
 | Attribute | Value |
 |---|---|
-| Triggered by | `/merge`, adding `ready`, `/test full`, or a new push to a PR that already has `ready` |
-| Runner | Buildkite agent that launches Modal GPU jobs |
+| Triggered by | `/merge`, adding `ready`, or a new push to a PR that already has `ready` |
+| Compute | Slinky Slurm only (`ci-runner` queue) |
 | Definition | `.buildkite/pipeline.yml` |
-| Entrypoint | `.buildkite/scripts/pr_test.sh` -> `fastvideo/tests/modal/pr_test.py` |
+| Entrypoint | `/opt/fastvideo-ci-runner/run-ci` (`run-unit` is a compatibility wrapper) |
 
-Full Suite is also path-filtered. It validates broader behavior before Mergify
-can merge a PR.
+Fastcheck is the universal six-lane baseline. The merge gate does not repeat
+those jobs: it classifies every changed path and adds only the relevant lanes
+from the fourteen-lane integration set below. Selected jobs are hard gates;
+there are no soft-fail hardware lanes. A documentation-only PR can therefore
+finish its merge build after the trusted uploader, while model-family changes
+typically add focused golden-gate and SSIM files and a training-only change
+adds only its owning training lane.
+
+`.github/scripts/plan_merge_ci.py` is the canonical path policy. It runs from
+the immutable base SHA under `pull_request_target`; PR code is never executed
+on the GitHub runner. The changed-file list includes both sides of renames. An
+API failure, truncated response, empty list, unknown build input, or unknown
+source path fails closed to all fourteen integration lanes.
 
 A `ready`-labeled PR does not hit Buildkite immediately:
 `ci-trigger-full-suite.yml` first runs `.github/scripts/gate_full_suite.sh`,
@@ -106,25 +138,100 @@ head. A red cheap check blocks the suite (fail closed; the next push re-arms
 it), while a GitHub outage or a >25 min wait lets it run anyway (fail open).
 `/test full` bypasses the gate.
 
-| Buildkite label | `TEST_TYPE` | Main watched paths |
-|---|---|---|
-| SSIM Tests | `ssim` | `fastvideo/**/*.py`, `pyproject.toml`, `docker/Dockerfile` |
-| LoRA Inference Tests | `inference_lora` | LoRA tests, loader, transformer tests, pipelines, LoRA layers |
-| LoRA Extraction Tests | `lora_extraction` | LoRA extraction scripts/tests, loader, training utilities, LoRA layers |
-| Training Tests | `training` | `fastvideo/**`, `pyproject.toml`, `docker/Dockerfile` |
-| Distillation DMD Tests | `distillation_dmd` | `fastvideo/training/*distillation_pipeline.py` |
-| Self-Forcing Tests | `self_forcing` | self-forcing distillation pipeline and tests |
-| LoRA Training Tests | `training_lora` | `fastvideo/**`, `pyproject.toml`, `docker/Dockerfile` |
-| Training Tests VSA | `training_vsa` | `fastvideo/**`, `fastvideo-kernel/**`, `pyproject.toml`, `docker/Dockerfile` |
-| Inference Tests VMoBA | `inference_vmoba` | `fastvideo-kernel/**`, `fastvideo/attention/backends/vmoba.py` |
-| Performance Tests | `performance` | DiTs, pipelines, attention, layers, worker, entrypoints, performance tests/configs |
-| API Server Tests | `api_server` | OpenAI entrypoints, serve CLI, OpenAI API integration test |
-| Train Framework Tests | `train_framework` | `fastvideo/train/**`, train model/method tests, model loader, DiTs |
-| Eval Metrics Tests | `eval` | `fastvideo/eval/**`, `fastvideo/tests/eval/**`, `pyproject.toml`, `docker/Dockerfile` |
+The complete static graph remains available through `/test full`; path
+selection never deletes or dynamically invents a Buildkite step.
+
+| Lane | Public `TEST_TYPE` | GPUs | Typical merge trigger |
+|---|---|---:|---|
+| Encoder | `encoder` | 1 | Universal Fastcheck |
+| VAE | `vae` | 1 | Universal Fastcheck |
+| Transformer | `transformer` | 1 | Universal Fastcheck |
+| Kernel | `kernel_tests` | 1 | Universal Fastcheck |
+| Unit | `unit_test` | 1 | Universal Fastcheck |
+| DreamVerse | `dreamverse_app` | 1 | Universal Fastcheck |
+| Golden gate | `golden_gate` | 1 | Model, pipeline, attention, layer, or output changes |
+| SSIM | `ssim` | 4 | Matching model/SSIM paths; focused files when possible |
+| LoRA inference | `inference_lora` | 1 | LoRA inference/shared LoRA paths |
+| LoRA extraction | `lora_extraction` | 1 | LoRA extraction/shared LoRA paths |
+| Vanilla training | `training` | 4 | Legacy vanilla/shared training paths |
+| DMD distillation | `distillation_dmd` | 2 | DMD/shared training paths |
+| Self-forcing | `self_forcing` | 2 | Self-forcing/shared training paths |
+| LoRA training | `training_lora` | 2 | LoRA/shared training paths |
+| VSA training | `training_vsa` | 2 | VSA/shared training paths |
+| VMoBA inference | `inference_vmoba` | 1 | VMoBA backend/config paths |
+| Performance | `performance` | 2 | Performance tests or benchmark policy |
+| API server | `api_server` | 1 | API, worker, or server entrypoint paths |
+| Modular train framework | `train_framework` | 1 | `fastvideo/train/` and its tests |
+| Eval metrics | `eval` | 1 | `fastvideo/eval/` and its tests |
+
+Golden-gate and SSIM selections are basenames, not arbitrary pytest arguments.
+The private host checks the comma-separated allowlist before staging, and the
+container checks it again before invoking pytest. Shared quality-harness
+changes still run the complete owning lane. The full SSIM matrix also runs on
+`main` every Sunday at 05:00 UTC through `ci-scheduled-ssim.yml`; `/test ssim`
+and `/test full` remain available for deliberate complete runs.
+
+The four Buildkite workers may accept multiple jobs concurrently. The
+agent-owned lease broker packs their requested GPU counts onto one persistent
+four-GPU Slurm allocation and waits when capacity is full. A four-GPU lane
+such as SSIM or vanilla training owns the whole tray; two two-GPU lanes or up
+to four one-GPU lanes can overlap without sharing devices. SSIM and vanilla
+training also share the `fastvideo/slinky/whole-tray` Buildkite concurrency
+group. That keeps the second whole-tray lane in Buildkite instead of consuming
+an agent and its command timeout while the first lane waits for four free GPUs.
+Because packed Enroot containers share the node network namespace, each GPU
+lease receives its own 100-port rendezvous range. Tests preserve the
+runner-assigned `MASTER_PORT`, and parallel SSIM tasks use distinct offsets
+inside that range.
 
 See [Performance Benchmarks](performance_benchmarks.md) for the performance
 lane's thresholds, rolling baseline, artifacts, and reseeding process.
 
+## `/merge` Request Flow
+
+The Buildkite agent and Slurm have deliberately separate responsibilities:
+
+```text
+/merge PR comment
+  -> GitHub verifies write permission and refreshes the `ready` label
+  -> base-branch `ci-trigger-full-suite` workflow fetches the PR file list,
+     computes MERGE_TEST_PLAN plus focused golden/SSIM basenames, and gates on
+     cheap checks
+  -> sends the PR SHA to pipeline `ci` with TEST_SCOPE=merge and FULL_SUITE=true
+  -> trusted `pipeline-upload` job on queue `ci-runner`
+       fetch exact-SHA .buildkite/pipeline.yml
+       normalize + validate the complete static 20-lane policy and conditions
+       upload static Buildkite steps
+  -> each accepted step reaches the host policy hook
+       validate org/repo/SHA/ref/step key/command/scope/timeout
+       skip checkout on the login plane
+       stage a mode-0600 request on Lustre
+       lease 1-4 GPUs and attach an `srun` step to the Slinky tray
+  -> Enroot worker container
+       clone and verify the exact PR SHA
+       install the lane's project extras and cached kernel
+       run the repository-owned lane script
+       write numeric exit status and approved artifacts
+  -> trusted host returns that status to Buildkite
+  -> Buildkite publishes `full-suite-passed` only when every selected lane passes
+```
+
+The pipeline uploader and dispatcher run on the Slurm login plane, but those
+are control-plane operations only. Python tests, model loading, CUDA kernels,
+Node/Playwright checks, inference, training, SSIM generation, and performance
+benchmarks all execute in Slurm allocations.
+
+PR-controlled values never become host commands. The host policy accepts only
+the pinned pipeline uploader or a known lane tuple. It rejects plugins,
+artifact globs, shell injection variables, non-immutable commits, and unknown
+commands before checkout. Hugging Face credentials are added only for lanes
+that declare them, passed through a mode-0600 request file, and removed before
+the PR payload starts. Active training lanes keep W&B offline and do not stage
+a W&B credential. The ARM64 image includes the pinned FA4 CuTe overlay validated
+on GB200. SSIM opts into FA4 to preserve its reference-video numerics; lanes
+with FA2 baselines keep `FASTVIDEO_FA4=0`. Performance artifacts are relayed
+afterward by the trusted host from an allowlisted directory and extension set.
+
 ## Slash Commands
 
 Slash commands are handled by `.github/workflows/ci-slash-commands.yml`.
@@ -132,7 +239,7 @@ Repository write permission is required.
 
 | Command | Effect |
 |---|---|
-| `/merge` | Adds `ready` and triggers Full Suite for the PR head branch. |
+| `/merge` | Adds `ready` and triggers the path-aware merge gate for the PR head. |
 | `/test full` | Runs the whole Full Suite with `TEST_SCOPE=full`. |
 | `/test fastcheck` | Runs the whole Fastcheck suite with `TEST_SCOPE=fastcheck`. |
 | `/test pre-commit` | Re-runs the pre-commit workflow on the PR merge ref. |
@@ -149,6 +256,7 @@ Valid direct test names:
 | `/test unit` | `unit_test` |
 | `/test dreamverse` | `dreamverse_app` |
 | `/test ssim` | `ssim` |
+| `/test golden-gate` | `golden_gate` |
 | `/test training` | `training` |
 | `/test lora-inference` | `inference_lora` |
 | `/test lora-training` | `training_lora` |
@@ -162,12 +270,18 @@ Valid direct test names:
 | `/test train-framework` | `train_framework` |
 | `/test eval` | `eval` |
 
+The temporary `<name>-ci` spellings remain accepted as compatibility aliases;
+they select the same Slurm lane and do not identify a second backend.
+
 When a direct test completes successfully, Buildkite posts
 `direct-test-completed`. `.github/workflows/ci-aggregate-status.yml` then reads
 the latest Buildkite statuses for the commit and updates `fastcheck-passed` or
 `full-suite-passed` if all jobs in that group are green.
 
-Skipped path-filtered jobs have no status entry and do not block the aggregate.
+Buildkite label emojis define the status namespace used by that aggregation:
+`:microscope:` is reserved for the six Fastcheck lanes, while Full-Suite-only
+lanes use `:test_tube:` or `:bar_chart:`. Each active lane has exactly one
+label and therefore one status context.
 
 ## Merge Protection
 
@@ -177,7 +291,7 @@ Mergify enforces these conditions before it squash-merges to `main`:
 |---|---|
 | `check-success~=pre-commit` | Tier 1 passed. |
 | `check-success=fastcheck-passed` | All triggered Fastcheck jobs passed. |
-| `check-success=full-suite-passed` | All triggered Full Suite jobs passed. |
+| `check-success=full-suite-passed` | The selected merge gate or explicit Full Suite passed. |
 | `#approved-reviews-by>=1` | At least one approving review. |
 | Valid title regex | PR title starts with an accepted `[type]` tag. |
 | `label=ready` | The PR has entered the merge flow. |
@@ -226,29 +340,40 @@ Process labels:
 
 | Label | Who sets it | Meaning |
 |---|---|---|
-| `ready` | `/merge` or maintainer action | Triggers/keeps Full Suite active and enables auto-merge. |
+| `ready` | `/merge` or maintainer action | Triggers/keeps the change-aware merge gate active and enables auto-merge. |
 | `needs-rebase` | Mergify | PR has merge conflicts. |
 | `do-not-merge` | Maintainer | Blocks merge regardless of CI status. |
 
-## Modal Test Entrypoints
+## Slurm Lane Entrypoints
 
-All Buildkite test jobs go through `.buildkite/scripts/pr_test.sh`, which:
+Every active test selection lives in `.buildkite/scripts/unit_test.sh` or a
+focused `.buildkite/scripts/lanes/<lane>.sh`. The private, agent-owned lane
+table binds each internal `*_ci` type to that script, its GPU count, wall-clock
+limit, dependency extras, kernel-build policy, secrets, and artifacts. The
+internal suffix is an implementation detail; there is only one active backend.
 
-1. Reads Buildkite secrets for Modal, Hugging Face, and W&B when needed.
-2. Selects a Modal function based on `TEST_TYPE`.
-3. Passes Buildkite metadata into the Modal container.
-4. Runs the selected test command from `fastvideo/tests/modal/pr_test.py` or
-   `fastvideo/tests/modal/ssim_test.py`.
-5. Uploads performance artifacts for `TEST_TYPE=performance`.
+SSIM uses `fastvideo/tests/ssim/ci_runner.py` inside a single four-GPU lease.
+It discovers `REQUIRED_GPUS` and `*_MODEL_TO_PARAMS` with AST parsing, then
+packs independent pytest subprocesses across the visible GPUs with fail-fast
+termination. Performance writes reports to a host-mounted artifact directory;
+the host uploads only `.md`, `.html`, `.json`, and `.csv` files after the
+container exits.
+
+The Modal launchers remain in the repository for manual rollback archaeology,
+but they are not CI entrypoints. `pr_test.sh` rejects Buildkite calls and needs
+`FASTVIDEO_ENABLE_LEGACY_MODAL_CI=1` even for a local manual invocation.
 
 If you add a new CI test category:
 
-1. Add the Modal function in `fastvideo/tests/modal/pr_test.py` or a focused
-   companion module.
-2. Add the `TEST_TYPE` case in `.buildkite/scripts/pr_test.sh`.
-3. Add the Buildkite direct-test step and any Fastcheck/Full Suite path filters
-   in `.buildkite/pipeline.yml`.
-4. Add or update the `/test` mapping in `.github/workflows/ci-slash-commands.yml`.
+1. Add an executable `.buildkite/scripts/lanes/<lane>.sh` containing the test
+   payload only.
+2. Add the static Buildkite step in `.buildkite/pipeline.yml`, its source/test
+   ownership in `.github/scripts/plan_merge_ci.py`, and the `/test` mapping in
+   `.github/workflows/ci-slash-commands.yml`.
+3. Extend `fastvideo/tests/contract/test_ci_test_collection.py`,
+   `test_merge_ci_plan.py`, and the trusted private runner's lane table plus
+   uploader policy in the same rollout.
+4. Validate on the target GB200 hardware before making the lane a merge gate.
 5. Document the lane here and link any domain-specific authoring guide.
 
 ## CD And Release Workflows
@@ -282,13 +407,18 @@ the explicit `py3.12-cuda13.0.0-latest` tag. This publication policy does not
 change the unparameterized `docker/Dockerfile` build defaults, which remain CUDA
 13 and `cu130`.
 
-Published amd64 development images keep their configured Hopper kernel wheel
-installed and also carry an immutable SM89 wheel under
-`/opt/fastvideo-kernel-prebuilt`. Modal PR and SSIM jobs select the exact
-source, ABI, and GPU-architecture match from that directory, so L40S jobs reuse
-the trusted image artifact while kernel-changing PRs still build locally. Once
-a kernel or artifact-key change reaches `main`, the image workflow republishes
-the matching trusted artifact before later jobs consume the updated image tag.
+Published development images carry architecture-specific kernel wheels under
+`/opt/fastvideo-kernel-prebuilt`. The Slurm worker selects the exact source,
+ABI, and GPU-architecture match, so normal lanes reuse the trusted artifact
+while kernel-changing PRs still build locally. Once a kernel or artifact-key
+change reaches `main`, the image workflow republishes the matching artifact
+before later jobs consume the updated image pin.
+
+The same workflow publishes a single-architecture ARM64, CUDA 13, SM100 image
+for the self-hosted CI runner under the
+`py3.12-cuda13.0.0-sm100-{latest,sha-*}` tags. It carries the matching prebuilt
+kernel wheel so runner jobs can validate and install the exact source and ABI
+match instead of recompiling it in every lane.
 
 The optional Dreamverse matrix builds backend and UI images for CUDA 12.6 and
 CUDA 13 on `amd64`. Dreamverse remains `amd64`-only because its FA4 dependency
@@ -320,15 +450,17 @@ The reusable implementation lives in
 | `.github/mergify.yml` | Merge protection, PR title validation, PR labels, conflict labels, auto-merge |
 | `.github/workflows/ci-precommit.yml` | Tier 1 pre-commit |
 | `.github/workflows/ci-slash-commands.yml` | `/merge` and `/test` handling |
-| `.github/workflows/ci-trigger-full-suite.yml` | Full Suite trigger for `ready` PRs and new pushes to ready PRs |
-| `.github/workflows/ci-aggregate-status.yml` | Aggregate Fastcheck/Full Suite commit statuses |
-| `.buildkite/pipeline.yml` | Buildkite test graph and path filters |
-| `.buildkite/scripts/pr_test.sh` | Buildkite-to-Modal test dispatcher |
-| `fastvideo/tests/modal/pr_test.py` | Modal functions for most GPU CI lanes |
-| `fastvideo/tests/modal/ssim_test.py` | Modal functions and partitioning for SSIM |
+| `.github/workflows/ci-trigger-full-suite.yml` | Change-aware merge-gate trigger for `ready` PRs and new pushes |
+| `.github/workflows/ci-scheduled-ssim.yml` | Weekly complete SSIM trigger on `main` |
+| `.github/scripts/plan_merge_ci.py` | Trusted changed-path to integration-lane planner |
+| `.github/workflows/ci-aggregate-status.yml` | Aggregate Fastcheck and explicit Full Suite direct-rerun statuses |
+| `.buildkite/pipeline.yml` | Static 20-lane Slurm Buildkite graph |
+| `.buildkite/scripts/unit_test.sh`, `.buildkite/scripts/lanes/*.sh` | Active Slurm lane payloads |
+| `fastvideo/tests/ssim/ci_runner.py` | Four-GPU Slurm SSIM scheduler |
+| `.buildkite/scripts/pr_test.sh`, `fastvideo/tests/modal/*.py` | Dormant manual Modal rollback path (disabled in Buildkite) |
 | `.buildkite/performance-benchmarks/tests/*.json` | Performance benchmark configs and thresholds |
 | `.github/workflows/infra-docs.yml` | Docs build and GitHub Pages deploy |
-| `.github/workflows/infra-build-image.yml` | Automatic CUDA matrix and manual Docker image builds |
+| `.github/workflows/infra-build-image.yml` | CUDA matrix, CI runner image, and manual Docker image builds |
 | `.github/workflows/publish-fastvideo.yml` | FastVideo PyPI publishing |
 | `.github/workflows/publish-kernel.yml` | FastVideo kernel PyPI publishing |
 | `.github/workflows/publish-comfyui.yml` | ComfyUI registry publishing |

docs/contributing/performance_benchmarks.md

@@ -23,8 +23,8 @@ It serves three audiences:
 pytest fastvideo/tests/performance/ -vs
 
 # Optional: compare against the rolling HF baseline.
-# PERF_REPORTS_DIR defaults to /root/data/perf_reports for Modal/CI, so
-# override it when running outside the container.
+# PERF_REPORTS_DIR defaults to /root/data/perf_reports in a CI container, so
+# override it for a local run.
 PERF_REPORTS_DIR=/tmp/fastvideo_perf_reports \
 python fastvideo/tests/performance/compare_baseline.py
 
@@ -459,16 +459,18 @@ FlashInfer, Cutlass DSL, SageAttention, Triton, and xFormers when installed.
 | `FASTVIDEO_FA4` | `0` | `test_inference_performance.py` | FlashAttention-4 toggle included in `software_profile_id`. |
 | `FASTVIDEO_PERFORMANCE_PROFILE_VERSION` | unset | `test_inference_performance.py` | Optional explicit software cohort/profile version included in `software_profile_id`. |
 | `IMAGE_VERSION` | unset | `test_inference_performance.py` | CI container image/profile version included in `software_profile_id` when available. |
-| `FASTVIDEO_CONTAINER_IMAGE_REF` | unset | `pr_test.py`, `launch_l40s_job.py`, `test_inference_performance.py` | Resolved CI container image ref or digest recorded in `environment_metadata` for audit without changing `software_profile_id`. |
+| `FASTVIDEO_CONTAINER_IMAGE_REF` | unset | Slurm runner, `test_inference_performance.py` | Pinned CI container image digest recorded in `environment_metadata` and `software_profile_id`. |
 | `FASTVIDEO_STAGE_LOGGING` | set by the pytest test | `test_inference_performance.py` | Enables pipeline stage timing capture for component metrics during benchmark runs. |
 
 ## CI integration
 
-The performance step can run on demand with `/test performance` and as part of
-the Full Suite (see [CI/CD Architecture](ci_architecture.md)). The Modal entry
-point is `fastvideo/tests/modal/pr_test.py:run_performance_tests` and the
-Buildkite artifact upload is in
-`.buildkite/scripts/pr_test.sh:upload_performance_artifacts`.
+The performance step can run on demand with `/test performance`, through a
+merge gate when performance tests or benchmark policy changed, and as part of
+an explicit `/test full` run (see [CI/CD Architecture](ci_architecture.md)).
+The weekly `fastvideo-performance-lane` schedule runs the same Slurm payload.
+The active entry point is `.buildkite/scripts/lanes/performance.sh`; the
+trusted host dispatcher relays its allowlisted reports to Buildkite after the
+isolated container exits.
 
 Each performance build runs pytest first. PR and direct runs only continue to
 `compare_baseline.py` when that fixed-threshold phase passes; if pytest fails,

docs/contributing/pull_requests.md

@@ -42,19 +42,20 @@ The important process labels are:
 
 | Label | Meaning |
 |---|---|
-| `ready` | The PR is ready for Full Suite and auto-merge consideration. |
+| `ready` | The PR is ready for the change-aware merge gate and auto-merge consideration. |
 | `needs-rebase` | The PR has merge conflicts with `main`. |
 | `do-not-merge` | A maintainer has blocked merge. |
 
 ## CI Summary
 
-FastVideo has three validation tiers:
+FastVideo has three routine validation tiers plus an explicit full diagnostic:
 
 | Tier | Runs when | What it does |
 |---|---|---|
 | Pre-commit | Pull requests and `/test pre-commit` | Formatting, linting, typing, spelling, Markdown, workflow syntax, filename checks |
-| Fastcheck | PR Buildkite builds | Path-filtered component and unit checks on Modal GPU runners |
-| Full Suite | `/merge`, `ready`, `/test full`, or new pushes to ready PRs | Path-filtered integration, SSIM, training, eval, API, and performance checks |
+| Fastcheck | PR Buildkite builds | Six component, kernel, unit, and app lanes on Slinky Slurm |
+| Merge gate | `/merge`, `ready`, or new pushes to ready PRs | Only path-relevant integration lanes on Slinky Slurm; Fastcheck remains the baseline |
+| Full Suite | `/test full` | Explicit all-twenty-lane diagnostic run on Slinky Slurm |
 
 See [CI/CD Architecture](ci_architecture.md#ci-tiers) for exact jobs, path
 filters, and workflow files.
@@ -66,10 +67,11 @@ filters, and workflow files.
 3. Fix pre-commit failures locally with `pre-commit run --all-files`.
 4. Wait for at least one approving review.
 5. When the PR is approved and ready, comment `/merge`.
-6. `/merge` adds `ready` and triggers the Full Suite for the PR branch.
+6. `/merge` adds `ready`, waits for cheap checks, and triggers the minimal
+   path-relevant integration lanes for the PR branch.
 7. If all required checks pass, Mergify squash-merges the PR to `main`.
-8. If Full Suite fails, fix the regression, push again, and re-run `/merge` or
-   the failed test.
+8. If the merge gate fails, fix the regression, push again, and re-run
+   `/merge`. Use a targeted `/test` command for diagnosis.
 
 Only contributors with repository write permission can use slash commands. If
 you are an external contributor, ask a maintainer to run `/merge` or add
@@ -128,7 +130,7 @@ git push --force-with-lease
 
 Mergify removes `needs-rebase` after conflicts are resolved.
 
-### Full Suite Fails
+### Merge Gate Or Full Suite Fails
 
 The failing Buildkite step is the source of truth. Common causes are:
 

docs/contributing/testing.md

@@ -138,47 +138,31 @@ pytest fastvideo/tests/ssim/ -vs
 
 Use a machine whose GPU and backend match the reference folder you are testing.
 
-## Modal Runs For SSIM
+## Slurm CI Runs For SSIM
 
-For CI-like SSIM execution, use `fastvideo/tests/modal/ssim_test.py`:
+Comment `/test ssim` on a pull request to run the canonical four-GPU SSIM
+lane on the Slinky Slurm cluster. `fastvideo/tests/ssim/ci_runner.py`
+discovers the suite without importing test modules, packs independent pytest
+processes across the four assigned GPUs, and stops the lane on the first
+failure.
+
+The change-aware `/merge` planner may run only the SSIM test basenames owned
+by the changed model family. Shared SSIM harness changes still select the
+complete lane. Independently, `main` runs the full SSIM matrix every Sunday at
+05:00 UTC so infrequently touched model families retain periodic coverage.
+
+For a focused developer run, invoke pytest directly and optionally select one
+model from a parameterized test through `FASTVIDEO_SSIM_MODEL_ID`:
 
 ```bash
-python -m modal run fastvideo/tests/modal/ssim_test.py::run_ssim_tests
+pytest fastvideo/tests/ssim/test_wan_t2v_similarity.py -vs
+
+FASTVIDEO_SSIM_MODEL_ID=Wan2.1-T2V-1.3B-Diffusers \
+pytest fastvideo/tests/ssim/test_wan_t2v_similarity.py -vs
 ```
 
-Target specific files or model ids:
-
-```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_API_KEY`, `HUGGINGFACE_HUB_TOKEN`, or `HF_TOKEN` is not set, 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 downloads into
-`./generated_videos_modal/<quality-tier>`. Convert those outputs into local
-references with `copy-local`:
-
-```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
-```
+The files under `fastvideo/tests/modal/` are retained only as a disabled
+manual rollback implementation. No active CI trigger invokes them.
 
 ### SSIM Bootstrap Mode
 
@@ -206,28 +190,30 @@ python fastvideo/tests/ssim/reference_videos_cli.py promote-draft \
 
 ## CI Integration
 
-FastVideo CI tests are orchestrated by Buildkite and run on Modal GPU
-instances. The main files are:
+FastVideo GPU CI is orchestrated by Buildkite and runs only on isolated Slinky
+Slurm workers. The main files are:
 
 | File | Purpose |
 |---|---|
-| `.buildkite/pipeline.yml` | Buildkite test graph and path filters. |
-| `.buildkite/scripts/pr_test.sh` | Dispatches `TEST_TYPE` to a Modal function. |
-| `fastvideo/tests/modal/pr_test.py` | Modal functions for most test lanes. |
-| `fastvideo/tests/modal/ssim_test.py` | Modal functions and partitioning for SSIM. |
+| `.buildkite/pipeline.yml` | Static, validated 20-lane Slurm test graph. |
+| `.github/scripts/plan_merge_ci.py` | Trusted path-to-lane and focused quality-test policy for `/merge`. |
+| `.buildkite/scripts/unit_test.sh`, `.buildkite/scripts/lanes/*.sh` | Repository-owned test payloads executed inside Slurm containers. |
+| `fastvideo/tests/ssim/ci_runner.py` | Four-GPU SSIM task discovery and scheduling. |
+| `.buildkite/scripts/pr_test.sh`, `fastvideo/tests/modal/*.py` | Dormant manual rollback path; rejected in Buildkite. |
 
-For exact tier membership, path filters, slash commands, and aggregate statuses,
+For exact tier membership, slash commands, runner isolation, and aggregate statuses,
 see [CI/CD Architecture](ci_architecture.md).
 
 ### Adding A New CI Test Category
 
 If a new test does not fit an existing lane:
 
-1. Add a Modal function in `fastvideo/tests/modal/pr_test.py` or a focused
-   companion module.
-2. Add a matching `TEST_TYPE` case in `.buildkite/scripts/pr_test.sh`.
-3. Add Buildkite direct-test and path-filter entries in `.buildkite/pipeline.yml`.
-4. Add the `/test` mapping in `.github/workflows/ci-slash-commands.yml`.
+1. Put the test payload in an executable `.buildkite/scripts/lanes/<lane>.sh`.
+2. Add its static step to `.buildkite/pipeline.yml`, its changed-path ownership
+   to `.github/scripts/plan_merge_ci.py`, and extend the CI contract tests.
+3. Add the `/test` mapping in `.github/workflows/ci-slash-commands.yml`.
+4. Coordinate the matching GPU, timeout, dependency, secret, and artifact
+   policy in the private Slurm runner allowlist.
 5. Document the new category in [CI/CD Architecture](ci_architecture.md) and add
    authoring notes here if contributors need them.
 

docs/cookbook/index.md

@@ -0,0 +1,44 @@
+# Inference Cookbook
+
+Choose a complete recipe maintained in the FastVideo repository. Each command
+runs its checked-in source directly, so coupled model, GPU, offload, and
+attention settings do not drift into unsupported combinations.
+
+The commands expect a local clone:
+
+```bash
+git clone https://github.com/hao-ai-lab/FastVideo.git
+cd FastVideo
+```
+
+<div class="cookbook-picker" data-cookbook data-recipes="../assets/cookbook-recipes.json">
+  <label for="cookbook-recipe"><strong>Recipe</strong></label>
+  <select id="cookbook-recipe" data-cookbook-recipe disabled>
+    <option>Loading recipes…</option>
+  </select>
+  <dl>
+    <dt>Model</dt>
+    <dd data-cookbook-model>Loading…</dd>
+    <dt>Source</dt>
+    <dd><a data-cookbook-source href="../inference/examples/basic/">Browse maintained examples</a></dd>
+  </dl>
+  <pre><code class="language-bash" data-cookbook-command>Loading…</code></pre>
+  <p class="cookbook-picker__status" role="status" aria-live="polite" data-cookbook-status></p>
+  <noscript>
+    JavaScript is needed for the recipe picker. Browse the
+    <a href="../inference/examples/examples_inference_index/">inference examples</a>
+    instead.
+  </noscript>
+</div>
+
+## Customize a recipe
+
+Start from the checked-in source, then change only the settings your model
+supports:
+
+- [Configuration](../inference/configuration.md) covers the Python and CLI
+  config surfaces.
+- [Optimizations](../inference/optimizations.md) covers attention backends,
+  compilation, and memory tradeoffs.
+- [Support matrix](../inference/support_matrix.md) lists supported models and
+  optimizations.

docs/design/apple_silicon_fast_mode.md

@@ -8,7 +8,7 @@ compute ~3.7×, so the wall-clock drops far more than 2×. RIFE (which estimates
 its own optical flow — no motion vectors needed) fills the dropped frames back
 in for ~1.4 s, and a light unsharp pass counters its softening.
 
-Measured on the 1.3B INT8 QAD model (fox, 480×832×81, M4): generate 41 + RIFE→81
+Measured on the 1.3B INT8 QAD model (480×832×81, M4): generate 41 + RIFE→81
 runs in ~35 s of denoise vs ~90 s full, at reconstruction MS-SSIM **0.97**.
 Reproduce with `python -m fastvideo.benchmarks.eval_metalfx_rife --mode int8`.
 
@@ -26,10 +26,11 @@ uv pip install -e ".[mlx]"   # RIFE ships vendored; this only needs MLX
 
 ```bash
 python examples/inference/basic/mlx_wan_prompt_to_video.py \
-  --mlx-checkpoint <FastWan2.1-T2V-1.3B-INT8-QAD> \
-  --prompt "A red fox trotting through a snowy pine forest at golden hour, cinematic" \
+  --model-root ./FastMetal-1.3B-QAD \
+  --mlx-checkpoint ./FastMetal-1.3B-QAD \
+  --prompt "A bird's-eye view of a misty forest valley at dawn." \
   --num-frames 81 --fast \
-  --output-path video_samples/fox_fast.mp4
+  --output-path video_samples/forest_fast.mp4
 ```
 
 `--num-frames` stays the *target* length; fast mode generates the smallest
@@ -57,9 +58,9 @@ of denoise. It composes with `--fast`; both together run the same clip in
 
 ```bash
 python examples/inference/basic/mlx_wan_prompt_to_video.py \
-  --prompt "A red fox trotting through a snowy pine forest at golden hour, cinematic" \
+  --prompt "A bird's-eye view of a misty forest valley at dawn." \
   --height 480 --width 832 --num-frames 81 --fast-spatial \
-  --output-path video_samples/fox_fast_spatial.mp4
+  --output-path video_samples/forest_fast_spatial.mp4
 ```
 
 | Flag | Default | Meaning |

docs/design/inference_schema_parity_inventory.yaml

@@ -21,6 +21,8 @@ surfaces:
       hsdp_shard_dim: generator.engine.parallelism.hsdp_shard_dim
       dist_timeout: generator.engine.parallelism.dist_timeout
       lora_path: generator.pipeline.components.lora_path
+      lora_nickname: generator.pipeline.components.lora_nickname
+      lora_strength: generator.pipeline.components.lora_strength
       dit_cpu_offload: generator.engine.offload.dit
       use_fsdp_inference: generator.engine.use_fsdp_inference
       dit_layerwise_offload: generator.engine.offload.dit_layerwise
@@ -72,10 +74,14 @@ surfaces:
     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."
+      VSA_tile_size: "VSA-H3 tile geometry request; model-specific optimization not yet represented in the typed public schema."
+      inference_torch_compile: "Regional inference compile opt-in currently carried through PipelineSelection.experimental rather than CompileConfig."
+      vae_parallel_decode: "MiniMax-H3 sequence-parallel VAE decode opt-in; model-specific optimization not yet represented in the typed public schema."
+      vae_parallel_encode: "MiniMax-H3 sequence-parallel reference VAE encode opt-in; model-specific optimization not yet represented in the typed public schema."
+      vae_parallel_decode_strategy: "Chunk-transport collective for vae_parallel_decode; model-specific optimization not yet represented in the typed public schema."
       attention_backend: "Process-wide default attention-backend request applied per component at load time; kernel-selection knob 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."
@@ -585,15 +591,31 @@ surfaces:
   openai_video_request:
     kept:
       model: "HTTP adapter model-routing field."
+      user: "OpenAI-compatible caller tracking field."
+      task: "SGLang-compatible MiniMax-H3 task selector validated against the startup pipeline."
+      quality: "vLLM-Omni-compatible quality-intent field; model-specific."
+      lora: "vLLM-Omni-compatible selector for the adapter fixed at server startup."
     moved:
       prompt: request.prompt
       input_reference: request.inputs.image_path
       reference_url: request.inputs.image_path
+      image_reference: request.inputs.image_path,last_image,references
+      video_reference: request.inputs.video_path,references
+      audio_reference: request.inputs.references
+      video_path: request.inputs.video_path
+      video_url: request.inputs.video_path
+      video_params: request.sampling.width,height,num_frames,fps
       size:
         target: request.sampling.width,height
         note: "Adapter parses OpenAI size strings as WIDTHxHEIGHT and forwards width then height."
+      width: request.sampling.width
+      height: request.sampling.height
       fps: request.sampling.fps
       num_frames: request.sampling.num_frames
+      aspect_ratio: request.sampling.width,height
+      short_edge: request.sampling.width,height
+      num_outputs_per_prompt: request.sampling.num_videos_per_prompt
+      n: request.sampling.num_videos_per_prompt
       seed: request.sampling.seed
       num_inference_steps: request.sampling.num_inference_steps
       guidance_scale: request.sampling.guidance_scale
@@ -601,11 +623,21 @@ surfaces:
       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
+      max_sequence_length: request.sampling.max_sequence_length
+      boundary_ratio: request.sampling.boundary_ratio
+      extra_params: request.extensions
     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."
+      start_time_seconds: "vLLM-Omni reference-video offset; rejected by pipelines that cannot represent it."
+      flow_shift: "vLLM-Omni request field; accepted only when the selected model exposes a matching request parameter."
+      generate_sound: "vLLM-Omni audio-output intent; accepted only by models with a matching request parameter."
+      sound_duration: "vLLM-Omni audio-duration intent; accepted only by models with a matching request parameter."
+      enable_frame_interpolation: "vLLM-Omni post-processing field; unavailable until FastVideo exposes a frame-interpolation stage."
+      frame_interpolation_exp: "vLLM-Omni post-processing field; unavailable until FastVideo exposes a frame-interpolation stage."
+      frame_interpolation_scale: "vLLM-Omni post-processing field; unavailable until FastVideo exposes a frame-interpolation stage."
+      frame_interpolation_model_path: "vLLM-Omni post-processing field; unavailable until FastVideo exposes a frame-interpolation stage."
 
 cli:
   notes:

docs/design/server_contracts/openai.md

@@ -1,116 +1,170 @@
-# OpenAI-compatible HTTP Contract
+# OpenAI-compatible HTTP contract
 
-The stateless FastVideo HTTP server lives at
-[`fastvideo/entrypoints/openai/`](https://github.com/hao-ai-lab/FastVideo/tree/main/fastvideo/entrypoints/openai).
-Launch: `fastvideo serve --config serve.yaml`.
+FastVideo exposes one model-agnostic REST engine for image and video models.
+Launch it from a typed serve config:
+
+```bash
+fastvideo serve --config examples/serving/openai_fasth3.yaml
+```
+
+All generation routes share one serialized engine. FastVideo pipelines mutate
+per-request sampling state, and some adapters merge weights at load time, so a
+single loaded pipeline is never entered concurrently by image and video
+requests. HTTP handling and job polling remain asynchronous.
 
 ## Endpoints
 
 | Method | Path | Description |
 | --- | --- | --- |
-| `POST` | `/v1/videos/generations` | Synchronous video generation |
-| `GET` | `/v1/videos` | List prior jobs held in the in-memory store |
-| `GET` | `/v1/videos/{id}` | Job status / result |
-| `GET` | `/v1/videos/{id}/content` | Download the MP4 once ready |
-| `POST` | `/v1/images/generations` | Synchronous image generation |
-| `GET` | `/v1/models` | Enumerate registered models |
+| `GET` | `/v1/models` | List the served model and optional startup adapter |
+| `GET` | `/v1/models/{model}` | Retrieve one served model card |
+| `POST` | `/v1/videos` | Submit an asynchronous video job |
+| `POST` | `/v1/videos/sync` | Generate and return an MP4 response directly |
+| `GET` | `/v1/videos` | List in-memory jobs with `after`, `limit`, and `order` |
+| `GET` | `/v1/videos/{id}` | Retrieve job status and metadata |
+| `GET` | `/v1/videos/{id}/content` | Download a completed MP4 |
+| `DELETE` | `/v1/videos/{id}` | Delete a job and its completed artifact |
+| `POST` | `/v1/images` | Generate an image |
+| `POST` | `/v1/images/edits` | Generate an image from image references |
+| `GET` | `/v1/images/{id}/content` | Download a generated image |
 | `GET` | `/health` | Liveness probe |
 
-## `VideoGenerationsRequest` shape
+`POST /v1/videos/generations` remains an alias for older FastVideo clients.
 
-Mirrors the OpenAI `POST /v1/videos/generations` shape:
+## Video requests
+
+The canonical shape follows vLLM-Omni and accepts SGLang's common flat
+extensions. Fields that FastVideo cannot represent for the loaded model fail
+at admission with HTTP 400 instead of creating a job that later fails.
 
 ```json
 {
-  "prompt": "a fox running through snow",
-  "size": "1024x1536",
-  "seconds": 5,
-  "fps": 24,
-  "num_frames": 121,
+  "model": "fasth3",
+  "prompt": "A fox runs through fresh snow.",
+  "seconds": "5",
+  "size": "1344x768",
+  "video_params": {
+    "fps": 24,
+    "num_frames": 124
+  },
   "seed": 42,
-  "num_inference_steps": 8,
+  "num_inference_steps": 5,
   "guidance_scale": 1.0,
-  "negative_prompt": "blurry, low quality",
-  "input_reference": "/path/to/init.png"
-}
-```
-
-SGLang-compatible extensions carried today:
-`num_inference_steps`, `guidance_scale`, `guidance_scale_2`,
-`true_cfg_scale`, `negative_prompt`, `enable_teacache`, `output_path`.
-
-## Merge precedence
-
-The server builds a `GenerationRequest` each call using three layers,
-highest first:
-
-1. **Request body (client-explicit)** — only fields carried in
-   `request.model_fields_set` (Pydantic v2). Unset fields do not count,
-   even if the Pydantic model has a schema default for them.
-2. **`ServeConfig.default_request` (operator-explicit)** — projected via
-   [`explicit_request_updates()`](https://github.com/hao-ai-lab/FastVideo/blob/main/fastvideo/api/compat.py);
-   only fields the operator actually wrote into the YAML count as
-   defaults. Every other field inherits the schema default rather than
-   being pinned.
-3. **Hardcoded fallback** — e.g. `fps = 24`.
-
-The gate matters: both surfaces carry schema defaults. Without
-`model_fields_set` / explicit-path tracking, schema defaults would
-masquerade as intent and silently shadow the other side.
-
-See [`video_api.py::_build_generation_kwargs`](https://github.com/hao-ai-lab/FastVideo/blob/main/fastvideo/entrypoints/openai/video_api.py)
-for the canonical implementation; the per-request assembly lives there,
-not in pipeline code.
-
-## Continuation state
-
-The stateless surface accepts an opaque `ContinuationState` round-trip.
-Clients that want continuation pass the prior `state` blob back on the
-next request, and receive a new one on the response when
-`request.output.return_state = true`.
-
-Shape:
-
-```json
-{
-  "state": {
-    "kind": "ltx2.v1",
-    "payload": { "schema_version": 1, "segment_index": 3, ... }
+  "image_reference": [
+    {"image_url": "https://example.com/first-frame.png"}
+  ],
+  "extra_params": {
+    "vsa_mode": "exempt"
   }
 }
 ```
 
-Payload is always JSON-serializable. Large tensors may live in an
-opaque blob-store reference the client simply round-trips; see
-[`LTX2ContinuationState`](https://github.com/hao-ai-lab/FastVideo/blob/main/fastvideo/pipelines/basic/ltx2/continuation.py).
+Resolution precedence matches vLLM-Omni:
 
-Continuation is not yet wired all the way through to
-`generator.generate_video(...)` — PR 7.6 (GPU pool upstream) is the
-pipeline-level consumer. PR 7 locked the envelope so this surface is
-stable ahead of that plumbing.
+1. `size`
+2. top-level `width` and `height`
+3. `video_params.width` and `video_params.height`
 
-## Error codes
+Top-level `fps` and `num_frames` similarly take precedence over the nested
+block. If `num_frames` is absent, `seconds * fps` is used. FastVideo also keeps
+the legacy `input_reference`, `reference_url`, `video_path`, and `video_url`
+spellings.
 
-| HTTP | Condition |
-| --- | --- |
-| `400 Bad Request` | Parse/validation failure (unknown field, type mismatch, incompatible preset/state) |
-| `404 Not Found` | `GET /v1/videos/{id}` for an unknown job |
-| `409 Conflict` | Job id already exists |
-| `500 Internal Server Error` | Pipeline raised; body mirrors upstream OpenAI error envelope |
-| `503 Service Unavailable` | No generator loaded, or shutdown in progress |
+Reference objects support URL or local-path strings through `image_url`,
+`video_url`, and `audio_url`. `file_id` references are schema-compatible but
+return HTTP 400 because FastVideo does not provide an OpenAI Files store.
+Image URLs, data URLs, local paths, and multipart `input_reference` uploads are
+materialized and decoded under the configured output directory during
+admission. Invalid media returns HTTP 400 before a job is created.
 
-Errors include a JSON body with
-`{"error": {"type": "...", "message": "..."}}` matching the OpenAI
-Python SDK's expectation.
+## Jobs and synchronous responses
 
-## What does not cross this boundary
+An asynchronous submission returns a `video` object in `queued` state. Its
+status advances through `in_progress` to `completed` or `failed`. Completed
+jobs expose `file_name`, the FastVideo compatibility extension `file_path`,
+timings, and peak-memory metadata when the pipeline reports them.
 
-* Flat legacy kwargs (`ltx2_refine_enabled`, `torch_compile_kwargs`,
-  etc.) — these are init-time, configured via `ServeConfig.generator`,
-  never per-request.
-* Private Dreamverse-only fields — those live in a private adapter on
-  the Dreamverse side; the public FastVideo surface never promises
-  backward compatibility for them.
-* Raw tensor payloads (`ltx2_audio_clean_latent` et al.) — these are
-  derived by the pipeline from `ContinuationState`, never shipped as
-  request fields.
+`POST /v1/videos/sync` returns `video/mp4` bytes. It includes
+`X-Request-Id`, `X-Model`, `X-Inference-Time-S`, `X-Stage-Durations`, and
+`X-Peak-Memory-MB` headers. Its temporary MP4 is removed after the response is
+streamed. Asynchronous artifacts remain available until their job is deleted.
+
+Output paths are controlled by the server. Clients cannot choose filesystem
+destinations; every video is written beneath `server.output_dir` with a unique
+request id.
+
+FastVideo's synchronous CUDA execution cannot be interrupted after launch.
+Deleting an in-progress resource removes it from the API immediately; the
+engine remains serialized until the call exits and then removes any artifact.
+
+## Model and LoRA selection
+
+`server.served_model_name` controls the public model id. If omitted, the
+checkpoint path is used. Requests that name another model fail with HTTP 400.
+
+LoRAs are configured under
+`generator.pipeline.components.{lora_path,lora_nickname,lora_strength}`. The
+startup adapter is the only model advertised by a LoRA server, and requests can
+select it by its model nickname or with a selector:
+
+```json
+{
+  "prompt": "A fox runs through fresh snow.",
+  "model": "fasth3-dense-datafree",
+  "lora": {
+    "name": "fasth3-dense-datafree",
+    "path": "/models/adapter_model.safetensors",
+    "scale": 1.0
+  }
+}
+```
+
+The selector must match the adapter already loaded at startup. FastH3 adapter
+files can contain dense replacement tensors and VSA gates in addition to
+low-rank factors, so swapping them inside concurrent requests would corrupt
+shared pipeline state. A mismatch is rejected with HTTP 400.
+
+## MiniMax-H3 and FastH3
+
+FastH3 uses the same general routes and adapter. `task` is accepted for
+SGLang-compatible H3 clients:
+
+- `t2va` uses text only.
+- `fl2va` takes one or two image references.
+- `ref2va` takes ordered image, video, and audio references and requires a
+  server started with `MiniMaxH3Ref2VAModularPipeline`.
+
+The released FastH3 pipeline generates one packed video/audio result per
+request, uses 24 fps, requires guidance scale 1, and accepts frame counts on
+its causal-VAE grid. The serving examples pin its five-point distilled sigma
+schedule (four DiT forwards).
+
+## Defaults and errors
+
+Incoming explicit fields override operator-explicit `default_request` fields,
+which override model preset defaults. Pydantic defaults do not masquerade as
+client intent; the transport uses `model_fields_set`, while typed config parsing
+tracks the exact paths written by the operator.
+
+Errors use the OpenAI envelope:
+
+```json
+{
+  "error": {
+    "message": "...",
+    "type": "invalid_request_error",
+    "param": null,
+    "code": 400
+  }
+}
+```
+
+Parse, model-selection, startup-LoRA, and unsupported-parameter failures are
+HTTP 400; missing resources are HTTP 404; generation failures are stored on
+asynchronous jobs and returned as HTTP 500 when that job is retrieved.
+Unknown top-level fields are rejected. `extra_params` accepts only the explicit
+request-batch passthrough fields supported by the typed request adapter.
+
+`GET /health` also verifies that the generation engine is open and all local
+multiprocess workers are alive. It returns HTTP 503 when the worker pool is no
+longer usable.

docs/generate_examples.py

@@ -2,6 +2,7 @@
 # adapted from vllm: https://github.com/vllm-project/vllm/blob/v0.7.3/docs/source/generate_examples.py
 
 import itertools
+import json
 import os
 import re
 from dataclasses import dataclass, field
@@ -19,6 +20,40 @@ GENERATED_DOC_PREFIXES = (
     "training/examples/",
     "distillation/examples/",
 )
+COOKBOOK_DATA = ROOT_DIR / "docs/assets/cookbook-recipes.json"
+COOKBOOK_SOURCE_ROOTS = (
+    ROOT_DIR / "examples/inference",
+    ROOT_DIR / "scripts/inference",
+)
+
+
+def validate_cookbook() -> None:
+    """Keep cookbook entries tied to checked-in runnable sources."""
+    recipes = json.loads(COOKBOOK_DATA.read_text(encoding="utf-8")).get("recipes")
+    if not isinstance(recipes, list) or not recipes:
+        raise ValueError(f"{COOKBOOK_DATA}: recipes must be a non-empty list")
+
+    seen: set[str] = set()
+    for recipe in recipes:
+        required = ("id", "task", "label", "model", "source", "command")
+        missing = {key for key in required if not recipe.get(key)}
+        if missing:
+            raise ValueError(f"Cookbook recipe is missing: {', '.join(sorted(missing))}")
+        if recipe["id"] in seen:
+            raise ValueError(f"Duplicate cookbook recipe id: {recipe['id']}")
+        seen.add(recipe["id"])
+
+        source = (ROOT_DIR / recipe["source"]).resolve()
+        if not any(source.is_relative_to(root.resolve()) for root in COOKBOOK_SOURCE_ROOTS):
+            raise ValueError(f"Cookbook source is outside an approved directory: {recipe['source']}")
+        if not source.is_file():
+            raise ValueError(f"Cookbook source does not exist: {recipe['source']}")
+
+        source_text = source.read_text(encoding="utf-8")
+        if recipe["model"] not in source_text:
+            raise ValueError(f"Cookbook model is not present in {recipe['source']}: {recipe['model']}")
+        if recipe["source"] not in recipe["command"]:
+            raise ValueError(f"Cookbook command does not invoke its source: {recipe['id']}")
 
 
 def fix_case(text: str) -> str:
@@ -536,6 +571,7 @@ def on_pre_build(config, **kwargs):
     MkDocs hook to generate examples before building the documentation.
     This function is called automatically by MkDocs' native hook system.
     """
+    validate_cookbook()
     print("Generating example documentation...")
     generate_examples(generate_main_index=True)
     print("Example documentation generated successfully!")
@@ -549,6 +585,7 @@ def on_page_context(context, page, **kwargs):
 
 
 if __name__ == "__main__":
+    validate_cookbook()
     print("Generating example documentation...")
     generate_examples(generate_main_index=True)
     print("Example documentation generated successfully!")

docs/getting_started/installation.md

@@ -65,6 +65,7 @@ uv pip install flash-attn --no-build-isolation -v
 
 ## Next Steps
 
-- [Quick Start Guide](quick_start.md) - Get started with your first video generation
+- [Quick Start](quick_start.md) - Generate your first video
+- [Inference Cookbook](../cookbook/index.md) - Choose a maintained recipe
 - [Configuration](../inference/configuration.md) - Learn about configuration options
-- [Examples](../inference/examples/examples_inference_index.md) - Explore example scripts and notebooks
+- [Examples](../inference/examples/examples_inference_index.md) - Explore scripts and notebooks

docs/getting_started/installation/mps.md

@@ -1,6 +1,11 @@
 # MPS (Apple Silicon)
 
-Instructions to install FastVideo for Apple Silicon.
+Install FastVideo on Apple Silicon and run FastMetal-QAD or FastH3 Preview.
+
+Apple Silicon uses the MLX runtime. FastMetal-QAD ships ready-to-run MLX
+checkpoints; FastH3 Preview currently requires a local MLX DiT conversion.
+See the [FastMetal-QAD blog](https://haoailab.com/blogs/fastmetal/) and the
+[FastMetal collection](https://huggingface.co/collections/FastVideo/fastmetal).
 
 ## Requirements
 
@@ -49,7 +54,7 @@ brew install ffmpeg
 
 ### Installation
 
-FastWan's native Apple Silicon runtime requires the `mlx` extra.
+FastMetal's native Apple Silicon runtime requires the `mlx` extra.
 
 #### With uv (recommended)
 
@@ -87,6 +92,98 @@ Alternative with Conda environment:
 uv pip install -e ".[mlx]"
 ```
 
+## Run FastMetal-QAD
+
+Each release is self-contained. Download one checkpoint and point both
+`--model-root` and `--mlx-checkpoint` at it (the example also auto-detects
+`mlx_dit.json` under `--model-root`).
+
+| Checkpoint | Script | Mac tier |
+| --- | --- | --- |
+| [`FastVideo/FastMetal-1.3B-QAD`](https://huggingface.co/FastVideo/FastMetal-1.3B-QAD) | `mlx_wan_prompt_to_video.py` | 16 GB+ |
+| [`FastVideo/FastMetal-5B-QAD`](https://huggingface.co/FastVideo/FastMetal-5B-QAD) | `mlx_wan22_generate.py` | 16 GB+ |
+| [`FastVideo/FastMetal-14B-QAD`](https://huggingface.co/FastVideo/FastMetal-14B-QAD) | `mlx_wan_prompt_to_video.py` | 36 GB+ |
+
+```bash
+hf download FastVideo/FastMetal-1.3B-QAD --local-dir ./FastMetal-1.3B-QAD
+
+python examples/inference/basic/mlx_wan_prompt_to_video.py \
+  --model-root ./FastMetal-1.3B-QAD \
+  --mlx-checkpoint ./FastMetal-1.3B-QAD \
+  --height 480 --width 832 --num-frames 81 \
+  --prompt "A bird's-eye view of a misty forest valley at dawn."
+```
+
+14B uses the same script. Point both flags at `./FastMetal-14B-QAD`. That repo also ships an EMA variant: keep `--model-root` at the repo root and set `--mlx-checkpoint ./FastMetal-14B-QAD/ema`.
+
+Wan2.2 5B uses a different latent layout, so it has its own entrypoint:
+
+```bash
+hf download FastVideo/FastMetal-5B-QAD --local-dir ./FastMetal-5B-QAD
+
+python examples/inference/basic/mlx_wan22_generate.py \
+  --mlx-checkpoint ./FastMetal-5B-QAD \
+  --text-encoder-root ./FastMetal-5B-QAD \
+  --vae-root ./FastMetal-5B-QAD/vae \
+  --height 704 --width 1280 --num-frames 81 \
+  --prompt "A cinematic portrait with soft neon lighting and smooth camera motion."
+```
+
+CUDA FastWan-QAD (`FastVideo/FastWan-QAD-1.3B`, `FastVideo/FastWan-QAD-FP8-1.3B`) is a separate NVIDIA release. The MLX examples look for FastMetal packed weights (`mlx_dit.json`).
+
+`basic_mps.py` is a generic PyTorch MPS demo. For local video on Mac, use the FastMetal commands above.
+
+## Run FastH3 Preview
+
+FastH3 Preview uses the existing MLX runtime for text-to-video-with-audio
+(T2VA). The runtime streams the Qwen3-VL text conditioner, loads one
+heavyweight component at a time, denoises synchronized video and audio
+latents with a converted INT8, INT6, or INT4 DiT, and decodes both modalities
+with native MLX VAEs.
+
+Download the FastH3 snapshot, then convert one or more DiT formats:
+
+```bash
+hf download FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2 \
+  --local-dir ./FastH3-Preview-v0.2
+
+python scripts/checkpoint_conversion/convert_minimax_h3_mlx.py \
+  --model-root ./FastH3-Preview-v0.2/transformer \
+  --out ./FastH3-MLX \
+  --formats "int6"
+```
+
+Run the baseline path:
+
+```bash
+python examples/inference/basic/mlx_fasth3.py \
+  --model-root ./FastH3-Preview-v0.2 \
+  --mlx-checkpoint ./FastH3-MLX/int6 \
+  --prompt "(S1) A presenter says <d>[English] Fast H3 is amazing.</d>" \
+  --height 480 --width 832 --num-frames 124 --seed 2026 \
+  --output-path ./outputs/fasth3_int6.mp4
+```
+
+Add `--fast` for temporal fast mode. It denoises a shorter video sequence,
+uses MLX RIFE to restore the requested frame count, and keeps the audio
+sequence at full duration:
+
+```bash
+python examples/inference/basic/mlx_fasth3.py \
+  --model-root ./FastH3-Preview-v0.2 \
+  --mlx-checkpoint ./FastH3-MLX/int6 \
+  --prompt "(S1) A presenter says <d>[English] Fast H3 is even faster.</d>" \
+  --height 720 --width 1280 --num-frames 124 --seed 2027 \
+  --fast \
+  --output-path ./outputs/fasth3_int6_fast_720p.mp4
+```
+
+!!! note "Current MLX scope"
+    This source runtime supports T2VA and temporal `--fast`. FL2VA, Ref2VA,
+    spatial fast mode, two-pass refinement, VSA, and `VideoGenerator`
+    registry dispatch are not wired yet. The checkpoint uses the MiniMax H3
+    Community License; review the model card before use or redistribution.
+
 ## Development Environment Setup
 
 If you're planning to contribute to FastVideo please see the following page:
@@ -94,9 +191,12 @@ If you're planning to contribute to FastVideo please see the following page:
 
 ## Hardware Requirements
 
-### For Basic Inference
-
-- Mac M1, M2, M3, or M4 (at least 32 GB RAM is preferable for high quality video generation)
+- **1.3B / 5B:** 16 GB unified memory and up (M1 and later)
+- **14B:** 36 GB unified memory and up
+- **FastH3 Preview:** validated on an M4 Max with 36 GB unified memory; use one
+  converted DiT format at a time and leave substantial free disk space for the
+  source snapshot plus the converted checkpoint
+- Fanless 13-inch MacBook Air can run 1.3B and 5B at the same resolutions
 
 ## Troubleshooting
 

docs/getting_started/installation/spark_performance.md

@@ -156,8 +156,11 @@ is power-cycled. To avoid it:
 
 - **Builds** (flash-attn, kernel): `nice -n 19`, `MAX_JOBS=2`, `nohup`. Never a
   bare foreground high-parallelism build.
-- Leave `*_cpu_offload` at the example defaults — "CPU" offload is the *same*
-  unified RAM on the GB10, so the win is tiling + sane resolution, not offloading.
+- FastVideo automatically disables DiT layerwise/CPU offload and encoder/VAE CPU
+  offload after each worker binds its GB10 device. Do not force those modes back
+  on: "CPU" offload uses the same unified RAM. Multi-GPU FSDP sharding remains
+  available because it partitions weights without parking them in a separate
+  host pool.
 
 ## Gotchas specific to the GB10
 

docs/getting_started/quick_start.md

@@ -23,61 +23,21 @@ Also optionally install flash-attn:
 uv pip install flash-attn --no-build-isolation -v
 ```
 
-## Basic Usage
+## Choose a maintained recipe
 
-### Text-to-Video Generation
+The cookbook selects complete, checked-in recipes instead of mixing model,
+parallelism, offload, and attention settings independently.
 
-```python
-from fastvideo import VideoGenerator
+[Open the inference cookbook](../cookbook/index.md){ .md-button .md-button--primary }
 
-def main():
-    # Create a video generator with a pre-trained model
-    generator = VideoGenerator.from_pretrained(
-        "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
-        num_gpus=1,  # Adjust based on your hardware
-    )
-
-    # Define a prompt for your video
-    prompt = "A curious raccoon peers through a vibrant field of yellow sunflowers, its eyes wide with interest."
-
-    # Generate the video
-    video = generator.generate_video(
-        prompt,
-        output_path="my_videos/",  # Controls where videos are saved
-        save_video=True
-    )
-
-if __name__ == '__main__':
-    main()
-```
-
-### Image-to-Video Generation
-
-```python
-from fastvideo import VideoGenerator, SamplingParam
-
-def main():
-    # Create the generator
-    model_name = "Wan-AI/Wan2.1-I2V-14B-480P-Diffusers"
-    generator = VideoGenerator.from_pretrained(model_name, num_gpus=1)
-
-    # Set up parameters with an initial image
-    sampling_param = SamplingParam.from_pretrained(model_name)
-    sampling_param.image_path = "https://huggingface.co/datasets/huggingface/documentation-images/resolve/main/diffusers/astronaut.jpg"
-    sampling_param.num_frames = 107
-
-    # Generate video based on the image
-    prompt = "A photograph coming to life with gentle movement"
-    generator.generate_video(prompt, sampling_param=sampling_param,
-                             output_path="my_videos/",
-                             save_video=True)
-
-if __name__ == '__main__':
-    main()
-```
+!!! tip "Need more control?"
+    Start from a maintained recipe, then use the
+    [configuration](../inference/configuration.md) and
+    [optimization](../inference/optimizations.md) guides for supported changes.
 
 ## Next Steps
 
+- [Inference Cookbook](../cookbook/index.md) - Choose a maintained recipe
 - [Installation Guide](installation.md) - Detailed installation instructions
 - [Configuration](../inference/configuration.md) - Learn about configuration options
 - [Examples](../inference/examples/examples_inference_index.md) - Explore more

docs/inference/offloading.md

@@ -14,6 +14,13 @@ vae_cpu_offload: bool = True
 pin_cpu_memory: bool = True
 ```
 
+On unified-memory accelerators such as NVIDIA GB10 and Apple silicon, FastVideo
+detects the selected device inside each worker and disables all five host-offload
+modes before loading modules. Host and accelerator allocations share one physical
+pool there, so offload adds transfers and duplicate residency instead of freeing
+memory. CUDA FSDP sharding remains enabled when requested; MPS continues to
+disable FSDP. `pin_cpu_memory` is not an offload mode and is left unchanged.
+
 ## Behavior Explanation
 
 !!! note

docs/inference/optimizations.md

@@ -88,6 +88,7 @@ runtime on some GPU/shape combinations. To use FA4, install the pinned
 `flash-attn-4` build (see the `flash-attn-4` source in `pyproject.toml`) and set:
 
 ```bash
+UV_TORCH_BACKEND=cu130 uv pip install -e ".[fasth3]"
 export FASTVIDEO_FA4=1
 ```
 
@@ -97,6 +98,21 @@ sm90+) and GQA attention (FA4's `pack_gqa` fails to JIT-compile below sm90).
 On sm90+ both run on FA4. If FA4 is unusable while `FASTVIDEO_FA4=1` is set,
 FastVideo fails loudly instead of silently falling back.
 
+MiniMax-H3 can additionally use FA4's packed-varlen entry point for its long,
+single-sequence dense DiT self-attention:
+
+```bash
+export FASTVIDEO_FA4=1
+export FASTVIDEO_MINIMAX_H3_FA4_PACKED_VARLEN=1
+```
+
+This route is inference-only and remains disabled by default. Runtime guards
+keep masked attention, batch sizes above one, unequal query/key lengths,
+grad-enabled calls, and NVFP4 on their established paths. It does not apply to
+the Preview checkpoint's sparse VSA blocks. Packed-varlen changes floating-point
+reduction order relative to fixed-length FA4, so treat it as a speed/quality
+evaluation option rather than an exact-parity mode.
+
 ### FP4 Flash Attention 4 (Blackwell only)
 
 **`FLASH_ATTN`** with **`--nvfp4_fa4`**
@@ -337,7 +353,42 @@ Only DiT submodules that declare `_compile_conditions` are compiled
 (most shipped models). The text encoder and VAE are not compiled by this
 flag.
 
-### What to expect
+### Regional fullgraph compile (experimental)
+
+`inference_torch_compile` is a stricter, kwargs-free variant that ports the
+training-side regional compile of
+[#1718](https://github.com/hao-ai-lab/FastVideo/pull/1718) to inference: the
+loader wraps each `_compile_conditions` block in
+`torch.compile(fullgraph=True)` with inductor
+`options={"emulate_precision_casts": True}` right after the transformer
+loads. The ordinary compile path keeps its historical compiler-disabled
+attention boundary by default; regional compile opts in only the compatible
+attention instances owned by this transformer. MiniMax-H3 VSA is supported
+only by the inference-only sm_100a tile-64 route
+(`FASTVIDEO_VSA_SM100A=1` and `VSA_tile_size=64`); the loader probes that
+route before capture and keeps the transformer eager when the kernel or
+device is unsupported. Legacy VSA, MiniMax-H3 tile-256 VSA, and the explicit
+`FASTVIDEO_DISABLE_ATTENTION_COMPILE=1` escape hatch keep the transformer
+eager with one warning instead of failing mid-denoise.
+
+```python
+generator = VideoGenerator.from_pretrained(
+    "MiniMaxAI/MiniMax-H3",
+    inference_torch_compile=True,  # or FASTVIDEO_INFERENCE_TORCH_COMPILE=1
+)
+```
+
+Do not combine it with `torch_compile_kwargs['mode']` (the loader injects
+inductor options, and torch.compile forbids mode+options); it is
+independent of `enable_torch_compile`, and when both are set the regional
+compile wins for the DiT.
+
+### What to expect from generic compile
+
+The Wan result below measures the existing generic
+`enable_torch_compile=True` path. It is useful evidence that compile can help,
+but it is **not** a benchmark or numerical gate for the stricter regional
+fullgraph path above.
 
 | Config | Effect |
 |---|---|
@@ -353,6 +404,24 @@ generations with the same input shapes. Always exclude the first
 (warmup) generation when measuring steady-state latency — measuring the
 warmup is the most common way to wrongly conclude "compile is slower".
 
+**Regional MiniMax-H3 accuracy caveat (job 2660).** On one GB200 at
+768×1344×124, the native 50-point schedule ran exactly 49 transformer
+forwards. After one warmup, three fixed-prompt/fixed-seed repeats averaged
+**185.08s → 157.01s end to end** and **174.90s → 147.12s denoising**. Each
+leg was independently pixel-deterministic, but compiled output did **not**
+match eager: mean absolute pixel error **20.247/255**, PSNR **16.67 dB**,
+mean SSIM **0.7108**, and mean MS-SSIM **0.6370** across 124 frames. Treat
+regional MiniMax-H3 compile as an opt-in performance experiment, not an
+eager-parity-safe mode.
+
+The same caveat applies to sparse MiniMax-H3 regional compile. Its mask
+compaction, sm_100a launch, trained compression gates, and inference-only H3
+fusions are fullgraph-compatible, but compilation can still change model
+numerics. The FastH3 `all` profile enables regional compile by default because
+it is the fastest measured route at 124, 243, and 345 frames. Use
+`--no-inference-torch-compile` when comparing against the eager sparse-DiT
+route.
+
 **Numerics.** Inductor's lowering is designed to preserve eager
 semantics within floating-point tolerance, but per-model equivalence is
 not asserted by any standing SSIM regression here — the SSIM tests in
@@ -371,9 +440,14 @@ MS-SSIM gate on *your* config, especially when combining
   hao-ai-lab/FastVideo#1365 — keep that fix to get a clean compiled
   region under the default offload path.
 - **`mode="reduce-overhead"` / CUDA graphs**: not yet supported
-  end-to-end. The attention dispatch is an untraceable custom op and
-  still breaks the graph, which CUDA-graph trees cannot span. Use the
-  default inductor mode (shown above) until that is resolved.
+  by regional compile because that path injects inductor `options`, and
+  PyTorch rejects `mode` together with `options`. The generic compile path
+  can accept `mode`, but CUDA-graph compatibility remains backend- and
+  shape-dependent. FA2/FA3 inference and FA4 expose traceable custom-op
+  boundaries. MiniMax-H3's sm_100a tile-64 inference route is the only VSA
+  path in the regional support envelope; other VSA paths and the FA3
+  grad-enabled path remain outside it. Use the default inductor mode shown
+  above unless your exact configuration has its own gate.
 
 Extra `torch.compile` options are passed through `torch_compile_kwargs`
 (a dict), accepted by `VideoGenerator.from_pretrained(...)` and by the

docs/inference/support_matrix.md

@@ -182,12 +182,17 @@ optimizations: absence means **untested**, not incompatible.
 
 | Release path | Model | Mode | Validated hardware | Status |
 | --- | --- | --- | --- | --- |
-| MLX FastWan T2V | FastWan-QAD-INT8-1.3B `[release model ID pending]` | 480x832, 81 frames, 3-step DMD, INT8 DiT + TAEHV decode | Apple M4 Max, 36 GB unified-memory class, MLX 0.31.2 | Release candidate; requires release-owner visual sign-off |
+| MLX FastMetal T2V 1.3B | [`FastVideo/FastMetal-1.3B-QAD`](https://huggingface.co/FastVideo/FastMetal-1.3B-QAD) | 480x832, 81 frames, 3-step DMD, INT8 DiT + TAEHV decode | Apple M4 Max, 16 GB+ unified memory | Released |
+| MLX FastMetal TI2V 5B | [`FastVideo/FastMetal-5B-QAD`](https://huggingface.co/FastVideo/FastMetal-5B-QAD) | 480p / 720p, 81 frames, 3-step DMD, INT8 DiT + TAEHV decode | Apple M4 Max, 16 GB+ unified memory | Released |
+| MLX FastMetal T2V 14B | [`FastVideo/FastMetal-14B-QAD`](https://huggingface.co/FastVideo/FastMetal-14B-QAD) | 480p / 720p, 81 frames, 3-step DMD, INT8 DiT + TAEHV decode | Apple M4 Max, 36 GB+ unified memory | Released |
+| MLX FastH3 Preview T2VA | [`FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2`](https://huggingface.co/FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2) + locally converted DiT | 480p / 720p, 124 frames, 4-step DMD2, INT8/INT6/INT4 DiT, native video + audio VAE; optional temporal RIFE fast mode | Apple M4 Max, 36 GB unified memory | Source runtime; T2VA only |
 
-This is a text-to-video-only source-install release. It is validated on the
-hardware listed above; MLX allocator caps are not evidence of support for a
-physical 16 GB Mac. See [Apple Silicon FastWan](../getting_started/installation/mps.md)
-for the supported command and release gates.
+Apple Silicon uses the native MLX runtime. FastMetal-QAD is the packaged Wan
+release, while FastH3 Preview currently uses a source checkout and local DiT
+conversion. CUDA FastWan-QAD (`FastVideo/FastWan-QAD-1.3B`,
+`FastVideo/FastWan-QAD-FP8-1.3B`) is the NVIDIA release. See the
+[Apple Silicon guide](../getting_started/installation/mps.md) and the
+[FastMetal-QAD blog](https://haoailab.com/blogs/fastmetal/).
 
 **Note**: Wan2.2 TI2V 5B has some quality issues when performing I2V generation. We are working on fixing this issue.
 
@@ -216,9 +221,10 @@ Per the installation guides:
   [GPU install guide](../getting_started/installation/gpu.md).
 - **NVIDIA DGX Spark (GB10, aarch64)** — CUDA 13, from-source kernel build; see
   the [DGX Spark install guide](../getting_started/installation/spark.md).
-- **Apple silicon (MPS)** — macOS 14 or newer; see the
-  [MPS install guide](../getting_started/installation/mps.md) and
-  [`basic_mps.py`](https://github.com/hao-ai-lab/FastVideo/blob/main/examples/inference/basic/basic_mps.py).
+- **Apple silicon** — macOS 14 or newer; FastMetal-QAD via the MLX runtime. See the
+  [Apple Silicon guide](../getting_started/installation/mps.md). The older
+  [`basic_mps.py`](https://github.com/hao-ai-lab/FastVideo/blob/main/examples/inference/basic/basic_mps.py)
+  demo is PyTorch MPS only.
 
 Optimization-specific hardware constraints (e.g. STA requiring Hopper) are
 listed under [Special requirements](#special-requirements).

examples/inference/basic/README.md

@@ -18,11 +18,44 @@ git clone https://github.com/hao-ai-lab/FastVideo.git && cd FastVideo
 python examples/inference/basic/basic.py
 ```
 
-For an example on Apple silicon: 
+### Apple Silicon (FastMetal-QAD)
+
+Use the MLX runtime with FastMetal-QAD. See the
+[Apple Silicon guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mps/).
+
+```bash
+hf download FastVideo/FastMetal-1.3B-QAD --local-dir ./FastMetal-1.3B-QAD
+
+python examples/inference/basic/mlx_wan_prompt_to_video.py \
+  --model-root ./FastMetal-1.3B-QAD \
+  --mlx-checkpoint ./FastMetal-1.3B-QAD \
+  --prompt "A bird's-eye view of a misty forest valley at dawn."
 ```
-python examples/inference/basic/basic_mps.py
+
+5B uses
+[`mlx_wan22_generate.py`](https://github.com/hao-ai-lab/FastVideo/blob/main/examples/inference/basic/mlx_wan22_generate.py)
+with
+`FastVideo/FastMetal-5B-QAD`.
+
+`examples/inference/basic/basic_mps.py` is the older PyTorch MPS demo.
+
+FastH3 Preview T2VA also runs through the native MLX runtime. Convert the DiT
+to INT8, INT6, or INT4 first, then run:
+
+```bash
+python examples/inference/basic/mlx_fasth3.py \
+  --model-root ./FastH3-Preview-v0.2 \
+  --mlx-checkpoint ./FastH3-MLX/int6 \
+  --prompt "(S1) A presenter says <d>[English] Fast H3 is amazing.</d>" \
+  --height 480 --width 832 --num-frames 124 \
+  --output-path ./outputs/fasth3_int6.mp4
 ```
 
+Pass `--fast` for temporal RIFE fast mode. This MLX entrypoint currently
+supports T2VA only; FL2VA, Ref2VA, spatial fast mode, and two-pass refinement
+remain follow-up work. The complete setup and conversion commands are in the
+[Apple Silicon guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mps/).
+
 For an example running DMD+VSA inference:
 ```
 python examples/inference/basic/basic_dmd.py
@@ -33,6 +66,96 @@ For the typed config/request path added during the inference API refactor:
 python examples/inference/basic/basic_dmd_new_api.py
 ```
 
+### FastH3 Preview
+
+The verified [basic FastH3 example](https://github.com/hao-ai-lab/FastVideo/blob/main/examples/inference/basic/basic_fasth3.py)
+runs the few-step (4-forward, DMD2-distilled) MiniMax-H3 preview, generating
+synchronized video and audio with its trained block-sparse VSA attention:
+
+```bash
+UV_TORCH_BACKEND=cu130 uv pip install -e ".[fasth3]"
+```
+
+This installs the pinned FA4 CuTe package and FastVideo kernel release used by
+the measured GB200 profile. Then run:
+
+```
+python examples/inference/basic/basic_fasth3.py --prompt "your prompt"
+```
+The default checkpoint, [FastH3 Preview v0.2](https://huggingface.co/FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2), is public on the Hub under the MiniMax H3 Community License. Review its model card and license before use or redistribution.
+
+The default `all` profile is the fastest measured four-GPU Preview recipe on GB200. It selects VSA sparsity 0.9 with 64-token tiles and the sm_100a sparse kernel, enables FA4 for eligible non-VSA paths, regionally compiles and replicates the sparse DiT, compiles and temporally parallelizes the video VAE with the `gather` strategy, and pins CPU-offloaded component memory. It also pins the benchmark protocol: five sigma-grid points (exactly four DiT forwards), one excluded seed-999 warmup, then three timed seed-1000 requests with distinct output paths.
+
+The equivalent explicit command is:
+
+```bash
+python examples/inference/basic/basic_fasth3.py \
+  --prompt "your prompt" \
+  --profile all \
+  --num-gpus 4 \
+  --steps 5 \
+  --vsa-sparsity 0.9 \
+  --vsa-tile-size 64 \
+  --vsa-kernel sm100a \
+  --compile-vae \
+  --parallel-vae \
+  --replicated-dit \
+  --pin-cpu-memory \
+  --fa4 \
+  --no-torch-compile \
+  --inference-torch-compile \
+  --ulysses-a2a off \
+  --warmup \
+  --repeats 3 \
+  --seed 1000 \
+  --warmup-seed 999
+```
+
+`all` enables the inference-only H3 fusions and regional compile. Both can change floating-point operation order, so this is a report-only performance profile rather than an exact-parity route. Use `--profile strict` to disable the H3 fusions while preserving regional compile, or `--profile strict --no-inference-torch-compile` for the eager strict route. Individual `--no-*` switches are available for portability and attribution; in particular, use `--vsa-kernel triton --no-fa4` if the Blackwell kernels are unavailable. The script preserves the warmup and each measured video under distinct paths, then prints per-request wall time plus a warmup-excluded median.
+
+One script covers each validated duration; regional compile is the fastest
+measured DiT route for all three:
+
+```bash
+# 5 s
+python examples/inference/basic/basic_fasth3.py \
+  --prompt "your prompt" --output outputs/fasth3_5s
+# 10 s
+python examples/inference/basic/basic_fasth3.py \
+  --prompt "your prompt" --num-frames 243 --output outputs/fasth3_10s
+# 15 s
+python examples/inference/basic/basic_fasth3.py \
+  --prompt "your prompt" --num-frames 345 --output outputs/fasth3_15s
+```
+
+Pass `--no-inference-torch-compile` to recover the eager sparse-DiT route.
+
+### FastH3 Preview LoRAs
+
+The LoRA release runs on top of `MiniMaxAI/MiniMax-H3` with the same default
+compile, fusion, FA4, VSA, and parallel-VAE profile as the full FastH3 example:
+
+```bash
+bash examples/inference/basic/run_fasth3_lora_preview_vsa_datafree.sh \
+  --prompt "your prompt"
+```
+
+The four release launchers are:
+
+- `run_fasth3_lora_preview_vsa_datafree.sh`
+- `run_fasth3_lora_preview_vsa_synthetic_step1300.sh`
+- `run_fasth3_lora_preview_vsa_synthetic_step1900.sh`
+- `run_fasth3_lora_preview_dense_datafree.sh`
+
+Each downloads its exact private adapter file from
+`FastVideo/FastVideo-FastH3-4-step-Preview-v1-LoRA`; authenticate with `hf auth
+login` first. Pass `--lora-strength 0.5` to interpolate every adapter payload at
+half strength. Strength `1` applies the published rank-64 adapter at its trained
+scale and approximates the full student; `0` removes its weight deltas. VSA
+launchers still use sparse attention at strength `0` and require FastVideo's
+tile-64 VSA kernel; the dense launcher selects FA4. Each launcher writes to its
+own variant directory by default so comparison outputs do not collide.
+
 ## 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_fasth3.py

@@ -0,0 +1,367 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Few-step video+audio generation with the DMD2-distilled MiniMax H3 preview.
+
+The default ``all`` profile reproduces the fastest measured FastH3 Preview
+recipe on four GB200 GPUs. It runs the checkpoint's native five-point sigma
+grid (exactly four DiT forwards), trained VSA policy, Blackwell sparse kernel,
+regional fullgraph DiT compile, compiled/parallel video VAE, and inference-only
+H3 fusions. One compile warmup is excluded before three measured requests.
+
+Both regional compile and the default fusions can change floating-point
+operation order, so ``all`` is a report-only performance profile.
+``--profile strict`` disables the H3 fusions but preserves regional compile;
+combine it with ``--no-inference-torch-compile`` for the eager strict route.
+"""
+
+from __future__ import annotations
+
+import argparse
+import importlib.util
+import os
+import statistics
+import time
+from collections.abc import Sequence
+from pathlib import Path
+
+from fastvideo import VideoGenerator
+from fastvideo.api import (
+    CompileConfig,
+    ComponentConfig,
+    EngineConfig,
+    GenerationRequest,
+    GeneratorConfig,
+    OffloadConfig,
+    OutputConfig,
+    ParallelismConfig,
+    PipelineSelection,
+    SamplingConfig,
+)
+
+DEFAULT_MODEL = "FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2"
+
+
+def build_parser(description: str | None = None) -> argparse.ArgumentParser:
+    """Build the shared FastH3 preview CLI used by full and LoRA checkpoints."""
+    parser = argparse.ArgumentParser(description=description or __doc__)
+    parser.add_argument("--model-path", default=DEFAULT_MODEL)
+    # The HF repo may require authentication while the MiniMax H3 Community
+    # License review completes. A local snapshot can be passed here instead.
+    parser.add_argument("--prompt", required=True)
+    parser.add_argument("--output", default="outputs/fasth3")
+    parser.add_argument("--profile",
+                        choices=("all", "strict"),
+                        default="all",
+                        help="all enables the fastest measured, non-parity H3 fusions; strict disables only them")
+    parser.add_argument("--height", type=int, default=768)
+    parser.add_argument("--width", type=int, default=1344)
+    parser.add_argument("--num-frames", type=int, default=124)
+    # num_inference_steps counts sigma-GRID POINTS. The distilled schedule is
+    # t=1000,750,500,250 -> 0: five points and exactly four DiT forwards.
+    parser.add_argument("--steps",
+                        type=int,
+                        default=5,
+                        help="sigma-grid points; N points run N-1 DiT forwards (the trained default is 5)")
+    parser.add_argument("--seed", type=int, default=1000, help="seed reused for every measured request")
+    parser.add_argument("--warmup-seed", type=int, default=999)
+    parser.add_argument("--repeats", type=int, default=3, help="number of measured requests after warmup")
+    parser.add_argument("--warmup",
+                        action=argparse.BooleanOptionalAction,
+                        default=True,
+                        help="run one excluded request before timing")
+    parser.add_argument("--num-gpus", type=int, default=4)
+    parser.add_argument("--vsa-sparsity",
+                        type=float,
+                        default=0.9,
+                        help="run-level VSA sparsity in [0, 1); 0.9 is the checkpoint's trained policy")
+    parser.add_argument("--vsa-tile-size",
+                        type=int,
+                        choices=(64, 256),
+                        default=64,
+                        help="VSA-H3 tile size; 64 is the checkpoint's trained and measured geometry")
+    parser.add_argument("--vsa-kernel",
+                        choices=("triton", "sm100a"),
+                        default="sm100a",
+                        help="tile-64 sparse kernel; sm100a is the measured GB200 route and requires a compatible "
+                        "fastvideo-kernel build")
+    parser.add_argument("--fa4",
+                        action=argparse.BooleanOptionalAction,
+                        default=True,
+                        help="use FA4 for eligible non-VSA attention paths")
+    parser.add_argument("--h3-fusions",
+                        action=argparse.BooleanOptionalAction,
+                        default=None,
+                        help="override the profile's H3 fusion policy (changes model numerics when enabled)")
+    parser.add_argument("--compile-vae",
+                        action=argparse.BooleanOptionalAction,
+                        default=True,
+                        help="compile the video VAE decoder independently of the DiT")
+    parser.add_argument("--parallel-vae",
+                        action=argparse.BooleanOptionalAction,
+                        default=True,
+                        help="round-robin VAE temporal chunks across sequence-parallel ranks")
+    parser.add_argument("--replicated-dit",
+                        action=argparse.BooleanOptionalAction,
+                        default=True,
+                        help="replicate DiT weights instead of FSDP-sharding them")
+    parser.add_argument("--pin-cpu-memory",
+                        action=argparse.BooleanOptionalAction,
+                        default=True,
+                        help="pin CPU-offloaded text-encoder and VAE weights")
+    parser.add_argument("--torch-compile",
+                        action=argparse.BooleanOptionalAction,
+                        default=False,
+                        help="compile the whole DiT path (off in the fastest FastH3 profile)")
+    parser.add_argument("--inference-torch-compile",
+                        action=argparse.BooleanOptionalAction,
+                        default=True,
+                        help="regionally compile DiT blocks (enabled in the fastest FastH3 profile)")
+    parser.add_argument("--ulysses-a2a",
+                        choices=("off", "auto"),
+                        default="off",
+                        help="sequence-parallel all-to-all route; off reproduces the fastest FastH3 profile, while "
+                        "auto opts into the fused NVLink kernel when the installed kernel package supports it")
+    parser.add_argument("--compile-mode",
+                        default=None,
+                        help='whole-DiT torch.compile mode, e.g. "reduce-overhead"; requires '
+                        "--no-inference-torch-compile")
+    return parser
+
+
+def validate_args(parser: argparse.ArgumentParser, args: argparse.Namespace) -> argparse.Namespace:
+    if args.repeats < 1:
+        parser.error("--repeats must be at least 1")
+    if args.num_gpus < 1:
+        parser.error("--num-gpus must be at least 1")
+    if not 0.0 <= args.vsa_sparsity < 1.0:
+        parser.error("--vsa-sparsity must be in [0, 1)")
+    if args.compile_mode is not None and args.inference_torch_compile:
+        parser.error("--compile-mode cannot be combined with regional compile; pass --no-inference-torch-compile")
+    return args
+
+
+def parse_args(argv: Sequence[str] | None = None) -> argparse.Namespace:
+    parser = build_parser()
+    return validate_args(parser, parser.parse_args(argv))
+
+
+def _uses_vsa(args: argparse.Namespace) -> bool:
+    """Full FastH3 checkpoints use VSA; LoRA previews may select dense attention."""
+    return bool(getattr(args, "vsa", True))
+
+
+def _h3_fusions_enabled(args: argparse.Namespace) -> bool:
+    if args.h3_fusions is not None:
+        return bool(args.h3_fusions)
+    return args.profile == "all"
+
+
+def profile_environment(args: argparse.Namespace) -> dict[str, str | None]:
+    """Return the complete boot-time environment for this profile.
+
+    ``None`` means the variable must be removed. Values are explicit even for
+    disabled features so a shell's inherited experiment settings cannot
+    silently change the advertised profile.
+    """
+    use_vsa = _uses_vsa(args)
+    return {
+        "FASTVIDEO_ATTENTION_BACKEND": "VIDEO_SPARSE_ATTN_H3" if use_vsa else "FLASH_ATTN",
+        "FASTVIDEO_VSA_SM100A": "1" if use_vsa and args.vsa_kernel == "sm100a" else "0",
+        "FASTVIDEO_VSA_CUTEDSL": "0",
+        # A non-empty output path enables the diagnostic probe.
+        "FASTVIDEO_H3_VSA_PROBE": None,
+        "FASTVIDEO_DISABLE_ATTENTION_COMPILE": "0",
+        "FASTVIDEO_FA4": "1" if args.fa4 else "0",
+        "FASTVIDEO_NVFP4_FA4": "0",
+        "FASTVIDEO_MINIMAX_H3_FA4_PACKED_VARLEN": "0",
+        "FASTVIDEO_MINIMAX_H3_FUSIONS": "all" if _h3_fusions_enabled(args) else "0",
+        "FASTVIDEO_INFERENCE_TORCH_COMPILE": "1" if args.inference_torch_compile else "0",
+        "FASTVIDEO_VAE_PARALLEL_DECODE": "1" if args.parallel_vae else "0",
+        "FASTVIDEO_VAE_PARALLEL_ENCODE": "0",
+        "FASTVIDEO_VAE_PARALLEL_DECODE_STRATEGY": "gather",
+        "FASTVIDEO_ULYSSES_A2A": args.ulysses_a2a,
+        "FASTVIDEO_STAGE_LOGGING": "1",
+    }
+
+
+def configure_environment(args: argparse.Namespace) -> dict[str, str | None]:
+    environment = profile_environment(args)
+    for name, value in environment.items():
+        if value is None:
+            os.environ.pop(name, None)
+        else:
+            os.environ[name] = value
+    return environment
+
+
+def _fa4_is_installed() -> bool:
+    try:
+        return importlib.util.find_spec("flash_attn.cute") is not None
+    except (ImportError, ModuleNotFoundError):
+        return False
+
+
+def _sm100a_kernel_is_installed() -> bool:
+    try:
+        from fastvideo_kernel import block_sparse_attn_sm100a
+    except ImportError:
+        return False
+    return bool(getattr(block_sparse_attn_sm100a, "_HAS_VSA_SM100A", False))
+
+
+def validate_profile_dependencies(args: argparse.Namespace) -> None:
+    """Fail before model loading when the selected measured route is absent."""
+    if args.fa4 and not _fa4_is_installed():
+        raise RuntimeError(
+            "FastH3's FA4 profile requires the pinned flash-attn-4 package. Install it with "
+            "`UV_TORCH_BACKEND=cu130 uv pip install -e \".[fasth3]\"`, or pass --no-fa4.")
+    if _uses_vsa(args) and args.vsa_kernel == "sm100a" and not _sm100a_kernel_is_installed():
+        raise RuntimeError(
+            "FastH3's sm100a profile requires fastvideo-kernel 0.3.4 built with the Blackwell VSA extension. "
+            "Install this checkout with `UV_TORCH_BACKEND=cu130 uv pip install -e \".[fasth3]\"` (or run "
+            "`cd fastvideo-kernel && ./build.sh`), or pass --vsa-kernel triton.")
+
+
+def build_generator_config(args: argparse.Namespace) -> GeneratorConfig:
+    use_vsa = _uses_vsa(args)
+    experimental: dict[str, object] = {
+        "attention_backend": "VIDEO_SPARSE_ATTN_H3" if use_vsa else "FLASH_ATTN",
+        "inference_torch_compile": args.inference_torch_compile,
+        "vae_parallel_decode": args.parallel_vae,
+        "vae_parallel_decode_strategy": "gather",
+    }
+    if use_vsa:
+        experimental.update({
+            "VSA_sparsity": args.vsa_sparsity,
+            "VSA_tile_size": args.vsa_tile_size,
+        })
+    return GeneratorConfig(
+        model_path=args.model_path,
+        pipeline=PipelineSelection(
+            components=ComponentConfig(
+                lora_path=getattr(args, "lora_path", None),
+                lora_strength=float(getattr(args, "lora_strength", 1.0)),
+            ),
+            experimental=experimental,
+        ),
+        engine=EngineConfig(
+            num_gpus=args.num_gpus,
+            use_fsdp_inference=args.num_gpus > 1 and not args.replicated_dit,
+            parallelism=ParallelismConfig(tp_size=1, sp_size=args.num_gpus),
+            offload=OffloadConfig(
+                dit=False,
+                dit_layerwise=False,
+                text_encoder=True,
+                vae=True,
+                pin_cpu_memory=args.pin_cpu_memory,
+            ),
+            compile=CompileConfig(
+                enabled=args.torch_compile,
+                mode=args.compile_mode,
+                vae_enabled=args.compile_vae,
+            ),
+        ),
+    )
+
+
+def build_request(args: argparse.Namespace, output_path: Path, seed: int) -> GenerationRequest:
+    return GenerationRequest(
+        prompt=args.prompt,
+        negative_prompt="",
+        sampling=SamplingConfig(
+            height=args.height,
+            width=args.width,
+            num_frames=args.num_frames,
+            fps=24,
+            num_inference_steps=args.steps,
+            # MiniMax-H3 is guidance-distilled; FastH3 inherits that contract.
+            guidance_scale=1.0,
+            batch_cfg=False,
+            seed=seed,
+        ),
+        output=OutputConfig(
+            output_path=str(output_path),
+            save_video=True,
+            return_frames=False,
+        ),
+    )
+
+
+def _actual_output_path(result: object, requested: Path) -> Path:
+    video_path = getattr(result, "video_path", None)
+    return Path(video_path) if video_path else requested
+
+
+def _denoise_seconds(result: object) -> float | None:
+    stages = getattr(getattr(result, "logging_info", None), "stages", None)
+    if not stages:
+        return None
+    for stage_name, metrics in stages.items():
+        if "denois" not in stage_name.lower():
+            continue
+        execution_time = metrics.get("execution_time")
+        return float(execution_time) if execution_time is not None else None
+    return None
+
+
+def run(args: argparse.Namespace) -> list[float]:
+    output_dir = Path(args.output)
+    output_dir.mkdir(parents=True, exist_ok=True)
+    environment = configure_environment(args)
+    validate_profile_dependencies(args)
+
+    print(f"Profile: {args.profile} ({'non-parity fusions' if _h3_fusions_enabled(args) else 'fusions off'})")
+    print(f"Output directory: {output_dir.resolve()}")
+    print("Denoising contract: 5 sigma points = 4 DiT forwards" if args.steps == 5 else
+          f"Denoising contract override: {args.steps} sigma points = {args.steps - 1} DiT forwards")
+    print("Profile environment: " + " ".join(f"{key}={value if value is not None else '<unset>'}"
+                                                  for key, value in environment.items()))
+
+    generator = VideoGenerator.from_config(build_generator_config(args))
+    measured_wall_times: list[float] = []
+    measured_denoise_times: list[float] = []
+    try:
+        if args.warmup:
+            warmup_path = output_dir / "_fasth3_warmup.mp4"
+            print(f"[warmup] generating (excluded from timing summary): {warmup_path}")
+            started = time.perf_counter()
+            warmup_result = generator.generate(build_request(args, warmup_path, args.warmup_seed))
+            warmup_wall = time.perf_counter() - started
+            actual_warmup_path = _actual_output_path(warmup_result, warmup_path)
+            print(f"[warmup] wall={warmup_wall:.3f}s (excluded)")
+            print(f"Warmup output written to: {actual_warmup_path}")
+
+        for index in range(1, args.repeats + 1):
+            requested_path = output_dir / f"fasth3_{args.profile}_run_{index:02d}.mp4"
+            print(f"[measured {index}/{args.repeats}] generating: {requested_path}")
+            started = time.perf_counter()
+            result = generator.generate(build_request(args, requested_path, args.seed))
+            wall = time.perf_counter() - started
+            measured_wall_times.append(wall)
+            actual_path = _actual_output_path(result, requested_path)
+            print(f"Output written to: {actual_path}")
+            print(f"E2E wall time: {wall:.3f}s")
+            generation_time = getattr(result, "generation_time", None)
+            if generation_time is not None:
+                print(f"Generation time: {float(generation_time):.3f}s")
+            denoise_time = _denoise_seconds(result)
+            if denoise_time is not None:
+                measured_denoise_times.append(denoise_time)
+                print(f"Denoising time: {denoise_time:.3f}s")
+
+        median = statistics.median(measured_wall_times)
+        print(f"Measured E2E wall times (n={len(measured_wall_times)}, warmup excluded): "
+              f"{[round(value, 3) for value in measured_wall_times]}")
+        print(f"Median E2E wall time: {median:.3f}s")
+        if measured_denoise_times:
+            print(f"Median denoising time: {statistics.median(measured_denoise_times):.3f}s")
+        return measured_wall_times
+    finally:
+        generator.shutdown()
+
+
+def main() -> None:
+    run(parse_args())
+
+
+if __name__ == "__main__":
+    main()

examples/inference/basic/basic_fasth3_lora_preview.py

@@ -0,0 +1,81 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Run a FastH3 four-step Preview LoRA with the measured FastVideo defaults.
+
+This is the LoRA counterpart of ``basic_fasth3.py``. Both routes share the
+same performance profile: four DiT forwards, regional fullgraph DiT compile,
+H3 fusions, compiled and sequence-parallel video VAE decode, replicated DiT,
+pinned CPU offload, FA4, and the sm100a tile-64 kernel for VSA adapters.
+
+The FastH3 adapters include low-rank factors plus exact dense deltas. Some also
+provide the VSA compression gate that is absent from the base checkpoint. Pass
+the adapter at construction so all three payload types receive the same
+``--lora-strength``. The attention backend is inferred from that payload unless
+``--vsa`` or ``--no-vsa`` is specified explicitly.
+"""
+
+from __future__ import annotations
+
+import argparse
+import math
+from collections.abc import Sequence
+
+try:
+    from . import basic_fasth3
+except ImportError:
+    # Direct script execution puts this directory, rather than ``examples``, on
+    # sys.path. Keep both ``python file.py`` and module/importlib use working.
+    import basic_fasth3  # type: ignore[no-redef]
+
+BASE_MODEL = "MiniMaxAI/MiniMax-H3"
+
+
+def parse_args(argv: Sequence[str] | None = None) -> argparse.Namespace:
+    parser = basic_fasth3.build_parser(description=__doc__)
+    parser.set_defaults(model_path=BASE_MODEL, output="outputs/fasth3_lora_preview")
+    parser.add_argument(
+        "--lora-path",
+        required=True,
+        help="FastH3 adapter safetensors file or local adapter directory",
+    )
+    parser.add_argument(
+        "--lora-strength",
+        type=float,
+        default=1.0,
+        help="adapter strength; 0 zeros its weights but keeps its backend, and 1 applies its published scale",
+    )
+    parser.add_argument(
+        "--vsa",
+        action=argparse.BooleanOptionalAction,
+        default=None,
+        help="select VSA explicitly; by default it is inferred from the adapter's compression-gate payload",
+    )
+    args = basic_fasth3.validate_args(parser, parser.parse_args(argv))
+    if not math.isfinite(args.lora_strength):
+        parser.error("--lora-strength must be finite")
+    return _resolve_attention_backend(parser, args)
+
+
+def _resolve_attention_backend(parser: argparse.ArgumentParser, args: argparse.Namespace) -> argparse.Namespace:
+    # Header-only inspection keeps the payload on disk. A replacement compression
+    # gate is an unambiguous VSA requirement; adapters without one default to dense.
+    from fastvideo.models.loader.lora_patch import DenseLoRAPatch
+
+    patch = DenseLoRAPatch.from_adapter(args.lora_path, strength=args.lora_strength)
+    needs_vsa = bool(patch and any("gate_compress" in name for name in patch.replacement_parameters))
+    if args.vsa is None:
+        args.vsa = needs_vsa
+    elif needs_vsa and not args.vsa:
+        parser.error(f"{args.lora_path} provides to_gate_compress and must run with VSA; drop --no-vsa")
+    return args
+
+
+def main() -> None:
+    args = parse_args()
+    print(f"FastH3 adapter: {args.lora_path}")
+    print(f"LoRA strength: {args.lora_strength:g}")
+    print(f"Attention: {'VSA-H3' if args.vsa else 'dense FA4'}")
+    basic_fasth3.run(args)
+
+
+if __name__ == "__main__":
+    main()

examples/inference/basic/basic_minimax_h3_fl2va.py

@@ -27,7 +27,8 @@ def parse_args() -> argparse.Namespace:
     parser.add_argument("--model-path", default="MiniMaxAI/MiniMax-H3")
     # Rank-reduced AdaLN checkpoint (-39% params, -23 GiB VRAM): pass
     #   --model-path noctuashap/MiniMax-H3-pruned-r16
-    # (or a local dir converted with tools/minimax_h3/fit_adaln_basis.py).
+    # (or a local dir produced by
+    #   scripts/checkpoint_conversion/convert_minimax_h3_adaln_rank.py).
     # adaln_rank is read from the checkpoint config; no other flags needed.
     # Rank-reduced checkpoints are inference-only: training needs the
     # full-rank release.

examples/inference/basic/basic_minimax_h3_ref2va.py

@@ -27,7 +27,8 @@ def parse_args() -> argparse.Namespace:
     parser.add_argument("--model-path", default="MiniMaxAI/MiniMax-H3")
     # Rank-reduced AdaLN checkpoint (-39% params, -23 GiB VRAM): pass
     #   --model-path noctuashap/MiniMax-H3-pruned-r16
-    # (or a local dir converted with tools/minimax_h3/fit_adaln_basis.py).
+    # (or a local dir produced by
+    #   scripts/checkpoint_conversion/convert_minimax_h3_adaln_rank.py).
     # adaln_rank is read from the checkpoint config; no other flags needed.
     # Rank-reduced checkpoints are inference-only: training needs the
     # full-rank release.

examples/inference/basic/basic_minimax_h3_t2v.py

@@ -15,6 +15,7 @@ from fastvideo.api import (
     OffloadConfig,
     OutputConfig,
     ParallelismConfig,
+    PipelineSelection,
     SamplingConfig,
 )
 
@@ -24,7 +25,8 @@ def parse_args() -> argparse.Namespace:
     parser.add_argument("--model-path", default="MiniMaxAI/MiniMax-H3")
     # Rank-reduced AdaLN checkpoint (-39% params, -23 GiB VRAM): pass
     #   --model-path noctuashap/MiniMax-H3-pruned-r16
-    # (or a local dir converted with tools/minimax_h3/fit_adaln_basis.py).
+    # (or a local dir produced by
+    #   scripts/checkpoint_conversion/convert_minimax_h3_adaln_rank.py).
     # adaln_rank is read from the checkpoint config; no other flags needed.
     # Rank-reduced checkpoints are inference-only: training needs the
     # full-rank release.
@@ -40,6 +42,12 @@ def parse_args() -> argparse.Namespace:
     parser.add_argument("--compile-mode",
                         default=None,
                         help='torch.compile mode, e.g. "reduce-overhead" for CUDA graphs')
+    parser.add_argument("--inference-torch-compile",
+                        action="store_true",
+                        help="regional fullgraph torch.compile of each DiT block after load (the #1718 "
+                        "training-port semantics: no kwargs; fullgraph + emulate_precision_casts injected). "
+                        "First generation pays the inductor JIT (~1-2 min); use --repeats >= 2 and time "
+                        "the last repeat. FASTVIDEO_INFERENCE_TORCH_COMPILE=1 is equivalent")
     parser.add_argument("--repeats",
                         type=int,
                         default=1,
@@ -53,9 +61,16 @@ def main() -> None:
     output_dir = Path(args.output)
     output_dir.mkdir(parents=True, exist_ok=True)
 
+    # Boot-time run configuration folded into FastVideoArgs (the same
+    # experimental-dict route basic_fasth3.py uses for the VSA knobs).
+    experimental: dict[str, object] = {}
+    if args.inference_torch_compile:
+        experimental["inference_torch_compile"] = True
+
     generator = VideoGenerator.from_config(
         GeneratorConfig(
             model_path=args.model_path,
+            pipeline=PipelineSelection(experimental=experimental),
             engine=EngineConfig(
                 num_gpus=args.num_gpus,
                 use_fsdp_inference=args.num_gpus > 1,

examples/inference/basic/mlx_fasth3.py

@@ -0,0 +1,108 @@
+# SPDX-License-Identifier: Apache-2.0
+"""End-to-end MiniMax-H3 (FastH3) generation with the Apple Silicon MLX runtime.
+
+Accepts a text prompt and produces an MP4 with H.264 video at 24 fps and
+stereo AAC audio at 32 kHz. One heavyweight model phase is resident at a time.
+
+    python examples/inference/basic/mlx_fasth3.py \
+      --model-root ~/models/FastH3-Preview-v0.2 \
+      --mlx-checkpoint ~/models/FastH3-MLX/int8 \
+      --prompt '(S1) A red panda says <d>[English] Fast H3 is amazing.</d>' \
+      --height 480 --width 832 --num-frames 124 --seed 2026 \
+      --output-path ~/fasth3_outputs/int8.mp4
+
+Conditioning uses the streamed Qwen3-VL text encoder on first use and caches
+the resulting embeddings under --prompt-cache-dir for instant reuse.
+
+``--fast`` is temporal fast mode. It keeps full-duration audio while
+denoising fewer video frames, then uses MLX RIFE 4.25 to reconstruct the
+requested frame count. A 1280x720 request runs on H3's 1280x736 grid and is
+center-cropped after decode.
+
+This entrypoint currently supports text-to-video-with-audio only. It does not
+yet wire FL2VA, Ref2VA, spatial fast mode, or two-pass refinement.
+"""
+
+from __future__ import annotations
+
+import argparse
+import json
+from pathlib import Path
+
+
+def parse_args() -> argparse.Namespace:
+    parser = argparse.ArgumentParser(description=__doc__,
+                                     formatter_class=argparse.RawDescriptionHelpFormatter)
+    parser.add_argument("--model-root", type=Path, default=Path.home() / "models/FastH3-Preview-v0.2",
+                        help="H3 snapshot root (vae/, audio_vae/, text_encoder/, tokenizer/)")
+    parser.add_argument("--mlx-checkpoint", type=Path, required=True,
+                        help="pre-quantized MLX DiT directory (int8/int6/int4 mlx_h3_dit format)")
+    parser.add_argument(
+        "--prompt",
+        required=True,
+        help="H3 text prompt; use (S1) and <d>[Language] words</d> for explicit dialogue",
+    )
+    parser.add_argument("--output-path", type=Path, required=True)
+    parser.add_argument("--height", type=int, default=480)
+    parser.add_argument("--width", type=int, default=832)
+    parser.add_argument("--num-frames", type=int, default=124)
+    parser.add_argument("--seed", type=int, default=0)
+    parser.add_argument("--steps", type=int, default=4, help="denoise steps (trained ladder = 4)")
+    parser.add_argument(
+        "--fast",
+        action=argparse.BooleanOptionalAction,
+        default=False,
+        help="denoise fewer video frames, then use MLX RIFE to restore the target frame count; audio stays full length",
+    )
+    parser.add_argument("--fast-factor", type=int, default=2,
+                        help="temporal reduction target for --fast (default: 2)")
+    parser.add_argument("--fast-sharpen", type=float, default=0.6,
+                        help="unsharp strength after RIFE interpolation (0 disables)")
+    parser.add_argument("--rife-weights-dir", type=Path, default=None,
+                        help="optional local mlx-community/RIFE-4.25 snapshot")
+    parser.add_argument("--vae-dtype", choices=("fp32", "fp16", "bf16"), default="fp32")
+    parser.add_argument("--prompt-cache-dir", type=Path, default=None,
+                        help="directory for reusable prompt embedding caches")
+    parser.add_argument(
+        "--tiled-video-decode",
+        action=argparse.BooleanOptionalAction,
+        default=True,
+        help="decode with the reference 256px overlapping VAE tiles (disable only for diagnostics)",
+    )
+    return parser.parse_args()
+
+
+def main() -> None:
+    args = parse_args()
+    from fastvideo.mlx_runtime.minimax_h3_pipeline import MiniMaxH3MLXPipeline
+
+    pipeline = MiniMaxH3MLXPipeline(
+        model_root=args.model_root,
+        mlx_dit_checkpoint=args.mlx_checkpoint,
+        vae_dtype=args.vae_dtype,
+        prompt_cache_dir=args.prompt_cache_dir,
+    )
+    result = pipeline.generate(
+        args.prompt,
+        output_path=args.output_path,
+        height=args.height,
+        width=args.width,
+        num_frames=args.num_frames,
+        seed=args.seed,
+        num_steps=args.steps,
+        tiled_video_decode=args.tiled_video_decode,
+        fast=args.fast,
+        fast_factor=args.fast_factor,
+        fast_sharpen=args.fast_sharpen,
+        rife_weights_dir=args.rife_weights_dir,
+    )
+    print(json.dumps({
+        "video_path": result.video_path,
+        "timings_s": {k: round(v, 2) for k, v in result.timings.items()},
+        "peak_memory_gib": {k: round(v, 2) for k, v in result.peak_memory_gib.items()},
+        "audio_samples": int(result.waveform.shape[-1]),
+    }, indent=2))
+
+
+if __name__ == "__main__":
+    main()

examples/inference/basic/mlx_wan22_generate.py

@@ -1,14 +1,18 @@
 # SPDX-License-Identifier: Apache-2.0
-"""End-to-end Wan2.2-TI2V-5B generation on Apple Silicon (MLX DiT + MLX TAEHV).
+"""End-to-end FastMetal-5B-QAD generation on Apple Silicon (MLX DiT + MLX TAEHV).
+
+This is the Wan2.2 TI2V entrypoint. Use FastVideo/FastMetal-5B-QAD:
+
+    hf download FastVideo/FastMetal-5B-QAD --local-dir ./FastMetal-5B-QAD
+    python examples/inference/basic/mlx_wan22_generate.py \\
+      --mlx-checkpoint ./FastMetal-5B-QAD \\
+      --text-encoder-root ./FastMetal-5B-QAD \\
+      --vae-root ./FastMetal-5B-QAD/vae
 
 Pipeline: torch/MPS UMT5 encode (shared with 1.3B) → MLXWan22DiT 3-step DMD
 (warped schedule, flow_shift=5) → MLX TAEHV decode (taew2_2.pth). Fully MLX
 on the heavy DiT + decode path.
 
-    PYTHONPATH=$PWD python examples/inference/basic/mlx_wan22_generate.py \
-      --prompt "A red fox trotting through a snowy pine forest at golden hour" \
-      --output-path video_samples/demo_5b/fox_5b_mlx.mp4
-
 Decoder backends: ``taehv`` (default, MLX, ~seconds), ``taehv-torch`` (parity),
 ``wan-vae`` (full AutoencoderKLWan on MPS, slow).
 """
@@ -31,6 +35,11 @@ from fastvideo.mlx_runtime.prompt_cache import (
     save_prompt_cache,
     text_encoder_fingerprint,
 )
+from fastvideo.mlx_runtime.checkpoint_compat import (
+    UnsupportedMLXCheckpointError,
+    raise_if_unsupported_mlx_checkpoint,
+    resolve_mlx_checkpoint,
+)
 from fastvideo.mlx_runtime.rife_interp import aligned_keyframe_count
 
 FASTWAN21_MODEL_ID = "FastVideo/FastWan2.1-T2V-1.3B-Diffusers"
@@ -165,7 +174,9 @@ def main() -> None:
         "--mlx-checkpoint",
         type=Path,
         default=None,
-        help="Pre-quantized MLX DiT checkpoint directory. Rewrapped with Wan2.2 per-token conditioning.",
+        help="Packed FastMetal-5B-QAD MLX DiT directory (mlx_dit.json + mlx_dit.safetensors). "
+        "If omitted, a FastMetal directory passed as --text-encoder-root is used when it "
+        "already contains those files.",
     )
     parser.add_argument("--vae-root", type=Path, default=None)
     parser.add_argument("--height", type=int, default=DEFAULT_HEIGHT)
@@ -241,6 +252,18 @@ def main() -> None:
     # latent never leaves the grid it was denoised on and the mode is usable.
     if args.refine and args.fast_spatial:
         print("[wan22] --refine takes precedence over --fast-spatial")
+
+    args.mlx_checkpoint = resolve_mlx_checkpoint(args.mlx_checkpoint, args.text_encoder_root)
+    if args.mlx_checkpoint is not None:
+        if args.text_encoder_root is None and (args.mlx_checkpoint / "text_encoder").is_dir():
+            args.text_encoder_root = args.mlx_checkpoint
+        if args.vae_root is None and (args.mlx_checkpoint / "vae").is_dir():
+            args.vae_root = args.mlx_checkpoint / "vae"
+    try:
+        raise_if_unsupported_mlx_checkpoint(args.mlx_checkpoint, args.dit_checkpoint)
+    except UnsupportedMLXCheckpointError as exc:
+        raise SystemExit(str(exc)) from exc
+
     args.text_encoder_root, args.dit_checkpoint, args.dit_config, args.vae_root = _resolve_model_paths(
         text_encoder_root=args.text_encoder_root,
         dit_checkpoint=args.dit_checkpoint,

examples/inference/basic/mlx_wan_prompt_to_video.py

@@ -1,11 +1,25 @@
-"""Generate a FastWan text-to-video clip with the Apple Silicon MLX runtime.
+"""Generate a FastMetal text-to-video clip with the Apple Silicon MLX runtime.
 
-This is the supported source-tree entrypoint for the FastWan-QAD-INT8-1.3B
-Apple release:
+This is the supported source-tree entrypoint for FastMetal-QAD (Wan2.1 1.3B
+and 14B). Use ``mlx_wan22_generate.py`` for FastMetal-5B-QAD.
+
+Download FastMetal-QAD and point ``--model-root`` / ``--mlx-checkpoint`` at it:
+
+    hf download FastVideo/FastMetal-1.3B-QAD --local-dir ./FastMetal-1.3B-QAD
+    python examples/inference/basic/mlx_wan_prompt_to_video.py \\
+      --model-root ./FastMetal-1.3B-QAD --mlx-checkpoint ./FastMetal-1.3B-QAD
+
+CUDA FastWan-QAD (``FastVideo/FastWan-QAD-1.3B``, ``FastVideo/FastWan-QAD-FP8-1.3B``)
+is a separate NVIDIA release.
+
+FastMetal-QAD Hugging Face repos ship ``mlx_dit.json`` + ``mlx_dit.safetensors``,
+not a Diffusers ``transformer/`` tree. Do not copy ``transformer/config.json``
+from Wan2.1 or other checkpoints; point ``--mlx-checkpoint`` at the FastMetal
+directory and the example reads the DiT config from ``mlx_dit.json``.
 
 - Hugging Face/torch encodes the prompt with UMT5 (bf16 by default: fp32
   exponent range without fp16 overflow risk, at fp16 memory cost).
-- MLX runs the FastWan DiT denoising loop (INT8 by default, compiled with
+- MLX runs the FastMetal DiT denoising loop (INT8 by default, compiled with
   ``mx.compile`` unless ``--no-mlx-compile``).
 - TAEHV (default, fast/low-memory) or the full Wan VAE (``--decode-backend
   wan-vae``, higher fidelity, bf16) decodes the final latents.
@@ -56,13 +70,18 @@ from fastvideo.mlx_runtime.prompt_cache import (
     save_prompt_cache,
     text_encoder_fingerprint,
 )
+from fastvideo.mlx_runtime.checkpoint_compat import (
+    UnsupportedMLXCheckpointError,
+    raise_if_unsupported_mlx_checkpoint,
+    resolve_mlx_checkpoint,
+)
 from fastvideo.mlx_runtime.rife_interp import aligned_keyframe_count
 
 if TYPE_CHECKING:  # pragma: no cover - typing only
     from fastvideo.mlx_runtime.fast_spatial import FastSpatialPlan
 
 
-DEFAULT_MODEL_ID = "FastVideo/FastWan2.1-T2V-1.3B-Diffusers"
+DEFAULT_MODEL_ID = "FastVideo/FastMetal-1.3B-QAD"
 
 # Legacy pinned-snapshot location, kept for callers that import it (the MLX
 # benchmark harness). New code should prefer resolve_model_root(None), which
@@ -99,6 +118,10 @@ def resolve_model_root(
         "tokenizer/*",
         "text_encoder/*",
         "vae/*",
+        "mlx_dit.json",
+        "mlx_dit.safetensors",
+        "ema/mlx_dit.json",
+        "ema/mlx_dit.safetensors",
         "transformer/*" if include_transformer else "transformer/config.json",
     ]
     return Path(snapshot_download(
@@ -138,6 +161,7 @@ def encode_prompt(
     text_encoder = UMT5EncoderModel.from_pretrained(
         model_root / "text_encoder",
         torch_dtype=dtype,
+        low_cpu_mem_usage=True,
         local_files_only=True,
     ).to(device)
     text_encoder.eval()
@@ -379,6 +403,7 @@ def decode_latents_to_video(
     vae = AutoencoderKLWan.from_pretrained(
         model_root / "vae",
         torch_dtype=dtype,
+        low_cpu_mem_usage=True,
         local_files_only=True,
     ).to(device)
     vae.eval()
@@ -471,11 +496,16 @@ def _rife_interpolate_video(*, video_path: Path, target_frames: int, factor: int
 
 
 def main() -> None:
-    parser = argparse.ArgumentParser(description="Prompt-to-video FastWan generation using MLX for the DiT")
+    parser = argparse.ArgumentParser(
+        description="Prompt-to-video FastMetal-QAD generation using the Apple Silicon MLX runtime")
     parser.add_argument("--model-root", type=Path, default=None,
-                        help=f"Model directory. Defaults to the local HF cache for {DEFAULT_MODEL_ID} "
+                        help="FastMetal-QAD directory (tokenizer, UMT5, VAE, packed MLX DiT). "
+                        f"Defaults to the local HF cache for {DEFAULT_MODEL_ID} "
                         "(downloading it if missing).")
-    parser.add_argument("--prompt", default="A paper boat sails through a shallow stream in a mossy forest.")
+    parser.add_argument(
+        "--prompt",
+        default="A bird's-eye view of a misty forest valley at dawn.",
+    )
     parser.add_argument("--output-path", type=Path, default=Path("video_samples/mlx_fastwan_prompt_to_video.mp4"))
     parser.add_argument("--height", type=int, default=480)
     parser.add_argument("--width", type=int, default=832)
@@ -639,9 +669,8 @@ def main() -> None:
                         "keyed by (model, prompt, length, dtype), so repeat runs skip "
                         "the text encoder entirely. Default: on.")
     parser.add_argument("--mlx-checkpoint", type=Path, default=None,
-                        help="Load the DiT from a pre-quantized MLX checkpoint directory "
-                        "(created with --save-mlx-checkpoint) instead of casting/quantizing "
-                        "the Diffusers weights on every run.")
+                        help="Packed FastMetal MLX DiT directory (mlx_dit.json + mlx_dit.safetensors). "
+                        "Defaults to --model-root when that directory already contains those files.")
     parser.add_argument("--save-mlx-checkpoint", type=Path, default=None,
                         help="After loading the DiT, save it (cast + quantized) as an MLX "
                         "checkpoint directory for fast reloads via --mlx-checkpoint.")
@@ -690,6 +719,12 @@ def main() -> None:
         np.save(args.encode_prompt_only, prompt_embeds.cpu().numpy())
         return
 
+    mlx_checkpoint = resolve_mlx_checkpoint(args.mlx_checkpoint, model_root)
+    try:
+        raise_if_unsupported_mlx_checkpoint(mlx_checkpoint or model_root)
+    except UnsupportedMLXCheckpointError as exc:
+        raise SystemExit(str(exc)) from exc
+
     import mlx.core as mx
     import torch
     from diffusers import UniPCMultistepScheduler
@@ -719,17 +754,21 @@ def main() -> None:
 
     config_path = model_root / "transformer/config.json"
     checkpoint_path = model_root / "transformer/diffusion_pytorch_model.safetensors"
-    config = json.loads(config_path.read_text())
-    # A pre-quantized MLX DiT can be paired with a lightweight asset root for
-    # UMT5/TAEHV.  In that case the model root is *not* the architecture
-    # authority: use the checkpoint's embedded transformer config for the
-    # sampler guard and latent geometry.
-    dit_config = config
-    if args.mlx_checkpoint is not None:
-        mlx_config_path = Path(args.mlx_checkpoint) / "mlx_dit.json"
-        if mlx_config_path.is_file():
-            mlx_checkpoint_config = json.loads(mlx_config_path.read_text())
-            dit_config = mlx_checkpoint_config.get("config", mlx_checkpoint_config)
+    # Packed FastMetal checkpoints are the architecture authority. Do not
+    # require transformer/config.json when mlx_dit.json is already present.
+    if mlx_checkpoint is not None:
+        mlx_checkpoint_config = json.loads((mlx_checkpoint / "mlx_dit.json").read_text())
+        dit_config = mlx_checkpoint_config.get("config", mlx_checkpoint_config)
+        config = dit_config
+    else:
+        if not config_path.is_file():
+            raise SystemExit(
+                f"No packed MLX DiT (mlx_dit.json) and no Diffusers transformer config at {config_path}. "
+                "FastMetal-QAD checkpoints intentionally omit transformer/; download "
+                "FastVideo/FastMetal-1.3B-QAD and pass --model-root / --mlx-checkpoint at that directory."
+            )
+        config = json.loads(config_path.read_text())
+        dit_config = config
     if int(dit_config.get("in_channels", 0)) == 48 and int(dit_config.get("out_channels", 0)) == 48:
         raise SystemExit(
             "Wan2.2-TI2V-5B uses 48-channel, per-token timestep conditioning. "
@@ -835,10 +874,10 @@ def main() -> None:
     load_start = time.perf_counter()
     mx.clear_cache()
     mx.reset_peak_memory()
-    if args.mlx_checkpoint is not None:
+    if mlx_checkpoint is not None:
         from fastvideo.mlx_runtime.checkpoint import load_mlx_dit_checkpoint
 
-        dit = load_mlx_dit_checkpoint(args.mlx_checkpoint, compile=args.mlx_compile)
+        dit = load_mlx_dit_checkpoint(mlx_checkpoint, compile=args.mlx_compile)
         config = dit.config
     else:
         dit = mlx_dit_from_diffusers_safetensors(

examples/inference/basic/run_fasth3_lora_preview_dense_datafree.sh

@@ -0,0 +1,13 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+repo="FastVideo/FastVideo-FastH3-4-step-Preview-v1-LoRA"
+adapter="dense-datafree/adapter_model.safetensors"
+adapter_path="$(hf download "$repo" "$adapter")"
+
+python examples/inference/basic/basic_fasth3_lora_preview.py \
+  --lora-path "$adapter_path" \
+  --lora-strength "${FASTH3_LORA_STRENGTH:-1.0}" \
+  --output "${FASTH3_LORA_OUTPUT:-outputs/fasth3_lora_preview/dense-datafree}" \
+  "$@" \
+  --no-vsa

examples/inference/basic/run_fasth3_lora_preview_vsa_datafree.sh

@@ -0,0 +1,13 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+repo="FastVideo/FastVideo-FastH3-4-step-Preview-v1-LoRA"
+adapter="vsa-datafree/adapter_model.safetensors"
+adapter_path="$(hf download "$repo" "$adapter")"
+
+python examples/inference/basic/basic_fasth3_lora_preview.py \
+  --lora-path "$adapter_path" \
+  --lora-strength "${FASTH3_LORA_STRENGTH:-1.0}" \
+  --output "${FASTH3_LORA_OUTPUT:-outputs/fasth3_lora_preview/vsa-datafree}" \
+  "$@" \
+  --vsa

examples/inference/basic/run_fasth3_lora_preview_vsa_synthetic_step1300.sh

@@ -0,0 +1,13 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+repo="FastVideo/FastVideo-FastH3-4-step-Preview-v1-LoRA"
+adapter="vsa-synthetic-step1300/adapter_model.safetensors"
+adapter_path="$(hf download "$repo" "$adapter")"
+
+python examples/inference/basic/basic_fasth3_lora_preview.py \
+  --lora-path "$adapter_path" \
+  --lora-strength "${FASTH3_LORA_STRENGTH:-1.0}" \
+  --output "${FASTH3_LORA_OUTPUT:-outputs/fasth3_lora_preview/vsa-synthetic-step1300}" \
+  "$@" \
+  --vsa

examples/inference/basic/run_fasth3_lora_preview_vsa_synthetic_step1900.sh

@@ -0,0 +1,13 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+repo="FastVideo/FastVideo-FastH3-4-step-Preview-v1-LoRA"
+adapter="vsa-synthetic-step1900/adapter_model.safetensors"
+adapter_path="$(hf download "$repo" "$adapter")"
+
+python examples/inference/basic/basic_fasth3_lora_preview.py \
+  --lora-path "$adapter_path" \
+  --lora-strength "${FASTH3_LORA_STRENGTH:-1.0}" \
+  --output "${FASTH3_LORA_OUTPUT:-outputs/fasth3_lora_preview/vsa-synthetic-step1900}" \
+  "$@" \
+  --vsa

examples/inference/gradio/lora_review.py

@@ -0,0 +1,184 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Review whether a FastH3 adapter reproduces the checkpoint it was extracted from.
+
+The question this page answers is not "is the video good" but "does base + adapter land
+where the real checkpoint lands". So each row is one checkpoint, and the two players in
+it are the checkpoint itself and base MiniMax-H3 with that checkpoint's adapter merged
+in. They share a seed and a prompt, so anything you can see between them is the
+adapter's approximation error and nothing else.
+
+The base model at four steps sits at the top as the floor. It is not distilled, so it
+should look clearly worse than everything below it -- if an adapter row looks like the
+floor instead of like its checkpoint, the adapter did not land.
+
+    python examples/inference/gradio/lora_review.py --runs /path/to/lora_review
+
+Expects one directory per arm, each holding ``<index>_<case_id>.mp4``:
+
+    <runs>/base/            <runs>/v1-true/     <runs>/v1-lora-r64/    ...
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import subprocess
+from functools import lru_cache
+from pathlib import Path
+
+import gradio as gr
+
+FLOOR_ARM = "base"
+
+
+def discover_pairs(arms: list[str]) -> tuple[list[tuple[str, str, str]], list[str]]:
+    """Split the arms present into checkpoint/adapter pairs and everything else.
+
+    Pairs are found by name -- ``<x>-true`` next to ``<x>-lora-<rank>`` -- rather than
+    listed, so adding an arm to the render directory is enough to get it on the page.
+    Arms that pair with nothing (a third-party adapter with no checkpoint to compare
+    against) still get shown, on their own row, instead of being silently dropped.
+    """
+    pairs, used = [], set()
+    for arm in sorted(arms):
+        if not arm.endswith("-true"):
+            continue
+        stem = arm[:-len("-true")]
+        partner = next((a for a in arms if a.startswith(f"{stem}-lora")), None)
+        if partner is None:
+            continue
+        pairs.append((stem, arm, partner))
+        used.update({arm, partner})
+    standalone = [a for a in sorted(arms) if a not in used and a != FLOOR_ARM]
+    return pairs, standalone
+
+
+def probe(path: Path) -> str:
+    """`WxH · Nf · Ds · MiB`, so a truncated or mis-sized render is visible as text."""
+    try:
+        out = subprocess.run(
+            ["ffprobe", "-v", "error", "-select_streams", "v:0", "-show_entries",
+             "stream=width,height,nb_read_packets,duration", "-count_packets", "-of", "json",
+             str(path)],
+            capture_output=True, text=True, check=True).stdout
+        stream = json.loads(out)["streams"][0]
+        frames = stream.get("nb_read_packets", "?")
+        duration = float(stream.get("duration", 0) or 0)
+        return (f"{stream['width']}x{stream['height']} · {frames}f · {duration:.1f}s · "
+                f"{path.stat().st_size / 2**20:.1f} MiB")
+    except (subprocess.CalledProcessError, KeyError, IndexError, json.JSONDecodeError):
+        return f"{path.stat().st_size / 2**20:.1f} MiB"
+
+
+class Runs:
+    """Which prompts rendered, and where each arm's clip for them lives."""
+
+    def __init__(self, runs_dir: Path, prompts_file: Path | None) -> None:
+        self.root = runs_dir
+        self.arms = sorted(d.name for d in runs_dir.iterdir() if d.is_dir())
+        self.prompts: dict[str, str] = {}
+        if prompts_file and prompts_file.exists():
+            with prompts_file.open() as handle:
+                for index, line in enumerate(handle):
+                    line = line.strip()
+                    if line:
+                        self.prompts[f"{index:03d}"] = json.loads(line).get("prompt", "")
+
+        self.clips: dict[str, dict[str, Path]] = {}
+        for arm in self.arms:
+            for mp4 in sorted((runs_dir / arm).glob("*.mp4")):
+                self.clips.setdefault(mp4.stem.split("_")[0], {})[arm] = mp4
+        if not self.clips:
+            raise SystemExit(f"no clips under {runs_dir}")
+
+    def label(self, index: str) -> str:
+        head = " ".join(self.prompts.get(index, "").split())[:90]
+        return f"[{index}] {head}..." if head else f"[{index}]"
+
+    def by_label(self, label: str) -> str:
+        return next(i for i in self.clips if self.label(i) == label)
+
+
+def build(runs: Runs, height: int) -> gr.Blocks:
+
+    @lru_cache(maxsize=512)
+    def cached_probe(path: str) -> str:
+        return probe(Path(path))
+
+    def player_update(index: str, arm: str, prefix: str):
+        mp4 = runs.clips.get(index, {}).get(arm)
+        if mp4 is None:
+            return gr.update(value=None, label=f"{prefix} — not rendered")
+        return gr.update(value=str(mp4), label=f"{prefix} — {cached_probe(str(mp4))}")
+
+    with gr.Blocks(title="FastH3 adapter review") as demo:
+        gr.Markdown(
+            "# FastH3 adapter review\n"
+            "Each row is one checkpoint: **left** is the real checkpoint, **right** is base "
+            "MiniMax-H3 with that checkpoint's rank-64 adapter merged in. Same prompt, same "
+            "seed, same sampler. Differences between the two are the adapter's approximation "
+            "error.\n\n"
+            "The top player is undistilled base MiniMax-H3 at four steps — the floor. Every "
+            "row below it should look clearly better than that; an adapter that landed looks "
+            "like its own left-hand player, not like the floor.")
+
+        prompt_dd = gr.Dropdown(choices=[runs.label(i) for i in runs.clips],
+                                value=runs.label(next(iter(runs.clips))),
+                                label="Prompt")
+        prompt_box = gr.Textbox(label="Prompt", lines=4, max_lines=6, interactive=False, show_copy_button=True)
+
+        with gr.Row():
+            floor = gr.Video(label="base MiniMax-H3, 4 steps (floor)", height=height, loop=True,
+                             autoplay=False, interactive=False)
+
+        pairs, standalone = discover_pairs(runs.arms)
+        players: list[tuple[gr.Video, str, str]] = []
+        for row_label, true_arm, lora_arm in pairs:
+            gr.Markdown(f"### {row_label}")
+            with gr.Row():
+                left = gr.Video(label=f"{row_label} — checkpoint", height=height, loop=True,
+                                autoplay=False, interactive=False)
+                right = gr.Video(label=f"{row_label} — base + adapter", height=height, loop=True,
+                                 autoplay=False, interactive=False)
+            players.append((left, true_arm, "checkpoint"))
+            players.append((right, lora_arm, "base + adapter"))
+
+        if standalone:
+            gr.Markdown("### Other adapters (no matching checkpoint to compare against)")
+            with gr.Row():
+                for arm in standalone:
+                    players.append((gr.Video(label=arm, height=height, loop=True, autoplay=False,
+                                             interactive=False), arm, arm))
+
+        def show(label: str):
+            index = runs.by_label(label)
+            return [
+                player_update(index, FLOOR_ARM, "base, 4 steps"),
+                gr.update(value=runs.prompts.get(index, "")),
+                *[player_update(index, arm, prefix) for _, arm, prefix in players],
+            ]
+
+        outputs = [floor, prompt_box, *[p for p, _, _ in players]]
+        gr.on(triggers=[prompt_dd.change, demo.load], fn=show, inputs=prompt_dd, outputs=outputs)
+
+    return demo
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    parser.add_argument("--runs", default="/mnt/lustre/vlm-s4duan/arena_arms/lora_review")
+    parser.add_argument("--prompts-file", default="/mnt/lustre/vlm-s4duan/FastVideo/prompts.jsonl")
+    parser.add_argument("--host", default="0.0.0.0")
+    parser.add_argument("--port", type=int, default=7865)
+    parser.add_argument("--share", action="store_true")
+    parser.add_argument("--video-height", type=int, default=420)
+    args = parser.parse_args()
+
+    runs = Runs(Path(args.runs).resolve(), Path(args.prompts_file))
+    print(f"arms: {runs.arms}")
+    print(f"prompts with output: {sorted(runs.clips)}")
+    build(runs, args.video_height).launch(server_name=args.host, server_port=args.port,
+                                          share=args.share, allowed_paths=[str(runs.root)])
+
+
+if __name__ == "__main__":
+    main()

examples/inference/lora/minimax_h3_lora_inference.py

@@ -0,0 +1,234 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Run a FastH3 distillation adapter on top of the base MiniMax-H3 checkpoint.
+
+The FastH3 checkpoints are full fine-tunes of MiniMax-H3 distilled to four steps under
+video sparse attention. Published as adapters they are three things at once, and all
+three have to land for the result to match the checkpoint:
+
+* low-rank factors for the attention, feed-forward, and AdaLN projections
+* exact ``.diff`` deltas for the norms and biases an SVD cannot usefully factor
+* ``.set_weight`` values for ``attn.to_gate_compress``, the VSA compression gate that
+  does not exist in the base model at all
+
+An adapter carrying that last one needs ``--vsa``: under any other attention backend the
+gate module is never constructed, so part of the distillation has nowhere to go. The
+requirement is read off the adapter rather than assumed, because community adapters
+built against the ComfyUI layout carry no gate and run fine either way -- run one of
+those with ``--no-vsa`` (see ``scripts/checkpoint_conversion/convert_minimax_h3_comfy_lora.py``
+for getting them into a layout this loads).
+
+Because a parameter the base lacks has to be supplied while weights are still unsharded,
+the adapter is passed at construction rather than swapped in afterwards.
+
+    python examples/inference/lora/minimax_h3_lora_inference.py \\
+        --lora-path /models/fasth3-loras-publish/FastH3-4-step-v1.1/rank-64 \\
+        --prompts-file prompts.jsonl --output outputs/v1.1-rank64
+
+Pass no ``--lora-path`` to render the unmodified base model as a control.
+"""
+
+from __future__ import annotations
+
+import argparse
+import json
+import math
+import os
+from collections.abc import Sequence
+from pathlib import Path
+
+# MiniMax-H3 generates 5-15 s at 24 fps, on a latent grid that only admits frame counts
+# of the form 17n + 5. 124 is the 5-second point the FastH3 profile is measured at.
+FRAMES_PER_CHUNK = 17
+LATENTS_PER_CHUNK = 5
+FPS = 24
+MIN_DURATION, MAX_DURATION = 5.0, 15.0
+
+
+def align_num_frames(num_frames: int) -> int:
+    """Round up to the next 17n + 5 the latent grid accepts."""
+    if num_frames <= LATENTS_PER_CHUNK:
+        return LATENTS_PER_CHUNK
+    chunks = -(-(num_frames - LATENTS_PER_CHUNK) // FRAMES_PER_CHUNK)
+    return LATENTS_PER_CHUNK + chunks * FRAMES_PER_CHUNK
+
+
+def parse_args(argv: Sequence[str] | None = None) -> argparse.Namespace:
+    parser = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    parser.add_argument("--model-path", default="MiniMaxAI/MiniMax-H3", help="the BASE checkpoint the adapter targets")
+    parser.add_argument("--lora-path", default=None, help="adapter file or directory; omit to render the base model")
+    parser.add_argument("--lora-nickname", default="fasth3")
+    parser.add_argument("--lora-strength", type=float, default=1.0)
+    parser.add_argument("--prompt", default=None)
+    parser.add_argument("--prompts-file", default=None, help="JSONL with a 'prompt' field per line")
+    parser.add_argument("--limit", type=int, default=None, help="use only the first N prompts")
+    parser.add_argument("--num-shards", type=int, default=1, help="split the prompt list across processes")
+    parser.add_argument("--shard", type=int, default=0)
+    parser.add_argument("--output", default="outputs/minimax_h3_lora")
+    parser.add_argument("--skip-existing", action="store_true", help="leave already-rendered clips alone")
+    parser.add_argument("--height", type=int, default=768)
+    parser.add_argument("--width", type=int, default=1344)
+    parser.add_argument("--num-frames", type=int, default=124)
+    # Counts sigma-GRID POINTS: N points run N-1 DiT forwards. The distilled ladder is
+    # t=1000,750,500,250 -> 0, which is five points and exactly four forwards.
+    parser.add_argument("--steps", type=int, default=5)
+    parser.add_argument("--seed", type=int, default=1000)
+    parser.add_argument("--num-gpus", type=int, default=4)
+    parser.add_argument("--vsa", action=argparse.BooleanOptionalAction, default=None,
+                        help="video sparse attention; inferred from the adapter when omitted")
+    parser.add_argument("--vsa-sparsity", type=float, default=0.9)
+    parser.add_argument("--vsa-tile-size", type=int, choices=(64, 256), default=64)
+    parser.add_argument("--vsa-kernel", choices=("triton", "sm100a"), default="sm100a")
+    parser.add_argument("--fa4", action=argparse.BooleanOptionalAction, default=True)
+    args = parser.parse_args(argv)
+    if not args.prompt and not args.prompts_file:
+        parser.error("pass --prompt or --prompts-file")
+    # Whether VSA is required is a property of the adapter, not of having one at all --
+    # community adapters built against the ComfyUI layout carry no gate. Checked in
+    # main(), once the path has been resolved.
+    aligned = align_num_frames(args.num_frames)
+    if not MIN_DURATION <= aligned / FPS <= MAX_DURATION:
+        parser.error(f"MiniMax-H3 generates {MIN_DURATION}-{MAX_DURATION}s at {FPS} fps; "
+                     f"aligned num_frames={aligned} is {aligned / FPS:.1f}s")
+    args.num_frames = aligned
+    if not math.isfinite(args.lora_strength):
+        parser.error("--lora-strength must be finite")
+    return args
+
+
+def configure_environment(args: argparse.Namespace) -> None:
+    """Set the boot-time backend selection explicitly, including what is off.
+
+    An inherited FASTVIDEO_* from an earlier experiment would otherwise silently change
+    which attention path the run actually took, which is the one thing this comparison
+    cannot afford to be vague about.
+    """
+    env: dict[str, str | None] = {
+        "FASTVIDEO_ATTENTION_BACKEND": "VIDEO_SPARSE_ATTN_H3" if args.vsa else "FLASH_ATTN",
+        "FASTVIDEO_VSA_SM100A": "1" if (args.vsa and args.vsa_kernel == "sm100a") else "0",
+        "FASTVIDEO_VSA_CUTEDSL": "0",
+        "FASTVIDEO_H3_VSA_PROBE": None,
+        "FASTVIDEO_DISABLE_ATTENTION_COMPILE": "0",
+        "FASTVIDEO_FA4": "1" if args.fa4 else "0",
+        "FASTVIDEO_NVFP4_FA4": "0",
+        "FASTVIDEO_MINIMAX_H3_FA4_PACKED_VARLEN": "0",
+        "FASTVIDEO_MINIMAX_H3_FUSIONS": "all",
+        "FASTVIDEO_INFERENCE_TORCH_COMPILE": "1",
+        "FASTVIDEO_VAE_PARALLEL_DECODE": "1",
+        "FASTVIDEO_VAE_PARALLEL_ENCODE": "0",
+        "FASTVIDEO_VAE_PARALLEL_DECODE_STRATEGY": "gather",
+        "FASTVIDEO_ULYSSES_A2A": "off",
+        "FASTVIDEO_STAGE_LOGGING": "1",
+    }
+    for name, value in env.items():
+        if value is None:
+            os.environ.pop(name, None)
+        else:
+            os.environ[name] = value
+
+
+def load_prompts(args: argparse.Namespace) -> list[dict]:
+    if args.prompt:
+        records = [{"id": "000", "prompt": args.prompt}]
+    else:
+        records = []
+        with open(args.prompts_file) as handle:
+            for index, line in enumerate(handle):
+                line = line.strip()
+                if not line:
+                    continue
+                item = json.loads(line)
+                records.append({
+                    "id": str(item.get("id", item.get("sample_id", f"{index:03d}"))),
+                    "prompt": item["prompt"],
+                })
+    if args.limit is not None:
+        records = records[:args.limit]
+    return [r for i, r in enumerate(records) if i % args.num_shards == args.shard]
+
+
+def main(argv: Sequence[str] | None = None) -> None:
+    args = parse_args(argv)
+    # Backend selection is finalized from the adapter before model construction.
+    from fastvideo.models.loader.lora_patch import DenseLoRAPatch
+    from fastvideo import VideoGenerator
+    from fastvideo.api import (CompileConfig, ComponentConfig, EngineConfig, GenerationRequest, GeneratorConfig,
+                               OffloadConfig, OutputConfig, ParallelismConfig, PipelineSelection, SamplingConfig)
+
+    # An adapter carrying to_gate_compress needs the VSA backend, because that is the
+    # only configuration in which the module exists. One that does not carry it runs
+    # fine either way, so the requirement is read off the file rather than assumed from
+    # the presence of an adapter at all.
+    patch = (DenseLoRAPatch.from_adapter(args.lora_path, strength=args.lora_strength)
+             if args.lora_path else None)
+    needs_vsa = bool(patch and any("gate_compress" in name for name in patch.replacement_parameters))
+    if args.vsa is None:
+        args.vsa = needs_vsa
+    if needs_vsa and not args.vsa:
+        raise SystemExit(f"{args.lora_path} carries to_gate_compress, which exists only under the VSA "
+                         "attention backend. Drop --no-vsa.")
+    if args.vsa and args.lora_path and not needs_vsa:
+        print(f"note: {args.lora_path} carries no VSA gate; running under VSA leaves the "
+              "compression branch at its zero-initialized value.")
+    configure_environment(args)
+
+    experimental: dict[str, object] = {
+        "inference_torch_compile": True,
+        "vae_parallel_decode": True,
+        "vae_parallel_decode_strategy": "gather",
+    }
+    if args.vsa:
+        experimental.update({
+            "attention_backend": "VIDEO_SPARSE_ATTN_H3",
+            "VSA_sparsity": args.vsa_sparsity,
+            "VSA_tile_size": args.vsa_tile_size,
+        })
+
+    config = GeneratorConfig(
+        model_path=args.model_path,
+        pipeline=PipelineSelection(
+            components=ComponentConfig(lora_path=args.lora_path, lora_strength=args.lora_strength),
+            experimental=experimental,
+        ),
+        engine=EngineConfig(
+            num_gpus=args.num_gpus,
+            use_fsdp_inference=False,
+            parallelism=ParallelismConfig(tp_size=1, sp_size=args.num_gpus),
+            offload=OffloadConfig(dit=False, dit_layerwise=False, text_encoder=True, vae=True, pin_cpu_memory=True),
+            compile=CompileConfig(enabled=False, vae_enabled=True),
+        ),
+    )
+
+    records = load_prompts(args)
+    out_dir = Path(args.output)
+    out_dir.mkdir(parents=True, exist_ok=True)
+    print(f"adapter: {args.lora_path or '(none, base model)'}")
+    print(f"prompts: {len(records)} (shard {args.shard}/{args.num_shards})")
+
+    generator = VideoGenerator.from_config(config)
+    for index, record in enumerate(records):
+        stem = f"{index:03d}_{record['id']}"
+        if args.skip_existing and (out_dir / f"{stem}.mp4").exists():
+            print(f"[{index}] skip {stem}")
+            continue
+        generator.generate(
+            GenerationRequest(
+                prompt=record["prompt"],
+                negative_prompt="",
+                sampling=SamplingConfig(
+                    height=args.height,
+                    width=args.width,
+                    num_frames=args.num_frames,
+                    fps=FPS,
+                    num_inference_steps=args.steps,
+                    # MiniMax-H3 is guidance-distilled; FastH3 inherits that contract.
+                    guidance_scale=1.0,
+                    batch_cfg=False,
+                    seed=args.seed,
+                ),
+                output=OutputConfig(output_path=str(out_dir / f"{stem}.mp4"), save_video=True, return_frames=False),
+            ))
+        print(f"[{index}] wrote {stem}")
+
+
+if __name__ == "__main__":
+    main()

examples/serving/README.md

@@ -0,0 +1,43 @@
+# OpenAI-compatible serving examples
+
+The REST serving engine is model-agnostic. Any model supported by
+`VideoGenerator` can use the same `/v1/models`, `/v1/videos`, and `/v1/images`
+surface; the two configs here are FastH3 validation profiles.
+
+Launch the full FastH3 checkpoint:
+
+```bash
+fastvideo serve --config examples/serving/openai_fasth3.yaml
+```
+
+Launch the dense FastH3 LoRA on the base MiniMax-H3 checkpoint:
+
+```bash
+adapter_path="$(hf download \
+  FastVideo/FastVideo-FastH3-4-step-Preview-v1-LoRA \
+  dense-datafree/adapter_model.safetensors)"
+
+fastvideo serve --config examples/serving/openai_fasth3_lora.yaml \
+  --generator.pipeline.components.lora_path "$adapter_path"
+```
+
+FastH3 adapters are hybrid startup patches: alongside low-rank factors they
+may contain dense deltas and a VSA compression-gate replacement. They must be
+selected when the server starts. A request may carry the vLLM-Omni `lora`
+selector, but its name, path, and scale must match that startup adapter. A VSA
+adapter also needs `attention_backend: VIDEO_SPARSE_ATTN_H3`, `VSA_sparsity`,
+and `VSA_tile_size` like the full-checkpoint config.
+
+Submit and poll an asynchronous job:
+
+```bash
+job_id="$(curl -sS http://localhost:8000/v1/videos \
+  -H 'content-type: application/json' \
+  -d '{"model":"fasth3","prompt":"A fox runs through fresh snow."}' \
+  | jq -r .id)"
+
+curl -sS "http://localhost:8000/v1/videos/$job_id"
+curl -o result.mp4 "http://localhost:8000/v1/videos/$job_id/content"
+```
+
+For a blocking call, `POST /v1/videos/sync` returns the MP4 body directly.

examples/serving/openai_fasth3.yaml

@@ -0,0 +1,49 @@
+# OpenAI-compatible FastH3 full-checkpoint server.
+generator:
+  model_path: FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2
+  engine:
+    num_gpus: 4
+    use_fsdp_inference: false
+    parallelism:
+      tp_size: 1
+      sp_size: 4
+    offload:
+      dit: false
+      dit_layerwise: false
+      text_encoder: true
+      image_encoder: true
+      vae: true
+      pin_cpu_memory: true
+    compile:
+      enabled: false
+      vae_enabled: false
+  pipeline:
+    workload_type: t2v
+    experimental:
+      attention_backend: VIDEO_SPARSE_ATTN_H3
+      VSA_sparsity: 0.9
+      VSA_tile_size: 64
+      inference_torch_compile: false
+      vae_parallel_decode: true
+      vae_parallel_decode_strategy: gather
+
+server:
+  host: 0.0.0.0
+  port: 8000
+  output_dir: outputs/openai_fasth3
+  served_model_name: fasth3
+
+default_request:
+  negative_prompt: ""
+  sampling:
+    height: 768
+    width: 1344
+    num_frames: 124
+    fps: 24
+    num_inference_steps: 5
+    guidance_scale: 1.0
+    batch_cfg: false
+    seed: 1000
+  output:
+    return_frames: false
+

examples/serving/openai_fasth3_lora.yaml

@@ -0,0 +1,51 @@
+# OpenAI-compatible FastH3 startup-LoRA server. Override lora_path at launch
+# with the exact file returned by `hf download`; see README.md.
+generator:
+  model_path: MiniMaxAI/MiniMax-H3
+  engine:
+    num_gpus: 4
+    use_fsdp_inference: false
+    parallelism:
+      tp_size: 1
+      sp_size: 4
+    offload:
+      dit: false
+      dit_layerwise: false
+      text_encoder: true
+      image_encoder: true
+      vae: true
+      pin_cpu_memory: true
+    compile:
+      enabled: false
+      vae_enabled: false
+  pipeline:
+    workload_type: t2v
+    components:
+      lora_path: /path/to/adapter_model.safetensors
+      lora_nickname: fasth3-dense-datafree
+      lora_strength: 1.0
+    experimental:
+      attention_backend: FLASH_ATTN
+      inference_torch_compile: false
+      vae_parallel_decode: true
+      vae_parallel_decode_strategy: gather
+
+server:
+  host: 0.0.0.0
+  port: 8000
+  output_dir: outputs/openai_fasth3_lora
+  served_model_name: fasth3-dense-datafree
+
+default_request:
+  negative_prompt: ""
+  sampling:
+    height: 768
+    width: 1344
+    num_frames: 124
+    fps: 24
+    num_inference_steps: 5
+    guidance_scale: 1.0
+    batch_cfg: false
+    seed: 1000
+  output:
+    return_frames: false

fastvideo-kernel/CMakeLists.txt

@@ -114,7 +114,63 @@ list(APPEND TORCH_INCLUDE_DIRS ${TORCH_INCLUDE_PATHS})
 # Find Torch package (still useful for libraries)
 find_package(Torch REQUIRED)
 
-# Include directories
+# The Ulysses kernel needs NCCL's 2.29 device API. Keep it optional so ROCm,
+# older NCCL installs, and minimal CUDA builders still produce a usable wheel.
+set(FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A "AUTO" CACHE STRING
+    "Build the NCCL-device-API Ulysses all-to-all: AUTO/ON/OFF")
+set_property(CACHE FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A PROPERTY STRINGS AUTO ON OFF)
+set(ENABLE_ULYSSES_A2A OFF)
+
+if(NOT GPU_BACKEND STREQUAL "ROCM" AND NOT FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A STREQUAL "OFF")
+    # nvidia.nccl is a namespace package, so find_spec rather than __file__.
+    execute_process(COMMAND "${Python_EXECUTABLE}" -c
+        "import importlib.util as u;s=u.find_spec('nvidia.nccl');print(list(s.submodule_search_locations)[0] if s else '')"
+        OUTPUT_VARIABLE NCCL_PIP_ROOT OUTPUT_STRIP_TRAILING_WHITESPACE ERROR_QUIET)
+    find_path(NCCL_INCLUDE_DIR nccl_device/coop.h
+        HINTS $ENV{NCCL_HOME}/include ${NCCL_PIP_ROOT}/include)
+    find_library(NCCL_LIBRARY NAMES nccl
+        HINTS $ENV{NCCL_HOME}/lib $ENV{NCCL_HOME}/lib64 ${NCCL_PIP_ROOT}/lib)
+    # PyPI's nvidia-nccl-cu* wheels ship the SONAME but not the development
+    # symlink (libnccl.so.2, no libnccl.so). Accept that exact versioned name;
+    # target_link_libraries can link an absolute SONAME path directly.
+    if(NOT NCCL_LIBRARY)
+        find_file(NCCL_VERSIONED_LIBRARY NAMES libnccl.so.2
+            HINTS $ENV{NCCL_HOME}/lib $ENV{NCCL_HOME}/lib64 ${NCCL_PIP_ROOT}/lib)
+        if(NCCL_VERSIONED_LIBRARY)
+            set(NCCL_LIBRARY "${NCCL_VERSIONED_LIBRARY}")
+        endif()
+    endif()
+
+    if(NCCL_INCLUDE_DIR AND NCCL_LIBRARY)
+        include(CheckCXXSourceCompiles)
+        set(_FASTVIDEO_REQUIRED_INCLUDES "${CMAKE_REQUIRED_INCLUDES}")
+        # nccl.h includes cuda_runtime.h, so a host-compiler feature probe
+        # needs the toolkit includes explicitly even though CUDA is enabled.
+        set(CMAKE_REQUIRED_INCLUDES "${NCCL_INCLUDE_DIR};${CUDAToolkit_INCLUDE_DIRS}")
+        unset(NCCL_HAS_REQUIRED_DEVICE_API CACHE)
+        check_cxx_source_compiles("\
+#define NCCL_HOSTLIB_ONLY
+#include <cstddef>
+#include <nccl_device.h>
+int main() {
+  ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
+  ncclCommProperties props = NCCL_COMM_PROPERTIES_INITIALIZER;
+  return (reqs.size == 0 || props.size == 0);
+}"
+            NCCL_HAS_REQUIRED_DEVICE_API)
+        set(CMAKE_REQUIRED_INCLUDES "${_FASTVIDEO_REQUIRED_INCLUDES}")
+    endif()
+
+    if(NCCL_HAS_REQUIRED_DEVICE_API)
+        set(ENABLE_ULYSSES_A2A ON)
+    elseif(FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A STREQUAL "ON")
+        message(FATAL_ERROR
+            "FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A=ON requires NCCL device headers/library "
+            "with NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER (NCCL 2.29+). "
+            "Resolved include='${NCCL_INCLUDE_DIR}', library='${NCCL_LIBRARY}'.")
+    endif()
+endif()
+
 include_directories(
     ${CMAKE_SOURCE_DIR}/include
     ${CMAKE_SOURCE_DIR}/include/cutlass/include
@@ -124,6 +180,9 @@ include_directories(
     ${CMAKE_SOURCE_DIR}/csrc/turbodiffusion
     ${TORCH_INCLUDE_DIRS}
 )
+if(ENABLE_ULYSSES_A2A)
+    include_directories(${NCCL_INCLUDE_DIR})
+endif()
 
 # ---------------------------
 # ThunderKittens (TK) toggles
@@ -157,6 +216,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}")
+message(STATUS "FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A: ${FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A}")
 
 set(ENABLE_TK_KERNELS OFF)
 if(FASTVIDEO_KERNEL_BUILD_TK STREQUAL "ON")
@@ -307,6 +367,9 @@ if(BUILD_CXX_KERNELS)
         csrc/turbodiffusion/norm/layernorm.cu
         csrc/turbodiffusion/quant/quant.cu
     )
+    if(ENABLE_ULYSSES_A2A)
+        list(APPEND EXTENSION_SOURCES csrc/comm/ulysses_all_to_all.cu)
+    endif()
     
     # Conditionally add TK kernels
     if(ENABLE_TK_KERNELS)
@@ -318,6 +381,38 @@ if(BUILD_CXX_KERNELS)
 
     # Combined FastVideo Extension
     # Using name 'fastvideo_kernel_ops' to distinguish from the python package namespace
+    # ---------------------------------------------------------------------------
+    # VSA block-sparse attention forward, Blackwell (sm_100a) only.
+    #
+    # NOTE the "a" suffix: -arch=sm_100a is NOT enough -- it emits a plain sm_100 target and
+    # ptxas rejects every tcgen05 / setmaxnreg instruction. The explicit gencode spelling
+    # below is required, and matches the 10.0a entry in TORCH_CUDA_ARCH_LIST.
+    #
+    # Built for 64-token sparse blocks -- FastVideo's default (4,4,4) tiling, so no
+    # tile-size change and no top-k granularity change is needed. VSA_BHSD
+    # selects [B, H, S, D]. Both are compile-time; the Python is_supported() checks incoming
+    # tensors against them so callers fall back to Triton rather than getting a wrong answer.
+    # Read the ENVIRONMENT as well as the cache variable. When TORCH_CUDA_ARCH_LIST is
+    # exported (build.sh, and `pip install` with it set) the branch above only prints it --
+    # the cmake variable stays empty, so testing that alone silently skips the kernel and
+    # leaves a build that succeeds with the op missing.
+    set(ENABLE_VSA_SM100A OFF)
+    set(_VSA_ARCH_LIST "${TORCH_CUDA_ARCH_LIST}")
+    if(NOT _VSA_ARCH_LIST AND DEFINED ENV{TORCH_CUDA_ARCH_LIST})
+        set(_VSA_ARCH_LIST "$ENV{TORCH_CUDA_ARCH_LIST}")
+    endif()
+    if(_VSA_ARCH_LIST MATCHES "(^|[; ,])(10\\.0a|100a|sm_100a)([; ,]|$)")
+        set(ENABLE_VSA_SM100A ON)
+    endif()
+    if(ENABLE_VSA_SM100A)
+        message(STATUS "fastvideo-kernel: building block_sparse_sm100a (Blackwell, 64- and 128-token blocks)")
+        list(APPEND EXTENSION_SOURCES csrc/attention/block_sparse_sm100a.cu
+                                      csrc/attention/block_sparse_blk128_sm100a.cu)
+        set_source_files_properties(csrc/attention/block_sparse_sm100a.cu
+                                    csrc/attention/block_sparse_blk128_sm100a.cu PROPERTIES
+            COMPILE_OPTIONS "-gencode;arch=compute_100a,code=sm_100a;-DVSA_BHSD=true")
+    endif()
+
     Python_add_library(fastvideo_kernel_ops MODULE USE_SABI ${SKBUILD_SABI_VERSION} WITH_SOABI
         ${EXTENSION_SOURCES}
     )
@@ -333,9 +428,16 @@ if(BUILD_CXX_KERNELS)
 
     # Build compile definitions list
     set(COMPILE_DEFS TORCH_EXTENSION_NAME=fastvideo_kernel_ops)
+    if(ENABLE_VSA_SM100A)
+        list(APPEND COMPILE_DEFS TK_COMPILE_BLOCK_SPARSE_VSA_SM100A)
+    endif()
     if(ENABLE_TK_KERNELS)
         list(APPEND COMPILE_DEFS TK_COMPILE_ST_ATTN TK_COMPILE_BLOCK_SPARSE)
     endif()
+    if(ENABLE_ULYSSES_A2A)
+        list(APPEND COMPILE_DEFS FASTVIDEO_KERNEL_COMPILE_ULYSSES_A2A)
+    endif()
+
 
     target_compile_definitions(fastvideo_kernel_ops PRIVATE ${COMPILE_DEFS})
 
@@ -346,6 +448,9 @@ if(BUILD_CXX_KERNELS)
     # Link against Torch libraries to avoid undefined symbols at import time
     # (e.g., torch::autograd vtables) when loading the extension module.
     target_link_libraries(fastvideo_kernel_ops PRIVATE ${TORCH_LIBRARIES})
+    if(ENABLE_ULYSSES_A2A)
+        target_link_libraries(fastvideo_kernel_ops PRIVATE "${NCCL_LIBRARY}")
+    endif()
 
     # Also link against libtorch_python to satisfy Python-binding symbols
     # (e.g., torch::PyWarningHandler) required by torch/extension.h.
@@ -449,6 +554,7 @@ message(STATUS "host / backend:       ${CMAKE_SYSTEM_PROCESSOR} / ${GPU_BACKEND}
 message(STATUS "TORCH_CUDA_ARCH_LIST: ${TORCH_CUDA_ARCH_LIST}")
 message(STATUS "fastvideo_kernel_ops: ON  (turbodiffusion int8-gemm/quant/rmsnorm/layernorm, all listed archs)")
 message(STATUS "  + TK sta/block_sparse (sm_90a only):          ${ENABLE_TK_KERNELS}")
+message(STATUS "  + Ulysses NCCL-device all-to-all:             ${ENABLE_ULYSSES_A2A}")
 message(STATUS "fp4attn/fp4quant (sm_120a only, CUDA >= 12.8):  ${ENABLE_ATTN_QAT_INFER}")
 message(STATUS "Triton fallbacks ship in python/fastvideo_kernel/triton_kernels regardless.")
 message(STATUS "============================================================")

fastvideo-kernel/README.md

@@ -10,6 +10,8 @@ Compiled CUDA extensions (CMake, see the build summary printed at the end of eve
 |---|---|---|---|---|
 | `fastvideo_kernel._C.fastvideo_kernel_ops` | TurboDiffusion INT8 GEMM, quant, RMSNorm, LayerNorm | `csrc/turbodiffusion/` | every arch in `TORCH_CUDA_ARCH_LIST` | always built |
 | same extension, optional part | ThunderKittens sliding-tile attention (`sta_fwd`) and VSA block-sparse (`block_sparse_fwd/bwd`) | `csrc/attention/*_h100.cu` | Hopper `sm_90a` only | `FASTVIDEO_KERNEL_BUILD_TK` (AUTO = ON iff `9.0a` is in the arch list; always OFF on aarch64 hosts — TK headers don't compile there) |
+| same extension, optional part | MiniMax-H3 block-sparse VSA forward (64- and 128-token blocks) | `csrc/attention/block_sparse*_sm100a.cu` | Blackwell `sm_100a` only | ON iff `10.0a` is in `TORCH_CUDA_ARCH_LIST` |
+| same extension, optional part | fused NVLink Ulysses all-to-all | `csrc/comm/ulysses_all_to_all.cu` | CUDA | `FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A` (AUTO = ON with NCCL 2.29+ device headers and library; always OFF on ROCm) |
 | `fp4attn_cuda`, `fp4quant_cuda` | FP4 attention + quantization ("attn_qat_infer", modified SageAttention3) | `attn_qat_infer/` | consumer Blackwell `sm_120a` only, CUDA ≥ 12.8 | `FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER` (AUTO = ON iff `12.0a` is in the arch list) |
 
 Runtime-JIT kernels (no build step, ship in every wheel/image):
@@ -22,20 +24,23 @@ Runtime-JIT kernels (no build step, ship in every wheel/image):
 
 ## What gets built where, and when
 
-| Surface | Trigger | Leg | `TORCH_CUDA_ARCH_LIST` | TK | FP4 |
-|---|---|---|---|---|---|
-| PyPI wheels (`.github/workflows/publish-kernel.yml`) | version bump in `fastvideo-kernel/pyproject.toml` on main, or manual dispatch | x86_64 cu126 | `9.0a` | ON | — (CUDA < 12.8) |
-| | | x86_64 cu130 | `9.0a;12.0a` | ON | ON |
-| | | aarch64 cu130 | `10.0a;12.0a` | — | ON |
-| Docker images `ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev` (`.github/workflows/infra-build-image.yml`) | `docker/Dockerfile` changes on main, or manual dispatch | amd64 cuda12.6.3 + cuda13.0.0 | `9.0a` | ON | — |
-| | | arm64 cuda12.6.3 (GH200) | `9.0a` | — (aarch64) | — |
-| | | arm64 cuda13.0.0 (GB10 / DGX Spark) | `12.1` | — | — |
-| Local `./build.sh` | manual | probes the visible GPU via torch | detected | ON iff sm_90 (non-aarch64 host) | ON iff sm_120 |
+| Surface | Trigger | Leg | `TORCH_CUDA_ARCH_LIST` | TK | Ulysses | FP4 |
+|---|---|---|---|---|---|---|
+| PyPI wheels (`.github/workflows/publish-kernel.yml`) | version bump in `fastvideo-kernel/pyproject.toml` on main, or manual dispatch | x86_64 cu126 | `9.0a` | ON | AUTO | — (CUDA < 12.8) |
+| | | x86_64 cu130 | `9.0a;10.0a;12.0a` | ON | AUTO | ON |
+| | | aarch64 cu130 | `10.0a;12.0a` | — | AUTO | ON |
+| Docker images `ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev` (`.github/workflows/infra-build-image.yml`) | `docker/Dockerfile` changes on main, or manual dispatch | amd64 cuda12.6.3 + cuda13.0.0 | `9.0a` | ON | AUTO | — |
+| | | arm64 cuda12.6.3 (GH200) | `9.0a` | — (aarch64) | AUTO | — |
+| | | arm64 cuda13.0.0 (GB10 / DGX Spark) | `12.1` | — | AUTO | — |
+| Local `./build.sh` | manual | probes the visible GPU via torch | detected | ON iff sm_90 (non-aarch64 host) | AUTO | ON iff sm_120 |
 
 Notes:
 
 - No Docker image ships the FP4 kernels; only the x86_64/aarch64 cu130 wheels do.
+- Both cu130 PyPI wheels ship the MiniMax-H3 sm_100a VSA forward.
 - On arm64 images (GH200 included) STA/VSA run on the Triton fallbacks, since TK never builds on aarch64.
+- Ulysses AUTO builds only when CMake finds a NCCL library and device-API headers with the 2.29 initializers. Use
+  `-DFASTVIDEO_KERNEL_BUILD_ULYSSES_A2A=ON` to require it or `OFF` to test the portable build.
 - Kernel tests run on Buildkite GPU CI for PRs touching `fastvideo-kernel/**` (see `.buildkite/pipeline.yml`).
 
 ## Installation

fastvideo-kernel/csrc/attention/block_sparse_blk128_sm100a.cu

@@ -0,0 +1,7 @@
+// block_sparse_blk128_sm100a.cu -- the 128-token-block instantiation of the torch binding.
+//
+// Same source as block_sparse_sm100a.cu with VSA_BLK128 set: the kernel and launch land in
+// namespace vsa_blk128 (distinct symbols, no ODR clash with the blk64 objects) and the
+// exported entry point becomes block_sparse_sm100a_blk128_fwd.
+#define VSA_BLK128 true
+#include "block_sparse_sm100a.cu"

fastvideo-kernel/csrc/attention/block_sparse_launch_sm100a.cuh

@@ -0,0 +1,201 @@
+#ifndef BLOCK_SPARSE_VSA_LAUNCH_SM100A_CUH
+#define BLOCK_SPARSE_VSA_LAUNCH_SM100A_CUH
+
+// Launch surface for the sm_100a VSA block-sparse FMHA forward.
+//
+// Everything a caller needs: a POD argument struct, a predicate saying whether this build can
+// run those arguments, and one launch entry point. The benchmark in
+// block_sparse_bench_sm100a.cu and the torch binding both go through here, so there is
+// one tensormap construction and one launch configuration rather than two that can drift.
+//
+// Two compile-time knobs select the four builds:
+//   VSA_BLK128  false -> 64-token sparse blocks, true -> 128-token
+//   VSA_BHSD    false -> [token][head][dim] (BSHD), true -> [batch][head][token][dim] (BHSD)
+
+#include "block_sparse_kernel_sm100a.cuh"
+
+namespace VSA_NAMESPACE {
+
+struct BlockSparseVsaArgs {
+  const __nv_bfloat16* q;
+  const __nv_bfloat16* k;
+  const __nv_bfloat16* v;      // natural layout; only blk128 reads it (blk64 still needs v_t)
+  const __nv_bfloat16* v_t;    // unused: kept so the bench's V_T buffer still binds
+  __nv_bfloat16* o;
+  float* lse;                  // [batch, num_heads, seqlen] fp32, or nullptr
+
+  const int* q2k_idx;              // [batch*num_heads*num_blocks, max_kv] int32
+  const int* q2k_num;              // [batch*num_heads*num_blocks] int32
+  const int* variable_block_sizes; // [num_blocks] int32, valid tokens per block
+
+  int batch;
+  int num_heads;
+  int seqlen;
+  int head_dim;
+  int num_blocks;
+  int max_kv;
+  float sm_scale;
+};
+
+// cudaSuccess iff this build can run `a`. Deliberately conservative: the caller is expected
+// to fall back to its own implementation rather than get a wrong answer.
+__host__ inline cudaError_t block_sparse_supported(const BlockSparseVsaArgs& a) {
+  if (a.head_dim != HEAD_DIM) return cudaErrorInvalidValue;      // compile-time in the kernel
+  if (a.num_blocks % 2 != 0) return cudaErrorInvalidValue;       // a CTA owns an adjacent pair
+  if (a.seqlen != a.num_blocks * BLOCK) return cudaErrorInvalidValue;
+  if (a.max_kv < 1 || a.num_blocks < 1) return cudaErrorInvalidValue;
+  if (a.q == nullptr || a.k == nullptr || a.o == nullptr) return cudaErrorInvalidValue;
+  if (a.q2k_idx == nullptr || a.q2k_num == nullptr) return cudaErrorInvalidValue;
+  // FastVideo always supplies this; without it padded keys would be attended as real zeros.
+  if (a.variable_block_sizes == nullptr) return cudaErrorInvalidValue;
+  // V is read MN-major at BOTH block sizes now, so no pre-transposed V_T is ever needed.
+  if (a.v == nullptr) return cudaErrorInvalidValue;
+  return cudaSuccess;
+}
+
+__host__ inline cudaError_t launch_block_sparse_sm100a(const BlockSparseVsaArgs& a,
+                                                           cudaStream_t stream) {
+  const cudaError_t sup = block_sparse_supported(a);
+  if (sup != cudaSuccess) return sup;
+
+  const int B = a.batch, H = a.num_heads, S = a.seqlen, hd = a.head_dim;
+  const int num_blocks = a.num_blocks, max_kv = a.max_kv;
+  const long tq = (long)B * S;
+  const int packed_mtiles_per_seq = num_blocks / 2;
+  const int total_work = B * H * packed_mtiles_per_seq;
+  constexpr bool BHSD = VSA_BHSD;
+
+  CUtensorMap tq_, tk_, tvt_, tv_, to_;
+  {
+    uint64_t gd[4] = { (uint64_t)SUB_COLS_BF16, BHSD ? (uint64_t)((long)B * H) : (uint64_t)H,
+                       BHSD ? (uint64_t)S : (uint64_t)tq, (uint64_t)Q_SUBTILES };
+    uint64_t gs[3] = { BHSD ? (uint64_t)((long)S * hd) * 2u : (uint64_t)hd * 2u,
+                       BHSD ? (uint64_t)hd * 2u : (uint64_t)((long)H * hd) * 2u,
+                       (uint64_t)SUB_COLS_BF16 * 2u };
+    uint32_t bd[4] = { (uint32_t)SUB_COLS_BF16, 1u, (uint32_t)M_TILE, (uint32_t)Q_SUBTILES };
+    uint32_t es[4] = { 1u, 1u, 1u, 1u };
+    if (cuTensorMapEncodeTiled(&tq_, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, 4,
+                               const_cast<__nv_bfloat16*>(a.q), gd, gs, bd, es,
+                               CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B,
+                               CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
+                               CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE) != CUDA_SUCCESS)
+      return cudaErrorInvalidValue;
+    if (cuTensorMapEncodeTiled(&to_, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, 4, a.o, gd, gs, bd, es,
+                               CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B,
+                               CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
+                               CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE) != CUDA_SUCCESS)
+      return cudaErrorInvalidValue;
+  }
+  {
+    uint64_t gd[4] = { (uint64_t)SUB_COLS_BF16,
+                       BHSD ? (uint64_t)S : (uint64_t)tq,
+                       BHSD ? (uint64_t)(hd / SUB_COLS_BF16)
+                            : (uint64_t)((long)H * hd / SUB_COLS_BF16),
+                       (uint64_t)((long)B * H) };
+    uint64_t gs[3] = { BHSD ? (uint64_t)hd * 2u : (uint64_t)((long)H * hd) * 2u,
+                       (uint64_t)SUB_COLS_BF16 * 2u,
+                       (uint64_t)((long)S * hd) * 2u };
+    uint32_t bd[4] = { (uint32_t)SUB_COLS_BF16, (uint32_t)BLOCK,
+                       BLK128 ? (uint32_t)K_SUBTILES : 1u, 1u };
+    uint32_t es[4] = { 1u, 1u, 1u, 1u };
+    if (cuTensorMapEncodeTiled(&tk_, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, BHSD ? 4 : 3,
+                               const_cast<__nv_bfloat16*>(a.k), gd, gs, bd, es,
+                               CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B,
+                               CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
+                               CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE) != CUDA_SUCCESS)
+      return cudaErrorInvalidValue;
+    // V map is byte-for-byte the K map over a.v: MN-major V needs no transpose (blk128).
+    const __nv_bfloat16* vbase = a.v ? a.v : a.k;
+    if (cuTensorMapEncodeTiled(&tv_, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, BHSD ? 4 : 3,
+                               const_cast<__nv_bfloat16*>(vbase), gd, gs, bd, es,
+                               CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B,
+                               CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
+                               CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE) != CUDA_SUCCESS)
+      return cudaErrorInvalidValue;
+  }
+  // V_T map: blk64 only. Unused at blk128 but must still be a valid tensormap to pass by value.
+  {
+    const __nv_bfloat16* vt = a.v_t ? a.v_t : a.k;
+    if constexpr (BLK128) {
+      uint64_t gd[3] = { (uint64_t)SUB_COLS_BF16, (uint64_t)((long)H * hd),
+                         (uint64_t)((long)tq / SUB_COLS_BF16) };
+      uint64_t gs[2] = { (uint64_t)tq * 2u, (uint64_t)SUB_COLS_BF16 * 2u };
+      uint32_t bd[3] = { (uint32_t)SUB_COLS_BF16, (uint32_t)hd, (uint32_t)V_SUBTILES };
+      uint32_t es[3] = { 1u, 1u, 1u };
+      if (cuTensorMapEncodeTiled(&tvt_, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, 3,
+                                 const_cast<__nv_bfloat16*>(vt), gd, gs, bd, es,
+                                 CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B,
+                                 CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
+                                 CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE) != CUDA_SUCCESS)
+        return cudaErrorInvalidValue;
+    } else {
+      if (make_tma_2d_tiled(&tvt_, const_cast<__nv_bfloat16*>(vt), (long)H * hd, (int)tq, hd,
+                            SUB_COLS_BF16, 2, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,
+                            CU_TENSOR_MAP_SWIZZLE_128B) != cudaSuccess)
+        return cudaErrorInvalidValue;
+    }
+  }
+
+  const size_t smem =
+        (size_t)2 * Q_TILE_BYTES + NUM_KV_STAGES * KV_RING_SLOT_BYTES
+      + (size_t)2 * M_TILE * HEAD_DIM * sizeof(__nv_bfloat16)
+      + (2 * NUM_KV_STAGES + 22) * 8
+      + (size_t)CLC_STAGES * (2 * 8 + 16) + 16
+      + 8
+      + (size_t)2 * STAT_REGIONS * STATS * sizeof(float)
+      + 256;
+
+#ifndef VSA_NAMED_BAR
+#define VSA_NAMED_BAR false
+#endif
+#ifndef VSA_THROTTLE
+#define VSA_THROTTLE false
+#endif
+#ifndef VSA_USE_CLC
+#define VSA_USE_CLC true
+#endif
+  constexpr bool FULL_NAMED_BAR = VSA_NAMED_BAR, EX2_EMU = true, SPLIT_P = true,
+                 SOFTMAX_THROTTLE = VSA_THROTTLE, USE_CLC = VSA_USE_CLC,
+                 Q_RASTER = true, MHA = true;
+  auto kfn = &fmha_context_bf16_gen_kernel<32, FULL_NAMED_BAR, EX2_EMU, SPLIT_P,
+                                           SOFTMAX_THROTTLE, USE_CLC, Q_RASTER, MHA,
+                                           /*RESCALE_THRESHOLD=*/8, /*BHSD=*/VSA_BHSD>;
+  cudaError_t e = cudaFuncSetAttribute(kfn, cudaFuncAttributeMaxDynamicSharedMemorySize, (int)smem);
+  if (e != cudaSuccess) return e;
+
+  const unsigned long long magic0 = make_magic((unsigned)(H * packed_mtiles_per_seq));
+  const unsigned long long magic1 = make_magic((unsigned)H);
+  const unsigned long long magic2 = make_magic((unsigned)packed_mtiles_per_seq);
+  const float scale_log2 = a.sm_scale * (float)M_LOG2E;
+
+  int numSM = 0;
+  e = cudaDeviceGetAttribute(&numSM, cudaDevAttrMultiProcessorCount, 0);
+  if (e != cudaSuccess) return e;
+  const int num_ctas = USE_CLC ? total_work : (total_work < numSM ? total_work : numSM);
+  dim3 grid(num_ctas, 1, 1), block(N_WARPS * 32, 1, 1);
+
+  if (USE_CLC) {
+    cudaLaunchConfig_t cfg = {};
+    cfg.gridDim = grid; cfg.blockDim = block; cfg.dynamicSmemBytes = smem; cfg.stream = stream;
+    cudaLaunchAttribute cfgAttr[1];
+    cfgAttr[0].id = cudaLaunchAttributeClusterDimension;
+    cfgAttr[0].val.clusterDim.x = 1; cfgAttr[0].val.clusterDim.y = 1;
+    cfgAttr[0].val.clusterDim.z = 1;
+    cfg.attrs = cfgAttr; cfg.numAttrs = 1;
+    return cudaLaunchKernelEx(&cfg, kfn, tq_, tk_, tvt_, tv_, to_, S, H, scale_log2, B,
+                              num_blocks, packed_mtiles_per_seq, max_kv, magic0, magic1, magic2,
+                              a.q2k_idx, a.q2k_num, a.variable_block_sizes, a.lse);
+  }
+  kfn<<<grid, block, smem, stream>>>(tq_, tk_, tvt_, tv_, to_, S, H, scale_log2, B, num_blocks,
+                                     packed_mtiles_per_seq, max_kv, magic0, magic1, magic2,
+                                     a.q2k_idx, a.q2k_num, a.variable_block_sizes, a.lse);
+  return cudaGetLastError();
+}
+
+}  // namespace VSA_NAMESPACE
+
+// Callers (the bench, the torch binding) keep using unqualified names; each translation unit
+// only ever sees the one configuration its VSA_BLK128 selected.
+using namespace VSA_NAMESPACE;
+
+#endif  // BLOCK_SPARSE_VSA_LAUNCH_SM100A_CUH

fastvideo-kernel/csrc/attention/block_sparse_sm100a.cu

@@ -0,0 +1,114 @@
+// block_sparse_sm100a.cu -- torch binding for the sm_100a VSA block-sparse FMHA forward.
+//
+// Forward only: returns (out, lse) so FastVideo's existing Triton backward keeps working
+// unchanged. lse is exactly the M tensor triton_block_sparse_attn_forward writes --
+// max(qk * qk_scale) + log2(l), [B, H, S] fp32 -- which is what lets
+// block_sparse_attn_backward_triton run against our forward untouched.
+//
+// The build is fixed at compile time by two flags, so one extension carries one configuration:
+//   VSA_BLK128  false -> 64-token sparse blocks, true -> 128-token
+//   VSA_BHSD    false -> [B, S, H, D], true -> [B, H, S, D]
+#include <torch/extension.h>
+
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
+
+#include "block_sparse_launch_sm100a.cuh"
+
+namespace {
+
+void check_qkv(const torch::Tensor& t, const char* name, int64_t B, int64_t H, int64_t S,
+               int64_t D) {
+  TORCH_CHECK(t.is_cuda(), name, " must be a CUDA tensor");
+  TORCH_CHECK(t.scalar_type() == at::kBFloat16, name, " must be bfloat16, got ", t.scalar_type());
+  TORCH_CHECK(t.dim() == 4, name, " must be 4-D, got ", t.dim(), " dims");
+  TORCH_CHECK(t.is_contiguous(), name, " must be contiguous");
+  if (VSA_BHSD) {
+    TORCH_CHECK(t.size(0) == B && t.size(1) == H && t.size(2) == S && t.size(3) == D, name,
+                " has shape ", t.sizes(), ", expected [", B, ",", H, ",", S, ",", D, "]");
+  } else {
+    TORCH_CHECK(t.size(0) == B && t.size(1) == S && t.size(2) == H && t.size(3) == D, name,
+                " has shape ", t.sizes(), ", expected [", B, ",", S, ",", H, ",", D, "]");
+  }
+}
+
+void check_index(const torch::Tensor& t, const char* name) {
+  TORCH_CHECK(t.is_cuda(), name, " must be a CUDA tensor");
+  TORCH_CHECK(t.scalar_type() == at::kInt, name, " must be int32, got ", t.scalar_type());
+  TORCH_CHECK(t.is_contiguous(), name, " must be contiguous");
+}
+
+}  // namespace
+
+// The exported symbol carries the block size: block_sparse_sm100a_fwd is the 64-token build,
+// block_sparse_sm100a_blk128_fwd the 128-token one (block_sparse_blk128_sm100a.cu re-includes
+// this file with VSA_BLK128 set). The python backend picks by the metadata's block size.
+#if VSA_BLK128
+#define BLOCK_SPARSE_SM100A_FWD block_sparse_sm100a_blk128_fwd
+#else
+#define BLOCK_SPARSE_SM100A_FWD block_sparse_sm100a_fwd
+#endif
+
+// Returns {out} or {out, lse}. Layout of out matches the inputs.
+std::vector<torch::Tensor> BLOCK_SPARSE_SM100A_FWD(torch::Tensor q, torch::Tensor k,
+                                                       torch::Tensor v,
+                                                       c10::optional<torch::Tensor> v_t,
+                                                       torch::Tensor q2k_idx,
+                                                       torch::Tensor q2k_num,
+                                                       torch::Tensor variable_block_sizes,
+                                                       double sm_scale, bool need_lse) {
+  const at::cuda::OptionalCUDAGuard guard(device_of(q));
+
+  const int64_t B = q.size(0);
+  const int64_t H = VSA_BHSD ? q.size(1) : q.size(2);
+  const int64_t S = VSA_BHSD ? q.size(2) : q.size(1);
+  const int64_t D = q.size(3);
+
+  check_qkv(q, "q", B, H, S, D);
+  check_qkv(k, "k", B, H, S, D);
+  check_qkv(v, "v", B, H, S, D);
+  check_index(q2k_idx, "q2k_idx");
+  check_index(q2k_num, "q2k_num");
+  check_index(variable_block_sizes, "variable_block_sizes");
+
+  const int64_t num_blocks = variable_block_sizes.numel();
+  const int64_t max_kv = q2k_idx.size(-1);
+  TORCH_CHECK(S == num_blocks * BLOCK, "seqlen ", S, " must equal num_blocks (", num_blocks,
+              ") * ", BLOCK, "; FastVideo pads the sequence up to whole blocks");
+
+  auto out = torch::empty_like(q);
+  torch::Tensor lse;
+  if (need_lse) lse = torch::empty({B, H, S}, q.options().dtype(torch::kFloat32));
+
+  BlockSparseVsaArgs a{};
+  a.q = reinterpret_cast<const __nv_bfloat16*>(q.data_ptr());
+  a.k = reinterpret_cast<const __nv_bfloat16*>(k.data_ptr());
+  a.v = reinterpret_cast<const __nv_bfloat16*>(v.data_ptr());
+  a.v_t = v_t.has_value() ? reinterpret_cast<const __nv_bfloat16*>(v_t->data_ptr()) : nullptr;
+  a.o = reinterpret_cast<__nv_bfloat16*>(out.data_ptr());
+  a.lse = need_lse ? lse.data_ptr<float>() : nullptr;
+  a.q2k_idx = q2k_idx.data_ptr<int>();
+  a.q2k_num = q2k_num.data_ptr<int>();
+  a.variable_block_sizes = variable_block_sizes.data_ptr<int>();
+  a.batch = (int)B;
+  a.num_heads = (int)H;
+  a.seqlen = (int)S;
+  a.head_dim = (int)D;
+  a.num_blocks = (int)num_blocks;
+  a.max_kv = (int)max_kv;
+  a.sm_scale = (float)sm_scale;
+
+  // Report an unsupported regime loudly rather than returning plausible-looking wrong values.
+  TORCH_CHECK(block_sparse_supported(a) == cudaSuccess,
+              "block_sparse_sm100a: unsupported configuration -- requires head_dim==",
+              HEAD_DIM, ", an even num_blocks, seqlen == num_blocks*", BLOCK,
+              ", and a variable_block_sizes tensor. Got head_dim=", D, " num_blocks=",
+              num_blocks, " seqlen=", S);
+
+  const cudaError_t err = launch_block_sparse_sm100a(a, at::cuda::getCurrentCUDAStream());
+  TORCH_CHECK(err == cudaSuccess,
+              "block_sparse_sm100a launch failed: ", cudaGetErrorString(err));
+
+  if (need_lse) return {out, lse};
+  return {out};
+}

fastvideo-kernel/csrc/attention/primitives.cuh

@@ -0,0 +1,877 @@
+// primitives.cuh -- device primitives for the sm_100a VSA block-sparse attention
+// forward: tcgen05 (alloc / mma / ld / st / commit / wait / fence), TMA load / store /
+// tensormap, mbarrier, cluster launch control, setmaxnreg, fast math, and the FMHA helpers.
+//
+// Generated and pruned to what the kernel reaches -- do not edit by hand.
+#pragma once
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cuda.h>
+#include <cuda_bf16.h>
+#include <cuda_fp16.h>
+#include <cuda_runtime.h>
+#include <cmath>
+#include <cassert>
+#include <cstring>
+#include <vector_types.h>
+
+#ifndef CUDA_CHECK
+#define CUDA_CHECK(stmt) do {                                                 \
+    cudaError_t _e = (stmt);                                                  \
+    if (_e != cudaSuccess) {                                                  \
+      fprintf(stderr, "CUDA error %s:%d: %s -> %s\n",                         \
+              __FILE__, __LINE__, #stmt, cudaGetErrorString(_e));             \
+      std::exit(1);                                                           \
+    }                                                                         \
+  } while (0)
+#endif
+
+__device__ __forceinline__
+uint64_t mbarrier_arrive(uint32_t mbar_smem) {
+  uint64_t state;
+  asm volatile("mbarrier.arrive.shared::cta.b64 %0, [%1];\n"
+               : "=l"(state) : "r"(mbar_smem) : "memory");
+  return state;
+}
+
+__device__ __forceinline__
+void mbarrier_arrive_nostate(uint32_t mbar_smem) {
+  asm volatile("mbarrier.arrive.shared::cta.b64 _, [%0];\n"
+               :: "r"(mbar_smem) : "memory");
+}
+
+__device__ __forceinline__
+void mbarrier_arrive_cluster_default(uint32_t cluster_smem_addr) {
+  asm volatile("mbarrier.arrive.shared::cluster.b64 _, [%0];\n"
+               :: "r"(cluster_smem_addr) : "memory");
+}
+
+__device__ __forceinline__
+void mbarrier_arrive_expect_tx(uint32_t mbar_smem, uint32_t expected_bytes) {
+  asm volatile("mbarrier.arrive.expect_tx.release.cta.shared::cta.b64 _, [%0], %1;\n"
+               :: "r"(mbar_smem), "r"(expected_bytes) : "memory");
+}
+
+__device__ __forceinline__
+void mbarrier_wait_parity_suspend(uint32_t mbar_smem, uint32_t phase_parity) {
+  asm volatile(
+    "{\n"
+    ".reg .pred P1;\n"
+    "LAB_WAIT:\n"
+    "mbarrier.try_wait.parity.shared::cta.b64 P1, [%0], %1, 10000000;\n"
+    "@!P1 bra.uni LAB_WAIT;\n"
+    "}\n"
+    :: "r"(mbar_smem), "r"(phase_parity) : "memory");
+}
+
+__device__ __forceinline__
+void mbarrier_wait_parity(uint32_t mbar_smem, uint32_t phase_parity) {
+  asm volatile(
+    "{\n"
+    ".reg .pred P1;\n"
+    "LAB_WAIT_HOT:\n"
+    "mbarrier.try_wait.parity.shared::cta.b64 P1, [%0], %1;\n"
+    "@!P1 bra.uni LAB_WAIT_HOT;\n"
+    "}\n"
+    :: "r"(mbar_smem), "r"(phase_parity) : "memory");
+}
+
+__device__ __forceinline__
+void fence_proxy_async_shared_cta() {
+  asm volatile("fence.proxy.async.shared::cta;\n" ::: "memory");
+}
+
+__device__ __forceinline__
+void fence_proxy_async_shared() {
+  asm volatile("fence.proxy.async.shared::cta;\n" ::: "memory");
+}
+
+__device__ __forceinline__ void clc_try_cancel_async(
+    uint32_t smem_dst, uint32_t mbar_smem) {
+  asm volatile(
+    "clusterlaunchcontrol.try_cancel.async.shared::cta.mbarrier::complete_tx::bytes.b128"
+    " [%0], [%1];\n"
+    :: "r"(smem_dst), "r"(mbar_smem) : "memory");
+}
+
+__device__ __forceinline__ void clc_load_response(
+    uint32_t smem_slot, uint32_t& r0, uint32_t& r1,
+    uint32_t& r2, uint32_t& r3) {
+  asm volatile("ld.shared::cta.v4.b32 {%0, %1, %2, %3}, [%4];\n"
+               : "=r"(r0), "=r"(r1), "=r"(r2), "=r"(r3)
+               : "r"(smem_slot));
+}
+
+template <int NUM_STAGES>
+struct MbarrierPhaseTracker {
+
+  uint32_t phase[NUM_STAGES];
+  int idx;
+
+  __device__ __forceinline__
+  void init() {
+    for (int i = 0; i < NUM_STAGES; ++i) phase[i] = 0;
+    idx = 0;
+  }
+
+  __device__ __forceinline__
+  uint32_t current_phase() const { return phase[idx]; }
+
+  __device__ __forceinline__
+  void advance() {
+    phase[idx] ^= 1u;
+    idx = (idx + 1) % NUM_STAGES;
+  }
+
+  __device__ __forceinline__
+  int stage() const { return idx; }
+};
+
+template <int NUM_STAGES>
+struct PhaseTracker {
+  int stage;
+  uint32_t phase;
+
+  __device__ __forceinline__
+  PhaseTracker() : stage(0), phase(0) {}
+
+  __device__ __forceinline__
+  void advance() {
+    stage++;
+    if (stage == NUM_STAGES) {
+      stage = 0;
+      phase ^= 1;
+    }
+  }
+
+  __device__ __forceinline__
+  int get_stage() const { return stage; }
+
+  __device__ __forceinline__
+  uint32_t get_phase() const { return phase; }
+};
+
+template <int NUM_STAGES>
+struct EmptyPhaseTracker {
+  int stage;
+  uint32_t phase;
+
+  __device__ __forceinline__
+  EmptyPhaseTracker() : stage(0), phase(1) {}
+
+  __device__ __forceinline__
+  void advance() {
+    stage++;
+    if (stage == NUM_STAGES) {
+      stage = 0;
+      phase ^= 1;
+    }
+  }
+
+  __device__ __forceinline__
+  int get_stage() const { return stage; }
+
+  __device__ __forceinline__
+  uint32_t get_phase() const { return phase; }
+};
+
+template <int STAGES>
+__device__ __forceinline__
+void advance_stage_phase(int& stage, uint32_t& phase) {
+  ++stage;
+  if (stage == STAGES) {
+    stage = 0;
+    phase ^= 1u;
+  }
+}
+
+static constexpr uint32_t SM100_CLC_PEER_MASK = 0xFEFFFFFF;
+
+struct ClcTileInfo {
+  int  m_tile;
+  int  n_tile;
+  bool valid;
+};
+
+enum class ClcRasterOrder { AlongN, AlongM };
+
+__device__ __forceinline__
+void clc_arrive_expect_tx_cta(uint32_t clc_full_local_addr, uint32_t tx_bytes) {
+  if ((threadIdx.x & 31) == 0) {
+    mbarrier_arrive_expect_tx(clc_full_local_addr, tx_bytes);
+  }
+}
+
+__device__ __forceinline__
+void clc_consumer_release(uint32_t clc_empty_local_addr) {
+  uint32_t peer0_addr = clc_empty_local_addr & SM100_CLC_PEER_MASK;
+  mbarrier_arrive_cluster_default(peer0_addr);
+}
+
+__device__ __forceinline__
+void clc_consumer_release_cta(uint32_t clc_empty_local_addr) {
+  mbarrier_arrive_nostate(clc_empty_local_addr);
+}
+
+template <int CLUSTER_SHAPE_M, int CLUSTER_SHAPE_N, ClcRasterOrder ORDER>
+__device__ __forceinline__
+ClcTileInfo clc_parse_response(uint32_t resp_smem_addr) {
+  uint32_t d0, d1, d2, d3;
+  fence_proxy_async_shared_cta();
+  clc_load_response(resp_smem_addr, d0, d1, d2, d3);
+  const int  ctaid_x = static_cast<int>(d0);
+  const int  ctaid_y = static_cast<int>(d1 & 0xFFFFu);
+  const bool valid   = (d2 & 1u) != 0u;
+  (void)d3;
+
+  ClcTileInfo info;
+  info.valid = valid;
+  if constexpr (ORDER == ClcRasterOrder::AlongN) {
+    info.m_tile = ctaid_y / CLUSTER_SHAPE_M;
+    info.n_tile = ctaid_x / CLUSTER_SHAPE_N;
+  } else {
+    info.m_tile = ctaid_x / CLUSTER_SHAPE_M;
+    info.n_tile = ctaid_y / CLUSTER_SHAPE_N;
+  }
+  return info;
+}
+
+template <int CLUSTER_SHAPE_M, int CLUSTER_SHAPE_N, ClcRasterOrder ORDER,
+          int CTA_GROUP = 2, bool SUSPEND = false>
+__device__ __forceinline__
+ClcTileInfo clc_fetch_next_tile(
+    uint64_t* clc_full_bar, uint64_t* clc_empty_bar, uint32_t* clc_response,
+    int clc_cons_stage, uint32_t clc_cons_phase, bool do_release) {
+  uint32_t full_addr = static_cast<uint32_t>(
+      __cvta_generic_to_shared(&clc_full_bar[clc_cons_stage]));
+  if constexpr (SUSPEND) mbarrier_wait_parity_suspend(full_addr, clc_cons_phase);
+  else                   mbarrier_wait_parity(full_addr, clc_cons_phase);
+  uint32_t resp_addr = static_cast<uint32_t>(
+      __cvta_generic_to_shared(&clc_response[clc_cons_stage * 4]));
+  ClcTileInfo t = clc_parse_response<
+      CLUSTER_SHAPE_M, CLUSTER_SHAPE_N, ORDER>(resp_addr);
+  if (do_release) {
+    uint32_t empty_local = static_cast<uint32_t>(
+        __cvta_generic_to_shared(&clc_empty_bar[clc_cons_stage]));
+    if constexpr (CTA_GROUP == 1) {
+      clc_consumer_release_cta(empty_local);
+    } else {
+      clc_consumer_release(empty_local);
+    }
+  }
+  return t;
+}
+
+template <int STAGES = 2>
+__device__ __forceinline__
+void clc_fetch_next_tile_advance(int& clc_cons_stage,
+                                 uint32_t& clc_cons_phase) {
+  advance_stage_phase<STAGES>(clc_cons_stage, clc_cons_phase);
+}
+
+__device__ __forceinline__ unsigned fdiv(unsigned n, unsigned long long pk) {
+  unsigned M = (unsigned)pk;
+  if (M == 0u) return n;
+  return __umulhi(n, M) >> (unsigned)(pk >> 32);
+}
+__host__ inline unsigned long long make_magic(unsigned d) {
+  if (d <= 1u) return 0ULL;
+  unsigned l = 0; while ((1u << (l + 1)) <= d) ++l;
+  unsigned p = 31u + l;
+  unsigned long long m = ((1ull << p) + (unsigned long long)d - 1ull) / d;
+  return (m & 0xffffffffULL) | ((unsigned long long)(p - 32u) << 32);
+}
+
+template <int BARRIER_ID>
+__device__ __forceinline__
+void bar_sync(uint32_t thread_count) {
+  static_assert(BARRIER_ID >= 0 && BARRIER_ID <= 15,
+                "bar.sync: BARRIER_ID must be in [0, 15]");
+  asm volatile("bar.sync %0, %1;\n"
+               :: "n"(BARRIER_ID), "r"(thread_count) : "memory");
+}
+
+template <int BARRIER_ID>
+__device__ __forceinline__
+void bar_arrive(uint32_t thread_count) {
+  static_assert(BARRIER_ID >= 0 && BARRIER_ID <= 15,
+                "bar.arrive: BARRIER_ID must be in [0, 15]");
+  asm volatile("bar.arrive %0, %1;\n"
+               :: "n"(BARRIER_ID), "r"(thread_count) : "memory");
+}
+
+__device__ __forceinline__ void full_bar_arrive(int m_tile, int band) {
+  switch (1 + m_tile * 4 + band) {
+    case 1: bar_arrive<1>(64); break;
+    case 2: bar_arrive<2>(64); break;
+    case 3: bar_arrive<3>(64); break;
+    case 4: bar_arrive<4>(64); break;
+    case 5: bar_arrive<5>(64); break;
+    case 6: bar_arrive<6>(64); break;
+    case 7: bar_arrive<7>(64); break;
+    case 8: bar_arrive<8>(64); break;
+  }
+}
+__device__ __forceinline__ void full_bar_wait(int m_tile, int band) {
+  switch (1 + m_tile * 4 + band) {
+    case 1: bar_sync<1>(64); break;
+    case 2: bar_sync<2>(64); break;
+    case 3: bar_sync<3>(64); break;
+    case 4: bar_sync<4>(64); break;
+    case 5: bar_sync<5>(64); break;
+    case 6: bar_sync<6>(64); break;
+    case 7: bar_sync<7>(64); break;
+    case 8: bar_sync<8>(64); break;
+  }
+}
+
+template <bool IS_CAUSAL, int K_TILE>
+__device__ __forceinline__ void mask_s_row_r2p(float* scores, int k_offset, int q_pos, int seqlen_k) {
+  int n_keep = seqlen_k - k_offset;
+  if constexpr (IS_CAUSAL) {
+    const int causal = q_pos - k_offset + 1;
+    n_keep = n_keep < causal ? n_keep : causal;
+  }
+  #pragma unroll
+  for (int s = 0; s < K_TILE / 32; ++s) {
+    int m = (s + 1) * 32 - n_keep;
+    m = m < 0 ? 0 : (m > 32 ? 32 : m);
+    const uint32_t keep = (m >= 32) ? 0u : (0xFFFFFFFFu >> m);
+    #pragma unroll
+    for (int i = 0; i < 32; ++i)
+      if (!(keep & (1u << i))) scores[s * 32 + i] = -INFINITY;
+  }
+}
+
+template <int CTA_GROUP>
+__device__ __forceinline__ void tcgen05_alloc(uint32_t smem_dst_ptr,
+                                              uint32_t n_cols) {
+  static_assert(CTA_GROUP == 1 || CTA_GROUP == 2,
+                "tcgen05_alloc: CTA_GROUP must be 1 or 2");
+  if constexpr (CTA_GROUP == 1) {
+    asm volatile(
+      "tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%0], %1;\n"
+      :: "r"(smem_dst_ptr), "r"(n_cols));
+  } else {
+    asm volatile(
+      "tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%0], %1;\n"
+      :: "r"(smem_dst_ptr), "r"(n_cols));
+  }
+}
+
+__device__ __forceinline__ void tcgen05_st_32x32b_x16(
+    uint32_t tmem_addr, const uint32_t (&r)[16]) {
+  asm volatile("tcgen05.st.sync.aligned.32x32b.x16.b32 "
+               "[%0], {%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,"
+               "%11,%12,%13,%14,%15,%16};\n"
+               :: "r"(tmem_addr),
+                  "r"(r[0]),"r"(r[1]),"r"(r[2]),"r"(r[3]),
+                  "r"(r[4]),"r"(r[5]),"r"(r[6]),"r"(r[7]),
+                  "r"(r[8]),"r"(r[9]),"r"(r[10]),"r"(r[11]),
+                  "r"(r[12]),"r"(r[13]),"r"(r[14]),"r"(r[15]));
+}
+
+__device__ __forceinline__ void tcgen05_st_32x32b_x32(
+    uint32_t tmem_addr, const uint32_t (&r)[32]) {
+  asm volatile("tcgen05.st.sync.aligned.32x32b.x32.b32 "
+               "[%0], {%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,"
+               "%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,"
+               "%21,%22,%23,%24,%25,%26,%27,%28,%29,%30,"
+               "%31,%32};\n"
+               :: "r"(tmem_addr),
+                  "r"(r[0]),"r"(r[1]),"r"(r[2]),"r"(r[3]),
+                  "r"(r[4]),"r"(r[5]),"r"(r[6]),"r"(r[7]),
+                  "r"(r[8]),"r"(r[9]),"r"(r[10]),"r"(r[11]),
+                  "r"(r[12]),"r"(r[13]),"r"(r[14]),"r"(r[15]),
+                  "r"(r[16]),"r"(r[17]),"r"(r[18]),"r"(r[19]),
+                  "r"(r[20]),"r"(r[21]),"r"(r[22]),"r"(r[23]),
+                  "r"(r[24]),"r"(r[25]),"r"(r[26]),"r"(r[27]),
+                  "r"(r[28]),"r"(r[29]),"r"(r[30]),"r"(r[31]));
+}
+
+__device__ __forceinline__ void tcgen05_commit1_lead(uint32_t lead, uint32_t mbar_smem_addr) {
+  asm volatile(
+    "{\n\t"
+    ".reg .pred q;\n\t"
+    "setp.ne.b32 q, %0, 0;\n\t"
+    "@q tcgen05.commit.cta_group::1.mbarrier::arrive::one.b64 [%1];\n\t"
+    "}\n"
+    :: "r"(lead), "r"(mbar_smem_addr));
+}
+
+__device__ __forceinline__ void tcgen05_wait_st() {
+  asm volatile("tcgen05.wait::st.sync.aligned;\n" ::: "memory");
+}
+
+__device__ __forceinline__ void tcgen05_fence_before_thread_sync() {
+  asm volatile("tcgen05.fence::before_thread_sync;\n" ::: "memory");
+}
+
+__device__ __forceinline__
+void tma_load_2d(uint32_t smem_dst, const void* tensormap_ptr,
+                 uint32_t mbar_smem, int coord_x, int coord_y) {
+  asm volatile(
+    "cp.async.bulk.tensor.2d.shared::cluster.global.tile"
+    ".mbarrier::complete_tx::bytes"
+    " [%0], [%1, {%3, %4}], [%2];\n"
+    :: "r"(smem_dst), "l"(tensormap_ptr),
+       "r"(mbar_smem), "r"(coord_x), "r"(coord_y)
+    : "memory");
+}
+
+__device__ __forceinline__
+void tma_load_3d(uint32_t smem_dst, const void* tensormap_ptr,
+                 uint32_t mbar_smem, int c0, int c1, int c2) {
+  asm volatile(
+    "cp.async.bulk.tensor.3d.shared::cluster.global.tile"
+    ".mbarrier::complete_tx::bytes"
+    " [%0], [%1, {%3, %4, %5}], [%2];\n"
+    :: "r"(smem_dst), "l"(tensormap_ptr),
+       "r"(mbar_smem), "r"(c0), "r"(c1), "r"(c2)
+    : "memory");
+}
+
+__device__ __forceinline__
+void tma_load_4d(uint32_t smem_dst, const void* tensormap_ptr,
+                 uint32_t mbar_smem, int c0, int c1, int c2, int c3) {
+  asm volatile(
+    "cp.async.bulk.tensor.4d.shared::cluster.global.tile"
+    ".mbarrier::complete_tx::bytes"
+    " [%0], [%1, {%3, %4, %5, %6}], [%2];\n"
+    :: "r"(smem_dst), "l"(tensormap_ptr),
+       "r"(mbar_smem), "r"(c0), "r"(c1), "r"(c2), "r"(c3)
+    : "memory");
+}
+
+template <int CTA_GROUP>
+__device__ __forceinline__ void tcgen05_dealloc(uint32_t tmem_addr,
+                                                uint32_t n_cols) {
+  static_assert(CTA_GROUP == 1 || CTA_GROUP == 2,
+                "tcgen05_dealloc: CTA_GROUP must be 1 or 2");
+  if constexpr (CTA_GROUP == 1) {
+    asm volatile(
+      "tcgen05.dealloc.cta_group::1.sync.aligned.b32 %0, %1;\n"
+      :: "r"(tmem_addr), "r"(n_cols));
+  } else {
+    asm volatile(
+      "tcgen05.dealloc.cta_group::2.sync.aligned.b32 %0, %1;\n"
+      :: "r"(tmem_addr), "r"(n_cols));
+  }
+}
+
+__device__ __forceinline__
+void tma_store_2d(const void* tensormap_ptr, int coord_x, int coord_y,
+                  uint32_t smem_src) {
+  asm volatile(
+    "cp.async.bulk.tensor.2d.global.shared::cta.tile.bulk_group"
+    " [%0, {%1, %2}], [%3];\n"
+    :: "l"(tensormap_ptr), "r"(coord_x), "r"(coord_y),
+       "r"(smem_src)
+    : "memory");
+}
+
+__device__ __forceinline__
+void tma_store_3d(const void* tensormap_ptr, int c0, int c1, int c2,
+                  uint32_t smem_src) {
+  asm volatile(
+    "cp.async.bulk.tensor.3d.global.shared::cta.tile.bulk_group"
+    " [%0, {%1, %2, %3}], [%4];\n"
+    :: "l"(tensormap_ptr), "r"(c0), "r"(c1), "r"(c2), "r"(smem_src)
+    : "memory");
+}
+
+__device__ __forceinline__
+void tma_store_4d(const void* tensormap_ptr, int c0, int c1, int c2, int c3,
+                  uint32_t smem_src) {
+  asm volatile(
+    "cp.async.bulk.tensor.4d.global.shared::cta.tile.bulk_group"
+    " [%0, {%1, %2, %3, %4}], [%5];\n"
+    :: "l"(tensormap_ptr), "r"(c0), "r"(c1), "r"(c2), "r"(c3), "r"(smem_src)
+    : "memory");
+}
+
+inline cudaError_t make_tma_2d_tiled(
+    CUtensorMap* out,
+    const void* ptr, int rows, int cols, int box_rows, int box_cols,
+    int elem_bytes, CUtensorMapDataType dtype,
+    CUtensorMapSwizzle swizzle = CU_TENSOR_MAP_SWIZZLE_128B,
+    CUtensorMapL2promotion l2 = CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
+    CUtensorMapFloatOOBfill oob = CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE) {
+  uint64_t globalDim[2]     = { (uint64_t)cols, (uint64_t)rows };
+  uint64_t globalStrides[1] = { (uint64_t)cols * (uint64_t)elem_bytes };
+  uint32_t boxDim[2]        = { (uint32_t)box_cols, (uint32_t)box_rows };
+  uint32_t elemStrides[2]   = { 1u, 1u };
+
+  CUresult r = cuTensorMapEncodeTiled(
+      out, dtype, 2,
+      const_cast<void*>(ptr), globalDim, globalStrides,
+      boxDim, elemStrides,
+      CU_TENSOR_MAP_INTERLEAVE_NONE,
+      swizzle, l2, oob);
+  return (r == CUDA_SUCCESS) ? cudaSuccess : cudaErrorInvalidValue;
+}
+
+__device__ __forceinline__
+void cp_async_bulk_commit_group() {
+  asm volatile("cp.async.bulk.commit_group;\n" ::: "memory");
+}
+
+template <int N>
+__device__ __forceinline__
+void cp_async_bulk_wait_group_read() {
+  asm volatile("cp.async.bulk.wait_group.read %0;\n" :: "n"(N) : "memory");
+}
+
+__device__ __forceinline__
+void mbarrier_init(uint32_t mbar_smem, uint32_t arrive_count) {
+  asm volatile("mbarrier.init.shared::cta.b64 [%0], %1;\n"
+               :: "r"(mbar_smem), "r"(arrive_count) : "memory");
+}
+
+template <int CTA_GROUP>
+__device__ __forceinline__ void tcgen05_relinquish_alloc_permit() {
+  static_assert(CTA_GROUP == 1 || CTA_GROUP == 2,
+                "tcgen05_relinquish_alloc_permit: CTA_GROUP must be 1 or 2");
+  if constexpr (CTA_GROUP == 1) {
+    asm volatile("tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;\n" ::);
+  } else {
+    asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;\n" ::);
+  }
+}
+
+__device__ __forceinline__
+void fence_mbarrier_init_release_cluster() {
+  asm volatile("fence.mbarrier_init.release.cluster;\n" ::: "memory");
+}
+
+__device__ __forceinline__ void tcgen05_mma_f16_ss_lead(uint32_t lead,
+    uint32_t tmem_c, uint64_t desc_a, uint64_t desc_b, uint32_t idesc,
+    bool enable_input_d) {
+  asm volatile(
+    "{\n\t"
+    ".reg .pred p, q;\n\t"
+    "setp.ne.b32 q, %0, 0;\n\t"
+    "setp.ne.b32 p, %5, 0;\n\t"
+    "@q tcgen05.mma.cta_group::1.kind::f16 [%1], %2, %3, %4, {%6, %7, %8, %9}, p;\n\t"
+    "}\n"
+    :: "r"(lead), "r"(tmem_c), "l"(desc_a), "l"(desc_b), "r"(idesc),
+       "r"(enable_input_d ? 1u : 0u), "r"(0u), "r"(0u), "r"(0u), "r"(0u));
+}
+
+__device__ __forceinline__ void tcgen05_mma_f16_ts_1sm_lead(uint32_t lead,
+    uint32_t tmem_c, uint32_t tmem_a, uint64_t desc_b, uint32_t idesc,
+    bool enable_input_d) {
+  asm volatile(
+    "{\n\t"
+    ".reg .pred p, q;\n\t"
+    "setp.ne.b32 q, %0, 0;\n\t"
+    "setp.ne.b32 p, %5, 0;\n\t"
+    "@q tcgen05.mma.cta_group::1.kind::f16 [%1], [%2], %3, %4, {%6, %7, %8, %9}, p;\n\t"
+    "}\n"
+    :: "r"(lead), "r"(tmem_c), "r"(tmem_a), "l"(desc_b), "r"(idesc),
+       "r"(enable_input_d ? 1u : 0u), "r"(0u), "r"(0u), "r"(0u), "r"(0u));
+}
+
+enum class SmemSwizzleBlackwell : uint32_t {
+  None = 0,
+  B128_32atom = 1,
+  B128 = 2,
+  B64  = 4,
+  B32  = 6,
+};
+
+__device__ __host__ __forceinline__ uint64_t build_smem_desc_blackwell(
+    uint32_t smem_addr,
+    uint32_t stride_byte_offset,
+    uint32_t leading_byte_offset,
+    SmemSwizzleBlackwell swizzle = SmemSwizzleBlackwell::B128,
+    uint32_t base_offset = 0) {
+  uint64_t d = 0;
+  d |= static_cast<uint64_t>((smem_addr >> 4) & 0x3FFF);
+  d |= static_cast<uint64_t>((leading_byte_offset >> 4) & 0x3FFF) << 16;
+  d |= static_cast<uint64_t>((stride_byte_offset >> 4) & 0x3FFF) << 32;
+  d |= static_cast<uint64_t>(1) << 46;
+  d |= (static_cast<uint64_t>(base_offset) & 0x7) << 49;
+  d |= static_cast<uint64_t>(static_cast<uint32_t>(swizzle) & 0x7) << 61;
+  return d;
+}
+
+__device__ __forceinline__
+uint32_t elect_one_sync() {
+  uint32_t elected;
+  asm volatile(
+    "{\n\t"
+    ".reg .pred p;\n\t"
+    "elect.sync %0|p, 0xffffffff;\n\t"
+    "selp.b32 %0, 1, 0, p;\n\t"
+    "}\n"
+    : "=r"(elected));
+  return elected;
+}
+
+__device__ __forceinline__
+uint32_t elect_one_sync(uint32_t membermask) {
+  uint32_t elected;
+  asm volatile(
+    "{\n\t"
+    ".reg .pred p;\n\t"
+    "elect.sync %0|p, %1;\n\t"
+    "selp.b32 %0, 1, 0, p;\n\t"
+    "}\n"
+    : "=r"(elected) : "r"(membermask));
+  return elected;
+}
+
+template <int N>
+__device__ __forceinline__
+void setmaxnreg_dec() {
+  static_assert(N >= 24 && N <= 256 && N % 8 == 0,
+                "setmaxnreg_dec: N must be in [24, 256] and a multiple of 8");
+  asm volatile("setmaxnreg.dec.sync.aligned.u32 %0;\n"
+               :: "n"(N) : "memory");
+}
+
+template <int N>
+__device__ __forceinline__
+void setmaxnreg_inc() {
+  static_assert(N >= 24 && N <= 256 && N % 8 == 0,
+                "setmaxnreg_inc: N must be in [24, 256] and a multiple of 8");
+  asm volatile("setmaxnreg.inc.sync.aligned.u32 %0;\n"
+               :: "n"(N) : "memory");
+}
+
+__device__ __forceinline__
+uint32_t cvt_f32x2_to_bf16x2(float a, float b) {
+  uint32_t r;
+  asm volatile("cvt.rn.bf16x2.f32 %0, %2, %1;\n"
+               : "=r"(r) : "f"(a), "f"(b));
+  return r;
+}
+
+namespace {
+__device__ __forceinline__ uint64_t f32x2_bits(float2 v) {
+  uint64_t b; __builtin_memcpy(&b, &v, 8); return b;
+}
+__device__ __forceinline__ float2 f32x2_make(uint64_t b) {
+  float2 v; __builtin_memcpy(&v, &b, 8); return v;
+}
+}
+
+__device__ __forceinline__ float2 fmul2(float2 a, float2 b) {
+  uint64_t d;
+  asm volatile("mul.f32x2 %0, %1, %2;\n" : "=l"(d) : "l"(f32x2_bits(a)), "l"(f32x2_bits(b)));
+  return f32x2_make(d);
+}
+
+__device__ __forceinline__ float2 fadd2(float2 a, float2 b) {
+  uint64_t d;
+  asm volatile("add.f32x2 %0, %1, %2;\n" : "=l"(d) : "l"(f32x2_bits(a)), "l"(f32x2_bits(b)));
+  return f32x2_make(d);
+}
+
+__device__ __forceinline__ float2 ffma2(float2 a, float2 b, float2 c) {
+  uint64_t d;
+  asm volatile("fma.rn.f32x2 %0, %1, %2, %3;\n"
+               : "=l"(d) : "l"(f32x2_bits(a)), "l"(f32x2_bits(b)), "l"(f32x2_bits(c)));
+  return f32x2_make(d);
+}
+
+__device__ __forceinline__ float2 f32x2_splat(float s) { return make_float2(s, s); }
+
+__device__ __forceinline__ float ex2_approx_f32(float z) {
+  float d;
+  asm volatile("ex2.approx.ftz.f32 %0, %1;\n" : "=f"(d) : "f"(z));
+  return d;
+}
+
+__device__ __forceinline__ float2 ex2_emu_f32x2(float x, float y) {
+  uint32_t ox, oy;
+  asm volatile(
+    "{\n\t"
+    ".reg .f32 f1,f2,f3,f4,f5,f6,f7;\n\t"
+    ".reg .b64 l1,l2,l3,l4,l5,l6,l7,l8,l9,l10;\n\t"
+    ".reg .s32 r1,r2,r3,r4,r5,r6,r7,r8;\n\t"
+    "max.f32 f1, %2, 0fC2FE0000;\n\t"
+    "max.f32 f2, %3, 0fC2FE0000;\n\t"
+    "mov.b64 l1, {f1, f2};\n\t"
+    "mov.f32 f3, 0f4B400000;\n\t"
+    "mov.b64 l2, {f3, f3};\n\t"
+    "add.rm.f32x2 l7, l1, l2;\n\t"
+    "sub.rn.f32x2 l8, l7, l2;\n\t"
+    "sub.rn.f32x2 l9, l1, l8;\n\t"
+    "mov.f32 f7, 0f3D9DF09D;\n\t"
+    "mov.b64 l6, {f7, f7};\n\t"
+    "mov.f32 f6, 0f3E6906A4;\n\t"
+    "mov.b64 l5, {f6, f6};\n\t"
+    "mov.f32 f5, 0f3F31F519;\n\t"
+    "mov.b64 l4, {f5, f5};\n\t"
+    "mov.f32 f4, 0f3F800000;\n\t"
+    "mov.b64 l3, {f4, f4};\n\t"
+    "fma.rn.f32x2 l10, l9, l6, l5;\n\t"
+    "fma.rn.f32x2 l10, l10, l9, l4;\n\t"
+    "fma.rn.f32x2 l10, l10, l9, l3;\n\t"
+    "mov.b64 {r1, r2}, l7;\n\t"
+    "mov.b64 {r3, r4}, l10;\n\t"
+    "shl.b32 r5, r1, 23;\n\t"
+    "add.s32 r7, r5, r3;\n\t"
+    "shl.b32 r6, r2, 23;\n\t"
+    "add.s32 r8, r6, r4;\n\t"
+    "mov.b32 %0, r7;\n\t"
+    "mov.b32 %1, r8;\n\t"
+    "}\n"
+    : "=r"(ox), "=r"(oy) : "f"(x), "f"(y));
+  float2 r; __builtin_memcpy(&r.x, &ox, 4); __builtin_memcpy(&r.y, &oy, 4);
+  return r;
+}
+
+__device__ __forceinline__ float rcp_approx_ftz_f32(float x) {
+  float d;
+  asm("rcp.approx.ftz.f32 %0, %1;\n" : "=f"(d) : "f"(x));
+  return d;
+}
+
+__device__ __forceinline__ uint32_t make_idesc_table44(
+    int M, int N,
+    uint32_t dtype, uint32_t atype, uint32_t btype,
+    bool transpose_a = false, bool transpose_b = false,
+    bool negate_a = false, bool negate_b = false) {
+  uint32_t idesc = 0;
+  idesc |= (dtype & 0x3) << 4;
+  idesc |= (atype & 0x7) << 7;
+  idesc |= (btype & 0x7) << 10;
+  idesc |= (negate_a ? 1u : 0u) << 13;
+  idesc |= (negate_b ? 1u : 0u) << 14;
+  idesc |= (transpose_a ? 1u : 0u) << 15;
+  idesc |= (transpose_b ? 1u : 0u) << 16;
+  idesc |= ((static_cast<uint32_t>(N) >> 3) & 0x3F) << 17;
+  idesc |= ((static_cast<uint32_t>(M) >> 4) & 0x1F) << 24;
+  return idesc;
+}
+
+__device__ __forceinline__ uint32_t make_idesc_bf16_f32(
+    int M, int N, bool ta = false, bool tb = false) {
+  return make_idesc_table44(M, N,  1,
+                              1,  1, ta, tb);
+}
+
+__device__ __forceinline__ void tcgen05_ld_32x32b_x16(
+    uint32_t tmem_addr, uint32_t (&r)[16]) {
+  asm volatile("tcgen05.ld.sync.aligned.32x32b.x16.b32 "
+               "{%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,"
+               "%10,%11,%12,%13,%14,%15}, [%16];\n"
+               : "=r"(r[0]),"=r"(r[1]),"=r"(r[2]),"=r"(r[3]),
+                 "=r"(r[4]),"=r"(r[5]),"=r"(r[6]),"=r"(r[7]),
+                 "=r"(r[8]),"=r"(r[9]),"=r"(r[10]),"=r"(r[11]),
+                 "=r"(r[12]),"=r"(r[13]),"=r"(r[14]),"=r"(r[15])
+               : "r"(tmem_addr));
+}
+
+__device__ __forceinline__ void tcgen05_ld_32x32b_x32(
+    uint32_t tmem_addr, uint32_t (&r)[32]) {
+  asm volatile("tcgen05.ld.sync.aligned.32x32b.x32.b32 "
+               "{%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,"
+               "%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,"
+               "%20,%21,%22,%23,%24,%25,%26,%27,%28,%29,"
+               "%30,%31}, [%32];\n"
+               : "=r"(r[0]),"=r"(r[1]),"=r"(r[2]),"=r"(r[3]),
+                 "=r"(r[4]),"=r"(r[5]),"=r"(r[6]),"=r"(r[7]),
+                 "=r"(r[8]),"=r"(r[9]),"=r"(r[10]),"=r"(r[11]),
+                 "=r"(r[12]),"=r"(r[13]),"=r"(r[14]),"=r"(r[15]),
+                 "=r"(r[16]),"=r"(r[17]),"=r"(r[18]),"=r"(r[19]),
+                 "=r"(r[20]),"=r"(r[21]),"=r"(r[22]),"=r"(r[23]),
+                 "=r"(r[24]),"=r"(r[25]),"=r"(r[26]),"=r"(r[27]),
+                 "=r"(r[28]),"=r"(r[29]),"=r"(r[30]),"=r"(r[31])
+               : "r"(tmem_addr));
+}
+
+__device__ __forceinline__ void tcgen05_ld_32x32b_x64(
+    uint32_t tmem_addr, uint32_t (&r)[64]) {
+  asm volatile("tcgen05.ld.sync.aligned.32x32b.x64.b32 "
+               "{%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,"
+               "%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,"
+               "%20,%21,%22,%23,%24,%25,%26,%27,%28,%29,"
+               "%30,%31,%32,%33,%34,%35,%36,%37,%38,%39,"
+               "%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,"
+               "%50,%51,%52,%53,%54,%55,%56,%57,%58,%59,"
+               "%60,%61,%62,%63}, [%64];\n"
+               : "=r"(r[0]),"=r"(r[1]),"=r"(r[2]),"=r"(r[3]),
+                 "=r"(r[4]),"=r"(r[5]),"=r"(r[6]),"=r"(r[7]),
+                 "=r"(r[8]),"=r"(r[9]),"=r"(r[10]),"=r"(r[11]),
+                 "=r"(r[12]),"=r"(r[13]),"=r"(r[14]),"=r"(r[15]),
+                 "=r"(r[16]),"=r"(r[17]),"=r"(r[18]),"=r"(r[19]),
+                 "=r"(r[20]),"=r"(r[21]),"=r"(r[22]),"=r"(r[23]),
+                 "=r"(r[24]),"=r"(r[25]),"=r"(r[26]),"=r"(r[27]),
+                 "=r"(r[28]),"=r"(r[29]),"=r"(r[30]),"=r"(r[31]),
+                 "=r"(r[32]),"=r"(r[33]),"=r"(r[34]),"=r"(r[35]),
+                 "=r"(r[36]),"=r"(r[37]),"=r"(r[38]),"=r"(r[39]),
+                 "=r"(r[40]),"=r"(r[41]),"=r"(r[42]),"=r"(r[43]),
+                 "=r"(r[44]),"=r"(r[45]),"=r"(r[46]),"=r"(r[47]),
+                 "=r"(r[48]),"=r"(r[49]),"=r"(r[50]),"=r"(r[51]),
+                 "=r"(r[52]),"=r"(r[53]),"=r"(r[54]),"=r"(r[55]),
+                 "=r"(r[56]),"=r"(r[57]),"=r"(r[58]),"=r"(r[59]),
+                 "=r"(r[60]),"=r"(r[61]),"=r"(r[62]),"=r"(r[63])
+               : "r"(tmem_addr));
+}
+
+__device__ __forceinline__ void tcgen05_ld_32x32b_x128(
+    uint32_t tmem_addr, uint32_t (&r)[128]) {
+  asm volatile("tcgen05.ld.sync.aligned.32x32b.x128.b32 "
+               "{%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,"
+               "%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,"
+               "%20,%21,%22,%23,%24,%25,%26,%27,%28,%29,"
+               "%30,%31,%32,%33,%34,%35,%36,%37,%38,%39,"
+               "%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,"
+               "%50,%51,%52,%53,%54,%55,%56,%57,%58,%59,"
+               "%60,%61,%62,%63,%64,%65,%66,%67,%68,%69,"
+               "%70,%71,%72,%73,%74,%75,%76,%77,%78,%79,"
+               "%80,%81,%82,%83,%84,%85,%86,%87,%88,%89,"
+               "%90,%91,%92,%93,%94,%95,%96,%97,%98,%99,"
+               "%100,%101,%102,%103,%104,%105,%106,%107,%108,%109,"
+               "%110,%111,%112,%113,%114,%115,%116,%117,%118,%119,"
+               "%120,%121,%122,%123,%124,%125,%126,%127}, [%128];\n"
+               : "=r"(r[  0]),"=r"(r[  1]),"=r"(r[  2]),"=r"(r[  3]),
+                 "=r"(r[  4]),"=r"(r[  5]),"=r"(r[  6]),"=r"(r[  7]),
+                 "=r"(r[  8]),"=r"(r[  9]),"=r"(r[ 10]),"=r"(r[ 11]),
+                 "=r"(r[ 12]),"=r"(r[ 13]),"=r"(r[ 14]),"=r"(r[ 15]),
+                 "=r"(r[ 16]),"=r"(r[ 17]),"=r"(r[ 18]),"=r"(r[ 19]),
+                 "=r"(r[ 20]),"=r"(r[ 21]),"=r"(r[ 22]),"=r"(r[ 23]),
+                 "=r"(r[ 24]),"=r"(r[ 25]),"=r"(r[ 26]),"=r"(r[ 27]),
+                 "=r"(r[ 28]),"=r"(r[ 29]),"=r"(r[ 30]),"=r"(r[ 31]),
+                 "=r"(r[ 32]),"=r"(r[ 33]),"=r"(r[ 34]),"=r"(r[ 35]),
+                 "=r"(r[ 36]),"=r"(r[ 37]),"=r"(r[ 38]),"=r"(r[ 39]),
+                 "=r"(r[ 40]),"=r"(r[ 41]),"=r"(r[ 42]),"=r"(r[ 43]),
+                 "=r"(r[ 44]),"=r"(r[ 45]),"=r"(r[ 46]),"=r"(r[ 47]),
+                 "=r"(r[ 48]),"=r"(r[ 49]),"=r"(r[ 50]),"=r"(r[ 51]),
+                 "=r"(r[ 52]),"=r"(r[ 53]),"=r"(r[ 54]),"=r"(r[ 55]),
+                 "=r"(r[ 56]),"=r"(r[ 57]),"=r"(r[ 58]),"=r"(r[ 59]),
+                 "=r"(r[ 60]),"=r"(r[ 61]),"=r"(r[ 62]),"=r"(r[ 63]),
+                 "=r"(r[ 64]),"=r"(r[ 65]),"=r"(r[ 66]),"=r"(r[ 67]),
+                 "=r"(r[ 68]),"=r"(r[ 69]),"=r"(r[ 70]),"=r"(r[ 71]),
+                 "=r"(r[ 72]),"=r"(r[ 73]),"=r"(r[ 74]),"=r"(r[ 75]),
+                 "=r"(r[ 76]),"=r"(r[ 77]),"=r"(r[ 78]),"=r"(r[ 79]),
+                 "=r"(r[ 80]),"=r"(r[ 81]),"=r"(r[ 82]),"=r"(r[ 83]),
+                 "=r"(r[ 84]),"=r"(r[ 85]),"=r"(r[ 86]),"=r"(r[ 87]),
+                 "=r"(r[ 88]),"=r"(r[ 89]),"=r"(r[ 90]),"=r"(r[ 91]),
+                 "=r"(r[ 92]),"=r"(r[ 93]),"=r"(r[ 94]),"=r"(r[ 95]),
+                 "=r"(r[ 96]),"=r"(r[ 97]),"=r"(r[ 98]),"=r"(r[ 99]),
+                 "=r"(r[100]),"=r"(r[101]),"=r"(r[102]),"=r"(r[103]),
+                 "=r"(r[104]),"=r"(r[105]),"=r"(r[106]),"=r"(r[107]),
+                 "=r"(r[108]),"=r"(r[109]),"=r"(r[110]),"=r"(r[111]),
+                 "=r"(r[112]),"=r"(r[113]),"=r"(r[114]),"=r"(r[115]),
+                 "=r"(r[116]),"=r"(r[117]),"=r"(r[118]),"=r"(r[119]),
+                 "=r"(r[120]),"=r"(r[121]),"=r"(r[122]),"=r"(r[123]),
+                 "=r"(r[124]),"=r"(r[125]),"=r"(r[126]),"=r"(r[127])
+               : "r"(tmem_addr));
+}
+
+__device__ __forceinline__
+uint32_t smem_ptr_u32(const void* ptr) {
+  return static_cast<uint32_t>(__cvta_generic_to_shared(ptr));
+}
+
+__device__ __forceinline__
+void sts_f32(uint32_t smem_addr, float val) {
+  asm volatile("st.shared.f32 [%0], %1;" :: "r"(smem_addr), "f"(val) : "memory");
+}
+

fastvideo-kernel/csrc/comm/ulysses_all_to_all.cu

@@ -0,0 +1,279 @@
+/*
+ * Copyright (c) 2025 by FlashInfer team.
+ *
+ * Adapted from flashinfer-ai/flashinfer @ 8a94642d83cba0939035868fb6c309b4474a13d6
+ * (PR #3820), csrc/ulysses_all_to_all.cu.
+ *
+ * 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.
+ */
+
+// Torch host bindings for the fused Ulysses all-to-all. Kernel in
+// include/comm/ulysses_all_to_all.cuh.
+//
+// The per-group context is an ncclDevComm plus a registered symmetric window,
+// both created here from the caller's ncclComm_t.
+
+#include <c10/cuda/CUDAGuard.h>
+#include <c10/cuda/CUDAStream.h>
+#include <torch/extension.h>
+
+#include <algorithm>
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+#include <nccl.h>
+#include <nccl_device.h>
+
+#include "comm/ulysses_all_to_all.cuh"
+
+namespace fi = fastvideo::comm::ulysses;
+
+namespace {
+
+// A symmetric window every rank can store into, plus the ncclDevComm the
+// kernel opens barrier sessions on.
+struct UlyssesContext {
+  ncclComm_t comm = nullptr;
+  ncclWindow_t win = nullptr;
+  void* buf = nullptr;
+  size_t nbytes = 0;
+  ncclDevComm devComm{};
+  bool dev_comm_created = false;
+  int device = -1;
+  int rank = 0;
+  int world = 0;
+};
+
+#define NCCL_TRY(expr, what)                                                   \
+  do {                                                                         \
+    ncclResult_t _r = (expr);                                                  \
+    TORCH_CHECK(_r == ncclSuccess, what " failed: rc=", static_cast<int>(_r)); \
+  } while (0)
+
+}  // namespace
+
+// Allocate the local half of a context. This is intentionally separate from
+// registration so Python can vote after local allocation: if one rank is OOM,
+// no peer enters a collective window registration alone.
+int64_t allocate_ulysses_a2a(int64_t nbytes, int64_t rank, int64_t world_size,
+                             int64_t device_index) {
+  TORCH_CHECK(world_size == 2 || world_size == 4 || world_size == 6 || world_size == 8,
+              "ulysses a2a only supports world size in (2, 4, 6, 8), got ", world_size);
+  TORCH_CHECK(rank >= 0 && rank < world_size, "invalid rank");
+  TORCH_CHECK(nbytes > 0, "nbytes must be positive");
+  TORCH_CHECK(device_index >= 0, "device index must be non-negative");
+
+  const at::cuda::CUDAGuard device_guard(
+      c10::Device(c10::DeviceType::CUDA, static_cast<c10::DeviceIndex>(device_index)));
+  auto ctx = std::make_unique<UlyssesContext>();
+  ctx->nbytes = static_cast<size_t>(nbytes);
+  ctx->device = static_cast<int>(device_index);
+  ctx->rank = static_cast<int>(rank);
+  ctx->world = static_cast<int>(world_size);
+
+  // The window must come from NCCL's allocator (4096B aligned per
+  // NCCL_WIN_REQUIRED_ALIGNMENT), which is why this is not a torch tensor.
+  NCCL_TRY(ncclMemAlloc(&ctx->buf, ctx->nbytes), "ncclMemAlloc");
+
+  return reinterpret_cast<int64_t>(ctx.release());
+}
+
+// Register the user window. Collective: every rank must call together.
+void register_ulysses_a2a_window(int64_t handle, int64_t comm_ptr) {
+  auto* ctx = reinterpret_cast<UlyssesContext*>(handle);
+  TORCH_CHECK(ctx != nullptr, "handle must come from allocate_ulysses_a2a");
+  TORCH_CHECK(ctx->buf != nullptr && ctx->win == nullptr && !ctx->dev_comm_created,
+              "ulysses a2a context is not in the allocated state");
+
+  const at::cuda::CUDAGuard device_guard(
+      c10::Device(c10::DeviceType::CUDA, static_cast<c10::DeviceIndex>(ctx->device)));
+  ctx->comm = reinterpret_cast<ncclComm_t>(comm_ptr);
+  NCCL_TRY(ncclCommWindowRegister(ctx->comm, ctx->buf, ctx->nbytes, &ctx->win,
+                                  NCCL_WIN_COLL_SYMMETRIC),
+           "ncclCommWindowRegister");
+}
+
+// Create the device communicator only after Python has voted that every rank
+// registered its window. This operation is collective as well.
+void create_ulysses_a2a_dev_comm(int64_t handle) {
+  auto* ctx = reinterpret_cast<UlyssesContext*>(handle);
+  TORCH_CHECK(ctx != nullptr, "handle must come from allocate_ulysses_a2a");
+  TORCH_CHECK(ctx->comm != nullptr && ctx->win != nullptr && !ctx->dev_comm_created,
+              "ulysses a2a context is not in the window-registered state");
+
+  const at::cuda::CUDAGuard device_guard(
+      c10::Device(c10::DeviceType::CUDA, static_cast<c10::DeviceIndex>(ctx->device)));
+
+  ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
+  reqs.lsaBarrierCount = fi::kMaxBlocks;
+  NCCL_TRY(ncclDevCommCreate(ctx->comm, &reqs, &ctx->devComm), "ncclDevCommCreate");
+  ctx->dev_comm_created = true;
+}
+
+static ncclResult_t first_error(ncclResult_t current, ncclResult_t next) {
+  return current == ncclSuccess ? next : current;
+}
+
+void dispose_ulysses_a2a(int64_t handle) {
+  auto* ctx = reinterpret_cast<UlyssesContext*>(handle);
+  if (ctx == nullptr) return;
+
+  const at::cuda::CUDAGuard device_guard(
+      c10::Device(c10::DeviceType::CUDA, static_cast<c10::DeviceIndex>(ctx->device)));
+  ncclResult_t result = ncclSuccess;
+  if (ctx->comm != nullptr && ctx->dev_comm_created) {
+    result = first_error(result, ncclDevCommDestroy(ctx->comm, &ctx->devComm));
+    ctx->dev_comm_created = false;
+  }
+  if (ctx->comm != nullptr && ctx->win != nullptr) {
+    result = first_error(result, ncclCommWindowDeregister(ctx->comm, ctx->win));
+    ctx->win = nullptr;
+  }
+  if (ctx->buf != nullptr) {
+    result = first_error(result, ncclMemFree(ctx->buf));
+    ctx->buf = nullptr;
+  }
+  delete ctx;
+  TORCH_CHECK(result == ncclSuccess, "ulysses a2a cleanup failed: rc=", static_cast<int>(result));
+}
+
+// Whether the whole group is load-store accessible. NCCL determined this at
+// ncclCommInitRank.
+bool ulysses_lsa_covers_group(int64_t comm_ptr, int64_t world_size) {
+  auto comm = reinterpret_cast<ncclComm_t>(comm_ptr);
+  ncclCommProperties properties = NCCL_COMM_PROPERTIES_INITIALIZER;
+  NCCL_TRY(ncclCommQueryProperties(comm, &properties), "ncclCommQueryProperties");
+  ncclTeam_t lsa = ncclTeamLsa(comm);
+  return properties.deviceApiSupport && lsa.nRanks == static_cast<int>(world_size);
+}
+
+// Fused-transpose Ulysses all-to-all.
+//   mode == 0: inp [B, S_local, H, D]        -> out [B, S_global, H_local, D]
+//   mode == 1: inp [B, S_global, H_local, D] -> out [B, S_local, H, D]
+// where H is the *global* head count and H_local = H / world_size.
+void ulysses_a2a(int64_t handle, torch::Tensor inp, torch::Tensor out, int64_t B, int64_t S_local,
+                 int64_t H, int64_t D, int64_t mode) {
+  auto* ctx = reinterpret_cast<UlyssesContext*>(handle);
+  TORCH_CHECK(ctx != nullptr, "handle must come from allocate_ulysses_a2a");
+
+  const at::cuda::CUDAGuard device_guard(inp.device());
+  auto stream = at::cuda::getCurrentCUDAStream();
+
+  TORCH_CHECK(inp.is_cuda() && out.is_cuda(), "inp and out must be CUDA tensors");
+  TORCH_CHECK(inp.is_contiguous() && out.is_contiguous(), "inp and out must be contiguous");
+  TORCH_CHECK(inp.device() == out.device(), "inp and out must be on the same device");
+  TORCH_CHECK(inp.scalar_type() == out.scalar_type(), "inp and out must share a dtype");
+  TORCH_CHECK(inp.numel() == out.numel(), "inp and out must have equal element counts");
+  TORCH_CHECK(inp.get_device() == ctx->device, "input is on CUDA device ", inp.get_device(),
+              " but the Ulysses context belongs to device ", ctx->device);
+  TORCH_CHECK(mode == 0 || mode == 1, "mode must be 0 or 1");
+  TORCH_CHECK(inp.dim() == 4 && out.dim() == 4, "inp and out must be 4-D");
+
+  const int W = ctx->world;
+  TORCH_CHECK(H % W == 0, "global head count must be divisible by world size");
+  const int H_local = static_cast<int>(H / W);
+
+  const torch::Tensor& local_op = (mode == 0) ? inp : out;   // [B, S_local, H, D]
+  const torch::Tensor& global_op = (mode == 0) ? out : inp;  // [B, S_global, H_local, D]
+  TORCH_CHECK(local_op.size(0) == B && local_op.size(1) == S_local && local_op.size(2) == H &&
+                  local_op.size(3) == D,
+              "the [B, S_local, H, D] operand of mode ", mode, " has shape (", local_op.size(0),
+              ", ", local_op.size(1), ", ", local_op.size(2), ", ", local_op.size(3),
+              "), expected (", B, ", ", S_local, ", ", H, ", ", D, ")");
+  TORCH_CHECK(global_op.size(0) == B && global_op.size(1) == W * S_local &&
+                  global_op.size(2) == H_local && global_op.size(3) == D,
+              "the [B, S_global, H_local, D] operand of mode ", mode, " has shape (",
+              global_op.size(0), ", ", global_op.size(1), ", ", global_op.size(2), ", ",
+              global_op.size(3), "), expected (", B, ", ", W * S_local, ", ", H_local, ", ", D,
+              ")");
+
+  const size_t out_bytes = out.numel() * out.element_size();
+  TORCH_CHECK(out_bytes <= ctx->nbytes, "operand of ", out_bytes,
+              " bytes exceeds the window capacity ", ctx->nbytes);
+
+  const int64_t num_rows = B * static_cast<int64_t>(W) * S_local;
+  const int blocks =
+      static_cast<int>(std::max<int64_t>(1, std::min<int64_t>(fi::kMaxBlocks, num_rows)));
+  const int threads = fi::kUlyssesThreads;
+
+#define LAUNCH_ULYSSES_A2A(T, NG, MODE)                                             \
+  fi::ulysses_a2a_kernel<T, NG, MODE><<<blocks, threads, 0, stream>>>(              \
+      reinterpret_cast<const T*>(inp.data_ptr()), ctx->devComm, ctx->win, /*off=*/0, \
+      ctx->rank, static_cast<int>(B), static_cast<int>(S_local), H_local, static_cast<int>(D))
+
+#define DISPATCH_NGPUS(T, MODE)                                                \
+  switch (W) {                                                                 \
+    case 2:                                                                    \
+      LAUNCH_ULYSSES_A2A(T, 2, MODE);                                          \
+      break;                                                                   \
+    case 4:                                                                    \
+      LAUNCH_ULYSSES_A2A(T, 4, MODE);                                          \
+      break;                                                                   \
+    case 6:                                                                    \
+      LAUNCH_ULYSSES_A2A(T, 6, MODE);                                          \
+      break;                                                                   \
+    case 8:                                                                    \
+      LAUNCH_ULYSSES_A2A(T, 8, MODE);                                          \
+      break;                                                                   \
+    default:                                                                   \
+      TORCH_CHECK(false, "ulysses_a2a only supports world size in (2,4,6,8)"); \
+  }
+
+#define DISPATCH_DTYPE(MODE)                                                              \
+  switch (out.scalar_type()) {                                                            \
+    case at::ScalarType::Float: {                                                         \
+      DISPATCH_NGPUS(float, MODE);                                                        \
+      break;                                                                              \
+    }                                                                                     \
+    case at::ScalarType::Half: {                                                          \
+      DISPATCH_NGPUS(half, MODE);                                                         \
+      break;                                                                              \
+    }                                                                                     \
+    case at::ScalarType::BFloat16: {                                                      \
+      DISPATCH_NGPUS(nv_bfloat16, MODE);                                                  \
+      break;                                                                              \
+    }                                                                                     \
+    default:                                                                              \
+      TORCH_CHECK(false, "ulysses_a2a only supports float32, float16 and bfloat16, got ", \
+                  out.scalar_type());                                                     \
+  }
+
+  if (mode == 0) {
+    DISPATCH_DTYPE(0);
+  } else {
+    DISPATCH_DTYPE(1);
+  }
+
+#undef DISPATCH_DTYPE
+#undef DISPATCH_NGPUS
+#undef LAUNCH_ULYSSES_A2A
+
+  TORCH_CHECK(cudaGetLastError() == cudaSuccess, "ulysses_a2a kernel launch failed");
+  // Copy this rank's completed result out of the window.
+  auto status = cudaMemcpyAsync(out.data_ptr(), ctx->buf, out_bytes, cudaMemcpyDeviceToDevice,
+                                stream);
+  TORCH_CHECK(status == cudaSuccess, "ulysses_a2a copy-out failed: ", cudaGetErrorString(status));
+}
+
+void register_ulysses_a2a(pybind11::module_& m) {
+  m.def("allocate_ulysses_a2a", &allocate_ulysses_a2a, "allocate a local ulysses a2a window");
+  m.def("register_ulysses_a2a_window", &register_ulysses_a2a_window,
+        "register the ulysses a2a window collectively");
+  m.def("create_ulysses_a2a_dev_comm", &create_ulysses_a2a_dev_comm,
+        "create the ulysses a2a device communicator collectively");
+  m.def("dispose_ulysses_a2a", &dispose_ulysses_a2a, "release a ulysses a2a context");
+  m.def("ulysses_lsa_covers_group", &ulysses_lsa_covers_group,
+        "whether the whole group is load-store accessible");
+  m.def("ulysses_a2a", &ulysses_a2a, "fused-transpose Ulysses all-to-all over NVLink");
+}

fastvideo-kernel/csrc/common_extension.cpp

@@ -22,16 +22,42 @@ extern std::vector<torch::Tensor> block_sparse_attention_backward(
 );
 #endif
 
+#ifdef FASTVIDEO_KERNEL_COMPILE_ULYSSES_A2A
+// Ulysses sequence-parallel all-to-all (csrc/comm/)
+void register_ulysses_a2a(pybind11::module_ &);
+#endif
+
 // TurboDiffusion kernels
 void register_quant(pybind11::module_ &);
 void register_rms_norm(pybind11::module_ &);
 void register_layer_norm(pybind11::module_ &);
 void register_gemm(pybind11::module_ &);
 
+#ifdef TK_COMPILE_BLOCK_SPARSE_VSA_SM100A
+extern std::vector<torch::Tensor> block_sparse_sm100a_fwd(
+    torch::Tensor q, torch::Tensor k, torch::Tensor v, c10::optional<torch::Tensor> v_t,
+    torch::Tensor q2k_idx, torch::Tensor q2k_num, torch::Tensor variable_block_sizes,
+    double sm_scale, bool need_lse);
+extern std::vector<torch::Tensor> block_sparse_sm100a_blk128_fwd(
+    torch::Tensor q, torch::Tensor k, torch::Tensor v, c10::optional<torch::Tensor> v_t,
+    torch::Tensor q2k_idx, torch::Tensor q2k_num, torch::Tensor variable_block_sizes,
+    double sm_scale, bool need_lse);
+#endif
+
 PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
     m.doc() = "FastVideo CUDA Kernels";
 
+#ifdef TK_COMPILE_BLOCK_SPARSE_VSA_SM100A
+    m.def("block_sparse_sm100a_fwd",
+          torch::wrap_pybind_function(block_sparse_sm100a_fwd),
+          "VSA block-sparse attention forward, 64-token blocks (Blackwell sm100a)");
+    m.def("block_sparse_sm100a_blk128_fwd",
+          torch::wrap_pybind_function(block_sparse_sm100a_blk128_fwd),
+          "VSA block-sparse attention forward, 128-token blocks (Blackwell sm100a)");
+#endif
+
 #ifdef TK_COMPILE_ST_ATTN
+
     m.def("sta_fwd", torch::wrap_pybind_function(sta_forward), "sliding tile attention (Hopper)");
 #endif
 
@@ -40,6 +66,11 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
     m.def("block_sparse_bwd", torch::wrap_pybind_function(block_sparse_attention_backward), "block sparse attention backward (Hopper)");
 #endif
 
+#ifdef FASTVIDEO_KERNEL_COMPILE_ULYSSES_A2A
+    // Ulysses sequence-parallel all-to-all
+    register_ulysses_a2a(m);
+#endif
+
     // TurboDiffusion
     register_quant(m);
     register_rms_norm(m);

fastvideo-kernel/include/comm/ulysses_all_to_all.cuh

@@ -0,0 +1,165 @@
+/*
+ * Copyright (c) 2025 by FlashInfer team.
+ *
+ * Adapted from flashinfer-ai/flashinfer @ 8a94642d83cba0939035868fb6c309b4474a13d6
+ * (PR #3820), which in turn adapted ThunderKittens' NVLink all-to-all:
+ * https://github.com/HazyResearch/ThunderKittens/blob/main/kernels/parallel/all_to_all/all_to_all.cu
+ *
+ * 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.
+ */
+
+// Fused-transpose Ulysses all-to-all over NVLink. Peer addresses come from
+// ncclGetLsaPointer and synchronization from ncclLsaBarrierSession; the index
+// math and slab decomposition below are upstream's.
+//
+// head_dim == 2 layout, uniform sequence splits. With
+//   W        = ulysses world size
+//   H_local  = H / W
+//   S_global = S_local * W
+//
+//   mode == 0 (input  a2a): [B, S_local, H,       D] -> [B, S_global, H_local, D]
+//       y_r[b, j*S_local + s, hl, d] = x_j[b, s, r*H_local + hl, d]
+//   mode == 1 (output a2a): [B, S_global, H_local, D] -> [B, S_local, H,       D]
+//       out_j[b, s, r*H_local + hl, d] = u_r[b, j*S_local + s, hl, d]
+//
+// In both modes the unit of transfer is a contiguous (H_local * D) block, so
+// every cross-GPU store is fully coalesced.
+
+#ifndef FASTVIDEO_COMM_ULYSSES_ALL_TO_ALL_CUH_
+#define FASTVIDEO_COMM_ULYSSES_ALL_TO_ALL_CUH_
+
+#include <cstdint>
+
+#include <nccl.h>
+#include <nccl_device.h>
+
+namespace fastvideo {
+namespace comm {
+namespace ulysses {
+
+constexpr int kUlyssesThreads = 512;
+// Deliberately modest: this is link-bandwidth bound, so a small grid leaves the
+// rest of the GPU free without costing throughput.
+constexpr int kMaxBlocks = 36;
+
+// Shared movement body for the fused-transpose all-to-all (no barriers).
+//
+// Rows are ordered ((b * W + peer) * S_local + s), so consecutive rows share a
+// (batch, peer) and are contiguous on the gather side of the transpose. Each
+// block takes a contiguous slab of rows and flattens its threads over the 16B
+// units in it, so consecutive lanes address one peer buffer back to back and the
+// remote writes coalesce into large bursts rather than (H_local * D)-sized
+// scattered ones.
+template <typename T, int NGPUS, int MODE>
+__device__ __forceinline__ void ulysses_a2a_move(const T* __restrict__ local_in,
+                                                 void* const* peer_ptrs, int rank, int B,
+                                                 int S_local, int H_local, int D) {
+  static_assert(MODE == 0 || MODE == 1, "MODE must be 0 or 1");
+  const int W = NGPUS;
+  const int64_t H = static_cast<int64_t>(H_local) * W;
+  const int64_t S_global = static_cast<int64_t>(S_local) * W;
+  const int64_t block_len = static_cast<int64_t>(H_local) * D;  // elements/row
+  const int64_t num_rows = static_cast<int64_t>(B) * W * S_local;
+
+  // 16B-vectorized fast path when every row is 16B aligned (the common case:
+  // contiguous bf16/fp16/fp32 tensors with block_len * sizeof(T) % 16 == 0).
+  using Vec = int4;
+  constexpr int kVecBytes = sizeof(Vec);
+  const int64_t row_bytes = block_len * static_cast<int64_t>(sizeof(T));
+  const bool vec_ok =
+      (row_bytes % kVecBytes) == 0 && (reinterpret_cast<uintptr_t>(local_in) % kVecBytes) == 0;
+
+  // Contiguous slab of rows for this block.
+  const int64_t rows_per_block = (num_rows + gridDim.x - 1) / gridDim.x;
+  const int64_t row_lo = static_cast<int64_t>(blockIdx.x) * rows_per_block;
+  int64_t row_hi = row_lo + rows_per_block;
+  if (row_hi > num_rows) row_hi = num_rows;
+  if (row_lo >= row_hi) return;
+
+  const int tid = threadIdx.x;
+  const int nthr = blockDim.x;
+
+  // Decode (b, peer, s) and compute src/dst element offsets for a given row.
+  auto offsets = [&](int64_t row, int64_t& src_off, int64_t& dst_off) {
+    const int64_t s = row % S_local;
+    const int64_t tmp = row / S_local;
+    const int64_t peer = tmp % W;
+    const int64_t b = tmp / W;
+    if constexpr (MODE == 0) {
+      src_off = ((b * S_local + s) * H + peer * H_local) * D;
+      dst_off = (b * S_global + static_cast<int64_t>(rank) * S_local + s) * block_len;
+    } else {
+      src_off = (b * S_global + peer * S_local + s) * block_len;
+      dst_off = ((b * S_local + s) * H + static_cast<int64_t>(rank) * H_local) * D;
+    }
+    return peer;
+  };
+
+  if (vec_ok) {
+    const int64_t units_per_row = row_bytes / kVecBytes;
+    const int64_t total_units = (row_hi - row_lo) * units_per_row;
+    for (int64_t u = tid; u < total_units; u += nthr) {
+      const int64_t local_row = u / units_per_row;
+      const int64_t unit = u - local_row * units_per_row;
+      const int64_t row = row_lo + local_row;
+      int64_t src_off, dst_off;
+      const int64_t peer = offsets(row, src_off, dst_off);
+      const Vec* s4 = reinterpret_cast<const Vec*>(local_in + src_off);
+      Vec* d4 = reinterpret_cast<Vec*>((T*)peer_ptrs[peer] + dst_off);
+      d4[unit] = s4[unit];
+    }
+  } else {
+    // Scalar fallback (unaligned / odd shapes).
+    for (int64_t row = row_lo; row < row_hi; ++row) {
+      int64_t src_off, dst_off;
+      const int64_t peer = offsets(row, src_off, dst_off);
+      const T* s_ptr = local_in + src_off;
+      T* d_ptr = (T*)peer_ptrs[peer] + dst_off;
+      for (int64_t i = tid; i < block_len; i += nthr) {
+        d_ptr[i] = s_ptr[i];
+      }
+    }
+  }
+}
+
+// The transfer mode is a compile-time template parameter so the address math
+// specializes and the coalesced slab decomposition is fully unrolled per mode.
+template <typename T, int NGPUS, int MODE>
+__global__ void __launch_bounds__(kUlyssesThreads, 1)
+    ulysses_a2a_kernel(const T* __restrict__ local_in, ncclDevComm devComm, ncclWindow_t win,
+                       size_t win_offset, int rank, int B, int S_local, int H_local, int D) {
+  // Resolved once; the movement loop would otherwise call this per 16B store.
+  void* peer_ptrs[NGPUS];
+#pragma unroll
+  for (int p = 0; p < NGPUS; ++p) {
+    peer_ptrs[p] = ncclGetLsaPointer(win, win_offset, p);
+  }
+
+  // Each CTA owns one barrier generation. Sharing index 0 across independently
+  // scheduled CTAs races the generation counter and is unsupported by NCCL's
+  // device API. The host reserves kMaxBlocks slots when it creates devComm.
+  ncclLsaBarrierSession<ncclCoopCta> bar(
+      ncclCoopCta(), devComm, ncclTeamTagLsa(), /*index=*/blockIdx.x);
+
+  // Every rank must have entered before anyone writes into peer buffers.
+  bar.sync(ncclCoopCta(), cuda::memory_order_relaxed);
+  ulysses_a2a_move<T, NGPUS, MODE>(local_in, peer_ptrs, rank, B, S_local, H_local, D);
+  // Release-acquire: all peer writes visible before a rank reads its window.
+  bar.sync(ncclCoopCta(), cuda::memory_order_acq_rel);
+}
+
+}  // namespace ulysses
+}  // namespace comm
+}  // namespace fastvideo
+
+#endif  // FASTVIDEO_COMM_ULYSSES_ALL_TO_ALL_CUH_

fastvideo-kernel/pyproject.toml

@@ -1,7 +1,7 @@
 [build-system]
 requires = [
     "scikit-build-core>=0.10",
-    "torch>=2.5.0",
+    "torch==2.12.0",
     "setuptools>=61.0.0",
     "wheel"
 ]
@@ -9,7 +9,7 @@ build-backend = "scikit_build_core.build"
 
 [project]
 name = "fastvideo-kernel"
-version = "0.3.2"
+version = "0.3.5"
 description = "Unified CUDA kernels for FastVideo"
 readme = "README.md"
 requires-python = ">=3.10"
@@ -22,7 +22,7 @@ classifiers = [
     "Environment :: GPU :: NVIDIA CUDA",
 ]
 dependencies = [
-    "torch>=2.5.0",
+    "torch==2.12.0",
     "triton>=2.0.0; sys_platform == 'linux'",
 ]
 

fastvideo-kernel/python/fastvideo_kernel/block_sparse_attn.py

@@ -49,6 +49,36 @@ def _force_tk() -> bool:
     return os.environ.get("FASTVIDEO_VSA_TK", "0") == "1"
 
 
+def _force_sm100a() -> bool:
+    """True iff the sm_100a (Blackwell) forward is explicitly opted into.
+
+    Opt-in only (same env the H3 backend honors): the sm_100a extension is
+    forward-only, so this routing pairs it with the Triton backward -- its lse
+    is already in Triton's M format. Honored only when
+    ``block_sparse_attn_sm100a.is_supported`` passes. Unsupported 64-token
+    metadata falls through to the default selection; unsupported 128-token
+    metadata raises because Triton has no compatible fallback.
+    ``FASTVIDEO_VSA_TRITON`` still wins.
+    """
+    return os.environ.get("FASTVIDEO_VSA_SM100A", "0") == "1"
+
+
+def _sm100a_is_supported(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> bool:
+    try:
+        from fastvideo_kernel import block_sparse_attn_sm100a as vsa_sm100a
+    except Exception:
+        return False
+    return vsa_sm100a.is_supported(q, variable_block_sizes)
+
+
+def _infer_block_size(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> int:
+    num_blocks = variable_block_sizes.numel()
+    seq_len = q.shape[2]
+    if num_blocks == 0 or seq_len % num_blocks != 0:
+        return 0
+    return seq_len // num_blocks
+
+
 # ---------------------------------------------------------------------------
 # Index helpers
 # ---------------------------------------------------------------------------
@@ -339,6 +369,76 @@ def _backward_sm90(ctx, grad_o, grad_lse):
 
 block_sparse_attn_sm90.register_autograd(_backward_sm90, setup_context=_setup_context_sm90)
 
+# ---------------------------------------------------------------------------
+# SM100A backend custom op (index-native)
+#
+# Forward runs the sm_100a CUDA extension; backward reuses the Triton kernels.
+# The sm_100a forward emits lse in exactly Triton's M format (max*log2e +
+# log2(l)), so the pairing needs no conversion. The Triton backward is
+# hardcoded to 64-token blocks, hence the block-size assert below.
+# ---------------------------------------------------------------------------
+
+
+@torch.library.custom_op(
+    "fastvideo_kernel::block_sparse_attn_sm100a",
+    mutates_args=(),
+    device_types="cuda",
+)
+def block_sparse_attn_sm100a_op(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q2k_idx: torch.Tensor,
+    q2k_num: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    from fastvideo_kernel.block_sparse_attn_sm100a import block_sparse_attn_sm100a
+
+    o, M = block_sparse_attn_sm100a(q, k, v, q2k_idx, q2k_num, variable_block_sizes,
+                                    need_lse=True)
+    return o, M
+
+
+@torch.library.register_fake("fastvideo_kernel::block_sparse_attn_sm100a")
+def _block_sparse_attn_sm100a_fake(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q2k_idx: torch.Tensor,
+    q2k_num: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    o = torch.empty_like(q)
+    M = torch.empty(
+        (q.shape[0], q.shape[1], q.shape[2]),
+        device=q.device,
+        dtype=torch.float32,
+    )
+    return o, M
+
+
+def _setup_context_sm100a(ctx, inputs, output):
+    q, k, v, q2k_idx, q2k_num, variable_block_sizes = inputs
+    o, M = output
+    ctx.save_for_backward(q, k, v, o, M, q2k_idx, q2k_num, variable_block_sizes)
+
+
+def _backward_sm100a(ctx, grad_o, grad_M):
+    q, k, v, o, M, q2k_idx, q2k_num, variable_block_sizes = ctx.saved_tensors
+    block = q.shape[2] // variable_block_sizes.numel()
+    if block != 64:
+        raise RuntimeError(
+            "block_sparse_attn_sm100a backward pairs the sm_100a forward with the "
+            f"Triton backward, which is hardcoded to 64-token blocks; got {block}. "
+            "Run 128-token-block metadata without grad, or use the Triton forward.")
+    dq, dk, dv = block_sparse_attn_backward_triton(grad_o, q, k, v, o, M, q2k_idx,
+                                                   q2k_num, variable_block_sizes)
+    return dq, dk, dv, None, None, None
+
+
+block_sparse_attn_sm100a_op.register_autograd(_backward_sm100a,
+                                              setup_context=_setup_context_sm100a)
+
 # ---------------------------------------------------------------------------
 # Public API
 # ---------------------------------------------------------------------------
@@ -364,11 +464,24 @@ def block_sparse_attn_from_indices(
 
     # Backend resolution:
     # - FASTVIDEO_VSA_TRITON forces Triton everywhere.
+    # - FASTVIDEO_VSA_SM100A opts into the sm_100a forward (Triton backward).
+    #   Unsupported 64-token metadata falls through; unsupported 128-token
+    #   metadata raises because Triton cannot consume it.
     # - FASTVIDEO_VSA_TK requests sm_90 TK; honored only when it's actually
     #   available (else falls through to the default below).
     # - Otherwise: TK on sm_90 if available, else Triton.
     if _force_triton():
         use_sm90 = False
+    elif _force_sm100a():
+        if _sm100a_is_supported(q, variable_block_sizes):
+            return block_sparse_attn_sm100a_op(q, k, v, q2k_idx, q2k_num,
+                                               variable_block_sizes)
+        if _infer_block_size(q, variable_block_sizes) == 128:
+            raise NotImplementedError(
+                "128-token block-sparse attention requires the sm_100a forward; "
+                "the Triton fallback only supports 64-token blocks, and the "
+                "sm_100a route is unavailable for this input.")
+        use_sm90 = sm90_available
     elif _force_tk():
         use_sm90 = sm90_available
     else:

fastvideo-kernel/python/fastvideo_kernel/block_sparse_attn_sm100a.py

@@ -0,0 +1,216 @@
+# SPDX-License-Identifier: Apache-2.0
+"""sm_100a (Blackwell) CUDA block-sparse VSA forward.
+
+A third backend behind the same VSA op as the Triton and CuTe-DSL paths. Forward only: it
+returns ``(out, lse)`` with ``lse`` in exactly the form ``triton_block_sparse_attn_forward``
+writes -- ``max(qk * qk_scale) + log2(l)``, ``[B, H, S]`` fp32 -- so
+``block_sparse_attn_backward_triton`` runs against it unchanged.
+
+The extension carries TWO instantiations of the kernel, for 64- and 128-token sparse blocks
+(tile volumes 64 and 128 in ``build_vsa_metadata``); the block size is inferred from the
+tensors and picks the op. Anything else falls back to Triton via ``is_supported``.
+"""
+
+from typing import Tuple
+
+import torch
+
+try:
+    # The pybind symbols live on fastvideo_kernel_ops, NOT on the _C package that contains it.
+    # `import fastvideo_kernel._C as _C` resolves to the namespace package, whose __init__ is
+    # empty, so hasattr() fails on a wheel install and the caller silently falls back with the
+    # kernel built and present.
+    from fastvideo_kernel._C import fastvideo_kernel_ops as _C
+    _FWD_BY_BLOCK = {
+        64: getattr(_C, "block_sparse_sm100a_fwd", None),
+        128: getattr(_C, "block_sparse_sm100a_blk128_fwd", None),
+    }
+    _HAS_VSA_SM100A = any(_FWD_BY_BLOCK.values())
+except ImportError:  # pragma: no cover - extension not built
+    _C = None
+    _FWD_BY_BLOCK = {}
+    _HAS_VSA_SM100A = False
+
+_SM100 = (10, 0)
+HEAD_DIM = 128
+# Must match the -DVSA_BHSD the extension was compiled with (see CMakeLists).
+BHSD = True
+
+
+def _block_size(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> int:
+    num_blocks = variable_block_sizes.numel()
+    seqlen = q.shape[2] if BHSD else q.shape[1]
+    return 0 if num_blocks == 0 or seqlen % num_blocks else seqlen // num_blocks
+
+
+def is_supported(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> bool:
+    """True iff this build can run these tensors; otherwise the caller uses Triton.
+
+    Static facts only -- shapes, dtypes, arch, layout. Deliberately NO reads of tensor
+    contents: the previous ``int(variable_block_sizes.min())`` was a GPU->CPU sync on every
+    call, and the kernel no longer needs it (see below). This predicate must stay cheap
+    enough to sit on a per-layer dispatch path.
+
+    What the kernel accepts (and is tested to handle):
+      * q/k/v: contiguous 4-D bf16 CUDA tensors on an sm_100 device, head_dim 128, laid out
+        as compiled (BHSD here); seqlen == num_blocks * block with an EVEN num_blocks (a CTA
+        owns an adjacent pair of query blocks) and a 64- or 128-token build present.
+      * q2k_num: any per-row counts in [0, max_kv], NON-uniform across rows included. Rows
+        with count 0 produce exactly-zero output rows (and a finite LSE sentinel) rather
+        than attending anywhere -- so no ``.min()`` floor is required of the caller.
+      * q2k_idx: rows only need valid entries (in [0, num_blocks)) BELOW that row's count;
+        padding past the count (e.g. map_to_index's -1 fill) is never dereferenced. max_kv
+        (= q2k_idx.shape[-1]) must be >= 1, which the host launcher re-checks.
+      * variable_block_sizes: per-KV-block valid-token counts in [0, block]; keys at or past
+        a block's count are masked. Integer metadata is converted to int32/contiguous by
+        ``block_sparse_attn_sm100a`` itself, so int64 inputs merely cost a cast.
+    """
+    if not _HAS_VSA_SM100A or not q.is_cuda:
+        return False
+    if torch.cuda.get_device_capability(q.device) != _SM100:
+        return False
+    if q.dtype != torch.bfloat16 or q.dim() != 4 or q.shape[-1] != HEAD_DIM:
+        return False
+    if not q.is_contiguous():
+        return False
+    if _FWD_BY_BLOCK.get(_block_size(q, variable_block_sizes)) is None:
+        return False
+    # A CTA owns an adjacent pair of query blocks.
+    if variable_block_sizes.numel() % 2 != 0:
+        return False
+    # Metadata must be integer-typed so the wrapper's int32 conversion is value-preserving.
+    if not variable_block_sizes.is_cuda or variable_block_sizes.dtype not in (torch.int32, torch.int64):
+        return False
+    return True
+
+
+@torch.library.custom_op(
+    "fastvideo_kernel::block_sparse_attn_sm100a_inference",
+    mutates_args=(),
+    device_types="cuda",
+)
+def _block_sparse_attn_sm100a_inference(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q2k_idx: torch.Tensor,
+    q2k_num: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> torch.Tensor:
+    """Opaque no-LSE launch used by the inference-only sm_100a route.
+
+    The extension is exposed as a raw pybind function rather than a dispatcher
+    op. Calling it directly makes Dynamo descend through a Python/C++ boundary
+    that has no fake implementation, so ``torch.compile(fullgraph=True)``
+    cannot capture a sparse H3 block. Keep that boundary inside this custom op;
+    its inputs have already been normalized by the public wrapper below.
+    """
+    fwd = _FWD_BY_BLOCK[_block_size(q, variable_block_sizes)]
+    sm_scale = 1.0 / (q.shape[-1]**0.5)
+    res = fwd(q, k, v, None, q2k_idx, q2k_num, variable_block_sizes, sm_scale, False)
+    return res[0]
+
+
+@torch.library.register_fake("fastvideo_kernel::block_sparse_attn_sm100a_inference")
+def _block_sparse_attn_sm100a_inference_fake(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q2k_idx: torch.Tensor,
+    q2k_num: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> torch.Tensor:
+    # The C++ binding allocates its output with torch::empty_like(q).
+    return torch.empty_like(q)
+
+
+@torch.library.custom_op(
+    "fastvideo_kernel::block_sparse_attn_sm100a_from_mask_inference",
+    mutates_args=(),
+    device_types="cuda",
+)
+def _block_sparse_attn_sm100a_from_mask_inference(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> torch.Tensor:
+    """Opaque mask compaction plus no-LSE sm_100a launch for inference.
+
+    H3 naturally produces a bool block map. Its Triton ``map_to_index`` call
+    must live behind the same opaque boundary as the raw pybind launch;
+    otherwise Dynamo sees that kernel before reaching the index-native custom
+    op and full-graph capture still fails.
+    """
+    from fastvideo_kernel.triton_kernels.index import map_to_index
+
+    q2k_idx, q2k_num = map_to_index(block_map)
+    return _block_sparse_attn_sm100a_inference(
+        q,
+        k,
+        v,
+        q2k_idx.to(torch.int32).contiguous(),
+        q2k_num.to(torch.int32).contiguous(),
+        variable_block_sizes,
+    )
+
+
+@torch.library.register_fake("fastvideo_kernel::block_sparse_attn_sm100a_from_mask_inference")
+def _block_sparse_attn_sm100a_from_mask_inference_fake(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> torch.Tensor:
+    return torch.empty_like(q)
+
+
+def block_sparse_attn_sm100a_from_mask(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> Tuple[torch.Tensor, None]:
+    """Inference forward from a bool block map, with compaction kept opaque."""
+    out = _block_sparse_attn_sm100a_from_mask_inference(
+        q.contiguous(),
+        k.contiguous(),
+        v.contiguous(),
+        block_map.to(torch.bool).contiguous(),
+        variable_block_sizes.to(torch.int32).contiguous(),
+    )
+    return out, None
+
+
+def block_sparse_attn_sm100a(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q2k_idx: torch.Tensor,
+    q2k_num: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+    need_lse: bool = True,
+) -> Tuple[torch.Tensor, torch.Tensor | None]:
+    """Forward pass. Returns ``(out, lse)``; ``out`` has q's layout."""
+    q = q.contiguous()
+    k = k.contiguous()
+    v = v.contiguous()
+    idx = q2k_idx.to(torch.int32).contiguous()
+    num = q2k_num.to(torch.int32).contiguous()
+    vbs = variable_block_sizes.to(torch.int32).contiguous()
+
+    if not need_lse:
+        # This is the production inference path. The custom op keeps the raw
+        # pybind launch opaque to Dynamo while its fake kernel carries output
+        # metadata through full-graph capture.
+        return _block_sparse_attn_sm100a_inference(q, k, v, idx, num, vbs), None
+
+    # Preserve the established LSE-producing path for correctness tests and
+    # any future forward/backward pairing; only inference needs the opaque op.
+    fwd = _FWD_BY_BLOCK[_block_size(q, vbs)]
+    sm_scale = 1.0 / (q.shape[-1]**0.5)
+    res = fwd(q, k, v, None, idx, num, vbs, sm_scale, True)
+    return res[0], res[1]

fastvideo-kernel/python/fastvideo_kernel/comm_ops.py

@@ -0,0 +1,86 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Ulysses sequence-parallel all-to-all ops.
+
+Thin wrappers over csrc/comm/ulysses_all_to_all.cu. The kernel stores directly
+into peers' memory through NCCL's device API, so the caller supplies an
+ncclComm_t for the group.
+"""
+
+import torch
+
+try:
+    from fastvideo_kernel._C import fastvideo_kernel_ops as _ops
+except ImportError:  # pragma: no cover - no compiled extension in this install
+    _ops = None
+
+_SUPPORTED_WORLD_SIZES = (2, 4, 6, 8)
+_REQUIRED_OPS = (
+    "allocate_ulysses_a2a",
+    "register_ulysses_a2a_window",
+    "create_ulysses_a2a_dev_comm",
+    "dispose_ulysses_a2a",
+    "ulysses_lsa_covers_group",
+    "ulysses_a2a",
+)
+
+
+def is_available() -> bool:
+    """Whether this wheel was built with the Ulysses all-to-all kernel."""
+    return _ops is not None and all(hasattr(_ops, name) for name in _REQUIRED_OPS)
+
+
+def _require() -> None:
+    if not is_available():
+        raise RuntimeError(
+            "the Ulysses all-to-all kernel is not present in this fastvideo-kernel build; "
+            "rebuild with ./build.sh or install a wheel that includes csrc/comm/")
+
+
+def lsa_covers_group(comm_ptr: int, world_size: int) -> bool:
+    """Whether every rank in the group is load-store accessible to every other."""
+    _require()
+    return bool(_ops.ulysses_lsa_covers_group(int(comm_ptr), int(world_size)))
+
+
+def allocate(nbytes: int, rank: int, world_size: int, device_index: int) -> int:
+    """Allocate one rank's local symmetric window without a collective."""
+    _require()
+    if world_size not in _SUPPORTED_WORLD_SIZES:
+        raise ValueError(f"ulysses a2a supports world sizes {_SUPPORTED_WORLD_SIZES}, "
+                         f"got {world_size}")
+    return int(_ops.allocate_ulysses_a2a(int(nbytes), int(rank), int(world_size), int(device_index)))
+
+
+def register_window(handle: int, comm_ptr: int) -> None:
+    """Register an allocated window with the supplied communicator.
+
+    Every rank in ``comm_ptr`` must call this together.
+    """
+    _require()
+    _ops.register_ulysses_a2a_window(int(handle), int(comm_ptr))
+
+
+def create_dev_comm(handle: int) -> None:
+    """Create the device communicator for a registered window collectively."""
+    _require()
+    _ops.create_ulysses_a2a_dev_comm(int(handle))
+
+
+def dispose(handle: int) -> None:
+    """Release a handle from :func:`allocate`. It is dangling afterwards."""
+    _require()
+    _ops.dispose_ulysses_a2a(int(handle))
+
+
+def all_to_all(handle: int, inp: torch.Tensor, out: torch.Tensor, B: int, S_local: int, H: int,
+               D: int, mode: int) -> None:
+    """Run one fused all-to-all on the current stream, writing into ``out``.
+
+    ``mode == 0``: ``[B, S_local, H, D] -> [B, S_global, H_local, D]``
+    ``mode == 1``: ``[B, S_global, H_local, D] -> [B, S_local, H, D]``
+
+    ``H`` is the global head count. Every rank must call with consistent
+    geometry in the same order.
+    """
+    _require()
+    _ops.ulysses_a2a(int(handle), inp, out, int(B), int(S_local), int(H), int(D), int(mode))

fastvideo-kernel/python/fastvideo_kernel/triton_kernels/block_sparse_attn_triton.py

@@ -16,14 +16,21 @@ import triton.language as tl
 import math  # small utility needed by the sparse wrapper
 # ──────────────────────────── SPARSE ADDITION END ─────────────────────────────
 
-# We don't run auto-tuning every time to keep the tutorial fast. Keeping
-# the code below and commenting out the equivalent parameters is convenient for
-# re-tuning.
+# BLOCK_M / BLOCK_N are fixed at 64 because they are structural, not tunable:
+# the kernel indexes the top-k list per BLOCK_M q-tile and addresses keys as
+# kv_idx * BLOCK_N, so both must match the granularity q2k_index and
+# variable_block_sizes were built at.
+#
+# num_stages / num_warps ARE free, and the previous {3, 4, 7} was inherited from
+# the upstream tutorial rather than tuned here. It skips 5 and 6; on Blackwell
+# (sm_121) the optimum is num_stages=5, so the search could not reach it. Both
+# block paths independently select 5 once it is available. Autotune still picks
+# per architecture, so other GPUs re-tune rather than inheriting this choice.
 configs = [
     triton.Config({'BLOCK_M': BM, 'BLOCK_N': BN}, num_stages=s, num_warps=w) \
     for BM in [64]\
     for BN in [64]\
-    for s in [3, 4, 7]\
+    for s in [2, 3, 4, 5, 6, 7]\
     for w in [4, 8]\
 ]
 
@@ -237,7 +244,12 @@ def _attn_bwd_dkdv(
         # Load m before computing qk to reduce pipeline stall.
         offs_m = start_m + block_sparse_offset + tl.arange(0, BLOCK_M1)
         m = tl.load(M + offs_m)
-        qkT = tl.dot(k, qT)
+        # Recompute logits exactly as the forward does: raw bf16 operands into
+        # the dot, fp32 scale after accumulation. A bf16 pre-scaled K perturbs
+        # the recomputed logits relative to the saved M by an error
+        # proportional to |logit|, which exp2 amplifies into arbitrarily wrong
+        # probabilities at large activations.
+        qkT = tl.dot(k, qT) * (sm_scale * 1.4426950408889634)
         pT = tl.math.exp2(qkT - m[None, :])
         mask = tl.arange(0, BLOCK_N1) < block_size
         pT = tl.where(mask[:, None], pT, 0.0)
@@ -268,6 +280,7 @@ def _attn_bwd_dq(
         do,
         m,
         D,
+        sm_scale,
         # shared by Q/K/V/DO.
         q2k_index,
         q2k_num,
@@ -315,7 +328,7 @@ def _attn_bwd_dq(
         block_sparse_offset = (kv_idx * 2 + half) * step_n * stride_tok
         kT = tl.load(kT_ptrs + block_sparse_offset)
         vT = tl.load(vT_ptrs + block_sparse_offset)
-        qk = tl.dot(q, kT)
+        qk = tl.dot(q, kT) * (sm_scale * 1.4426950408889634)
         p = tl.math.exp2(qk - m)
         offs_in_block = half * step_n + tl.arange(0, BLOCK_N2)
         mask = offs_in_block < block_size
@@ -324,8 +337,7 @@ def _attn_bwd_dq(
         dp = tl.dot(do, vT).to(tl.float32)
         ds = p * (dp - Di[:, None])
         ds = ds.to(tl.bfloat16)
-        # Compute dQ.
-        # NOTE: We need to de-scale dq in the end, because kT was pre-scaled.
+        # Compute dQ (kT is raw; the caller applies sm_scale once at the end).
         dq += tl.dot(ds, tl.trans(kT))
         # Increment pointers.
     return dq
@@ -453,6 +465,7 @@ def _attn_bwd(
         do,
         m,
         D,  #
+        sm_scale,
         q2k_index,
         q2k_num,
         max_kv_blks,
@@ -470,7 +483,7 @@ def _attn_bwd(
     )
     # Write back dQ.
     dq_ptrs = DQ + offs_m[:, None] * stride_tok + offs_k[None, :] * stride_d
-    dq *= LN2
+    dq *= sm_scale
     tl.store(dq_ptrs, dq)
 
 
@@ -591,6 +604,7 @@ def _attn_bwd_dq_kernel(
         Q,
         K,
         V,
+        sm_scale,
         DO,  #
         DQ,
         M,
@@ -663,6 +677,7 @@ def _attn_bwd_dq_kernel(
         do,
         m,
         D,
+        sm_scale,
         q2k_index,
         q2k_num,
         max_kv_blks,
@@ -680,7 +695,7 @@ def _attn_bwd_dq_kernel(
     )
 
     dq_ptrs = DQ + offs_m[:, None] * stride_tok + offs_k[None, :] * stride_d
-    dq_acc *= LN2
+    dq_acc *= sm_scale
     tl.store(dq_ptrs, dq_acc)
 
 
@@ -748,9 +763,11 @@ def triton_block_sparse_attn_backward(do, q, k, v, o, M, q2k_index, q2k_num, k2q
     dv = torch.empty_like(v)
     BATCH, N_HEAD = q.shape[:2]
     BLOCK_M1, BLOCK_N1, BLOCK_M2, BLOCK_N2 = 32, 64, 64, 32
-    RCP_LN2 = 1.4426950408889634  # = 1.0 / ln(2)
+    # K stays raw: the backward kernels apply sm_scale in fp32 after the dot,
+    # matching the forward's rounding exactly. (A bf16 pre-scaled K perturbs
+    # the recomputed logits vs the saved M; exp2 turns that into unboundedly
+    # wrong probabilities at large activations.)
     arg_k = k
-    arg_k = arg_k * (sm_scale * RCP_LN2)
     PRE_BLOCK = 64
     assert Tq % PRE_BLOCK == 0
     pre_grid = (Tq // PRE_BLOCK, BATCH * N_HEAD)
@@ -813,6 +830,7 @@ def triton_block_sparse_attn_backward(do, q, k, v, o, M, q2k_index, q2k_num, k2q
         q,
         arg_k,
         v,
+        sm_scale,
         do,
         dq,
         M,

fastvideo-kernel/python/fastvideo_kernel/version.py

@@ -1 +1 @@
-__version__ = "0.3.2"
+__version__ = "0.3.5"

fastvideo-kernel/python/fastvideo_kernel/vsa_utils.py

@@ -14,7 +14,9 @@ import math
 import torch
 
 VSA_TILE_SIZE = (4, 4, 4)
-_SUPPORTED_VSA_BLOCK_VOLUMES = (64, 256)
+# 128 is served by the sm_100a CUDA backend (block_sparse_attn_sm100a); 64 and 256 by
+# Triton and the CuTe-DSL path. A volume here only needs a backend that accepts it.
+_SUPPORTED_VSA_BLOCK_VOLUMES = (64, 128, 256)
 
 
 def _canonicalize_device(device: torch.device | str) -> torch.device:

fastvideo-kernel/tests/test_vsa_triton_backward_scale.py

@@ -0,0 +1,104 @@
+"""Regression: Triton block-sparse backward gradient parity at realistic activation scale.
+
+The backward used to fold ``sm_scale / ln(2)`` into K in bf16 before the
+exp2-based logit recompute. The bf16 rounding error on the pre-scaled K grows
+proportionally to |logit| and exp2 amplifies it into exponentially wrong
+probabilities, so dQ/dK/dV were correct at unit scale (every pre-existing test)
+but off by orders of magnitude at real activation magnitudes.
+
+This test sweeps the input scale and checks the Triton kernel's gradients
+against an fp32 masked-dense SDPA reference. The unit-scale case is the
+control (it passed even with the broken kernel); the large-scale cases are
+the regression.
+"""
+
+import pytest
+import torch
+
+from fastvideo_kernel.block_sparse_attn import _map_to_index, block_sparse_attn_triton
+
+from .utils import generate_block_sparse_mask_for_function
+
+BLOCK = 64
+
+
+@pytest.fixture(autouse=True)
+def _seed_rng():
+    """Pin the RNG so these cases do not depend on what ran before them.
+
+    Same convention as test_vsa_varlen.py: every tensor here comes from the
+    global torch RNG and the checks use tight thresholds, so an unseeded run
+    would shift inputs whenever an earlier test file draws a different number
+    of randoms.
+    """
+    torch.manual_seed(0)
+    torch.cuda.manual_seed_all(0)
+
+
+def _dense_reference(q, k, v, block_mask):
+    """fp32 masked-dense SDPA over the token-expanded block mask.
+
+    q/k/v: [B, H, S, D]; block_mask: [B, H, S // BLOCK, S // BLOCK] bool.
+    """
+    qf, kf, vf = q.float(), k.float(), v.float()
+    token_mask = block_mask.repeat_interleave(BLOCK, dim=-2).repeat_interleave(BLOCK, dim=-1)
+    logits = torch.matmul(qf, kf.transpose(-2, -1)) * (q.shape[-1]**-0.5)
+    logits = logits.masked_fill(~token_mask, float("-inf"))
+    return torch.matmul(logits.softmax(dim=-1), vf)
+
+
+@pytest.mark.cuda
+@pytest.mark.parametrize("scale", [1.0, 4.0, 16.0])
+def test_triton_backward_grad_parity_across_input_scales(scale: float) -> None:
+    """Kernel dQ/dK/dV must stay within a few percent of the fp32 reference
+    regardless of input magnitude.
+
+    With the bf16 K pre-scaling bug, scale<=4.0 passes at this geometry while
+    scale=16.0 fails (measured on GB200: dq relative L2 error 5.9e-1 vs 6.9e-3
+    fixed); at larger geometries and real activation magnitudes the broken
+    kernel is off by orders of magnitude. The passing unit-scale case is
+    exactly how the bug survived the original test suite.
+    """
+    if not torch.cuda.is_available():
+        pytest.skip("CUDA is required")
+
+    device = torch.device("cuda")
+    dtype = torch.bfloat16
+    batch, heads, dim = 1, 4, 128
+    num_blocks = 8
+    seq = num_blocks * BLOCK
+
+    q = torch.randn(batch, heads, seq, dim, device=device, dtype=dtype) * scale
+    k = torch.randn(batch, heads, seq, dim, device=device, dtype=dtype) * scale
+    v = torch.randn(batch, heads, seq, dim, device=device, dtype=dtype)
+    grad_out = torch.randn_like(q)
+
+    block_mask = generate_block_sparse_mask_for_function(heads, num_blocks, num_blocks, k=3,
+                                                         device=device).unsqueeze(0)
+    q2k_idx, q2k_num = _map_to_index(block_mask)
+    variable_block_sizes = torch.full((num_blocks, ), BLOCK, dtype=torch.int32, device=device)
+
+    q_ker, k_ker, v_ker = (t.detach().clone().requires_grad_(True) for t in (q, k, v))
+    out_ker, _ = block_sparse_attn_triton(q_ker, k_ker, v_ker, q2k_idx, q2k_num, variable_block_sizes)
+    (out_ker.float() * grad_out.float()).sum().backward()
+
+    q_ref, k_ref, v_ref = (t.detach().clone().requires_grad_(True) for t in (q, k, v))
+    out_ref = _dense_reference(q_ref, k_ref, v_ref, block_mask)
+    (out_ref * grad_out.float()).sum().backward()
+
+    # Forward is exact at any scale; this pins the harness itself.
+    fwd_rel = ((out_ker.float() - out_ref).norm() / out_ref.norm()).item()
+    assert fwd_rel < 2e-2, f"scale={scale}: forward rel err {fwd_rel:.3e}"
+
+    for name, g_ker, g_ref in (
+        ("dq", q_ker.grad, q_ref.grad),
+        ("dk", k_ker.grad, k_ref.grad),
+        ("dv", v_ker.grad, v_ref.grad),
+    ):
+        assert torch.isfinite(g_ker).all().item(), f"scale={scale}: non-finite {name}"
+        ref_norm = g_ref.float().norm()
+        rel = ((g_ker.float() - g_ref.float()).norm() / ref_norm.clamp_min(1e-12)).item()
+        ratio = (g_ker.float().norm() / ref_norm.clamp_min(1e-12)).item()
+        print(f"scale={scale} {name}: rel_l2={rel:.4e} norm_ratio={ratio:.4f}")
+        assert rel < 5e-2, f"scale={scale}: {name} rel l2 err {rel:.3e} >= 5e-2"
+        assert 0.98 < ratio < 1.02, f"scale={scale}: {name} grad-norm ratio {ratio:.4f}"

fastvideo/api/compat.py

@@ -44,6 +44,16 @@ _LEGACY_REQUEST_ALIASES = {
 _REQUEST_PIPELINE_OVERRIDE_FIELDS = frozenset({
     "embedded_cfg_scale",
 })
+REQUEST_BATCH_EXTRA_PASSTHROUGH_FIELDS = (
+    "ltx2_audio_latents",
+    "ltx2_audio_clean_latent",
+    "ltx2_audio_denoise_mask",
+    "audio_num_frames",
+    "video_position_offset_sec",
+    "vsa_mode",
+    "vsa_dense_first_n_steps",
+    "vsa_dense_layers",
+)
 # torch.compile kwargs that map to first-class CompileConfig fields.
 _COMPILE_TYPED_KEYS = ("backend", "fullgraph", "mode", "dynamic")
 # LTX-2 refine flat kwargs (init + per-request) known to FastVideoArgs.
@@ -165,6 +175,10 @@ def legacy_from_pretrained_to_config(
             pipeline["workload_type"] = value
         elif key == "lora_path":
             components["lora_path"] = value
+        elif key == "lora_nickname":
+            components["lora_nickname"] = value
+        elif key == "lora_strength":
+            components["lora_strength"] = value
         elif key == "override_pipeline_cls_name":
             components["override_pipeline_cls_name"] = value
         elif key == "override_transformer_cls_name":
@@ -283,6 +297,8 @@ def generator_config_to_fastvideo_args(config: GeneratorConfig | Mapping[str, An
         kwargs["pipeline_config"] = components.pipeline_config_path
     if components.lora_path is not None:
         kwargs["lora_path"] = components.lora_path
+        kwargs["lora_nickname"] = components.lora_nickname
+        kwargs["lora_strength"] = components.lora_strength
     if components.override_pipeline_cls_name is not None:
         kwargs["override_pipeline_cls_name"] = components.override_pipeline_cls_name
     if components.override_transformer_cls_name is not None:
@@ -368,7 +384,7 @@ def request_to_sampling_param(
     for key, value in updates.items():
         if hasattr(sampling_param, key):
             setattr(sampling_param, key, deepcopy(value))
-        elif key in _REQUEST_PIPELINE_OVERRIDE_FIELDS:
+        elif key in _REQUEST_PIPELINE_OVERRIDE_FIELDS or key in REQUEST_BATCH_EXTRA_PASSTHROUGH_FIELDS:
             continue
         elif value == _SCHEMA_DEFAULT_UPDATES.get(key, _MISSING):
             # Schema-default field that isn't on SamplingParam; tolerated
@@ -467,6 +483,14 @@ def request_to_pipeline_overrides(request: GenerationRequest) -> dict[str, Any]:
     return overrides
 
 
+def request_to_batch_extra(request: GenerationRequest) -> dict[str, Any]:
+    """Extract typed-request extensions consumed through ``ForwardBatch.extra``."""
+    return {
+        key: deepcopy(value)
+        for key, value in explicit_request_updates(request).items() if key in REQUEST_BATCH_EXTRA_PASSTHROUGH_FIELDS
+    }
+
+
 def explicit_request_updates(request: GenerationRequest) -> dict[str, Any]:
     """Project a ``GenerationRequest`` down to *explicitly set* fields only.
 
@@ -640,6 +664,7 @@ def _validate_batched_input_length(
 
 
 __all__ = [
+    "REQUEST_BATCH_EXTRA_PASSTHROUGH_FIELDS",
     "explicit_request_updates",
     "generator_config_to_fastvideo_args",
     "legacy_from_pretrained_to_config",
@@ -648,6 +673,7 @@ __all__ = [
     "normalize_generation_request",
     "normalize_generator_config",
     "register_continuation_kind",
+    "request_to_batch_extra",
     "request_to_pipeline_overrides",
     "request_to_sampling_param",
 ]

fastvideo/api/schema.py

@@ -10,6 +10,7 @@ class ServerConfig:
     host: str = "0.0.0.0"
     port: int = 8000
     output_dir: str = "outputs/"
+    served_model_name: str | None = None
 
 
 @dataclass
@@ -94,6 +95,8 @@ class ComponentConfig:
     vae_weights: str | None = None
     upsampler_weights: str | None = None
     lora_path: str | None = None
+    lora_nickname: str = "default"
+    lora_strength: float = 1.0
     override_pipeline_cls_name: str | None = None
     override_transformer_cls_name: str | None = None
 

fastvideo/attention/backends/flash_attn.py

@@ -10,6 +10,14 @@ from fastvideo.attention.utils.flash_attn_default import (
     flash_attn_func_compilable,
 )
 
+if fa_version == "4":
+    # The FA4 varlen wrapper is already a compile-safe custom op. Keep the
+    # import conditional so FA2/FA3 environments do not need flash_attn.cute.
+    from fastvideo.attention.utils.flash_attn_cute import (
+        flash_attn_varlen_func as flash_attn_varlen_func_compilable, )
+else:
+    flash_attn_varlen_func_compilable = None
+
 from fastvideo.attention.backends.abstract import (
     AttentionBackend,
     AttentionImpl,
@@ -19,7 +27,12 @@ from fastvideo.attention.backends.abstract import (
 from fastvideo.logger import init_logger
 
 logger = init_logger(__name__)
-logger.info("Using FlashAttention-%s backend", fa_version)
+# Every worker records the loaded FlashAttention implementation so a
+# distributed profiling log contains one backend receipt per rank.
+logger.info("Worker %s Using FlashAttention-%s backend",
+            os.environ.get("RANK", "0"),
+            fa_version,
+            local_main_process_only=False)
 
 # FP4 FA4 support: quantize Q/K to NVFP4 E2M1 for block-scaled MMA on Blackwell.
 # Requires: flash-attention-fp4, flashinfer, cutlass-dsl. Enable via nvfp4_fa4=True kwarg.
@@ -33,6 +46,21 @@ except ImportError:
     _FA4_FP4_AVAILABLE = False
 
 _FA4_QUANT_OPS: tuple | None = None
+_FA4_PACKED_VARLEN_CONFIG_LOGGED = False
+
+
+def _log_fa4_packed_varlen_config() -> None:
+    """Emit one configuration receipt per worker process.
+
+    This intentionally does not claim that a particular forward used the
+    kernel: masks, gradients, batch size, unequal Q/K lengths, and NVFP4 are
+    runtime guards evaluated later in ``_forward_impl``.
+    """
+    global _FA4_PACKED_VARLEN_CONFIG_LOGGED
+    if not _FA4_PACKED_VARLEN_CONFIG_LOGGED:
+        logger.info("MiniMax-H3 dense attention: FA4 packed-varlen route configured (runtime guards apply)",
+                    local_main_process_only=False)
+        _FA4_PACKED_VARLEN_CONFIG_LOGGED = True
 
 
 def _import_fa4_quant_ops() -> tuple:
@@ -223,7 +251,21 @@ class FlashAttentionImpl(AttentionImpl):
     ) -> None:
         self.causal = causal
         self.softmax_scale = softmax_scale
-        self.nvfp4_fa4 = extra_impl_args.get("nvfp4_fa4", False) or os.environ.get("FASTVIDEO_NVFP4_FA4", "0") == "1"
+        # MiniMax-H3's dense DiT explicitly enables this faster FA4 entry
+        # point. It remains off for every other model and for the H3 text
+        # refiner; grad-enabled calls stay on the established fixed path.
+        self.fa4_packed_varlen = bool(extra_impl_args.get("fa4_packed_varlen", False))
+        if self.fa4_packed_varlen and fa_version == "4":
+            _log_fa4_packed_varlen_config()
+        # An explicit ``nvfp4_fa4`` impl arg wins over the process-wide
+        # FASTVIDEO_NVFP4_FA4 env opt-in, so precision-sensitive layers (e.g.
+        # the FP32-pinned H3 VAE attention) can force-disable FP4 Q/K
+        # quantization while the DiT keeps it. When the arg is absent the env
+        # keeps its previous semantics.
+        nvfp4_fa4 = extra_impl_args.get("nvfp4_fa4")
+        if nvfp4_fa4 is None:
+            nvfp4_fa4 = os.environ.get("FASTVIDEO_NVFP4_FA4", "0") == "1"
+        self.nvfp4_fa4 = bool(nvfp4_fa4)
         if self.nvfp4_fa4:
             cap = torch.cuda.get_device_capability()
             assert cap in [(10, 0), (10, 3)], (f"NVFP4 FA4 requires Blackwell (sm100a/sm103a), got sm{cap[0]}{cap[1]}")
@@ -318,6 +360,29 @@ class FlashAttentionImpl(AttentionImpl):
         elif self.nvfp4_fa4:
             output = self._forward_nvfp4(query, key, value)
 
+        elif (self.fa4_packed_varlen and fa_version == "4" and not torch.is_grad_enabled() and query.shape[0] == 1
+              and query.shape[1] == key.shape[1] == value.shape[1]):
+            # FA4's packed-varlen entry point is materially faster for H3's
+            # long, single-document self-attention. Flatten only the batch
+            # dimension and describe that one sequence with CUDA int32
+            # cumulative lengths; the existing custom-op wrapper keeps this
+            # route traceable under torch.compile(fullgraph=True).
+            assert flash_attn_varlen_func_compilable is not None
+            sequence_length = query.shape[1]
+            cu_seqlens = torch.arange(2, dtype=torch.int32, device=query.device) * sequence_length
+            output = flash_attn_varlen_func_compilable(
+                query.squeeze(0),
+                key.squeeze(0),
+                value.squeeze(0),
+                cu_seqlens,
+                cu_seqlens,
+                sequence_length,
+                sequence_length,
+                dropout_p=0.0,
+                softmax_scale=self.softmax_scale,
+                causal=self.causal,
+            ).unsqueeze(0)
+
         else:
             # Route through the compilable wrapper so dynamo sees a
             # registered op (no graph break) for FA2/FA3; identical

fastvideo/attention/backends/video_sparse_attn_h3.py

@@ -5,13 +5,17 @@ H3 runs one joint bidirectional attention over
 ``[text | condition keyframes | audio | generated video]``, so this
 backend differs from the Wan-tuned ``video_sparse_attn``:
 
-- Tiles are ``[segment-pure prefix chunks] + [3D (4,8,8) video tiles]``;
-  prefix tiles never straddle segment boundaries.
+- Tiles are ``[segment-pure prefix chunks] + [3D video tiles]``; prefix
+  tiles never straddle segment boundaries. The tile size is selectable at
+  metadata build time: 256 tokens ``(4,8,8)`` (default) or 64 tokens
+  ``(4,4,4)`` (see ``VSA_H3_TILE_SHAPES``).
 - Selection is pure Python on pooled tile scores; the block-sparse kernel
   consumes an explicit bool mask, so no kernel changes are needed.
-- The compression branch is gated by ``to_gate_compress``, which the H3
-  checkpoint does not carry: the loader zero-initializes it, so untrained
-  inference is exactly pure sparse and finetuning can learn the gate.
+- The compression branch is gated by ``to_gate_compress``, which the base
+  H3 checkpoint does not carry: the loader zero-initializes it, so
+  untrained inference is exactly pure sparse and finetuning can learn the
+  gate. VSA-distilled students (e.g. FastVideo-Minimax-H3-Preview) ship
+  trained gates, which load and activate the branch.
 - Non-video *queries* are always dense. Non-video *keys* are either
   always-selected for every query ("exempt", default) or compete in
   top-k under a FLOP-matched budget ("compete") — the ablation axis,
@@ -20,22 +24,50 @@ backend differs from the Wan-tuned ``video_sparse_attn``:
   (``vsa_dense_first_n_steps``, ``vsa_dense_layers``) let mixed schedules
   run the diffuse steps/layers dense while pushing the rest harder.
 
-Targets sm10.x through the FA4 CuTe 256-tile path
+At tile 256 this targets sm10.x through the FA4 CuTe 256-tile path
 (``FASTVIDEO_VSA_CUTEDSL=1``); the Triton 256→64 expansion is the
-fallback and keeps identical mask semantics.
+fallback and keeps identical mask semantics. At tile 64 the block map is
+already at the kernels' native 64-token granularity, so both forward and
+backward run the Triton block-sparse kernels directly (no expansion,
+``FASTVIDEO_VSA_CUTEDSL`` does not apply). A third, opt-in route exists
+for the tile-64 FORWARD only: ``FASTVIDEO_VSA_SM100A=1`` sends no-grad
+forwards through the sm_100a CUDA block-sparse kernel
+(``fastvideo_kernel.block_sparse_attn_sm100a``, upstream PR #1719 plus
+our per-q-tile ``q2k_num`` fix) when the extension is built, the device
+is sm_100, and the geometry qualifies. The CUDA kernel assigns adjacent
+pairs of query tiles to CTAs, so an odd logical tile count receives one
+internal, zero-valid partner tile for the no-grad call only. Score search,
+the trained mask, gate-compress, and the returned packed sequence remain on
+the original logical tiles. Grad-tracking forwards and every backward stay
+on Triton unchanged. If the env is set but a precondition fails, the route
+logs one warning and falls back.
 """
 
 import functools
 import math
+import os
 from dataclasses import dataclass
 from typing import Any
 
 import torch
 
 try:
+    from fastvideo_kernel.block_sparse_attn import block_sparse_attn as block_sparse_attn_64_bhsd
     from fastvideo_kernel.block_sparse_attn_256 import block_sparse_attn_256_bshd
+    from fastvideo_kernel.triton_kernels.index import map_to_index
 except ImportError:
+    block_sparse_attn_64_bhsd = None
     block_sparse_attn_256_bshd = None
+    map_to_index = None
+
+try:
+    # Optional: only present in fastvideo_kernel builds that carry the sm_100a
+    # CUDA block-sparse forward (upstream PR #1719). The module itself imports
+    # fine without the compiled symbols (`_HAS_VSA_SM100A` is then False and
+    # `is_supported` says no), so this only guards *module* availability.
+    from fastvideo_kernel import block_sparse_attn_sm100a as _sm100a
+except ImportError:
+    _sm100a = None
 
 from fastvideo.attention.backends.abstract import (AttentionBackend, AttentionImpl, AttentionMetadata,
                                                    AttentionMetadataBuilder, layer_idx_from_prefix)
@@ -43,51 +75,174 @@ from fastvideo.attention.backends.video_sparse_attn import (compute_topk, constr
                                                             get_non_pad_index, get_tile_partition_indices,
                                                             scatter_into_tile_buf)
 from fastvideo.attention.backends.video_sparse_attn_h3_probe import probe_enabled, record_probe
+from fastvideo.logger import init_logger
 
-VSA_H3_TILE_SIZE = (4, 8, 8)  # 256 elements -> FA4 CuTe fastpath on sm10.x
+logger = init_logger(__name__)
+
+# Opt-in switch for the sm_100a CUDA forward on the tile-64 no-grad path.
+VSA_SM100A_ENV = "FASTVIDEO_VSA_SM100A"
+
+VSA_H3_TILE_SIZE = (4, 8, 8)  # 256 elements -> FA4 CuTe fastpath on sm10.x (default)
 _TILE_ELEMS = math.prod(VSA_H3_TILE_SIZE)
+# Selectable tile geometries, keyed by element count (= the build-time
+# ``tile_size``). 64 runs the native 64-token Triton block-sparse kernels for
+# forward AND backward — the block map is already at kernel granularity, so no
+# 256->64 mask expansion is involved and FASTVIDEO_VSA_CUTEDSL does not apply.
+VSA_H3_TILE_SHAPES: dict[int, tuple[int, int, int]] = {
+    _TILE_ELEMS: VSA_H3_TILE_SIZE,
+    64: (4, 4, 4),
+}
 
 
-def token_tile_and_valid(variable_block_sizes: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor]:
+@torch.library.custom_op(
+    "fastvideo::h3_vsa_sm100a_from_mask_compat",
+    mutates_args=(),
+    device_types="cuda",
+)
+def _h3_vsa_sm100a_from_mask_compat(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> torch.Tensor:
+    """Compile-safe mask adapter for kernel wheels predating the mask API."""
+    if _sm100a is None or map_to_index is None:
+        raise RuntimeError("The sm100a compatibility route requires the raw kernel and map_to_index")
+    q2k_idx, q2k_num = map_to_index(block_map)
+    out, _ = _sm100a.block_sparse_attn_sm100a(
+        q,
+        k,
+        v,
+        q2k_idx.to(torch.int32).contiguous(),
+        q2k_num.to(torch.int32).contiguous(),
+        variable_block_sizes.to(torch.int32).contiguous(),
+        need_lse=False,
+    )
+    return out
+
+
+@torch.library.register_fake("fastvideo::h3_vsa_sm100a_from_mask_compat")
+def _h3_vsa_sm100a_from_mask_compat_fake(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> torch.Tensor:
+    del k, v, block_map, variable_block_sizes
+    return torch.empty_like(q)
+
+
+def _sm100a_has_compile_safe_mask_route(sm100a_mod: Any) -> bool:
+    return (callable(getattr(sm100a_mod, "block_sparse_attn_sm100a_from_mask", None))
+            or (callable(getattr(sm100a_mod, "block_sparse_attn_sm100a", None)) and map_to_index is not None))
+
+
+def _sm100a_from_mask(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    block_map: torch.Tensor,
+    variable_block_sizes: torch.Tensor,
+) -> tuple[torch.Tensor, None]:
+    """Use the native mask entry when installed, otherwise the local adapter."""
+    native = getattr(_sm100a, "block_sparse_attn_sm100a_from_mask", None)
+    if callable(native):
+        return native(q, k, v, block_map, variable_block_sizes)
+    return _h3_vsa_sm100a_from_mask_compat(q, k, v, block_map, variable_block_sizes), None
+
+
+def token_tile_and_valid(variable_block_sizes: torch.Tensor,
+                         tile_elems: int = _TILE_ELEMS) -> tuple[torch.Tensor, torch.Tensor]:
     """Per padded-token tile id and pad-validity mask.
 
     The single encoding of the padding contract, shared by the probe and the
     test oracle so they cannot drift from the backend's tile geometry.
+    ``tile_elems`` must match the metadata the sizes came from
+    (``MiniMaxH3VSAMetadata.tile_elems``).
     """
     device = variable_block_sizes.device
-    token_tile = torch.arange(variable_block_sizes.numel(), device=device).repeat_interleave(_TILE_ELEMS)
-    token_valid = (torch.arange(_TILE_ELEMS, device=device)[None, :] < variable_block_sizes[:, None]).reshape(-1)
+    token_tile = torch.arange(variable_block_sizes.numel(), device=device).repeat_interleave(tile_elems)
+    token_valid = (torch.arange(tile_elems, device=device)[None, :] < variable_block_sizes[:, None]).reshape(-1)
     return token_tile, token_valid
 
 
+def _validate_h3_tile_geometry(
+    prefix_segments: tuple[int, ...],
+    dit_seq_shape: tuple[int, int, int],
+    variable_block_sizes: torch.Tensor,
+    untile_combined_index: torch.Tensor,
+    tile_elems: int = _TILE_ELEMS,
+) -> None:
+    """Fail synchronously on out-of-bounds tile geometry.
+
+    Invariants the block-sparse kernel trusts without checking:
+    every tile's valid size is in (0, tile_elems]; the sizes sum to the
+    packed sequence length; and ``untile_combined_index`` maps each packed
+    row to exactly one non-pad slot of the padded tile buffer. A violation
+    would surface only as an async device fault at some later kernel or
+    collective (e.g. an FSDP all-gather), which is unattributable — so raise
+    here, once per cached geometry, with the numbers in hand.
+    """
+    total = sum(prefix_segments) + math.prod(dit_seq_shape)
+    n_pad = variable_block_sizes.numel() * tile_elems
+    sizes_min = int(variable_block_sizes.min())
+    sizes_max = int(variable_block_sizes.max())
+    sizes_sum = int(variable_block_sizes.sum())
+    if sizes_min < 1 or sizes_max > tile_elems or sizes_sum != total:
+        raise ValueError(f"VSA-H3 tile sizes out of bounds for prefix={prefix_segments}, video={dit_seq_shape}, "
+                         f"tile_elems={tile_elems}: min={sizes_min}, max={sizes_max}, sum={sizes_sum}, "
+                         f"expected sum={total}.")
+    if untile_combined_index.numel() != total:
+        raise ValueError(f"VSA-H3 untile index has {untile_combined_index.numel()} entries for a packed "
+                         f"sequence of {total} rows (prefix={prefix_segments}, video={dit_seq_shape}).")
+    idx_min = int(untile_combined_index.min())
+    idx_max = int(untile_combined_index.max())
+    if idx_min < 0 or idx_max >= n_pad:
+        # Range first: the pad-slot gather below would itself index out of
+        # bounds (the very async fault this guard exists to preempt).
+        raise ValueError(f"VSA-H3 untile index is not an injective map into non-pad slots: range "
+                         f"[{idx_min}, {idx_max}] vs padded length {n_pad} "
+                         f"(prefix={prefix_segments}, video={dit_seq_shape}).")
+    in_tile_offset = untile_combined_index % tile_elems
+    maps_into_pad = bool((in_tile_offset >= variable_block_sizes[untile_combined_index // tile_elems]).any())
+    if maps_into_pad or int(torch.unique(untile_combined_index).numel()) != total:
+        raise ValueError(f"VSA-H3 untile index is not an injective map into non-pad slots: "
+                         f"pad-slot hit={maps_into_pad} "
+                         f"(prefix={prefix_segments}, video={dit_seq_shape}).")
+
+
 @functools.lru_cache(maxsize=10)
 def _h3_tile_geometry(
     prefix_segments: tuple[int, ...],
     dit_seq_shape: tuple[int, int, int],
     device: torch.device,
+    tile_shape: tuple[int, int, int] = VSA_H3_TILE_SIZE,
 ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, int, int]:
     """Tile the packed sequence: segment-pure prefix chunks, then video tiles.
 
     Returns (tile_partition_indices, variable_block_sizes,
     untile_combined_index, num_prefix_tiles, num_video_tiles).
     """
+    tile_elems = math.prod(tile_shape)
     prefix_len = sum(prefix_segments)
 
     prefix_sizes: list[int] = []
     for segment in prefix_segments:
-        full, rem = divmod(segment, _TILE_ELEMS)
-        prefix_sizes.extend([_TILE_ELEMS] * full)
+        full, rem = divmod(segment, tile_elems)
+        prefix_sizes.extend([tile_elems] * full)
         if rem:
             prefix_sizes.append(rem)
     num_prefix_tiles = len(prefix_sizes)
 
-    ts_t, ts_h, ts_w = VSA_H3_TILE_SIZE
+    ts_t, ts_h, ts_w = tile_shape
     t, h, w = dit_seq_shape
     num_tiles = (math.ceil(t / ts_t), math.ceil(h / ts_h), math.ceil(w / ts_w))
-    video_sizes = construct_variable_block_sizes(dit_seq_shape, num_tiles, device, VSA_H3_TILE_SIZE)
+    video_sizes = construct_variable_block_sizes(dit_seq_shape, num_tiles, device, tile_shape)
     num_video_tiles = int(video_sizes.numel())
 
-    video_indices = get_tile_partition_indices(dit_seq_shape, VSA_H3_TILE_SIZE, device) + prefix_len
+    video_indices = get_tile_partition_indices(dit_seq_shape, tile_shape, device) + prefix_len
     tile_partition_indices = torch.cat([
         torch.arange(prefix_len, device=device, dtype=torch.long),
         video_indices,
@@ -100,9 +255,11 @@ def _h3_tile_geometry(
 
     # get_non_pad_index is lru-cached on tensor identity; variable_block_sizes
     # is itself cached by this function, so the identity stays stable.
-    non_pad_index = get_non_pad_index(variable_block_sizes, _TILE_ELEMS)
+    non_pad_index = get_non_pad_index(variable_block_sizes, tile_elems)
 
     untile_combined_index = non_pad_index[torch.argsort(tile_partition_indices)]
+    # One-time (lru-cached) synchronous bounds check; see _validate_h3_tile_geometry.
+    _validate_h3_tile_geometry(prefix_segments, dit_seq_shape, variable_block_sizes, untile_combined_index, tile_elems)
     return (tile_partition_indices, variable_block_sizes, untile_combined_index, num_prefix_tiles, num_video_tiles)
 
 
@@ -131,6 +288,14 @@ class MiniMaxH3VSABackend(AttentionBackend):
         return MiniMaxH3VSAMetadataBuilder
 
 
+class _MiniMaxH3VSATileBufferHolder:
+    """Builder-owned no-grad tile scratch and its active geometry."""
+
+    def __init__(self) -> None:
+        self.buffer: torch.Tensor | None = None
+        self.untile_geometry: torch.Tensor | None = None
+
+
 @dataclass
 class MiniMaxH3VSAMetadata(AttentionMetadata):
     total_seq_length: int
@@ -139,44 +304,54 @@ class MiniMaxH3VSAMetadata(AttentionMetadata):
     exempt: bool
     variable_block_sizes: torch.Tensor
     untile_combined_index: torch.Tensor
+    # Device-side copy of ``dense_layers``. Regional fullgraph capture uses
+    # this tensor with each implementation's tensor-valued layer index so the
+    # shared block code does not specialize once per Python ``layer_idx``.
+    dense_layers_tensor: torch.Tensor
+    # tokens per tile (256 or 64); selects the tile geometry AND the kernel
+    # route in forward() (256 -> VSA-256 CuTe/Triton, 64 -> native Triton)
+    tile_elems: int = _TILE_ELEMS
     # layers forced dense regardless of sparsity (probe-guided opt-outs)
     dense_layers: tuple[int, ...] = ()
-    # Single-slot holder for the padded tile buffer, owned by the BUILDER so
-    # one buffer serves the whole denoising loop (pad slots stay zero and
-    # every non-pad slot is fully overwritten per tile(), so cross-step reuse
-    # is valid; saves a ~1.4 GB alloc+memset per step at 720p). VSA-H3 runs
-    # eager today — revisit the reuse if it ever goes under cudagraphs.
-    tile_buf_holder: list = None  # type: ignore[assignment]
+    # Builder-owned padded tile buffer. It records the geometry that last
+    # populated the allocation so a same-shaped geometry change can clear
+    # stale pad rows once while steady-state denoising reuses the buffer.
+    tile_buf_holder: _MiniMaxH3VSATileBufferHolder | None = None
 
 
 class MiniMaxH3VSAMetadataBuilder(AttentionMetadataBuilder):
 
     def __init__(self) -> None:
-        self._tile_buf_holder: list = [None]
+        self._tile_buf_holder = _MiniMaxH3VSATileBufferHolder()
 
     def prepare(self) -> None:
         pass
 
     def build(  # type: ignore
-            self,
-            current_timestep: int,
-            raw_latent_shape: tuple[int, int, int],
-            patch_size: tuple[int, int, int],
-            VSA_sparsity: float,
-            prefix_segments: tuple[int, ...],
-            device: torch.device,
-            exempt: bool = True,
-            dense_layers: tuple[int, ...] = (),
-            **kwargs: dict[str, Any],
+        self,
+        current_timestep: int,
+        raw_latent_shape: tuple[int, int, int],
+        patch_size: tuple[int, int, int],
+        VSA_sparsity: float,
+        prefix_segments: tuple[int, ...],
+        device: torch.device,
+        exempt: bool = True,
+        dense_layers: tuple[int, ...] = (),
+        tile_size: int = _TILE_ELEMS,
+        **kwargs: dict[str, Any],
     ) -> MiniMaxH3VSAMetadata:
+        tile_shape = VSA_H3_TILE_SHAPES.get(int(tile_size))
+        if tile_shape is None:
+            raise ValueError(f"VSA-H3 tile_size must be one of {sorted(VSA_H3_TILE_SHAPES)}, got {tile_size!r}")
         dit_seq_shape = (raw_latent_shape[0] // patch_size[0], raw_latent_shape[1] // patch_size[1],
                          raw_latent_shape[2] // patch_size[2])
         prefix_segments = tuple(int(s) for s in prefix_segments if s > 0)
         total_seq_length = sum(prefix_segments) + math.prod(dit_seq_shape)
 
         (_tile_partition_indices, variable_block_sizes, untile_combined_index, num_prefix_tiles,
-         num_video_tiles) = _h3_tile_geometry(prefix_segments, dit_seq_shape, device)
+         num_video_tiles) = _h3_tile_geometry(prefix_segments, dit_seq_shape, device, tile_shape)
 
+        dense_layers = tuple(int(layer) for layer in dense_layers)
         return MiniMaxH3VSAMetadata(
             current_timestep=current_timestep,
             VSA_sparsity=VSA_sparsity,
@@ -186,13 +361,15 @@ class MiniMaxH3VSAMetadataBuilder(AttentionMetadataBuilder):
             exempt=exempt,
             variable_block_sizes=variable_block_sizes,
             untile_combined_index=untile_combined_index,
-            dense_layers=tuple(int(layer) for layer in dense_layers),
+            tile_elems=int(tile_size),
+            dense_layers=dense_layers,
+            dense_layers_tensor=torch.tensor(dense_layers, device=device, dtype=torch.int64),
             tile_buf_holder=self._tile_buf_holder,
         )
 
 
-def _pool_tiles(x: torch.Tensor, variable_block_sizes: torch.Tensor) -> torch.Tensor:
-    """fp32 mean over each 256-token tile. x: [B, S_pad, H, D] -> [B, H, n_tiles, D].
+def _pool_tiles(x: torch.Tensor, variable_block_sizes: torch.Tensor, tile_elems: int = _TILE_ELEMS) -> torch.Tensor:
+    """fp32 mean over each tile_elems-token tile. x: [B, S_pad, H, D] -> [B, H, n_tiles, D].
 
     Pad positions in the tile buffer are guaranteed zero (zeros-init, never
     written), so a plain sum with fp32 accumulation needs no validity mask
@@ -200,8 +377,8 @@ def _pool_tiles(x: torch.Tensor, variable_block_sizes: torch.Tensor) -> torch.Te
     the masked mean exactly.
     """
     batch, seq_len, heads, dim = x.shape
-    n_tiles = seq_len // _TILE_ELEMS
-    pooled = x.view(batch, n_tiles, _TILE_ELEMS, heads, dim).sum(dim=2, dtype=torch.float32)
+    n_tiles = seq_len // tile_elems
+    pooled = x.view(batch, n_tiles, tile_elems, heads, dim).sum(dim=2, dtype=torch.float32)
     pooled = pooled / variable_block_sizes.view(1, -1, 1, 1)
     return pooled.permute(0, 2, 1, 3)
 
@@ -232,6 +409,24 @@ def _build_block_mask(
     return mask
 
 
+def _sm100a_unavailable_reason(sm100a_mod: Any, query_bhsd: torch.Tensor, variable_block_sizes: torch.Tensor,
+                               grad_mode: bool) -> str | None:
+    """Why the opt-in sm_100a forward route cannot run here, or None if it can.
+
+    Pure decision logic, split out so the routing is unit-testable without a
+    GPU or the compiled extension (tests substitute ``sm100a_mod``). Order
+    matters only for the message: the cheapest, most actionable reason first.
+    """
+    if sm100a_mod is None:
+        return "fastvideo_kernel.block_sparse_attn_sm100a is not installed"
+    if grad_mode:
+        return "inputs require grad and the sm_100a kernel is forward-only; grad paths keep Triton"
+    if not sm100a_mod.is_supported(query_bhsd, variable_block_sizes):
+        return ("block_sparse_attn_sm100a.is_supported returned False (needs an sm_100 device, a built "
+                "extension, bf16, head_dim 128, an even tile count, and integer tile sizes)")
+    return None
+
+
 class MiniMaxH3VSAImpl(AttentionImpl):
 
     def __init__(
@@ -246,26 +441,109 @@ class MiniMaxH3VSAImpl(AttentionImpl):
     ) -> None:
         self.prefix = prefix
         self.layer_idx = layer_idx_from_prefix(prefix, default=-1)
+        self.head_size = head_size
+        # None means the regional-compile preparation hook has not run.  The
+        # eager path deliberately ignores this cache and preserves its
+        # request-time env/probe/fallback behavior; only Dynamo capture reads
+        # the prepared, static route.
+        self._regional_compile_sm100a_enabled: bool | None = None
+        self._regional_compile_layer_idx: torch.Tensor | None = None
+
+    def prepare_for_regional_compile(self, device: torch.device) -> str | None:
+        """Resolve the inference-only sm_100a route before fullgraph capture.
+
+        The ordinary eager route probes the environment, extension, device,
+        and tensor contract at every call so it can warn and fall back.  Those
+        Python/device-capability checks are not safe inside a regional
+        ``fullgraph=True`` block.  Probe one representative tile-64 input on
+        the loaded model's device now, then let ``forward`` specialize on the
+        resulting plain bool while Dynamo is compiling.
+        """
+        requested = os.environ.get(VSA_SM100A_ENV, "0") == "1"
+        enabled = False
+        reason = None if requested else f"{VSA_SM100A_ENV}=1 is required for compile-safe VSA-H3 attention"
+        if requested:
+            if _sm100a is None:
+                reason = "fastvideo_kernel.block_sparse_attn_sm100a is not installed"
+            elif not _sm100a_has_compile_safe_mask_route(_sm100a):
+                reason = ("neither a native block_sparse_attn_sm100a_from_mask entry nor the raw sm100a "
+                          "kernel plus map_to_index compatibility route is installed")
+            else:
+                # Two 64-token blocks exercise the exact sm_100a inference
+                # specialization while keeping the one-time probe tiny.  The
+                # kernel predicate checks extension presence, CUDA capability,
+                # dtype/layout, head size, block size, and even block count
+                # without reading metadata tensor contents.
+                probe_query = torch.empty((1, 1, 128, self.head_size), device=device, dtype=torch.bfloat16)
+                probe_block_sizes = torch.full((2, ), 64, device=device, dtype=torch.int32)
+                reason = _sm100a_unavailable_reason(
+                    _sm100a,
+                    probe_query,
+                    probe_block_sizes,
+                    grad_mode=False,
+                )
+                enabled = reason is None
+
+        self._regional_compile_sm100a_enabled = enabled
+        # Keep this marker unset when preparation fails. Generic/training
+        # torch.compile must retain the established Triton attention route.
+        self._regional_compile_layer_idx = (torch.tensor(self.layer_idx, device=device, dtype=torch.int64)
+                                            if enabled else None)
+        if enabled:
+            route = ("native fastvideo-kernel mask entry" if callable(
+                getattr(_sm100a, "block_sparse_attn_sm100a_from_mask", None)) else
+                     "FastVideo compatibility mask adapter")
+            logger.info_once(f"VSA-H3 regional compile mask route: {route}")
+        if requested and reason is not None:
+            logger.warning_once(f"VSA-H3 regional compile is unavailable and will stay eager: {reason}")
+        return reason
 
     def tile(self, x: torch.Tensor, attn_metadata: MiniMaxH3VSAMetadata) -> torch.Tensor:
         """Scatter rows into the padded tile buffer (pad positions stay zero).
 
         The returned tensor aliases the builder-owned buffer; callers must
         consume it before the next ``tile()`` (both call sites in
-        ``forward()`` read it immediately).
+        ``forward()`` read it immediately). Odd tile-64 no-grad sm100a
+        requests carry one additional all-zero tile internally; metadata and
+        all observable outputs retain the logical geometry.
         """
         if x.shape[1] != attn_metadata.total_seq_length:
             raise ValueError(f"VSA-H3 metadata was built for sequence length {attn_metadata.total_seq_length}, "
                              f"got {x.shape[1]}. A non-packed sequence (e.g. the token refiner) is "
                              "routed to the VSA-H3 backend; exclude it from the supported backends.")
         n_tiles = attn_metadata.variable_block_sizes.numel()
-        target_shape = (x.shape[0], n_tiles * _TILE_ELEMS, x.shape[-2], x.shape[-1])
+        grad_mode = torch.is_grad_enabled() and x.requires_grad
+        compiling = torch.compiler.is_compiling()
+        regional_compiling = compiling and self._regional_compile_layer_idx is not None
+        if regional_compiling:
+            sm100a_requested = bool(self._regional_compile_sm100a_enabled)
+        elif compiling:
+            # Training/generic compile keeps the long-standing Triton route.
+            sm100a_requested = False
+        else:
+            sm100a_requested = os.environ.get(VSA_SM100A_ENV, "0") == "1"
+        needs_sm100a_pair = (attn_metadata.tile_elems == 64 and n_tiles % 2 != 0 and not grad_mode and sm100a_requested)
+        kernel_tiles = n_tiles + int(needs_sm100a_pair)
+        target_shape = (x.shape[0], kernel_tiles * attn_metadata.tile_elems, x.shape[-2], x.shape[-1])
 
-        # single scatter: untile_combined_index maps original row i to its
-        # padded slot, so this is exactly the inverse of postprocess_output
+        # ``untile_combined_index`` maps each packed row to a logical tile
+        # slot. Different geometries can share one transport shape; clear a
+        # reused allocation once when the mapping identity changes so no old
+        # valid row can survive as padding.
         holder = attn_metadata.tile_buf_holder
-        holder[0] = scatter_into_tile_buf(x, target_shape, attn_metadata.untile_combined_index, holder[0])
-        return holder[0]
+        if holder is None:
+            raise RuntimeError("VSA-H3 metadata has no builder-owned tile buffer holder")
+        buffer_matches = (holder.buffer is not None and holder.buffer.shape == target_shape
+                          and holder.buffer.dtype == x.dtype and holder.buffer.device == x.device)
+        if buffer_matches and holder.untile_geometry is not attn_metadata.untile_combined_index:
+            holder.buffer.zero_()
+        holder.buffer = scatter_into_tile_buf(x, target_shape, attn_metadata.untile_combined_index, holder.buffer)
+        holder.untile_geometry = attn_metadata.untile_combined_index
+        if needs_sm100a_pair:
+            # A prior even geometry can reuse this allocation and may have
+            # written the last tile as logical data.
+            holder.buffer[:, n_tiles * attn_metadata.tile_elems:].zero_()
+        return holder.buffer
 
     def preprocess_qkv(self, qkv: torch.Tensor, attn_metadata: MiniMaxH3VSAMetadata) -> torch.Tensor:
         return self.tile(qkv, attn_metadata)
@@ -281,24 +559,71 @@ class MiniMaxH3VSAImpl(AttentionImpl):
         gate_compress: torch.Tensor | None,
         attn_metadata: MiniMaxH3VSAMetadata,
     ) -> torch.Tensor:
-        if block_sparse_attn_256_bshd is None:
+        compiling = torch.compiler.is_compiling()
+        regional_compiling = compiling and self._regional_compile_layer_idx is not None
+
+        tile_elems = attn_metadata.tile_elems
+        if regional_compiling and tile_elems != 64:
+            raise RuntimeError("VSA-H3 regional fullgraph compile requires 64-token tiles; disable "
+                               "inference_torch_compile for tile-256/CuTe runs.")
+        if tile_elems == 64:
+            if block_sparse_attn_64_bhsd is None:
+                raise NotImplementedError("fastvideo_kernel.block_sparse_attn is not installed")
+        elif block_sparse_attn_256_bshd is None:
             raise NotImplementedError("fastvideo_kernel.block_sparse_attn_256 is not installed")
 
-        # probe-guided per-layer opt-out: diffuse layers run dense (all-True
-        # mask) while the rest keep the configured sparsity
-        layer_sparsity = 0.0 if self.layer_idx in attn_metadata.dense_layers else attn_metadata.VSA_sparsity
-        probe_dir = probe_enabled()
+        # Probe recording performs filesystem writes and host synchronizations,
+        # so the loader keeps probe-enabled runs eager. Avoid even reading that
+        # environment switch while Dynamo captures a regional full graph.
+        # The metadata always describes the trained logical geometry.
+        # ``tile()`` may append exactly one transport-only partner for an odd
+        # tile-64 sm100a call. Keep score selection and the gate branch on the
+        # logical prefix, and reject every other shape before a kernel sees it.
+        n_tiles = attn_metadata.variable_block_sizes.numel()
+        logical_seq_len = n_tiles * tile_elems
+        pair_pad_seq_len = logical_seq_len + tile_elems
+        pair_pad_is_valid = tile_elems == 64 and n_tiles % 2 != 0
+        allowed_seq_lengths = (logical_seq_len, pair_pad_seq_len) if pair_pad_is_valid else (logical_seq_len, )
+        if query.shape[1] not in allowed_seq_lengths:
+            expected = (f"the logical length {logical_seq_len} or one sm100a partner tile "
+                        f"({pair_pad_seq_len})" if pair_pad_is_valid else f"the logical length {logical_seq_len}")
+            raise ValueError(f"VSA-H3 tiled query has length {query.shape[1]}, expected {expected}.")
+        has_sm100a_pair = query.shape[1] == pair_pad_seq_len
+        for name, tensor in (("key", key), ("value", value)):
+            if tensor.shape[1] != query.shape[1]:
+                raise ValueError(f"VSA-H3 tiled {name} length {tensor.shape[1]} does not match query "
+                                 f"length {query.shape[1]}.")
+        if gate_compress is not None and gate_compress.shape[1] != query.shape[1]:
+            raise ValueError(f"VSA-H3 tiled gate length {gate_compress.shape[1]} does not match query "
+                             f"length {query.shape[1]}.")
+
+        logical_query = query[:, :logical_seq_len]
+        logical_key = key[:, :logical_seq_len]
+        logical_value = value[:, :logical_seq_len]
+        logical_gate = gate_compress[:, :logical_seq_len] if gate_compress is not None else None
+
+        # Probe-guided per-layer opt-out: diffuse layers run dense (all-True
+        # mask) while the rest keep the configured sparsity. During regional
+        # capture, keep the layer decision tensor-valued so the 50 block
+        # instances reuse one graph instead of specializing on layer_idx.
+        force_dense = None
+        if regional_compiling:
+            assert self._regional_compile_layer_idx is not None
+            force_dense = (attn_metadata.dense_layers_tensor == self._regional_compile_layer_idx).any()
+            layer_sparsity = attn_metadata.VSA_sparsity
+        else:
+            layer_sparsity = 0.0 if self.layer_idx in attn_metadata.dense_layers else attn_metadata.VSA_sparsity
+        probe_dir = None if compiling else probe_enabled()
 
         scores = None
         if layer_sparsity > 0.0 or gate_compress is not None or probe_dir is not None:
-            q_pooled = _pool_tiles(query, attn_metadata.variable_block_sizes)
-            k_pooled = _pool_tiles(key, attn_metadata.variable_block_sizes)
+            q_pooled = _pool_tiles(logical_query, attn_metadata.variable_block_sizes, tile_elems)
+            k_pooled = _pool_tiles(logical_key, attn_metadata.variable_block_sizes, tile_elems)
             scores = torch.matmul(q_pooled, k_pooled.transpose(-2, -1)) / (query.shape[-1]**0.5)
             if probe_dir is not None:
-                record_probe(probe_dir, self.layer_idx, query, key, scores, attn_metadata)
+                record_probe(probe_dir, self.layer_idx, logical_query, logical_key, scores, attn_metadata)
 
         if scores is None:
-            n_tiles = attn_metadata.variable_block_sizes.numel()
             mask = torch.ones(query.shape[0], query.shape[2], n_tiles, n_tiles, dtype=torch.bool, device=query.device)
         else:
             mask = _build_block_mask(
@@ -308,24 +633,134 @@ class MiniMaxH3VSAImpl(AttentionImpl):
                 layer_sparsity,
                 attn_metadata.exempt,
             )
+        if force_dense is not None:
+            # A scalar bool tensor broadcasts over the block map. This exactly
+            # preserves the eager dense-layer contract without a Python branch.
+            mask = mask | force_dense
 
-        out, _ = block_sparse_attn_256_bshd(query, key, value, mask, attn_metadata.variable_block_sizes)
+        if tile_elems == 64:
+            # Native 64-token path: the block map is already at the kernels'
+            # granularity. Both 64-token entries take BHSD ([B, H, S_pad, D]);
+            # mirror block_sparse_attn_256_bshd's Triton branch and transpose
+            # around the call.
+            q_bhsd = query.transpose(1, 2).contiguous()
+            k_bhsd = key.transpose(1, 2).contiguous()
+            v_bhsd = value.transpose(1, 2).contiguous()
 
-        if gate_compress is not None:
+            sm100a_mask = mask
+            sm100a_variable_block_sizes = attn_metadata.variable_block_sizes
+            if has_sm100a_pair:
+                # The synthetic tile is neither a logical query nor key. Its
+                # all-False row yields q2k_num=0, the all-False column keeps it
+                # out of real rows, and vbs=0 masks all of its key slots.
+                sm100a_mask = torch.nn.functional.pad(mask, (0, 1, 0, 1), value=False)
+                sm100a_variable_block_sizes = torch.nn.functional.pad(
+                    attn_metadata.variable_block_sizes,
+                    (0, 1),
+                    value=0,
+                )
+
+            # Opt-in sm_100a CUDA forward (upstream PR #1719 + per-q-tile
+            # q2k_num fix). Forward-only: grad-tracking calls stay on Triton
+            # so autograd keeps the Triton fwd+bwd pairing untouched. The
+            # kernel does return an LSE in Triton's M format, so a future
+            # fwd/bwd pairing is possible, but it is not built here.
+            grad_mode = torch.is_grad_enabled() and (query.requires_grad or key.requires_grad or value.requires_grad)
+            use_sm100a = False
+            if regional_compiling:
+                if self._regional_compile_sm100a_enabled is None:
+                    raise RuntimeError(
+                        "VSA-H3 sm_100a routing was not resolved before torch.compile; "
+                        "call prepare_for_regional_compile(device) on every MiniMaxH3VSAImpl after loading weights.")
+                # The preparation probe established module/device/kernel
+                # support.  Keep only static tensor/geometry facts here; no
+                # env access, device-capability query, or is_supported call may
+                # enter the Dynamo graph.
+                if not (self._regional_compile_sm100a_enabled and not grad_mode and q_bhsd.dtype == torch.bfloat16
+                        and q_bhsd.shape[-1] == 128 and sm100a_variable_block_sizes.numel() % 2 == 0):
+                    raise RuntimeError(
+                        "VSA-H3 regional fullgraph compile requires the prepared sm_100a BF16/head-128 route "
+                        "on a supported device; disable inference_torch_compile for this request.")
+                use_sm100a = True
+            elif not compiling and os.environ.get(VSA_SM100A_ENV, "0") == "1":
+                reason = _sm100a_unavailable_reason(_sm100a, q_bhsd, sm100a_variable_block_sizes, grad_mode)
+                if reason is None and map_to_index is None:
+                    reason = "fastvideo_kernel.triton_kernels.index (map_to_index) is not importable"
+                if reason is None:
+                    use_sm100a = True
+                else:
+                    logger.warning_once(f"{VSA_SM100A_ENV}=1 but falling back to the Triton-64 kernels: {reason}")
+
+            if use_sm100a:
+                # Regional preparation emits the compile-route receipt before
+                # capture. Logging from this branch would itself break a
+                # ``fullgraph=True`` forward.
+                if not compiling:
+                    logger.info_once("MiniMax-H3 VSA tile-64 forward: using the sm100a CUDA block-sparse kernel")
+                if regional_compiling:
+                    # The compile-safe wrapper keeps both Triton mask
+                    # compaction and the raw pybind launch behind one
+                    # fake-backed custom-op boundary.
+                    out_bhsd, _ = _sm100a_from_mask(
+                        q_bhsd,
+                        k_bhsd,
+                        v_bhsd,
+                        sm100a_mask,
+                        sm100a_variable_block_sizes,
+                    )
+                else:
+                    # Preserve the established eager/index-native route and
+                    # compatibility with older kernel wheels. Per-row counts
+                    # are non-uniform (prefix queries are dense; video queries
+                    # run prefix+top-k), which the fixed kernel supports.
+                    q2k_idx, q2k_num = map_to_index(sm100a_mask)
+                    out_bhsd, _ = _sm100a.block_sparse_attn_sm100a(
+                        q_bhsd,
+                        k_bhsd,
+                        v_bhsd,
+                        q2k_idx,
+                        q2k_num,
+                        sm100a_variable_block_sizes.to(torch.int32),
+                        need_lse=False,
+                    )
+            else:
+                if has_sm100a_pair:
+                    q_bhsd = q_bhsd[:, :, :logical_seq_len].contiguous()
+                    k_bhsd = k_bhsd[:, :, :logical_seq_len].contiguous()
+                    v_bhsd = v_bhsd[:, :, :logical_seq_len].contiguous()
+                out_bhsd, _ = block_sparse_attn_64_bhsd(
+                    q_bhsd,
+                    k_bhsd,
+                    v_bhsd,
+                    mask,
+                    attn_metadata.variable_block_sizes,
+                )
+            if has_sm100a_pair and use_sm100a:
+                out_bhsd = out_bhsd[:, :, :logical_seq_len]
+            out = out_bhsd.transpose(1, 2).contiguous()
+        else:
+            out, _ = block_sparse_attn_256_bshd(
+                logical_query,
+                logical_key,
+                logical_value,
+                mask,
+                attn_metadata.variable_block_sizes,
+            )
+
+        if logical_gate is not None:
             # Wan-style compression branch: dense attention over pooled tiles,
             # broadcast to each tile's rows, scaled by the learned gate
             # (zero-initialized for H3 => branch contributes nothing until
             # finetuned; the model layer skips it entirely for all-zero gates).
-            v_pooled = _pool_tiles(value, attn_metadata.variable_block_sizes)
+            v_pooled = _pool_tiles(logical_value, attn_metadata.variable_block_sizes, tile_elems)
             out_c = torch.matmul(torch.softmax(scores, dim=-1), v_pooled)  # [B, H, n_tiles, D]
             out_c = out_c.permute(0, 2, 1, 3).to(out.dtype)  # [B, n_tiles, H, D]
             batch, seq_len, heads, dim = out.shape
-            n_tiles = attn_metadata.variable_block_sizes.numel()
             # Out-of-place: on the CuTe backend ``out`` is the tensor FA4's
             # autograd node saved for its backward, so an in-place add here
             # bumps its version counter and backward dies with "one of the
             # variables needed for gradient computation has been modified".
-            out_tiled = out.view(batch, n_tiles, _TILE_ELEMS, heads, dim)
-            gate_tiled = gate_compress.view(batch, n_tiles, _TILE_ELEMS, heads, dim)
+            out_tiled = out.view(batch, n_tiles, tile_elems, heads, dim)
+            gate_tiled = logical_gate.view(batch, n_tiles, tile_elems, heads, dim)
             out = (out_tiled + out_c.unsqueeze(2) * gate_tiled).view(batch, seq_len, heads, dim)
         return out

fastvideo/attention/backends/video_sparse_attn_h3_probe.py

@@ -60,7 +60,7 @@ def record_probe(
     gen = torch.Generator(device="cpu").manual_seed(step * 1000 + layer)
     # sample among video rows in the PADDED/tiled domain that are non-pad
     from fastvideo.attention.backends.video_sparse_attn_h3 import token_tile_and_valid
-    token_tile, token_valid = token_tile_and_valid(attn_metadata.variable_block_sizes)
+    token_tile, token_valid = token_tile_and_valid(attn_metadata.variable_block_sizes, attn_metadata.tile_elems)
     video_rows = torch.nonzero((token_tile >= P) & token_valid, as_tuple=False).flatten()
     idx = video_rows[torch.randint(0, video_rows.numel(), (_TRUE_ROWS, ), generator=gen).to(query.device)]