[bugfix]: select architecture-specific Blackwell kernel targets

- Detect GB200 and GB300 as 10.0a and 10.3a - Build future GB200 CI images with 10.0a - Add regression tests for detection, cache keys, and build arguments
[feat]: add selectable Modal and GB200 Kubernetes GPU CI
2026-10-05 15:42:48 -07:00 · 2026-10-01 14:44:14 -07:00 · 2026-09-29 16:57:22 -07:00 · 2026-09-29 14:36:16 -07:00 · 2026-09-29 13:03:18 -07:00 · 2026-09-29 12:50:45 -07:00
277 changed files with 13995 additions and 2592 deletions
@@ -0,0 +1,76 @@
+---
+name: env-var-conventions
+description: Add, read, rename, or remove an environment variable in FastVideo, or change the environment-variable policy. Use before touching fastvideo/envs.py, os.environ, os.getenv, or monkeypatch.setenv in fastvideo/, and when fastvideo/tests/contract/test_env_policy.py fails.
+---
+
+# Environment Variable Conventions
+
+## Purpose
+
+FastVideo registers its environment variables as typed fields in
+`fastvideo/envs.py`. The policy that governs them is
+`docs/contributing/env_vars.md`, and the contract test
+`fastvideo/tests/contract/test_env_policy.py` enforces the policy in the unit
+CI lane. This skill routes an environment-variable change through that policy.
+The policy doc is the single source of the rules; read it instead of relying
+on a summary here.
+
+## Prerequisites
+
+- Read `docs/contributing/env_vars.md` in full.
+- Decide whether the setting belongs in an environment variable or an argument
+  (rule 5 in the policy doc). Settings that users change per deployment are
+  arguments; add them through `fastvideo/fastvideo_args.py` instead.
+
+## Inputs
+
+| Parameter  | Required | Description                                                    |
+| ---------- | -------- | -------------------------------------------------------------- |
+| `change`   | Yes      | Add, read, rename, or remove a variable, or change the policy. |
+| `variable` | Yes      | The variable name, with the `FASTVIDEO_` prefix.               |
+
+## Steps
+
+1. **Declare or edit the variable in `fastvideo/envs.py`.**
+   - Pick the field type and category that the policy doc lists.
+   - Write a description that states what the variable does and its units.
+   - To rename, keep the old name in `deprecated_names`. To remove, add the
+     name to `DEPRECATED_VARIABLES`. Update the uses in `examples/`,
+     `scripts/`, `docs/`, `apps/`, and the tests.
+2. **Read the variable with `envs.NAME.get()` inside a function.**
+   - In tests, change the value with `envs.NAME.override(value)`.
+   - Do not call `os.environ`, `os.getenv`, or `monkeypatch.setenv` for a
+     FastVideo variable.
+   - To set a variable that another tool reads, call `envs.set_external`,
+     `envs.setdefault_external`, or `envs.unset_external`.
+3. **Regenerate the table in the policy doc.**
+   - Run `python fastvideo/tests/contract/test_env_policy.py`.
+4. **Run the contract test.**
+   - Run `pytest fastvideo/tests/contract/test_env_policy.py`.
+   - When the test reports a fixed known violation, delete or lower its entry
+     in `KNOWN_VIOLATIONS`. Never add an entry to `KNOWN_VIOLATIONS`.
+5. **When the policy itself changes, update the policy doc and the contract
+   test in the same pull request.**
+   - The rules in `docs/contributing/env_vars.md`, the checks and allowlist in
+     `fastvideo/tests/contract/test_env_policy.py`, and this skill must agree.
+
+## Outputs
+
+- A registry entry in `fastvideo/envs.py` and call sites that use
+  `envs.NAME.get()`.
+- A regenerated table in `docs/contributing/env_vars.md`.
+- A passing `fastvideo/tests/contract/test_env_policy.py`.
+
+## Example Usage
+
+```
+Add a FASTVIDEO_DEBUG_MY_STAGE switch that logs MyStage inputs.
+```
+
+## References
+
+- `docs/contributing/env_vars.md`: the policy, the field types, and the
+  violation kinds that the contract test reports.
+- `fastvideo/envs.py`: the registry.
+- `fastvideo/tests/contract/test_env_policy.py`: the contract test,
+  `EXTERNAL_ALLOWLIST`, and `KNOWN_VIOLATIONS`.
@@ -22,6 +22,12 @@ else
  export PERF_UPLOAD_POLICY=never
 fi

+# Alternate GPU backends compare against references without publishing records.
+# Their worker has read-only Hub credentials; publication is an operator task.
+if [ "${FASTVIDEO_CI_LOCAL_ONLY:-0}" = 1 ]; then
+  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 &
@@ -16,6 +16,7 @@ exec pytest \
  ./fastvideo/tests/worker/ \
  ./fastvideo/tests/training/test_trackers.py \
  ./fastvideo/tests/attention/test_sdpa_metadata_mask_contract.py \
+  ./fastvideo/tests/attention/test_vsa_h3_tile_grad_safety.py \
  ./fastvideo/tests/modal/test_kernel_build_cache.py \
  ./fastvideo/tests/modal/test_pr_test.py \
  ./fastvideo/tests/modal/test_ssim_test.py \
@@ -11,6 +11,7 @@ jobs:
    if: >-
      github.event.context == 'direct-test-completed'
      && github.event.state == 'success'
+      && (vars.CI_GPU_BACKEND == '' || vars.CI_GPU_BACKEND == 'slurm')
    runs-on: ubuntu-latest
    steps:
      - name: Check and update aggregate status
@@ -0,0 +1,62 @@
+name: Promote Selected GPU Backend Status
+
+on:
+  status:
+
+permissions:
+  statuses: write
+
+concurrency:
+  group: gpu-ci-status-${{ github.event.sha }}-${{ vars.CI_GPU_BACKEND }}
+  cancel-in-progress: false
+
+jobs:
+  promote:
+    if: >-
+      (vars.CI_GPU_BACKEND == 'modal' || vars.CI_GPU_BACKEND == 'vllm')
+      && (github.event.context == format('gpu-ci/{0}/fastcheck-passed', vars.CI_GPU_BACKEND)
+          || github.event.context == format('gpu-ci/{0}/full-suite-passed', vars.CI_GPU_BACKEND))
+    runs-on: ubuntu-latest
+    env:
+      SELECTED_BACKEND: ${{ vars.CI_GPU_BACKEND }}
+    steps:
+      - name: Mirror the selected backend's latest suite results
+        uses: actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea # v7.0.1
+        with:
+          script: |
+            const backend = process.env.SELECTED_BACKEND;
+            if (!['modal', 'vllm'].includes(backend)) {
+              throw new Error('Unsupported selected GPU backend');
+            }
+            const sha = context.payload.sha;
+            // Read current state after entering the serialized workflow. A
+            // delayed event must not overwrite a newer failure with success.
+            const statuses = await github.paginate(github.rest.repos.listCommitStatusesForRef, {
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              ref: sha,
+              per_page: 100,
+            });
+            for (const suffix of ['fastcheck-passed', 'full-suite-passed']) {
+              const sourceContext = `gpu-ci/${backend}/${suffix}`;
+              const matches = statuses.filter(status => status.context === sourceContext);
+              matches.sort((a, b) =>
+                Date.parse(b.updated_at) - Date.parse(a.updated_at) || b.id - a.id
+              );
+              const latest = matches[0];
+              const state = latest ? latest.state : 'pending';
+              if (!['pending', 'success', 'failure', 'error'].includes(state)) {
+                throw new Error(`Unsupported status state for ${sourceContext}`);
+              }
+              await github.rest.repos.createCommitStatus({
+                owner: context.repo.owner,
+                repo: context.repo.repo,
+                sha,
+                context: suffix,
+                state,
+                description: latest
+                  ? `${backend} ${suffix}: ${state}`
+                  : `Waiting for ${backend} ${suffix}`,
+                ...(latest && latest.target_url ? {target_url: latest.target_url} : {}),
+              });
+            }
@@ -92,6 +92,7 @@ jobs:
            fastvideo/tests/mlx/test_frame_upsample.py \
            fastvideo/tests/mlx/test_mlx_fast_spatial.py \
            fastvideo/tests/mlx/test_mlx_refine.py \
+            fastvideo/tests/mlx/test_mlx_prompt_enhance.py \
            fastvideo/tests/mlx/test_mlx_prompt_to_video_decode.py \
            fastvideo/tests/mlx/test_mlx_wan22_prompt_cache_fingerprint.py \
            fastvideo/tests/mlx/test_wan22_sample.py \
@@ -156,6 +157,7 @@ jobs:
            fastvideo/tests/mlx/test_frame_upsample.py \
            fastvideo/tests/mlx/test_mlx_fast_spatial.py \
            fastvideo/tests/mlx/test_mlx_refine.py \
+            fastvideo/tests/mlx/test_mlx_prompt_enhance.py \
            fastvideo/tests/mlx/test_mlx_prompt_to_video_decode.py \
            fastvideo/tests/mlx/test_mlx_wan22_prompt_cache_fingerprint.py \
            fastvideo/tests/mlx/test_wan22_sample.py \
@@ -203,7 +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 targets sm_100 rather than sm_121.
+  # The CI runner is ARM64 like DGX Spark, but targets sm_100a rather than sm_121.
+  # The architecture-specific target includes the GB200 VSA CUDA extensions.
  # 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:
@@ -219,7 +220,7 @@ jobs:
        PYTHON_VERSION=3.12
        CUDA_VERSION=13.0.0
        UV_TORCH_BACKEND=cu130
-        TORCH_CUDA_ARCH_LIST=10.0
+        TORCH_CUDA_ARCH_LIST=10.0a
        CMAKE_BUILD_PARALLEL_LEVEL=1
        FLASH_ATTN_WHEEL_TAG=cu130torch2.12
        FLASH_ATTN_WHEEL_RELEASE_ARM64=https://github.com/mjun0812/flash-attention-prebuild-wheels/releases/download/v0.9.22
@@ -176,11 +176,11 @@ jobs:
          # the main extension for the full arch list. CMAKE_BUILD_PARALLEL_LEVEL caps
          # Ninja so heavy CUTLASS/TK template TUs don't OOM the 16 GB runner (exit 143).
          if [ "${{ matrix.platform.arch }}" = "aarch64" ]; then
-            export TORCH_CUDA_ARCH_LIST="10.0a;12.0a"
+            export TORCH_CUDA_ARCH_LIST="10.0a;10.3a;12.0a"
            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;10.0a;12.0a"
+            export TORCH_CUDA_ARCH_LIST="9.0a;10.0a;10.3a;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
@@ -66,14 +66,18 @@ Local guidance lives next to the code. Read the in-scope file before editing:
 | `fastvideo/AGENTS.md` | Core package map, public API, registry-driven model dispatch |
 | `fastvideo/configs/AGENTS.md` | Arch + pipeline config dataclasses, `param_names_mapping` |
 | `fastvideo/models/AGENTS.md` | DiT / VAE / encoder / scheduler / loader layout (pre-commit excluded) |
+| `fastvideo/models/wan/AGENTS.md` | Wan family-local transformers, VAE, configs, and the SP sharding invariant |
 | `fastvideo/layers/AGENTS.md` | Tensor-parallel linear/attention layer rules for ports |
 | `fastvideo/attention/AGENTS.md` | Backend registry + env-var override |
 | `fastvideo/pipelines/AGENTS.md` | Stage ABC, `basic/<model>/`, `preprocess/`, presets |
+| `fastvideo/pipelines/basic/wan/AGENTS.md` | Wan sampling stages, first-frame conditioning, DMD/causal boundaries |
+| `fastvideo/pipelines/basic/magi_human/AGENTS.md` | MagiHuman umbrella repo, lazy-loaded components, packing invariants |
 | `fastvideo/training/AGENTS.md` | Legacy monolithic pipelines (frozen for existing models) |
 | `fastvideo/train/AGENTS.md` | New modular trainer (methods × models × callbacks, YAML) |
 | `fastvideo/tests/AGENTS.md` | Test taxonomy, conftest, pre-commit-excluded path |
 | `fastvideo/tests/ssim/AGENTS.md` | GPU SSIM regression authoring + reference video sync |
 | `scripts/checkpoint_conversion/AGENTS.md` | Adding a converter for a new HF/official checkpoint |
+| `apps/dreamverse/AGENTS.md` | DreamVerse app structure and conventions |

 ## Critical: Two Training Stacks Coexist

@@ -9,9 +9,10 @@
 **FastVideo is a unified post-training and real-time inference framework for accelerated video generation.**

 ## NEWS
+- `2026/09/15`: Release [FastH3 8-Step V2](https://huggingface.co/FastVideo/FastVideo-FastH3-8-Step-V2), an eight-forward data-free DMD2 checkpoint distilled from MiniMax-H3 with 80% Video Sparse Attention. Run it with `examples/inference/basic/basic_fasth3_8step.py` or the [FastH3 8-Step V2 recipe](https://haoailab.com/FastVideo/cookbook/minimax-h3/).
 - `2026/09/01`: FastH3 now runs locally on Apple Silicon through MLX and on NVIDIA DGX Spark through CUDA 13, including two-Spark inference. Follow the [FastH3 recipes](https://haoailab.com/FastVideo/cookbook/minimax-h3/) and read the [Blog](https://haoailab.com/blogs/fasth3-local/).
 - `2026/08/27`: [FastH3 Preview v1](https://haoailab.com/blogs/fasth3-preview/) is an open-weight 4-step sparse-distilled MiniMax-H3 model for synchronized video-and-audio generation, developed in collaboration with [Nuva Lab](https://nuvalab.ai/) and the [NVIDIA FastGen team](https://github.com/NVlabs/FastGen). Download the recommended [VSA / Data-Free weights](https://huggingface.co/FastVideo/FastVideo-FastH3-4-step-Preview-v1-VSA-DataFree), or see the [full FastH3 collection](https://huggingface.co/collections/FastVideo/fastvideo-fasth3).
- `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/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 [MLX install guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mlx/) 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/).
@@ -63,13 +64,12 @@ 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/).
+[MLX install guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mlx/).

-> **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/).
+> **On an Apple Silicon Mac?** Install with `uv pip install -e '.[mlx]'` from
+> a clone, then pick a recipe in the
+> [cookbook](https://haoailab.com/FastVideo/cookbook/). See the
+> [MLX install guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mlx/).

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

@@ -87,7 +87,7 @@ Install FastVideo (https://github.com/hao-ai-lab/FastVideo) into a fresh uv virt
   https://hao-ai-lab.github.io/FastVideo/getting_started/installation/):
     - NVIDIA GPU, x86_64                         -> docs/getting_started/installation/gpu.md
     - NVIDIA DGX Spark / GB10, aarch64, CUDA 13  -> docs/getting_started/installation/spark.md
-     - Apple Silicon, macOS                       -> docs/getting_started/installation/mps.md
+     - Apple Silicon, macOS                       -> docs/getting_started/installation/mlx.md
 3. Use uv for every step. If a command fails, debug it and tell me what you changed.
 4. Verify the result:
     python -c "import fastvideo, torch; print('cuda', torch.cuda.is_available())"
@@ -331,11 +331,11 @@ def test_rewrite_prompt_sequence_accepts_numbered_prose_output():
    ]


-def test_enhance_prompt_prefers_cerebras_before_groq_fallback():
+def test_enhance_prompt_uses_groq_when_it_returns_first():
    enhancer = _build_staged_enhancer(
        cerebras_payload=_chat_payload_with_content('{"prompt":"Cerebras prompt"}'),
        groq_payload=_chat_payload_with_content('{"prompt":"Groq prompt"}'),
-        cerebras_delay_s=0.01,
+        cerebras_delay_s=0.08,
        groq_delay_s=0.01,
    )

@@ -346,12 +346,12 @@ def test_enhance_prompt_prefers_cerebras_before_groq_fallback():

    assert result.fallback_used is False
    assert result.error is None
-    assert result.provider == "cerebras"
+    assert result.provider == "groq"
    assert result.model == "gpt-test"
-    assert result.prompt == "Cerebras prompt"
+    assert result.prompt == "Groq prompt"
    assert enhancer.get_provider_success_counts() == {
-        "cerebras": 1,
-        "groq": 0,
+        "cerebras": 0,
+        "groq": 1,
    }


@@ -0,0 +1,63 @@
+import { expect, test } from '@playwright/test';
+
+import { skipWithoutMock } from './helpers';
+
+test.describe('create job interactions', () => {
+  skipWithoutMock();
+
+  for (const jobType of ['inference', 'finetuning', 'distillation']) {
+    test(`${jobType} remains interactive after repeated dialog dismissals`, async ({ page }) => {
+      await page.goto(`/${jobType}`);
+      const trigger = page.getByRole('button', { name: 'Create Job', exact: true });
+      const dialog = page.getByRole('dialog');
+
+      // Exercise both dismissal paths and reopen without reloading the page.
+      for (const closeWithEscape of [false, true]) {
+        await trigger.click();
+        await page.getByRole('menuitem').first().click();
+        await expect(dialog).toBeVisible();
+        if (closeWithEscape) {
+          await page.keyboard.press('Escape');
+        } else {
+          await dialog.getByRole('button', { name: 'Close', exact: true }).click();
+        }
+        await expect(dialog).toBeHidden();
+        await expect(page.locator('body')).toHaveCSS('pointer-events', 'auto');
+        await expect(trigger).toBeFocused();
+      }
+
+      await page.getByRole('link', { name: 'Datasets', exact: true }).click();
+      await expect(page).toHaveURL(/\/datasets$/);
+    });
+  }
+
+  test('preserves keyboard menu dismissal and dialog focus trapping', async ({ page }) => {
+    await page.goto('/inference');
+    const trigger = page.getByRole('button', { name: 'Create Job', exact: true });
+    await trigger.focus();
+    await page.keyboard.press('Enter');
+    const firstItem = page.getByRole('menuitem').first();
+    await expect(firstItem).toBeFocused();
+    await page.keyboard.press('Escape');
+    await expect(page.getByRole('menu')).toBeHidden();
+    await expect(trigger).toBeFocused();
+    await expect(page.locator('body')).toHaveCSS('pointer-events', 'auto');
+
+    await page.keyboard.press('Enter');
+    await expect(firstItem).toBeFocused();
+    await page.keyboard.press('Enter');
+    const dialog = page.getByRole('dialog');
+    await expect(dialog).toBeVisible();
+    await expect(dialog.getByLabel('Name (optional)')).toBeFocused();
+
+    // Shift+Tab from the first field wraps to Close, then Tab wraps back.
+    await page.keyboard.press('Shift+Tab');
+    await expect(dialog.getByRole('button', { name: 'Close', exact: true })).toBeFocused();
+    await page.keyboard.press('Tab');
+    await expect(dialog.getByLabel('Name (optional)')).toBeFocused();
+    await page.keyboard.press('Escape');
+    await expect(dialog).toBeHidden();
+    await expect(trigger).toBeFocused();
+    await expect(page.locator('body')).toHaveCSS('pointer-events', 'auto');
+  });
+});
@@ -1,6 +1,6 @@
 import { expect, test } from '@playwright/test';

-import { skipWithoutMock } from './helpers';
+import { API_BASE, skipWithoutMock } from './helpers';

 /**
 * Create-job flow: open the Create Job modal on /inference, fill the prompt
@@ -10,7 +10,8 @@ import { skipWithoutMock } from './helpers';
 test.describe('create inference job', () => {
  skipWithoutMock();

-  test('creates a T2V job and shows it in the queue', async ({ page }) => {
+  test('creates a T2V job and starts it without refreshing', async ({ page, request }) => {
+    await request.put(`${API_BASE}/settings`, { data: { autoStartJob: false } });
    await page.goto('/inference');

    // The trigger opens a real menu on click, so this path works for touch,
@@ -38,5 +39,20 @@ test.describe('create inference job', () => {
    // Modal closes and the queue refreshes with the newly created job.
    await expect(dialog).toBeHidden();
    await expect(page.getByText(prompt)).toBeVisible();
+    await expect(page.locator('body')).toHaveCSS('pointer-events', 'auto');
+
+    const card = page.getByRole('article').filter({ hasText: prompt });
+    await expect(card.getByText('pending', { exact: true })).toBeVisible();
+    const started = page.waitForResponse((response) =>
+      response.url().startsWith(`${API_BASE}/jobs/`) &&
+      response.url().endsWith('/start') &&
+      response.request().method() === 'POST',
+    );
+    await card.getByRole('button', { name: 'Start', exact: true }).click();
+    expect((await started).ok()).toBe(true);
+    await expect(card.getByText('running', { exact: true })).toBeVisible();
+
+    await page.getByRole('link', { name: 'Datasets', exact: true }).click();
+    await expect(page).toHaveURL(/\/datasets$/);
  });
 });
@@ -17,20 +17,11 @@
 		"start:all": "concurrently --kill-others-on-fail \"npm:start:api\" \"npm:start:web\""
 	},
 	"dependencies": {
-		"@radix-ui/react-dialog": "^1.1.0",
-		"@radix-ui/react-dropdown-menu": "^2.1.24",
-		"@radix-ui/react-label": "^2.1.8",
-		"@radix-ui/react-scroll-area": "^1.2.10",
-		"@radix-ui/react-select": "^2.2.6",
-		"@radix-ui/react-separator": "^1.1.8",
-		"@radix-ui/react-slider": "^1.2.0",
-		"@radix-ui/react-slot": "^1.2.4",
-		"@radix-ui/react-switch": "^1.1.0",
-		"@radix-ui/react-tabs": "^1.1.0",
 		"class-variance-authority": "^0.7.1",
 		"clsx": "^2.1.1",
 		"lucide-react": "^0.577.0",
 		"next": "15.5.18",
+		"radix-ui": "^1.6.7",
 		"react": "^19.1.0",
 		"react-dom": "^19.1.0",
 		"sonner": "^2.0.7",
@@ -1,24 +1,26 @@
-import { render, screen } from '@testing-library/react';
+import { render, screen, waitFor, within } from '@testing-library/react';
 import userEvent from '@testing-library/user-event';
-import { describe, expect, it, vi } from 'vitest';
+import { beforeEach, describe, expect, it, vi } from 'vitest';

 import CreateJobButton from './CreateJobButton';
+import { getDatasets, getModels } from '@/lib/api';

-vi.mock('./CreateJobModal', () => ({
-  default: ({
-    isOpen,
-    workloadType,
-  }: {
-    isOpen: boolean;
-    workloadType: string;
-  }) =>
-    isOpen ? (
-      <div role="dialog" data-workload-type={workloadType}>
-        Create job form
-      </div>
-    ) : null,
+vi.mock('@/lib/api', () => ({
+  createJob: vi.fn(),
+  getModels: vi.fn(),
+  getDatasets: vi.fn(),
+  uploadImage: vi.fn(),
+  getSettings: vi.fn(),
+  updateSettings: vi.fn(),
 }));

+beforeEach(() => {
+  vi.mocked(getModels).mockResolvedValue([
+    { id: 'wan/t2v-1.3b', label: 'Wan T2V' },
+  ]);
+  vi.mocked(getDatasets).mockResolvedValue([]);
+});
+
 describe('CreateJobButton', () => {
  it('opens the workload menu on click and selects an item', async () => {
    const user = userEvent.setup();
@@ -27,10 +29,9 @@ describe('CreateJobButton', () => {
    await user.click(screen.getByRole('button', { name: 'Create Job' }));
    await user.click(screen.getByRole('menuitem', { name: /I2V/i }));

-    expect(screen.getByRole('dialog')).toHaveAttribute(
-      'data-workload-type',
-      'i2v',
-    );
+    expect(
+      screen.getByRole('dialog', { name: 'New Inference Job (I2V)' }),
+    ).toBeInTheDocument();
  });

  it('opens and operates the workload menu from the keyboard', async () => {
@@ -45,9 +46,42 @@ describe('CreateJobButton', () => {
    expect(firstItem).toHaveFocus();
    await user.keyboard('{Enter}');

-    expect(screen.getByRole('dialog')).toHaveAttribute(
-      'data-workload-type',
-      't2v',
+    expect(
+      screen.getByRole('dialog', { name: 'New Inference Job (T2V)' }),
+    ).toBeInTheDocument();
+    await user.keyboard('{Escape}');
+    await waitFor(() =>
+      expect(screen.queryByRole('dialog')).not.toBeInTheDocument(),
    );
+    await waitFor(() =>
+      expect(document.body.style.pointerEvents).not.toBe('none'),
+    );
+    expect(trigger).toHaveFocus();
  });
+
+  it.each(['inference', 'finetuning', 'distillation'] as const)(
+    'restores page interaction after closing the real %s dialog',
+    async (jobType) => {
+      const user = userEvent.setup();
+      render(<CreateJobButton jobType={jobType} />);
+      const trigger = screen.getByRole('button', { name: 'Create Job' });
+
+      // Keep the real Dialog mounted: mocking it hides conflicting Radix layers.
+      for (let attempt = 0; attempt < 2; attempt++) {
+        await user.click(trigger);
+        await user.click(screen.getAllByRole('menuitem')[0]);
+        const dialog = screen.getByRole('dialog');
+        await user.click(
+          within(dialog).getByRole('button', { name: 'Close' }),
+        );
+        await waitFor(() =>
+          expect(screen.queryByRole('dialog')).not.toBeInTheDocument(),
+        );
+        await waitFor(() =>
+          expect(document.body.style.pointerEvents).not.toBe('none'),
+        );
+        expect(trigger).toHaveFocus();
+      }
+    },
+  );
 });
@@ -2,7 +2,7 @@

 import * as React from 'react';
 import { ChevronDown } from 'lucide-react';
-import * as DropdownMenu from '@radix-ui/react-dropdown-menu';
+import { DropdownMenu } from 'radix-ui';

 import CreateJobModal from '@/components/jobs/CreateJobModal';
 import { Button } from '@/components/ui/button';
@@ -16,6 +16,7 @@ interface CreateJobButtonProps {

 export default function CreateJobButton({ jobType }: CreateJobButtonProps) {
  const options = WORKLOAD_OPTIONS[jobType] ?? [];
+  const triggerRef = React.useRef<HTMLButtonElement>(null);

  const [modalOpen, setModalOpen] = React.useState(false);
  const [workloadType, setWorkloadType] = React.useState(
@@ -36,7 +37,7 @@ export default function CreateJobButton({ jobType }: CreateJobButtonProps) {
    <>
      <DropdownMenu.Root>
        <DropdownMenu.Trigger asChild>
-          <Button type="button" className="gap-1.5">
+          <Button ref={triggerRef} type="button" className="gap-1.5">
            Create Job
            <ChevronDown className="size-3.5 opacity-85" aria-hidden />
          </Button>
@@ -66,6 +67,11 @@ export default function CreateJobButton({ jobType }: CreateJobButtonProps) {
      <CreateJobModal
        isOpen={modalOpen}
        onClose={() => setModalOpen(false)}
+        onCloseAutoFocus={(event) => {
+          // This dialog opens from a menu item, so it has no DialogTrigger.
+          event.preventDefault();
+          triggerRef.current?.focus();
+        }}
        onSuccess={handleSuccess}
        jobType={jobType}
        workloadType={workloadType}
@@ -55,6 +55,9 @@ import { jobToFormFields, type JobLike } from '@/lib/jobToFields';
 export interface CreateJobModalProps {
  isOpen: boolean;
  onClose: () => void;
+  onCloseAutoFocus?: React.ComponentProps<
+    typeof DialogContent
+  >['onCloseAutoFocus'];
  onSuccess: () => void;
  jobType: JobType;
  workloadType: string;
@@ -67,6 +70,7 @@ export interface CreateJobModalProps {
 export default function CreateJobModal({
  isOpen,
  onClose,
+  onCloseAutoFocus,
  onSuccess,
  jobType,
  workloadType,
@@ -644,6 +648,7 @@ export default function CreateJobModal({
    >
      <DialogContent
        className="max-h-[90vh] w-[90vw] max-w-[850px] overflow-y-auto"
+        onCloseAutoFocus={onCloseAutoFocus}
        onEscapeKeyDown={(e) => {
          if (isSubmitting) e.preventDefault();
        }}
@@ -1,5 +1,7 @@
+import * as React from 'react';
 import { render, screen } from '@testing-library/react';
-import { describe, expect, it } from 'vitest';
+import userEvent from '@testing-library/user-event';
+import { describe, expect, it, vi } from 'vitest';

 import { Button } from './button';
 import { Input } from './input';
@@ -8,6 +10,24 @@ import { Slider } from './slider';
 import { Switch } from './switch';

 describe('shared control accessibility', () => {
+  it('forwards refs and click handlers to the asChild button', async () => {
+    const user = userEvent.setup();
+    const ref = React.createRef<HTMLButtonElement>();
+    const onClick = vi.fn();
+    render(
+      <Button asChild ref={ref} onClick={onClick}>
+        <button type="button">Slotted action</button>
+      </Button>,
+    );
+
+    const button = screen.getByRole('button', { name: 'Slotted action' });
+    expect(screen.getAllByRole('button')).toHaveLength(1);
+    expect(ref.current).toBe(button);
+    await user.click(button);
+    expect(onClick).toHaveBeenCalledTimes(1);
+    expect(button).toHaveFocus();
+  });
+
  it('keeps button, input, and select targets at least 44px tall', () => {
    render(
      <>
@@ -1,7 +1,7 @@
 "use client";

 import * as React from "react";
-import { Slot } from "@radix-ui/react-slot";
+import { Slot } from "radix-ui";
 import { cva, type VariantProps } from "class-variance-authority";

 import { cn } from "@/lib/utils";
@@ -37,7 +37,7 @@ export interface ButtonProps extends React.ButtonHTMLAttributes<HTMLButtonElemen
 }

 const Button = React.forwardRef<HTMLButtonElement, ButtonProps>(({ className, variant, size, asChild = false, ...props }, ref) => {
-	const Comp = asChild ? Slot : "button";
+	const Comp = asChild ? Slot.Root : "button";
 	return <Comp className={cn(buttonVariants({ variant, size, className }))} ref={ref} {...props} />;
 });
 Button.displayName = "Button";
@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as DialogPrimitive from '@radix-ui/react-dialog';
+import { Dialog as DialogPrimitive } from 'radix-ui';
 import { X } from 'lucide-react';

 import { cn } from '@/lib/utils';
@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as LabelPrimitive from '@radix-ui/react-label';
+import { Label as LabelPrimitive } from 'radix-ui';
 import { cva, type VariantProps } from 'class-variance-authority';

 import { cn } from '@/lib/utils';
@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as ScrollAreaPrimitive from '@radix-ui/react-scroll-area';
+import { ScrollArea as ScrollAreaPrimitive } from 'radix-ui';

 import { cn } from '@/lib/utils';

@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as SelectPrimitive from '@radix-ui/react-select';
+import { Select as SelectPrimitive } from 'radix-ui';
 import { Check, ChevronDown, ChevronUp } from 'lucide-react';

 import { cn } from '@/lib/utils';
@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as SeparatorPrimitive from '@radix-ui/react-separator';
+import { Separator as SeparatorPrimitive } from 'radix-ui';

 import { cn } from '@/lib/utils';

@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as SliderPrimitive from '@radix-ui/react-slider';
+import { Slider as SliderPrimitive } from 'radix-ui';

 import { cn } from '@/lib/utils';

@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as SwitchPrimitives from '@radix-ui/react-switch';
+import { Switch as SwitchPrimitives } from 'radix-ui';

 import { cn } from '@/lib/utils';

@@ -1,7 +1,7 @@
 'use client';

 import * as React from 'react';
-import * as TabsPrimitive from '@radix-ui/react-tabs';
+import { Tabs as TabsPrimitive } from 'radix-ui';

 import { cn } from '@/lib/utils';

@@ -1,5 +1,5 @@
 {
-  "version": 8,
+  "version": 11,
  "recipes": [
    {
      "id": "fastwan21-t2v",
@@ -445,13 +445,13 @@
    {
      "id": "fasth3-preview-cuda",
      "group": "fasth3-preview",
-      "group_label": "FastH3 Preview",
+      "group_label": "FastH3 V1",
      "group_task": "4-step text to video + audio",
      "family": "minimax_h3",
      "stage": "inference",
      "task": "Few-step text to video (with audio)",
-      "label": "FastH3 Preview on CUDA",
-      "summary": "Run the DMD2-distilled FastH3 Preview with four DiT forwards, trained H3 sparse attention, compiled decode, and synchronized audio.",
+      "label": "FastH3 V1 on CUDA",
+      "summary": "Run FastH3 V1 with four DiT forwards, trained H3 sparse attention, compiled decode, and synchronized audio.",
      "model": "FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2",
      "source": "examples/inference/basic/basic_fasth3.py",
      "serving": {
@@ -484,13 +484,13 @@
        "prepare": "hf download FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2 --local-dir ./FastH3-Preview-v0.2\npython scripts/checkpoint_conversion/convert_minimax_h3_mlx.py --model-root ./FastH3-Preview-v0.2/transformer --out ./FastH3-MLX --formats \"int6\""
      },
      "group": "fasth3-preview",
-      "group_label": "FastH3 Preview",
+      "group_label": "FastH3 V1",
      "group_task": "4-step text to video + audio",
      "family": "minimax_h3",
      "stage": "inference",
      "task": "Few-step text to video (with audio)",
-      "label": "FastH3 Preview on MLX",
-      "summary": "Run FastH3 Preview on Apple Silicon with a locally converted INT6 DiT, streamed Qwen3-VL conditioning, and native MLX video and audio VAEs.",
+      "label": "FastH3 V1 on MLX",
+      "summary": "Run FastH3 V1 on Apple Silicon with a locally converted INT6 DiT, streamed Qwen3-VL conditioning, and native MLX video and audio VAEs.",
      "model": "FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2",
      "source": "examples/inference/basic/mlx_fasth3.py",
      "command": "hf download FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2 --local-dir ./FastH3-Preview-v0.2\npython scripts/checkpoint_conversion/convert_minimax_h3_mlx.py --model-root ./FastH3-Preview-v0.2/transformer --out ./FastH3-MLX --formats \"int6\"\npython examples/inference/basic/mlx_fasth3.py --model-root ./FastH3-Preview-v0.2 --mlx-checkpoint ./FastH3-MLX/int6 --prompt \"(S1) A presenter says <d>[English] FastVideo runs FastH3.</d>\" --height 480 --width 832 --num-frames 124 --seed 2026 --output-path ./outputs/fasth3_int6.mp4",
@@ -516,13 +516,13 @@
    {
      "id": "fasth3-preview-spark",
      "group": "fasth3-preview",
-      "group_label": "FastH3 Preview",
+      "group_label": "FastH3 V1",
      "group_task": "4-step text to video + audio",
      "family": "minimax_h3",
      "stage": "inference",
      "task": "Few-step text to video (with audio)",
-      "label": "FastH3 Preview on one DGX Spark",
-      "summary": "Run FastH3 Preview on one GB10 with Triton VSA, FA4 off, and lazy module load. Height, width, frames, and steps in the YAML are examples.",
+      "label": "FastH3 V1 on one DGX Spark",
+      "summary": "Run FastH3 V1 on one GB10 with Triton VSA, FA4 off, and lazy module load. Height, width, frames, and steps in the YAML are examples.",
      "model": "FastVideo/FastVideo-FastH3-4-step-Preview-v1-VSA-DataFree",
      "source": "examples/inference/basic/basic_fasth3_spark.yaml",
      "serving": {
@@ -551,12 +551,12 @@
    {
      "id": "fasth3-spark-pair",
      "group": "fasth3-preview",
-      "group_label": "FastH3 Preview",
+      "group_label": "FastH3 V1",
      "group_task": "4-step text to video + audio",
      "family": "minimax_h3",
      "stage": "inference",
      "task": "Few-step text to video (with audio)",
-      "label": "FastH3 Preview on two DGX Sparks",
+      "label": "FastH3 V1 on two DGX Sparks",
      "summary": "Run one FastH3 clip across two GB10s with Ray sequence parallel over QSFP RoCE. Sequential load and lazy module load stay on because SP replicates the DiT on each node.",
      "model": "FastVideo/FastVideo-FastH3-4-step-Preview-v1-VSA-DataFree",
      "source": "examples/inference/basic/basic_fasth3_spark_pair.yaml",
@@ -578,6 +578,83 @@
        "Height, width, frames, and steps in the YAML are examples. Edit them or pass CLI flags. See docs/getting_started/installation/spark_pair.md."
      ]
    },
+    {
+      "id": "fasth3-8step-v2-cuda",
+      "group": "fasth3-8step-v2",
+      "group_label": "FastH3 V2",
+      "group_task": "8-step text to video + audio",
+      "family": "minimax_h3",
+      "stage": "inference",
+      "task": "Few-step text to video (with audio)",
+      "label": "FastH3 V2 on CUDA",
+      "summary": "Run FastH3 V2, the eight-forward checkpoint (video/audio shifts 10/3, VSA 0.8, 64-token tiles) with the trained DMD ladder loaded from the checkpoint's fastvideo_inference.json.",
+      "model": "FastVideo/FastVideo-FastH3-8-Step-V2",
+      "source": "examples/inference/basic/basic_fasth3_8step.py",
+      "serving": {
+        "source": "examples/serving/openai_fasth3_8step.yaml",
+        "install": "UV_TORCH_BACKEND=cu130 uv pip install -e \".[fasth3]\""
+      },
+      "command": "UV_TORCH_BACKEND=cu130 uv pip install -e \".[fasth3]\"\npython examples/inference/basic/basic_fasth3_8step.py --prompt \"(S1) A presenter says <d>[English] FastVideo runs FastH3.</d>\" --profile strict --no-inference-torch-compile --no-compile-vae",
+      "gpu_types": ["NVIDIA"],
+      "hardware": {
+        "platform": "cuda",
+        "gpu_count": 4,
+        "accelerator": "NVIDIA GB200",
+        "evidence": "validated",
+        "evidence_url": "https://github.com/hao-ai-lab/FastVideo/pull/1852"
+      },
+      "evidence": "Verified",
+      "expected_artifact": "Warmup and measured MP4 files under outputs/fasth3_8step/",
+      "modes": ["T2VA", "8-step FastH3"],
+      "knobs": [
+        {"key": "num_gpus", "label": "GPUs", "hint": "Sequence-parallel degree", "flag": "--num-gpus", "options": [1, 2, 4, 8], "default": 4},
+        {"key": "video_decode_backend", "label": "VAE decode", "hint": "Fidelity vs. speed", "flag": "--video-decode-backend", "options": [{"value": "h3-vae", "label": "Full H3 VAE"}, {"value": "taeh3", "label": "TAEH3 preview"}], "default": "h3-vae"}
+      ],
+      "limitations": [
+        "Nine sigma-grid points (eight transformer forwards) are fixed by the checkpoint's trained ladder; the example rejects any other --steps.",
+        "Validated with the eager strict route (--profile strict --no-inference-torch-compile --no-compile-vae). The compiled all profile has not been measured for this checkpoint.",
+        "T2AV only; no FL2VA/Ref2VA distillation and no matching LoRA. V2 uses eight forwards rather than V1's four."
+      ]
+    },
+    {
+      "id": "fasth3-8step-v2-mlx",
+      "serving": {
+        "source": "examples/serving/mlx_fasth3_8step.yaml",
+        "install": "uv pip install -e \".[mlx]\"",
+        "prepare": "hf download FastVideo/FastVideo-FastH3-8-Step-V2 --local-dir ./FastH3-8-Step-V2\npython scripts/checkpoint_conversion/convert_minimax_h3_mlx.py --model-root ./FastH3-8-Step-V2/transformer --out ./FastH3-8-Step-V2-MLX --formats \"int8\" --include-vsa"
+      },
+      "group": "fasth3-8step-v2",
+      "group_label": "FastH3 V2",
+      "group_task": "8-step text to video + audio",
+      "family": "minimax_h3",
+      "stage": "inference",
+      "task": "Few-step text to video (with audio)",
+      "label": "FastH3 V2 on MLX",
+      "summary": "Run FastH3 V2 on Apple Silicon with a locally converted INT8 DiT, the checkpoint's DMD contract ladder, and trained VSA (sparsity 0.8, 64-token tiles).",
+      "model": "FastVideo/FastVideo-FastH3-8-Step-V2",
+      "source": "examples/inference/basic/mlx_fasth3_8step.py",
+      "command": "hf download FastVideo/FastVideo-FastH3-8-Step-V2 --local-dir ./FastH3-8-Step-V2\npython scripts/checkpoint_conversion/convert_minimax_h3_mlx.py --model-root ./FastH3-8-Step-V2/transformer --out ./FastH3-8-Step-V2-MLX --formats \"int8\" --include-vsa\npython examples/inference/basic/mlx_fasth3_8step.py --model-root ./FastH3-8-Step-V2 --mlx-checkpoint ./FastH3-8-Step-V2-MLX/int8 --prompt \"(S1) A presenter says <d>[English] FastVideo runs FastH3.</d>\" --height 480 --width 832 --num-frames 124 --seed 2026 --output-path ./outputs/fasth3_8step_int8.mp4",
+      "gpu_types": ["Apple Silicon"],
+      "hardware": {
+        "platform": "mlx",
+        "accelerator": "Apple M4 Max",
+        "system_memory": "36 GB unified memory",
+        "peak_memory": "27.39 GiB peak MLX memory during denoising",
+        "evidence": "validated",
+        "evidence_url": "https://github.com/hao-ai-lab/FastVideo/pull/1863"
+      },
+      "evidence": "Verified",
+      "expected_artifact": "MP4 with H.264 video and stereo AAC audio at outputs/fasth3_8step_int8.mp4",
+      "modes": ["T2VA", "8-step FastH3", "trained VSA"],
+      "knobs": [
+        {"key": "video_decode_backend", "label": "VAE decode", "hint": "Fidelity vs. speed", "flag": "--video-decode-backend", "options": [{"value": "h3-vae", "label": "Full H3 VAE"}, {"value": "taeh3", "label": "TAEH3 preview"}], "default": "h3-vae"}
+      ],
+      "limitations": [
+        "Convert with --include-vsa. mlx_fasth3_8step.py turns VSA on (sparsity 0.8, tile 64). A dense export fails at configure_vsa.",
+        "--steps 8 or 9 both run the eight trained forwards. Reuse the preview VAE, audio VAE, text encoder, and tokenizer if those directories already exist.",
+        "The MLX path supports T2VA only. FL2VA, Ref2VA, and two-pass refinement are not wired."
+      ]
+    },
    {
      "id": "minimax-h3-fl2va",
      "family": "minimax_h3",
@@ -136,7 +136,7 @@
    if (platform === "mps") {
      return {
        id: "mps",
-        label: "Apple Silicon · MPS",
+        label: "Apple Silicon · PyTorch MPS",
        hint: recipe.hardware?.minimum_memory || recipe.hardware?.system_memory || "Memory not recorded",
      };
    }
@@ -519,10 +519,10 @@
          option.setAttribute("aria-pressed", String(selected));
        });
        servingAvailability.textContent = profile
-          ? "The playground and API clients share one server process. Both workflows can run on your own machine."
+          ? "The playground and the OpenAI Python client share one server process. Both workflows can run on your own machine."
          : servingLoadFailed
            ? "Server examples could not be loaded. Open the H3 server guide below, or use Python directly."
-            : "This recipe uses Python directly. For the playground and API clients, choose FastH3 Preview with CUDA, MLX, or one Spark.";
+            : "This recipe uses Python directly. FastH3 V1 and FastH3 V2 can also run a local server for the playground and the OpenAI Python client.";
        servingPanel.hidden = !useServer;
        commandBlock.hidden = useServer;
        root.querySelector("[data-cookbook-python-note]").hidden = useServer;
@@ -537,7 +537,7 @@
            ? "Start once, then change prompts in the playground or your app. On a DGX Spark, lazy module load still reloads Qwen3-VL and the DiT between phases of each request, so later prompts are not a free hot cache."
            : "Start once, then change prompts in the playground or your app. CUDA requests reuse the loaded model. The Python SDK can also reuse a generator within one process.";
        servingPanel.querySelector("[data-cookbook-install-guide]").href = isMLX
-          ? "../../getting_started/installation/mps/#run-fasth3-preview"
+          ? "../../getting_started/installation/mlx/"
          : isSpark
            ? "../../getting_started/installation/spark/"
            : "../../getting_started/installation/gpu/";
@@ -574,7 +574,7 @@
      });

      description.textContent = useServer
-        ? `FastH3 Preview generates video with audio. This server profile uses the checked-in ${runtime.label} configuration.`
+        ? `${recipe.group_label || recipe.label} generates video with audio. Start the local server, then use the playground or the OpenAI Python client. This profile uses the checked-in ${runtime.label} configuration.`
        : recipe.summary;
      label.textContent = useServer ? `${recipe.group_label || recipe.label} · Server` : recipe.label;
      model.textContent = recipe.model;
@@ -18,8 +18,8 @@ the CUDA `fastvideo-kernel` package:
 - **Dense-only checkpoints** (the default converter) drop the 50 gate
  matrices and keep fused SDPA. They remain valid for dense inference.
 - **VSA-capable checkpoints** retain those gates, quantize them on the same
-  affine grid, and record `vsa.capable` in `mlx_h3_dit.json`. Runtime VSA is
-  still off until you pass `--vsa`.
+  affine grid, and record `vsa.capable` in `mlx_h3_dit.json`. Preview leaves
+  runtime VSA off until you pass `--vsa`. `mlx_fasth3_8step.py` turns it on.
 - **Tile sizes** 64 `(4, 4, 4)` and 256 `(4, 8, 8)`. Prefix keys can be
  `exempt` or `compete`. `--vsa-dense-first-n-steps` and `--vsa-dense-layers`
  force dense SDPA on the selected steps or blocks.
@@ -32,9 +32,12 @@ the CUDA `fastvideo-kernel` package:
  but does not yet match reference video. `--vsa-impl reference` is the same
  as `auto`.

-See the [Apple Silicon guide](../../getting_started/installation/mps.md) for
-conversion and `mlx_fasth3.py` flags. Do not enable VSA on a dense-only
-checkpoint; reconvert with `--include-vsa` first.
+See the [MLX install guide](../../getting_started/installation/mlx.md)
+and the [MiniMax H3 cookbook](../../cookbook/minimax-h3.md) for conversion
+and `mlx_fasth3.py` / `mlx_fasth3_8step.py` flags. Do not enable
+VSA on a dense-only checkpoint; reconvert with `--include-vsa` first. V1
+VSA is opt-in. FastH3 V2 converts with `--include-vsa` and turns VSA on
+by default.

 H3 uses fused MLX RMSNorm by default, including dense inference. This can
 change BF16 rounding relative to the older explicit normalization path.
@@ -30,7 +30,7 @@ FastVideo and the reference model first produce different numbers?"
 | General logging | `init_logger(__name__)` |
 | Per-stage timing | `FASTVIDEO_STAGE_LOGGING` |
 | Profiling kernel timings | `FASTVIDEO_TORCH_PROFILER_DIR` (see [Profiling](profiling.md)) |
-| Function-call tracing | `FASTVIDEO_TRACE_FUNCTION` (heavy) |
+| Function-call tracing | `fastvideo.logger.enable_trace_function_call()` (heavy) |

 ## Quickstart

@@ -4,6 +4,11 @@ This is the canonical reference for FastVideo's CI/CD system. Contributor-facing
 PR steps live in [Pull Requests](pull_requests.md), and test-authoring guidance
 lives in [Testing](testing.md).

+The existing Slurm route below remains the default. Operators can also install
+the [selectable GPU dispatcher](gpu_ci_backends.md) to run the same lane scripts
+on Modal or Kubernetes in the `vllm` namespace. That opt-in path enforces two
+active PRs, four GPUs per PR, and eight total, with separate backend statuses.
+
 ## Overview

 FastVideo splits validation across GitHub Actions, Buildkite, Slinky Slurm,
@@ -435,8 +440,11 @@ 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.
+kernel wheel compiled with `TORCH_CUDA_ARCH_LIST=10.0a` to include the GB200
+VSA CUDA extensions. Runtime kernel detection uses the same target, so runner
+jobs can validate and install the exact source and ABI match instead of
+recompiling it in every lane. Older artifacts built for `10.0` have a different
+cache key and trigger a local rebuild when the worker detects `10.0a`.

 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
@@ -0,0 +1,242 @@
+# Environment Variables
+
+FastVideo reads environment variables for expert switches, debugging, profiling, and the settings that launchers such
+as `torchrun` provide. This page is the policy for those variables. The contract test
+`fastvideo/tests/contract/test_env_policy.py` enforces the policy in the unit CI lane, and the coding-agent skill
+`.agents/skills/env-var-conventions/SKILL.md` points here. When the policy changes, update this page and the contract
+test in the same pull request.
+
+## Rules
+
+1. **Register every FastVideo variable in `fastvideo/envs.py`.** Each entry declares a type, a default, a category,
+   and a description. Variables that other tools own (CUDA, NCCL, PyTorch, launchers) are not registered; code reads
+   them directly with `os.environ.get("NAME")`, and the name must be in the external-variable allowlist
+   (`EXTERNAL_ALLOWLIST` in the contract test). When FastVideo sets such a variable for the other tool, it calls
+   `envs.set_external`, `envs.setdefault_external`, or `envs.unset_external`, and the name must be in
+   `EXTERNAL_WRITE_ALLOWLIST`.
+2. **Read with `envs.NAME.get()`, write with `envs.NAME.set()`, and change a value in tests with
+   `envs.NAME.override()`.** Each type has one parsing rule. A value that the rule rejects raises
+   `fastvideo.envs.EnvVarError` instead of falling back to the default.
+3. **Name FastVideo variables with the `FASTVIDEO_` prefix.** The second word states the purpose where one applies:
+   `ENABLE_`, `DISABLE_`, `USE_`, `FORCE_`, `DEBUG_`, `TEST_`.
+4. **Keep a renamed variable as a deprecated alias until the next minor release.** Setting the old name logs a
+   warning. Delete a variable that no code reads, and list it in `DEPRECATED_VARIABLES` so that setting it logs a
+   warning.
+5. **Give each setting one source: an argument or an environment variable.** Settings that users change per
+   deployment are arguments (CLI or YAML). Expert switches, emergency off switches, and debugging and test switches
+   are environment variables.
+6. **Read variables inside functions.** `envs.NAME.get()` runs when the function runs, so a changed value takes
+   effect without re-importing a module. Module level, class bodies, decorators, and default argument values run at
+   import time.
+7. **Do not write the environment to pass values between parts of FastVideo.** Pass an argument instead. Tests use
+   `envs.NAME.override()`.
+
+## Field types
+
+| Class        | Value type       | Parsing rule                                                                        |
+| ------------ | ---------------- | ----------------------------------------------------------------------------------- |
+| `EnvBool`    | `bool`           | `1`, `true`, `yes`, `on` are true; `0`, `false`, `no`, `off`, and `""` are false.   |
+|              |                  | Case-insensitive; surrounding whitespace is ignored.                                |
+| `EnvInt`     | `int`            | `int(value)`                                                                        |
+| `EnvFloat`   | `float`          | `float(value)`                                                                      |
+| `EnvStr`     | `str` or `None`  | The raw string. A `None` default means that the variable has no default.            |
+| `EnvPath`    | `str` or `None`  | The raw string with a leading `~` expanded.                                         |
+| `EnvChoice`  | `str`            | Stripped and lower-cased, then checked against the declared `choices`.              |
+
+A default can be a zero-argument function; `get()` calls it on each read while the variable is unset. The path roots
+use this to follow `XDG_CONFIG_HOME` and `XDG_CACHE_HOME`.
+
+Using a field without a method, as in `if envs.FASTVIDEO_FA4:`, raises `TypeError`.
+
+## Add a variable
+
+1. Declare the variable in the matching section of `fastvideo/envs.py`:
+
+    ```python
+    FASTVIDEO_DEBUG_MY_STAGE = EnvBool(False, category="debug", doc="Log the inputs of MyStage.")
+    ```
+
+    The category is one of the values in `envs.CATEGORIES`.
+
+2. Read the variable inside a function:
+
+    ```python
+    import fastvideo.envs as envs
+
+    def forward(self, batch):
+        if envs.FASTVIDEO_DEBUG_MY_STAGE.get():
+            logger.info("MyStage inputs: %s", batch.keys())
+    ```
+
+3. Regenerate the table at the end of this page:
+
+    ```bash
+    python fastvideo/tests/contract/test_env_policy.py
+    ```
+
+4. Run the contract test:
+
+    ```bash
+    pytest fastvideo/tests/contract/test_env_policy.py
+    ```
+
+In a test, change the value with `override`, which restores the previous value on exit:
+
+```python
+with envs.FASTVIDEO_DEBUG_MY_STAGE.override(True):
+    run_stage()
+```
+
+## Rename or remove a variable
+
+To rename a variable, declare it under the new name and list the old name in `deprecated_names`:
+
+```python
+FASTVIDEO_LTX2_USE_DISTILLED_SIGMAS = EnvBool(True,
+                                              category="sampling",
+                                              doc="...",
+                                              deprecated_names=("LTX2_USE_DISTILLED_SIGMAS", ))
+```
+
+`get()` reads an old name only when the new name is unset, and logs a warning once. Update the uses of the old name
+in `examples/`, `scripts/`, `docs/`, `apps/`, and the tests in the same pull request. Delete the old name in the next
+minor release.
+
+To remove a variable that no code reads, delete its entry and add the name to `DEPRECATED_VARIABLES` in
+`fastvideo/envs.py` with a reason. `FastVideoArgs` calls `envs.warn_deprecated_variables()`, which logs a warning for
+each listed variable that is set. Delete the entry in the next minor release.
+
+## What the contract test checks
+
+The test parses every Python file under `fastvideo/`, including `fastvideo/tests/`, with Python's `ast` module. It
+skips `fastvideo/third_party/`, which is copied from upstream projects, and the registry `fastvideo/envs.py`. It does
+not check `apps/`, `examples/`, `scripts/`, `fastvideo-kernel/`, or `docs/`.
+
+It reports each violation as `<path>: <kind> <name>`:
+
+| Kind               | Code that triggers it                                        | Fix                                          |
+| ------------------ | ------------------------------------------------------------ | -------------------------------------------- |
+| `read`             | `os.getenv`, `os.environ.get`, `os.environ[...]`, or         | Register the variable and call               |
+|                    | `"NAME" in os.environ` with a name outside the allowlist, or | `envs.NAME.get()`. For a variable that       |
+|                    | with a name built at runtime (`<dynamic>`)                   | another tool owns, add it to                 |
+|                    |                                                              | `EXTERNAL_ALLOWLIST` with a reason.          |
+| `write`            | `os.environ[...] = ...`, `setdefault`, `pop`, `del`,         | Pass an argument instead. In tests, use      |
+|                    | `os.putenv`, `os.unsetenv`, `monkeypatch.setenv`/`delenv`,   | `envs.NAME.override()`. For a variable that  |
+|                    | or an `envs.*_external` call with a name outside             | another tool reads, call an                  |
+|                    | `EXTERNAL_WRITE_ALLOWLIST`                                   | `envs.*_external` helper and add the name to |
+|                    |                                                              | `EXTERNAL_WRITE_ALLOWLIST` with a reason.    |
+| `whole-environ`    | `os.environ.copy()`, `dict(os.environ)`, iteration,          | Read the specific variables that the code    |
+|                    | `mock.patch.dict(os.environ, ...)`, `os.environ.update`      | needs.                                       |
+| `bare-field`       | A registry field used without calling one of its methods,    | Call `envs.NAME.get()`.                      |
+|                    | as in `envs.NAME == "auto"` or `getter = envs.NAME.get`      |                                              |
+| `import-time-read` | `envs.NAME.get()` outside a function                         | Move the read into the function that uses    |
+|                    |                                                              | the value.                                   |
+| `prefix`           | A registry entry without the `FASTVIDEO_` prefix             | Rename the variable and keep the old name in |
+|                    |                                                              | `deprecated_names`.                          |
+| `unread`           | A registry entry that no code reads with `get()` or          | Delete the variable and add it to            |
+|                    | `is_set()`                                                   | `DEPRECATED_VARIABLES`.                      |
+
+The test recognizes `os` imported under another name, `from os import environ, getenv`, and a name held in a
+module-level string constant. Code that reaches the environment through `importlib` or `getattr(os, "environ")` is
+left to code review.
+
+The test also checks that every registry entry has a category from `envs.CATEGORIES` and a description, and that the
+table at the end of this page matches the registry.
+
+**Known violations.** `KNOWN_VIOLATIONS` in the contract test lists the violations that existed when the policy was
+introduced. The list only shrinks. A violation that is not in the list fails the test. A listed violation that no
+longer exists also fails the test, so the fixing pull request deletes its entry.
+
+## Registered variables
+
+<!-- BEGIN GENERATED ENV TABLE: python fastvideo/tests/contract/test_env_policy.py -->
+| Variable                                                   | Type                           | Default                                               | Category    | Description                                                                                                                                                                                                                                              |
+| ---------------------------------------------------------- | ------------------------------ | ----------------------------------------------------- | ----------- | -------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- |
+| `FASTVIDEO_CONFIG_ROOT`                                    | path                           | computed                                              | path        | Root directory for FastVideo configuration files, at runtime and at installation. Defaults to ~/.config/fastvideo, or $XDG_CONFIG_HOME/fastvideo when XDG_CONFIG_HOME is set.                                                                            |
+| `FASTVIDEO_CACHE_ROOT`                                     | path                           | computed                                              | path        | Root directory for FastVideo cache files. Defaults to ~/.cache/fastvideo, or $XDG_CACHE_HOME/fastvideo when XDG_CACHE_HOME is set.                                                                                                                       |
+| `FASTVIDEO_REASON1_WEIGHTS_PATH`                           | str                            | unset                                                 | path        | Local path or Hugging Face id of Reason1 weights to load instead of the checkpoint's own.                                                                                                                                                                |
+| `FASTVIDEO_HOST_IP`                                        | str                            | `""`                                                  | distributed | IP address of this node when the node has several network interfaces. Set it on each node for multi-node inference.                                                                                                                                      |
+| `FASTVIDEO_LOOPBACK_IP`                                    | str                            | `""`                                                  | distributed | Loopback IP address to use instead of the detected one.                                                                                                                                                                                                  |
+| `FASTVIDEO_RAY_PER_WORKER_GPUS`                            | float                          | `1.0`                                                 | distributed | GPUs per Ray worker. A fraction lets Ray schedule several actors on one GPU, so other actors can share the GPUs with FastVideo.                                                                                                                          |
+| `FASTVIDEO_NCCL_SO_PATH`                                   | str                            | unset                                                 | distributed | Path to the NCCL library file. Needed because the nccl&gt;=2.19 that PyTorch ships has a bug (https://github.com/NVIDIA/nccl/issues/1234).                                                                                                               |
+| `FASTVIDEO_HCCL_SO_PATH`                                   | str                            | unset                                                 | distributed | Path to the HCCL library file on Ascend NPUs. Deprecated names: `HCCL_SO_PATH`.                                                                                                                                                                          |
+| `FASTVIDEO_WORKER_MULTIPROC_METHOD`                        | one of spawn, fork, forkserver | `spawn`                                               | distributed | Multiprocessing start method for worker processes.                                                                                                                                                                                                       |
+| `FASTVIDEO_ULYSSES_A2A`                                    | one of off, auto               | `off`                                                 | distributed | Sequence-parallel all-to-all backend. off uses the NCCL path in DistributedAutograd.AllToAll4D. auto uses the fused NVLink kernel when the group is a load-store accessible mesh of 2, 4, 6, or 8 ranks in eager execution, and the NCCL path otherwise. |
+| `FASTVIDEO_CONFIGURE_LOGGING`                              | bool                           | `1`                                                   | logging     | Configure logging at import. When true, FastVideo uses its default logging configuration or the file in FASTVIDEO_LOGGING_CONFIG_PATH.                                                                                                                   |
+| `FASTVIDEO_LOGGING_CONFIG_PATH`                            | str                            | unset                                                 | logging     | Path to a JSON logging configuration file.                                                                                                                                                                                                               |
+| `FASTVIDEO_LOGGING_LEVEL`                                  | str                            | `INFO`                                                | logging     | Default logging level.                                                                                                                                                                                                                                   |
+| `FASTVIDEO_LOGGING_PREFIX`                                 | str                            | `""`                                                  | logging     | Prefix prepended to every log message.                                                                                                                                                                                                                   |
+| `FASTVIDEO_STAGE_LOGGING`                                  | bool                           | `0`                                                   | logging     | Log the time that each pipeline stage takes.                                                                                                                                                                                                             |
+| `FASTVIDEO_ATTENTION_BACKEND`                              | str                            | unset                                                 | attention   | Attention backend, as an AttentionBackendEnum name such as TORCH_SDPA, FLASH_ATTN, VIDEO_SPARSE_ATTN, SAGE_ATTN, or SAGE_ATTN_THREE. FastVideoArgs uses it when FastVideoArgs.attention_backend is unset.                                                |
+| `FASTVIDEO_FA4`                                            | bool                           | `0`                                                   | attention   | The FLASH_ATTN backend uses FlashAttention-4 (flash_attn.cute) instead of FA3 or FA2.                                                                                                                                                                    |
+| `FASTVIDEO_MINIMAX_H3_FA4_PACKED_VARLEN`                   | bool                           | `0`                                                   | attention   | MiniMax-H3 dense DiT self-attention uses the FlashAttention-4 packed-varlen entry point. This changes the floating-point reduction order, so it is an inference-only opt-in.                                                                             |
+| `FASTVIDEO_VSA_SM100A`                                     | bool                           | `0`                                                   | attention   | VIDEO_SPARSE_ATTN_H3 sends no-grad tile-64 forwards to the data-center Blackwell (sm_100a) kernel. fastvideo-kernel reads the same variable with the same rule.                                                                                          |
+| `FASTVIDEO_NVFP4_FA4`                                      | bool                           | `0`                                                   | attention   | FlashAttention-4 quantizes Q and K to NVFP4. An explicit nvfp4_fa4 attention implementation argument takes precedence.                                                                                                                                   |
+| `FASTVIDEO_DISABLE_ATTENTION_COMPILE`                      | bool                           | `1`                                                   | attention   | Keep attention forward out of torch.compile graphs (torch.compiler.disable). Set it to 0 to let attention constructed under that setting be traced. Setting it explicitly to true also blocks regional compile.                                          |
+| `FASTVIDEO_MLX_WINDOW`                                     | int                            | `0`                                                   | attention   | MLX FastWan windowed attention size in tokens. 0 uses full attention.                                                                                                                                                                                    |
+| `FASTVIDEO_MLX_WINDOW_SINK`                                | int                            | `0`                                                   | attention   | Number of sink tokens that MLX windowed attention always attends to.                                                                                                                                                                                     |
+| `FASTVIDEO_INFERENCE_TORCH_COMPILE`                        | bool                           | `0`                                                   | performance | Compile each DiT transformer block with fullgraph torch.compile at inference. Same as FastVideoArgs.inference_torch_compile=True.                                                                                                                        |
+| `FASTVIDEO_VAE_PARALLEL_DECODE`                            | bool                           | `0`                                                   | performance | MiniMax-H3 VAE decode splits its temporal chunks across the sequence-parallel ranks instead of running serially on the output rank. Same as FastVideoArgs.vae_parallel_decode=True.                                                                      |
+| `FASTVIDEO_VAE_PARALLEL_ENCODE`                            | bool                           | `0`                                                   | performance | MiniMax-H3 reference-video VAE encode splits its temporal chunks across the sequence-parallel ranks. Same as FastVideoArgs.vae_parallel_encode=True.                                                                                                     |
+| `FASTVIDEO_VAE_PARALLEL_DECODE_STRATEGY`                   | str                            | unset                                                 | performance | Collective that moves chunks in parallel VAE decode: gather (used when unset) or all_gather.                                                                                                                                                             |
+| `FASTVIDEO_MINIMAX_H3_FUSIONS`                             | str                            | `""`                                                  | performance | MiniMax-H3 inference-only Triton fusions: all, 1, or a comma-separated subset of modulate,qknorm_rope,swiglu. Empty, 0, or none keeps the eager implementation.                                                                                          |
+| `FASTVIDEO_FSDP2_AUTOWRAP`                                 | bool                           | `0`                                                   | performance | FSDP2 shards modules by parameter count instead of the model's shard conditions. Not supported by self-forcing distillation.                                                                                                                             |
+| `FASTVIDEO_FSDP2_MIN_PARAMS`                               | int                            | `10000000`                                            | performance | Minimum parameter count of a module that FASTVIDEO_FSDP2_AUTOWRAP shards.                                                                                                                                                                                |
+| `FASTVIDEO_MLX_COMPILE`                                    | bool                           | `0`                                                   | performance | Compile the MLX DiT forward with mx.compile.                                                                                                                                                                                                             |
+| `FASTVIDEO_MLX_FAST_NORM`                                  | bool                           | `0`                                                   | performance | Use MLX fast normalization kernels.                                                                                                                                                                                                                      |
+| `FASTVIDEO_MLX_DQ_GEMM`                                    | str                            | `1`                                                   | performance | MLX dequantized GEMM for affine-quantized weights: 0 turns it off, 1 uses the measured minimum row count, and an integer sets the minimum row count.                                                                                                     |
+| `FASTVIDEO_LTX2_VAE_CHANNELS_LAST_3D`                      | bool                           | `1`                                                   | performance | LTX-2 VAE uses the channels_last_3d memory format.                                                                                                                                                                                                       |
+| `FASTVIDEO_LTX2_DISABLE_AUDIO_AUTOCAST`                    | bool                           | `1`                                                   | performance | LTX-2 audio decoding runs without CUDA autocast. Deprecated names: `LTX2_DISABLE_AUDIO_AUTOCAST`.                                                                                                                                                        |
+| `FASTVIDEO_FLUX2_DISABLE_BF16_REDUCED_PRECISION_REDUCTION` | bool                           | `0`                                                   | performance | Flux denoising disables reduced-precision reductions in bf16 matmuls, which tightens accumulation for the 4-step Klein model.                                                                                                                            |
+| `FASTVIDEO_FFMPEG_BIN`                                     | str                            | `ffmpeg`                                              | output      | ffmpeg executable used to save video with audio.                                                                                                                                                                                                         |
+| `FASTVIDEO_VIDEO_CODEC`                                    | str                            | `libx264`                                             | output      | ffmpeg video codec for saved videos.                                                                                                                                                                                                                     |
+| `FASTVIDEO_NVENC_PRESET`                                   | str                            | `p1`                                                  | output      | NVENC preset when the codec is an \*_nvenc codec.                                                                                                                                                                                                        |
+| `FASTVIDEO_NVENC_TUNE`                                     | str                            | `ull`                                                 | output      | NVENC tune option.                                                                                                                                                                                                                                       |
+| `FASTVIDEO_NVENC_RC`                                       | str                            | `constqp`                                             | output      | NVENC rate-control mode.                                                                                                                                                                                                                                 |
+| `FASTVIDEO_NVENC_QP`                                       | str                            | `28`                                                  | output      | NVENC quantization parameter.                                                                                                                                                                                                                            |
+| `FASTVIDEO_NVENC_BF`                                       | str                            | `0`                                                   | output      | NVENC number of B-frames.                                                                                                                                                                                                                                |
+| `FASTVIDEO_X264_PRESET`                                    | str                            | `ultrafast`                                           | output      | x264 preset for non-NVENC codecs.                                                                                                                                                                                                                        |
+| `FASTVIDEO_OUTPUT_PIX_FMT`                                 | str                            | `yuv420p`                                             | output      | ffmpeg pixel format for saved videos.                                                                                                                                                                                                                    |
+| `FASTVIDEO_NVTX_PROFILE`                                   | bool                           | `0`                                                   | profiling   | Emit NVTX ranges for external profilers such as Nsight Systems.                                                                                                                                                                                          |
+| `FASTVIDEO_TORCH_PROFILER_DIR`                             | path                           | unset                                                 | profiling   | Enables the torch profiler and sets the directory for its traces. Must be an absolute path.                                                                                                                                                              |
+| `FASTVIDEO_TORCH_PROFILER_RECORD_SHAPES`                   | bool                           | `0`                                                   | profiling   | Torch profiler records shapes.                                                                                                                                                                                                                           |
+| `FASTVIDEO_TORCH_PROFILER_WITH_PROFILE_MEMORY`             | bool                           | `0`                                                   | profiling   | Torch profiler profiles memory.                                                                                                                                                                                                                          |
+| `FASTVIDEO_TORCH_PROFILER_WITH_STACK`                      | bool                           | `0`                                                   | profiling   | Torch profiler captures stacks. Costs about 1.5x runtime and 1.4x trace size.                                                                                                                                                                            |
+| `FASTVIDEO_TORCH_PROFILER_WITH_FLOPS`                      | bool                           | `0`                                                   | profiling   | Torch profiler profiles FLOPs.                                                                                                                                                                                                                           |
+| `FASTVIDEO_TORCH_PROFILE_REGIONS`                          | str                            | `""`                                                  | profiling   | Comma-separated profiler regions to record. The torch profiler requires at least one region.                                                                                                                                                             |
+| `FASTVIDEO_TRACE_ACTIVATIONS`                              | bool                           | `0`                                                   | debug       | Enable activation trace hooks.                                                                                                                                                                                                                           |
+| `FASTVIDEO_TRACE_LAYERS`                                   | str                            | `""`                                                  | debug       | Regex filter for traced module names. Empty means all.                                                                                                                                                                                                   |
+| `FASTVIDEO_TRACE_STATS`                                    | str                            | `abs_mean,sum`                                        | debug       | Comma-separated activation statistics dumped for each output tensor.                                                                                                                                                                                     |
+| `FASTVIDEO_TRACE_OUTPUT`                                   | str                            | `/tmp/fv_trace_<pid>.jsonl`                           | debug       | JSONL path for activation traces. The literal &lt;pid&gt; is replaced at runtime.                                                                                                                                                                        |
+| `FASTVIDEO_TRACE_STEPS`                                    | str                            | `""`                                                  | debug       | Comma-separated denoising step indices. Empty means all.                                                                                                                                                                                                 |
+| `FASTVIDEO_H3_VSA_PROBE`                                   | str                            | unset                                                 | debug       | Output directory for the VSA-H3 attention-mass probe, which writes one .pt file per step, layer, and rank. Keeps the model out of regional compile.                                                                                                      |
+| `FASTVIDEO_LTX2_GEMMA_LOG`                                 | str                            | `""`                                                  | debug       | Log file for LTX-2 Gemma text-encoder hidden states, used by parity tests. Deprecated names: `LTX2_FASTVIDEO_GEMMA_LOG`.                                                                                                                                 |
+| `FASTVIDEO_COSMOS25_LOG_KNOBS`                             | bool                           | `0`                                                   | debug       | Log the Cosmos 2.5 latent-preparation conditioning inputs.                                                                                                                                                                                               |
+| `FASTVIDEO_CFG_GATE_STEP`                                  | float                          | `1.0`                                                 | sampling    | CFG gating fraction in [0, 1]. Steps before len(timesteps) \* X run the conditional and unconditional forwards; later steps reuse the cached difference. 1.0 disables gating.                                                                            |
+| `FASTVIDEO_LTX2_USE_DISTILLED_SIGMAS`                      | bool                           | `1`                                                   | sampling    | LTX-2 uses the distilled sigma schedule when FastVideoArgs.ltx2_use_distilled_sigmas is also true. Deprecated names: `LTX2_USE_DISTILLED_SIGMAS`.                                                                                                        |
+| `FASTVIDEO_EVAL_CACHE`                                     | path                           | computed                                              | eval        | Cache directory for evaluation models and datasets. Defaults to $FASTVIDEO_CACHE_ROOT/eval.                                                                                                                                                              |
+| `FASTVIDEO_PHYSICS_IQ_BUCKET_URL`                          | str                            | `https://storage.googleapis.com/physics-iq-benchmark` | eval        | Base URL of the Physics-IQ benchmark bucket.                                                                                                                                                                                                             |
+| `FASTVIDEO_VBENCH_FULL_INFO_JSON`                          | str                            | unset                                                 | eval        | Path to VBench_full_info.json, used instead of the vendored copy. Deprecated names: `VBENCH_FULL_INFO_JSON`.                                                                                                                                             |
+| `FASTVIDEO_FVD_REF_FEATURES`                               | str                            | unset                                                 | eval        | Cached reference-feature file for the FVD metric.                                                                                                                                                                                                        |
+| `FASTVIDEO_FAD_REF_FEATURES`                               | str                            | unset                                                 | eval        | Cached reference-feature file for the audio Frechet distance metric.                                                                                                                                                                                     |
+| `FASTVIDEO_TEST_LTX2_OVERFIT_DATA_DIR`                     | str                            | `data/cats`                                           | test        | Raw data directory for preprocess_ltx2_overfit.py. Deprecated names: `LTX2_OVERFIT_DATA_DIR`.                                                                                                                                                            |
+| `FASTVIDEO_TEST_LTX2_OVERFIT_CAPTION_JSON`                 | str                            | `videos2caption_1_sample.json`                        | test        | Caption file, relative to the raw data directory. Deprecated names: `LTX2_OVERFIT_CAPTION_JSON`.                                                                                                                                                         |
+| `FASTVIDEO_TEST_LTX2_OVERFIT_VIDEO_SUBDIR`                 | str                            | `video`                                               | test        | Video subdirectory, relative to the raw data directory. Deprecated names: `LTX2_OVERFIT_VIDEO_SUBDIR`.                                                                                                                                                   |
+| `FASTVIDEO_TEST_LTX2_OVERFIT_OUTPUT_DIR`                   | str                            | `data/ltx2_overfit_preprocessed`                      | test        | Output directory for preprocess_ltx2_overfit.py. Deprecated names: `LTX2_OVERFIT_OUTPUT_DIR`.                                                                                                                                                            |
+| `FASTVIDEO_TEST_LTX2_OVERFIT_MODEL`                        | str                            | `FastVideo/LTX2-Distilled-Diffusers`                  | test        | Model repository whose encoders preprocess_ltx2_overfit.py uses. Deprecated names: `LTX2_OVERFIT_MODEL`.                                                                                                                                                 |
+| `FASTVIDEO_TEST_LTX2_OVERFIT_NUM_COPIES`                   | int                            | `4`                                                   | test        | Number of copies of the overfit sample in the parquet file. Deprecated names: `LTX2_OVERFIT_NUM_COPIES`.                                                                                                                                                 |
+| `FASTVIDEO_TEST_KANDINSKY5_OVERFIT_DATA_DIR`               | str                            | `data/kandinsky5_overfit`                             | test        | Raw data directory for preprocess_kandinsky5_overfit.py. Deprecated names: `KANDINSKY5_OVERFIT_DATA_DIR`.                                                                                                                                                |
+| `FASTVIDEO_TEST_KANDINSKY5_OVERFIT_OUTPUT_DIR`             | str                            | `data/kandinsky5_overfit_preprocessed`                | test        | Output directory for preprocess_kandinsky5_overfit.py. Deprecated names: `KANDINSKY5_OVERFIT_OUTPUT_DIR`.                                                                                                                                                |
+
+Variables that FastVideo no longer reads; setting one logs a warning:
+
+| Deprecated variable                       | Reason           |
+| ----------------------------------------- | ---------------- |
+| `FASTVIDEO_TARGET_DEVICE`                 | no code reads it |
+| `FASTVIDEO_USE_PRECOMPILED`               | no code reads it |
+| `FASTVIDEO_RINGBUFFER_WARNING_INTERVAL`   | no code reads it |
+| `FASTVIDEO_ENGINE_ITERATION_TIMEOUT_S`    | no code reads it |
+| `FASTVIDEO_SERVER_DEV_MODE`               | no code reads it |
+| `FASTVIDEO_TEST_DYNAMO_FULLGRAPH_CAPTURE` | no code reads it |
+| `FASTVIDEO_TRACE_FUNCTION`                | no code reads it |
+<!-- END GENERATED ENV TABLE -->
@@ -0,0 +1,252 @@
+# Selectable GPU CI Backends
+
+GPU CI can use the existing Slurm dispatcher, Modal, or GB200 Kubernetes Jobs
+in the `vllm` namespace. GitHub and Buildkite remain the trigger and reporting
+systems. `vllm` names the cluster namespace here; tests run FastVideo's existing
+lane scripts rather than a vLLM inference server.
+
+The default remains Slurm. The existing `.buildkite/pipeline.yml`, private
+Slurm uploader, and dormant Modal launchers are preserved. Deploying the new
+trusted uploader is an operator step; merging these files does not change the
+live Buildkite pipelines or create cluster resources. See
+[CI/CD Architecture](ci_architecture.md) for the existing installation.
+
+## Execution and Limits
+
+```text
+GitHub PR / slash command / schedule
+  -> trusted Buildkite bootstrap
+     -> slurm: existing validated graph and private dispatcher
+     -> modal or vllm: one trusted suite coordinator
+        -> shared PR admission and GPU reservations
+        -> isolated GPU worker per existing lane
+        -> lane results, logs, and backend-specific suite status
+```
+
+All new Modal and `vllm` coordinators share one admission database. It enforces:
+
+- At most two distinct PRs with admitted GPU work.
+- At most four reserved GPUs per PR across its concurrent builds and lanes.
+- At most eight reserved GPUs across both backends combined.
+- Additional PRs wait without creating GPU workers.
+
+A PR is identified by repository and PR number, so Fastcheck, merge builds,
+manual reruns, and separate attempts share its allowance. Its slot remains
+occupied between lanes until all admitted builds for that PR finish. A
+non-PR build, including a schedule on `main`, consumes its own slot and the
+same GPU budget. The preserved Slurm dispatcher has its existing independent
+limits; this new admission database does not control legacy Slurm jobs.
+
+Lanes retain their existing one-, two-, or four-GPU requirements. Four-GPU
+SSIM and training lanes wait for that PR's other lanes to release capacity.
+Selected integration lanes wait for the golden gate and are skipped if it
+fails. CPU-only GitHub checks are unaffected. Pending admission is bounded by
+`queue_timeout_seconds`, initially six hours. The coordinator's Buildkite
+command timeout is eight hours, including admission and lane waits after the
+command starts; time waiting for a Buildkite agent is outside that timeout.
+
+Reservations persist before a worker is created. Cancellation releases them
+only after worker termination is confirmed. A lost coordinator, uncertain
+create response, or unavailable backend retains its reservations until
+recovery confirms cleanup. There is no heartbeat expiry that silently frees
+GPUs while an old worker could still be running.
+
+## Operator Installation
+
+Use an operator-reviewed immutable checkout under
+`/opt/fastvideo-gpu-ci/source`. Install the supplied `scripts/gpu_ci/run` and
+`scripts/gpu_ci/upload` wrappers as `/opt/fastvideo-gpu-ci/run` and
+`/opt/fastvideo-gpu-ci/upload`. They invoke the reviewed Python entrypoint
+with isolated import mode. Install the supplied files instead of generating
+shell wrappers from build metadata:
+
+```bash
+install -m 755 /opt/fastvideo-gpu-ci/source/scripts/gpu_ci/run /opt/fastvideo-gpu-ci/run
+install -m 755 /opt/fastvideo-gpu-ci/source/scripts/gpu_ci/upload /opt/fastvideo-gpu-ci/upload
+```
+
+The controller needs Python 3.10 or newer, PyYAML,
+the Buildkite agent, and `kubectl`; Modal additionally needs the reviewed
+Modal SDK and a controller-side Modal credential.
+
+Copy [the configuration example](../../scripts/gpu_ci/config.example.json)
+to `/etc/fastvideo-gpu-ci.json`. Keep the installation and configuration
+operator-owned and unwritable by workers. Configure the actual legacy
+`slurm_uploader` command, replace each enabled backend's image placeholder
+with a reviewed registry digest, and leave `default_backend` as `slurm`
+during canaries. The checked-in placeholders deliberately cannot start jobs.
+
+The permanent dispatcher must reach the Kubernetes API independently of a
+developer laptop, SSH tunnel, or Tailscale session. The existing CPU development
+pod is useful for investigating access, but is not a production CI service.
+Use a dedicated controller identity and a separate `gpu-ci-dispatch` Buildkite
+queue. The trusted bootstrap also needs access to the existing Slurm uploader
+while that backend remains enabled.
+
+Run all coordinators on **one host**, with `state_path` and `artifacts_dir` on
+its persistent local disk. The implementation uses SQLite transactions and
+local file locks. Do not place the database or locks on Lustre/NFS, run
+independent database copies, or scale controller replicas across nodes. Those
+configurations would invalidate the global limits. Multiple Buildkite agent
+processes on the same host can share the installation and ledger; enough
+agents are needed for concurrent suite coordinators and queued work.
+
+Configure agent-owned hooks to skip repository checkout and reject commands
+other than the trusted uploader and coordinator. Disable repository hooks
+and plugins on this queue. Do not use a PR checkout to load `scripts/gpu_ci`,
+its configuration, or its worker entrypoint. Build metadata is validated by
+the dispatcher; arbitrary environment variables are not forwarded into GPU
+workers. Buildkite, Kubernetes, and Modal control-plane credentials stay on
+the dispatcher.
+
+For Kubernetes, provision namespace-scoped permissions to create/get/delete
+Jobs and read Pods and their logs. Worker Pods disable service-account token
+mounting and request the exact GPU count on ARM64 GB200 nodes. The image must
+include the expected `/opt/venv` runtime, CUDA/SM100 support, and the reviewed
+FA4 dependencies. Use a separate AMD64 image digest for Modal. The Modal
+profile preserves H100 for encoder, custom-kernel, and VSA lanes, and L40S
+for the remaining lanes. Its reviewed image must support both SM89 and
+SM90a. Attention settings are selected per backend and lane to preserve the
+existing Modal and GB200 FA4 profiles; do not replace them with one global
+attention override. Validate the images against their respective hardware
+before enabling merge gates.
+
+Optional Kubernetes configuration includes `context`, `hf_secret`,
+`hf_secret_key`, `cache_pvc`, `cache_subpath`, `artifacts_pvc`, and
+`artifacts_subpath`. Only provide a read-only Hugging Face credential when
+private/gated downloads require it. PR workers are untrusted and can access
+any credential supplied to them, so never provide a reference-publication
+or other write-capable token. Prefer a pre-populated read-only cache.
+
+Use dedicated CI PVCs and relative CI subpaths. The personal
+`lustre-pvc-vllm` and `nfs-pvc-vllm` claims are rejected. The cache is mounted
+read-only; mutable references and locks stay inside the worker. The artifact
+init container prepares each Job's dedicated artifact subdirectory before
+the worker mounts it. An
+operator quota or admission policy scoped to CI can provide a second GPU
+ceiling; do not apply an eight-GPU quota to the shared `vllm` namespace if it
+would also cap other users' work.
+
+## Wire Every Buildkite Entry Pipeline
+
+Set each pipeline's operator-owned bootstrap command to
+`/opt/fastvideo-gpu-ci/upload`. Apply this to all three entry pipelines:
+
+| Pipeline | Trigger and scope |
+|---|---|
+| `pr-fastcheck` | Automatic PR webhook; `TEST_SCOPE=fastcheck` or unset. |
+| `ci` | Existing API triggers for Fastcheck, full, merge, direct, and scheduled SSIM. Keep its incoming PR webhook disabled to avoid duplicate builds. |
+| `fastvideo-performance-lane` | Existing schedule; `TEST_SCOPE=direct`, `TEST_TYPE=performance`. |
+
+The wrapper resolves `CI_GPU_BACKEND` from the build environment, falling
+back to `default_backend` in the trusted configuration. Allowed values are
+exactly `slurm`, `modal`, and `vllm`; unknown values fail. For Slurm, upload
+delegates to the unchanged private uploader. For Modal or `vllm`, it uploads
+one fixed `/opt/fastvideo-gpu-ci/run` command on the dedicated queue, with
+the selected backend and scope pinned in step environment.
+
+Set `CI_GPU_BACKEND=vllm` in the Buildkite build environment for a rack canary,
+or `modal` for a Modal canary. The existing API build payload can carry the
+same string in its `env` object. Automatic PR webhook builds inherit the
+operator's configured default unless pipeline/build configuration explicitly
+overrides it. Do not put backend selection inside a PR-controlled command.
+
+Keep the existing exact `BUILDKITE_COMMIT`, repository, PR identity, and
+`TEST_SCOPE` metadata. Direct runs also need an allowlisted `TEST_TYPE`.
+Merge runs need the trusted base-branch planner's `MERGE_TEST_PLAN`,
+`MERGE_GOLDEN_TESTS`, and `MERGE_SSIM_TESTS`. Full/direct/scheduled quality
+runs keep their complete matrices. The worker fetches and verifies the
+immutable commit before installing dependencies or invoking a lane script.
+
+The new Modal adapter uses bounded one-to-four-GPU sandboxes and the same
+lane scripts as Kubernetes. It does not reactivate `pr_test.sh` or the old
+Modal SSIM fan-out, which cannot enforce this shared four-GPU-per-PR budget.
+The legacy files remain available for their existing manual workflows.
+
+## Statuses and Default-Backend Cutover
+
+The selected adapter reports separate suite contexts:
+
+| Scope | GitHub context |
+|---|---|
+| Fastcheck | `gpu-ci/<backend>/fastcheck-passed` |
+| Merge or explicit full suite | `gpu-ci/<backend>/full-suite-passed` |
+| Direct lane | `gpu-ci/<backend>/direct-test-completed` |
+| Scheduled SSIM | `gpu-ci/<backend>/scheduled-ssim-passed` |
+
+The repository variable `CI_GPU_BACKEND` controls which new backend can
+publish the existing required `fastcheck-passed` and `full-suite-passed`
+contexts. Empty or `slurm` preserves the existing status behavior. `modal`
+or `vllm` enables `ci-gpu-backend-status.yml`, which reads the latest statuses
+for that one backend and mirrors both required contexts. A missing result
+becomes pending; failure and error remain failures. It never combines one
+backend's Fastcheck with another backend's full-suite result. Late status
+events re-read current state instead of replaying stale event payloads.
+
+Direct tests are diagnostic and do not promote a whole suite on the new
+backends. Rerun the matching Fastcheck or merge/full suite to clear its gate.
+Per-build tests on the other new backend remain separate diagnostics. The
+legacy direct-test aggregation workflow is disabled while a new backend is
+selected. These workflows retain the existing trust assumption that only
+authorized status-writing integrations can publish CI status contexts.
+
+For production cutover, drain existing Slurm builds: its preserved pipeline
+still emits canonical status contexts. The new uploader rejects a Slurm
+override when `default_backend` is `modal` or `vllm`, preventing later legacy
+builds from overwriting the promoted backend's checks. Ensure no old bootstrap
+bypasses the new uploader. Synchronize the operator `default_backend` and
+the GitHub repository variable, then clear or rerun required checks for all
+open PRs. Old green contexts do not become new-backend validation merely
+because a setting changed. Run both Fastcheck and the merge/full suite on
+the promoted backend before allowing merge. Apply the same drain and rerun
+procedure when rolling back to Slurm.
+
+## Validation and Recovery
+
+Inspect the rendered opt-in pipeline without submitting a build:
+
+```bash
+CI_GPU_BACKEND=vllm TEST_SCOPE=fastcheck \
+  /opt/fastvideo-gpu-ci/venv/bin/python -I \
+  /opt/fastvideo-gpu-ci/source/scripts/gpu_ci/entrypoint.py render \
+  --config /etc/fastvideo-gpu-ci.json
+```
+
+Validate a one-GPU lane, then a two-/four-GPU lane and the complete Fastcheck
+suite. Exercise two PRs plus a third waiter, concurrent builds of the same
+PR, cancellation, retries, and coordinator restart. Confirm observed GPU
+reservations never exceed two PRs, four per PR, or eight in total. A GPU
+reservation includes a pending worker, so unavailable nodes cannot cause
+the controller to submit more work than its allowance.
+
+Run SSIM, training, and performance canaries separately. References must
+match the effective GPU/runtime/attention backend; do not silently reuse
+L40S performance results as GB200 baselines or reseed references as part of
+routine CI. Workers keep W&B offline and disable reference publication.
+
+Inspect the persistent ledger and recover an abandoned attempt with:
+
+```bash
+/opt/fastvideo-gpu-ci/venv/bin/python -I \
+  /opt/fastvideo-gpu-ci/source/scripts/gpu_ci/entrypoint.py status \
+  --config /etc/fastvideo-gpu-ci.json
+
+/opt/fastvideo-gpu-ci/venv/bin/python -I \
+  /opt/fastvideo-gpu-ci/source/scripts/gpu_ci/entrypoint.py recover \
+  --config /etc/fastvideo-gpu-ci.json --build-id BUILD_ID.JOB_ID
+```
+
+Use the exact ledger ID from `status`; each Buildkite retry has a distinct
+job ID. Recovery refuses a live coordinator, persists cancellation, stops
+owned resources, and releases reservations only after confirming termination.
+If creation or cleanup is ambiguous, investigate the recorded handle on its
+backend and retain the reservation until the outcome is known. Do not delete
+the database or manually zero counters to unblock the queue.
+
+The dispatcher uploads controller logs, per-lane numeric results, request
+metadata, and the suite summary to Buildkite. These are the sources for its
+exit status. Generated videos and JUnit files stay in the worker unless a
+dedicated Kubernetes artifact PVC is configured; automatic publication of
+those worker files and Modal worker artifacts is not implemented. This is
+a deployment limitation to account for before replacing existing artifact
+review workflows.
@@ -5,7 +5,7 @@ hide:

 # Cosmos recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="cosmos" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="cosmos" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # FLUX recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="flux" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="flux" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # GLM-Image recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="glm_image" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="glm_image" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # Hunyuan recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="hunyuan" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="hunyuan" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -41,7 +41,7 @@ hide:
        <span class="cookbook-family-tile__footer">
          <span class="cookbook-family-tile__footer-top">
            <span><strong>MiniMax H3</strong><small>Video + stereo audio</small></span>
-            <span class="cookbook-count">8 recipes</span>
+            <span class="cookbook-count">9 recipes</span>
          </span>
          <ul class="cookbook-mode-row">
            <li>T2VA</li>
@@ -5,7 +5,7 @@ hide:

 # Kandinsky 5 recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="kandinsky5" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="kandinsky5" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # LongCat recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="longcat" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="longcat" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # LTX recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="ltx2" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="ltx2" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # Matrix Game recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="matrixgame" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="matrixgame" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,13 @@ hide:

 # MiniMax H3 recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="minimax_h3" data-default-recipe="fasth3-preview-cuda" data-recipes="../../assets/cookbook-recipes.json?v=8">
+FastH3 is two distilled MiniMax-H3 checkpoints. **V1** is the four-step
+launch. Some Hub repo names still say Preview. That name is historical. V1 is
+a full model, not a demo. **V2** is the eight-step checkpoint. More forwards
+is why V2 is the higher-quality FastH3. The V2 schedule contract is in
+[FastH3 distilled checkpoint schedules](../inference/fasth3-distilled.md).
+
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="minimax_h3" data-default-recipe="fasth3-preview-cuda" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -16,7 +22,7 @@ hide:
        <p class="cookbook-eyebrow">Primary focus · Inference</p>
        <h2>MiniMax H3 recipes</h2>
        <p>Generate video and audio with H3. Run a server on CUDA, one DGX Spark, or Apple Silicon MLX to iterate on prompts, or call the pipeline directly from Python.</p>
-        <span class="cookbook-count" data-cookbook-count>8 maintained recipes</span>
+        <span class="cookbook-count" data-cookbook-count>9 maintained recipes</span>
      </div>
    </div>
    <div class="cookbook-lifecycle" aria-label="Lifecycle stages">
@@ -41,9 +47,9 @@ hide:
    <h2 id="h3-modes-heading">Supported modes</h2>
    <p>
      CUDA covers T2VA, FL2VA, and Ref2VA on the full checkpoint, plus FastH3
-      Preview and FastH3 LoRA. FastH3 Preview also has a DGX Spark runtime with
-      a 1-Spark or 2-Spark device row. MLX is T2VA only. Temporal <code>--fast</code>,
-      spatial <code>--fast-spatial</code>, and opt-in VSA are flags on the same
+      V1 and FastH3 V2. FastH3 V1 also has a DGX Spark runtime with
+      a 1-Spark or 2-Spark device row. MLX is T2VA only: V1 and V2.
+      Temporal <code>--fast</code>, spatial <code>--fast-spatial</code>, and opt-in VSA are flags on the same
      MLX script, not extra recipes.
    </p>
    <div class="cookbook-modes__table-wrap">
@@ -58,8 +64,8 @@ hide:
        <tbody>
          <tr>
            <td>T2VA</td>
-            <td>Full H3, FastH3 Preview, FastH3 LoRA</td>
-            <td>FastH3 Preview after a local DiT conversion</td>
+            <td>Full H3, FastH3 V1, FastH3 LoRA, FastH3 V2</td>
+            <td>FastH3 V1 or FastH3 V2 after a local DiT conversion</td>
          </tr>
          <tr>
            <td>FL2VA</td>
@@ -93,7 +99,7 @@ hide:
          </tr>
          <tr>
            <td>DGX Spark</td>
-            <td>FastH3 Preview on one GB10, or two Sparks with Ray sequence parallel (<code>sp_size=2</code>) over QSFP RoCE. Select NVIDIA DGX Spark, then 1 Spark or 2 Sparks.</td>
+            <td>FastH3 V1 on one GB10, or two Sparks with Ray sequence parallel (<code>sp_size=2</code>) over QSFP RoCE. Select NVIDIA DGX Spark, then 1 Spark or 2 Sparks.</td>
            <td>Not wired</td>
          </tr>
        </tbody>
@@ -251,7 +257,7 @@ cd FastVideo</code></pre>
      <pre><code>UV_TORCH_BACKEND=cu130 uv pip install -e ".[fasth3]"</code></pre>
      <p class="cookbook-eyebrow">Apple Silicon</p>
      <pre><code>uv pip install -e ".[mlx]"</code></pre>
-      <p>Follow the <a href="../../getting_started/installation/mps/#run-fasth3-preview">Apple Silicon guide</a> for the download, conversion, and storage requirements.</p>
+      <p>Follow the <a href="../../getting_started/installation/mlx/">MLX install guide</a> for the extra, <code>ffmpeg</code>, and a clone. Then pick FastH3 V1 or V2 in the builder above.</p>
      <p class="cookbook-eyebrow">NVIDIA DGX Spark</p>
      <pre><code>UV_TORCH_BACKEND=cu130 uv pip install -e .</code></pre>
      <p>Follow the <a href="../../getting_started/installation/spark/">DGX Spark install guide</a> for ARM64 CUDA 13. One Spark is a local process. Two Sparks need Ray on the QSFP link:</p>
@@ -266,10 +272,10 @@ cd FastVideo</code></pre>
      <ul>
        <li>The full CUDA H3 examples request four GPUs by default. Their sources do not claim a GPU model or memory minimum.</li>
        <li>The FastH3 CUDA performance profile was measured on four GB200 GPUs. Use its strict profile when exact operation order matters more than the measured performance configuration.</li>
-        <li>The MLX source runtime supports T2VA, optional temporal <code>--fast</code>, optional spatial <code>--fast-spatial</code>, and opt-in VSA on <code>--include-vsa</code> checkpoints. FL2VA, Ref2VA, and two-pass refinement are not wired.</li>
+        <li>The MLX source runtime supports T2VA, optional temporal <code>--fast</code>, optional spatial <code>--fast-spatial</code>, and opt-in VSA on <code>--include-vsa</code> checkpoints. FastH3 V2 MLX converts with <code>--include-vsa</code> and runs eight forwards. FL2VA, Ref2VA, and two-pass refinement are not wired.</li>
        <li>GPU count and VAE decode backend are configurable in the builder above for FastH3 CUDA recipes. Only the value shown by default has a recorded run; other supported values are unmeasured here.</li>
-        <li>DGX Spark is a runtime on FastH3 Preview, not a separate family card. Select NVIDIA DGX Spark, then 1 Spark or 2 Sparks. The CUDA GPU-count knob does not apply to Spark.</li>
-        <li>GB10 has no FA4 / sm_100a VSA kernel. Keep <code>FASTVIDEO_FA4=0</code> and <code>FASTVIDEO_VSA_SM100A=0</code>. Legal <code>num_frames</code> values are <code>17n+5</code>, capped at 345 (15 s). A 345-frame request on one Spark can OOM.</li>
+        <li>DGX Spark is a runtime on FastH3 V1, not a separate family card. Select NVIDIA DGX Spark, then 1 Spark or 2 Sparks. The CUDA GPU-count knob does not apply to Spark.</li>
+        <li>GB10 has no FA4 / sm_100a VSA kernel. Keep <code>FASTVIDEO_FA4=0</code> and <code>FASTVIDEO_VSA_SM100A=0</code>. Legal <code>num_frames</code> values are <code>17n+5</code>, capped at 362 (15.08 s). A 345-frame request on one Spark can OOM.</li>
        <li>Gated or missing checkpoints: run <code>huggingface-cli login</code> and confirm you accepted the model's license on Hugging Face.</li>
      </ul>
  </div>
@@ -5,7 +5,7 @@ hide:

 # MMAudio recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="mmaudio" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="mmaudio" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -8,8 +8,9 @@ JavaScript client examples use that interface, but requests go to FastVideo.
 You do not need an OpenAI account or cloud key.

 The [H3 recipe selector](minimax-h3.md) provides the same workflow with runtime
-selection. This guide covers FastH3 Preview text-to-video/audio. Other H3
-recipes keep their direct Python commands.
+selection. This guide covers FastH3 V1 (four forwards) and FastH3 V2
+(eight forwards). Start a server, then iterate in the playground or with the
+OpenAI Python client. Other H3 recipes keep their direct Python commands.

 CUDA requests reuse one loaded `VideoGenerator`. The Python SDK can do the same
 when you reuse the generator across `generate()` calls. MLX keeps one
@@ -36,6 +37,14 @@ advertises it as `fasth3`. It configures four CUDA GPUs but does not record
 a GPU model or VRAM requirement. This is a source-backed server profile, not
 the measured GB200 Python performance profile. Compilation is disabled.

+For FastH3 V2, use the same install and the 8-step config (nine sigma
+points, eight DiT forwards, VSA sparsity 0.8):
+
+```bash
+UV_TORCH_BACKEND=cu130 uv pip install -e ".[fasth3]"
+fastvideo serve --config examples/serving/openai_fasth3_8step.yaml --server.host 127.0.0.1
+```
+
 Keep the server running. In another terminal, check readiness:

 ```bash
@@ -58,7 +67,7 @@ FASTVIDEO_VSA_SM100A=0 FASTVIDEO_FA4=0 FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_
 The configuration loads `FastVideo/FastVideo-FastH3-4-step-Preview-v1-VSA-DataFree`
 on one GB10 and advertises it as `fasth3`. Lazy module load still reloads
 Qwen3-VL and the DiT between phases of each request. Legal `num_frames` values
-are `17n+5`, capped at 345 (15 s); a 345-frame request on one Spark can OOM.
+are `17n+5`, capped at 362 (15.08 s); a 345-frame request on one Spark can OOM.
 There is no cookbook server for two Sparks; use the generate YAML after
 [pairing two Sparks](../getting_started/installation/spark_pair.md).

@@ -72,7 +81,7 @@ After model loading completes, the response is `{"status":"ok"}`.

 ### Apple Silicon MLX

-Complete the [Apple Silicon installation](../getting_started/installation/mps.md#run-fasth3-preview),
+Complete the [MLX install](../getting_started/installation/mlx.md),
 including `ffmpeg`. From your FastVideo clone, install the MLX extra:

 ```bash
@@ -111,6 +120,22 @@ The MLX server has no recorded device or unified-memory requirement. The
 direct Python recipe's M4 Max measurements are not a server benchmark or a
 minimum-memory claim.

+For FastH3 V2, convert that checkpoint's transformer with `--include-vsa`
+and start the 8-step config. Do not overwrite a V1 export. The
+[MiniMax H3 cookbook](minimax-h3.md) has the same download and convert
+commands as the Python recipe.
+
+```bash
+hf download FastVideo/FastVideo-FastH3-8-Step-V2 --local-dir ./FastH3-8-Step-V2
+python scripts/checkpoint_conversion/convert_minimax_h3_mlx.py --model-root ./FastH3-8-Step-V2/transformer --out ./FastH3-8-Step-V2-MLX --formats "int8" --include-vsa
+python -m fastvideo.entrypoints.openai.mlx_server --config examples/serving/mlx_fasth3_8step.yaml
+```
+
+That adapter passes `num_steps=8` when the HTTP field is `num_inference_steps=9`,
+and it enables the trained VSA recipe (sparsity 0.8, 64-token tiles). Reuse the
+preview VAE, audio VAE, text encoder, and tokenizer if those directories already
+exist; edit the YAML paths if your files live elsewhere.
+
 ## Open the playground

 Open [the local H3 playground](http://127.0.0.1:8000/playground/) after startup.
@@ -136,9 +161,11 @@ or manage a GPU server for you.
 ## Generate with cURL or an SDK

 These examples use the server's resolution, frame count, and sampling defaults.
-Do not copy Sora-specific durations or resolutions onto H3. Both server configs
-use 124 frames, 24 fps, and the five-point distilled sigma schedule with four
-DiT forwards. CUDA and one Spark use 1344 × 768; MLX uses 832 × 480.
+Do not copy Sora-specific durations or resolutions onto H3. V1 configs use
+124 frames, 24 fps, and the five-point distilled sigma schedule with four DiT
+forwards. V2 configs use nine sigma points and eight DiT forwards. CUDA
+and one Spark use 1344 × 768; MLX uses 832 × 480. The OpenAI Python client is
+the same for every FastH3 server that advertises `fasth3`.

 Each client submits a job, checks for completion or failure, and downloads an
 MP4 named after the job ID. Polling stops after 30 minutes; a timeout does not
@@ -5,7 +5,7 @@ hide:

 # Stable Audio recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="stable_audio" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="stable_audio" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # Stable Diffusion recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="sd35" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="sd35" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # TurboDiffusion recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="turbodiffusion" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="turbodiffusion" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -5,7 +5,7 @@ hide:

 # Wan recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="wan" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="wan" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -186,7 +186,7 @@ cd FastVideo</code></pre>
      <ul>
        <li>Out of memory on the A14B recipes: the checked-in sources already enable CPU offload; see <a href="../../inference/configuration/">Configuration</a> for the offload surface before reducing resolution or frames.</li>
        <li>The FastWan2.1 recipe requires <code>VIDEO_SPARSE_ATTN</code>; confirm the environment variable in the command was set in the same shell.</li>
-        <li>FastMetal MLX: install with <code>uv pip install -e ".[mlx]"</code>, then follow the <a href="../../getting_started/installation/mps/">Apple Silicon guide</a>. CUDA FastWan-QAD checkpoints are refused on the MLX runtime.</li>
+        <li>FastMetal MLX: install with the <a href="../../getting_started/installation/mlx/">MLX install guide</a>, then pick a FastMetal recipe in the builder. CUDA FastWan-QAD checkpoints are refused on the MLX runtime.</li>
        <li>FastMetal 5B uses <code>mlx_wan22_generate.py</code>. 1.3B and 14B use <code>mlx_wan_prompt_to_video.py</code>.</li>
        <li>Gated or missing checkpoints: run <code>huggingface-cli login</code> and confirm you accepted the model's license on Hugging Face.</li>
      </ul>
@@ -5,7 +5,7 @@ hide:

 # Z-Image recipes

-<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="zimage" data-recipes="../../assets/cookbook-recipes.json?v=8">
+<div class="cookbook-shell cookbook-family-page" data-cookbook data-family="zimage" data-recipes="../../assets/cookbook-recipes.json?v=11">
  <header class="cookbook-family-header">
    <a class="cookbook-back-link" href="../"><span aria-hidden="true">←</span> All model families</a>
    <div class="cookbook-family-header__body">
@@ -25,11 +25,13 @@ requests. HTTP handling and job polling remain asynchronous.
 | `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/generations` | OpenAI-compatible alias for image generation |
 | `POST` | `/v1/images/edits` | Generate an image from image references |
 | `GET` | `/v1/images/{id}/content` | Download a generated image |
 | `GET` | `/health` | Liveness probe |

-`POST /v1/videos/generations` remains an alias for older FastVideo clients.
+`POST /v1/videos/generations` remains an alias for older FastVideo clients, and
+`POST /v1/images/generations` is the OpenAI Python client's image generation path.
 The OpenAI Python and JavaScript clients can create, retrieve, list, download,
 and delete video jobs. Use the [H3 server cookbook](../../cookbook/openai-api.md)
 for pinned client versions and executable examples. Download variants other
@@ -7,7 +7,7 @@ FastVideo supports the following hardware platforms:
 - **NVIDIA DGX Spark / GB10 (ARM64 + CUDA 13)** — [install](installation/spark.md),
  [performance](installation/spark_performance.md),
  [pair two Sparks](installation/spark_pair.md)
- [Apple silicon](installation/mps.md)
+- [Apple silicon (MLX)](installation/mlx.md)

 ## Quick Installation

@@ -15,7 +15,7 @@ FastVideo supports the following hardware platforms:

 Use uv as the default environment manager for faster and more stable installs.
 The commands below target NVIDIA CUDA 12; use `UV_TORCH_BACKEND=cu130` on
-CUDA 13. Apple silicon users should follow the [MPS guide](installation/mps.md).
+CUDA 13. Apple silicon users should follow the [MLX install guide](installation/mlx.md).

 ```bash
 # Create and activate a new uv environment
@@ -0,0 +1,65 @@
+# Install FastVideo with MLX
+
+Install FastVideo on a Mac, then generate from the cookbook. Local video uses
+the native MLX runtime, not CUDA, and not the old PyTorch MPS demo at
+`examples/inference/basic/basic_mps.py`.
+
+## Requirements
+
+- macOS 14 or newer
+- Python 3.12
+- `ffmpeg` (`brew install ffmpeg`)
+
+## Install
+
+Cookbook commands run from a clone.
+
+```bash
+git clone https://github.com/hao-ai-lab/FastVideo.git && cd FastVideo
+uv venv --python 3.12 --seed
+source .venv/bin/activate
+brew install ffmpeg
+uv pip install -e ".[mlx]"
+```
+
+Conda is optional. After you activate a Conda env, still install with
+`uv pip` as above.
+
+`uv pip install "fastvideo[mlx]"` from PyPI installs the extra only. It does
+not ship the example scripts the cookbook copies.
+
+## Generate a video
+
+Open the cookbook. Select Apple Silicon as the runtime. Each recipe has a
+Python command. FastH3 also has a server path for the playground and the
+OpenAI Python client.
+
+- [Wan recipes](../../cookbook/wan.md) for FastMetal 1.3B, 5B, and 14B
+- [MiniMax H3 recipes](../../cookbook/minimax-h3.md) for FastH3 V1 and FastH3 V2
+- [H3 server guide](../../cookbook/openai-api.md) for the playground, cURL, and SDKs
+
+FastH3 is two distilled MiniMax-H3 checkpoints. V1 is the four-step launch.
+Some Hub repo names still say Preview. That name is historical. V1 is a full
+model, not a demo. V2 is the eight-step checkpoint. More forwards is why V2
+is the higher-quality FastH3.
+
+Recorded shapes and evidence live in the
+[support matrix](../../inference/support_matrix.md#apple-silicon-native-runtime).
+
+## Hardware
+
+- FastMetal 1.3B and 5B: 16 GB unified memory and up
+- FastMetal 14B: 36 GB unified memory and up
+- FastH3 V1 and V2: validated on an M4 Max with 36 GB unified memory
+
+## Troubleshooting
+
+- **`basic_mps.py` is the wrong path.** That script is PyTorch MPS. Use an
+  Apple Silicon recipe in the cookbook.
+- **Muxing fails.** Install `ffmpeg` with Homebrew.
+- **A cookbook command cannot find a script.** Run it from the FastVideo
+  clone after `uv pip install -e ".[mlx]"`.
+
+If that does not match what you see, open an issue on the
+[GitHub repository](https://github.com/hao-ai-lab/FastVideo) or ask in the
+[Slack community](https://join.slack.com/t/fastvideo/shared_invite/zt-3f4lao1uq-u~Ipx6Lt4J27AlD2y~IdLQ).
@@ -1,260 +0,0 @@
-# MPS (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
-
- **OS: macOS 14 or newer**
- **Python: 3.12.4**
-
-## Set up using Python
-
-### Create a new Python environment
-
-#### uv
-Recommended default: use [uv](https://docs.astral.sh/uv/) for faster and more stable environment setup.
-
-Please follow the [documentation](https://docs.astral.sh/uv/#getting-started) to install `uv`. After installing `uv`, create a new environment using:
-
-```console
-# (Recommended) Create a new uv environment. Use `--seed` to install `pip` and `setuptools`.
-uv venv --python 3.12 --seed
-source .venv/bin/activate
-```
-
-#### Conda (alternative)
-
-You can also create a Python environment using [Conda](https://docs.conda.io/projects/conda/en/stable/user-guide/getting-started.html).
-
-##### 1. Install Miniconda (if not already installed)
-
-```bash
-wget https://repo.anaconda.com/miniconda/Miniconda3-latest-MacOSX-arm64.sh
-bash Miniconda3-latest-MacOSX-arm64.sh
-source ~/.zshrc
-```
-
-##### 2. Create and activate a Conda environment for FastVideo
-
-```bash
-conda create -n fastvideo python=3.12.4 -y
-conda activate fastvideo
-```
-
-### Dependencies
-
-```
-brew install ffmpeg
-```
-
-### Installation
-
-FastMetal's native Apple Silicon runtime requires the `mlx` extra.
-
-#### With uv (recommended)
-
-```bash
-uv pip install "fastvideo[mlx]"
-```
-
-#### With Conda environment (alternative)
-
-`uv` works inside an active conda env too, so prefer `uv pip` for the actual install:
-
-```bash
-uv pip install "fastvideo[mlx]"
-```
-
-### Installation from Source
-
-#### 1. Clone the FastVideo repository
-
-```bash
-git clone https://github.com/hao-ai-lab/FastVideo.git && cd FastVideo
-```
-
-#### 2. Install FastVideo
-
-Basic installation:
-
-```bash
-uv pip install -e ".[mlx]"
-```
-
-Alternative with Conda environment:
-
-```bash
-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"
-```
-
-Dense conversion drops the trained VSA gate projections. To keep them (INT6
-weight-only, same affine grid as the other linear matrices) write a **new**
-directory:
-
-```bash
-python scripts/checkpoint_conversion/convert_minimax_h3_mlx.py \
-  --model-root ./FastH3-Preview-v0.2/transformer \
-  --out ./FastH3-MLX-vsa \
-  --formats "int6" \
-  --include-vsa
-```
-
-Do not overwrite an existing dense export such as `./FastH3-MLX/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
-```
-
-Add `--fast-spatial` for spatial fast mode, `--fast`'s spatial twin. It
-denoises and decodes on the smallest 32px-aligned canvas covering the
-requested size divided by `--fast-spatial-scale` (a 480x832 request runs on a
-256x416 canvas), then resamples the decoded frames up to the requested size
-in pixel space. It composes with `--fast`. This is a speed/quality trade-off
-and stays off by default: the output carries the reduced canvas's detail
-budget, so it reads softer than a native-resolution render, with the unsharp
-pass countering some but not all of the difference:
-
-```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 fastest.</d>" \
-  --height 480 --width 832 --num-frames 124 --seed 2028 \
-  --fast --fast-spatial \
-  --output-path ./outputs/fasth3_int6_fast_spatial.mp4
-```
-
-VSA is off by default. A dense-only checkpoint (no `--include-vsa`) keeps the
-existing fused-SDPA path. After converting with `--include-vsa`, enable the
-sparse path explicitly:
-
-```bash
-python examples/inference/basic/mlx_fasth3.py \
-  --model-root ./FastH3-Preview-v0.2 \
-  --mlx-checkpoint ./FastH3-MLX-vsa/int6 \
-  --vsa --vsa-sparsity 0.9 --vsa-tile-size 64 --vsa-prefix-mode exempt \
-  --prompt "(S1) A presenter says <d>[English] Fast H3 is amazing.</d>" \
-  --height 720 --width 1280 --num-frames 124 --seed 2026 \
-  --output-path ./outputs/fasth3_int6_vsa_720p.mp4
-```
-
-`--vsa-impl auto` uses the chunked gather+SDPA **reference** path.
-`--vsa-impl simd` is an opt-in SIMD-group kernel (tile 64, head dim 128) that
-falls back to reference on unsupported shapes. It is not the default.
-`--vsa-impl reference` is the same as `auto`.
-
-!!! note "Current MLX scope"
-    This source runtime supports T2VA, temporal `--fast`, spatial
-    `--fast-spatial`, and opt-in VSA. FL2VA, Ref2VA, two-pass refinement, and
-    `VideoGenerator` registry dispatch are not wired yet. INT8/INT6/INT4
-    quantization is **weight-only**; VSA attention Q/K/V stay BF16. Old dense
-    MLX checkpoints remain valid for dense inference and raise a reconvert
-    error if `--vsa` is set. 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:
-[Contributor Guide](../../contributing/overview.md)
-
-## Hardware Requirements
-
- **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
-
-If you encounter any issues during installation, please open an issue on our [GitHub repository](https://github.com/hao-ai-lab/FastVideo).
-
-You can also join our [Slack community](https://join.slack.com/t/fastvideo/shared_invite/zt-3f4lao1uq-u~Ipx6Lt4J27AlD2y~IdLQ) for additional support.
@@ -147,7 +147,7 @@ that can't move here.
 Two Sparks with QSFP cables: [Pair two NVIDIA DGX Sparks](spark_pair.md) for
 one FastH3 clip across both GPUs (`sp_size=2` over Ray). Copy-paste commands
 for one or two Sparks also live on the
-[MiniMax H3 cookbook](../../cookbook/minimax-h3.md): pick FastH3 Preview,
+[MiniMax H3 cookbook](../../cookbook/minimax-h3.md): pick FastH3 V1,
 then NVIDIA DGX Spark, then 1 Spark or 2 Sparks.

 ## Development Environment Setup
@@ -7,7 +7,7 @@ by the QSFP ConnectX-7 cables can run **one clip faster** and can hold a
 **longer clip** (up to the FastH3 15 s cap).

 Copy-paste commands for both counts are on the
-[MiniMax H3 cookbook](../../cookbook/minimax-h3.md): FastH3 Preview → NVIDIA DGX
+[MiniMax H3 cookbook](../../cookbook/minimax-h3.md): FastH3 V1 → NVIDIA DGX
 Spark → 1 Spark or 2 Sparks.

 This is FastVideo sequence parallel (`sp_size=2`) over Ray, not a third-party
@@ -109,7 +109,7 @@ Run the driver on the **head**, same venv, same QSFP IP.
 `basic_fasth3.py` defaults target a four-GPU GB200 profile: 768×1344, `sm100a`
 VSA, FA4, four GPUs. On Sparks you must override the kernel flags. Height,
 width, frames, steps, seed, and prompt are yours. Change them. Legal
-`num_frames` values are `17n+5`, capped at 345.
+`num_frames` values are `17n+5`, capped at 362.

 GB10 has no FA4 / sm_100a VSA kernel, so `--vsa-kernel triton --no-fa4` stays
 required on this box. `--execution-backend ray` is optional when `RAY_ADDRESS`
@@ -152,8 +152,8 @@ Stop the cluster when you are done: `ray stop` on both nodes.
 ## FastH3 frame counts

 H3 is 24 fps. Legal `num_frames` values are `17n+5`. The pipeline rejects
-clips longer than **15 s**. The longest legal length is **345 frames**
-(14.375 s). 360 frames aligns to 362 and fails the duration check.
+clips longer than **15 s**. The longest legal length is **362 frames**
+(15.083 s). 360 frames aligns to 362 and is accepted.

 ## Measured on two GB10s (2026-08-31)

@@ -179,7 +179,7 @@ VAE, same 4-step schedule:
 | Two Sparks, SP=2 | 2 | 124 | **215.2 s** | 72.4 s |

 Those medians used `--height` / `--width` / `--num-frames` as CLI flags. Swap
-them. Native 480p on this model is 480×832, 124 frames. The 15 s cap is 345
+them. Native 480p on this model is 480×832, 124 frames. The 15 s cap is 362
 frames.

 The first VAE decode still pays `torch.compile`. Later `generate()` calls in
@@ -181,6 +181,9 @@ is power-cycled. To avoid it:
  and frame counts are valid. Weights stay replicated, so lazy module load
  (auto on GB10) is still required on each box. Bring-up and knobs:
  [Pair two NVIDIA DGX Sparks](spark_pair.md).
+- A worker's SIGTERM log and traceback show where it was interrupted, not why it
+  was selected; confirm the cause in the `earlyoom` service or system logs. A
+  later SIGKILL or kernel OOM kill cannot be caught and reported by Python.

 ## Gotchas specific to the GB10

@@ -0,0 +1,90 @@
+# FastH3 distilled checkpoint schedules
+
+Base MiniMax-H3 still uses the scheduler shifts in its checkpoint (video 12,
+audio 3), BF16 text encoding, and the existing uniform schedule. The default
+`basic_fasth3.py` example still targets the four-forward preview. Selecting a
+shift-10 eight-forward checkpoint is an explicit choice of model and recipe;
+it does not change either default or enable NVFP4.
+
+## Eight-forward T2AV recipe
+
+The public checkpoint is
+[`FastVideo/FastVideo-FastH3-8-Step-V2`](https://huggingface.co/FastVideo/FastVideo-FastH3-8-Step-V2)
+(MiniMax H3 Community License), trained with video/audio shifts 10/3, VSA
+sparsity 0.8, 64-token tiles, and the DMD rungs
+`[999, 874, 749, 624, 500, 375, 250, 125]`. `basic_fasth3_8step.py` pins that
+checkpoint and recipe as defaults; it shares the preview example's CLI, so every
+other flag works unchanged:
+
+```bash
+python examples/inference/basic/basic_fasth3_8step.py \
+  --prompt 'A slow cinematic drone shot glides over a coastal town; gulls call over the harbor.' \
+  --num-gpus 4 --vsa-kernel sm100a \
+  --profile strict --no-inference-torch-compile --no-compile-vae \
+  --height 768 --width 1344 --num-frames 124 \
+  --output outputs/fasth3-8step
+```
+
+Pass `--model-path` to use a local snapshot of the full export (not just its
+`transformer` subdirectory). `--steps` is the number of sigma-grid points,
+including the terminal zero; nine points run exactly eight transformer
+forwards, and the script rejects any other value because the checkpoint's
+ladder has eight rungs. The rungs are unshifted noise levels on the 1000-step training
+clock; each scheduler applies its own shift once, and the transformer receives
+H3 clean-time values (`1 - sigma`). A uniform nine-point grid is not a substitute
+for those rungs.
+
+Compilation and H3 fusions are disabled above to establish an eager reference;
+they can be evaluated separately. On hardware without the sm100a extension,
+use `--vsa-kernel triton`; compare outputs and performance before adopting that
+backend. This recipe is T2AV-only, not a distilled `transformer_ref` model.
+
+## Export metadata and validation
+
+The export's `fastvideo_inference.json` supplies the trained ladder. The
+schedule fields of `fasth3-inference-contract-v1` are:
+
+```json
+{
+  "schema_version": "fasth3-inference-contract-v1",
+  "dmd_denoising_steps": [999, 874, 749, 624, 500, 375, 250, 125],
+  "num_inference_steps": 9,
+  "transformer_forwards": 8,
+  "video_scheduler_shift": 10.0,
+  "audio_scheduler_shift": 3.0
+}
+```
+
+The loader keeps this file when downloading the selected H3 components from
+Hugging Face. It checks that the two declared shifts agree with
+`scheduler/scheduler_config.json` and `audio_scheduler/scheduler_config.json`.
+Missing/invalid rungs, inconsistent counts, or an explicit conflicting ladder
+are errors. The denoiser rejects a request with the wrong number of grid points.
+The metadata does not silently change request dimensions, step count, attention
+backend, sparsity, precision, or offload/compile settings: set those explicitly
+as above.
+
+For exports without this sidecar, an explicit ladder is supported via
+`MiniMaxH3PipelineConfig.dmd_denoising_steps`, or through the typed API's
+`PipelineSelection(experimental={"dmd_denoising_steps": [...]})`. The shifts
+still come from the checkpoint scheduler configs. Keep generic `flow_shift`
+unset: H3 has separate video and audio shifts, not one shared shift.
+
+This documents execution support for the published checkpoint. It is not a
+quality claim: compare video/audio output against base MiniMax-H3 on your own
+prompts before adopting it.
+
+## Apple Silicon
+
+`mlx_fasth3.py` stays on FastH3 V1 and its uniform AdaLN cache.
+`mlx_fasth3_8step.py` is the eight-forward MLX recipe for FastH3 V2. It
+reads the same `fastvideo_inference.json` rungs and shifts, and it expects an
+MLX DiT whose AdaLN cache was converted from that contract. Reuse the V1
+snapshot's VAE, audio VAE, text encoder, and tokenizer; only the DiT and the
+sidecar change. Rank-reduced AdaLN checkpoints are unchanged and are not
+produced by the MLX converter.
+
+Install is in the
+[MLX install guide](../getting_started/installation/mlx.md). Generation and
+serving are in the [MiniMax H3 cookbook](../cookbook/minimax-h3.md) and the
+[H3 server guide](../cookbook/openai-api.md).
@@ -5,10 +5,10 @@ This page contains step-by-step instructions to get you quickly started with vid
 ## Requirements

 - **OS**: Linux (tested on Ubuntu 22.04+), or macOS on Apple silicon via the
-  [MPS installation guide](../getting_started/installation/mps.md)
+  [MLX install guide](../getting_started/installation/mlx.md)
 - **Python**: 3.10-3.12
 - **CUDA**: 12.6 or 13.0 (NVIDIA GPUs)
- **GPU**: At least one NVIDIA GPU, or an Apple silicon chip with MPS
+- **GPU**: At least one NVIDIA GPU, or an Apple silicon chip with the MLX runtime

 ## Installation

@@ -10,7 +10,9 @@ that tradeoff. It is not a lossless acceleration of the full VAE.

 ## Generate a video

-Use your existing MLX FastH3 environment and converted checkpoint:
+Use your existing MLX FastH3 environment and converted checkpoint. The same
+`--video-decode-backend taeh3` flag works on `mlx_fasth3.py` (V1) and
+`mlx_fasth3_8step.py` (V2):

 ```bash
 python examples/inference/basic/mlx_fasth3.py \
@@ -294,6 +294,9 @@ automatically.
 ### Requirements

 - **GPU**: sm89+ (H100, L40S, RTX 4090, or newer) for hardware FP8 compute
+- **ROCm**: CDNA4 (MI350X / MI355X, gfx950) runs the FP8 `_scaled_mm` path through
+  hipBLASLt (OCP e4m3fn). MI300X (gfx942) only exposes the `fnuz` FP8 formats and
+  takes the bf16 dequant fallback like a pre-sm89 GPU.
 - No additional packages required beyond the base FastVideo install

 ### Usage
@@ -185,13 +185,14 @@ optimizations: absence means **untested**, not incompatible.
 | 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 T2V 5B | [`FastVideo/FastMetal-5B-QAD`](https://huggingface.co/FastVideo/FastMetal-5B-QAD) | 480p / 720p, 81 frames, 3-step DMD, INT8 DiT + TAEHV decode; optional `--fast`, `--fast-spatial`, `--refine`. The checked-in example is T2V. CUDA Wan2.2 TI2V 5B is the image-capable path. | 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 **weight-only** DiT, native video + audio VAE; optional temporal RIFE fast mode; optional spatial fast mode; optional VSA (tile 64/256, exempt/compete) on `--include-vsa` checkpoints | Apple M4 Max, 36 GB unified memory | Source runtime; T2VA only |
+| MLX FastH3 V1 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 **weight-only** DiT, native video + audio VAE; optional temporal RIFE fast mode; optional spatial fast mode; optional VSA (tile 64/256, exempt/compete) on `--include-vsa` checkpoints | Apple M4 Max, 36 GB unified memory | Source runtime; T2VA only |
+| MLX FastH3 V2 T2VA | [`FastVideo/FastVideo-FastH3-8-Step-V2`](https://huggingface.co/FastVideo/FastVideo-FastH3-8-Step-V2) + locally converted DiT | 480p / 720p, 124 frames, 8-step DMD2, INT8 **weight-only** DiT with `--include-vsa`, trained VSA 0.8 / tile 64, native video + audio VAE; AdaLN cache from `fastvideo_inference.json` | Apple M4 Max, 36 GB unified memory | Source runtime; T2VA only |

 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`,
+release. FastH3 V1 and FastH3 V2 use 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
+[MLX install guide](../getting_started/installation/mlx.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.
@@ -223,8 +224,9 @@ Per the installation guides:
  the [DGX Spark install guide](../getting_started/installation/spark.md).
  Two Sparks over QSFP use Ray sequence parallel; see
  [Pair two NVIDIA DGX Sparks](../getting_started/installation/spark_pair.md).
- **Apple silicon** — macOS 14 or newer; FastMetal-QAD via the MLX runtime. See the
-  [Apple Silicon guide](../getting_started/installation/mps.md). The older
+- **Apple silicon** — macOS 14 or newer; MLX runtime for FastMetal-QAD, FastH3
+  V1, and FastH3 V2. See the
+  [MLX install guide](../getting_started/installation/mlx.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.

@@ -21,7 +21,7 @@ python examples/inference/basic/basic.py
 ### 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/).
+[MLX install guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mlx/).

 ```bash
 hf download FastVideo/FastMetal-1.3B-QAD --local-dir ./FastMetal-1.3B-QAD
@@ -39,7 +39,7 @@ with

 `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
+FastH3 V1 T2VA also runs through the native MLX runtime. Convert the DiT
 to INT8, INT6, or INT4 first, then run:

 ```bash
@@ -51,17 +51,21 @@ python examples/inference/basic/mlx_fasth3.py \
  --output-path ./outputs/fasth3_int6.mp4
 ```

+FastH3 V2 is a separate checkpoint and script. Convert with
+`--include-vsa` and run `mlx_fasth3_8step.py`. Do not reuse a V1 DiT.
+
 Pass `--fast` for temporal RIFE fast mode and `--fast-spatial` for spatial
 fast mode (reduced-canvas denoise + pixel-space upsample); the two compose.
-VSA is opt-in: convert with `--include-vsa` and pass `--vsa` (see the
-[Apple Silicon guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mps/)).
-This MLX entrypoint currently supports T2VA only; FL2VA, Ref2VA, and
+V1 VSA is opt-in: convert with `--include-vsa` and pass `--vsa`. V2
+turns VSA on by default. See the
+[MLX install guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mlx/).
+This MLX entrypoint supports T2VA only; FL2VA, Ref2VA, and
 two-pass refinement remain follow-up work. INT6/INT8/INT4 are
-weight-only; VSA attention activations stay BF16. Dense-only checkpoints keep
-working for dense inference.
+weight-only; VSA attention activations stay BF16. Dense-only V1
+checkpoints keep working for dense inference.

 The complete setup and conversion commands are in the
-[Apple Silicon guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mps/).
+[MLX install guide](https://hao-ai-lab.github.io/FastVideo/getting_started/installation/mlx/).

 For an example running DMD+VSA inference:
 ```
@@ -0,0 +1,55 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Eight-forward video+audio generation with the FastH3 8-Step V2 checkpoint.
+
+``FastVideo/FastVideo-FastH3-8-Step-V2`` is a DMD-distilled MiniMax H3 T2AV
+student trained with video/audio scheduler shifts 10/3, VSA sparsity 0.8 and
+64-token tiles, on the explicit rung ladder ``[999, 874, 749, 624, 500, 375,
+250, 125]``. Its ``fastvideo_inference.json`` sidecar carries that ladder; the
+pipeline loads it, checks the declared shifts against the checkpoint scheduler
+configs, and runs exactly eight transformer forwards from nine sigma-grid
+points. This script only sets those defaults; it does not change the
+four-forward preview example.
+
+``--steps`` must stay at 9 for this checkpoint: the trained ladder has eight
+rungs, and a different grid is not a valid recipe for it.
+"""
+
+from __future__ import annotations
+
+import argparse
+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]
+
+MODEL = "FastVideo/FastVideo-FastH3-8-Step-V2"
+GRID_POINTS = 9  # nine sigma-grid points = eight DiT forwards
+VSA_SPARSITY = 0.8
+
+
+def parse_args(argv: Sequence[str] | None = None) -> argparse.Namespace:
+    parser = basic_fasth3.build_parser(description=__doc__)
+    parser.set_defaults(
+        model_path=MODEL,
+        output="outputs/fasth3_8step",
+        steps=GRID_POINTS,
+        vsa_sparsity=VSA_SPARSITY,
+        vsa_tile_size=64,
+    )
+    args = basic_fasth3.validate_args(parser, parser.parse_args(argv))
+    if args.steps != GRID_POINTS:
+        parser.error(f"--steps must be {GRID_POINTS} for {MODEL}: the checkpoint's trained ladder has "
+                     f"{GRID_POINTS - 1} rungs and the pipeline rejects any other grid")
+    return args
+
+
+def main() -> None:
+    basic_fasth3.run(parse_args())
+
+
+if __name__ == "__main__":
+    main()
@@ -28,7 +28,8 @@ detail budget and reads softer than a native-resolution render, so it stays
 off by default.

 This entrypoint currently supports text-to-video-with-audio only. It does not
-yet wire FL2VA, Ref2VA, or two-pass refinement.
+yet wire FL2VA, Ref2VA, or two-pass refinement. The four-step preview stays
+the default. FastH3 8-Step V2 uses ``mlx_fasth3_8step.py``.

 VSA is off by default; existing dense MLX checkpoints remain supported.
 H3 uses fused MLX RMSNorm, which can change BF16 rounding compared with the
@@ -51,7 +52,7 @@ def _dense_layers(value: str) -> tuple[int, ...]:
    return layers


-def parse_args() -> argparse.Namespace:
+def build_parser() -> argparse.ArgumentParser:
    parser = argparse.ArgumentParser(description=__doc__,
                                     formatter_class=argparse.RawDescriptionHelpFormatter)
    parser.add_argument("--model-root", type=Path, default=Path.home() / "models/FastH3-Preview-v0.2",
@@ -164,11 +165,20 @@ def parse_args() -> argparse.Namespace:
            "(simd is opt-in)"
        ),
    )
-    return parser.parse_args()
+    parser.add_argument("--vae-tile-height", type=int, default=256,
+                        help="overlapping video-VAE tile height in pixels")
+    parser.add_argument("--vae-tile-width", type=int, default=256,
+                        help="overlapping video-VAE tile width in pixels")
+    parser.add_argument("--inter-step-cooldown-s", type=float, default=0.0,
+                        help="optional pause between denoise steps")
+    return parser


-def main() -> None:
-    args = parse_args()
+def parse_args(argv: list[str] | None = None) -> argparse.Namespace:
+    return build_parser().parse_args(argv)
+
+
+def run(args: argparse.Namespace) -> None:
    from fastvideo.mlx_runtime.minimax_h3_pipeline import MiniMaxH3MLXPipeline

    pipeline = MiniMaxH3MLXPipeline(
@@ -189,6 +199,9 @@ def main() -> None:
        seed=args.seed,
        num_steps=args.steps,
        tiled_video_decode=args.tiled_video_decode,
+        vae_tile_height=args.vae_tile_height,
+        vae_tile_width=args.vae_tile_width,
+        inter_step_cooldown_s=args.inter_step_cooldown_s,
        fast=args.fast,
        fast_factor=args.fast_factor,
        fast_sharpen=args.fast_sharpen,
@@ -215,5 +228,9 @@ def main() -> None:
    }, indent=2))


+def main() -> None:
+    run(parse_args())
+
+
 if __name__ == "__main__":
    main()
@@ -0,0 +1,43 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Eight-forward FastH3 8-Step V2 generation on Apple Silicon.
+
+``FastVideo/FastVideo-FastH3-8-Step-V2`` trains eight DMD forwards at video
+shift 10 and audio shift 3, VSA sparsity 0.8, and 64-token tiles. Point
+``--model-root`` at a snapshot that contains that checkpoint's
+``fastvideo_inference.json`` plus the shared VAE, audio VAE, text encoder, and
+tokenizer. Point ``--mlx-checkpoint`` at an MLX DiT converted from this
+checkpoint (the AdaLN cache must be the contract ladder, not the four-step
+uniform grid).
+
+``--steps`` accepts 8 transformer forwards or 9 sigma-grid points. Either
+value runs the trained rungs. The four-step preview entrypoint is unchanged.
+"""
+
+from __future__ import annotations
+
+from collections.abc import Sequence
+
+try:
+    from . import mlx_fasth3
+except ImportError:
+    import mlx_fasth3  # type: ignore[no-redef]
+
+FORWARDS = 8
+GRID_POINTS = FORWARDS + 1
+
+
+def parse_args(argv: Sequence[str] | None = None):
+    parser = mlx_fasth3.build_parser()
+    parser.set_defaults(steps=FORWARDS, vsa=True, vsa_sparsity=0.8, vsa_tile_size=64)
+    args = parser.parse_args(argv)
+    if args.steps not in (FORWARDS, GRID_POINTS):
+        parser.error(f"--steps must be {FORWARDS} or {GRID_POINTS} for FastH3 8-Step V2")
+    return args
+
+
+def main() -> None:
+    mlx_fasth3.run(parse_args())
+
+
+if __name__ == "__main__":
+    main()
@@ -37,7 +37,8 @@ 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.
+# of the form 17n + 5. 124 is the 5-second point the FastH3 profile is measured at;
+# the 15-second cap pads to 362 frames.
 FRAMES_PER_CHUNK = 17
 LATENTS_PER_CHUNK = 5
 FPS = 24
@@ -52,6 +53,10 @@ def align_num_frames(num_frames: int) -> int:
    return LATENTS_PER_CHUNK + chunks * FRAMES_PER_CHUNK


+MIN_ALIGNED_FRAMES = align_num_frames(int(MIN_DURATION * FPS))
+MAX_ALIGNED_FRAMES = align_num_frames(int(MAX_DURATION * FPS))
+
+
 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")
@@ -86,9 +91,10 @@ def parse_args(argv: Sequence[str] | None = None) -> argparse.Namespace:
    # 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")
+    if not MIN_ALIGNED_FRAMES <= aligned <= MAX_ALIGNED_FRAMES:
+        parser.error(f"MiniMax-H3 generates {MIN_DURATION:g}-{MAX_DURATION:g}s at {FPS} fps; "
+                     f"aligned num_frames={aligned} is {aligned / FPS:.1f}s, outside the accepted "
+                     f"{MIN_ALIGNED_FRAMES}-{MAX_ALIGNED_FRAMES} frame range")
    args.num_frames = aligned
    if not math.isfinite(args.lora_strength):
        parser.error("--lora-strength must be finite")
@@ -2,8 +2,10 @@

 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 FastH3 configs here are source-backed examples, not recorded
-serving benchmarks. Both configure four CUDA GPUs without a GPU model or VRAM claim.
+surface. The two CUDA FastH3 configs here are source-backed examples, not
+recorded serving benchmarks. Both configure four CUDA GPUs without a GPU model
+or VRAM claim. The Wan configs are copied from the working generate config or
+example for each checkpoint.

 Use the [H3 server cookbook](https://haoailab.com/FastVideo/cookbook/openai-api/)
 for the guided install, server, and client workflow.
@@ -39,10 +41,39 @@ and API clients share the same server; both can run locally.

 For native Apple Silicon MLX, use
 `python -m fastvideo.entrypoints.openai.mlx_server --config examples/serving/mlx_fasth3.yaml`
-after preparing the weights and editing their paths in that configuration.
-The same playground and clients work with MLX. Its pipeline still loads and
-releases components between phases to limit unified-memory use. See the
-cookbook for setup and the supported text-to-video/audio request fields.
+(Preview) or `examples/serving/mlx_fasth3_8step.yaml` (8-Step V2) after preparing
+the weights and editing their paths in that configuration. The same playground
+and OpenAI Python client work with MLX. Its pipeline still loads and releases
+components between phases to limit unified-memory use. See the cookbook for
+setup and the supported text-to-video/audio request fields.
+
+CUDA FastH3 8-Step V2 serving uses `examples/serving/openai_fasth3_8step.yaml`.
+
+## Wan configs
+
+Four Wan CUDA configs mirror the working generate config or example for each
+checkpoint:
+
+| Config | Model | GPUs | Served alias |
+| --- | --- | --- | --- |
+| `openai_fastwan21_1_3b.yaml` | FastWan2.1 T2V 1.3B (DMD, VSA) | 1 | `fastwan21-1.3b` |
+| `openai_wan21_i2v_14b.yaml` | Wan2.1 I2V 14B 480P | 2 | `wan21-i2v-14b` |
+| `openai_wan22_t2v_a14b.yaml` | Wan2.2 T2V A14B | 2 | `wan22-t2v-a14b` |
+| `openai_wan22_ti2v_5b.yaml` | Wan2.2 TI2V 5B | 1 | `wan22-ti2v-5b` |
+
+The 1.3B config needs the VSA attention backend:
+
+```bash
+FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_ATTN \
+  fastvideo serve --config examples/serving/openai_fastwan21_1_3b.yaml
+```
+
+The I2V 14B and TI2V 5B servers accept image-conditioned requests. Supply the
+source image through the OpenAI-compatible reference fields, for example
+`{"input_reference": "/path/to/first-frame.png"}` or
+`{"image_reference": [{"image_url": "https://example.com/first-frame.png"}]}`.
+The I2V 14B server requires one on every request; the TI2V 5B server treats
+its absence as text-to-video.

 Or submit, poll, and download with the OpenAI Python client:

@@ -53,7 +84,7 @@ python examples/serving/clients/video.py

 The `clients/` directory also contains cURL and OpenAI JavaScript examples.
 They default to model `fasth3`; set `FASTVIDEO_MODEL` to the advertised alias
-when using the LoRA config. Set `FASTVIDEO_BASE_URL` for another endpoint.
+when using the LoRA or Wan configs. Set `FASTVIDEO_BASE_URL` for another endpoint.
 These clients call your FastVideo server, not OpenAI's cloud. The SDK key
 `local` is a placeholder, not authentication. Keep the server on loopback or
 use an authenticated proxy for remote access.
@@ -0,0 +1,26 @@
+# Native MLX FastH3 8-Step V2 server. Convert the DiT with --include-vsa first.
+runtime: mlx
+generator:
+  model_path: FastVideo/FastVideo-FastH3-8-Step-V2
+  model_root: ./FastH3-8-Step-V2
+  mlx_checkpoint: ./FastH3-8-Step-V2-MLX/int8
+  prompt_cache_dir: ./outputs/h3_prompt_cache
+  vae_dtype: fp32
+  vsa: true
+  vsa_sparsity: 0.8
+  vsa_tile_size: 64
+server:
+  host: 127.0.0.1
+  port: 8000
+  served_model_name: fasth3
+  output_dir: outputs/mlx_fasth3_8step
+default_request:
+  sampling:
+    height: 480
+    width: 832
+    num_frames: 124
+    fps: 24
+    # HTTP convention: nine sigma points, eight DiT forwards. Native MLX uses num_steps=8.
+    num_inference_steps: 9
+    guidance_scale: 1.0
+    seed: 2026
@@ -0,0 +1,48 @@
+# OpenAI-compatible FastH3 8-Step V2 server.
+generator:
+  model_path: FastVideo/FastVideo-FastH3-8-Step-V2
+  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.8
+      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_8step
+  served_model_name: fasth3
+
+default_request:
+  negative_prompt: ""
+  sampling:
+    height: 768
+    width: 1344
+    num_frames: 124
+    fps: 24
+    num_inference_steps: 9
+    guidance_scale: 1.0
+    batch_cfg: false
+    seed: 1000
+  output:
+    return_frames: false
@@ -0,0 +1,44 @@
+# OpenAI-compatible server for FastWan2.1 T2V 1.3B (DMD-distilled, VSA sparse attention).
+# Requires: FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_ATTN
+#   FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_ATTN fastvideo serve --config examples/serving/openai_fastwan21_1_3b.yaml
+generator:
+  model_path: FastVideo/FastWan2.1-T2V-1.3B-Diffusers
+  engine:
+    num_gpus: 1
+    parallelism:
+      tp_size: 1
+      sp_size: 1
+    offload:
+      dit: false
+      vae: false
+      pin_cpu_memory: false
+    compile:
+      enabled: true
+  pipeline:
+    experimental:
+      VSA_sparsity: 0.8
+      dmd_denoising_steps: [1000, 757, 522]
+
+server:
+  host: 0.0.0.0
+  port: 8000
+  output_dir: outputs/openai_fastwan21_1_3b
+  served_model_name: fastwan21-1.3b
+
+default_request:
+  negative_prompt: >-
+    Bright tones, overexposed, static, blurred details, subtitles, style,
+    works, paintings, images, static, overall gray, worst quality, low
+    quality, JPEG compression residue, ugly, incomplete, extra fingers,
+    poorly drawn hands, poorly drawn faces, deformed, disfigured, misshapen
+    limbs, fused fingers, still picture, messy background, three legs, many
+    people in the background, walking backwards
+  sampling:
+    seed: 1024
+    num_frames: 81
+    height: 480
+    width: 832
+    fps: 16
+    num_inference_steps: 3
+  output:
+    return_frames: false
@@ -0,0 +1,48 @@
+# OpenAI-compatible server for Wan2.1 I2V 14B 480P. Each request must supply
+# its own source image through the OpenAI-compatible reference fields, e.g.
+# {"input_reference": "/path/to/first-frame.png"} or
+# {"image_reference": [{"image_url": "https://example.com/first-frame.png"}]};
+# there is no default image.
+#   fastvideo serve --config examples/serving/openai_wan21_i2v_14b.yaml
+generator:
+  model_path: Wan-AI/Wan2.1-I2V-14B-480P-Diffusers
+  engine:
+    num_gpus: 2
+    parallelism:
+      tp_size: 2
+      sp_size: 2
+    offload:
+      dit: true
+      vae: true
+      text_encoder: true
+      image_encoder: true
+      pin_cpu_memory: true
+  pipeline:
+    workload_type: i2v
+    experimental:
+      flow_shift: 3.0
+
+server:
+  host: 0.0.0.0
+  port: 8000
+  output_dir: outputs/openai_wan21_i2v_14b
+  served_model_name: wan21-i2v-14b
+
+default_request:
+  negative_prompt: >-
+    Bright tones, overexposed, static, blurred details, subtitles, style,
+    works, paintings, images, static, overall gray, worst quality, low
+    quality, JPEG compression residue, ugly, incomplete, extra fingers,
+    poorly drawn hands, poorly drawn faces, deformed, disfigured, misshapen
+    limbs, fused fingers, still picture, messy background, three legs, many
+    people in the background, walking backwards
+  sampling:
+    seed: 1024
+    num_frames: 77
+    height: 480
+    width: 832
+    fps: 16
+    num_inference_steps: 40
+    guidance_scale: 5.0
+  output:
+    return_frames: false
@@ -0,0 +1,29 @@
+# OpenAI-compatible server for Wan2.2 T2V A14B (mixture-of-experts DiT).
+#   fastvideo serve --config examples/serving/openai_wan22_t2v_a14b.yaml
+generator:
+  model_path: Wan-AI/Wan2.2-T2V-A14B-Diffusers
+  engine:
+    num_gpus: 2
+    use_fsdp_inference: false
+    parallelism:
+      tp_size: 1
+      sp_size: 2
+    offload:
+      dit: true
+      vae: false
+      text_encoder: true
+      pin_cpu_memory: true
+
+server:
+  host: 0.0.0.0
+  port: 8000
+  output_dir: outputs/openai_wan22_t2v_a14b
+  served_model_name: wan22-t2v-a14b
+
+default_request:
+  sampling:
+    num_frames: 81
+    height: 720
+    width: 1280
+  output:
+    return_frames: false
@@ -0,0 +1,27 @@
+# OpenAI-compatible server for Wan2.2 TI2V 5B. Supports both text-to-video
+# (omit the image reference) and image-to-video (set input_reference or
+# image_reference) from the same model.
+#   fastvideo serve --config examples/serving/openai_wan22_ti2v_5b.yaml
+generator:
+  model_path: Wan-AI/Wan2.2-TI2V-5B-Diffusers
+  engine:
+    num_gpus: 1
+    use_fsdp_inference: false
+    parallelism:
+      tp_size: 1
+      sp_size: 1
+    offload:
+      dit: true
+      vae: false
+      text_encoder: true
+      pin_cpu_memory: true
+
+server:
+  host: 0.0.0.0
+  port: 8000
+  output_dir: outputs/openai_wan22_ti2v_5b
+  served_model_name: wan22-ti2v-5b
+
+default_request:
+  output:
+    return_frames: false
@@ -27,7 +27,7 @@ layout) with:
 ```bash
 # from the repo root; reads data/kandinsky5_overfit and writes
 # data/kandinsky5_overfit_preprocessed by default -- override with
-# KANDINSKY5_OVERFIT_DATA_DIR / KANDINSKY5_OVERFIT_OUTPUT_DIR env vars
+# FASTVIDEO_TEST_KANDINSKY5_OVERFIT_DATA_DIR / FASTVIDEO_TEST_KANDINSKY5_OVERFIT_OUTPUT_DIR env vars
 python -m fastvideo.pipelines.preprocess.preprocess_kandinsky5_overfit
 ```

@@ -32,7 +32,7 @@
 # (see docs/training/data_preprocess.md;
 # fastvideo/pipelines/preprocess/preprocess_ltx2_overfit.py is a minimal
 # LTX-2 reference to adapt — set
-# LTX2_OVERFIT_MODEL=FastVideo/LTX-2.3-Distilled-Diffusers). Data paths,
+# FASTVIDEO_TEST_LTX2_OVERFIT_MODEL=FastVideo/LTX-2.3-Distilled-Diffusers). Data paths,
 # step counts, and learning rate below are placeholders you must adapt
 # to your workload.
 #
@@ -12,7 +12,7 @@
 # (see docs/training/data_preprocess.md;
 # fastvideo/pipelines/preprocess/preprocess_ltx2_overfit.py is a minimal
 # LTX-2 reference to adapt — set
-# LTX2_OVERFIT_MODEL=FastVideo/LTX-2.3-Distilled-Diffusers).
+# FASTVIDEO_TEST_LTX2_OVERFIT_MODEL=FastVideo/LTX-2.3-Distilled-Diffusers).
 #
 # Run:
 #   NUM_GPUS=4 bash examples/train/run.sh \
@@ -8,8 +8,8 @@
 #
 # Preprocess data first (writes data/ltx2_3_overfit_preprocessed):
 #   CUDA_VISIBLE_DEVICES=0 \
-#   LTX2_OVERFIT_MODEL=FastVideo/LTX-2.3-Distilled-Diffusers \
-#   LTX2_OVERFIT_OUTPUT_DIR=data/ltx2_3_overfit_preprocessed \
+#   FASTVIDEO_TEST_LTX2_OVERFIT_MODEL=FastVideo/LTX-2.3-Distilled-Diffusers \
+#   FASTVIDEO_TEST_LTX2_OVERFIT_OUTPUT_DIR=data/ltx2_3_overfit_preprocessed \
 #   python fastvideo/pipelines/preprocess/preprocess_ltx2_overfit.py
 #
 # Run:
@@ -16,14 +16,55 @@ if(NOT GPU_BACKEND)
    set(GPU_BACKEND "CUDA")
 endif()

-if(GPU_BACKEND STREQUAL "ROCM")
-    enable_language(HIP)
+# Compiled extension: AUTO/ON/OFF. On ROCm, AUTO falls back to the Python + Triton
+# package when HIP cannot be configured (#1850: pip-SDK images without hip-lang).
+if(NOT DEFINED FASTVIDEO_KERNEL_BUILD_EXTENSION AND DEFINED ENV{FASTVIDEO_KERNEL_BUILD_EXTENSION})
+    set(FASTVIDEO_KERNEL_BUILD_EXTENSION "$ENV{FASTVIDEO_KERNEL_BUILD_EXTENSION}")
+endif()
+set(FASTVIDEO_KERNEL_BUILD_EXTENSION "AUTO" CACHE STRING "Build the compiled extension: AUTO/ON/OFF")
+set_property(CACHE FASTVIDEO_KERNEL_BUILD_EXTENSION PROPERTY STRINGS AUTO ON OFF)
+string(TOUPPER "${FASTVIDEO_KERNEL_BUILD_EXTENSION}" FASTVIDEO_KERNEL_BUILD_EXTENSION)
+if(NOT FASTVIDEO_KERNEL_BUILD_EXTENSION MATCHES "^(AUTO|ON|OFF)$")
+    message(FATAL_ERROR
+        "fastvideo-kernel: FASTVIDEO_KERNEL_BUILD_EXTENSION must be AUTO, ON or OFF "
+        "(got '${FASTVIDEO_KERNEL_BUILD_EXTENSION}').")
+endif()
+
+set(BUILD_CXX_KERNELS ON)
+if(FASTVIDEO_KERNEL_BUILD_EXTENSION STREQUAL "OFF")
+    set(BUILD_CXX_KERNELS OFF)
+elseif(GPU_BACKEND STREQUAL "ROCM")
+    include(CheckLanguage)
+    check_language(HIP)
+    if(CMAKE_HIP_COMPILER)
+        enable_language(HIP)
+    elseif(FASTVIDEO_KERNEL_BUILD_EXTENSION STREQUAL "ON")
+        message(FATAL_ERROR
+            "fastvideo-kernel: FASTVIDEO_KERNEL_BUILD_EXTENSION=ON, but CMake found no usable HIP "
+            "toolchain (it needs the hip-lang CMake package under ROCM_PATH or CMAKE_HIP_COMPILER). "
+            "Install the ROCm CMake packages, or build the Python + Triton package with "
+            "FASTVIDEO_KERNEL_BUILD_EXTENSION=OFF (build.sh --rocm --python-only).")
+    else()
+        message(STATUS
+            "fastvideo-kernel: no usable HIP toolchain; building the Python + Triton package "
+            "without the compiled extension (FASTVIDEO_KERNEL_BUILD_EXTENSION=AUTO)")
+        set(BUILD_CXX_KERNELS OFF)
+    endif()
 else()
    enable_language(CUDA)
    # Ensure CUDA toolkit targets (CUDA::cudart, CUDA::cuda_driver, etc.) are available.
    find_package(CUDAToolkit REQUIRED)
 endif()

+if(NOT BUILD_CXX_KERNELS)
+    message(STATUS "============== fastvideo-kernel build summary ==============")
+    message(STATUS "host / backend:       ${CMAKE_SYSTEM_PROCESSOR} / ${GPU_BACKEND}")
+    message(STATUS "fastvideo_kernel_ops: OFF (Python + Triton package only; FASTVIDEO_KERNEL_BUILD_EXTENSION=${FASTVIDEO_KERNEL_BUILD_EXTENSION})")
+    message(STATUS "Triton fallbacks ship in python/fastvideo_kernel/triton_kernels regardless.")
+    message(STATUS "============================================================")
+    return()
+endif()
+
 # Import common utils if needed, but we keep it simple for now

 # Find Python and Torch
@@ -307,9 +348,6 @@ else()
    endif()
 endif()

-# Always try to build the extension if CUDA is available, but conditionally add sources/flags
-set(BUILD_CXX_KERNELS ON)
-
 # ---------------------------------------------------------------------------
 # Per-arch split for the Blackwell FP4 (attn_qat_infer) build
 # ---------------------------------------------------------------------------
@@ -396,27 +434,43 @@ if(BUILD_CXX_KERNELS)
    # 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_SM100_FAMILY OFF)
+    set(ENABLE_VSA_SM100A OFF)
+    set(ENABLE_VSA_SM103A 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\\.(0|3)a|10(0|3)a|sm_10(0|3)a)([; ,]|$)")
-        set(ENABLE_VSA_SM100_FAMILY ON)
+    if(_VSA_ARCH_LIST MATCHES "(^|[; ,])(10\\.0a|100a|sm_100a)([; ,]|$)")
+        set(ENABLE_VSA_SM100A ON)
    endif()
-    if(ENABLE_VSA_SM100_FAMILY)
-        message(STATUS "fastvideo-kernel: building block_sparse_sm100a (sm_100a/sm_103a, 64- and 128-token blocks)")
+    if(_VSA_ARCH_LIST MATCHES "(^|[; ,])(10\\.3a|103a|sm_103a)([; ,]|$)")
+        set(ENABLE_VSA_SM103A ON)
+    endif()
+    set(_VSA_FWD_GENCODE "")
+    set(_VSA_FWD_ARCHS "")
+    if(ENABLE_VSA_SM100A)
+        list(APPEND _VSA_FWD_GENCODE "-gencode;arch=compute_100a,code=sm_100a")
+        list(APPEND _VSA_FWD_ARCHS "sm_100a")
+    endif()
+    if(ENABLE_VSA_SM103A)
+        list(APPEND _VSA_FWD_GENCODE "-gencode;arch=compute_103a,code=sm_103a")
+        list(APPEND _VSA_FWD_ARCHS "sm_103a")
+    endif()
+    if(_VSA_FWD_GENCODE)
+        message(STATUS "fastvideo-kernel: building block_sparse_sm100a (${_VSA_FWD_ARCHS}; 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;-gencode;arch=compute_103a,code=sm_103a;-DVSA_BHSD=true")
-        # VSA block-sparse attention BACKWARD, 64-token blocks. sm_100a only for now: validated
-        # on GB200, not yet on B300/GB300, so no sm_103a image is built and the Python side keeps
-        # the Triton backward for sm_103a devices.
-        list(APPEND EXTENSION_SOURCES csrc/attention/block_sparse_bwd_sm100a.cu)
-        set_source_files_properties(csrc/attention/block_sparse_bwd_sm100a.cu PROPERTIES
-            COMPILE_OPTIONS "-gencode;arch=compute_100a,code=sm_100a;-DVSA_BHSD=true")
+            COMPILE_OPTIONS "${_VSA_FWD_GENCODE};-DVSA_BHSD=true")
+    endif()
+    # VSA block-sparse attention BACKWARD, 64- and 128-token blocks: the same images as the forward.
+    if(_VSA_FWD_GENCODE)
+        list(APPEND EXTENSION_SOURCES csrc/attention/block_sparse_bwd_sm100a.cu
+                                      csrc/attention/block_sparse_bwd_blk128_sm100a.cu)
+        set_source_files_properties(csrc/attention/block_sparse_bwd_sm100a.cu
+                                    csrc/attention/block_sparse_bwd_blk128_sm100a.cu PROPERTIES
+            COMPILE_OPTIONS "${_VSA_FWD_GENCODE};-DVSA_BHSD=true")
    endif()

    Python_add_library(fastvideo_kernel_ops MODULE USE_SABI ${SKBUILD_SABI_VERSION} WITH_SOABI
@@ -434,7 +488,7 @@ if(BUILD_CXX_KERNELS)

    # Build compile definitions list
    set(COMPILE_DEFS TORCH_EXTENSION_NAME=fastvideo_kernel_ops)
-    if(ENABLE_VSA_SM100_FAMILY)
+    if(ENABLE_VSA_SM100A OR ENABLE_VSA_SM103A)
        list(APPEND COMPILE_DEFS TK_COMPILE_BLOCK_SPARSE_VSA_SM100A)
    endif()
    if(ENABLE_TK_KERNELS)
@@ -69,6 +69,20 @@ cd fastvideo-kernel
 ./build.sh --rocm
 ```

+The compiled extension needs the HIP CMake toolchain (`hip-lang`). Images that install
+ROCm as a pip SDK, such as the `rocm/pytorch` 7.14 images, ship `hipcc` but not that
+package, so `build.sh --rocm` falls back to the Python + Triton package there and says so
+in the configure summary. That package is all the video sparse attention path needs on
+ROCm. To ask for it explicitly, or to fail instead of falling back:
+
+```bash
+./build.sh --rocm --python-only                                    # Python + Triton only
+CMAKE_ARGS="-DFASTVIDEO_KERNEL_BUILD_EXTENSION=ON" ./build.sh --rocm  # error if HIP cannot be configured
+```
+
+`FASTVIDEO_KERNEL_BUILD_EXTENSION` (AUTO/ON/OFF) is also read from the environment by a
+plain `pip install --no-build-isolation .`.
+
 ### Optional: FA4 CuTe block-sparse backend (VSA-128/256 fastpath)

 The VSA-128/256 fastpaths (tile volume 128 or 256, on NVIDIA Blackwell / sm_100) route to the
@@ -145,8 +159,11 @@ python benchmarks/benchmark_attn_qat_train.py
 ```

 The benchmark reports both conventional attention FLOPs and the extra matrix
-multiplications executed by the QAT straight-through path. Override
-`--peak-tflops` when running on a GPU other than RTX 5090.
+multiplications executed by the QAT straight-through path. It resolves dense
+BF16 peak throughput for the full-GPU variants listed in
+`benchmarks/device_specs.py`. Unknown or partitioned devices still report
+timing and achieved TFLOPS, with MFU shown as `N/A`; pass a positive, finite
+`--peak-tflops` value to report MFU for those devices.

 The QAT kernel is entirely Triton and routes by architecture at runtime. SM100
 uses a large-tile forward and split 64x64 backward for the production
@@ -14,7 +14,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 """
 import torch
-import fp4quant
+import fp4quant_cuda as fp4quant
 from triton.tools.mxfp import MXFP4Tensor

 from bench_utils import bench_kineto
@@ -14,7 +14,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 """
 import torch
-import fp4quant
+import fp4quant_cuda as fp4quant
 from triton.tools.mxfp import MXFP4Tensor

 from bench_utils import bench_kineto
@@ -14,7 +14,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 """
 import torch
-import fp4quant
+import fp4quant_cuda as fp4quant
 from triton.tools.mxfp import MXFP4Tensor

 from bench_utils import bench_kineto
@@ -14,9 +14,9 @@ from collections.abc import Callable

 import torch

-from fastvideo_kernel.triton_kernels.attn_qat_train import attention
+from device_specs import format_tflops_with_mfu, resolve_peak_tflops

-RTX_5090_DENSE_BF16_TFLOPS = 209.5
+from fastvideo_kernel.triton_kernels.attn_qat_train import attention


 def _qat_attention(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor) -> torch.Tensor:
@@ -62,14 +62,15 @@ def _format_result(
    timing_ms: tuple[float, float, float],
    algorithmic_flops: int,
    executed_matmul_flops: int,
-    peak_tflops: float,
+    peak_tflops: float | None,
 ) -> str:
    median_ms, min_ms, max_ms = timing_ms
    algorithmic_tflops = algorithmic_flops / (median_ms * 1e9)
    executed_tflops = executed_matmul_flops / (median_ms * 1e9)
+    algorithmic_result = format_tflops_with_mfu(algorithmic_tflops, peak_tflops)
+    executed_result = format_tflops_with_mfu(executed_tflops, peak_tflops)
    return (f"{label}: {median_ms:.3f} ms (min={min_ms:.3f}, max={max_ms:.3f}), "
-            f"algorithmic={algorithmic_tflops:.2f} TFLOPS/{100 * algorithmic_tflops / peak_tflops:.2f}% MFU, "
-            f"executed_matmul={executed_tflops:.2f} TFLOPS/{100 * executed_tflops / peak_tflops:.2f}% MFU")
+            f"algorithmic={algorithmic_result}, executed_matmul={executed_result}")


 def main() -> None:
@@ -84,12 +85,16 @@ def main() -> None:
    parser.add_argument(
        "--peak-tflops",
        type=float,
-        default=RTX_5090_DENSE_BF16_TFLOPS,
-        help="Dense BF16 Tensor TFLOPS with FP32 accumulation; default is RTX 5090 boost-clock peak.",
+        default=None,
+        help="Dense BF16 Tensor TFLOPS with FP32 accumulation. Defaults to the "
+        "device table in device_specs.py. Unknown devices report MFU as N/A.",
    )
    args = parser.parse_args()
    kv_length = args.kv_length or args.query_length

+    device_name = torch.cuda.get_device_name()
+    peak_tflops, peak_source = resolve_peak_tflops(device_name, args.peak_tflops)
+
    torch.manual_seed(0)
    q_shape = (args.batch_size, args.heads, args.query_length, args.head_dim)
    kv_shape = (args.batch_size, args.heads, kv_length, args.head_dim)
@@ -114,10 +119,14 @@ def main() -> None:
    # Conventional attention FLOPs are 4*base forward and 10*base backward.
    # QAT additionally computes the STE high-precision P@V path in forward and
    # the quantized-P dV path in backward, for 6*base and 14*base matmul FLOPs.
-    print(f"device: {torch.cuda.get_device_name()}")
+    print(f"device: {device_name}")
+    if peak_tflops is None:
+        print(f"peak: unavailable ({peak_source})")
+    else:
+        print(f"peak: {peak_tflops:.2f} dense BF16 TFLOPS ({peak_source})")
    print(f"q: {q_shape}; k/v: {kv_shape}; compile+first-forward: {compile_seconds:.3f} s")
-    print(_format_result("forward", forward_ms, 4 * base_flops, 6 * base_flops, args.peak_tflops))
-    print(_format_result("backward", backward_ms, 10 * base_flops, 14 * base_flops, args.peak_tflops))
+    print(_format_result("forward", forward_ms, 4 * base_flops, 6 * base_flops, peak_tflops))
+    print(_format_result("backward", backward_ms, 10 * base_flops, 14 * base_flops, peak_tflops))


 if __name__ == "__main__":
@@ -0,0 +1,56 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Per-device peak throughput used to turn benchmark timings into %MFU.
+
+Values are dense (non-sparse) BF16 Tensor Core TFLOPS with FP32 accumulation.
+Only unambiguous full-GPU names with published peaks are listed. Timings and
+achieved TFLOPS remain useful on other devices, so unknown devices report MFU
+as unavailable unless the user supplies an explicit peak.
+"""
+
+from __future__ import annotations
+
+import math
+
+# Exact torch.cuda.get_device_name() values for supported full GPUs. H100 form
+# factors have different published peaks, and exact matching also prevents MIG,
+# vGPU, or future variants from silently inheriting a full-GPU denominator.
+DENSE_BF16_TFLOPS: dict[str, float] = {
+    "NVIDIA H100 80GB HBM3": 989.0,
+    "NVIDIA H100 NVL": 835.5,
+    "NVIDIA H100 PCIe": 756.5,
+    "NVIDIA GeForce RTX 5090": 209.5,  # Matches the value this benchmark historically used.
+    "NVIDIA A100-SXM4-80GB": 312.0,
+    "NVIDIA A100-SXM4-40GB": 312.0,
+    "NVIDIA A100-PCIE-80GB": 312.0,
+    "NVIDIA A100-PCIE-40GB": 312.0,
+    "NVIDIA L40S": 362.05,
+}
+
+
+def resolve_peak_tflops(device_name: str, override: float | None = None) -> tuple[float | None, str]:
+    """Return (peak_tflops, source_description) for *device_name*.
+
+    An explicit positive, finite *override* always wins. Unknown and partitioned
+    devices return ``None`` so the benchmark can still report timing and
+    achieved TFLOPS without presenting a misleading MFU value.
+    """
+    if override is not None:
+        if not math.isfinite(override) or override <= 0:
+            raise ValueError("--peak-tflops must be finite and greater than zero")
+        return float(override), "explicit --peak-tflops override"
+
+    normalized_name = device_name.strip()
+    if "MIG" in normalized_name.upper():
+        return None, "MIG/partitioned device; pass --peak-tflops for this partition to report MFU"
+
+    if normalized_name in DENSE_BF16_TFLOPS:
+        return DENSE_BF16_TFLOPS[normalized_name], f"exact device table entry {normalized_name!r}"
+
+    return None, f"no dense BF16 peak for {device_name!r}; pass --peak-tflops to report MFU"
+
+
+def format_tflops_with_mfu(achieved_tflops: float, peak_tflops: float | None) -> str:
+    """Format achieved throughput, including MFU only when its denominator is known."""
+    if peak_tflops is None:
+        return f"{achieved_tflops:.2f} TFLOPS/MFU N/A"
+    return f"{achieved_tflops:.2f} TFLOPS/{100 * achieved_tflops / peak_tflops:.2f}% MFU"
@@ -5,9 +5,11 @@ set -ex
 # Usage:
 #   ./build.sh  # local build (torch-based arch detection, TK only on SM90)
 #   ./build.sh --wheel-dir /tmp/wheels  # build + install a reusable wheel
+#   ./build.sh --rocm --python-only     # Python + Triton package, no compiled extension
 # Environment overrides (if set, they win over auto-detection):
 #   TORCH_CUDA_ARCH_LIST
-#   CMAKE_ARGS (for FASTVIDEO_KERNEL_BUILD_TK / CMAKE_CUDA_ARCHITECTURES / GPU_BACKEND)
+#   CMAKE_ARGS (for FASTVIDEO_KERNEL_BUILD_TK / CMAKE_CUDA_ARCHITECTURES / GPU_BACKEND /
+#               FASTVIDEO_KERNEL_BUILD_EXTENSION)

 echo "Building fastvideo-kernel..."

@@ -40,42 +42,19 @@ if [[ -n "${CONDA_PREFIX:-}" ]]; then
    unset _need_clean _host_arch
 fi

-# Ensure only the kernel's required headers are initialized. A repository-wide
-# update also clones the unrelated VBench evaluation submodule. Skip outside a
-# git checkout (e.g. Docker contexts that exclude .git), where the submodule
-# contents must already be present.
-if git rev-parse --git-dir >/dev/null 2>&1; then
-    git submodule update --init --recursive include/cutlass include/tk
-fi
-# Fail fast with a clear message if the headers are still missing (e.g. a
-# Docker context that excluded .git AND the submodule contents) instead of
-# dying later in a wall of nvcc include errors. CUTLASS is consumed by the
-# always-built turbodiffusion sources, so it is a hard error; ThunderKittens
-# only feeds the TK-gated Hopper kernels, so a missing tree just warns (the
-# TK gate resolves later, and non-SM90/ROCm builds never touch it).
-if [ ! -d include/cutlass/include ]; then
-    echo "ERROR: include/cutlass/include is missing. Outside a git checkout the" >&2
-    echo "       CUTLASS sources must already be present (run" >&2
-    echo "       'git submodule update --init --recursive include/cutlass include/tk'" >&2
-    echo "       in the source checkout, or include them in the build context)." >&2
-    exit 1
-fi
-if [ ! -d include/tk/include ]; then
-    echo "WARNING: include/tk/include is missing; ThunderKittens (Hopper sm_90a)" >&2
-    echo "         kernels cannot be built. Fine for non-SM90/ROCm targets." >&2
-fi
-
-# Install build dependencies
-uv pip install scikit-build-core cmake ninja
-
 GPU_BACKEND=CUDA
 WHEEL_DIR=""
+PYTHON_ONLY=0
 while [[ $# -gt 0 ]]; do
    case "$1" in
        --rocm)
            GPU_BACKEND=ROCM
            shift
            ;;
+        --python-only)
+            PYTHON_ONLY=1
+            shift
+            ;;
        --wheel-dir)
            if [[ $# -lt 2 ]]; then
                echo "ERROR: --wheel-dir requires a directory path" >&2
@@ -96,6 +75,55 @@ has_cmake_arg() {
    [[ "${CMAKE_ARGS:-}" =~ (^|[[:space:]])-D${key}(=|$) ]]
 }

+extension_mode_from_cmake_args() {
+    [[ "${CMAKE_ARGS:-}" =~ (^|[[:space:]])-DFASTVIDEO_KERNEL_BUILD_EXTENSION=([^[:space:]]*) ]] \
+        && printf '%s' "${BASH_REMATCH[2]^^}"
+}
+if (( PYTHON_ONLY )); then
+    _explicit_mode="$(extension_mode_from_cmake_args)"
+    if [[ -n "${_explicit_mode}" && "${_explicit_mode}" != "OFF" ]]; then
+        echo "ERROR: --python-only conflicts with CMAKE_ARGS -DFASTVIDEO_KERNEL_BUILD_EXTENSION=${_explicit_mode}" >&2
+        exit 2
+    fi
+    if ! has_cmake_arg "FASTVIDEO_KERNEL_BUILD_EXTENSION"; then
+        CMAKE_ARGS="${CMAKE_ARGS:-} -DFASTVIDEO_KERNEL_BUILD_EXTENSION=OFF"
+    fi
+elif [[ "$(extension_mode_from_cmake_args)" == "OFF" ]]; then
+    PYTHON_ONLY=1
+fi
+
+# The vendored headers only feed the compiled extension, so a Python-only build
+# neither initializes nor checks them.
+if ! (( PYTHON_ONLY )); then
+    # Ensure only the kernel's required headers are initialized. A repository-wide
+    # update also clones the unrelated VBench evaluation submodule. Skip outside a
+    # git checkout (e.g. Docker contexts that exclude .git), where the submodule
+    # contents must already be present.
+    if git rev-parse --git-dir >/dev/null 2>&1; then
+        git submodule update --init --recursive include/cutlass include/tk
+    fi
+    # Fail fast with a clear message if the headers are still missing (e.g. a
+    # Docker context that excluded .git AND the submodule contents) instead of
+    # dying later in a wall of nvcc include errors. CUTLASS is consumed by the
+    # always-built turbodiffusion sources, so it is a hard error; ThunderKittens
+    # only feeds the TK-gated Hopper kernels, so a missing tree just warns (the
+    # TK gate resolves later, and non-SM90/ROCm builds never touch it).
+    if [ ! -d include/cutlass/include ]; then
+        echo "ERROR: include/cutlass/include is missing. Outside a git checkout the" >&2
+        echo "       CUTLASS sources must already be present (run" >&2
+        echo "       'git submodule update --init --recursive include/cutlass include/tk'" >&2
+        echo "       in the source checkout, or include them in the build context)." >&2
+        exit 1
+    fi
+    if [ ! -d include/tk/include ]; then
+        echo "WARNING: include/tk/include is missing; ThunderKittens (Hopper sm_90a)" >&2
+        echo "         kernels cannot be built. Fine for non-SM90/ROCm targets." >&2
+    fi
+fi
+
+# Install build dependencies
+uv pip install scikit-build-core cmake ninja
+
 detect_with_torch() {
    # Prefer the active venv's python directly over `uv run --active --no-project`,
    # which on some uv versions provisions its own interpreter and misses packages
@@ -113,7 +141,7 @@ mj, mn = torch.cuda.get_device_capability(0)
 print(f'{mj}.{mn}')"
 }

-if [ "${GPU_BACKEND}" = "CUDA" ]; then
+if [ "${GPU_BACKEND}" = "CUDA" ] && ! (( PYTHON_ONLY )); then
    # Compute capability drives the arch/TK defaults below. Prefer an explicit
    # TORCH_CUDA_ARCH_LIST (works on GPU-less build machines such as CI/Docker);
    # only probe a live GPU via torch when no arch was provided.
@@ -0,0 +1,7 @@
+// block_sparse_bwd_blk128_sm100a.cu -- the 128-token-block instantiation of the torch binding.
+//
+// Same source as block_sparse_bwd_sm100a.cu with VSA_BLK128 set: the kernel and launch land in
+// namespace vsa_bwd_blk128 (distinct symbols, no ODR clash with the blk64 objects) and the
+// exported entry point becomes block_sparse_sm100a_blk128_bwd.
+#define VSA_BLK128 true
+#include "block_sparse_bwd_sm100a.cu"
@@ -123,7 +123,9 @@ __global__ void __cluster_dims__(1, 1, 1) __launch_bounds__(N_WARPS * 32, 1) vsa
    const int* __restrict__ k2q_num, const int* __restrict__ workitem_remap,
    const int* __restrict__ variable_block_sizes, int max_q_blocks, int num_samples, int num_heads,
    int seqlen, float scale_log2, float sm_scale) {
-#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL))
+#if !defined(__CUDA_ARCH__) || \
+    ((__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL)) || \
+     (__CUDA_ARCH__ == 1030 && defined(__CUDA_ARCH_FEAT_SM103_ALL)))
  using DQ                        = DQConfig<DQ_DTYPE>;
  const int num_kv_blocks_per_seq = seqlen / BLOCK;

@@ -863,7 +865,9 @@ __global__ void __launch_bounds__(ORDER_THREADS, 1)
    vsa_bwd_order_kernel(const int* __restrict__ k2q_idx, const int* __restrict__ k2q_num,
                         int max_q_blocks, int num_kv_blocks_per_seq, int order_bin, bool snake,
                         int* __restrict__ order_out) {
-#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL))
+#if !defined(__CUDA_ARCH__) || \
+    ((__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL)) || \
+     (__CUDA_ARCH__ == 1030 && defined(__CUDA_ARCH_FEAT_SM103_ALL)))
  extern __shared__ int order_smem[];
  int* sbin      = order_smem;
  int* smid      = order_smem + num_kv_blocks_per_seq;
@@ -908,7 +912,9 @@ __global__ void __launch_bounds__(256, 1)
                              __nv_bfloat16* __restrict__ dk, __nv_bfloat16* __restrict__ dv,
                              const int* __restrict__ k2q_num, int num_samples, int num_heads,
                              int seqlen) {
-#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL))
+#if !defined(__CUDA_ARCH__) || \
+    ((__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL)) || \
+     (__CUDA_ARCH__ == 1030 && defined(__CUDA_ARCH_FEAT_SM103_ALL)))
  __shared__ __align__(128) __nv_bfloat16 tile[PRE_TOKENS][SUB_COLS_BF16 + 4];
  const int num_kv_blocks_per_seq = seqlen / BLOCK;
  const int token_block_id        = (int)blockIdx.x;
@@ -1038,7 +1044,9 @@ template <bool BHSD = false, typename DQ_DTYPE = float>
 __global__ void __launch_bounds__(128, 1)
    vsa_bwd_postprocess_kernel(const DQ_DTYPE* __restrict__ dqaccum, __nv_bfloat16* __restrict__ dq,
                               int num_heads, int seqlen, float sm_scale) {
-#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL))
+#if !defined(__CUDA_ARCH__) || \
+    ((__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL)) || \
+     (__CUDA_ARCH__ == 1030 && defined(__CUDA_ARCH_FEAT_SM103_ALL)))
  const int q_block_id = (int)blockIdx.x;
  const int batch_head = (int)blockIdx.y;
  const int batch = batch_head / num_heads, head = batch_head % num_heads;
@@ -0,0 +1,231 @@
+// block_sparse_bwd_launch_sm100a_blk128.cuh -- host surface of the 128-token VSA block-sparse
+// backward drop: argument struct, workspace sizes, the support predicate and the stream-chained
+// launch (preprocess -> main -> postprocess, each a programmatic dependent launch of the one
+// before). Tensor maps are encoded per call (a torch caller hands us fresh pointers every time).
+#ifndef BLOCK_SPARSE_VSA_BWD_LAUNCH_SM100A_BLK128_CUH
+#define BLOCK_SPARSE_VSA_BWD_LAUNCH_SM100A_BLK128_CUH
+
+#include <cmath>
+#include "block_sparse_bwd_kernel_sm100a_blk128.cuh"
+
+#ifndef VSA_BHSD
+#define VSA_BHSD false
+#endif
+
+namespace vsa_bwd_blk128 {
+
+using dq_accum_t = float;
+
+struct BlockSparseVsaBwdArgs {
+  // Activations are bf16, contiguous, [B, H, S, 128] under VSA_BHSD, else [B*S, H, 128].
+  // nb below = num_kv_blocks_per_seq = S / 128.
+
+  // Forward operands and results.
+  const __nv_bfloat16* q;
+  const __nv_bfloat16* k;
+  const __nv_bfloat16* v;
+  const __nv_bfloat16* o;
+  // Gradient of the forward output.
+  const __nv_bfloat16* dout;
+  // [B, H, S] fp32 log-sum-exp in Triton's M form: max(qk * sm_scale * log2e) + log2(l).
+  const float* lse;
+
+  // Sparsity metadata, FastVideo's invert_indices layout, in 128-token blocks.
+  // [B*H*nb, max_q_blocks] int32: q blocks selecting each kv block; entries past the count unread.
+  const int* k2q_idx;
+  // [B*H*nb] int32: valid entries per k2q_idx row (0 allowed: the kernel writes zero dk/dv rows).
+  const int* k2q_num;
+  // [nb] int32: valid kv tokens per block (<= 128); kv rows at or past the count are masked.
+  const int* variable_block_sizes;
+
+  // Outputs, inputs' layout; every row is written, including the rows of unselected kv blocks.
+  __nv_bfloat16* dq;
+  __nv_bfloat16* dk;
+  __nv_bfloat16* dv;
+
+  // Scratch, caller-allocated; byte sizes from the block_sparse_bwd_*_bytes helpers below.
+  // [B*H*S*128] fp32, drain-native; preprocess zeroes, main reduce-adds, postprocess reads.
+  dq_accum_t* dqaccum;
+  // [B*H*S] fp32 rowsum(bf16(o) * dout), written by the preprocess.
+  float* delta;
+
+  int batch;
+  int num_heads;
+  // S = nb * 128.
+  int seqlen;
+  // Must be 128.
+  int head_dim;
+  // nb = seqlen / 128.
+  int num_kv_blocks_per_seq;
+  // k2q_idx row stride (FastVideo passes nb).
+  int max_q_blocks;
+  // Softmax scale; dq and dk carry it, dv does not.
+  float sm_scale;
+};
+
+__host__ inline size_t block_sparse_bwd_dqaccum_bytes(int batch, int num_heads, int seqlen) {
+  return (size_t)batch * num_heads * seqlen * HEAD_DIM * sizeof(dq_accum_t);
+}
+__host__ inline size_t block_sparse_bwd_delta_bytes(int batch, int num_heads, int seqlen) {
+  return (size_t)batch * num_heads * seqlen * sizeof(float);
+}
+
+__host__ inline cudaError_t block_sparse_bwd_supported(const BlockSparseVsaBwdArgs& args) {
+  if (args.head_dim != HEAD_DIM) {
+    return cudaErrorInvalidValue;
+  }
+  if (args.num_kv_blocks_per_seq < 1) {
+    return cudaErrorInvalidValue;
+  }
+  if (args.seqlen != args.num_kv_blocks_per_seq * BLOCK) {
+    return cudaErrorInvalidValue;
+  }
+  if (args.batch < 1 || args.num_heads < 1 || args.max_q_blocks < 1) {
+    return cudaErrorInvalidValue;
+  }
+  // The main grid is (kv block, batch * head): gridDim.y is limited to 65535.
+  if ((long)args.batch * args.num_heads > 65535) {
+    return cudaErrorInvalidValue;
+  }
+  if (!std::isfinite(args.sm_scale)) {
+    return cudaErrorInvalidValue;
+  }
+  if (!args.q || !args.k || !args.v || !args.o || !args.dout || !args.lse) {
+    return cudaErrorInvalidValue;
+  }
+  if (!args.dq || !args.dk || !args.dv) {
+    return cudaErrorInvalidValue;
+  }
+  if (!args.k2q_idx || !args.k2q_num || !args.variable_block_sizes) {
+    return cudaErrorInvalidValue;
+  }
+  if (!args.dqaccum || !args.delta) {
+    return cudaErrorInvalidValue;
+  }
+  return cudaSuccess;
+}
+
+// Q, K, V, dO, dK, dV tensor maps, one box_tokens x 64-hd box per TMA (two per tile):
+//   BSHD: 3D [64 hd, B*S tokens, H*2 hd units], strides {H*128*2, 128} bytes.
+//   BHSD: 4D [64 hd, S tokens, 2 hd units, B*H], strides {128*2, 128, S*128*2} bytes.
+__host__ inline cudaError_t make_tma_tile_units(CUtensorMap* map, const __nv_bfloat16* ptr, int B,
+                                                int H, int S, int box_tokens) {
+  CUresult r;
+  if (VSA_BHSD) {
+    uint64_t gd[4] = {(uint64_t)SUB_COLS_BF16, (uint64_t)S, (uint64_t)HD_SUBTILES, (uint64_t)B * H};
+    uint64_t gs[3] = {(uint64_t)HEAD_DIM * 2, (uint64_t)SUB_COLS_BYTES, (uint64_t)S * HEAD_DIM * 2};
+    uint32_t bd[4] = {(uint32_t)SUB_COLS_BF16, (uint32_t)box_tokens, 1u, 1u};
+    uint32_t es[4] = {1u, 1u, 1u, 1u};
+    r              = cuTensorMapEncodeTiled(
+        map, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, 4, const_cast<__nv_bfloat16*>(ptr), 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);
+  } else {
+    uint64_t gd[3] = {(uint64_t)SUB_COLS_BF16, (uint64_t)B * S, (uint64_t)H * HD_SUBTILES};
+    uint64_t gs[2] = {(uint64_t)H * HEAD_DIM * 2, (uint64_t)SUB_COLS_BYTES};
+    uint32_t bd[3] = {(uint32_t)SUB_COLS_BF16, (uint32_t)box_tokens, 1u};
+    uint32_t es[3] = {1u, 1u, 1u};
+    r              = cuTensorMapEncodeTiled(
+        map, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, 3, const_cast<__nv_bfloat16*>(ptr), 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);
+  }
+  return (r == CUDA_SUCCESS) ? cudaSuccess : cudaErrorInvalidValue;
+}
+
+// Every launch carries the programmatic-stream-serialization attribute when the kernels are
+// built with KERNEL_PDL (their griddepcontrol.wait guards the predecessor's data).
+__host__ inline cudaLaunchConfig_t pdl_launch_config(dim3 grid, dim3 block, size_t smem,
+                                                     cudaStream_t stream, cudaLaunchAttribute* at) {
+  cudaLaunchConfig_t cfg = {};
+  cfg.gridDim            = grid;
+  cfg.blockDim           = block;
+  cfg.dynamicSmemBytes   = smem;
+  cfg.stream             = stream;
+  at->id                 = cudaLaunchAttributeProgrammaticStreamSerialization;
+  at->val.programmaticStreamSerializationAllowed = 1;
+  cfg.attrs                                      = at;
+  cfg.numAttrs                                   = KERNEL_PDL ? 1 : 0;
+  return cfg;
+}
+
+__host__ inline cudaError_t launch_block_sparse_bwd_sm100a(const BlockSparseVsaBwdArgs& args,
+                                                           cudaStream_t stream) {
+  const cudaError_t supported = block_sparse_bwd_supported(args);
+  if (supported != cudaSuccess) {
+    return supported;
+  }
+  const int B = args.batch, H = args.num_heads, S = args.seqlen;
+
+  CUtensorMap tq, tk, tv, tdo, tdk, tdv;
+  if (make_tma_tile_units(&tq, args.q, B, H, S, Q_TILE) != cudaSuccess ||
+      make_tma_tile_units(&tk, args.k, B, H, S, KV_TILE) != cudaSuccess ||
+      make_tma_tile_units(&tv, args.v, B, H, S, KV_TILE) != cudaSuccess ||
+      make_tma_tile_units(&tdo, args.dout, B, H, S, Q_TILE) != cudaSuccess ||
+      make_tma_tile_units(&tdk, args.dk, B, H, S, KV_TILE) != cudaSuccess ||
+      make_tma_tile_units(&tdv, args.dv, B, H, S, KV_TILE) != cudaSuccess) {
+    return cudaErrorInvalidValue;
+  }
+
+  auto main_kernel        = vsa_bwd_main_kernel<VSA_BHSD, Sched::NON_PERSISTENT>;
+  auto postprocess_kernel = vsa_bwd_postprocess_kernel<VSA_BHSD>;
+  constexpr int POST_SMEM_BYTES = DQ::DQ_BLOCK_ELEMS * (int)sizeof(float);
+  cudaError_t e =
+      cudaFuncSetAttribute(main_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, SMEM_TOTAL);
+  if (e != cudaSuccess) {
+    return e;
+  }
+  e = cudaFuncSetAttribute(postprocess_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize,
+                           POST_SMEM_BYTES);
+  if (e != cudaSuccess) {
+    return e;
+  }
+
+  // Preprocess: dqaccum zero + Delta, one CTA per (q block, batch * head).
+  cudaLaunchAttribute pre_at[1];
+  cudaLaunchConfig_t pre_cfg = pdl_launch_config(
+      dim3((unsigned)(S / Q_TILE), (unsigned)(B * H), 1), dim3(256, 1, 1), 0, stream, pre_at);
+  e = cudaLaunchKernelEx(&pre_cfg, vsa_bwd_preprocess_kernel<VSA_BHSD>, args.o, args.dout,
+                         args.delta, args.dqaccum, H, S);
+  if (e != cudaSuccess) {
+    return e;
+  }
+
+  // Main: one CTA per (kv block, batch * head).
+  const float scale_log2 = args.sm_scale * 1.4426950408889634f;
+  cudaLaunchAttribute main_at[2];
+  cudaLaunchConfig_t main_cfg = {};
+  main_cfg.gridDim            = dim3((unsigned)args.num_kv_blocks_per_seq, (unsigned)(B * H), 1);
+  main_cfg.blockDim           = dim3(N_WARPS * 32, 1, 1);
+  main_cfg.dynamicSmemBytes   = SMEM_TOTAL;
+  main_cfg.stream             = stream;
+  main_at[0].id               = cudaLaunchAttributeClusterDimension;
+  main_at[0].val.clusterDim.x = 1;
+  main_at[0].val.clusterDim.y = 1;
+  main_at[0].val.clusterDim.z = 1;
+  main_at[1].id               = cudaLaunchAttributeProgrammaticStreamSerialization;
+  main_at[1].val.programmaticStreamSerializationAllowed = 1;
+  main_cfg.attrs                                        = main_at;
+  main_cfg.numAttrs                                     = KERNEL_PDL ? 2 : 1;
+  const int* workitem_remap                             = nullptr;
+  e = cudaLaunchKernelEx(&main_cfg, main_kernel, tq, tk, tv, tdo, tdk, tdv, args.dqaccum, args.lse,
+                         args.delta, args.k2q_idx, args.k2q_num, args.max_q_blocks, workitem_remap,
+                         args.variable_block_sizes, B, H, S, scale_log2, args.sm_scale);
+  if (e != cudaSuccess) {
+    return e;
+  }
+
+  // Postprocess: drain-native dqaccum -> bf16 dq, one CTA per (q block, batch * head).
+  cudaLaunchAttribute post_at[1];
+  cudaLaunchConfig_t post_cfg =
+      pdl_launch_config(dim3((unsigned)(S / Q_TILE), (unsigned)(B * H), 1), dim3(128, 1, 1),
+                        POST_SMEM_BYTES, stream, post_at);
+  return cudaLaunchKernelEx(&post_cfg, postprocess_kernel, args.dqaccum, args.dq, H, S,
+                            args.sm_scale);
+}
+
+}  // namespace vsa_bwd_blk128
+
+using namespace vsa_bwd_blk128;
+
+#endif  // BLOCK_SPARSE_VSA_BWD_LAUNCH_SM100A_BLK128_CUH
@@ -10,6 +10,9 @@
 //
 // The layout is fixed at compile time: VSA_BHSD true -> [B, H, S, 128] (FastVideo's build),
 // false -> [B, S, H, 128] (repo native; the kernel addresses it as [B*S tokens, H, 128]).
+// VSA_BLK128 picks the kernel: false -> 64-token sparse blocks (block_sparse_sm100a_bwd), true ->
+// 128-token (block_sparse_sm100a_blk128_bwd; block_sparse_bwd_blk128_sm100a.cu re-includes this
+// file with it set). The python backend picks by the metadata's block size.
 #include <torch/extension.h>

 #include <ATen/cuda/CUDAContext.h>
@@ -17,7 +20,16 @@

 #include <vector>

+#ifndef VSA_BLK128
+#define VSA_BLK128 false
+#endif
+#if VSA_BLK128
+#include "block_sparse_bwd_launch_sm100a_blk128.cuh"
+#define BLOCK_SPARSE_SM100A_BWD block_sparse_sm100a_blk128_bwd
+#else
 #include "block_sparse_bwd_launch_sm100a.cuh"
+#define BLOCK_SPARSE_SM100A_BWD block_sparse_sm100a_bwd
+#endif

 namespace {

@@ -49,7 +61,7 @@ __nv_bfloat16* bf16_ptr(const torch::Tensor& t) {
 }  // namespace

 // Returns {dq, dk, dv}: bf16, each with the shape and layout of q, k, v respectively.
-std::vector<torch::Tensor> block_sparse_sm100a_bwd(torch::Tensor grad_o, torch::Tensor q,
+std::vector<torch::Tensor> BLOCK_SPARSE_SM100A_BWD(torch::Tensor grad_o, torch::Tensor q,
                                                   torch::Tensor k, torch::Tensor v,
                                                   torch::Tensor o, torch::Tensor lse,
                                                   torch::Tensor k2q_idx, torch::Tensor k2q_num,
@@ -105,20 +117,23 @@ std::vector<torch::Tensor> block_sparse_sm100a_bwd(torch::Tensor grad_o, torch::
  auto dv = torch::empty_like(v);

  // Workspace: torch::empty is enough. The preprocess kernel zeroes dqaccum and fully writes
-  // qt, dot and delta before the main kernel reads them; it also zeroes the dk/dv rows of kv
-  // blocks that no q block selects, so the empty_like outputs above come back fully defined.
+  // delta (and, at 64-token blocks, qt and dot) before the main kernel reads them. The dk/dv rows
+  // of kv blocks that no q block selects are zeroed by the preprocess (64) or by the main kernel
+  // (128), so the empty_like outputs above come back fully defined.
  const auto bytes = q.options().dtype(at::kByte);
  const int b = (int)B, h = (int)H, s = (int)S;
  auto dqaccum = torch::empty({(int64_t)block_sparse_bwd_dqaccum_bytes(b, h, s)}, bytes);
-  auto qt      = torch::empty({(int64_t)block_sparse_bwd_transposed_bytes(b, h, s)}, bytes);
-  auto dot     = torch::empty({(int64_t)block_sparse_bwd_transposed_bytes(b, h, s)}, bytes);
  auto delta   = torch::empty({(int64_t)block_sparse_bwd_delta_bytes(b, h, s)}, bytes);
+#if !VSA_BLK128
+  auto qt  = torch::empty({(int64_t)block_sparse_bwd_transposed_bytes(b, h, s)}, bytes);
+  auto dot = torch::empty({(int64_t)block_sparse_bwd_transposed_bytes(b, h, s)}, bytes);
  torch::Tensor order;
  const bool device_order = num_kv_blocks_per_seq >= ORDER_MIN_KV_BLOCKS;
  if (device_order) {
    order = torch::empty({(int64_t)block_sparse_bwd_order_bytes(b, h, (int)num_kv_blocks_per_seq)},
                         bytes);
  }
+#endif

  BlockSparseVsaBwdArgs a{};
  a.q                     = bf16_ptr(q);
@@ -133,12 +148,14 @@ std::vector<torch::Tensor> block_sparse_sm100a_bwd(torch::Tensor grad_o, torch::
  a.k2q_idx               = k2q_idx.data_ptr<int>();
  a.k2q_num               = k2q_num.data_ptr<int>();
  a.variable_block_sizes  = variable_block_sizes.data_ptr<int>();
-  a.workitem_remap        = nullptr;
-  a.order_workspace       = device_order ? reinterpret_cast<int*>(order.data_ptr()) : nullptr;
  a.dqaccum               = reinterpret_cast<dq_accum_t*>(dqaccum.data_ptr());
-  a.qt                    = bf16_ptr(qt);
-  a.dot                   = bf16_ptr(dot);
  a.delta                 = reinterpret_cast<float*>(delta.data_ptr());
+#if !VSA_BLK128
+  a.workitem_remap  = nullptr;
+  a.order_workspace = device_order ? reinterpret_cast<int*>(order.data_ptr()) : nullptr;
+  a.qt              = bf16_ptr(qt);
+  a.dot             = bf16_ptr(dot);
+#endif
  a.batch                 = b;
  a.num_heads             = h;
  a.seqlen                = s;
@@ -149,16 +166,17 @@ std::vector<torch::Tensor> block_sparse_sm100a_bwd(torch::Tensor grad_o, torch::

  // Report an unsupported regime loudly rather than returning plausible-looking wrong values.
  TORCH_CHECK(block_sparse_bwd_supported(a) == cudaSuccess,
-              "block_sparse_sm100a_bwd: unsupported configuration -- requires head_dim==", HEAD_DIM,
+              "block_sparse_sm100a_bwd (", BLOCK,
+              "-token blocks): unsupported configuration -- requires head_dim==", HEAD_DIM,
              ", seqlen == num_kv_blocks_per_seq*", BLOCK,
-              " with seqlen % 128 == 0, "
+              " with seqlen % 128 == 0, batch*heads <= 65535, "
              "max_q_blocks >= 1 and a finite sm_scale. Got head_dim=",
              D, " num_kv_blocks_per_seq=", num_kv_blocks_per_seq, " seqlen=", S,
              " max_q_blocks=", max_q_blocks, " sm_scale=", sm_scale);

  const cudaError_t err = launch_block_sparse_bwd_sm100a(a, at::cuda::getCurrentCUDAStream());
-  TORCH_CHECK(err == cudaSuccess,
-              "block_sparse_sm100a_bwd launch failed: ", cudaGetErrorString(err));
+  TORCH_CHECK(err == cudaSuccess, "block_sparse_sm100a_bwd (", BLOCK,
+              "-token blocks) launch failed: ", cudaGetErrorString(err));

  return {dq, dk, dv};
 }
@@ -985,3 +985,30 @@ uint64_t make_l2cache_policy_fractional_evict_last_unchanged(float fraction_f32)
      : "f"(fraction_f32));
  return policy;
 }
+
+__device__ __forceinline__
+void cp_async_cg_16(uint32_t smem_dst, const void* gmem_src) {
+  asm volatile("cp.async.cg.shared.global [%0], [%1], 16;\n"
+               :: "r"(smem_dst), "l"(gmem_src) : "memory");
+}
+
+__device__ __forceinline__
+void cp_async_commit_group() {
+  asm volatile("cp.async.commit_group;\n" ::: "memory");
+}
+
+template <int N>
+__device__ __forceinline__
+void cp_async_wait_group() {
+  asm volatile("cp.async.wait_group %0;\n" :: "n"(N) : "memory");
+}
+
+__device__ __forceinline__
+void griddepcontrol_wait() {
+  asm volatile("griddepcontrol.wait;\n" ::: "memory");
+}
+
+__device__ __forceinline__
+void griddepcontrol_launch_dependents() {
+  asm volatile("griddepcontrol.launch_dependents;\n" ::: "memory");
+}
@@ -46,6 +46,10 @@ extern std::vector<torch::Tensor> block_sparse_sm100a_bwd(
    torch::Tensor grad_o, torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o,
    torch::Tensor lse, torch::Tensor k2q_idx, torch::Tensor k2q_num,
    torch::Tensor variable_block_sizes, double sm_scale);
+extern std::vector<torch::Tensor> block_sparse_sm100a_blk128_bwd(
+    torch::Tensor grad_o, torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o,
+    torch::Tensor lse, torch::Tensor k2q_idx, torch::Tensor k2q_num,
+    torch::Tensor variable_block_sizes, double sm_scale);
 #endif

 PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
@@ -60,7 +64,10 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
          "VSA block-sparse attention forward, 128-token blocks (Blackwell sm100a/sm103a)");
    m.def("block_sparse_sm100a_bwd",
          torch::wrap_pybind_function(block_sparse_sm100a_bwd),
-          "VSA block-sparse attention backward, 64-token blocks (Blackwell sm100a)");
+          "VSA block-sparse attention backward, 64-token blocks (Blackwell sm100a/sm103a)");
+    m.def("block_sparse_sm100a_blk128_bwd",
+          torch::wrap_pybind_function(block_sparse_sm100a_blk128_bwd),
+          "VSA block-sparse attention backward, 128-token blocks (Blackwell sm100a/sm103a)");
 #endif

 #ifdef TK_COMPILE_ST_ATTN
@@ -31,6 +31,11 @@ auto quant(
  std::optional<torch::Tensor>& Output_S
 ) {

+  TORCH_CHECK(
+    Input.scalar_type() == torch::kHalf || Input.scalar_type() == torch::kBFloat16,
+    "Unsupported input data type for quant_cuda (int8 quantization)."
+  );
+
  using ElementOut = int8_t;
  static constexpr int BlockSize = 128;
  static constexpr int NumThrPerCta = 256;
@@ -62,8 +67,7 @@ auto quant(
    }

    default: {
-      std::cerr << "Observing: " << Input.scalar_type() << " for the input datatype which is invalid";
-      throw std::runtime_error("Unsupported input data type for quantize_to_fp4.");
+      TORCH_CHECK(false, "Unsupported input data type for quant_cuda (int8 quantization).");
    }
  }
  
@@ -52,15 +52,31 @@ def _force_tk() -> bool:
 def _force_sm100a() -> bool:
    """True iff the data-center Blackwell forward is explicitly opted into.

-    Opt-in only (same legacy-named env the H3 backend honors): the 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.
+    Opt-in only (same legacy-named env the H3 backend honors). Honored only when
+    ``block_sparse_attn_sm100a.is_supported`` passes; the backward then runs the
+    sm_100a/sm_103a CUDA backward when ``block_sparse_attn_bwd_sm100a.is_supported``
+    passes and Triton otherwise (the forward's lse is already in Triton's M
+    format). 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"
+    return _env_bool("FASTVIDEO_VSA_SM100A")
+
+
+def _env_bool(name: str) -> bool:
+    """Parse a boolean environment variable with the rule of ``fastvideo.envs.EnvBool``.
+
+    fastvideo-kernel cannot import fastvideo, so the rule is repeated here to keep
+    the kernel and the package in agreement: 1, true, yes, on are true; 0, false,
+    no, off, and the empty string are false; case-insensitive; anything else raises.
+    """
+    value = os.environ.get(name, "0").strip().lower()
+    if value in ("1", "true", "yes", "on"):
+        return True
+    if value in ("0", "false", "no", "off", ""):
+        return False
+    raise ValueError(f"Invalid value {value!r} for {name}: expected 1, true, yes, on, 0, false, no, off, "
+                     "or an empty string")


 def _sm100a_is_supported(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> bool:
@@ -373,11 +389,12 @@ block_sparse_attn_sm90.register_autograd(_backward_sm90, setup_context=_setup_co
 # Data-center Blackwell backend custom op (index-native; legacy sm100a API name)
 #
 # Forward runs the sm_100a/sm_103a CUDA extension. Backward runs the sm_100a CUDA
-# backward when block_sparse_attn_bwd_sm100a.is_supported passes (64-token blocks,
-# sm_100a device, extension built with the op) and the Triton kernels otherwise.
-# The native forward emits lse in exactly Triton's M format (max*log2e +
-# log2(l)), so either pairing needs no conversion. Both backwards are
-# hardcoded to 64-token blocks, hence the block-size assert below.
+# backward when block_sparse_attn_bwd_sm100a.is_supported passes (64- or 128-token
+# blocks, data-center Blackwell device, extension built with that block's op) and
+# the Triton kernels otherwise. The native forward emits lse in exactly Triton's M
+# format (max*log2e + log2(l)), so either pairing needs no conversion. The Triton
+# backward is hardcoded to 64-token blocks, so unsupported 128-token metadata
+# raises instead of falling back.
 # ---------------------------------------------------------------------------


@@ -477,16 +494,17 @@ def _setup_context_sm100a(ctx, inputs, output):

 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/sm_103a forward with a "
-            f"backward that is hardcoded to 64-token blocks; got {block}. "
-            "Run 128-token-block metadata without grad, or use the Triton forward.")
    if _sm100a_backward_is_supported(q, variable_block_sizes):
        dq, dk, dv = block_sparse_attn_backward_sm100a(grad_o, q, k, v, o, M, q2k_idx,
                                                       q2k_num, variable_block_sizes)
    else:
+        block = q.shape[2] // variable_block_sizes.numel()
+        if block != 64:
+            raise RuntimeError(
+                "block_sparse_attn_sm100a backward: no sm_100a/sm_103a backward for "
+                f"{block}-token blocks on this build/device, and the Triton backward is "
+                "hardcoded to 64-token blocks. Run this metadata without grad, or build "
+                "the extension with block_sparse_sm100a_blk128_bwd.")
        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
@@ -520,8 +538,9 @@ def block_sparse_attn_from_indices(

    # Backend resolution:
    # - FASTVIDEO_VSA_TRITON forces Triton everywhere.
-    # - FASTVIDEO_VSA_SM100A opts into the data-center Blackwell forward
-    #   (Triton backward). The environment name is retained for compatibility.
+    # - FASTVIDEO_VSA_SM100A opts into the data-center Blackwell forward (CUDA
+    #   backward when supported, else Triton). The environment name is retained
+    #   for compatibility.
    #   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
@@ -1,11 +1,13 @@
 # SPDX-License-Identifier: Apache-2.0
-"""sm_100a (Blackwell) CUDA block-sparse VSA backward.
+"""sm_100a/sm_103a (data-center Blackwell) CUDA block-sparse VSA backward.

 Companion of ``block_sparse_attn_sm100a`` (the forward): consumes the forward's ``lse`` in the
 Triton "M format" (``max(qk * sm_scale * log2e) + log2(l)``, ``[B, H, S]`` fp32) unchanged and
 FastVideo's k2q index metadata, returns ``(dq, dk, dv)`` in bf16 with the inputs' layout and the
-Triton backward's scaling (dq and dk carry sm_scale, dv does not). 64-token blocks only; any
-other configuration falls back to Triton via ``is_supported``.
+Triton backward's scaling (dq and dk carry sm_scale, dv does not). Two kernels, picked by the
+metadata's block size: 64-token blocks (``block_sparse_sm100a_bwd``) and 128-token blocks
+(``block_sparse_sm100a_blk128_bwd``); any other configuration falls back to Triton via
+``is_supported``.
 """

 from typing import Tuple
@@ -16,62 +18,83 @@ try:
    # The pybind symbols live on fastvideo_kernel_ops, NOT on the _C package that contains it
    # (its __init__ is empty, so hasattr on the package fails with the kernel built and present).
    from fastvideo_kernel._C import fastvideo_kernel_ops as _C
-    _BWD = getattr(_C, "block_sparse_sm100a_bwd", None)
-    _HAS_VSA_BWD_SM100A = _BWD is not None
+    _BWD_BY_BLOCK = {
+        64: getattr(_C, "block_sparse_sm100a_bwd", None),
+        128: getattr(_C, "block_sparse_sm100a_blk128_bwd", None),
+    }
+    _HAS_VSA_BWD_SM100A = any(_BWD_BY_BLOCK.values())
 except ImportError:  # pragma: no cover - extension not built
    _C = None
-    _BWD = None
+    _BWD_BY_BLOCK = {}
    _HAS_VSA_BWD_SM100A = False

-_SM100 = (10, 0)
+_SUPPORTED_COMPUTE_CAPABILITIES = {(10, 0), (10, 3)}
 HEAD_DIM = 128
-BLOCK = 64
+BLOCKS = (64, 128)
 # Must match the -DVSA_BHSD the extension was compiled with (FastVideo builds with true).
 BHSD = True


 def set_extension(module) -> None:
-    """Use an already-loaded extension module exposing ``block_sparse_sm100a_bwd``.
+    """Use an already-loaded extension module exposing ``block_sparse_sm100a_bwd`` and / or
+    ``block_sparse_sm100a_blk128_bwd``.

-    A standalone build of ``block_sparse_bwd_sm100a.cu`` (for example through
+    A standalone build of the binding .cu files (for example through
    ``torch.utils.cpp_extension.load`` with a ten-line pybind wrapper) can be injected here, so
    the backend can be exercised without rebuilding the fastvideo_kernel wheel.
    """
-    global _C, _BWD, _HAS_VSA_BWD_SM100A
+    global _C, _BWD_BY_BLOCK, _HAS_VSA_BWD_SM100A
    _C = module
-    _BWD = getattr(module, "block_sparse_sm100a_bwd", None)
-    _HAS_VSA_BWD_SM100A = _BWD is not None
+    _BWD_BY_BLOCK = {
+        64: getattr(module, "block_sparse_sm100a_bwd", None),
+        128: getattr(module, "block_sparse_sm100a_blk128_bwd", None),
+    }
+    _HAS_VSA_BWD_SM100A = any(_BWD_BY_BLOCK.values())


 def _seqlen(q: torch.Tensor) -> int:
    return q.shape[2] if BHSD else q.shape[1]


+def _block_size(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> int:
+    """The block size the metadata implies (seqlen / num_blocks), 0 when it is not integral.
+
+    A plain int: under dynamic-shape torch.compile the shapes are SymInts, which cannot key
+    ``_BWD_BY_BLOCK``; ``int()`` specializes the value with a shape guard.
+    """
+    num_blocks = variable_block_sizes.numel()
+    seqlen = _seqlen(q)
+    return 0 if num_blocks == 0 or seqlen % num_blocks else int(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), never tensor contents, so it is cheap
-    enough for a per-layer dispatch path. The kernel is fixed at 64-token blocks with
-    head_dim 128 and needs seqlen == 64 * num_blocks with an even num_blocks (its preprocess
-    works in 128-token blocks). Per-row k2q counts may be anything in [0, num_q_blocks],
-    including 0: unselected kv blocks get exactly-zero dk/dv rows.
+    enough for a per-layer dispatch path. Both kernels take head_dim 128 and seqlen == block *
+    num_blocks; the 64-token kernel also needs an even num_blocks (its preprocess works in
+    128-token blocks); batch * heads is the launch grid's y extent, so it must be <= 65535.
+    Per-row k2q counts may be anything in [0, num_q_blocks], including 0: unselected kv blocks
+    get exactly-zero dk/dv rows.
    """
    if not _HAS_VSA_BWD_SM100A or not q.is_cuda:
        return False
-    if torch.cuda.get_device_capability(q.device) != _SM100:
+    if torch.cuda.get_device_capability(q.device) not in _SUPPORTED_COMPUTE_CAPABILITIES:
        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 q.shape[0] * (q.shape[1] if BHSD else q.shape[2]) > 65535:
+        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
-    num_blocks = variable_block_sizes.numel()
-    if num_blocks == 0 or num_blocks % 2 != 0:
+    block = _block_size(q, variable_block_sizes)
+    if _BWD_BY_BLOCK.get(block) is None:
        return False
-    if _seqlen(q) != BLOCK * num_blocks:
+    if block == 64 and variable_block_sizes.numel() % 2 != 0:
        return False
    return True

@@ -89,15 +112,21 @@ def block_sparse_attn_backward_sm100a_from_k2q(
 ) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
    """Backward from k2q metadata already in hand (``invert_indices`` layout).

-    ``k2q_idx`` is ``[B, H, num_kv_blocks, max_q_blocks]`` (or the flat 2-D view) of LOCAL q64
-    block ids, ``k2q_num`` ``[B, H, num_kv_blocks]``; entries past a row's count are never read.
+    ``k2q_idx`` is ``[B, H, num_kv_blocks, max_q_blocks]`` (or the flat 2-D view) of LOCAL q
+    block ids (64- or 128-token blocks, as the metadata's block size), ``k2q_num``
+    ``[B, H, num_kv_blocks]``; entries past a row's count are never read.
    """
+    block = _block_size(q, variable_block_sizes)
+    backward = _BWD_BY_BLOCK.get(block)
+    if backward is None:
+        raise RuntimeError(f"block_sparse_attn_backward_sm100a: no kernel for {block}-token "
+                           f"blocks (built: {sorted(b for b, f in _BWD_BY_BLOCK.items() if f)})")
    sm_scale = 1.0 / (q.shape[-1]**0.5)
    idx = k2q_idx.to(torch.int32).contiguous()
    num = k2q_num.to(torch.int32).contiguous()
    vbs = variable_block_sizes.to(torch.int32).contiguous()
-    res = _BWD(grad_o.contiguous(), q.contiguous(), k.contiguous(), v.contiguous(),
-               o.contiguous(), lse.contiguous(), idx, num, vbs, sm_scale)
+    res = backward(grad_o.contiguous(), q.contiguous(), k.contiguous(), v.contiguous(),
+                   o.contiguous(), lse.contiguous(), idx, num, vbs, sm_scale)
    return res[0], res[1], res[2]


@@ -8,7 +8,7 @@ A third backend behind the same VSA op as the Triton and CuTe-DSL paths. This mo
 forward: 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 both ``block_sparse_attn_backward_triton`` and the sm_100a CUDA backward
-(``block_sparse_attn_bwd_sm100a``, 64-token blocks, sm_100a devices only) run against it
+(``block_sparse_attn_bwd_sm100a``, 64- and 128-token blocks, sm_100a / sm_103a) run against it
 unchanged.

 The extension carries TWO instantiations of the kernel, for 64- and 128-token sparse blocks
@@ -43,9 +43,10 @@ BHSD = True


 def _block_size(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> int:
+    # A plain int: dynamic-shape torch.compile hands SymInt shapes, which cannot key _FWD_BY_BLOCK.
    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
+    return 0 if num_blocks == 0 or seqlen % num_blocks else int(seqlen // num_blocks)


 def is_supported(q: torch.Tensor, variable_block_sizes: torch.Tensor) -> bool:
@@ -1054,6 +1054,15 @@ class _attention(torch.autograd.Function):
        use_global_sf_P=True,
        use_global_sf_QKV=True,
    ):
+        if smooth_q:
+            # Not implemented, not merely disabled: only the backward kernels
+            # carry a SMOOTH_Q path (_attn_fwd has none), so restoring the
+            # triton_group_mean wiring would yield silently wrong gradients
+            # rather than a feature. The flag is kept for signature
+            # compatibility only.
+            raise NotImplementedError("smooth_q=True is not supported: the forward kernel has no "
+                                      "SMOOTH_Q path, so enabling it would produce incorrect gradients")
+
        # shape constraints
        HEAD_DIM_Q, HEAD_DIM_K = q.shape[-1], k.shape[-1]
        # when v is in float8_e5m2 it is transposed.
@@ -1364,10 +1373,8 @@ class _attention(torch.autograd.Function):
                                       BLOCK_M=PRE_BLOCK,
                                       HEAD_DIM=ctx.HEAD_DIM)

+        # smooth_q is rejected in forward(); ctx.smooth_q is always False here.
        q_m = None
-        if ctx.smooth_q:
-            # _, q_m = triton_group_mean(q)
-            q_m = q_m.repeat_interleave(q.shape[2] // q_m.shape[2], dim=2)  # B,H,L,D

        sm100_optimized_backward = (getattr(ctx, "sm100_optimized", False) and ctx.use_qat_qkv_backward
                                    and not ctx.smooth_k and not ctx.smooth_q and N_CTX_KV % 16 == 0)
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Satyam Srivastava	0ef47caf86	[bugfix]: select architecture-specific Blackwell kernel targets - Detect GB200 and GB300 as 10.0a and 10.3a - Build future GB200 CI images with 10.0a - Add regression tests for detection, cache keys, and build arguments	2026-10-05 15:42:48 -07:00
Satyam Srivastava	e71d01648c	[feat]: add selectable Modal and GB200 Kubernetes GPU CI - Preserve existing CI and add configurable GPU backend routing - Limit execution to 2 active PRs, 4 GPUs per PR, and 8 GPUs total - Add persistent admission, cancellation recovery, and backend statuses - Document deployment and add infrastructure regression tests	2026-10-01 14:44:14 -07:00
lpc0220 Pengcheng Li Claude Fable 5.1	e1f3904799	[kernel] sm_100a CUDA backward for VSA block-sparse attention, 128-token blocks (#1866 ) Co-authored-by: Pengcheng Li <pengchengl@ptyche0203.ptyche.clusters.nvidia.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>	2026-09-29 16:57:22 -07:00
Junda Su	02f1ce11ae	[misc] Move environment var reads into env registry and rename unprefixed variables (#1897 )	2026-09-29 14:36:16 -07:00
Aryan KumarandAryan Kumar	7f03e03dc6	[docs]: FastH3 V2 cookbook serving and MLX install guide (#1884 ) Co-authored-by: Aryan Kumar <aryan5v@users.noreply.github.com>	2026-09-29 13:03:18 -07:00
Junda Su	e3b88bb12a	[feat] Add a typed environment-variable registry, policy doc, and contract test (#1896 )	2026-09-29 12:50:45 -07:00
Junda Su	cb66acd400	[bugfix] Fix environment-variable bugs in LTX-2 debug logging, HF token lookup, and attention backend reads (#1895 )	2026-09-29 11:48:31 -07:00
li-lizhe	442e2d2e18	[bugfix] fix(metrics): set SceneMetric device_map for any non-CPU device (#1817 )	2026-09-28 10:24:00 -07:00
Max LI	dd35763ad6	[bugfix] Fix MLX prompt enhancement sampler compatibility (#1891 )	2026-09-27 14:53:25 -07:00
Keith	e90be598e5	[perf] MiniMax H3: return uint8 frames from the decode worker (#1828 )	2026-09-24 18:08:56 -07:00
Leleand武垚乐	ba5e81083c	[bugfix] Preserve video frames when decoded audio is shorter (#1857 ) Signed-off-by: 武垚乐 <wuyaole@mininglamp.com> Co-authored-by: 武垚乐 <wuyaole@mininglamp.com>	2026-09-24 17:20:37 -07:00
YZJF	76ce9c7fd6	[bugfix] Reject mismatched model names on image generation routes (#1881 )	2026-09-24 17:19:29 -07:00
Junda Su	08d99c089e	[perf] Reuse cached VAE offload for H3 conditioning (#1882 )	2026-09-24 17:18:39 -07:00
Jerry-XY	20751a21aa	[bugfix] Keep the VSA-H3 tile buffer out of the autograd graph (#1858 )	2026-09-24 17:16:49 -07:00
alanhuangyooandSolitaryThinker	9dd2a837f4	[bugfix] encoders: keep Qwen2.5-VL importable on transformers 5 (#1791 ) Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com>	2026-09-24 17:13:23 -07:00
alanhuangyoo	93aab45ac2	[bugfix] tests: include ltx2_3_base in local LTX-2 preset set (#1750 )	2026-09-24 17:11:31 -07:00
Yogya Mehrotra	017ce6602d	[bugfix] Fix dead NVFP4 QAT plumbing: bench imports, int8 error string, smooth_q crash (#1723 )	2026-09-24 17:09:33 -07:00
Yogya Mehrotra	a575055eec	[perf] benchmark_attn_qat_train: device peak-TFLOPS table instead of silent RTX 5090 default (#1724 )	2026-09-24 17:08:03 -07:00
Kyle Hu	81f3fec7fd	[bugfix]: workers killed by a signal now log the reason (#1725 )	2026-09-24 17:06:55 -07:00
Raghav K	d265a454bf	[bugfix]: Fall back when large pinned output allocation fails (#1759 )	2026-09-24 17:05:59 -07:00
Xingjian Guo Gxj230958 SolitaryThinker	c100c66578	[bugfix]: isolate cancelled streaming requests from GPU result readers (#1848 ) Co-authored-by: Gxj230958 <222823329+Gxj230958@users.noreply.github.com> Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com>	2026-09-24 17:04:59 -07:00
Aryan Kumar	8760eb7a06	[feat] Add FastH3 8-Step V2 MLX inference (#1863 ) Keep the four-step uniform AdaLN cache, and use explicit DMD rungs only when the snapshot declares its own scheduler shifts.	2026-09-23 12:42:03 -07:00
Aryan Kumar	361f919c88	[feat]: add a narrow FastH3 8-step MLX ladder Keep the four-step uniform AdaLN cache, and use explicit DMD rungs only when the snapshot declares its own scheduler shifts.	2026-09-23 11:06:55 -07:00
Aryan Kumar	8b5377aab2	Merge remote-tracking branch 'origin/main' into aryan/pr-1863-fasth3-8step	2026-09-23 10:52:20 -07:00
Yixing Wangandleo	d995516da0	[bugfix] Restore page interaction after dismissing Create Job or successfully creating a job (#1869 ) Co-authored-by: leo <yixingwang@YIXINGs-MacBook-Pro.local>	2026-09-21 12:03:54 -07:00
Keith	f47ad3f5b7	[bugfix] fastvideo-kernel: install the Python + Triton package when the HIP toolchain cannot be configured (#1871 )	2026-09-21 11:59:13 -07:00
Hexu ZhaoandClaude Opus 5	10bdf5e076	[bugfix] frozen offload: no host copy for an already resident module `load` took the host copies before checking which tensors were missing from the device, so with `vae_cpu_offload=False` — where the loader builds the VAEs on CUDA and `unload` is never called — it pulled about 11 GB per rank back over PCIe once and kept a pinned mirror of it for the life of the process, for a module that never leaves the device. It now asks what is absent first and returns early when nothing is. The offloaded path is unchanged: weights still round-trip and come back bit identical. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-09-20 22:46:13 -07:00
Hexu ZhaoandClaude Opus 5	3e26db40b0	[perf] H3 decode: pin the frozen VAEs' host copies With the device-to-host copy gone, the remaining host-to-device copy is the stage's cost, and from pageable memory it is staged through a bounce buffer at roughly 2.6 GB/s instead of full PCIe speed. The host copies are made once, so pinning them costs nothing per request. It does cost the module's size in non-pageable host memory (10.4 GB for the H3 video VAE, 0.6 GB for the audio VAE, per rank) for as long as the pipeline is alive. It follows --pin-cpu-memory, which is on by default; --no-pin-cpu- memory keeps the copies pageable and only loses the bandwidth. Note that this flag did not already reach the VAE: its other use is FSDP's own CPUOffloadPolicy and the VAE does not go through FSDP. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01ATw7g6cxq9N2LtnMratsru	2026-09-20 22:46:13 -07:00
Hexu ZhaoandClaude Opus 5	c73dd0ab55	[perf] H3 decode: offload the frozen VAEs without the device-to-host copy The decode stages move the 10.4 GB fp32 video VAE (and the 0.6 GB audio VAE) with module.to(device) / module.to("cpu") on every request. The weights are frozen -- the whole stage runs under no_grad -- so the copy back to the host returns bytes the host already has. Keep one host copy per tensor, made the first time the module is loaded, and on unload point .data back at it instead of copying. One H2D per request, no D2H at all; the module still lives on the host between requests. Where the old path allocated a fresh host tensor per cycle and freed the previous one, this reuses one buffer, so peak host memory can only go down. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01ATw7g6cxq9N2LtnMratsru	2026-09-20 22:46:13 -07:00
primorLee William Lin SolitaryThinker	430e52154e	[bugfix]: Restore exact Qwen3-VL vision interpolation (#1737 ) Co-authored-by: William Lin <8941107+SolitaryThinker@users.noreply.github.com> Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com>	2026-09-18 05:00:05 -07:00
vanch	384eee8aef	feat(mlx): add FastH3-8-Step-V2 DMD schedule and thermal-safe MLX inference - Support official 8-step DMD schedule from fastvideo_inference.json in MLX pipeline - Implement MiniMaxH3SchedulerState.from_dmd_steps for exact sigma ladder calculation - Add auto-detection of fastvideo_inference.json in checkpoint converter AdaLN precomputation - Add inter_step_cooldown_s and per-step MLX graph/cache cleanup to prevent thermal throttling - Support configurable VAE tile sizes to guarantee 256px tiled decode without grid artifacts - Skip non-existent shard files gracefully in Qwen3-VL conditioner	2026-09-17 16:12:02 +08:00
KeithandSolitaryThinker	c4824c7764	[bugfix] FP8: gate the ROCm _scaled_mm path on CDNA4 (gfx950) instead of the capability tuple (#1859 ) Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com>	2026-09-16 19:58:55 -07:00
Raghav K	9b0e57fe4b	[ci] Make Dreamverse provider race test deterministic (#1729 )	2026-09-15 14:13:56 -07:00
William Lin	0100218594	[feat] Support the FastH3 8-Step V2 checkpoint: checkpoint-defined shifts, explicit DMD schedule, new example (#1852 )	2026-09-15 14:13:45 -07:00
li-lizhe	39718cd54d	[bugfix] fix(cosmos): make AdaLayerNorm autocast device-agnostic (#1818 )	2026-09-15 07:44:51 -07:00
IshanandSolitaryThinker	37d06a832f	[feat] Add fastvideo serve configs for Wan CUDA models (#1801 ) Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com>	2026-09-14 18:23:35 -07:00
sudhirpol522	8839ba8d4d	[bugfix] Add OpenAI-compatible image generation endpoint (#1840 )	2026-09-14 17:55:56 -07:00
sudhirpol522andSolitaryThinker	61b91220c0	[bugfix] Validate image response format before generation (#1841 ) Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com>	2026-09-14 17:01:40 -07:00
Lele	316f3876c2	[bugfix]: allow MiniMax H3 frame padding at the 15-second limit Accept the causal-VAE-aligned 362-frame bucket (15.083 s) for 15-second H3 requests. - Hoist MINIMAX_H3_MIN/MAX_ALIGNED_FRAMES next to align_num_frames and reuse them in the CUDA stage, the MLX runtime, and the LoRA example so the bound is consistent across entry points. - State the accepted frame range in the error messages. - Cover seconds="15" -> 362 at OpenAI admission and the MLX resolve_geometry bound. - Update the Spark/openai cookbook frame caps from 345 to 362. Known gaps: no GPU-lane coverage for the 362 bucket (SSIM runs 124 frames; the golden gate is a fixed-geometry fingerprint), and explicit num_frames=360 still hits the pre-existing grid gate.	2026-09-14 16:54:41 -07:00
Yaegaki1Erika	614b59543c	[bugfix]: respect serialized tensor dtype in parquet dataloader (#1843 )	2026-09-14 16:34:25 -07:00
William Lin	1c14afd559	[docs]: refresh AGENTS.md maps and add Wan SP/I2V and CI test notes (#1845 )	2026-09-14 16:06:53 -07:00
Yaegaki1Erika	9a3c45779c	[bugfix]: fix Wan I2V DMD conditioning under sequence parallelism (#1844 ) The pipeline pre-sharded only the image conditioning along the temporal axis while the noise input stayed full-length, so concatenation could not match for sp_world_size > 1. WanTransformer3DModel shards the flattened token sequence after patch embedding, so pass full-length mask+latent conditioning and let the transformer shard. Also reads temporal_compression_ratio from the VAE config, simplifies the conditioning mask, concatenates in the transformer's (bs, c, t, h, w) layout, and adds unit coverage. Co-authored-by: Yaegaki1Erika <70182590+Yaegaki1Erika@users.noreply.github.com>	2026-09-14 15:58:13 -07:00
Kyle Hu	bfc9c01797	[feat]: convert the MiniMax H3 text encoder to NVFP4 (#1838 )	2026-09-12 18:48:07 -07:00
Kyle Hu	3a3ad3d209	[feat]: NVFP4 text encoder for MiniMax H3 (#1837 )	2026-09-12 18:31:59 -07:00
lpc0220andClaude Fable 5.1	aef4e9b3b1	[kernel] VSA kernel: one sm_100a / sm_103a image per listed arch; un-gate backward on sm_103a (#1833 ) Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>	2026-09-12 16:49:36 -07:00