[ci]: collect fastvideo/tests/stages in the unit-test lane (was never run by any CI job)

Co-authored-by: SolitaryThinker <wlsaidhi@gmail.com>

.buildkite/performance-benchmarks/tests/wan-t2v-1.3b.json

@@ -1,5 +1,9 @@
 {
   "benchmark_id": "wan-t2v-1.3b-2gpu",
+  "config_schema_version": 2,
+  "workload_id": "wan-t2v-1.3b",
+  "variant_id": "canonical",
+  "benchmark_version": 1,
   "description": "Wan2.1 T2V 1.3B inference performance",
   "model": {
     "model_path": "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",

.buildkite/pipeline.yml

@@ -114,6 +114,17 @@ steps:
             limit: 2
       agents:
         queue: "default"
+    - label: ":test_tube: LoRA Extraction Tests"
+      if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "lora_extraction"
+      command: "timeout 90m .buildkite/scripts/pr_test.sh"
+      retry:
+        automatic:
+          - exit_status: 128
+            limit: 3
+          - exit_status: -1
+            limit: 2
+      agents:
+        queue: "default"
     - label: ":test_tube: Training Tests"
       if: build.env("TEST_SCOPE") == "direct" && build.env("TEST_TYPE") == "training"
       command: "timeout 90m .buildkite/scripts/pr_test.sh"
@@ -371,6 +382,21 @@ steps:
                   - TEST_TYPE=inference_lora
                 agents:
                   queue: "default"
+            - path:
+                - "scripts/lora_extraction/**"
+                - "fastvideo/tests/lora_extraction/**"
+                - "fastvideo/models/loader/**"
+                - "fastvideo/training/training_utils.py"
+                - "fastvideo/layers/lora/**"
+                - "pyproject.toml"
+                - "docker/Dockerfile"
+              config:
+                command: "timeout 90m .buildkite/scripts/pr_test.sh"
+                label: ":test_tube: LoRA Extraction Tests"
+                env:
+                  - TEST_TYPE=lora_extraction
+                agents:
+                  queue: "default"
             - path:
                 - "fastvideo/**"
                 - "pyproject.toml"

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

@@ -125,7 +125,7 @@ jobs:
           set -euo pipefail
           TEST_NAME=$(echo "$COMMENT" | grep -oP '(?<=/test\s)\S+' | head -1 || true)
 
-          VALID="encoder vae transformer kernel unit dreamverse ssim training lora-inference lora-training distillation self-forcing vsa vmoba performance api train-framework eval full fastcheck pre-commit"
+          VALID="encoder vae transformer kernel unit dreamverse ssim training lora-inference lora-training lora-extraction distillation self-forcing vsa vmoba performance api train-framework eval full fastcheck pre-commit"
           if [ -z "$TEST_NAME" ] || ! echo "$VALID" | grep -qw "$TEST_NAME"; then
             echo "Unknown test: '$TEST_NAME'. Valid: $VALID"
             exit 1
@@ -136,6 +136,7 @@ jobs:
             [kernel]=kernel_tests [unit]=unit_test [dreamverse]=dreamverse_app
             [ssim]=ssim [training]=training
             [lora-inference]=inference_lora [lora-training]=training_lora
+            [lora-extraction]=lora_extraction
             [distillation]=distillation_dmd [self-forcing]=self_forcing
             [vsa]=training_vsa [vmoba]=inference_vmoba
             [performance]=performance [api]=api_server

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

@@ -13,12 +13,33 @@ on:
         required: false
         default: false
         type: boolean
+  # Auto-rebuild the CUDA images when their Dockerfile changes on main. The CUDA
+  # matrix is the only lane that builds from docker/Dockerfile, so a path-scoped
+  # push trigger is a sufficient change detector on its own -- no separate
+  # detect-changes/paths-filter job is needed now that there is a single
+  # in-scope Dockerfile. Dreamverse (apps/dreamverse/docker/Dockerfile) and the
+  # rocm Dockerfile stay manual-dispatch only.
+  push:
+    branches: [main]
+    paths:
+      - 'docker/Dockerfile'
 
 
 permissions:
   contents: read
   packages: write
 
+# One static group, no cancellation: every run of this workflow writes the same
+# mutable registry tags (latest, py3.12-latest, ...), so runs must serialize —
+# concurrent push/dispatch runs would race on those tags, and cancelling a run
+# mid-publish can strand the cu126/cu130 tag families at different commits. An
+# in-flight superseded build wastes its runner time, but its tags are then
+# overwritten by the newer queued run. GitHub keeps a single pending run per
+# group: the newest queued run replaces any older queued one.
+concurrency:
+  group: infra-build-image
+  cancel-in-progress: false
+
 jobs:
   # CUDA matrix: Python 3.12 x {12.6.3, 13.0.0} x {amd64, arm64}. Each architecture
   # builds natively and pushes only by digest; publish-cuda-manifests is the sole
@@ -28,7 +49,11 @@ jobs:
   # aliases; 13.0.0/cu130 is published under explicit versioned tags. Flash-attn
   # 2.8.3 comes from the architecture-specific prebuilt releases.
   build-cuda-images:
-    if: ${{ github.event.inputs.build_cuda_matrix == 'true' }}
+    # Runs on a manual dispatch when build_cuda_matrix is set, or automatically
+    # on a push that changed docker/Dockerfile (inputs are null on push). The
+    # repository guard keeps fork syncs from auto-building; manual dispatch
+    # still works in forks.
+    if: ${{ (github.event_name == 'push' && github.repository == 'hao-ai-lab/FastVideo') || github.event.inputs.build_cuda_matrix == 'true' }}
     strategy:
       fail-fast: false
       matrix:
@@ -75,10 +100,10 @@ jobs:
     secrets: inherit
 
   publish-cuda-manifests:
-    # !cancelled(): a failed sibling build leg must not skip the manifests for a
-    # CUDA lane whose own digests all exist; the digest-count check below fails
-    # the incomplete lane loudly instead.
-    if: ${{ !cancelled() && github.event.inputs.build_cuda_matrix == 'true' }}
+    # !cancelled(): publish lanes whose digests exist even if a sibling build
+    # leg failed (the digest-count check fails incomplete lanes); it also
+    # bypasses skipped-needs propagation, hence the explicit skipped check.
+    if: ${{ !cancelled() && needs.build-cuda-images.result != 'skipped' }}
     needs: build-cuda-images
     runs-on: ubuntu-latest
     permissions:

apps/dreamverse/docker/Dockerfile

@@ -84,9 +84,12 @@ RUN source /opt/venv/bin/activate \
        FFMPEG_NATIVE_CXX=/usr/bin/g++ \
        bash /opt/FastVideo/apps/dreamverse/scripts/install_native_ffmpeg.sh
 
+# FASTVIDEO_FA4: FA4 (flash_attn.cute) is opt-in; this image installs it via
+# the dreamverse extra and is validated with it, so enable it here.
 ENV FASTVIDEO_DREAMVERSE_HOME=/var/lib/dreamverse \
     STREAM_MODE=av_fmp4 \
     FASTVIDEO_ENABLE_PROMPT_SAFETY=0 \
+    FASTVIDEO_FA4=1 \
     HF_HOME=/root/.cache/huggingface
 
 RUN mkdir -p /var/lib/dreamverse

docker/Dockerfile

@@ -170,12 +170,14 @@ RUN --mount=type=cache,target=/opt/uv/cache \
 # rmtree clears the wheel's stale cute files first to avoid an install conflict.
 # Then verify both survive so a broken overlay fails the build instead of shipping
 # an FA2-less image. x86 only: the FA4 stack (quack-kernels etc.) is unvalidated on
-# arm64 / GB10 (sm_121), so there we skip the overlay and FA4 falls back to FA2.
+# arm64 / GB10 (sm_121), so there we skip the overlay; FA4 is opt-in
+# (FASTVIDEO_FA4=1) and errors if set without the overlay, so leave it unset on
+# arm64 and the image runs FA3/FA2 as usual.
 RUN --mount=type=cache,target=/opt/uv/cache \
     source $HOME/.local/bin/env && \
     source /opt/venv/bin/activate && \
     if [ "${TARGETARCH}" = "arm64" ]; then \
-        echo "Skipping FA4 cute overlay on arm64 (FA4 stack unvalidated there; FA4 falls back to FA2)"; \
+        echo "Skipping FA4 cute overlay on arm64 (FA4 stack unvalidated there; do not set FASTVIDEO_FA4)"; \
     else \
         python -c "import glob, shutil; [shutil.rmtree(d, ignore_errors=True) for d in glob.glob('/opt/venv/lib/python*/site-packages/flash_attn/cute')]" && \
         uv pip install "flash-attn-4 @ git+https://github.com/Dao-AILab/flash-attention.git@${FA4_CUTE_REF}#subdirectory=flash_attn/cute" && \

docs/contributing/ci_architecture.md

@@ -103,6 +103,7 @@ can merge a PR.
 |---|---|---|
 | SSIM Tests | `ssim` | `fastvideo/**/*.py`, `pyproject.toml`, `docker/Dockerfile` |
 | LoRA Inference Tests | `inference_lora` | LoRA tests, loader, transformer tests, pipelines, LoRA layers |
+| LoRA Extraction Tests | `lora_extraction` | LoRA extraction scripts/tests, loader, training utilities, LoRA layers |
 | Training Tests | `training` | `fastvideo/**`, `pyproject.toml`, `docker/Dockerfile` |
 | Distillation DMD Tests | `distillation_dmd` | `fastvideo/training/*distillation_pipeline.py` |
 | Self-Forcing Tests | `self_forcing` | self-forcing distillation pipeline and tests |
@@ -144,6 +145,7 @@ Valid direct test names:
 | `/test training` | `training` |
 | `/test lora-inference` | `inference_lora` |
 | `/test lora-training` | `training_lora` |
+| `/test lora-extraction` | `lora_extraction` |
 | `/test distillation` | `distillation_dmd` |
 | `/test self-forcing` | `self_forcing` |
 | `/test vsa` | `training_vsa` |

docs/contributing/performance_benchmarks.md

@@ -172,6 +172,45 @@ agent skill to advance the rolling median.
 
 ## Schemas
 
+### Benchmark config (`.buildkite/performance-benchmarks/tests/*.json`)
+
+Benchmark configs without `config_schema_version` are treated as legacy v1
+configs and remain loadable. New or migrated configs should use
+`config_schema_version: 2` and include explicit comparable identity fields:
+
+```jsonc
+{
+  "benchmark_id": "wan-t2v-1.3b-2gpu",
+  "config_schema_version": 2,
+  "workload_id": "wan-t2v-1.3b",
+  "variant_id": "canonical",
+  "benchmark_version": 1
+}
+```
+
+`benchmark_id` is still required in this phase because raw artifact names,
+generated-video directories, normalized record paths, and the current rolling
+baseline comparator still depend on it. The v2 identity fields are config
+metadata that make the measured workload explicit:
+
+| Field | Purpose |
+|---|---|
+| `workload_id` | Stable benchmark family, such as `wan-t2v-1.3b`. |
+| `variant_id` | Intentional recipe family, such as `canonical`. |
+| `benchmark_version` | Version of the measurement protocol and comparison policy. |
+
+If a config declares `config_schema_version: 2`, loading fails clearly when any
+required v2 identity field is missing. If v2 identity or metadata fields are
+added without `config_schema_version: 2`, loading also fails so partial
+migrations do not silently run as v1 configs. Optional v2 metadata fields
+reserved for follow-up work, such as `recipe`, `metric_threshold_policy`, and
+`quality_metadata`, must be JSON objects when present.
+
+Recipe fingerprinting, hardware/software profile IDs, exact-identity
+comparison, metric-specific threshold policy behavior, promoted baselines, and
+dashboard regrouping are separate follow-up changes. Until those land, rolling
+baseline comparison remains keyed by `(model_id, gpu_type)`.
+
 ### Raw record (`results/perf_*.json`)
 
 Written by `test_inference_performance.py`. One file per benchmark run.
@@ -179,6 +218,10 @@ Written by `test_inference_performance.py`. One file per benchmark run.
 ```jsonc
 {
   "benchmark_id": "wan-t2v-1.3b-2gpu",
+  "config_schema_version": 2,
+  "workload_id": "wan-t2v-1.3b",
+  "variant_id": "canonical",
+  "benchmark_version": 1,
   "model_short_name": "Wan2.1-T2V-1.3B-Diffusers",
   "device": "NVIDIA L40S",
   "num_gpus": 2,
@@ -279,11 +322,16 @@ observability. When pytest passes, the rolling-baseline phase emits:
 ## Adding a new benchmark
 
 1. Drop a new JSON config into
-   `.buildkite/performance-benchmarks/tests/<name>.json`. Required keys:
+   `.buildkite/performance-benchmarks/tests/<name>.json`. New configs should
+   use v2 identity fields:
 
    ```json
    {
      "benchmark_id": "<unique-id>",
+     "config_schema_version": 2,
+     "workload_id": "<stable-workload-id>",
+     "variant_id": "canonical",
+     "benchmark_version": 1,
      "model": { "model_path": "...", "model_short_name": "..." },
      "init_kwargs": { "num_gpus": 1, ... },
      "generation_kwargs": { "num_frames": 45, ... },
@@ -303,6 +351,10 @@ observability. When pytest passes, the rolling-baseline phase emits:
    }
    ```
 
+   Legacy v1 configs without `config_schema_version` still load, but should not
+   gain v2 identity or metadata fields until they are migrated to
+   `config_schema_version: 2`.
+
 2. The pytest test auto-discovers all configs — no test code needed. CI
    picks it up on the next `/test performance` run.
 

docs/inference/optimizations.md

@@ -74,6 +74,23 @@ uv pip install ninja
 python setup.py install
 ```
 
+### Flash Attention 4 (opt-in)
+
+FastVideo never auto-selects FlashAttention-4 (`flash_attn.cute`) just because it
+is installed: its CuTeDSL kernels JIT-compile per shape family and can fail at
+runtime on some GPU/shape combinations. To use FA4, install the pinned
+`flash-attn-4` build (see the `flash-attn-4` source in `pyproject.toml`) and set:
+
+```bash
+export FASTVIDEO_FA4=1
+```
+
+On GPUs below sm90 a capability gate routes to FlashAttention-2 the calls FA4
+cannot serve there: grad-enabled (training) attention (FA4's backward requires
+sm90+) and GQA attention (FA4's `pack_gqa` fails to JIT-compile below sm90).
+On sm90+ both run on FA4. If FA4 is unusable while `FASTVIDEO_FA4=1` is set,
+FastVideo fails loudly instead of silently falling back.
+
 ### FP4 Flash Attention 4 (Blackwell only)
 
 **`FLASH_ATTN`** with **`--nvfp4_fa4`**

fastvideo-kernel/CMakeLists.txt

@@ -12,6 +12,9 @@ set(_FASTVIDEO_USER_CUDA_ARCH "${CMAKE_CUDA_ARCHITECTURES}")
 if(NOT DEFINED GPU_BACKEND AND DEFINED ENV{GPU_BACKEND})
     set(GPU_BACKEND "$ENV{GPU_BACKEND}")
 endif()
+if(NOT GPU_BACKEND)
+    set(GPU_BACKEND "CUDA")
+endif()
 
 if(GPU_BACKEND STREQUAL "ROCM")
     enable_language(HIP)
@@ -50,7 +53,16 @@ if(NOT GPU_BACKEND STREQUAL "ROCM")
         if(_FASTVIDEO_USER_CUDA_ARCH)
             # Caller pinned -DCMAKE_CUDA_ARCHITECTURES (which torch ignores); translate it
             # to the TORCH_CUDA_ARCH_LIST spelling: "121" -> "12.1", "90a" -> "9.0a".
+            # Only numeric spellings translate; keywords like "native"/"all" would
+            # otherwise be mangled into nonsense ("nativ.e").
             foreach(_fv_arch IN LISTS _FASTVIDEO_USER_CUDA_ARCH)
+                if(NOT _fv_arch MATCHES "^[0-9]+[af]?$")
+                    message(FATAL_ERROR
+                        "fastvideo-kernel: CMAKE_CUDA_ARCHITECTURES='${_fv_arch}' is not "
+                        "supported. Use a numeric arch (e.g. 90a, 121), set "
+                        "TORCH_CUDA_ARCH_LIST directly (e.g. 9.0a), or unset both to "
+                        "auto-detect from the visible GPU.")
+                endif()
                 string(REGEX MATCH "[af]$" _fv_suffix "${_fv_arch}")
                 string(REGEX REPLACE "[af]$" "" _fv_num "${_fv_arch}")
                 string(REGEX REPLACE "(.)$" ".\\1" _fv_num "${_fv_num}")    # dot before the last digit
@@ -173,6 +185,14 @@ else()
     endif()
 endif()
 
+# ThunderKittens headers don't compile on aarch64 hosts: plain char is unsigned
+# there, and tk's base_types.cuh brace-initializes signed-char vector members
+# from char (narrowing error). Skip TK until upstream is aarch64-clean.
+if(ENABLE_TK_KERNELS AND CMAKE_SYSTEM_PROCESSOR MATCHES "^(aarch64|arm64)$")
+    message(STATUS "ThunderKittens kernels: forced OFF on ${CMAKE_SYSTEM_PROCESSOR} (tk headers are not aarch64-clean)")
+    set(ENABLE_TK_KERNELS OFF)
+endif()
+
 if(ENABLE_TK_KERNELS)
     message(STATUS "ThunderKittens kernels: ENABLED")
 else()
@@ -263,11 +283,6 @@ set(CUDA_FLAGS
     "--expt-relaxed-constexpr"
     "-Xcompiler=-fno-strict-aliasing"
     "-Xcompiler=-fPIC"
-    # ARM/aarch64 defaults `char` to unsigned, but ThunderKittens headers assume the
-    # x86 signed-char behavior (else base_types.cuh hits "narrowing conversion from
-    # char to signed char"). Force signed char so TK compiles on Grace Hopper; this
-    # is a no-op on x86_64, where char is already signed.
-    "-Xcompiler=-fsigned-char"
     "-DTORCH_COMPILE"
     "-Xnvlink=--verbose"
     "-Xptxas=--verbose"
@@ -426,3 +441,14 @@ if(ENABLE_ATTN_QAT_INFER)
     install(TARGETS fp4attn_cuda LIBRARY DESTINATION .)
     install(TARGETS fp4quant_cuda LIBRARY DESTINATION .)
 endif()
+
+# One-look answer to "what is this build producing?" — kept last so it is the
+# final thing configure prints. The per-kernel matrix lives in README.md.
+message(STATUS "============== fastvideo-kernel build summary ==============")
+message(STATUS "host / backend:       ${CMAKE_SYSTEM_PROCESSOR} / ${GPU_BACKEND}")
+message(STATUS "TORCH_CUDA_ARCH_LIST: ${TORCH_CUDA_ARCH_LIST}")
+message(STATUS "fastvideo_kernel_ops: ON  (turbodiffusion int8-gemm/quant/rmsnorm/layernorm, all listed archs)")
+message(STATUS "  + TK sta/block_sparse (sm_90a only):          ${ENABLE_TK_KERNELS}")
+message(STATUS "fp4attn/fp4quant (sm_120a only, CUDA >= 12.8):  ${ENABLE_ATTN_QAT_INFER}")
+message(STATUS "Triton fallbacks ship in python/fastvideo_kernel/triton_kernels regardless.")
+message(STATUS "============================================================")

fastvideo-kernel/README.md

@@ -2,6 +2,42 @@
 
 CUDA kernels for FastVideo video generation.
 
+## Kernel inventory
+
+Compiled CUDA extensions (CMake, see the build summary printed at the end of every configure):
+
+| Extension | Kernels | Sources | GPU arch | Build gate |
+|---|---|---|---|---|
+| `fastvideo_kernel._C.fastvideo_kernel_ops` | TurboDiffusion INT8 GEMM, quant, RMSNorm, LayerNorm | `csrc/turbodiffusion/` | every arch in `TORCH_CUDA_ARCH_LIST` | always built |
+| same extension, optional part | ThunderKittens sliding-tile attention (`sta_fwd`) and VSA block-sparse (`block_sparse_fwd/bwd`) | `csrc/attention/*_h100.cu` | Hopper `sm_90a` only | `FASTVIDEO_KERNEL_BUILD_TK` (AUTO = ON iff `9.0a` is in the arch list; always OFF on aarch64 hosts — TK headers don't compile there) |
+| `fp4attn_cuda`, `fp4quant_cuda` | FP4 attention + quantization ("attn_qat_infer", modified SageAttention3) | `attn_qat_infer/` | consumer Blackwell `sm_120a` only, CUDA ≥ 12.8 | `FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER` (AUTO = ON iff `12.0a` is in the arch list) |
+
+Runtime-JIT kernels (no build step, ship in every wheel/image):
+
+| Kernels | Where | Used when |
+|---|---|---|
+| Triton: STA, VSA block-sparse, SLA, fused compress+topk, FP4 QAT training, quant/norm utils | `python/fastvideo_kernel/triton_kernels/` | automatic fallback when the matching C++ op is absent (`ops.py`, `turbodiffusion_ops.py`) |
+| FA4 CuTe-DSL block-sparse (VSA-256 fastpath on `sm_100`) | `block_sparse_attn_cute_fwd.py` | optional `flash_attn.cute` dependency, see below |
+| VMoBA `moba_attn_varlen` | `vmoba.py` | wraps flash-attn varlen |
+
+## What gets built where, and when
+
+| Surface | Trigger | Leg | `TORCH_CUDA_ARCH_LIST` | TK | FP4 |
+|---|---|---|---|---|---|
+| PyPI wheels (`.github/workflows/publish-kernel.yml`) | version bump in `fastvideo-kernel/pyproject.toml` on main, or manual dispatch | x86_64 cu126 | `9.0a` | ON | — (CUDA < 12.8) |
+| | | x86_64 cu130 | `9.0a;12.0a` | ON | ON |
+| | | aarch64 cu130 | `10.0a;12.0a` | — | ON |
+| Docker images `ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev` (`.github/workflows/infra-build-image.yml`) | `docker/Dockerfile` changes on main, or manual dispatch | amd64 cuda12.6.3 + cuda13.0.0 | `9.0a` | ON | — |
+| | | arm64 cuda12.6.3 (GH200) | `9.0a` | — (aarch64) | — |
+| | | arm64 cuda13.0.0 (GB10 / DGX Spark) | `12.1` | — | — |
+| Local `./build.sh` | manual | probes the visible GPU via torch | detected | ON iff sm_90 (non-aarch64 host) | ON iff sm_120 |
+
+Notes:
+
+- No Docker image ships the FP4 kernels; only the x86_64/aarch64 cu130 wheels do.
+- On arm64 images (GH200 included) STA/VSA run on the Triton fallbacks, since TK never builds on aarch64.
+- Kernel tests run on Buildkite GPU CI for PRs touching `fastvideo-kernel/**` (see `.buildkite/pipeline.yml`).
+
 ## Installation
 
 ### Standard Installation (Local Development)

fastvideo-kernel/build.sh

@@ -46,6 +46,23 @@ fi
 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

fastvideo/attention/backends/flash_attn.py

@@ -1,15 +1,39 @@
 # SPDX-License-Identifier: Apache-2.0
 
+import importlib.util
 import os
+
 import torch
 import torch.nn.functional as F
 from dataclasses import dataclass
 
-try:
-    from fastvideo.attention.utils.flash_attn_cute import flash_attn_func
+from fastvideo import envs
+from fastvideo.attention.backends.abstract import (
+    AttentionBackend,
+    AttentionImpl,
+    AttentionMetadata,
+    AttentionMetadataBuilder,
+)
+from fastvideo.logger import init_logger
 
+logger = init_logger(__name__)
+
+# FA4 (flash_attn.cute) is explicit opt-in via FASTVIDEO_FA4=1, mirroring the
+# kernel package's FASTVIDEO_VSA_CUTEDSL: its CuTeDSL kernels JIT-compile per
+# shape family and can fail at runtime on some arch/shape combinations, so it
+# is never auto-selected just because it is installed. Below sm90 a capability
+# gate in flash_attn_cute routes to FA2 the calls FA4 cannot serve there:
+# grad-enabled (its backward asserts sm90+) and GQA (pack_gqa fails CuTeDSL
+# JIT, observed on sm_89).
+if envs.FASTVIDEO_FA4:
+    try:
+        from fastvideo.attention.utils.flash_attn_cute import flash_attn_func
+    except ImportError as e:
+        raise RuntimeError(f"FASTVIDEO_FA4=1 but flash_attn.cute (FA4) is not usable ({e}); "
+                           "fix the FA4 install (see the flash-attn-4 pin in pyproject.toml) "
+                           "or unset FASTVIDEO_FA4.") from e
     fa_version = "4"
-except ImportError:
+else:
     try:
         from flash_attn_interface import flash_attn_func as flash_attn_3_func
 
@@ -21,6 +45,12 @@ except ImportError:
         from flash_attn import flash_attn_func as flash_attn_2_func
         flash_attn_func = flash_attn_2_func
         fa_version = "2"
+    try:
+        if importlib.util.find_spec("flash_attn.cute") is not None:
+            logger.info("flash_attn.cute (FA4) is installed but not enabled; "
+                        "set FASTVIDEO_FA4=1 to use it for inference.")
+    except ImportError:
+        pass
 
 # torch.compile traceability: the FA4/cute path (fa_version=="4") is
 # already a registered torch.library custom op, so dynamo treats it as a
@@ -87,9 +117,10 @@ if fa_version in ("2", "3"):
             return _fa_default(q, k, v, softmax_scale=softmax_scale, causal=causal)
         return torch.ops.fastvideo._flash_attn_default_forward(q, k, v, softmax_scale, causal)
 elif fa_version == "4":
-    # FA4 path: `flash_attn_func` is already a torch.library custom op
-    # (registered in `fastvideo.attention.utils.flash_attn_cute`), so a
-    # passthrough is enough — no extra registration needed.
+    # FA4 path: `flash_attn_func` (from `flash_attn_cute`) goes through a
+    # registered torch.library custom op (with an FA4 backward on sm90+;
+    # grad-enabled and GQA calls below sm90 route to FA2), so a passthrough
+    # is enough — no extra registration needed.
     def flash_attn_func_compilable(q, k, v, softmax_scale=None, causal=False):
         return flash_attn_func(q, k, v, softmax_scale=softmax_scale, causal=causal)
 else:
@@ -99,17 +130,6 @@ else:
     raise RuntimeError(f"Unsupported FlashAttention version: {fa_version!r} — expected "
                        f"'2', '3', or '4' from the import probe above.")
 
-from fastvideo.attention.backends.abstract import (
-    AttentionBackend,
-    AttentionImpl,
-    AttentionMetadata,
-    AttentionMetadataBuilder,
-)
-from fastvideo.logger import init_logger
-
-logger = init_logger(__name__)
-
-_WARNED_NON_FA_DTYPE = False
 logger.info("Using FlashAttention-%s backend", fa_version)
 
 # FP4 FA4 support: quantize Q/K to NVFP4 E2M1 for block-scaled MMA on Blackwell.
@@ -271,12 +291,8 @@ class FlashAttentionImpl(AttentionImpl):
         # SP). Cast through bf16 and restore, matching TORCH_SDPA's tolerance.
         orig_dtype = query.dtype
         if orig_dtype not in (torch.float16, torch.bfloat16):
-            global _WARNED_NON_FA_DTYPE
-            if not _WARNED_NON_FA_DTYPE:
-                _WARNED_NON_FA_DTYPE = True
-                logger.warning(
-                    "FLASH_ATTN received %s inputs; casting to bfloat16 for the "
-                    "kernel and restoring on output.", orig_dtype)
+            logger.warning_once(f"FLASH_ATTN received {orig_dtype} inputs; casting to "
+                                f"bfloat16 for the kernel and restoring on output.")
             query = query.to(torch.bfloat16)
             key = key.to(torch.bfloat16)
             value = value.to(torch.bfloat16)

fastvideo/attention/utils/flash_attn_cute.py

@@ -4,10 +4,9 @@ import functools
 from collections.abc import Callable
 
 import torch
-from flash_attn import flash_attn_func as _flash_attn_2_func
-from flash_attn import flash_attn_varlen_func as _flash_attn_2_varlen_func
 
 from fastvideo.logger import init_logger
+from fastvideo.platforms import current_platform
 
 logger = init_logger(__name__)
 
@@ -16,7 +15,8 @@ if torch.cuda.is_available():
         from flash_attn.cute.interface import _flash_attn_bwd, _flash_attn_fwd
     except ImportError:
         # flash_attn.cute (FA4) is simply not installed -- expected on builds
-        # without it; callers fall back to FA3/FA2 quietly.
+        # without it; callers handle the ImportError (the FASTVIDEO_FA4 gate in
+        # flash_attn.py raises, the FP4 probe treats FA4 as unavailable).
         raise
     except Exception as e:
         # flash_attn.cute IS installed but failed to import -- almost always an
@@ -24,23 +24,57 @@ if torch.cuda.is_available():
         # 'cutlass.cute.core' has no attribute 'ThrMma'" (an AttributeError, not
         # ImportError). This is fixable by pinning a compatible
         # nvidia-cutlass-dsl, so warn loudly, then re-raise as ImportError so
-        # callers fall back to FA3/FA2 instead of crashing worker init.
+        # callers can handle it uniformly.
         logger.warning(
-            "flash_attn.cute (FA4) is installed but failed to import (%r); "
-            "falling back to FA3/FA2. This is usually an nvidia-cutlass-dsl "
-            "version mismatch -- pin a compatible nvidia-cutlass-dsl to "
-            "restore FA4.", e)
+            "flash_attn.cute (FA4) is installed but failed to import (%r). "
+            "This is usually an nvidia-cutlass-dsl version mismatch -- pin a "
+            "compatible nvidia-cutlass-dsl to restore FA4.", e)
         raise ImportError(f"flash_attn.cute (FA4) import failed: {e!r}") from e
 else:
     # This error will be caught in flash_attn.py or flash_attn_no_pad.py
     raise ImportError("flash_attn.cute is only available on CUDA devices; this error must be handled internally")
 
+try:
+    # FA2 serves the calls FA4 cute cannot on pre-sm90 GPUs (backward, GQA).
+    # Optional so FA4-only installs can still import this module.
+    from flash_attn import flash_attn_func as _flash_attn_2_func
+    from flash_attn import flash_attn_varlen_func as _flash_attn_2_varlen_func
+except ImportError:
+    _flash_attn_2_func = None
+    _flash_attn_2_varlen_func = None
+
 
 def _check_dropout(dropout_p: float) -> None:
     if dropout_p != 0.0:
         raise NotImplementedError(f"flash_attn.cute does not support dropout (got dropout_p={dropout_p})")
 
 
+@functools.cache
+def _sm90_or_newer() -> bool:
+    return current_platform.has_device_capability(90)
+
+
+def _use_fa2(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor) -> bool:
+    if _sm90_or_newer():
+        return False
+    # Pre-sm90 FA4 cute limitations, both served by FA2 (deterministic
+    # capability gate, not a runtime fallback):
+    #   * the backward asserts sm90+ (L40S/sm_89 dies on its arch check);
+    #   * GQA fails CuTeDSL JIT in pack_gqa ("ValueError: Operation creation
+    #     failed", observed on sm_89 with HunyuanGameCraft/LTX2).
+    if q.shape[-2] != k.shape[-2]:
+        return True
+    return torch.is_grad_enabled() and any(t.requires_grad for t in (q, k, v))
+
+
+def _fa2_or_raise(fa2_func: Callable | None) -> Callable:
+    if fa2_func is None:
+        raise RuntimeError("this attention call cannot run on FA4 cute below sm90 (its backward and "
+                           "GQA support require sm90+) and flash-attn 2, which serves it there, is "
+                           "not installed.")
+    return fa2_func
+
+
 @torch.library.custom_op(
     "fastvideo::_flash_attn_cute_forward",
     mutates_args=(),
@@ -243,70 +277,6 @@ torch.library.register_autograd(
 )
 
 
-# FA4's CuTeDSL kernels JIT-compile per shape family, and some configurations
-# fail MLIR op creation at runtime even though the import succeeded (observed:
-# GQA models on sm_89 dying in pack_gqa with "ValueError: Operation creation
-# failed"). Degrade to FA2 once, process-wide, instead of crashing inference.
-class _FA4Policy:
-    """Per-call gate for the FA4 cute fast path, with FA2 as the fallback.
-
-    FA4 is skipped when:
-      * a previous call failed at runtime -- CuTeDSL JIT compilation is
-        shape-dependent, so the first failure disables FA4 for the rest of
-        the process instead of retrying a broken JIT on every call; or
-      * the call needs autograd -- FA4's backward asserts sm90+ (L40S/sm_89
-        dies on its arch check) and is unvalidated for training in this repo
-        (its lse is not even allocated through our inference-shaped custom
-        op), so training keeps the pre-FA4 behavior: FA2 on every device.
-    """
-
-    def __init__(self) -> None:
-        self.broken = False
-
-    def use_fa4(self, *tensors: torch.Tensor) -> bool:
-        if self.broken:
-            return False
-        return not (torch.is_grad_enabled() and any(t.requires_grad for t in tensors))
-
-    def mark_broken(self, error: Exception) -> None:
-        if not self.broken:
-            self.broken = True
-            logger.warning(
-                "flash_attn.cute (FA4) failed at runtime (%r); falling back "
-                "to FA2 for the rest of this process.", error)
-
-
-_FA4 = _FA4Policy()
-
-
-def _with_fa2_fallback(fa2_func: Callable) -> Callable:
-    """Pair an FA4 cute wrapper with its FA2 twin of the same signature.
-
-    The decorated body runs only when ``_FA4`` allows it; otherwise (or after
-    the first FA4 runtime failure) the call is served by ``fa2_func``.
-    ``NotImplementedError`` is a contract error (e.g. dropout), not a JIT
-    failure, so it propagates without disabling FA4.
-    """
-
-    def decorator(fa4_func: Callable) -> Callable:
-
-        @functools.wraps(fa4_func)
-        def wrapper(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, *args, **kwargs) -> torch.Tensor:
-            if _FA4.use_fa4(q, k, v):
-                try:
-                    return fa4_func(q, k, v, *args, **kwargs)
-                except NotImplementedError:
-                    raise
-                except Exception as e:  # CuTeDSL compile errors surface as ValueError
-                    _FA4.mark_broken(e)
-            return fa2_func(q, k, v, *args, **kwargs)
-
-        return wrapper
-
-    return decorator
-
-
-@_with_fa2_fallback(_flash_attn_2_func)
 def flash_attn_func(
     q: torch.Tensor,
     k: torch.Tensor,
@@ -317,6 +287,16 @@ def flash_attn_func(
     deterministic: bool = False,
 ) -> torch.Tensor:
     """Only returns the output, not the lse."""
+    if _use_fa2(q, k, v):
+        return _fa2_or_raise(_flash_attn_2_func)(
+            q,
+            k,
+            v,
+            dropout_p=dropout_p,
+            softmax_scale=softmax_scale,
+            causal=causal,
+            deterministic=deterministic,
+        )
     _check_dropout(dropout_p)
     out, _ = torch.ops.fastvideo._flash_attn_cute_forward(q, k, v, softmax_scale, causal, deterministic)
     return out
@@ -392,7 +372,6 @@ def flash_attn_fp4_func(
     return torch.ops.fastvideo._flash_attn_cute_fp4_forward(q, k, v, sfq, sfk, softmax_scale, causal)
 
 
-@_with_fa2_fallback(_flash_attn_2_varlen_func)
 def flash_attn_varlen_func(
     q: torch.Tensor,
     k: torch.Tensor,
@@ -407,6 +386,20 @@ def flash_attn_varlen_func(
     deterministic: bool = False,
 ) -> torch.Tensor:
     """Only returns the output, not the lse."""
+    if _use_fa2(q, k, v):
+        return _fa2_or_raise(_flash_attn_2_varlen_func)(
+            q,
+            k,
+            v,
+            cu_seqlens_q,
+            cu_seqlens_k,
+            max_seqlen_q,
+            max_seqlen_k,
+            dropout_p=dropout_p,
+            softmax_scale=softmax_scale,
+            causal=causal,
+            deterministic=deterministic,
+        )
     _check_dropout(dropout_p)
     out, _ = torch.ops.fastvideo._flash_attn_cute_varlen_forward(
         q,

fastvideo/attention/utils/flash_attn_no_pad.py

@@ -21,24 +21,35 @@ from einops import rearrange
 from flash_attn import flash_attn_varlen_qkvpacked_func
 from flash_attn.bert_padding import pad_input, unpad_input
 
+from fastvideo import envs
+
 
 def _resolve_flash_attn_varlen_func() -> Any:
-    try:
-        from fastvideo.attention.utils.flash_attn_cute import (
-            flash_attn_varlen_func as flash_attn_varlen_func_cute, )
+    if envs.FASTVIDEO_FA4:
+        # FA4 cute is explicit opt-in (see fastvideo/attention/backends/
+        # flash_attn.py); with FASTVIDEO_FA4=1 an unimportable FA4 build must
+        # fail loudly here rather than fall through to FA3/FA2. RuntimeError,
+        # not ImportError: importers like bsa_attn.py treat ImportError as
+        # "flash-attn not installed" and silently degrade to reference kernels.
+        try:
+            from fastvideo.attention.utils.flash_attn_cute import (
+                flash_attn_varlen_func as flash_attn_varlen_func_cute, )
+        except ImportError as e:
+            raise RuntimeError(f"FASTVIDEO_FA4=1 but flash_attn.cute (FA4) is not usable ({e}); "
+                               "fix the FA4 install (see the flash-attn-4 pin in pyproject.toml) "
+                               "or unset FASTVIDEO_FA4.") from e
 
         return flash_attn_varlen_func_cute
+    try:
+        from flash_attn_interface import (
+            flash_attn_varlen_func as flash_attn_varlen_func_interface, )
+
+        return flash_attn_varlen_func_interface
     except ImportError:
-        try:
-            from flash_attn_interface import (
-                flash_attn_varlen_func as flash_attn_varlen_func_interface, )
+        from flash_attn import (
+            flash_attn_varlen_func as flash_attn_varlen_func_flash, )
 
-            return flash_attn_varlen_func_interface
-        except ImportError:
-            from flash_attn import (
-                flash_attn_varlen_func as flash_attn_varlen_func_flash, )
-
-            return flash_attn_varlen_func_flash
+        return flash_attn_varlen_func_flash
 
 
 flash_attn_varlen_func_impl = _resolve_flash_attn_varlen_func()

fastvideo/envs.py

@@ -20,6 +20,7 @@ if TYPE_CHECKING:
     FASTVIDEO_LOGGING_CONFIG_PATH: str | None = None
     FASTVIDEO_TRACE_FUNCTION: int = 0
     FASTVIDEO_ATTENTION_BACKEND: str | None = None
+    FASTVIDEO_FA4: bool = False
     FASTVIDEO_WORKER_MULTIPROC_METHOD: str = "spawn"
     FASTVIDEO_TARGET_DEVICE: str = "cuda"
     MAX_JOBS: str | None = None
@@ -207,9 +208,18 @@ environment_variables: dict[str, Callable[[], Any]] = {
     # - "VIDEO_SPARSE_ATTN": use Video Sparse Attention
     # - "SAGE_ATTN": use Sage Attention
     # - "SAGE_ATTN_THREE": use Sage Attention 3
+    # FLASH_ATTN uses FlashAttention-3/2; to run FlashAttention-4 set
+    # FASTVIDEO_FA4=1 as well (see below).
     "FASTVIDEO_ATTENTION_BACKEND":
     lambda: os.getenv("FASTVIDEO_ATTENTION_BACKEND", None),
 
+    # If set (=1), the FLASH_ATTN backend uses FlashAttention-4
+    # (flash_attn.cute). FA4 is opt-in and never auto-selected just because it
+    # is installed. Below sm90, grad-enabled and GQA calls are routed to FA2
+    # (FA4's backward asserts sm90+ and its pack_gqa fails to JIT there).
+    "FASTVIDEO_FA4":
+    lambda: os.getenv("FASTVIDEO_FA4", "0") != "0",
+
     # Use dedicated multiprocess context for workers.
     "FASTVIDEO_WORKER_MULTIPROC_METHOD":
     lambda: os.getenv("FASTVIDEO_WORKER_MULTIPROC_METHOD", "spawn"),

fastvideo/models/dits/causal_wanvideo.py

@@ -323,9 +323,9 @@ class CausalWanTransformerBlock(nn.Module):
         value, _ = self.to_v(norm_hidden_states)
 
         if self.norm_q is not None:
-            query = self.norm_q.forward_native(query)
+            query = self.norm_q(query)
         if self.norm_k is not None:
-            key = self.norm_k.forward_native(key)
+            key = self.norm_k(key)
 
         query = query.squeeze(1).unflatten(2, (self.num_attention_heads, -1))
         key = key.squeeze(1).unflatten(2, (self.num_attention_heads, -1))

fastvideo/models/dits/matrixgame2/causal_model.py

@@ -537,9 +537,9 @@ class CausalMatrixGame2TransformerBlock(nn.Module):
         value, _ = self.to_v(norm_hidden_states)
 
         if self.norm_q is not None:
-            query = self.norm_q.forward_native(query)
+            query = self.norm_q(query)
         if self.norm_k is not None:
-            key = self.norm_k.forward_native(key)
+            key = self.norm_k(key)
 
         query = query.squeeze(1).unflatten(2, (self.num_attention_heads, -1))
         key = key.squeeze(1).unflatten(2, (self.num_attention_heads, -1))

fastvideo/tests/attention/test_flash_attn_no_pad_resolver.py

@@ -3,15 +3,16 @@
 
 Landed in PR #1225 slice 5 (Attn-QAT 5/12). The resolver centralises the
 varlen-flash-attn import-fallback logic that several backends
-(``bsa_attn.py``, ``video_sparse_attn.py``) used to duplicate. The fallback
+(``bsa_attn.py``, ``video_sparse_attn.py``) used to duplicate. The resolution
 order is:
 
-    1. ``fastvideo.attention.utils.flash_attn_cute``
+    1. ``fastvideo.attention.utils.flash_attn_cute`` -- only when
+       ``FASTVIDEO_FA4=1`` (explicit opt-in), and then it must import or the
+       resolver raises RuntimeError instead of falling through
     2. ``flash_attn_interface``
     3. ``flash_attn``
 
-These tests verify that the resolver picks the highest-priority impl
-available and falls through cleanly on ``ImportError``. CPU-only, no
+These tests verify the opt-in gate and the FA3/FA2 fallthrough. CPU-only, no
 flash-attn install required.
 """
 
@@ -35,8 +36,36 @@ def _reload_resolver_module():
     return importlib.import_module("fastvideo.attention.utils.flash_attn_no_pad")
 
 
-def test_resolver_falls_back_when_cute_unavailable(monkeypatch) -> None:
-    """When ``flash_attn_cute`` is unimportable, resolver tries the next impl."""
+def test_resolver_skips_cute_without_opt_in(monkeypatch) -> None:
+    """Without ``FASTVIDEO_FA4=1`` the resolver must not even attempt the cute
+    import."""
+    monkeypatch.delenv("FASTVIDEO_FA4", raising=False)
+    attempted: list[str] = []
+    real_import = builtins.__import__
+
+    def spying_import(name, globals=None, locals=None, fromlist=(), level=0):
+        attempted.append(name)
+        return real_import(name, globals, locals, fromlist, level)
+
+    monkeypatch.setattr(builtins, "__import__", spying_import)
+
+    mod = _reload_resolver_module()
+    resolved = mod._resolve_flash_attn_varlen_func()
+    assert resolved is not None
+    assert resolved.__name__ == "flash_attn_varlen_func"
+    assert "fastvideo.attention.utils.flash_attn_cute" not in attempted
+
+
+def test_resolver_raises_when_opted_in_but_cute_unavailable(monkeypatch) -> None:
+    """With ``FASTVIDEO_FA4=1`` an unimportable cute build fails loudly instead
+    of silently falling through to FA3/FA2.
+
+    The resolver runs at module import time, so the reload itself must raise.
+    It raises RuntimeError (not ImportError) so importers that treat
+    ImportError as "flash-attn not installed" (``bsa_attn.py``) cannot swallow
+    the opted-in failure.
+    """
+    monkeypatch.setenv("FASTVIDEO_FA4", "1")
     real_import = builtins.__import__
 
     def patched_import(name, globals=None, locals=None, fromlist=(), level=0):
@@ -46,21 +75,17 @@ def test_resolver_falls_back_when_cute_unavailable(monkeypatch) -> None:
 
     monkeypatch.setattr(builtins, "__import__", patched_import)
 
-    mod = _reload_resolver_module()
-    resolved = mod._resolve_flash_attn_varlen_func()
-    assert resolved is not None
-    assert resolved.__name__ == "flash_attn_varlen_func"
+    with pytest.raises(RuntimeError, match="cute disabled for test"):
+        _reload_resolver_module()
 
 
-def test_resolver_returns_flash_attn_when_cute_and_interface_unavailable(monkeypatch) -> None:
+def test_resolver_returns_flash_attn_when_interface_unavailable(monkeypatch) -> None:
     """The terminal fallback is the plain ``flash_attn`` import."""
+    monkeypatch.delenv("FASTVIDEO_FA4", raising=False)
     real_import = builtins.__import__
 
     def patched_import(name, globals=None, locals=None, fromlist=(), level=0):
-        if name in {
-                "fastvideo.attention.utils.flash_attn_cute",
-                "flash_attn_interface",
-        }:
+        if name == "flash_attn_interface":
             raise ImportError(f"{name} disabled for test")
         return real_import(name, globals, locals, fromlist, level)
 

fastvideo/tests/layers/test_rmsnorm_forward_dispatch.py

@@ -0,0 +1,203 @@
+# SPDX-License-Identifier: Apache-2.0
+
+from __future__ import annotations
+
+import argparse
+import json
+import os
+import subprocess
+from pathlib import Path
+from typing import Any
+
+import pytest
+import torch
+import torch.distributed as dist
+from torch.distributed import init_device_mesh
+from torch.distributed.fsdp import CPUOffloadPolicy, fully_shard
+from torch.distributed.tensor import DTensor
+
+from fastvideo.layers.layernorm import RMSNorm
+
+WORLD_SIZE = 2
+HIDDEN_SIZE = 8
+SEED = 1379
+REPO_ROOT = Path(__file__).resolve().parents[3]
+
+
+def _run_torchrun(script_path: Path, mode: str, output_path: Path) -> None:
+    # --standalone binds the rendezvous port atomically, avoiding the
+    # free-port-probe race a hand-picked --master_port would have.
+    cmd = [
+        "torchrun",
+        "--standalone",
+        "--nproc_per_node",
+        str(WORLD_SIZE),
+        str(script_path),
+        "--rmsnorm-fsdp-worker",
+        "--mode",
+        mode,
+        "--output",
+        str(output_path),
+    ]
+    env = os.environ.copy()
+    env.setdefault("TORCHDYNAMO_DISABLE", "1")
+    try:
+        process = subprocess.run(
+            cmd,
+            capture_output=True,
+            text=True,
+            env=env,
+            timeout=120,
+        )
+    except subprocess.TimeoutExpired as error:
+        raise RuntimeError(
+            f"{mode} worker timed out after 120 seconds\n"
+            f"STDOUT:\n{error.stdout}\n"
+            f"STDERR:\n{error.stderr}"
+        ) from error
+    if process.returncode != 0:
+        raise RuntimeError(
+            f"{mode} worker failed with code {process.returncode}\n"
+            f"STDOUT:\n{process.stdout}\n"
+            f"STDERR:\n{process.stderr}"
+        )
+
+
+def _summarize_tensor(tensor: torch.Tensor | Any) -> dict[str, Any]:
+    return {
+        "type": type(tensor).__name__,
+        "is_dtensor": isinstance(tensor, DTensor),
+        "shape": list(tensor.shape) if hasattr(tensor, "shape") else None,
+        "device": str(tensor.device) if hasattr(tensor, "device") else None,
+        "dtype": str(tensor.dtype) if hasattr(tensor, "dtype") else None,
+    }
+
+
+def _run_worker(mode: str, output_path: Path) -> None:
+    if mode not in {
+        "module_no_offload",
+        "direct_no_offload",
+        "module_cpu_offload",
+        "direct_cpu_offload",
+    }:
+        raise ValueError(f"Unsupported mode: {mode}")
+
+    dist.init_process_group("nccl")
+    rank = dist.get_rank()
+    world_size = dist.get_world_size()
+    local_rank = int(os.environ.get("LOCAL_RANK", "0"))
+    device = torch.device(f"cuda:{local_rank}")
+    torch.cuda.set_device(device)
+    torch.manual_seed(SEED + rank)
+
+    try:
+        mesh = init_device_mesh("cuda", (world_size,))
+        norm = RMSNorm(HIDDEN_SIZE, eps=1e-6, has_weight=True).to(device)
+        with torch.no_grad():
+            norm.weight.fill_(1.0)
+
+        fsdp_kwargs: dict[str, Any] = {"mesh": mesh}
+        if mode.endswith("cpu_offload"):
+            fsdp_kwargs["offload_policy"] = CPUOffloadPolicy(pin_memory=False)
+        # fully_shard is applied to the bare RMSNorm to make the hook bypass
+        # observable. Production sharding (fsdp_load.shard_model) only wraps
+        # whole transformer blocks, whose pre-forward all-gather localizes norm
+        # weights before the qk-norm call sites run, so this pins the dispatch
+        # invariant rather than reproducing a production topology.
+        fully_shard(norm, **fsdp_kwargs)
+
+        x = torch.randn(2, 3, HIDDEN_SIZE, device=device, dtype=torch.bfloat16)
+        call_kind = "direct" if mode.startswith("direct") else "module"
+
+        try:
+            if call_kind == "direct":
+                output = norm.forward_native(x)
+            else:
+                output = norm(x)
+            torch.cuda.synchronize(device)
+            result = {
+                "rank": rank,
+                "ok": True,
+                "mode": mode,
+                "weight": _summarize_tensor(norm.weight),
+                "output": _summarize_tensor(output),
+            }
+        except Exception as exc:
+            result = {
+                "rank": rank,
+                "ok": False,
+                "mode": mode,
+                "error_type": type(exc).__name__,
+                "error": str(exc),
+                "weight": _summarize_tensor(norm.weight),
+            }
+
+        gathered = [None for _ in range(world_size)] if rank == 0 else None
+        dist.gather_object(result, object_gather_list=gathered, dst=0)
+        if rank == 0:
+            output_path.write_text(json.dumps(gathered, indent=2), encoding="utf-8")
+        dist.barrier()
+    finally:
+        dist.destroy_process_group()
+
+
+@pytest.mark.parametrize(
+    ("mode", "expect_ok"),
+    [
+        ("module_no_offload", True),
+        ("direct_no_offload", False),
+        ("module_cpu_offload", True),
+        ("direct_cpu_offload", False),
+    ],
+)
+def test_rmsnorm_forward_native_bypasses_fsdp_hooks(mode: str, expect_ok: bool, tmp_path: Path) -> None:
+    if not torch.cuda.is_available():
+        pytest.skip("This test requires CUDA.")
+    if torch.cuda.device_count() < WORLD_SIZE:
+        pytest.skip(f"This test requires at least {WORLD_SIZE} CUDA devices.")
+
+    output_path = tmp_path / f"{mode}.json"
+    _run_torchrun(Path(__file__).resolve(), mode, output_path)
+    results = json.loads(output_path.read_text(encoding="utf-8"))
+    print(f"\n{mode} results:\n{json.dumps(results, indent=2)}")
+
+    if expect_ok:
+        failures = [result for result in results if not result["ok"]]
+        assert not failures, json.dumps(results, indent=2)
+        return
+
+    successes = [result for result in results if result["ok"]]
+    assert not successes, json.dumps(results, indent=2)
+    error_text = "\n".join(result.get("error", "") for result in results)
+    # Pin the specific bypassed-hook failure: "got mixed torch.Tensor and
+    # DTensor" ("Tensor" alone is a substring of "DTensor", so it adds nothing).
+    assert "mixed" in error_text and "DTensor" in error_text, json.dumps(results, indent=2)
+
+
+def test_no_direct_forward_native_calls_in_models() -> None:
+    """Direct .forward_native(...) calls bypass nn.Module.__call__ and FSDP
+    hooks (issue #1379); model code must use module dispatch instead."""
+    models_dir = REPO_ROOT / "fastvideo" / "models"
+    offenders = [
+        str(path.relative_to(REPO_ROOT))
+        for path in sorted(models_dir.rglob("*.py"))
+        if ".forward_native(" in path.read_text(encoding="utf-8")
+    ]
+    assert not offenders, f"Replace .forward_native(...) with module dispatch in: {offenders}"
+
+
+def _parse_args() -> argparse.Namespace:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--rmsnorm-fsdp-worker", action="store_true")
+    parser.add_argument("--mode", type=str, default=None)
+    parser.add_argument("--output", type=str, default=None)
+    return parser.parse_args()
+
+
+if __name__ == "__main__":
+    args = _parse_args()
+    if not args.rmsnorm_fsdp_worker:
+        raise SystemExit("This module is intended to be run by pytest.")
+    if args.mode is None or args.output is None:
+        raise SystemExit("--mode and --output are required in worker mode.")
+    _run_worker(mode=args.mode, output_path=Path(args.output))

fastvideo/tests/modal/launch_l40s_job.py

@@ -32,7 +32,8 @@ import modal
 
 sys.path.insert(0, os.path.dirname(os.path.abspath(__file__)))
 try:
-    from modal_image_utils import resolve_image_ref  # noqa: E402
+    from modal_image_utils import (  # noqa: E402
+        resolve_image_ref, resolve_uv_torch_backend)
 except ModuleNotFoundError:
     # Remote Modal containers re-import this module but mount only the
     # entrypoint file; the digest resolution already happened at local
@@ -40,6 +41,9 @@ except ModuleNotFoundError:
     def resolve_image_ref(image_ref: str) -> str:
         return image_ref
 
+    def resolve_uv_torch_backend(image_tag: str) -> str | None:
+        return os.environ.get("UV_TORCH_BACKEND")
+
 app = modal.App("fastvideo-gpu-job")
 
 REPO_DIR = "/FastVideo"
@@ -72,12 +76,7 @@ local_secrets = modal.Secret.from_dict({
 # Mutable tags inherit the registry image's baked backend, including custom
 # FASTVIDEO_MODAL_IMAGE overrides. Explicit CUDA tags also work with older
 # images that predate the baked setting, and a caller override always wins.
-uv_torch_backend_override = os.environ.get("UV_TORCH_BACKEND")
-if not uv_torch_backend_override:
-    if "cuda13" in IMAGE_TAG.lower():
-        uv_torch_backend_override = "cu130"
-    elif "cuda12.6" in IMAGE_TAG.lower():
-        uv_torch_backend_override = "cu126"
+uv_torch_backend_override = resolve_uv_torch_backend(IMAGE_TAG)
 
 image = (
     modal.Image.from_registry(IMAGE_REF, add_python="3.12")
@@ -98,6 +97,9 @@ image = (
         "TOKENIZERS_PARALLELISM": "false",
         **({"UV_TORCH_BACKEND": uv_torch_backend_override} if uv_torch_backend_override else {}),
         "FASTVIDEO_ATTENTION_BACKEND": os.environ.get("FASTVIDEO_ATTENTION_BACKEND", "FLASH_ATTN"),
+        # FA4 is opt-in (FASTVIDEO_FA4); keep CI parity with the seeded
+        # references. Caller override wins.
+        "FASTVIDEO_FA4": os.environ.get("FASTVIDEO_FA4", "1"),
     })
 )
 

fastvideo/tests/modal/modal_image_utils.py

@@ -5,6 +5,7 @@ through the ``fastvideo`` package, so they keep working on thin CI hosts that
 have ``modal`` but not torch.
 """
 import json
+import os
 import urllib.request
 
 _REGISTRY = "ghcr.io"
@@ -55,3 +56,23 @@ def resolve_image_ref(image_ref: str) -> str:
         print(f"WARNING: could not resolve {image_ref} to a digest ({error}); "
               "Modal may reuse a stale cached image for this tag.")
     return image_ref
+
+
+def resolve_uv_torch_backend(image_tag: str) -> str | None:
+    """UV_TORCH_BACKEND for a launcher image tag.
+
+    A caller-set UV_TORCH_BACKEND always wins. Otherwise sniff the CUDA
+    version from an explicit image tag (cuda13 -> cu130, cuda12.6 -> cu126)
+    so uv resolves torch against the image's toolkit. Mutable tags (e.g.
+    py3.12-latest) return None and inherit the registry image's baked
+    backend, which keeps a latest-tag CUDA transition safe.
+    """
+    override = os.environ.get("UV_TORCH_BACKEND")
+    if override:
+        return override
+    tag = image_tag.lower()
+    if "cuda13" in tag:
+        return "cu130"
+    if "cuda12.6" in tag:
+        return "cu126"
+    return None

fastvideo/tests/modal/pr_test.py

@@ -5,7 +5,8 @@ import modal
 
 sys.path.insert(0, os.path.dirname(os.path.abspath(__file__)))
 try:
-    from modal_image_utils import resolve_image_ref  # noqa: E402
+    from modal_image_utils import (  # noqa: E402
+        resolve_image_ref, resolve_uv_torch_backend)
 except ModuleNotFoundError:
     # Remote Modal containers re-import this module but mount only the
     # entrypoint file; the digest resolution already happened at local
@@ -13,6 +14,9 @@ except ModuleNotFoundError:
     def resolve_image_ref(image_ref: str) -> str:
         return image_ref
 
+    def resolve_uv_torch_backend(image_tag: str) -> str | None:
+        return os.environ.get("UV_TORCH_BACKEND")
+
 app = modal.App()
 
 model_vol = modal.Volume.from_name("hf-model-weights")
@@ -24,12 +28,7 @@ print(f"Using image: {image_ref}")
 # Mutable tags inherit the registry image's baked backend, keeping a latest-tag
 # transition safe. Explicit CUDA tags also work with older images that predate
 # the baked setting, and a caller override always wins.
-uv_torch_backend_override = os.environ.get("UV_TORCH_BACKEND")
-if not uv_torch_backend_override:
-    if "cuda13" in image_tag.lower():
-        uv_torch_backend_override = "cu130"
-    elif "cuda12.6" in image_tag.lower():
-        uv_torch_backend_override = "cu126"
+uv_torch_backend_override = resolve_uv_torch_backend(image_tag)
 
 image = (modal.Image.from_registry(
     image_ref, add_python="3.12"
@@ -61,6 +60,10 @@ image = (modal.Image.from_registry(
     **({
         "UV_TORCH_BACKEND": uv_torch_backend_override
     } if uv_torch_backend_override else {}),
+    # FA4 is opt-in (FASTVIDEO_FA4); CI lanes keep it enabled to match the
+    # SSIM/perf baselines. Caller override wins.
+    "FASTVIDEO_FA4":
+    os.environ.get("FASTVIDEO_FA4", "1"),
     "HF_REPO_ID":
     "FastVideo/performance-tracking",
 }))
@@ -273,7 +276,7 @@ def run_self_forcing_tests():
 @app.function(gpu="L40S:1", image=image, timeout=900)
 def run_unit_test():
     run_test(
-        "pytest ./fastvideo/tests/api/ ./fastvideo/tests/contract/ ./fastvideo/tests/dataset/ ./fastvideo/tests/workflow/ ./fastvideo/tests/entrypoints/ ./fastvideo/tests/train/ --ignore=./fastvideo/tests/entrypoints/test_openai_api_integration.py --ignore=./fastvideo/tests/train/models --ignore=./fastvideo/tests/train/methods -vs"
+        "pytest ./fastvideo/tests/api/ ./fastvideo/tests/contract/ ./fastvideo/tests/dataset/ ./fastvideo/tests/workflow/ ./fastvideo/tests/entrypoints/ ./fastvideo/tests/train/ ./fastvideo/tests/stages/ --ignore=./fastvideo/tests/entrypoints/test_openai_api_integration.py --ignore=./fastvideo/tests/train/models --ignore=./fastvideo/tests/train/methods -vs"
     )
 
 

fastvideo/tests/modal/ssim_test.py

@@ -13,7 +13,8 @@ import modal
 
 sys.path.insert(0, os.path.dirname(os.path.abspath(__file__)))
 try:
-    from modal_image_utils import resolve_image_ref  # noqa: E402
+    from modal_image_utils import (  # noqa: E402
+        resolve_image_ref, resolve_uv_torch_backend)
 except ModuleNotFoundError:
     # Remote Modal containers re-import this module but mount only the
     # entrypoint file; the digest resolution already happened at local
@@ -21,6 +22,9 @@ except ModuleNotFoundError:
     def resolve_image_ref(image_ref: str) -> str:
         return image_ref
 
+    def resolve_uv_torch_backend(image_tag: str) -> str | None:
+        return os.environ.get("UV_TORCH_BACKEND")
+
 app = modal.App()
 
 model_vol = modal.Volume.from_name("hf-model-weights")
@@ -32,12 +36,7 @@ print(f"Using image: {image_ref}")
 # Mutable tags inherit the registry image's baked backend, keeping a latest-tag
 # transition safe. Explicit CUDA tags also work with older images that predate
 # the baked setting, and a caller override always wins.
-uv_torch_backend_override = os.environ.get("UV_TORCH_BACKEND")
-if not uv_torch_backend_override:
-    if "cuda13" in image_tag.lower():
-        uv_torch_backend_override = "cu130"
-    elif "cuda12.6" in image_tag.lower():
-        uv_torch_backend_override = "cu126"
+uv_torch_backend_override = resolve_uv_torch_backend(image_tag)
 
 image = (
     modal.Image.from_registry(image_ref, add_python="3.12")
@@ -64,6 +63,9 @@ image = (
             "BUILDKITE_PULL_REQUEST": os.environ.get("BUILDKITE_PULL_REQUEST", ""),
             "IMAGE_VERSION": image_version,
             **({"UV_TORCH_BACKEND": uv_torch_backend_override} if uv_torch_backend_override else {}),
+            # FA4 is opt-in (FASTVIDEO_FA4); the SSIM references were seeded
+            # with FA4 inference, so keep it enabled in CI. Caller override wins.
+            "FASTVIDEO_FA4": os.environ.get("FASTVIDEO_FA4", "1"),
         }
     )
 )

fastvideo/tests/performance/test_benchmark_config_validation.py

@@ -0,0 +1,116 @@
+# SPDX-License-Identifier: Apache-2.0
+
+import pytest
+
+from fastvideo.tests.performance.test_inference_performance import (
+    _benchmark_display_id,
+    _config_identity_metadata,
+    _is_v2_config,
+    _validate_benchmark_config,
+)
+
+
+def _v2_config():
+    return {
+        "benchmark_id": "wan-t2v-1.3b-2gpu",
+        "config_schema_version": 2,
+        "workload_id": "wan-t2v-1.3b",
+        "variant_id": "canonical",
+        "benchmark_version": 1,
+    }
+
+
+def test_v1_benchmark_config_without_schema_version_validates():
+    cfg = {
+        "benchmark_id": "legacy-benchmark",
+    }
+
+    _validate_benchmark_config(cfg, "legacy.json")
+
+    assert _is_v2_config(cfg) is False
+    assert _config_identity_metadata(cfg) == {}
+    assert _benchmark_display_id(cfg) == "legacy-benchmark"
+
+
+def test_v2_benchmark_config_identity_validates_and_is_preserved():
+    cfg = _v2_config()
+    cfg["quality_metadata"] = {"some": "data"}
+
+    _validate_benchmark_config(cfg, "wan.json")
+
+    assert _is_v2_config(cfg) is True
+    assert _config_identity_metadata(cfg) == {
+        "config_schema_version": 2,
+        "workload_id": "wan-t2v-1.3b",
+        "variant_id": "canonical",
+        "benchmark_version": 1,
+        "quality_metadata": {"some": "data"},
+    }
+
+
+def test_v2_benchmark_config_missing_identity_fields_fails_clearly():
+    cfg = _v2_config()
+    del cfg["variant_id"]
+    del cfg["benchmark_version"]
+
+    expected = "wan.json: missing required v2 identity fields: variant_id, benchmark_version"
+    with pytest.raises(ValueError, match=expected):
+        _validate_benchmark_config(cfg, "wan.json")
+
+
+@pytest.mark.parametrize(
+    ("field", "value"),
+    [
+        ("workload_id", {}),
+        ("workload_id", ""),
+        ("workload_id", "   "),
+        ("variant_id", []),
+        ("variant_id", ""),
+        ("variant_id", "   "),
+    ],
+)
+def test_v2_benchmark_config_rejects_invalid_string_identity_values(field, value):
+    cfg = _v2_config()
+    cfg[field] = value
+
+    expected = f"wan.json: v2 identity field {field!r} must be a non-empty string"
+    with pytest.raises(ValueError, match=expected):
+        _validate_benchmark_config(cfg, "wan.json")
+
+
+@pytest.mark.parametrize("value", [None, "1", 1.5, True])
+def test_v2_benchmark_config_rejects_invalid_benchmark_version_values(value):
+    cfg = _v2_config()
+    cfg["benchmark_version"] = value
+
+    expected = "wan.json: v2 identity field 'benchmark_version' must be an integer"
+    with pytest.raises(ValueError, match=expected):
+        _validate_benchmark_config(cfg, "wan.json")
+
+
+def test_partial_v2_identity_requires_schema_version():
+    cfg = {
+        "benchmark_id": "wan-t2v-1.3b-2gpu",
+        "workload_id": "wan-t2v-1.3b",
+        "variant_id": "canonical",
+        "benchmark_version": 1,
+    }
+
+    expected = "wan.json: v2 benchmark identity fields require config_schema_version=2"
+    with pytest.raises(ValueError, match=expected):
+        _validate_benchmark_config(cfg, "wan.json")
+
+
+def test_optional_v2_metadata_fields_must_be_objects():
+    cfg = {
+        "benchmark_id": "wan-t2v-1.3b-2gpu",
+        "config_schema_version": 2,
+        "workload_id": "wan-t2v-1.3b",
+        "variant_id": "canonical",
+        "benchmark_version": 1,
+        "quality_metadata": ["not", "an", "object"],
+    }
+
+    expected = "wan.json: optional v2 metadata field 'quality_metadata' must be an object"
+    with pytest.raises(ValueError, match=expected):
+        _validate_benchmark_config(cfg, "wan.json")

fastvideo/tests/performance/test_inference_performance.py

@@ -28,6 +28,17 @@ STAGE_METRIC_MAP: dict[str, str] = {
     "DmdDenoisingStage": "dit_time_s",
     "DecodingStage": "vae_decode_time_s",
 }
+V2_CONFIG_SCHEMA_VERSION = 2
+V2_REQUIRED_IDENTITY_FIELDS = (
+    "workload_id",
+    "variant_id",
+    "benchmark_version",
+)
+V2_OPTIONAL_METADATA_FIELDS = (
+    "recipe",
+    "metric_threshold_policy",
+    "quality_metadata",
+)
 
 # -- Config discovery -------------------------------------------------------
 
@@ -42,6 +53,71 @@ _BENCHMARKS_DIR = os.path.join(
 )
 
 
+def _has_v2_fields(cfg):
+    v2_fields = V2_REQUIRED_IDENTITY_FIELDS + V2_OPTIONAL_METADATA_FIELDS
+    return any(field in cfg for field in v2_fields)
+
+
+def _is_v2_config(cfg):
+    return cfg.get("config_schema_version") == V2_CONFIG_SCHEMA_VERSION
+
+
+def _validate_non_empty_string(value, field, path):
+    if not isinstance(value, str) or not value.strip():
+        raise ValueError(f"{path}: v2 identity field {field!r} must be a non-empty string")
+
+
+def _validate_integer(value, field, path):
+    if isinstance(value, bool) or not isinstance(value, int):
+        raise ValueError(f"{path}: v2 identity field {field!r} must be an integer")
+
+
+def _validate_benchmark_config(cfg, path="<memory>"):
+    missing_common = [field for field in ("benchmark_id",) if field not in cfg]
+    if missing_common:
+        raise ValueError(f"{path}: missing required benchmark config fields: {', '.join(missing_common)}")
+
+    schema_version = cfg.get("config_schema_version")
+    if schema_version is None:
+        if _has_v2_fields(cfg):
+            raise ValueError(f"{path}: v2 benchmark identity fields require config_schema_version=2")
+        return
+
+    if schema_version != V2_CONFIG_SCHEMA_VERSION:
+        raise ValueError(f"{path}: unsupported benchmark config_schema_version={schema_version!r}")
+
+    missing_v2 = [field for field in V2_REQUIRED_IDENTITY_FIELDS if field not in cfg]
+    if missing_v2:
+        raise ValueError(f"{path}: missing required v2 identity fields: {', '.join(missing_v2)}")
+
+    _validate_non_empty_string(cfg["workload_id"], "workload_id", path)
+    _validate_non_empty_string(cfg["variant_id"], "variant_id", path)
+    _validate_integer(cfg["benchmark_version"], "benchmark_version", path)
+
+    for field in V2_OPTIONAL_METADATA_FIELDS:
+        if field in cfg and not isinstance(cfg[field], Mapping):
+            raise ValueError(f"{path}: optional v2 metadata field {field!r} must be an object")
+
+
+def _config_identity_metadata(cfg):
+    if not _is_v2_config(cfg):
+        return {}
+    metadata = {
+        "config_schema_version": cfg["config_schema_version"],
+        "workload_id": cfg["workload_id"],
+        "variant_id": cfg["variant_id"],
+        "benchmark_version": cfg["benchmark_version"],
+    }
+    for field in V2_OPTIONAL_METADATA_FIELDS:
+        if field in cfg:
+            metadata[field] = cfg[field]
+    return metadata
+
+
+def _benchmark_display_id(cfg):
+    return cfg["benchmark_id"]
+
+
 def _discover_benchmarks():
     """Glob benchmark JSON configs and return list of (id, config) tuples."""
     pattern = os.path.join(_BENCHMARKS_DIR, "*.json")
@@ -49,6 +125,7 @@ def _discover_benchmarks():
     for path in sorted(glob.glob(pattern)):
         with open(path) as f:
             cfg = json.load(f)
+        _validate_benchmark_config(cfg, path)
         configs.append(cfg)
     return configs
 
@@ -219,6 +296,7 @@ def _run_benchmark(cfg):
 
     results = {
         "benchmark_id": cfg["benchmark_id"],
+        **_config_identity_metadata(cfg),
         "model_short_name": model_info.get("model_short_name", ""),
         "device": device_name,
         "num_gpus": init_kwargs.get("num_gpus", 1),
@@ -275,7 +353,7 @@ def _run_benchmark(cfg):
 @pytest.mark.parametrize(
     "cfg",
     _BENCHMARK_CONFIGS,
-    ids=[c["benchmark_id"] for c in _BENCHMARK_CONFIGS],
+    ids=[_benchmark_display_id(c) for c in _BENCHMARK_CONFIGS],
 )
 def test_inference_performance(cfg):
     """Measure generation latency, peak GPU memory, and component-level timings

fastvideo/tests/stages/test_text_encoding.py

@@ -21,6 +21,26 @@ class FakeTokenizer:
             "attention_mask": torch.ones(B, seq_len, dtype=torch.long),
         })
 
+
+class FakeChatTokenizer:
+    def __init__(self):
+        self.last_messages = None
+        self.last_kwargs = None
+
+    def apply_chat_template(self, messages, **kwargs):
+        self.last_messages = messages
+        self.last_kwargs = kwargs
+        assert isinstance(messages[0], list)
+        assert messages[0][0]["role"] == "system"
+        assert messages[0][1]["role"] == "user"
+        B = len(messages)
+        seq_len = int(kwargs.get("max_length", 4))
+        return TensorDict({
+            "input_ids": torch.arange(B * seq_len).view(B, seq_len),
+            "attention_mask": torch.ones(B, seq_len, dtype=torch.long),
+        })
+
+
 class FakeTextEncoder(torch.nn.Module):
     def __init__(self, hidden_size=8):
         super().__init__()
@@ -38,6 +58,14 @@ class FakeTextEncoder(torch.nn.Module):
 def id_preprocess(x: str) -> str:
     return x
 
+
+def chat_list_preprocess(x: str):
+    return [
+        {"role": "system", "content": "Describe the video."},
+        {"role": "user", "content": x if x else " "},
+    ]
+
+
 def take_mean_postprocess(outputs: BaseEncoderOutput) -> torch.Tensor:
     # [B, T, H] -> [B, H]
     return outputs.last_hidden_state.mean(dim=1)
@@ -156,3 +184,32 @@ def test_encode_text_does_not_force_hidden_states_for_ltx2_prefix():
     stage.encode_text("a", fastvideo_args, encoder_index=[0])
 
     assert stage.text_encoders[0].last_output_hidden_states is False
+
+
+def test_chat_list_preprocess_output_is_not_stripped():
+    fastvideo_args, hidden = make_args(num_encoders=1, text_len=5, hidden_size=8)
+    encoder_config = fastvideo_args.pipeline_config.text_encoder_configs[0]
+    encoder_config.is_chat_model = True
+    encoder_config.treat_empty_as_dot = True
+    fastvideo_args.pipeline_config.preprocess_text_funcs = (chat_list_preprocess, )
+
+    tokenizer = FakeChatTokenizer()
+    stage = TextEncodingStage(
+        text_encoders=[FakeTextEncoder(hidden_size=hidden)],
+        tokenizers=[tokenizer],
+    )
+
+    embeds, masks = stage.encode_text(
+        "a robotic arm welding a metal structure",
+        fastvideo_args,
+        encoder_index=[0],
+        return_attention_mask=True,
+    )
+
+    assert embeds[0].shape == (1, hidden)
+    assert masks[0].shape == (1, 5)
+    assert tokenizer.last_messages == [[
+        {"role": "system", "content": "Describe the video."},
+        {"role": "user", "content": "a robotic arm welding a metal structure"},
+    ]]
+    assert tokenizer.last_kwargs["return_tensors"] == "pt"

fastvideo/tests/train/methods/test_wan_causal_dfsft.py

@@ -127,4 +127,12 @@ def test_wan_causal_dfsft_single_train_step(
 
     # 5a-ii: device-keyed grad-norm regression on top of the same harness.
     # Skips when the current GPU has no seeded reference.
-    check_grad_norm_regression("test_wan_causal_dfsft", model.transformer)
+    # rtol above the harness default: the causal model compiles flex_attention
+    # with max-autotune (required for Wan 1.3B's head config), and the
+    # timing-based kernel selection is bimodal across L40S containers —
+    # observed 3.2562 vs 3.5860 (10.13% apart) with identical code, straddling
+    # the default 10%. 12% covers both winners; real wiring breakage (dead
+    # grads, scale bugs) still lands far outside it.
+    check_grad_norm_regression("test_wan_causal_dfsft",
+                               model.transformer,
+                               rtol=0.12)