Merge branch 'main' into codex-split-attention-offload-guides

[Bugfix][XPU][Tests]Scope XPU pytest to explicit paths to fix collection crash (#6175 )
Signed-off-by: Joshna Medisetty <joshna.medisetty@intel.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Chendi.Xue <chendi.xue@intel.com>
2026-08-15 07:00:22 +08:00 · 2026-08-14 15:23:41 -05:00 · 2026-08-14 19:13:00 +00:00 · 2026-08-14 11:47:17 +00:00 · 2026-08-14 09:49:26 +00:00 · 2026-08-14 15:51:23 +08:00
125 changed files with 10924 additions and 2198 deletions
@@ -1,153 +1,103 @@
 #!/bin/bash
-# vllm-omni customized version
-# Based on: vllm/.buildkite/scripts/hardware_ci/run-amd-test.sh
-# Last synced: 2025-12-15
-# Modifications: docker image name for vllm-omni

-# This script runs test inside the corresponding ROCm docker container.
-set -o pipefail
+# Run vLLM-Omni ROCm tests directly in the MI300 Kubernetes pod. The pod's
+# container image is selected by test-template-amd-omni.j2; MI300 has no DinD.
+set -euo pipefail

-# Export Python path
-export PYTHONPATH=".."
+: "${PYTHONFAULTHANDLER:=1}"
+: "${HF_HOME:=/home/buildkite-agent/huggingface}"
+: "${HF_HUB_DOWNLOAD_TIMEOUT:=300}"
+: "${HF_HUB_ETAG_TIMEOUT:=60}"
+: "${MIOPEN_DEBUG_CONV_DIRECT:=0}"
+: "${MIOPEN_DEBUG_CONV_GEMM:=0}"
+: "${VLLM_ROCM_USE_AITER:=0}"
+export PYTHONFAULTHANDLER HF_HOME HF_HUB_DOWNLOAD_TIMEOUT HF_HUB_ETAG_TIMEOUT
+export MIOPEN_DEBUG_CONV_DIRECT MIOPEN_DEBUG_CONV_GEMM VLLM_ROCM_USE_AITER
+export PYTORCH_ROCM_ARCH=""
+export PYTHONPATH="${PYTHONPATH:-..}"
+
+if [[ "${VLLM_CI_DOCKER_DISABLED:-0}" != "1" ]]; then
+    echo "Error: MI300 CI must run natively with Docker disabled." >&2
+    exit 1
+fi
+
+if [[ -n "${TEST_COMMANDS:-}" ]]; then
+    commands="${TEST_COMMANDS}"
+else
+    commands="$*"
+fi
+if [[ -z "${commands}" ]]; then
+    echo "Error: No test commands provided." >&2
+    exit 1
+fi
+
+echo "--- Native in-pod ROCm CI"
+
+job_id="${BUILDKITE_JOB_ID:-${BUILDKITE_PARALLEL_JOB:-local}}"
+job_id="${job_id//[^A-Za-z0-9_.-]/_}"
+native_cache_root="/tmp/vllm-omni-native-${job_id}"
+TMPDIR="${native_cache_root}/tmp"
+TORCHINDUCTOR_CACHE_DIR="${native_cache_root}/torchinductor"
+TRITON_CACHE_DIR="${native_cache_root}/triton"
+VLLM_CACHE_ROOT="${native_cache_root}/vllm"
+XDG_CACHE_HOME="${native_cache_root}/xdg"
+HF_DATASETS_CACHE="${native_cache_root}/huggingface/datasets"
+export TMPDIR TORCHINDUCTOR_CACHE_DIR TRITON_CACHE_DIR VLLM_CACHE_ROOT
+export XDG_CACHE_HOME HF_DATASETS_CACHE
+mkdir -p "${TMPDIR}" "${TORCHINDUCTOR_CACHE_DIR}" "${TRITON_CACHE_DIR}" \
+    "${VLLM_CACHE_ROOT}" "${XDG_CACHE_HOME}" "${HF_HOME}" \
+    "${HF_DATASETS_CACHE}"
+
+if [[ "${VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE:-0}" == "1" ]]; then
+    if ! command -v findmnt >/dev/null 2>&1; then
+        echo "Error: findmnt is required to verify the Hugging Face cache mount." >&2
+        exit 1
+    fi
+    hf_mount=$(findmnt -n -T "${HF_HOME}" -o TARGET 2>/dev/null || true)
+    if [[ -z "${hf_mount}" || "${hf_mount}" == "/" ]]; then
+        echo "Error: MI300 CI requires a persistent volume at or above ${HF_HOME}." >&2
+        exit 1
+    fi
+fi

-# Print ROCm version
-echo "--- ROCm info"
 rocminfo

-# cleanup older docker images
-cleanup_docker() {
-  # Get Docker's root directory
-  docker_root=$(docker info -f '{{.DockerRootDir}}')
-  if [ -z "$docker_root" ]; then
-    echo "Failed to determine Docker root directory."
-    exit 1
-  fi
-  echo "Docker root directory: $docker_root"
-  # Check disk usage of the filesystem where Docker's root directory is located
-  disk_usage=$(df "$docker_root" | tail -1 | awk '{print $5}' | sed 's/%//')
-  # Define the threshold
-  threshold=70
-  if [ "$disk_usage" -gt "$threshold" ]; then
-    echo "Disk usage is above $threshold%. Cleaning up Docker images and volumes..."
-    # Remove dangling images (those that are not tagged and not used by any container)
-    docker image prune -f
-    # Remove unused volumes / force the system prune for old images as well.
-    docker volume prune -f && docker system prune --force --filter "until=72h" --all
-    echo "Docker images and volumes cleanup completed."
-  else
-    echo "Disk usage is below $threshold%. No cleanup needed."
-  fi
-}
+expected_gpus="${VLLM_CI_EXPECTED_GPU_COUNT:-1}"
+python3 - <<'PY'
+import os

-# Call the cleanup docker function
-cleanup_docker
+import torch

-echo "--- Pulling container"
-## Temporary change to use AMD Docker Hub to store the vllm-omni image
-# to bypass the rate limit issue with ECR Public Gallery.
-# Images are now stored in a separate repository for vllm-omni, instead of vllm-ci.
-# TODO: @tjtanaa point back to ECR Public Gallery
-# once the amd agents are configured to use ECR Public Gallery.
-# image_name="public.ecr.aws/q9t5s3a7/vllm-ci-test-repo:${BUILDKITE_COMMIT}-rocm-omni"
-image_name="rocm/vllm-omni:${BUILDKITE_COMMIT}"
-container_name="rocm_${BUILDKITE_COMMIT}_$(tr -dc A-Za-z0-9 < /dev/urandom | head -c 10; echo)"
+expected = int(os.environ.get("VLLM_CI_EXPECTED_GPU_COUNT", "1"))
+assert torch.version.hip, "PyTorch is not a ROCm build"
+assert torch.cuda.is_available(), "ROCm GPU is not available to PyTorch"
+actual = torch.cuda.device_count()
+assert actual == expected, f"Expected {expected} ROCm GPU(s), found {actual}"
+PY

-# TODO: @tjtanaa uncomment this once the amd agents are configured to use ECR Public Gallery.
-# # Install AWS CLI to authenticate to ECR Public Gallery to get higher rate limit for pulling images
-# sudo apt-get update && sudo apt-get install -y awscli
-##  Use safe docker login helper to prevent race conditions
-# source "$(dirname "${BASH_SOURCE[0]}")/../../common/scripts/docker_login_ecr_public.sh"
-# safe_docker_login_ecr_public
-
-## Pull the container from AMD Docker Hub
-
-docker pull "${image_name}"
-
-remove_docker_container() {
-   docker ps -aq --filter "name=${container_name}" | xargs -r docker rm -f || true
-   docker image rm -f "${image_name}" || true
-}
-trap remove_docker_container EXIT
-
-echo "--- Running container"
-
-HF_CACHE="$(realpath ~)/huggingface"
-mkdir -p "${HF_CACHE}"
-HF_MOUNT="/root/.cache/huggingface"
-
-if [[ -n "${TEST_COMMAND:-}" ]]; then
-    commands="$TEST_COMMAND"
-else
-    commands="$@"
-fi
-echo "Commands:$commands"
-
-PARALLEL_JOB_COUNT=8
-MYPYTHONPATH=".."
-
-# Test that we're launching on the machine that has
-# proper access to GPUs
-render_gid=$(getent group render | cut -d: -f3)
-if [[ -z "$render_gid" ]]; then
-  echo "Error: 'render' group not found. This is required for GPU access." >&2
-  exit 1
-fi
-
-# check if the command contains shard flag, we will run all shards in parallel because the host have 8 GPUs.
-# TODO: @tjtanaa reenable to run VLLM_ROCM_USE_AITER=1 when AITER is shipped with prebuilt kernels.
-if [[ $commands == *"--shard-id="* ]]; then
-  # assign job count as the number of shards used
-  commands=$(echo "$commands" | sed -E "s/--num-shards[[:blank:]]*=[[:blank:]]*[0-9]*/--num-shards=${PARALLEL_JOB_COUNT} /g" | sed 's/ \\ / /g')
-  for GPU in $(seq 0 $(($PARALLEL_JOB_COUNT-1))); do
-    # assign shard-id for each shard
-    commands_gpu=$(echo "$commands" | sed -E "s/--shard-id[[:blank:]]*=[[:blank:]]*[0-9]*/--shard-id=${GPU} /g" | sed 's/ \\ / /g')
-    echo "Shard ${GPU} commands:$commands_gpu"
-    echo "Render devices: $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES"
-    docker run \
-        --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
-        --network=host \
-        --shm-size=16gb \
-        --group-add "$render_gid" \
-        -e MIOPEN_DEBUG_CONV_DIRECT=0 \
-        -e MIOPEN_DEBUG_CONV_GEMM=0 \
-        -e VLLM_ROCM_USE_AITER=0 \
-        -e HIP_VISIBLE_DEVICES="${GPU}" \
-        -e HF_TOKEN \
-        -v "${HF_CACHE}:${HF_MOUNT}" \
-        -e "HF_HOME=${HF_MOUNT}" \
-        -e "PYTHONPATH=${MYPYTHONPATH}" \
-        --name "${container_name}_${GPU}" \
-        "${image_name}" \
-        /bin/bash -c "${commands_gpu}" \
-        |& while read -r line; do echo ">>Shard $GPU: $line"; done &
-    PIDS+=($!)
-  done
-  #wait for all processes to finish and collect exit codes
-  for pid in "${PIDS[@]}"; do
-    wait "${pid}"
-    STATUS+=($?)
-  done
-  for st in "${STATUS[@]}"; do
-    if [[ ${st} -ne 0 ]]; then
-      echo "One of the processes failed with $st"
-      exit "${st}"
+echo "Commands:${commands}"
+if /bin/bash -o pipefail -c '
+set -E
+test_status=0
+trap '\''
+    command_status=$?
+    # Keep running subsequent commands, but retain a failure for the final exit.
+    # Prefer a test failure over pytest "no tests collected" when both occur.
+    if (( test_status == 0 || (test_status == 5 && command_status != 5) )); then
+        test_status=${command_status}
    fi
-  done
+'\'' ERR
+eval "$1"
+exit "${test_status}"
+' _ "${commands}"; then
+    exit 0
 else
-  echo "Render devices: $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES"
-  docker run \
-          --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
-          --network=host \
-          --shm-size=16gb \
-          --group-add "$render_gid" \
-          -e MIOPEN_DEBUG_CONV_DIRECT=0 \
-          -e MIOPEN_DEBUG_CONV_GEMM=0 \
-          -e VLLM_ROCM_USE_AITER=0 \
-          -e HF_TOKEN \
-          -v "${HF_CACHE}:${HF_MOUNT}" \
-          -e "HF_HOME=${HF_MOUNT}" \
-          -e "PYTHONPATH=${MYPYTHONPATH}" \
-          --name "${container_name}" \
-          "${image_name}" \
-          /bin/bash -c "${commands}"
+    exit_code=$?
 fi
+
+if [[ ${exit_code} -eq 5 && "${VLLM_CI_ALLOW_NO_TESTS:-0}" == "1" ]]; then
+    echo "Pytest collected no tests; VLLM_CI_ALLOW_NO_TESTS=1, treating exit code 5 as success."
+    exit 0
+fi
+
+exit "${exit_code}"
@@ -7,54 +7,64 @@ steps:

 - group: ":card_index_dividers: Simple Test"
  steps:
-    - label: "Simple · Diffusion & Model Executor Test (1/2)"
-      agent_pool: mi325_1
+    - label: "Simple · Model Executor Test · Shard %N/%t"
+      agent_pool: mi300_1
+      parallelism: 2
      depends_on: amd-build
-      timeout_in_minutes: 70
      mirror_hardwares: [amdproduction]
      grade: Blocking
      commands:
        - export VLLM_ROCM_USE_AITER=0
        # ignore test_teacache_extractors.py because it use rocm gemm kernel from vLLM
        # that is not supported on CPU
-        # CPU tests on ROCm gpu machines are very slow, so we give them more time.
-        - "timeout 60m pytest -sv tests/diffusion -m 'core_model and cpu' --ignore=tests/diffusion/cache/test_teacache_extractors.py --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"
+        - "pytest -sv tests/model_executor -m 'core_model and cpu' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml:coverage-model-executor-$$BUILDKITE_PARALLEL_JOB.xml"

-    - label: "Simple · Diffusion & Model Executor Test (2/2)"
-      agent_pool: mi325_1
+    - label: "Simple · Diffusion Test · Shard %N/%t"
+      agent_pool: mi300_1
+      parallelism: 2
      depends_on: amd-build
-      timeout_in_minutes: 70
      mirror_hardwares: [amdproduction]
      grade: Blocking
      commands:
        - export VLLM_ROCM_USE_AITER=0
-        # CPU tests on ROCm gpu machines are very slow, so we give them more time.
-        - "timeout 60m pytest -sv tests/model_executor -m 'core_model and cpu' --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"
+        # ignore test_teacache_extractors.py because it use rocm gemm kernel from vLLM
+        # that is not supported on CPU
+        - "pytest -sv tests/diffusion -m 'core_model and cpu' --ignore=tests/diffusion/cache/test_teacache_extractors.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml:coverage-diffusion-$$BUILDKITE_PARALLEL_JOB.xml"
+
+    - label: "Simple · Engine&Entrypoints Test"
+      agent_pool: mi300_1
+      depends_on: amd-build
+      mirror_hardwares: [amdproduction]
+      grade: Blocking
+      commands:
+        - export VLLM_ROCM_USE_AITER=0
+        - "pytest -sv tests/entrypoints tests/engine -m 'core_model and cpu' --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"

    - label: "Simple · Other Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
-      timeout_in_minutes: 50
      mirror_hardwares: [amdproduction]
      grade: Blocking
      commands:
        - export VLLM_ROCM_USE_AITER=0
        # ignore test_teacache_extractors.py because it use rocm gemm kernel from vLLM
        # that is not supported on CPU
-        - "timeout 40m pytest -sv tests/diffusion tests/model_executor -m 'core_model and cpu' --ignore=tests/diffusion/cache/test_teacache_extractors.py --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"
+        - "pytest -sv tests/ -m 'core_model and cpu' --ignore=tests/diffusion --ignore=tests/model_executor --ignore=tests/entrypoints --ignore=tests/engine  --ignore=tests/diffusion/cache/test_teacache_extractors.py --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"

 - group: ":card_index_dividers: Diffusion Test"
  steps:
    - label: "Diffusion · LoRA Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
+      artifact_paths:
+        - lora-server.log
      commands:
-        - timeout 20m pytest -s -v tests/diffusion/lora/test_images_generations_lora.py
+        - pytest -sv tests/diffusion/lora/ -m 'advanced_model and cuda' --run-level "advanced_model" 2>&1 | tee lora-server.log

    - label: "Diffusion Model CPU offloading Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -62,7 +72,7 @@ steps:
        - export VLLM_LOGGING_LEVEL=DEBUG
        - export VLLM_WORKER_MULTIPROC_METHOD=spawn
        - |
-          timeout 20m bash -c '
+          bash -c '
            set +e
            pytest -s -v tests/diffusion/offloader/test_diffusion_cpu_offload.py
            EXIT1=\$?
@@ -72,7 +82,7 @@ steps:
          '

    - label: "Diffusion Sequence Parallelism Test (Need 4 GPUs)"
-      agent_pool: mi325_4
+      agent_pool: mi300_4
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -84,7 +94,7 @@ steps:

    # merge-only tests
    - label: "Diffusion Tensor Parallelism Test"
-      agent_pool: mi325_2
+      agent_pool: mi300_2
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -92,7 +102,7 @@ steps:
        - pytest -s -v tests/diffusion/distributed/test_tensor_parallel.py

    - label: "Diffusion Model Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -100,7 +110,7 @@ steps:
        - timeout 30m pytest -s -v tests/e2e/offline_inference/test_z_image.py

    - label: "Diffusion GPU Worker Test"
-      agent_pool: mi325_2
+      agent_pool: mi300_2
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -108,7 +118,7 @@ steps:
        - timeout 20m pytest -s -v tests/diffusion/test_diffusion_worker.py

    - label: "Diffusion · Wan22 Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -121,7 +131,7 @@ steps:

 ## ISSUE depends on `diffusers` package: https://github.com/huggingface/diffusers/issues/13274
 # - label: "Audio Generation Model Test"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -132,7 +142,7 @@ steps:
 #     - timeout 20m pytest -s -v tests/e2e/offline_inference/test_stable_audio_expansion.py -m "advanced_model and diffusion and L4" --run-level advanced_model

 - label: "Engine Test"
-  agent_pool: mi325_1
+  agent_pool: mi300_1
  depends_on: amd-build
  mirror_hardwares: [amdproduction]
  grade: Blocking
@@ -147,7 +157,7 @@ steps:
 - group: ":card_index_dividers: E2E Test"
  steps:
    - label: "Omni Model Test Qwen2-5-Omni"
-      agent_pool: mi325_4
+      agent_pool: mi300_4
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      commands:
@@ -155,7 +165,7 @@ steps:
        - pytest -s -v tests/e2e/*/test_qwen2_5_omni_expansion.py --run-level "core_model"

    - label: "Omni Model Test Qwen3-Omni"
-      agent_pool: mi325_2
+      agent_pool: mi300_2
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -163,22 +173,23 @@ steps:
        - export VLLM_WORKER_MULTIPROC_METHOD=spawn
        - pytest -s -v tests/e2e/offline_inference/test_qwen3_omni.py tests/e2e/online_serving/test_qwen3_omni.py -m 'advanced_model and cuda' --run-level 'advanced_model'

-    - label: "Qwen3-TTS CustomVoice E2E Test"
-      agent_pool: mi325_1
+    - label: "Qwen3-TTS CustomVoice E2E Test · Shard %N/%t"
+      agent_pool: mi300_1
+      parallelism: 2
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
      commands:
        - |
-          timeout 50m bash -c '
+          bash -c '
            export VLLM_LOGGING_LEVEL=DEBUG
            export VLLM_WORKER_MULTIPROC_METHOD=spawn
            export VLLM_ALLOW_LONG_MAX_MODEL_LEN="1"
-            pytest -s -v tests/e2e/online_serving/test_qwen3_tts_customvoice.py tests/e2e/offline_inference/test_qwen3_tts_customvoice.py -m "advanced_model" --run-level "advanced_model"
+            pytest -s -v tests/e2e/online_serving/test_qwen3_tts_customvoice.py tests/e2e/offline_inference/test_qwen3_tts_customvoice.py -m "advanced_model" --run-level "advanced_model" --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
          '

    - label: "Qwen3-TTS Base E2E Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -196,7 +207,7 @@ steps:
 # Need to debug before reneable numerical changes across large PRs
 # # split Bagel Model Test with H100 (Real Weights) into three tests
 # - label: "Bagel Text2Img Model Test (1/3)"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -208,7 +219,7 @@ steps:
 #     - timeout 30m pytest -s -v tests/e2e/offline_inference/test_bagel.py -m "advanced_model" --run-level "advanced_model" -k "shared_memory" -k "rocm"

 # - label: "Bagel Img2Img Model Test (2/3)"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -220,7 +231,7 @@ steps:
 #     - timeout 30m pytest -s -v tests/e2e/offline_inference/test_bagel.py -m "advanced_model" --run-level "advanced_model" -k "rocm"

 # - label: "Bagel Online Serving Test (3/3)"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -7,8 +7,9 @@ steps:

 - group: ":card_index_dividers: Simple Test"
  steps:
-    - label: "Simple · Diffusion & Model Executor Test"
-      agent_pool: mi325_1
+    - label: "Simple · Model Executor Test · Shard %N/%t"
+      agent_pool: mi300_1
+      parallelism: 2
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -16,10 +17,30 @@ steps:
        - export VLLM_ROCM_USE_AITER=0
        # ignore test_teacache_extractors.py because it use rocm gemm kernel from vLLM
        # that is not supported on CPU
-        - "timeout 40m pytest -sv tests/diffusion tests/model_executor -m 'core_model and cpu' --ignore=tests/diffusion/cache/test_teacache_extractors.py --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"
+        - "pytest -sv tests/model_executor -m 'core_model and cpu' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml:coverage-model-executor-$$BUILDKITE_PARALLEL_JOB.xml"
+
+    - label: "Simple · Diffusion Test"
+      agent_pool: mi300_1
+      depends_on: amd-build
+      mirror_hardwares: [amdproduction]
+      grade: Blocking
+      commands:
+        - export VLLM_ROCM_USE_AITER=0
+        # ignore test_teacache_extractors.py because it use rocm gemm kernel from vLLM
+        # that is not supported on CPU
+        - "pytest -sv tests/diffusion -m 'core_model and cpu'  --ignore=tests/diffusion/cache/test_teacache_extractors.py --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"
+
+    - label: "Simple · Engine&Entrypoints Test"
+      agent_pool: mi300_1
+      depends_on: amd-build
+      mirror_hardwares: [amdproduction]
+      grade: Blocking
+      commands:
+        - export VLLM_ROCM_USE_AITER=0
+        - "pytest -sv tests/entrypoints tests/engine -m 'core_model and cpu' --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"

    - label: "Simple · Other Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -27,19 +48,19 @@ steps:
        - export VLLM_ROCM_USE_AITER=0
        # ignore test_teacache_extractors.py because it use rocm gemm kernel from vLLM
        # that is not supported on CPU
-        - "timeout 40m pytest -sv tests/diffusion tests/model_executor -m 'core_model and cpu' --ignore=tests/diffusion/cache/test_teacache_extractors.py --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"
+        - "pytest -sv tests/ -m 'core_model and cpu' --ignore=tests/diffusion --ignore=tests/model_executor --ignore=tests/entrypoints --ignore=tests/engine  --ignore=tests/diffusion/cache/test_teacache_extractors.py --cov=vllm_omni --cov-report=term-missing:skip-covered --cov-report=xml"

 - label: "Custom Pipeline Test"
-  agent_pool: mi325_1
+  agent_pool: mi300_1
  depends_on: amd-build
  mirror_hardwares: [amdproduction]
  grade: Blocking
  commands:
-    - timeout 20m pytest -s -v tests/e2e/features/custom_pipeline/ -m "core_model"
+    - pytest -s -v tests/e2e/features/custom_pipeline/ -m "core_model"


 - label: "Voxtral TTS CUDA Unit Test"
-  agent_pool: mi325_1
+  agent_pool: mi300_1
  depends_on: amd-build
  mirror_hardwares: [amdproduction]
  grade: Blocking
@@ -49,33 +70,25 @@ steps:
 - group: ":card_index_dividers: Diffusion Test"
  steps:
    - label: "Diffusion · Batch Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
      commands:
-        - timeout 20m pytest -sv tests/diffusion/batching -m 'core_model and cuda' --run-level "core_model"
+        - pytest -sv tests/diffusion/batching -m 'core_model and cuda' --run-level "core_model"

    - label: "Diffusion · Offloader Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
      commands:
        - export VLLM_LOGGING_LEVEL=DEBUG
        - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-        - |
-          timeout 20m bash -c '
-            set +e
-            pytest -s -v tests/diffusion/offloader/test_diffusion_cpu_offload.py
-            EXIT1=\$?
-            pytest -s -v tests/diffusion/offloader/test_diffusion_layerwise_offload.py
-            EXIT2=\$?
-            exit \$((EXIT1 | EXIT2))
-          '
+        - pytest -s -v tests/diffusion/offloader -m "core_model and cuda" --run-level "core_model"

    - label: "Diffusion · Cache Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -83,10 +96,10 @@ steps:
        - export GPU_ARCHS=gfx942
        - export VLLM_LOGGING_LEVEL=DEBUG
        - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-        - timeout 15m pytest -s -v -m "core_model and cache and diffusion and not distributed_cuda and L4"
+        - pytest -sv tests/diffusion/cache -m 'core_model and cuda' --run-level "core_model"

    - label: "Diffusion Sequence Parallelism Test"
-      agent_pool: mi325_2
+      agent_pool: mi300_2
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -96,7 +109,7 @@ steps:
        - timeout 20m pytest -s -v tests/diffusion/distributed/test_sequence_parallel.py -m core_model

    - label: "Diffusion GPU Worker Test"
-      agent_pool: mi325_2
+      agent_pool: mi300_2
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -104,7 +117,7 @@ steps:
        - timeout 20m pytest -s -v tests/diffusion/test_diffusion_worker.py

    - label: "Diffusion · Model Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -112,7 +125,7 @@ steps:
        - timeout 15m pytest -sv tests/diffusion/models/ -m 'core_model and cuda' --run-level "core_model"

    - label: "Diffusion · Wan22 Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -126,7 +139,7 @@ steps:

 ## ISSUE depends on `diffusers` package: https://github.com/huggingface/diffusers/issues/13274
 # - label: "Audio Generation Model Test"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -137,7 +150,7 @@ steps:
 #     - timeout 20m pytest -s -v tests/e2e/offline_inference/test_stable_audio_expansion.py -m "advanced_model and diffusion and L4" --run-level advanced_model

 - label: "Engine Test"
-  agent_pool: mi325_1
+  agent_pool: mi300_1
  depends_on: amd-build
  mirror_hardwares: [amdproduction]
  grade: Blocking
@@ -149,7 +162,7 @@ steps:
      '

 # - label: "Omni Model Test Qwen3-Omni"
-#   agent_pool: mi325_2
+#   agent_pool: mi300_2
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -162,7 +175,7 @@ steps:
 - group: ":card_index_dividers: E2E Test"
  steps:
    - label: "Qwen3-TTS E2E Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -173,7 +186,7 @@ steps:
        - timeout 30m pytest -s -v tests/e2e/online_serving/test_qwen3_tts_customvoice.py -m "core_model" --run-level "core_model"

    - label: "CosyVoice3-TTS E2E Test"
-      agent_pool: mi325_1
+      agent_pool: mi300_1
      depends_on: amd-build
      mirror_hardwares: [amdproduction]
      grade: Blocking
@@ -186,7 +199,7 @@ steps:
 # TODO: Bagel test on ROCm is very unstable. @tjtanaa
 # Need to debug before reneable numerical changes across large PRs
 # - label: "Bagel Text2Img Model Test"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -198,7 +211,7 @@ steps:
 #     - timeout 30m pytest -s -v tests/e2e/offline_inference/test_bagel.py -m "core_model" --run-level "core_model" -k "rocm"

 # - label: "Bagel Img2Img Model Test"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -210,7 +223,7 @@ steps:
 #     - timeout 30m pytest -s -v tests/e2e/offline_inference/test_bagel.py -m "core_model" --run-level "core_model" -k "rocm"

 # - label: "Bagel Online Serving Test"
-#   agent_pool: mi325_1
+#   agent_pool: mi300_1
 #   depends_on: amd-build
 #   mirror_hardwares: [amdproduction]
 #   grade: Blocking
@@ -1,11 +1,65 @@
 {# vllm-omni customized version
   Based on: https://github.com/vllm-project/ci-infra/blob/main/buildkite/test-template-amd.j2
-   Last synced: 2025-12-15
-   Modifications: Removed unused CUDA/NVIDIA logic, keeping only AMD tests
+   Last synced: 2026-08-04
+   Modifications: Use the shared MI300 queues and Omni command-passing contract
 #}
 {% set docker_image_amd = "rocm/vllm-omni:$BUILDKITE_COMMIT" %}
 {% set default_working_dir = "/app/vllm-omni" %}

+{# MI300 runs tests natively in Kubernetes; Docker-in-Docker is unavailable. #}
+{% macro amd_native_k8s_plugin(indent, step) -%}
+{{ indent }}plugins:
+{{ indent }}  - kubernetes:
+{{ indent }}      podSpecPatch:
+{{ indent }}        automountServiceAccountToken: false
+{{ indent }}        securityContext:
+{{ indent }}          seccompProfile:
+{{ indent }}            type: RuntimeDefault
+{{ indent }}        imagePullSecrets:
+{{ indent }}          - name: docker-config
+{{ indent }}        containers:
+{{ indent }}          - name: container-0
+{{ indent }}            image: {{ docker_image_amd }}
+{{ indent }}            imagePullPolicy: Always
+{{ indent }}            securityContext:
+{{ indent }}              allowPrivilegeEscalation: false
+{{ indent }}              capabilities:
+{{ indent }}                add:
+{{ indent }}                  - IPC_LOCK
+{{ indent }}            resources:
+{{ indent }}              limits:
+{{ indent }}                amd.com/gpu: "{{ step.agent_pool[-1:] }}"
+{{ indent }}              requests:
+{{ indent }}                amd.com/gpu: "{{ step.agent_pool[-1:] }}"
+{{ indent }}            volumeMounts:
+{{ indent }}              - name: devshm
+{{ indent }}                mountPath: /dev/shm
+{{ indent }}            env:
+{{ indent }}              - name: AMD_CI_RUNTIME
+{{ indent }}                value: native
+{{ indent }}              - name: NATIVE_CI
+{{ indent }}                value: "true"
+{{ indent }}              - name: VLLM_CI_DOCKER_DISABLED
+{{ indent }}                value: "1"
+{{ indent }}              - name: VLLM_CI_EXPECTED_GPU_COUNT
+{{ indent }}                value: "{{ step.agent_pool[-1:] }}"
+{{ indent }}              - name: HF_HOME
+{{ indent }}                value: /home/buildkite-agent/huggingface
+{{ indent }}              - name: MIOPEN_DEBUG_CONV_DIRECT
+{{ indent }}                value: "0"
+{{ indent }}              - name: MIOPEN_DEBUG_CONV_GEMM
+{{ indent }}                value: "0"
+{{ indent }}              - name: VLLM_ROCM_USE_AITER
+{{ indent }}                value: "0"
+{{ indent }}              - name: PYTORCH_ROCM_ARCH
+{{ indent }}                value: ""
+{{ indent }}        volumes:
+{{ indent }}          - name: devshm
+{{ indent }}            emptyDir:
+{{ indent }}              medium: Memory
+{{ indent }}              sizeLimit: 16Gi
+{%- endmacro %}
+
 {% if env %}
 env:
 {% for name in env %}
@@ -53,19 +107,31 @@ steps:
          {% if step.agent_pool %}
          queue: amd_{{ step.agent_pool }}
          {% else %}
-          queue: amd_mi325_1
+          queue: amd_mi300_1
          {% endif %}
+{{ amd_native_k8s_plugin("        ", step) }}
 {% set cmd_body = (step.command or (step.commands | join("\n"))) | trim %}
 {% set indented_cmd = cmd_body | replace("\n", "\n            ") %}
        command: bash .buildkite/amd/scripts/run-amd-test.sh
        env:
-          DOCKER_BUILDKIT: "1"
+          AMD_CI_RUNTIME: "native"
+          NATIVE_CI: "true"
+          VLLM_CI_DOCKER_DISABLED: "1"
+          VLLM_CI_EXPECTED_GPU_COUNT: "{{ step.agent_pool[-1:] }}"
+          VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE: "1"
+          HF_HOME: "/home/buildkite-agent/huggingface"
+          HF_HUB_DOWNLOAD_TIMEOUT: "300"
+          HF_HUB_ETAG_TIMEOUT: "60"
+          MIOPEN_DEBUG_CONV_DIRECT: "0"
+          MIOPEN_DEBUG_CONV_GEMM: "0"
+          VLLM_ROCM_USE_AITER: "0"
+          PYTORCH_ROCM_ARCH: ""
 {% if step.env %}
 {% for name in step.env %}
          {{ name }}: "{{ step.env[name] }}"
 {% endfor %}
 {% endif %}
-          TEST_COMMAND: |-
+          TEST_COMMANDS: |-
            (command rocm-smi || true) && cd {{ (step.working_dir or default_working_dir) | safe }}
 {% if env %}
 {% for name in env %}
@@ -82,6 +148,15 @@ steps:
 {% endif %}
            {{ indented_cmd | safe }}
        priority: 100
+{% if step.parallelism %}
+        parallelism: {{ step.parallelism }}
+{% endif %}
+{% if step.artifact_paths %}
+        artifact_paths:
+{% for path in step.artifact_paths %}
+          - "{{ path }}"
+{% endfor %}
+{% endif %}
        {% if step.timeout_in_minutes %}
        timeout_in_minutes: {{ step.timeout_in_minutes }}
        {% endif %}
@@ -100,19 +175,31 @@ steps:
      {% if step.agent_pool %}
      queue: amd_{{ step.agent_pool }}
      {% else %}
-      queue: amd_mi325_1
+      queue: amd_mi300_1
      {% endif %}
+{{ amd_native_k8s_plugin("    ", step) }}
 {% set cmd_body = (step.command or (step.commands | join("\n"))) | trim %}
 {% set indented_cmd = cmd_body | replace("\n", "\n        ") %}
    command: bash .buildkite/amd/scripts/run-amd-test.sh
    env:
-      DOCKER_BUILDKIT: "1"
+      AMD_CI_RUNTIME: "native"
+      NATIVE_CI: "true"
+      VLLM_CI_DOCKER_DISABLED: "1"
+      VLLM_CI_EXPECTED_GPU_COUNT: "{{ step.agent_pool[-1:] }}"
+      VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE: "1"
+      HF_HOME: "/home/buildkite-agent/huggingface"
+      HF_HUB_DOWNLOAD_TIMEOUT: "300"
+      HF_HUB_ETAG_TIMEOUT: "60"
+      MIOPEN_DEBUG_CONV_DIRECT: "0"
+      MIOPEN_DEBUG_CONV_GEMM: "0"
+      VLLM_ROCM_USE_AITER: "0"
+      PYTORCH_ROCM_ARCH: ""
 {% if step.env %}
 {% for name in step.env %}
      {{ name }}: "{{ step.env[name] }}"
 {% endfor %}
 {% endif %}
-      TEST_COMMAND: |-
+      TEST_COMMANDS: |-
        (command rocm-smi || true) && cd {{ (step.working_dir or default_working_dir) | safe }}
 {% if env %}
 {% for name in env %}
@@ -129,6 +216,15 @@ steps:
 {% endif %}
        {{ indented_cmd | safe }}
    priority: 100
+{% if step.parallelism %}
+    parallelism: {{ step.parallelism }}
+{% endif %}
+{% if step.artifact_paths %}
+    artifact_paths:
+{% for path in step.artifact_paths %}
+      - "{{ path }}"
+{% endfor %}
+{% endif %}
    {% if step.timeout_in_minutes %}
    timeout_in_minutes: {{ step.timeout_in_minutes }}
    {% endif %}
@@ -9,6 +9,11 @@
 #   a2b3_npu_1  — ascend-a2b3, resource_class npu-1 (A2/B3 CI image)
 #   a2b3_npu_4  — ascend-a2b3, resource_class npu-4
 #   a3_npu_2    — ascend-a3, resource_class npu-2
+#
+# The ``l4_*`` presets run under the ``docker`` plugin, which does not put
+# ``buildkite-agent`` on PATH inside the container the way the ``kubernetes``
+# presets do, so they set ``mount-buildkite-agent: true``; without it any step
+# calling ``buildkite-agent`` (artifact upload, annotate, meta-data) fails there.

 x-npu-k8s-plugins: &npu_k8s_plugins
  - kubernetes:
@@ -26,6 +31,7 @@ mirror_hardwares:
          image: public.ecr.aws/q9t5s3a7/vllm-ci-test-repo:$BUILDKITE_COMMIT
          always-pull: true
          propagate-environment: true
+          mount-buildkite-agent: true
          shm-size: "8gb"
          environment:
            - HF_HOME=/fsx/hf_cache
@@ -41,6 +47,7 @@ mirror_hardwares:
          image: public.ecr.aws/q9t5s3a7/vllm-ci-test-repo:$BUILDKITE_COMMIT
          always-pull: true
          propagate-environment: true
+          mount-buildkite-agent: true
          shm-size: "8gb"
          environment:
            - HF_HOME=/fsx/hf_cache
@@ -0,0 +1,82 @@
+#!/usr/bin/env bash
+# Run a model's E2E tests once per entry mode as separate coverage runs and upload
+# the reports, so per-mode attribution survives (see vllm-project/vllm-omni#5332).
+#
+# Usage: run_cov_split.sh --model-id <id> [--offline <paths>] [--online <paths>] \
+#                         -- <pytest args...>
+#
+#   --model-id  Names the artifacts: coverage-<id>-{offline,online}-<step id>.xml.gz.
+#               Use the model's directory name under vllm_omni/*/models/.
+#   --offline   Pytest path(s) for the offline half. Repeatable.
+#   --online    Pytest path(s) for the online half. Repeatable.
+#   --          Everything after this is forwarded verbatim to every pytest run,
+#               quoting preserved: the job's own -m expression, --run-level, and
+#               anything else it needs (e.g. --test-config-file).
+#
+# At least one of --offline/--online is required, so a job with only one mode runs
+# just that half. Total runtime is bounded by the step's timeout_in_minutes.
+#
+# Exits non-zero if any half or the upload failed.
+set -uo pipefail
+
+MODEL_ID=""
+OFFLINE=()
+ONLINE=()
+PYTEST_ARGS=()
+
+while [[ $# -gt 0 ]]; do
+    case "$1" in
+        --model-id) MODEL_ID="$2"; shift 2 ;;
+        --offline)  OFFLINE+=("$2"); shift 2 ;;
+        --online)   ONLINE+=("$2"); shift 2 ;;
+        --)         shift; PYTEST_ARGS=("$@"); break ;;
+        *) echo "run_cov_split.sh: unknown argument: $1" >&2; exit 2 ;;
+    esac
+done
+
+if [[ -z "${MODEL_ID}" ]]; then
+    echo "run_cov_split.sh: missing --model-id" >&2
+    exit 2
+fi
+
+if ((${#OFFLINE[@]} == 0 && ${#ONLINE[@]} == 0)); then
+    echo "run_cov_split.sh: need at least one of --offline / --online" >&2
+    exit 2
+fi
+
+# Two steps can test the same model, and artifacts are scoped to the build (a
+# nightly one carries the ready, merge and nightly tiers at once), so unsuffixed
+# reports would share a path -- a glob download then races to rename onto one file.
+# Not BUILDKITE_JOB_ID: the step id survives retries, so attempts collapse under
+# Buildkite's latest-attempt lookup instead of accumulating. A parallel step would
+# need BUILDKITE_PARALLEL_JOB as well; no pilot is one.
+REPORT_SUFFIX="${BUILDKITE_STEP_ID:-local}"
+OFFLINE_REPORT="coverage-${MODEL_ID}-offline-${REPORT_SUFFIX}.xml"
+ONLINE_REPORT="coverage-${MODEL_ID}-online-${REPORT_SUFFIX}.xml"
+
+run_half() {
+    local mode="$1"
+    shift
+    local report="coverage-${MODEL_ID}-${mode}-${REPORT_SUFFIX}.xml"
+    echo "--- coverage: ${MODEL_ID} ${mode}"
+    pytest -s -v "$@" "${PYTEST_ARGS[@]}" --cov=vllm_omni --cov-report="xml:${report}"
+    local status=$?
+    # Cobertura lists every statement in vllm_omni whether or not it ran, so the
+    # report is ~7MB no matter how little the job covers. The repetition is what
+    # dominates it, and gzip takes that to ~400KB.
+    gzip -9 -f "${report}" || status=1
+    return "${status}"
+}
+
+# Clear both modes, not just the ones being run: the upload glob matches both, so
+# a report left by an earlier run would otherwise be published as if the mode it
+# belongs to had run this time.
+rm -f "${OFFLINE_REPORT}" "${ONLINE_REPORT}" "${OFFLINE_REPORT}.gz" "${ONLINE_REPORT}.gz"
+
+EXIT=0
+((${#OFFLINE[@]})) && { run_half offline "${OFFLINE[@]}" || EXIT=1; }
+((${#ONLINE[@]})) && { run_half online "${ONLINE[@]}" || EXIT=1; }
+
+buildkite-agent artifact upload "coverage-${MODEL_ID}-*-${REPORT_SUFFIX}.xml.gz" || EXIT=1
+
+exit "${EXIT}"
@@ -94,13 +94,19 @@ steps:
          - vllm_omni/model_executor/models/qwen3_tts/
          - vllm_omni/model_executor/stage_input_processors/qwen3_tts.py
          - vllm_omni/deploy/qwen3_tts.yaml
+          - .buildkite/common/scripts/run_cov_split.sh
+          - pyproject.toml
+        # 40 -> 60: the coverage split runs the model twice.
+        timeout_in_minutes: 60
        commands:
          - |
-            timeout 40m bash -c "
-              export VLLM_LOGGING_LEVEL=DEBUG
-              export VLLM_ALLOW_LONG_MAX_MODEL_LEN=1
-              pytest -s -v tests/e2e/online_serving/test_qwen3_tts_base.py tests/e2e/offline_inference/test_qwen3_tts_base.py -m 'advanced_model and cuda' --run-level 'advanced_model'
-            "
+            export VLLM_LOGGING_LEVEL=DEBUG
+            export VLLM_ALLOW_LONG_MAX_MODEL_LEN=1
+            .buildkite/common/scripts/run_cov_split.sh \
+              --model-id qwen3_tts \
+              --offline tests/e2e/offline_inference/test_qwen3_tts_base.py \
+              --online tests/e2e/online_serving/test_qwen3_tts_base.py \
+              -- -m 'advanced_model and cuda' --run-level 'advanced_model'
        mirror_hardwares: l4_1

      - label: "TTS · Higgs-Audio-V3 Test"
@@ -203,9 +209,17 @@ steps:
          - vllm_omni/diffusion/models/bagel/
          - vllm_omni/model_executor/stage_input_processors/bagel.py
          - vllm_omni/deploy/bagel.yaml
-        timeout_in_minutes: 20
+          - .buildkite/common/scripts/run_cov_split.sh
+          - pyproject.toml
+        # 20 -> 40: the coverage split runs the model twice.
+        timeout_in_minutes: 40
        commands:
-          - pytest -s -v tests/e2e/offline_inference/test_bagel.py tests/e2e/online_serving/test_bagel.py -m 'advanced_model and cuda' --run-level 'advanced_model'
+          - |
+            .buildkite/common/scripts/run_cov_split.sh \
+              --model-id bagel \
+              --offline tests/e2e/offline_inference/test_bagel.py \
+              --online tests/e2e/online_serving/test_bagel.py \
+              -- -m 'advanced_model and cuda' --run-level 'advanced_model'
        mirror_hardwares: h100_1

      - label: "Diffusion · Qwen Image Test"
@@ -456,7 +456,7 @@ steps:
        timeout_in_minutes: 90
        commands:
          - pytest -s -v tests/e2e/offline_inference/test_wan22_autoround_w4a16_expansion.py -m "full_model and cuda" --run-level "full_model"
-          - pytest -s -v tests/e2e/online_serving/test_wan_2_1_vace_expansion.py tests/e2e/online_serving/test_ltx2_expansion.py -m "full_model and cuda" --run-level "full_model"
+          - pytest -s -v tests/e2e/online_serving/test_wan_2_1_vace_expansion.py tests/e2e/online_serving/test_ltx.py -m "full_model and cuda" --run-level "full_model"
        mirror_hardwares: h100_2

      - label: ":full_moon: Diffusion X2V · Wan2.2 Accuracy Test"
@@ -477,10 +477,16 @@ steps:
        commands:
          - pytest -s -v tests/e2e/accuracy/hunyuanvideo15_t2v/test_hunyuanvideo15_t2v_video_similarity.py -m full_model --run-level full_model
          - pytest -s -v tests/e2e/accuracy/hunyuanvideo15_i2v/test_hunyuanvideo15_i2v_video_similarity.py -m full_model --run-level full_model
-          - pytest -s -v 'tests/e2e/accuracy/ltx/test_ltx_official_similarity.py::test_ltx_one_stage_matches_official[ltx2_3_i2v]' -m full_model --run-level full_model
          - pytest -s -v tests/e2e/accuracy/test_video_streaming_output_similarity.py -m full_model --run-level full_model
        mirror_hardwares: h100_2

+      - label: ":full_moon: Diffusion X2V · LTX Accuracy Test"
+        timeout_in_minutes: 180
+        commands:
+          - pytest -s -v 'tests/e2e/accuracy/ltx/test_ltx_official_similarity.py::test_ltx_matches_official[ltx2_3_i2v]' -m full_model --run-level full_model
+          - pytest -s -v 'tests/e2e/accuracy/ltx/test_ltx_official_similarity.py::test_ltx_matches_official[ltx23_two_stage_layer_fused_i2v]' -m full_model --run-level full_model
+        mirror_hardwares: h100_1
+
      - label: ":full_moon: Diffusion X2V · Perf Test"
        key: nightly-diffusion-x2v-performance
        timeout_in_minutes: 180
@@ -87,8 +87,9 @@ time timeout -k 30 30m docker run \
    echo $ZE_AFFINITY_MASK
    pip install tblib==3.1.0
    cd /workspace/vllm-omni
-    pytest -v -s -m "core_model and xpu and B60"
+    XPU_TEST_PATHS="tests/diffusion tests/dfx tests/e2e"
+    pytest -v -s $XPU_TEST_PATHS -m "core_model and xpu and B60"
    pytest -v -s tests/diffusion/quantization/test_mxfp8_config.py
-    pytest -v -s -m "advanced_model and xpu and B60"
-    pytest -v -s -m "omni and xpu and B60"
+    pytest -v -s $XPU_TEST_PATHS -m "advanced_model and xpu and B60"
+    pytest -v -s $XPU_TEST_PATHS -m "omni and xpu and B60"
 '
@@ -46,6 +46,8 @@ htmlcov/
 .cache
 nosetests.xml
 coverage.xml
+coverage-*.xml
+coverage-*.xml.gz
 *.cover
 *.py,cover
 .hypothesis/
@@ -16,6 +16,7 @@ Easy, fast, and cheap omni-modality model serving for everyone
 ---

 *Latest News* 🔥
+- [2026/08] [VeRL-Omni](https://github.com/verl-project/verl-omni) `v0.2.0` is released: faster diffusion RL powered by vLLM-Omni (request-level/step-wise batching with FA3), rebuilt Qwen3-Omni multimodal training (DPO & GSPO), plus LTX-2.3, Qwen-Image-Edit support and more. See the [release notes](https://github.com/verl-project/verl-omni/releases/tag/v0.2.0).
 - [2026/08] We released [0.26.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.26.0) - aligned with the vLLM 0.26 release line, featuring [MiniMax H3](recipes/MiniMaxAI/MiniMax-H3.md) joint video/audio generation, an experimental full-duplex realtime runtime for [MiniCPM-o 4.5](recipes/OpenBMB/MiniCPM-o-4_5.md), distributed layerwise diffusion offload, and broader model, hardware, streaming, TTS, and quantization support.
 - [2026/07] We released [0.24.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.24.0) - aligned with the vLLM 0.24 release line, expanding production-ready coverage across TTS, speech, diffusion, image/video generation, and robot-policy serving, with major Omni stage runtime refactoring, diffusion request-level batching, async output materialization, quantization/cache/memory improvements, and broad CUDA/ROCm/XPU/NPU support.
 - [2026/06] Starting with [0.14.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.14.0), vLLM-Omni publishes a stable release aligned with every even-numbered upstream vLLM minor version. [0.16.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.16.0), [0.18.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.18.0), [0.20.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.20.0), and [0.22.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.22.0) continued this cadence, expanding omni and world-model support with [NVIDIA Cosmos3](recipes/cosmos3/Cosmos3-Nano.md) and DreamZero, adding models such as MiniCPM-o 4.5, MOSS-TTS, and Lance, and advancing TTS, diffusion, distributed execution, quantization, RL integration through [VeRL-Omni](https://github.com/verl-project/verl-omni), and CUDA/ROCm/MUSA/NPU/XPU coverage.
@@ -1,4 +1,4 @@
-ARG BASE_IMAGE=vllm/vllm-openai-rocm:v0.26.0
+ARG BASE_IMAGE=vllm/vllm-openai-rocm:v0.27.0
 FROM ${BASE_IMAGE} AS base

 # Declare a variable to know if we want to use the nightly build or the stable build.
@@ -10,7 +10,7 @@ FROM ${BASE_IMAGE} AS base
 #    we should swap over to use stable release ASAP.
 #    We should avoid relying on custom commits.
 ARG USE_NIGHTLY_BUILD=0
-ARG VLLM_VERSION_OR_COMMIT_HASH=v0.26.0
+ARG VLLM_VERSION_OR_COMMIT_HASH=v0.27.0
 ARG ARG_PYTORCH_ROCM_ARCH
 ENV PYTORCH_ROCM_ARCH=${ARG_PYTORCH_ROCM_ARCH:-${PYTORCH_ROCM_ARCH}}

@@ -79,9 +79,16 @@ RUN printf '%s\n' \
    '    os._exit(_exit_code)' \
    > ${COMMON_WORKDIR}/vllm-omni/conftest.py

+# Some optional dependencies pull in the CPU ONNX Runtime distribution. The
+# CPU and ROCm wheels own the same import path and cannot coexist, so remove
+# the CPU wheel and reinstall the ROCm wheel after dependency resolution.
 RUN cd ${COMMON_WORKDIR}/vllm-omni && \
    uv pip install --python "$(python3 -c 'import sys; print(sys.executable)')" \
        --no-cache-dir ".[dev]" --no-build-isolation && \
+    python3 -m pip uninstall -y onnxruntime && \
+    python3 -m pip install --no-cache-dir --force-reinstall --no-deps \
+        "onnxruntime-rocm>=1.22.2" && \
+    python3 -c "import importlib.metadata as m; import onnxruntime as ort; assert 'ROCMExecutionProvider' in ort.get_available_providers(), ort.get_available_providers(); assert m.version('onnxruntime-rocm'); assert not any(d.metadata['Name'] == 'onnxruntime' for d in m.distributions())" && \
    python3 -c "from torchcodec.decoders import VideoDecoder"

 RUN ln -sf /usr/bin/python3 /usr/bin/python
@@ -60,7 +60,11 @@ nav:
      - Overview and Compatibility:
        - Overview: user_guide/diffusion_features.md
        - Feature Compatibility: user_guide/feature_compatibility.md
-      - CPU Offloading: user_guide/diffusion/cpu_offload.md
+      - CPU Offloading:
+        - Overview: user_guide/diffusion/cpu_offload.md
+        - Model-Level: user_guide/diffusion/offloader/cpu_offload.md
+        - Layerwise: user_guide/diffusion/offloader/layerwise_offload.md
+        - Distributed Layerwise: user_guide/diffusion/offloader/distributed_layerwise_offload.md
      - Cache Acceleration:
        - TeaCache: user_guide/diffusion/cache_acceleration/teacache.md
        - Cache-DiT: user_guide/diffusion/cache_acceleration/cache_dit.md
@@ -71,8 +75,14 @@ nav:
        - Hybrid Sharded Data Parallel: user_guide/diffusion/parallelism/hsdp.md
        - Sequence Parallel: user_guide/diffusion/parallelism/sequence_parallel.md
        - Tensor Parallel: user_guide/diffusion/parallelism/tensor_parallel.md
-        - VAE Parallel: user_guide/diffusion/parallelism/vae_parallel.md
-      - Attention Backends: user_guide/diffusion/attention_backends.md
+        - VAE Parallelism: user_guide/diffusion/parallelism/vae_parallelism.md
+      - Attention Backends:
+        - Overview: user_guide/diffusion/attention_backends.md
+        - Dense Backends: user_guide/diffusion/attention_backends/dense_backends.md
+        - TRTLLM: user_guide/diffusion/attention_backends/trtllm.md
+        - SageAttention: user_guide/diffusion/attention_backends/sage.md
+        - Hugging Face Hub: user_guide/diffusion/attention_backends/huggingface_hub.md
+        - RainFusion: user_guide/diffusion/attention_backends/rainfusion.md
      - Regional Compilation: user_guide/diffusion/regional_compilation.md
      - Frame Interpolation: user_guide/diffusion/frame_interpolation.md
      - Startup and Loading: user_guide/diffusion/startup_and_loading.md
@@ -123,13 +133,18 @@ nav:
          - Pipeline Parallel: design/feature/pipeline_parallel.md
          - Sequence Parallel: design/feature/sequence_parallel.md
          - Tensor Parallel: design/feature/tensor_parallel.md
-          - VAE Parallel: design/feature/vae_parallel.md
+          - VAE Patch Parallelism: design/feature/vae_parallel.md
        - Attention Optimization:
+          - Backend Selection: design/feature/attention_backend_selection.md
          - Skip-Softmax: design/feature/skip_softmax.md
+        - CPU Offloading:
+          - Overview: design/feature/offloader/README.md
+          - Model-Level: design/feature/offloader/cpu_offload.md
+          - Layerwise: design/feature/offloader/layerwise_offload.md
+          - Distributed Layerwise: design/feature/offloader/distributed_layerwise_offload.md
        - Cache-DiT: design/feature/cache_dit.md
        - TeaCache: design/feature/teacache.md
        - Diffusion Continuous Batching: design/feature/diffusion_continuous_batching.md
-        - Distributed Layerwise Offload: design/feature/distributed_layerwise_offload.md
    - Infrastructure and Performance:
      - Prometheus Metrics: design/metrics.md
      - Speech Generation Performance Optimizations: design/qwen3_omni_tts_performance_optimization.md
@@ -327,6 +327,96 @@ On PR builds, `upload_pipeline.py` logs `skip '…' (no changes under …)` for
 - Use **groups** for dashboard readability; keep **E2E Test** as the group name CUDA diff filtering expects for `--e2e` nightly runs on main.
 - Prefix labels by model domain: **Omni ·**, **TTS ·**, **Diffusion ·**, **Simple ·**, etc., matching existing steps.

+### Per-model coverage
+
+Pilot (v1): two Merge-tier (L3) jobs in `.buildkite/cuda/test-merge.yml` upload
+per-model, per-entry-mode coverage as Buildkite artifacts — "Diffusion · Bagel
+Test" and "TTS · Qwen3-TTS Base Test". Both run on a single GPU so the pilot is
+cheap to reproduce. This is a pilot, not a rollout: no dashboard/visualization
+lives in this repo, and nothing is gated on coverage.
+
+**Naming convention**: `coverage-<model_id>-<mode>-<step_id>.xml.gz`, where
+`<model_id>` is the model's directory name under
+`vllm_omni/diffusion/models/<model_id>/` or
+`vllm_omni/model_executor/models/<model_id>/` (e.g. `z_image`, `qwen3_omni`),
+`<mode>` is `online` or `offline`, and `<step_id>` is `BUILDKITE_STEP_ID`
+(`local` outside Buildkite). `gunzip` before feeding a report to a Cobertura
+consumer.
+
+Reports are gzipped because `--cov=vllm_omni` sets coverage's `source`, which
+makes it walk the package and emit one `<line>` element per statement for every
+file — including the ones a single model's tests never import. That inventory,
+not the coverage, is the size: a Bagel report is 7.1 MB, of which 96.5% is
+`<line>` elements and 77% belongs to the 840 of 1169 files that run never
+imported. A run covering zero lines produces the same 7 MB. gzip takes it to
+roughly 410 KB and changes nothing about the data.
+
+`<step_id>` is what keeps two steps that cover the same model apart. Artifacts are
+scoped to the build, and a nightly build carries the ready, merge and nightly tiers
+at once, so several same-model steps land in one namespace — `qwen3_tts` already has
+both a Base and a CustomVoice job, and `minicpmo_4_5` a base and a duplex one. Two
+uploads on one path do not fail loudly: an exact-name download can report the path
+as ambiguous, and a glob download fetches every match concurrently and renames each
+onto the same destination, so whichever finishes last wins. The step id remains the
+logical step identifier across retries, while Buildkite's default artifact lookup
+selects only the latest attempt — `BUILDKITE_JOB_ID` would instead give every
+attempt its own filename, leaving a failed attempt's partial report to be picked up
+alongside the real one. A step using `parallelism`/matrix would need
+`BUILDKITE_PARALLEL_JOB` added, since its jobs share a step id; neither pilot is
+one.
+
+**Opting in a new model**: replace the job's combined `pytest` command with
+[`run_cov_split.sh`](https://github.com/vllm-project/vllm-omni/blob/main/.buildkite/common/scripts/run_cov_split.sh),
+which runs each mode, names the reports, uploads them, and fails the step if
+either half or the upload failed:
+
+```yaml
+commands:
+  - |
+    .buildkite/common/scripts/run_cov_split.sh \
+      --model-id <model_id> \
+      --offline <offline test path> \
+      --online <online test path> \
+      -- <the job's existing pytest flags, e.g. -m '...' --run-level '...'>
+```
+
+Everything after `--` is forwarded verbatim to both runs with quoting preserved,
+so an expression like `-m 'advanced_model and cuda'` survives intact and a job can
+pass anything else it needs. `--offline`/`--online` are repeatable and only one is
+required, so a job with tests in a single mode runs just that half. Runtime is
+bounded solely by the step's `timeout_in_minutes` — raise it, since the split
+loads the model once per mode.
+
+Before splitting, confirm each half actually collects at least one test under the job's forwarded filters — pytest exits with code 5 ("no tests collected") if a half is empty, which the script reports as a failure, red-ing a job that used to pass as one combined invocation.
+
+Also check the tests' `num_cards` against the job's `mirror_hardwares` preset.
+`cuda_marks` turns `num_cards > 1` into `skipif(device_count() < num_cards)`, so a
+test asking for more GPUs than the preset provides is silently skipped and its
+coverage artifact comes back empty. Splitting such a job produces a file that
+looks fine and measures nothing — "Diffusion · Z Image Test" is the current
+example (online tests want 4 cards, the job runs on `l4_1`), which is why it is
+not a coverage pilot.
+
+Online-mode coverage depends on the `[tool.coverage.run]` `patch`/`sigterm`
+settings in `pyproject.toml`: online tests launch the server with
+`subprocess.Popen` and stop it with SIGTERM, and without those settings the XML
+reflects only the pytest parent process, not the server's code paths.
+
+The upload depends on `buildkite-agent` being callable inside the container. The
+`kubernetes` presets (`h100_*`, `*_npu_*`) provide it; the `docker` ones (`l4_*`)
+only do because they set `mount-buildkite-agent: true`. A new docker preset that
+runs a coverage job needs the same.
+
+List both `run_cov_split.sh` and `pyproject.toml` in every opted-in job's
+`source_file_dependencies` — both change what the job measures, so without them a
+change there is filtered out of normal PR builds and only surfaces in a later
+nightly. `tests/buildkite/test_upload_pipeline.py` asserts this for the pilots.
+Editing only the surrounding CI YAML still does not schedule them, so a PR that
+touches just the wiring needs a full E2E run (or the commands run on a GPU host)
+to produce artifacts. When checking a new model's
+artifacts, compare `lines-covered` between the online and offline XML rather than
+just confirming both files exist.
+
 ### Validation checklist

 | Check | Command / location |
@@ -398,7 +398,7 @@ the source for configuration, compatibility, and implementation details.
 * **Composable parallelism:** Diffusion stages can combine [CFG-Parallel](feature/cfg_parallel.md),
  [Expert Parallel](feature/expert_parallel.md), [HSDP](feature/hsdp.md),
  [Pipeline Parallel](feature/pipeline_parallel.md), [Sequence Parallel](feature/sequence_parallel.md),
-  [Tensor Parallel](feature/tensor_parallel.md), and [VAE Parallel](feature/vae_parallel.md)
+  [Tensor Parallel](feature/tensor_parallel.md), and [VAE Patch Parallelism](feature/vae_parallel.md)
  according to the pipeline and hardware topology.
 * **Attention and cache acceleration:** [Skip-Softmax](feature/skip_softmax.md),
  [Cache-DiT](feature/cache_dit.md), and [TeaCache](feature/teacache.md)
@@ -406,7 +406,7 @@ the source for configuration, compatibility, and implementation details.
  stage contract.
 * **Quantization and memory efficiency:** [Quantization](feature/quantization.md)
  resolves per-pipeline or per-component quantization configurations, while
-  [Distributed Layerwise Offload](feature/distributed_layerwise_offload.md)
+  [Distributed Layerwise Offload](feature/offloader/distributed_layerwise_offload.md)
  streams diffusion blocks from host memory within the existing parallel
  topology.

@@ -0,0 +1,89 @@
+# Diffusion Attention Backend Selection
+
+This document defines the selection and extension contract for diffusion
+attention backends. User-facing backend choices, installation, and tuning are
+in the [attention backend guides](../../user_guide/diffusion/attention_backends.md).
+
+## Scope
+
+The contract applies to DiT and other diffusion attention layers. It is
+separate from vLLM's autoregressive attention selector.
+
+The selection path has four responsibilities:
+
+1. normalize user configuration into `AttentionConfig` and `AttentionSpec`;
+2. resolve a spec for an attention role;
+3. ask the active platform to validate an explicit backend or choose a
+   hardware default; and
+4. load the selected `AttentionBackend` class from the registry path.
+
+## Resolution contract
+
+`get_attn_backend_for_role()` resolves configuration in this order:
+
+1. exact `per_role[role]`;
+2. category `per_role[role_category]`;
+3. `default`;
+4. platform default.
+
+An explicit resolution returns both the backend class and its
+`AttentionSpec`. A platform-default resolution returns the class and `None`.
+Layers must therefore treat the spec as optional and must not infer that a
+missing spec means SDPA.
+
+Class resolution is cached by backend name, head size, and whether TRTLLM may
+be selected automatically. Logging is cached separately by role and source so
+the startup record identifies why each role received its backend.
+
+## Registry and platform boundary
+
+`DiffusionAttentionBackendEnum` maps stable configuration names to default
+qualified class paths. `register_diffusion_backend()` may replace a path at
+runtime without changing the public enum value.
+
+The active `OmniPlatform` owns compatibility policy through
+`get_diffusion_attn_backend_cls()`. It must:
+
+- validate explicit selections and fail with an actionable error when the
+  requested kernel cannot run;
+- choose only an available, compatible default when no backend is explicit;
+- account for head size and other platform-visible constraints; and
+- return a qualified class path implementing `AttentionBackend`.
+
+The selector must not duplicate device capability or package-availability
+policy that belongs to the platform.
+
+## Typed backend options
+
+Backend-specific settings remain typed fields on `AttentionSpec` rather than
+unstructured keyword dictionaries:
+
+- `quant` is consumed by FlashInfer and TRTLLM with backend-specific value
+  validation;
+- `skip_softmax` is consumed by TRTLLM; and
+- `block_sparse` is consumed by block-sparse backends such as RainFusion.
+
+A backend reads only the fields it owns and rejects incompatible values. New
+options should be added to a shared typed spec only when more than one backend
+shares their semantics; otherwise add a dedicated typed block.
+
+## Model integration contract
+
+Each diffusion attention call declares a stable role and, when appropriate, a
+broader category. Model-specific roles permit precise overrides; categories
+preserve common `self` or `cross` policies. A model also decides whether its
+path is compatible with automatic TRTLLM selection, because only the model
+knows whether masking and packing satisfy the kernel contract.
+
+## Adding or changing a backend
+
+An implementation change is complete only when it:
+
+1. implements the `AttentionBackend` interface and exposes a stable enum name;
+2. adds platform validation and default routing where appropriate;
+3. defines and validates any typed configuration;
+4. tests explicit selection, default selection, and incompatible paths;
+5. documents installation, fallback behavior, and quality implications in the
+   corresponding user guide; and
+6. updates the overview matrix without moving algorithm details back into the
+   landing page.
@@ -0,0 +1,80 @@
+# CPU Offloading
+
+This document defines the shared architecture for diffusion CPU offloading.
+The three strategies have separate user and design pages:
+
+| Strategy | User guide | Design contract |
+| --- | --- | --- |
+| Model-level | [Guide](../../../user_guide/diffusion/offloader/cpu_offload.md) | [Design](cpu_offload.md) |
+| Layerwise | [Guide](../../../user_guide/diffusion/offloader/layerwise_offload.md) | [Design](layerwise_offload.md) |
+| Distributed layerwise | [Guide](../../../user_guide/diffusion/offloader/distributed_layerwise_offload.md) | [Design](distributed_layerwise_offload.md) |
+
+## Strategy selection
+
+`OffloadConfig.from_od_config()` converts diffusion configuration into one
+`OffloadStrategy`. If callers set multiple flags, selection is deterministic:
+
+1. `DISTRIBUTED_LAYERWISE`;
+2. `LAYERWISE`;
+3. `MODEL_LEVEL`; and
+4. `NONE`.
+
+The factory logs when a higher-priority strategy disables another. It derives
+parallel and HSDP state from `DiffusionParallelConfig`; callers do not provide
+a separate offload group size.
+
+`get_offload_backend()` then validates platform offload support, resolves the
+device, and creates exactly one backend. Returning `None` means offloading is
+disabled or unsupported and must not leave partially installed hooks.
+
+## Shared lifecycle
+
+Every backend implements `OffloadBackend`:
+
+- `enable(pipeline)` discovers modules, establishes initial residency, and
+  installs hooks;
+- `disable()` removes owned hooks and resources; and
+- `is_enabled()` reports lifecycle state.
+
+`disable()` does not promise to restore the pipeline's original device
+placement. The caller owns any subsequent rematerialization.
+
+Hooks are registered through `HookRegistry` and `ModelHook`; offload backends
+must use distinct hook names and remove only hooks they own.
+
+## Discovery and topology
+
+Pipeline component discovery prefers `SupportsComponentDiscovery` declarations:
+
+- `_dit_modules`;
+- `_encoder_modules`;
+- `_vae_modules`; and
+- `_resident_modules`.
+
+Dotted paths are supported. Legacy pipelines may use the fallback scan of
+well-known attribute names, but new integrations should declare components
+explicitly.
+
+Ordinary layerwise offload obtains topology from declared
+`_layerwise_offload_blocks_attrs` (with a deprecated singular-name
+compatibility path). Distributed layerwise offload may instead use a pipeline
+`OffloadPlan`, then falls back to declared block attributes and its documented
+compatibility discovery. Discovery metadata describes structure only; the
+backend remains responsible for transfer, synchronization, and storage
+ownership.
+
+## Cross-strategy invariants
+
+- At most one strategy owns offload hooks for a pipeline.
+- A parameter has one authoritative host representation while offloaded.
+- Device storage is not freed until dependent compute or transfer work is
+  complete.
+- Non-persistent and model-specific buffers remain correct after movement or
+  checkpoint rematerialization.
+- Unsupported parallel or loading combinations fail before hooks mutate the
+  model.
+- Platform streams, events, synchronization, and cache management go through
+  the vLLM-Omni platform abstraction.
+
+The diffusion [Offloader module design](../../module/diffusion/offloader.md)
+describes how these feature contracts fit into the larger diffusion runtime.
@@ -0,0 +1,56 @@
+# Model-Level Offload
+
+Model-level offload enforces mutual exclusion between pipeline component
+groups. For commands and model support, see the
+[model-level user guide](../../../user_guide/diffusion/offloader/cpu_offload.md).
+
+## Generic hook path
+
+`ModelLevelOffloadBackend.enable()` discovers DiTs, encoders, and VAEs. It
+makes encoders and VAEs device resident for initialization, then installs
+`SequentialOffloadHook` instances:
+
+- each DiT hook offloads all encoders and other DiTs before loading itself;
+- each encoder hook offloads every DiT before loading itself; and
+- VAEs are not part of the mutual-exclusion hook set.
+
+The hook moves parameters and buffers without recursively calling
+`module.to()`. This avoids recursion through modules that retain references to
+other transformer blocks. Host tensors are pinned when configured; DTensors
+are not pinned through the ordinary tensor path.
+
+Before forward returns control to the model, the hook synchronizes the active
+platform so compute cannot race the transfer. XPU uses blocking device-to-host
+copies because its allocator cannot safely evict storage while an asynchronous
+copy is pending.
+
+## Split-model delegation
+
+If a pipeline provides callable `enable_omni_model_cpu_offload`, the backend
+delegates instead of installing generic encoder/DiT hooks. This supports models
+whose mutually exclusive components are nested inside one transformer, such
+as a reasoner and generator.
+
+The delegated pipeline owns its internal contexts, but must preserve the same
+contract:
+
+- only the executing component is device resident;
+- enable and disable are idempotent;
+- transfers use the supplied device, pinning, and HSDP settings; and
+- disabling removes all model-local offload hooks or contexts.
+
+The delegation is currently duck typed. If another independent implementation
+is added, formalize the enable/disable pair as a protocol rather than expanding
+attribute-name heuristics.
+
+## Extension invariants
+
+- New ordinary pipelines declare both DiT and encoder groups through
+  `SupportsComponentDiscovery`.
+- Models with multiple DiTs offload the inactive DiTs as well as encoders.
+- Empty or parameterless modules do not trigger transfers.
+- Hook removal does not move modules; callers must not assume original
+  residency is restored.
+
+Shared selection and lifecycle behavior is defined in the
+[CPU Offloading design](README.md).
@@ -7,7 +7,7 @@ either shard those host-side weights across an existing parallel group or keep
 the standard loader's rank-local weights and avoid an additional collective.

 For user-facing commands, see the
-[distributed layerwise offloading guide](../../user_guide/diffusion/cpu_offload.md)
+[distributed layerwise offloading guide](../../../user_guide/diffusion/offloader/distributed_layerwise_offload.md)
 and the [Cosmos3 recipe](https://github.com/vllm-project/vllm-omni/blob/main/recipes/cosmos3/Cosmos3-DistOffload.md).

 ## Status
@@ -56,7 +56,7 @@ AllGather:                [full N+1]             [full N+2]
 Buffers:    [current slot]       [prefetch slot]       [current slot]
 ```

-![DLO double-buffer prefetch pipeline](../figures/dlo/dlo_pipeline.gif)
+![DLO double-buffer prefetch pipeline](../../figures/dlo/dlo_pipeline.gif)

 The backend uses two shared device buffers, so accelerator weight residency is
 bounded by the largest streamed blocks rather than the complete model.
@@ -0,0 +1,57 @@
+# Layerwise Offload
+
+Layerwise offload streams transformer blocks through one accelerator. For
+commands and model support, see the
+[layerwise user guide](../../../user_guide/diffusion/offloader/layerwise_offload.md).
+
+## Hook ring
+
+`LayerWiseOffloadBackend` discovers every streamable block and registers one
+`LayerwiseOffloadHook` per block. Hooks form a ring: each hook owns the pinned
+host representation of its registered block and a reference to the next block
+it must prefetch; the final hook prefetches the first block for the next
+denoising iteration.
+
+During initialization, parameters and buffers are grouped by dtype and copied
+into contiguous pinned CPU tensors. Original tensor storage is replaced with
+an empty or meta placeholder while identity and module structure are retained.
+Metadata records names, offsets, shapes, and dtypes for rematerialization.
+
+## Forward lifecycle
+
+For each block:
+
+1. the previous hook's completion event ensures this block is materialized;
+2. the copy stream asynchronously allocates and fills the next block's device
+   tensors;
+3. parameter and buffer storage is rebound to views of those tensors;
+4. the current block computes on the normal stream; and
+5. its post-forward hook replaces device storage with placeholders after all
+   dependent work is safe.
+
+The copy stream waits on the compute stream before reusing storage, and the
+compute stream waits on the per-block prefetch event before forward. These
+dependencies, not incidental global synchronization, define correctness.
+
+## Topology contract
+
+A DiT declares ordered block-container attributes using
+`_layerwise_offload_blocks_attrs` or an `OffloadPlan`. Each resolved item must
+be an executable `nn.Module`; ordering must match forward execution. Multiple
+containers are concatenated in execution order.
+
+Non-block DiT modules, encoders, VAEs, and declared resident modules are not
+streamed by this backend.
+
+## Invariants and limitations
+
+- Each offloaded tensor has one pinned host backing store.
+- Parameter identity is preserved when storage is rebound.
+- Blocks with heterogeneous shapes and dtypes use their recorded metadata;
+  implementations must not assume identical blocks.
+- The ring is initialized before the first forward so block zero is available.
+- This backend is single-device; distributed weight reconstruction belongs to
+  [Distributed Layerwise Offload](distributed_layerwise_offload.md).
+
+Shared selection and lifecycle behavior is defined in the
+[CPU Offloading design](README.md).
@@ -186,7 +186,7 @@ Distributed layerwise offload has an additional boundary: its sharded mmap
 AllGather path is not compatible with online quantization because the path
 expects already materialized checkpoint weights. Use the rank-local
 `--dlo-no-use-allgather` path or a compatible pre-quantized checkpoint when
-combining these features. See the [DLO design document](distributed_layerwise_offload.md)
+combining these features. See the [DLO design document](offloader/distributed_layerwise_offload.md)
 for the full compatibility matrix.

 ## Adding a quantization backend
@@ -2,8 +2,10 @@

 Skip-Softmax is the sparse-attention mode of the `TRTLLM_ATTN` backend. Usage — the config keys
 and how to pick an operating point — is in
-[Diffusion Attention Backends](../../user_guide/diffusion/attention_backends.md). This page explains
-the algorithm.
+[TRTLLM Attention](../../user_guide/diffusion/attention_backends/trtllm.md#skip-softmax).
+The shared selector contract is documented in
+[Diffusion Attention Backend Selection](attention_backend_selection.md). This
+page explains the algorithm.

 ## Motivation

@@ -1,57 +1,42 @@
-# VAE Parallel
+# VAE Patch Parallelism

-This document defines the VAE parallelism strategy contract and describes how to
-add a strategy to a diffusion model. Tile/patch parallelism is the portable
-baseline. Spatial-shard parallelism is a model-specific decode strategy that can
-be selected explicitly or through a request-aware policy.
-
-We use **Qwen-Image** as the reference implementation for tile-parallel decode,
-**Wan2.2** for tile-parallel encode, and **Wan2.2** as the first spatial-shard
-decode implementation. Other VAE families can add spatial strategies later
-without changing the public mode meanings defined here.
+This document describes how to add **VAE Patch Parallelism** support to a diffusion model.
+We use **Qwen-Image** as the reference implementation for decode parallel, and **Wan2.2** for encode parallel.

 ---

 ## Table of Contents

 - [Overview](#overview)
- [Strategy Selection Contract](#strategy-selection-contract)
- [Tile/Patch Implementation (Decode)](#tilepatch-implementation-decode)
- [Tile/Patch Encode Implementation](#tilepatch-encode-implementation)
+- [Step-by-Step Implementation (Decode)](#step-by-step-implementation-decode)
+- [Encode Parallel Implementation](#encode-parallel-implementation)
 - [Testing](#testing)
 - [Reference Implementations](#reference-implementations)
- [Spatially-Sharded Decode](#spatially-sharded-decode)
- [Adding Another VAE Strategy](#adding-another-vae-strategy)
 - [Summary](#summary)

 ---

 ## Overview

-### What is VAE parallelism?
+### What is Vae Patch parallel?

-VAE parallelism distributes VAE **encoding** or **decoding** across multiple
-ranks. The implementation currently has two strategy families:
+**VAE Patch Parallelism** is an acceleration technique for both **encoding** and **decoding**. Instead of processing the entire tensor at once, the tensor is:

-| Strategy | Unit of work | Communication | Current scope |
-|----------|--------------|---------------|---------------|
-| Tile/patch | Independent overlapping tiles | Gather, stitch, and optionally broadcast | Encode and decode across supported distributed VAEs |
-| Spatial shard | One feature map sharded along a spatial axis | Per-layer halo exchange followed by output gather | Wan decode |
+ Split into multiple spatial tiles

-Both strategies preserve the model-specific receptive field and reconstruct the
-same logical output as single-rank execution. They differ in how work is divided:
-tile parallelism assigns complete overlapping tiles to ranks, while spatial
-sharding keeps one logical feature map partitioned throughout the decoder.
+ Distributed across multiple ranks

-This approach can:
+ Encoded/Decoded in parallel

-+ Distribute computation across multiple devices
-+ Reduce peak memory usage per device
-+ Accelerate encoding or decoding latency
+ Merged to reconstruct the final output

-The actual benefit depends on shape, topology, communication cost, and VAE
-architecture. A strategy must not become the global default without parity and
-performance evidence on its supported models.
+This approach:
+
+ Distributes computation across multiple devices
+
+ Reduces peak memory usage per device
+
+ Accelerates encoding/decoding latency

 ### When to Use Encode vs Decode Parallel

@@ -60,16 +45,8 @@ performance evidence on its supported models.
 | **Decode Parallel** | Text-to-Image, Text-to-Video | Latent → Image/Video |
 | **Encode Parallel** | Image-to-Video (I2V) | Image → Latent (for conditioning) |

-### Architecture and ownership
-
-`DiffusionParallelConfig.vae_parallel_mode` carries the public strategy request.
-The shared registry forwards that value to the VAE, but each distributed VAE
-adapter owns its supported modes, request-shape selection, topology validation,
-and fallback behavior. Model-specific policy must not be added to the shared
-runner or registry.
-
-**DistributedVaeExecutor** is the model-agnostic core for tile-parallel VAE
-encoding and decoding.
+### Architecture
+We introduce **DistributedVaeExecutor** as the core component responsible for distributed VAE encoding/decoding.

 The executor is model-agnostic and accepts three function parameters:

@@ -99,11 +76,6 @@ This design separates:

 + Model-specific tiling and merging logic

-Spatial-shard implementations use the same configured VAE process group but do
-not use the tile task executor. They keep the feature map sharded across decoder
-layers and own their halo-exchange and final-gather contracts inside the model's
-distributed autoencoder implementation.
-
 #### Why split / exec / merge is necessary?

 The latent tensor cannot be arbitrarily partitioned.
@@ -120,34 +92,7 @@ Therefore:

 + Merge must perform blending to avoid seams

-## Strategy Selection Contract
-
-`vae_parallel_mode="tile"` is the compatibility-preserving global default.
-Other values are stable strategy requests, but support remains model-specific.
-
-| Mode | Meaning | Current implementation | Unsupported or ineligible request |
-|------|---------|------------------------|-----------------------------------|
-| `tile` | Use the existing tile/patch executor | All distributed VAE adapters | Not applicable |
-| `auto` | Let the VAE adapter select a validated strategy for each request | Wan decode selects height or width spatial sharding | Wan uses tile when ineligible; unsupported VAE families reject or document their fallback |
-| `spatial_shard_height` | Force spatial-shard decode along height | Wan decode | Wan warns and uses tile when ineligible; unsupported VAE families reject |
-| `spatial_shard_width` | Force spatial-shard decode along width | Wan decode | Wan warns and uses tile when ineligible; unsupported VAE families reject |
-
-The shared configuration layer validates the mode name. It does not decide
-whether a model supports that mode. A VAE adapter adding a mode must:
-
-1. Keep `tile` as an exact feature-off path.
-2. Select a strategy per request rather than from stale process-global state.
-3. Validate input rank/shape and process-group membership before collectives.
-4. Reject or explicitly fall back from unsupported combinations; never silently
-   reinterpret one forced strategy as another.
-5. Preserve direct encode/decode behavior when Diffusers does not select a tiled
-   path.
-
-`auto` is a policy hook, not a universal heuristic. Each VAE family may select
-from only the strategies it has validated. Adding another family must not change
-Wan's selector or the meanings of the explicit modes.
-
-## Tile/Patch Implementation (Decode)
+## Step-by-Step Implementation (Decode)

 ### Step 1: Implement DistributedAutoencoderKLQwenImage
 `QwenImagePipeline` use `AutoencoderKLQwenImage` for vae, so implement a distributed version:
@@ -306,7 +251,7 @@ class YourModelPipeline(nn.Module):
 +       ).to(self.device)
 ```

-## Tile/Patch Encode Implementation
+## Encode Parallel Implementation

 For models that require VAE encoding (e.g., Image-to-Video), you can also parallelize the encode operation. We use **Wan2.2** as the reference implementation.

@@ -504,24 +449,18 @@ Complete examples in the codebase:

 ---

-## Spatially-Sharded Decode
+## Spatially-Sharded Decode (Wan)

-The tile-parallel executor above assigns independent spatial **tiles** to ranks.
-A spatial-shard backend instead partitions one global decoder feature map along
-height or width. Wan is the first supported implementation.
-
-`DiffusionParallelConfig.vae_parallel_mode="tile"` remains the default. For Wan,
-opt into request-shape selection with `"auto"`, or force an axis with
-`"spatial_shard_height"` or `"spatial_shard_width"`.
+The tile-parallel executor above assigns independent spatial **tiles** to ranks. The Wan VAE additionally supports a **spatially-sharded decode** backend, selected through `DiffusionParallelConfig.vae_parallel_mode` (`"spatial_shard_height"` or `"spatial_shard_width"`; default is `"tile"`).

 ### How it differs from tile parallel

-| Aspect | Tile parallel (`"tile"`) | Automatic (`"auto"`) | Spatially-sharded (`"spatial_shard_height"`/`"spatial_shard_width"`) |
-|--------|--------------------------|----------------------|-----------------------------------------------|
-| Unit of work | Independent overlapping tiles | Chooses tile or a spatial shard per request | A single global feature map sharded along H or W |
-| Cross-rank communication | Gather tiles to rank 0, stitch + blend | Depends on the selected path | Per-conv **halo exchange** of boundary rows/cols (P2P) |
-| Output assembly | Blend overlapping tiles | Depends on the selected path | All-gather shards on rank 0, trim padding (matches `broadcast_result=False`) |
-| Scope | Decode + encode | Wan decode selection; encode remains tiled | Wan decode only |
+| Aspect | Tile parallel (`"tile"`) | Spatially-sharded (`"spatial_shard_height"`/`"spatial_shard_width"`) |
+|--------|--------------------------|-----------------------------------------------|
+| Unit of work | Independent overlapping tiles | A single global feature map sharded along H or W |
+| Cross-rank communication | Gather tiles to rank 0, stitch + blend | Per-conv **halo exchange** of boundary rows/cols (P2P) |
+| Output assembly | Blend overlapping tiles | All-gather shards on rank 0, trim padding (matches `broadcast_result=False`) |
+| Scope | Decode + encode | Decode only |

 Spatial-shard decode swaps the decoder's spatial convolutions/padding for halo-exchanging variants (`WanDistConv2d`, `WanDistCausalConv3d`, `WanDistZeroPad2d`) so each rank only holds a shard of the activations but still sees the correct receptive field at shard boundaries. Implementation lives in `vllm_omni/diffusion/distributed/autoencoders/wan_spatial_shard.py`.

@@ -537,65 +476,21 @@ serve.py (--vae-parallel-mode) / OmniEngineArgs
  -> DistributedAutoencoderKLWan.tiled_decode dispatch
 ```

-`DistributedAutoencoderKLWan._spatial_shard_decode_enabled()` gates the path: it
-requires distributed decode to be enabled, a 5D latent, and
-`vae_patch_parallel_size == DiT group size`. The method is reached only after
-Diffusers has selected tiled decode for the request. In `auto` mode the longer
-axis that can cover every rank is selected, with width winning ties. If neither
-axis can cover the group, auto mode uses tile decode. This policy is cached only
-through normal kernel shape caches; it is reevaluated for every request.
+`DistributedAutoencoderKLWan._spatial_shard_decode_enabled()` gates the path: it requires distributed decode to be enabled, a 5D latent, and `vae_patch_parallel_size == DiT group size`. Otherwise it logs a warning and falls back to tile-parallel decode.

 ### Notes

- The decoder is patched **in place** the first time spatial-shard decode runs. Its wrappers consult a context-local split dimension and retain an exact direct fallback, allowing later requests to select tile, height, or width safely.
+- The decoder is patched **in place** the first time spatial-shard decode runs and is bound to a single split dimension for the lifetime of the VAE instance.
 - Numerical correctness vs. single-GPU decode is covered by `tests/diffusion/distributed/test_wan_spatial_shard.py::test_spatial_shard_decode_matches_reference` (multi-GPU, nightly `full_model` + `distributed_cuda`).

-## Adding Another VAE Strategy
-
-Add future strategies through the distributed autoencoder that owns the model
-semantics. Do not add model-name dispatch to the shared runner, registry, or
-`DistributedVaeExecutor`.
-
-For a new VAE family:
-
-1. **Declare capability**: document which modes, operations, shapes, dtypes, and
-   process-group layouts the adapter supports.
-2. **Implement model-local selection**: keep `auto` policy and explicit-mode
-   dispatch in the distributed autoencoder. The selector may use request shape,
-   but must be deterministic and side-effect free.
-3. **Preserve the baseline**: retain exact tile/direct fallbacks and state-dict
-   compatibility. Structural wrappers must not register duplicate parameters or
-   alter checkpoint names.
-4. **Define communication**: derive shards from the configured VAE group, keep
-   collectives symmetric, specify padding/halo ownership, and return the same
-   output placement as the tile path.
-5. **Bound request state**: make the selected axis and scratch buffers
-   request-scoped, release temporary buffers on success and failure, and support
-   alternating modes in one long-lived worker.
-6. **Validate independently**: compare each explicit strategy against the
-   single-rank reference, then test `auto`, tile fallback, unsupported topology,
-   and a mixed request sequence. Record per-rank peak memory and decode latency
-   before documenting a preferred mode.
-
-Promote a selector into shared infrastructure only after at least two VAE
-families use the same inputs, invariants, fallback behavior, and lifecycle. Until
-then, shared code owns transport and configuration while each model adapter owns
-policy.
-
 ---

 ## Summary

-Adding VAE parallel support to a diffusion model:
+Adding VAE Patch Parallel support to diffusion model:

-1. **Implement the distributed VAE** - Inherit from the base VAE class and
-   `DistributedVaeMixin`.
-2. **Keep tile as the baseline** - For tile decode, refactor `tiled_decode` into
-   `tile_split`/`tile_exec`/`tile_merge`.
-3. **Add encode parallelism when needed** - Implement
-   `encode_tile_split`/`encode_tile_exec`/`encode_tile_merge` for I2V models.
-4. **Add model-specific strategies locally** - Own selectors, topology checks,
-   communication, and fallback in the distributed autoencoder.
-5. **Wire the pipeline** - Replace the pipeline VAE with its distributed version.
-6. **Test strategy and lifecycle parity** - Compare one rank with every supported
-   strategy, including feature-off, fallback, and mixed-request sequences.
+1. **Implement Distributed VAE** - Inherit from base VAE class and `DistributedVaeMixin`
+2. **Decode Parallel** - Refactor `tiled_decode` into `tile_split`/`tile_exec`/`tile_merge`
+3. **Encode Parallel** (optional) - Implement `encode_tile_split`/`encode_tile_exec`/`encode_tile_merge` for I2V models
+4. **Change VAE model in pipeline** - Use the distributed version
+5. **Test** - Verify numerical consistency with `vae_patch_parallel_size=1` vs `N`
@@ -48,20 +48,27 @@ implementation contract; it is not, by itself, a general support claim.
 - [Pipeline Parallel](feature/pipeline_parallel.md)
 - [Sequence Parallel](feature/sequence_parallel.md)
 - [Tensor Parallel](feature/tensor_parallel.md)
- [VAE Parallel](feature/vae_parallel.md)
+- [VAE Patch Parallelism](feature/vae_parallel.md)

 #### Attention optimization

 The [Diffusion Attention Backends](../user_guide/diffusion/attention_backends.md)
-guide lists all selectable backends and their platform defaults. Skip-Softmax
-is the backend-related optimization with a standalone feature-design contract:
+guides list selectable backends, platform defaults, installation, and tuning.
+The design contracts separate selection mechanics from backend algorithms:

+- [Attention Backend Selection](feature/attention_backend_selection.md)
 - [Skip-Softmax](feature/skip_softmax.md)

+#### CPU offloading
+
+- [Overview and Shared Contracts](feature/offloader/README.md)
+- [Model-Level Offload](feature/offloader/cpu_offload.md)
+- [Layerwise Offload](feature/offloader/layerwise_offload.md)
+- [Distributed Layerwise Offload](feature/offloader/distributed_layerwise_offload.md)
+
 - [Cache-DiT](feature/cache_dit.md)
 - [TeaCache](feature/teacache.md)
 - [Diffusion Continuous Batching](feature/diffusion_continuous_batching.md)
- [Distributed Layerwise Offload](feature/distributed_layerwise_offload.md)

 ## Infrastructure and Performance

@@ -48,10 +48,10 @@ defines the shared configuration and backend extension points.
 | Goal | User guide | Related design contract |
 | --- | --- | --- |
 | Compare acceleration methods and supported combinations | [Overview](../user_guide/diffusion_features.md), [Feature Compatibility](../user_guide/feature_compatibility.md) | [Diffusion acceleration designs](../design/index.md#diffusion-acceleration) |
-| Move weights between host and device memory | [CPU Offloading](../user_guide/diffusion/cpu_offload.md) | [Distributed Layerwise Offload](../design/feature/distributed_layerwise_offload.md) |
+| Move weights between host and device memory | [CPU Offloading](../user_guide/diffusion/cpu_offload.md) | [CPU Offloading](../design/feature/offloader/README.md) |
 | Reuse denoising computation | [Cache Acceleration](../user_guide/diffusion/cache_acceleration/cache_dit.md) | [Cache-DiT](../design/feature/cache_dit.md), [TeaCache](../design/feature/teacache.md) |
 | Distribute diffusion work across devices | [Parallelism](../user_guide/diffusion/parallelism/overview.md) | [Parallelism designs](../design/index.md#parallelism) |
-| Select dense, sparse, or quantized attention paths | [Attention Backends](../user_guide/diffusion/attention_backends.md) | [Skip-Softmax](../design/feature/skip_softmax.md) |
+| Select dense, sparse, or quantized attention paths | [Attention Backends](../user_guide/diffusion/attention_backends.md) | [Attention Backend Selection](../design/feature/attention_backend_selection.md) |
 | Compile repeated diffusion regions | [Regional Compilation](../user_guide/diffusion/regional_compilation.md) | User-facing optimization guide |
 | Add generated video frames | [Frame Interpolation](../user_guide/diffusion/frame_interpolation.md) | User-facing extension guide |
 | Reduce diffusion model startup time | [Startup and Loading](../user_guide/diffusion/startup_and_loading.md) | User-facing loading guide |
@@ -42,7 +42,8 @@ th {
 | `Wan22VACEPipeline` | Wan2.1-VACE | `Wan-AI/Wan2.1-VACE-1.3B-diffusers`, `Wan-AI/Wan2.1-VACE-14B-diffusers` | ✅︎ | ✅︎ | ✅︎ | ✅︎ | — |
 | `Wan22VACEPipeline` | Wan2.2-VACE | `Pyros13/Wan2.2-VACE-Fun-A14B-Diffusers` | ✅︎ |   |   |   | — |
 | `LTX2Pipeline` | LTX-2 / LTX-2.3 one-stage T2V and I2V | `Lightricks/LTX-2`, `diffusers/LTX-2.3-Diffusers` | ✅︎ | ✅︎ | | | — |
-| `LTX2DistilledPipeline` | LTX-2 distilled two-stage T2V and I2V | `rootonchair/LTX-2-19b-distilled` | ✅︎ | ✅︎ | | | — |
+| `LTX2TwoStagePipeline` | LTX-2 / LTX-2.3 ordinary two-stage T2V and I2V | `Lightricks/LTX-2`, `diffusers/LTX-2.3-Diffusers` + matching Lightricks LoRA and upsampler | ✅︎ | ✅︎ | | | — |
+| `LTX2DistilledPipeline` | LTX-2 / LTX-2.3 full-distilled two-stage T2V and I2V | `rootonchair/LTX-2-19b-distilled`, `diffusers/LTX-2.3-Distilled-Diffusers` + `Lightricks/LTX-2.3` upsampler | ✅︎ | ✅︎ | | | — |
 | `LingBotVideoPipeline` | LingBot-Video dense and MoE T2I, T2V, TI2V | `robbyant/lingbot-video-dense-1.3b`, `robbyant/lingbot-video-moe-30b-a3b` | ✅︎ | | | | — |
 | `MiniMaxH3Pipeline` | MiniMax H3 T2VA, first/last-keyframe FL2VA, and mixed-reference Ref2VA | `MiniMaxAI/MiniMax-H3` | ✅︎ | ✅︎ | | | [Published](https://recipes.vllm.ai/MiniMaxAI/MiniMax-H3) |
 | `DreamZeroPipeline` | DreamZero-DROID | `GEAR-Dreams/DreamZero-DROID` | ✅︎ | ✅︎ | ✅︎ | ✅︎ | — |
@@ -58,6 +59,7 @@ th {
 | `CosyVoice3Model` | CosyVoice3 | `FunAudioLLM/Fun-CosyVoice3-0.5B-2512` | ✅︎ | ✅︎ | | ✅︎ | — |
 | `OmniVoiceModel` | OmniVoice | `k2-fsa/OmniVoice` | ✅︎ | | | | — |
 | `VoxCPM2TalkerForConditionalGeneration` | VoxCPM2 | `openbmb/VoxCPM2` | ✅︎ | | | | — |
+| `DotsTTSForConditionalGeneration` | dots.tts | `rednote-hilab/dots.tts-soar` | ✅︎ | | | | — |
 | `MammothModa2ForConditionalGeneration` | MammothModa2-Preview | `bytedance-research/MammothModa2-Preview` | ✅︎ | ✅︎ | | | — |
 | `MammothModa2ForConditionalGeneration` | MammothModa2-Dev (AR-only image understanding) | `bytedance-research/MammothModa2-Dev` | ✅︎ | | | | — |
 | `Flux2KleinPipeline` | FLUX.2-klein | `black-forest-labs/FLUX.2-klein-4B`, `black-forest-labs/FLUX.2-klein-9B` | ✅︎ | ✅︎ | ✅︎ | ✅︎ | — |
@@ -1,494 +1,192 @@
 # Diffusion Attention Backends

-This document describes the diffusion attention backends available in vLLM-Omni, how to select them globally and per-role, the per-platform defaults, and how to use SageAttention.
+Use this page to select and configure a diffusion attention backend. Backend-
+specific installation, tuning, and compatibility details live in separate
+guides so that the selection contract stays easy to scan.

-## Overview
+Diffusion attention backend selection applies to DiT attention in image and
+video generation models. It does **not** change autoregressive LLM attention,
+which uses vLLM's attention selector.

-Diffusion attention backend selection is resolved in `vllm_omni.diffusion.attention.selector`. It looks up the backend from a structured `AttentionConfig` carried on `OmniDiffusionConfig` and falls back to the platform default when nothing is configured.
+For the internal selector, registry, and platform contract, see
+[Attention Backend Selection](../../design/feature/attention_backend_selection.md).

-This backend is used by diffusion attention layers such as the DiT attention in video and image generation models. It does **not** affect autoregressive (LLM) attention paths — those go through vLLM's own attention backend selector.
+## Choose a guide

-The full set of backends and their platform defaults is in the **Backend Options** and **Platform Defaults** sections below. If no attention backend is configured, vLLM-Omni asks the current platform to choose the default.
+| Need | Guide |
+| --- | --- |
+| Select a conservative or platform-native dense kernel | [Dense Backends](attention_backends/dense_backends.md) |
+| Use FlashInfer trtllm-gen FMHA, Skip-Softmax, or TRTLLM SAGE quantization | [TRTLLM Attention](attention_backends/trtllm.md) |
+| Install and use SageAttention 2.2 or SageAttention3 | [SageAttention](attention_backends/sage.md) |
+| Match training or rollout kernels loaded from Hugging Face | [Hugging Face Hub Backends](attention_backends/huggingface_hub.md) |
+| Use block-sparse video attention on Ascend NPU | [RainFusion](attention_backends/rainfusion.md) |

-## Backend Options
-
-| Value | Notes |
-|---|---|
-| `TRTLLM_ATTN` | FlashInfer's trtllm-gen FMHA (TensorRT-LLM's generated kernels, vendored by FlashInfer). Dense BF16, GQA native, `head_dim=128`. Datacenter Blackwell only (sm_100 / sm_103). Packed paths can provide `cu_seqlens` directly. Supports optional **Skip-Softmax** sparse attention — see below. Requires `flashinfer`. |
-| `FLASH_ATTN` | Wraps FlashAttention 4 on Blackwell when `flash-attn-4` is installed, then falls back to FlashAttention 3/2. Default on Hopper / Ada / Ampere when a compatible FlashAttention package is installed. |
-| `CUDNN_ATTN` | Pins `sdpa_kernel([CUDNN_ATTENTION])`. Default on Blackwell (sm_10x / sm_12x) with cuDNN ≥ 9.5. Wins on mask-heavy DiTs (HunyuanVideo-1.5: 2× e2e vs SDPA). |
-| `FLASHINFER_ATTN` | Uses FlashInfer's batch-prefill wrapper and supports mixed Q/K and V input dtypes through backend-specific configuration. Used as Blackwell fallback when cuDNN is unavailable. Requires `flashinfer` >= 0.6.16rc1 for mixed-dtype configurations. |
-| `TORCH_SDPA` | PyTorch `scaled_dot_product_attention` with the default backend dispatcher. Most conservative; always available. |
-| `SAGE_ATTN` | SageAttention 2.2 — INT8-quantized attention with FP16 accumulation. Lossy but typically visually indistinguishable on diffusion outputs. Requires `sageattention`. |
-| `SAGE_ATTN_3` | Requires `sageattn3` from `SageAttention/sageattention3_blackwell`. CUDA only, intended for Blackwell GPUs, with GQA/MQA requests falling back to PyTorch SDPA. |
-| `FLASH_ATTN_HUB` | FlashAttention 2 from HuggingFace `kernels` library. Useful for train/rollout alignment. |
-| `FLASH_ATTN_3_HUB` | FlashAttention 3 from HuggingFace `kernels` library. CUDA Hopper (sm_90+) only; falls back to `FLASH_ATTN_HUB` on older GPUs. |
-| `RAINFUSION_ATTN` | MindIE-SD **RainFusion** block-sparse video attention — see [below](#rainfusion_attn-backend-and-block-sparse-video-attention). Ascend NPU only; requires `mindiesd`. Delegates to `FLASH_ATTN` for anything that is not a packed video sequence. |
+## Backend options

+| Value | Family | Primary use | Detail |
+| --- | --- | --- | --- |
+| `TORCH_SDPA` | Dense | Conservative reference; always available | [Dense Backends](attention_backends/dense_backends.md#torch_sdpa) |
+| `FLASH_ATTN` | Dense | FlashAttention 4/3/2 depending on the installed package and GPU | [Dense Backends](attention_backends/dense_backends.md#flash_attn) |
+| `CUDNN_ATTN` | Dense | Mask-heavy DiTs on Blackwell with cuDNN 9.5 or newer | [Dense Backends](attention_backends/dense_backends.md#cudnn_attn) |
+| `FLASHINFER_ATTN` | Dense or quantized | FlashInfer batch prefill; optional mixed Q/K and V dtypes | [Dense Backends](attention_backends/dense_backends.md#flashinfer_attn) |
+| `TRTLLM_ATTN` | Dense, sparse, or quantized | Datacenter Blackwell with `head_dim=128` and compatible packed paths | [TRTLLM Attention](attention_backends/trtllm.md) |
+| `SAGE_ATTN` | Quantized | SageAttention 2.2 INT8 attention | [SageAttention](attention_backends/sage.md#sage_attn) |
+| `SAGE_ATTN_3` | Quantized | SageAttention3 on Blackwell | [SageAttention](attention_backends/sage.md#sage_attn_3) |
+| `FLASH_ATTN_HUB` | Hub kernel | FlashAttention 2 from Hugging Face `kernels` | [Hugging Face Hub Backends](attention_backends/huggingface_hub.md) |
+| `FLASH_ATTN_3_HUB` | Hub kernel | FlashAttention 3 from Hugging Face `kernels` on Hopper or newer | [Hugging Face Hub Backends](attention_backends/huggingface_hub.md) |
+| `RAINFUSION_ATTN` | Block sparse | MindIE-SD RainFusion video attention on Ascend NPU | [RainFusion](attention_backends/rainfusion.md) |

 ## Configuration

-Diffusion attention backends can be configured three ways, in priority order:
+Backend selection is resolved in this order:

-1. **`--diffusion-attention-config`** — structured per-role config (highest priority).
-2. **`--diffusion-attention-backend` / `DIFFUSION_ATTENTION_BACKEND` env var** — global shorthand that sets the default backend.
-3. **Platform default** — used when nothing is configured.
+1. `--diffusion-attention-config` per-role configuration.
+2. `--diffusion-attention-backend` or `DIFFUSION_ATTENTION_BACKEND` as a
+   global default.
+3. The current platform's default.

-`--diffusion-attention-backend` is shorthand for `--diffusion-attention-config.default.backend`. It may be combined with `--diffusion-attention-config.per_role.*` overrides, but is mutually exclusive with `--diffusion-attention-config.default.backend`.
+`--diffusion-attention-backend` is shorthand for
+`--diffusion-attention-config.default.backend`. Do not pass it together with an
+explicit `default.backend` in the structured configuration.

 ### Global default

-Set the default backend for every diffusion attention layer:
-
 ```bash
-# CLI flag
-vllm-omni serve <model> --diffusion-attention-backend SAGE_ATTN
+vllm-omni serve <model> --diffusion-attention-backend FLASH_ATTN

-# Environment variable (also recognized for backwards compatibility)
-export DIFFUSION_ATTENTION_BACKEND=SAGE_ATTN
+# Backwards-compatible environment variable
+export DIFFUSION_ATTENTION_BACKEND=FLASH_ATTN
 ```

 ### Per-role configuration

-Roles are free-form strings declared by each diffusion model. The two common categories are `"self"` and `"cross"`; model-specific roles (e.g. `"ltx2.audio_to_video"`) may also be declared. A role string is matched in this order:
+Roles are declared by each diffusion model. Common categories are `self` and
+`cross`; a model may also use a more specific role such as
+`ltx2.audio_to_video`. Resolution order is:

-1. Exact `per_role[role]` match
-2. `per_role[role_category]` fallback (e.g. `"ltx2.audio_to_video"` → `"cross"`)
-3. `default`
-4. Platform default
-
-Use vLLM-style dotted flags or one JSON blob:
+1. Exact `per_role[role]` match.
+2. Category `per_role[role_category]` match.
+3. `default`.
+4. Platform default.

 ```bash
 # Dotted flags
 vllm-omni serve <model> \
-    --diffusion-attention-config.default.backend FLASH_ATTN \
-    --diffusion-attention-config.per_role.cross.backend TORCH_SDPA
+  --diffusion-attention-config.default.backend FLASH_ATTN \
+  --diffusion-attention-config.per_role.cross.backend TORCH_SDPA

-# JSON
+# Equivalent JSON
 vllm-omni serve <model> \
-    --diffusion-attention-config '{"default":{"backend":"FLASH_ATTN"},"per_role":{"cross":{"backend":"TORCH_SDPA"}}}'
+  --diffusion-attention-config \
+  '{"default":{"backend":"FLASH_ATTN"},"per_role":{"cross":{"backend":"TORCH_SDPA"}}}'
 ```

-A backend that needs configuration exposes it as a typed field on the spec:
-`TRTLLM_ATTN`'s Skip-Softmax (see [below](#trtllm_attn-backend-and-skip-softmax)) and
-`RAINFUSION_ATTN`'s RainFusion block (see [below](#rainfusion_attn-backend-and-block-sparse-video-attention)):
-
-```bash
--diffusion-attention-config.default.backend TRTLLM_ATTN \
--diffusion-attention-config.default.skip_softmax.target_sparsity 0.5
-```
-
-### Programmatic API
-
-When constructing `OmniDiffusionConfig` directly:
+### Python API

 ```python
-from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, AttnQuantSpec, OmniDiffusionConfig
+from vllm_omni.diffusion.data import (
+    AttentionConfig,
+    AttentionSpec,
+    OmniDiffusionConfig,
+)

 config = OmniDiffusionConfig(
    diffusion_attention_config=AttentionConfig(
        default=AttentionSpec(backend="FLASH_ATTN"),
-        per_role={
-            "cross": AttentionSpec(backend="TORCH_SDPA"),
-        },
+        per_role={"cross": AttentionSpec(backend="TORCH_SDPA")},
    ),
    ...,
 )
 ```

-A plain dict is also accepted and normalized to `AttentionConfig`.
+Backend-specific typed blocks are documented with their consumers:

-#### FlashInfer QK16/V8 quantized attention
+- `quant`: [FlashInfer](attention_backends/dense_backends.md#flashinfer-quantized-attention)
+  and [TRTLLM SAGE](attention_backends/trtllm.md#sage-quantization).
+- `skip_softmax`: [TRTLLM Skip-Softmax](attention_backends/trtllm.md#skip-softmax).
+- `block_sparse`: [RainFusion](attention_backends/rainfusion.md#configuration).

-`FLASHINFER_ATTN` accepts quantization options through `AttentionSpec.quant`. For QK16/V8:
-
-```python
-config = OmniDiffusionConfig(
-    diffusion_attention_config=AttentionConfig(
-        default=AttentionSpec(
-            backend="FLASHINFER_ATTN",
-            quant=AttnQuantSpec(
-                dtype_qk="bfloat16",
-                dtype_vo="fp8_e4m3",
-            ),
-        ),
-    ),
-    ...,
-)
-```
-
-`dtype_qk` configures the Q and K input dtypes. `dtype_vo` configures the V input dtype.
-
-## Platform Defaults
+## Platform defaults

 ### Blackwell (sm_100 / sm_103 / sm_120 / sm_121)

-Auto-route preference, in order:
+The CUDA auto-route preference is:

-1. `TRTLLM_ATTN` — on **datacenter** Blackwell (sm_100 / sm_103) when `flashinfer` is installed, the model's `head_dim` is 128, **and the model declares a compatible packed/mask-free path**
-2. `CUDNN_ATTN` — when cuDNN ≥ 9.5 is available (ships in PyTorch 2.5+ wheels)
-3. `FLASHINFER_ATTN` — when `flashinfer` is installed but cuDNN < 9.5
-4. `FLASH_ATTN` — when `flash-attn` is installed with the Blackwell CUTE kernel
-5. `TORCH_SDPA` — last resort
+1. `TRTLLM_ATTN` on datacenter Blackwell (sm_100/sm_103) when FlashInfer is
+   available, `head_dim=128`, and the model declares a compatible packed or
+   mask-free path.
+2. `CUDNN_ATTN` when cuDNN 9.5 or newer is available.
+3. `FLASHINFER_ATTN` when FlashInfer is available but cuDNN is too old.
+4. `FLASH_ATTN` when a compatible package is installed.
+5. `TORCH_SDPA`.

-`TRTLLM_ATTN` is skipped on workstation Blackwell (sm_120 / sm_121) and for any `head_dim != 128`, so those GPUs keep the `CUDNN_ATTN` route described below.
+`TRTLLM_ATTN` is not auto-selected on workstation Blackwell
+(sm_120/sm_121), for other head dimensions, or for paths that require an
+unsupported mask.

-`TRTLLM_ATTN` outranks `CUDNN_ATTN` on datacenter Blackwell for compatible packed/mask-free pipelines. Workstation Blackwell (sm_120 / sm_121) and pipelines that require attention masks retain their normal fallback.
+### Hopper, Ada, and Ampere

-The startup log line `Defaulting to diffusion attention backend CUDNN_ATTN (Blackwell sm_120, cuDNN 91002)` confirms the route.
+The CUDA auto-route uses `FLASH_ATTN` when available and otherwise falls back
+to `TORCH_SDPA`. `CUDNN_ATTN` and `FLASHINFER_ATTN` remain explicit options.

-**Why CUDNN_ATTN by default**: on mask-heavy diffusion models (HunyuanVideo-1.5, Qwen-Image), cuDNN's pinned FMHA kernel sidesteps a PyTorch SDPA dispatch quirk where the unpinned dispatcher picks `EFFICIENT_ATTENTION` (~25 ms) for masked calls instead of cuDNN (~11 ms). The pin gives 2× e2e on HV-1.5 with no regression on lighter models.
+Other platforms validate an explicit backend and choose their own default
+through the platform implementation. Check the startup log to confirm the
+resolved backend.

-### Hopper (sm_90) / Ada (sm_89) / Ampere (sm_80–sm_86)
+## Choosing a backend manually

-Auto-route preference:
+Override the platform default when you need:

-1. `FLASH_ATTN` — when `flash-attn` is installed
-2. `TORCH_SDPA` — fallback
+- a correctness reference (`TORCH_SDPA`);
+- a backend-specific workaround;
+- training/rollout kernel alignment (Hub backends); or
+- an explicitly validated sparse or quantized speedup.

-`CUDNN_ATTN` and `FLASHINFER_ATTN` are still selectable via env var on these GPUs but are not in the auto-route — FlashAttention 2 is the well-tuned path on pre-Blackwell hardware.
+The startup log prints the resolved backend and whether it came from explicit
+configuration or platform defaulting. If no resolution message appears, check
+earlier logs for diffusion-stage initialization failures.

-## TRTLLM_ATTN Backend and Skip-Softmax
+## Reference benchmark

-`TRTLLM_ATTN` runs FlashInfer's trtllm-gen FMHA. Dense it is a BF16 backend like the others; on top of
-that it can enable **Skip-Softmax**, a sparse-attention mode that trades a little fidelity for speed
-(algorithm: [design doc](../../design/feature/skip_softmax.md)).
+The following BF16 results were measured on an sm_120 RTX Pro 6000 Blackwell
+with the same prompt and seed across runs. Treat them as reference results, not
+portable guarantees.

-Enable it through the typed `skip_softmax` block on the attention spec:
-
-| Key | Valid values | Meaning |
-|---|---|---|
-| `target_sparsity` | finite, `[0, 1]` | Operating point on a calibrated curve. Needs a calibration for the model. |
-| `threshold` | finite, `≥ 0` | Direct threshold, no calibration needed. Mutually exclusive with `target_sparsity`. |
-| `disabled_until_timestep` | finite, `[0, 1]` | Holds early, high-noise steps dense; skip turns on once the normalized timestep `t ≤ D`. |
-
-```bash
-vllm-omni serve Wan-AI/Wan2.2-T2V-A14B-Diffusers \
-  --diffusion-attention-config '{"default": {"backend": "TRTLLM_ATTN",
-      "skip_softmax": {"target_sparsity": 0.65, "disabled_until_timestep": 0.86}}}'
-```
-
-Programmatically the same block is a typed `SkipSoftmaxSpec` (values validated at construction):
-
-```python
-from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, SkipSoftmaxSpec
-
-AttentionConfig(
-    default=AttentionSpec(
-        backend="TRTLLM_ATTN",
-        skip_softmax=SkipSoftmaxSpec(target_sparsity=0.65, disabled_until_timestep=0.86),
-    ),
-)
-```
-
-**Start at `target_sparsity=0.65, disabled_until_timestep=0.86`** and raise `target_sparsity` only
-as far as your quality bar allows — the output diverges from dense as it climbs. The gain grows with
-sequence length, since Skip-Softmax only accelerates attention. Dense BF16 is `TRTLLM_ATTN` with no
-`skip_softmax` block.
-
-Requires datacenter Blackwell with `head_dim == 128`; elsewhere, or without FlashInfer, selecting
-`TRTLLM_ATTN` raises rather than silently degrading.
-
-## TRTLLM_ATTN SAGE Quantization
-
-On top of dense/Skip-Softmax, `TRTLLM_ATTN` can run **SAGE** (SageAttention) — the trtllm-gen
-FMHA kernel with per-block quantized Q/K and per-channel fp8 V, so the two attention matmuls run
-in low precision. Enable it through the typed `quant` block on the attention spec:
-
-| Key | Valid values | Meaning |
-|---|---|---|
-| `dtype_qk` | `int8`, `fp8_e4m3` | Quantization dtype for Q and K. Absent ⇒ dense (SAGE off). |
-| `q_block_size` | `1`, `4`, `16` | Per-token block size for Q scales (default `1`). |
-| `k_block_size` | `1`, `4`, `16` | Per-token block size for K scales (default `16`). |
-
-V is always quantized per-channel to fp8_e4m3 and K-smoothing is applied inside the routine, so
-neither is a user knob. Only block sizes `1`, `4`, `16` have compiled kernels.
-
-The `quant` block is a shared `AttnQuantSpec` consumed by more than one backend; each reads the
-fields that apply and rejects the rest. `TRTLLM_ATTN` (SAGE) reads `dtype_qk` (`int8` / `fp8_e4m3`),
-`q_block_size`, and `k_block_size`; `FLASHINFER_ATTN` reads `dtype_qk` (`float16` / `bfloat16`),
-`dtype_vo`, and `flashinfer_backend` (see its section). A `dtype_qk` from the wrong set raises.
-
-```bash
-vllm-omni serve Wan-AI/Wan2.2-T2V-A14B-Diffusers \
-  --diffusion-attention-config '{"default": {"backend": "TRTLLM_ATTN",
-      "quant": {"dtype_qk": "fp8_e4m3", "q_block_size": 1, "k_block_size": 16}}}'
-```
-
-```python
-from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, AttnQuantSpec
-
-AttentionConfig(
-    default=AttentionSpec(
-        backend="TRTLLM_ATTN",
-        quant=AttnQuantSpec(dtype_qk="fp8_e4m3", q_block_size=1, k_block_size=16),
-    ),
-)
-```
-
-**Requirements.** Needs FlashInfer ≥ 0.6.16rc1. Kernel
-availability is arch-dependent: `fp8_e4m3` QK has kernels on both **SM100** (B200) and **SM103**
-(B300); `int8` QK kernels are compiled for **SM100 only**.
-
-## RAINFUSION_ATTN Backend and Block-Sparse Video Attention
-
-`RAINFUSION_ATTN` runs MindIE-SD's RainFusion (`rf_v2`) kernel on Ascend NPU. It pools each
-128-token block of keys, ranks them per query block, and computes attention against only the
-top-scoring ones. Video tokens are rearranged into `(t, h, w)` order first, so the kept blocks are
-spatiotemporal neighbours rather than arbitrary rows of the packed sequence.
-
-Only the video segment is sparsified. The prefix rows (text, visual conditions, audio) and the
-first-frame blocks are always kept dense, which is why the realized sparsity is lower than the
-nominal `sparsity` you configure. Anything the kernel cannot handle — a warmup denoise step, an
-exempt layer, a sequence with no published video geometry, a video segment under 32 blocks —
-delegates to `FLASH_ATTN`, so a model can select this backend unconditionally.
-
-Enable it through the typed `block_sparse` block on the attention spec, shared by every
-block-sparse backend:
-
-| Key | Valid values | Meaning |
-|---|---|---|
-| `sparsity` | finite, `[0, 1]` | Nominal fraction of key blocks dropped per query block. `0` disables sparsity. Defaults to `0.8`. |
-| `start_step` | `≥ 0` | Keep the first N denoise steps dense. Layout is decided early, so these steps dominate structural fidelity. |
-| `skip_layers` | index selector, e.g. `"0-3,38"` | DiT blocks that always stay dense. |
-
-```bash
-vllm-omni serve MiniMaxAI/MiniMax-H3 \
-  --diffusion-attention-config '{"default": {"backend": "RAINFUSION_ATTN",
-      "block_sparse": {"sparsity": 0.8, "start_step": 0}}}'
-```
-
-Programmatically the same block is a typed `BlockSparseSpec` (values validated at construction):
-
-```python
-from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, BlockSparseSpec
-
-AttentionConfig(
-    default=AttentionSpec(
-        backend="RAINFUSION_ATTN",
-        block_sparse=BlockSparseSpec(sparsity=0.8, start_step=0, skip_layers="0-1"),
-    ),
-)
-```
-
-**Tune in the order `start_step` → `sparsity` → `skip_layers`.** Raise `start_step` first: it is the
-cheapest way to recover structure, because the early high-noise steps decide global layout while the
-later steps only refine texture. Then walk `sparsity` up to your quality bar. Reach for
-`skip_layers` last, once an A/B against dense at the same seed points at specific blocks.
-
-Requires Ascend NPU with `mindiesd`; selecting it on any other platform raises. It is also
-incompatible with ring sequence parallelism, since `rf_v2` needs the whole key sequence to rank
-blocks — use Ulysses SP (`ring_degree=1`).
-
-### Which geometries run sparse
-
-MindIE-SD rf_v2 handles arbitrary video grids by spatially rearranging video first and, when
-necessary, promoting a real-video suffix to its always-kept prefix segment. It never pads: rf_v2
-does not consume an attention mask for padded keys. The remaining video rows stay sparse, while the
-promoted rows remain visible to every query and key.
-
-For example, MiniMax-H3 at 1344x768 has grid `(62, 24, 42)`: its 62496 video rows leave a 32-row
-128-block residue, and its two non-8-wide latent columns form a 2928-row rearranged suffix.
-MindIE-SD promotes 2976 rows, leaving 59520 sparse rows (`465 × 128`). This keeps the block mask
-and kernel tiling aligned, at the cost of retaining more blocks.
-
-8x8-aligned spatial grids remain preferable for performance. A latent `h` or `w` not divisible by 8
-creates an always-kept residual suffix and lowers realized sparsity. "Multiple of 8" on the latent
-grid means **width and height that are multiples of 256**. This path requires a MindIE-SD release
-that implements protected video tails for rf_v2.
-
-## End-to-End Benchmark (BF16, sm_120 RTX Pro 6000 Blackwell)
-
-Same prompt and seed across runs. `Total generation time` from `text_to_video.py` / `text_to_image.py`.
-
-| Model | Shape | TORCH_SDPA | CUDNN_ATTN | FLASHINFER_ATTN |
-|---|---|---|---|---|
+| Model | Shape | `TORCH_SDPA` | `CUDNN_ATTN` | `FLASHINFER_ATTN` |
+| --- | --- | ---: | ---: | ---: |
 | HunyuanVideo-1.5 (T2V) | 480p / 33f / 50 steps | 147.05 s | **73.02 s** | 127.84 s |
 | Wan 2.2 14B (T2V) | 480p / 33f / 40 steps | 117.75 s | 117.17 s | **115.07 s** |
 | Qwen-Image (T2I) | 1024² / 50 steps | 17.41 s | **15.14 s** | 16.02 s |
 | FLUX.2-dev (T2I) | 1024² / 50 steps, TP=2 | 53.62 s | **53.30 s** | 54.94 s |

-Pattern: mask-heavy DiTs (HV-1.5, Qwen-Image) favor `CUDNN_ATTN`; lighter-mask DiTs and TP-saturated configs (Wan 2.2, FLUX.2 TP=2) tie within noise.
+Mask-heavy DiTs favored `CUDNN_ATTN`; lighter-mask or TP-saturated workloads
+were close enough that users should benchmark their exact model and shape.

-## Known Limitations
+## Compatibility anchors

-### LTX-2.0: `CUDNN_ATTN` crashes under torch.compile
+The following headings preserve links to sections that moved into dedicated
+guides.

-LTX-2's audio attention has a symbolic head_dim under torch.compile tracing. cuDNN's SDPA backend selector rejects symbolic dims and Dynamo aborts compilation. Tracked in [#3121](https://github.com/vllm-project/vllm-omni/issues/3121).
+## TRTLLM_ATTN Backend and Skip-Softmax

-**Workaround**: explicitly select `FLASHINFER_ATTN` or `TORCH_SDPA` for LTX-2.0:
+See [TRTLLM Attention](attention_backends/trtllm.md#skip-softmax).

-```bash
-DIFFUSION_ATTENTION_BACKEND=FLASHINFER_ATTN python examples/offline_inference/text_to_video/text_to_video.py \
-    --model Lightricks/LTX-2 ...
-```
+## TRTLLM_ATTN SAGE Quantization

-### FlashAttention-4 on Blackwell
+See [TRTLLM Attention](attention_backends/trtllm.md#sage-quantization).

-Install the optional CUDA 13 extra to use the CuTe-based FA4 path:
+## RAINFUSION_ATTN Backend and Block-Sparse Video Attention

-```bash
-pip install 'vllm-omni[fa4]'
-```
-
-`FLASH_ATTN` prefers `flash_attn.cute` on Blackwell and falls back to FA3/FA2
-when it is unavailable. Version `4.0.0b18` is required; earlier beta
-wheels had known JIT failures on Blackwell.
-
-## Choosing a Backend Manually
-
-### When to override the default
-
- **Quality validation**: compare a new backend against `TORCH_SDPA` as the reference, since SDPA's default dispatcher is the most extensively tested.
- **Lossy speedup hunting**: try `SAGE_ATTN` (INT8 quantized) on diffusion outputs — typically indistinguishable visually but always validate.
- **Workaround for known issues**: see Known Limitations above.
-
-### Verifying which backend is in use
-
-The startup log prints one of:
-
-```
-Using diffusion attention backend 'CUDNN_ATTN'           # explicit override
-Defaulting to diffusion attention backend CUDNN_ATTN ... # auto-route
-Defaulting to diffusion attention backend SDPA           # nothing else available
-```
-
-If you don't see one of these, the model didn't reach diffusion stage init — check earlier logs for failures.
+See [RainFusion](attention_backends/rainfusion.md).

 ## SageAttention Installation

-vLLM-Omni expects SageAttention to be installed into the same Python environment as vLLM-Omni.
-
-Build from source:
-
-```bash
-git clone https://github.com/thu-ml/SageAttention.git
-cd SageAttention
-
-export EXT_PARALLEL=4 NVCC_APPEND_FLAGS="--threads 8" MAX_JOBS=32
-pip install . --no-build-isolation
-```
-
-Quick check:
-
-```bash
-python -c "import sageattention; print(sageattention.__file__)"
-```
+See [SageAttention](attention_backends/sage.md#installation).

 ## SageAttention3 Installation

-vLLM-Omni expects SageAttention3 to be installed into the same Python environment as vLLM-Omni.
-
-Build from source:
-
-```bash
-git clone https://github.com/thu-ml/SageAttention.git
-cd SageAttention/sageattention3_blackwell
-python setup.py install
-```
-
-Quick check:
-
-```bash
-python -c "import sageattn3; print(sageattn3.__file__)"
-```
-
-Notes:
-
- `SAGE_ATTN_3` is only selected on CUDA when `sageattn3` is importable and the GPU is Blackwell-class.
- SageAttention3's Blackwell kernel assumes `Hq == Hkv`. In vLLM-Omni, GQA/MQA diffusion requests fall back to PyTorch SDPA for correctness.
+See [SageAttention](attention_backends/sage.md#sageattention3-installation).

 ## HuggingFace Kernels Hub Backends

-To achieve perfect numerical consistency between **training** (typically using HuggingFace Diffusers with Hub kernels) and **serving/rollout** (using vLLM-Omni), you can use the Hub-based attention backends. This eliminates numerical drift / sampling divergence caused by executing different local kernel versions during rollout.
-
-The following backend options are supported:
- `FLASH_ATTN_HUB` (HuggingFace `kernels-community/flash-attn2`)
- `FLASH_ATTN_3_HUB` (HuggingFace `kernels-community/flash-attn3`, Hopper sm_90+ only)
-
-### Installation
-
-To use these backends, you must install the `kernels` library:
-
-```bash
-pip install kernels==0.14.1
-```
-
-If the `kernels` library is not available in the environment, vLLM-Omni will log a warning and fall back gracefully to the corresponding local backend implementations (`FLASH_ATTN`). On CUDA GPUs below Hopper (compute capability < 9.0), `FLASH_ATTN_3_HUB` falls back to `FLASH_ATTN_HUB`.
-
-### Usage
-
-Select a Hub backend using the global CLI flag or environment variables:
-
-```bash
-# Environment variable
-export DIFFUSION_ATTENTION_BACKEND=FLASH_ATTN_HUB
-
-# CLI flag
-vllm-omni serve <model> --diffusion-attention-backend FLASH_ATTN_HUB
-```
-
-## Usage Examples
-
-### Default (auto-route)
-
-```bash
-python examples/offline_inference/text_to_video/text_to_video.py \
-    --model hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v \
-    --prompt "A dog running across a field of golden wheat." \
-    --height 480 --width 832 --num-frames 33 \
-    --num-inference-steps 50 --seed 42 --guidance-scale 6.0 \
-    --output hv15.mp4
-```
-
-On Blackwell this picks `CUDNN_ATTN` automatically. Check the log for the `Defaulting to ...` line.
-
-### Explicit backend selection
-
-```bash
-DIFFUSION_ATTENTION_BACKEND=FLASHINFER_ATTN python examples/offline_inference/text_to_video/text_to_video.py \
-    --model Lightricks/LTX-2 \
-    --prompt "A dog running across a field of golden wheat." \
-    --height 480 --width 832 --num-frames 33 \
-    --num-inference-steps 40 --seed 42 --guidance-scale 4.0 \
-    --output ltx2.mp4
-```
-
-### SageAttention (lossy)
-
-```bash
-DIFFUSION_ATTENTION_BACKEND=SAGE_ATTN python examples/offline_inference/text_to_video/text_to_video.py \
-    --model hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v \
-    --prompt "A dog running across a field of golden wheat." \
-    --height 480 --width 832 --num-frames 33 \
-    --num-inference-steps 30 --seed 42 --guidance-scale 6.0 \
-    --tensor-parallel-size 2 \
-    --output hv15_sage.mp4
-```
-
-Example: Wan2.2 TI2V 5B
-
-```bash
-DIFFUSION_ATTENTION_BACKEND=SAGE_ATTN python examples/offline_inference/text_to_video/text_to_video.py \
-    --model Wan-AI/Wan2.2-TI2V-5B-Diffusers \
-    --prompt "A dog running across a field of golden wheat." \
-    --height 704 --width 1280 --num-frames 49 \
-    --num-inference-steps 30 --seed 42 --guidance-scale 5.0 \
-    --tensor-parallel-size 2 \
-    --output outputs/wan22_sage.mp4
-```
-
-### Enable SageAttention3
-
-Example:
-
-```bash
-DIFFUSION_ATTENTION_BACKEND=SAGE_ATTN_3 python examples/offline_inference/text_to_video/text_to_video.py \
-    --model hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v \
-    --prompt "A dog running across a field of golden wheat." \
-    --height 480 --width 832 --num-frames 33 \
-    --num-inference-steps 30 --seed 42 --guidance-scale 6.0 \
-    --tensor-parallel-size 2 \
-    --output outputs/hv15_sage3.mp4
-```
+See [Hugging Face Hub Backends](attention_backends/huggingface_hub.md).
@@ -0,0 +1,110 @@
+# Dense Attention Backends
+
+Use these backends when you want dense attention without backend-specific
+sparsity. Start with `TORCH_SDPA` for correctness comparisons, then benchmark
+the fastest compatible kernel for your model, shape, and hardware.
+
+For selection precedence and per-role configuration, see the
+[attention backend overview](../attention_backends.md).
+
+## `TORCH_SDPA`
+
+`TORCH_SDPA` calls PyTorch `scaled_dot_product_attention` and lets PyTorch's
+dispatcher choose the implementation. It is always available and is the most
+conservative reference when validating another backend.
+
+```bash
+vllm-omni serve <model> --diffusion-attention-backend TORCH_SDPA
+```
+
+## `FLASH_ATTN`
+
+`FLASH_ATTN` uses the installed FlashAttention implementation. On Blackwell it
+prefers the CuTe-based FlashAttention 4 path and falls back to
+FlashAttention 3 or 2 when that path is unavailable. On Hopper, Ada, and
+Ampere it is the preferred automatic route when a compatible package is
+installed.
+
+```bash
+vllm-omni serve <model> --diffusion-attention-backend FLASH_ATTN
+```
+
+### FlashAttention 4 on Blackwell
+
+Install the optional CUDA 13 extra:
+
+```bash
+pip install 'vllm-omni[fa4]'
+```
+
+Version `4.0.0b18` is required; earlier beta wheels had known JIT failures on
+Blackwell. If the CuTe path is unavailable, the backend falls back to the
+compatible FlashAttention 3 or 2 path.
+
+## `CUDNN_ATTN`
+
+`CUDNN_ATTN` pins PyTorch SDPA to `CUDNN_ATTENTION`. It is particularly useful
+for mask-heavy DiTs and is automatically preferred on Blackwell when cuDNN
+9.5 or newer is available and the higher-priority TRTLLM route is not
+compatible.
+
+```bash
+vllm-omni serve <model> --diffusion-attention-backend CUDNN_ATTN
+```
+
+### LTX-2.0 limitation
+
+LTX-2 audio attention has a symbolic head dimension during `torch.compile`
+tracing. The cuDNN SDPA selector rejects that symbolic dimension and Dynamo
+aborts compilation. This is tracked in
+[issue #3121](https://github.com/vllm-project/vllm-omni/issues/3121).
+
+Use `FLASHINFER_ATTN` or `TORCH_SDPA` as a workaround:
+
+```bash
+DIFFUSION_ATTENTION_BACKEND=FLASHINFER_ATTN \
+  python examples/offline_inference/text_to_video/text_to_video.py \
+  --model Lightricks/LTX-2 ...
+```
+
+## `FLASHINFER_ATTN`
+
+`FLASHINFER_ATTN` uses FlashInfer's batch-prefill wrapper. It is an explicit
+option on CUDA platforms and an automatic Blackwell fallback when FlashInfer
+is installed but cuDNN is too old for `CUDNN_ATTN`.
+
+```bash
+vllm-omni serve <model> --diffusion-attention-backend FLASHINFER_ATTN
+```
+
+### FlashInfer quantized attention
+
+The backend accepts an `AttentionSpec.quant` block. For QK16/V8, keep Q and K
+in FP16 or BF16 and use FP8 E4M3 for V:
+
+```python
+from vllm_omni.diffusion.data import (
+    AttentionConfig,
+    AttentionSpec,
+    AttnQuantSpec,
+    OmniDiffusionConfig,
+)
+
+config = OmniDiffusionConfig(
+    diffusion_attention_config=AttentionConfig(
+        default=AttentionSpec(
+            backend="FLASHINFER_ATTN",
+            quant=AttnQuantSpec(
+                dtype_qk="bfloat16",
+                dtype_vo="fp8_e4m3",
+            ),
+        ),
+    ),
+    ...,
+)
+```
+
+`dtype_qk` controls Q and K; `dtype_vo` controls V. Mixed-dtype
+configurations require FlashInfer 0.6.16rc1 or newer. The shared quantization
+schema is also consumed by TRTLLM, but each backend validates its own allowed
+fields and values; see [TRTLLM SAGE quantization](trtllm.md#sage-quantization).
@@ -0,0 +1,32 @@
+# Hugging Face Hub Attention Backends
+
+Hub backends load the same published kernels commonly used with Hugging Face
+Diffusers. They are useful when training and serving must use matching kernel
+implementations to reduce numerical drift and sampling divergence.
+
+| Backend | Kernel | Platform |
+| --- | --- | --- |
+| `FLASH_ATTN_HUB` | `kernels-community/flash-attn2` | Compatible CUDA GPUs |
+| `FLASH_ATTN_3_HUB` | `kernels-community/flash-attn3` | Hopper sm_90 or newer |
+
+## Installation
+
+```bash
+pip install kernels==0.14.1
+```
+
+If `kernels` is unavailable, vLLM-Omni warns and falls back to the
+corresponding local FlashAttention implementation. On CUDA devices older than
+Hopper, `FLASH_ATTN_3_HUB` falls back to `FLASH_ATTN_HUB`.
+
+## Usage
+
+```bash
+export DIFFUSION_ATTENTION_BACKEND=FLASH_ATTN_HUB
+
+# Equivalent CLI selection
+vllm-omni serve <model> --diffusion-attention-backend FLASH_ATTN_HUB
+```
+
+For per-role kernel alignment, configure the Hub backend on only the relevant
+role; see the [attention backend overview](../attention_backends.md#per-role-configuration).
@@ -0,0 +1,72 @@
+# RainFusion Attention
+
+`RAINFUSION_ATTN` runs MindIE-SD RainFusion (`rf_v2`) block-sparse video
+attention on Ascend NPU. It pools 128-token key blocks, ranks them per query
+block, and attends to the highest-scoring blocks after arranging video tokens
+in `(t, h, w)` order.
+
+Only the video segment is sparse. Prefix rows such as text, visual conditions,
+and audio, plus first-frame blocks, remain dense. Unsupported calls—including
+warmup steps, skipped layers, missing video geometry, and video segments below
+32 blocks—delegate to `FLASH_ATTN`, so compatible models can select RainFusion
+globally.
+
+## Configuration
+
+| Key | Valid values | Meaning |
+| --- | --- | --- |
+| `sparsity` | finite, `[0, 1]` | Nominal dropped-key-block fraction; default `0.8`; `0` disables sparsity |
+| `start_step` | integer, `>= 0` | Number of early denoise steps kept dense |
+| `skip_layers` | selector such as `"0-3,38"` | DiT blocks kept dense |
+
+```bash
+vllm-omni serve MiniMaxAI/MiniMax-H3 \
+  --diffusion-attention-config '{"default":{"backend":"RAINFUSION_ATTN",\
+    "block_sparse":{"sparsity":0.8,"start_step":0}}}'
+```
+
+```python
+from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, BlockSparseSpec
+
+config = AttentionConfig(
+    default=AttentionSpec(
+        backend="RAINFUSION_ATTN",
+        block_sparse=BlockSparseSpec(
+            sparsity=0.8,
+            start_step=0,
+            skip_layers="0-1",
+        ),
+    ),
+)
+```
+
+Tune `start_step`, then `sparsity`, then `skip_layers`. Increasing the dense
+early-step window is usually the cheapest way to recover global structure;
+use layer exclusions only after a same-seed dense comparison identifies
+sensitive blocks.
+
+## Requirements and compatibility
+
+RainFusion requires Ascend NPU and `mindiesd`. Selecting it on another
+platform raises. It is incompatible with ring sequence parallelism because
+the kernel needs the complete key sequence for ranking; use Ulysses sequence
+parallelism with `ring_degree=1`.
+
+## Geometry handling
+
+RainFusion handles arbitrary video grids by rearranging video spatially and,
+when necessary, promoting a real-video suffix to the always-kept prefix. It
+does not pad because `rf_v2` does not consume a padding attention mask.
+
+For MiniMax-H3 at 1344x768, grid `(62, 24, 42)` produces 62,496 video rows.
+The implementation promotes 2,976 rows and leaves 59,520 sparse rows
+(`465 x 128`) so the mask and kernel tiling remain aligned.
+
+Spatial grids aligned to 8x8 are still preferable. A latent height or width
+not divisible by 8 creates an always-kept suffix and reduces realized
+sparsity. At the image level, that means width and height divisible by 256.
+Protected video tails require a compatible MindIE-SD release.
+
+For common configuration and selector behavior, see the
+[attention backend overview](../attention_backends.md) and the
+[backend selection design](../../../design/feature/attention_backend_selection.md).
@@ -0,0 +1,63 @@
+# SageAttention
+
+SageAttention backends provide lossy low-precision attention for diffusion
+models. Validate output quality against `TORCH_SDPA` at the same seed before
+using either backend in production.
+
+## `SAGE_ATTN`
+
+`SAGE_ATTN` uses SageAttention 2.2 with INT8-quantized attention and FP16
+accumulation.
+
+### Installation
+
+Install SageAttention into the same environment as vLLM-Omni:
+
+```bash
+git clone https://github.com/thu-ml/SageAttention.git
+cd SageAttention
+export EXT_PARALLEL=4 NVCC_APPEND_FLAGS="--threads 8" MAX_JOBS=32
+pip install . --no-build-isolation
+```
+
+Verify the installation:
+
+```bash
+python -c "import sageattention; print(sageattention.__file__)"
+```
+
+Select it globally:
+
+```bash
+vllm-omni serve <model> --diffusion-attention-backend SAGE_ATTN
+```
+
+## `SAGE_ATTN_3`
+
+`SAGE_ATTN_3` uses the SageAttention3 Blackwell implementation.
+
+### SageAttention3 installation
+
+```bash
+git clone https://github.com/thu-ml/SageAttention.git
+cd SageAttention/sageattention3_blackwell
+python setup.py install
+```
+
+Verify the installation:
+
+```bash
+python -c "import sageattn3; print(sageattn3.__file__)"
+```
+
+```bash
+vllm-omni serve <model> --diffusion-attention-backend SAGE_ATTN_3
+```
+
+`SAGE_ATTN_3` requires CUDA, an importable `sageattn3`, and a Blackwell-class
+GPU. Its kernel assumes the query-head count equals the key/value-head count.
+GQA and MQA diffusion calls therefore fall back to PyTorch SDPA for
+correctness.
+
+For common configuration and platform routing, see the
+[attention backend overview](../attention_backends.md).
@@ -0,0 +1,103 @@
+# TRTLLM Attention
+
+`TRTLLM_ATTN` runs FlashInfer's trtllm-gen FMHA kernels. It supports dense
+BF16 attention plus optional Skip-Softmax sparsity or SAGE quantization.
+
+For common selection and per-role configuration, see the
+[attention backend overview](../attention_backends.md).
+
+## Requirements
+
+Selecting `TRTLLM_ATTN` requires:
+
+- datacenter Blackwell, sm_100 or sm_103;
+- `head_dim=128`;
+- FlashInfer; and
+- a model-declared packed or mask-free compatible path.
+
+It is not supported on workstation Blackwell (sm_120/sm_121). An incompatible
+explicit selection raises instead of silently changing the backend.
+
+## Dense mode
+
+Select `TRTLLM_ATTN` without `skip_softmax` or `quant` to run dense BF16:
+
+```bash
+vllm-omni serve <model> --diffusion-attention-backend TRTLLM_ATTN
+```
+
+## Skip-Softmax
+
+Skip-Softmax trades some fidelity for speed by skipping low-contribution
+attention work. See the [feature design](../../../design/feature/skip_softmax.md)
+for the algorithm and implementation contract.
+
+| Key | Valid values | Meaning |
+| --- | --- | --- |
+| `target_sparsity` | finite, `[0, 1]` | Point on a model-calibrated curve |
+| `threshold` | finite, `>= 0` | Direct threshold; mutually exclusive with `target_sparsity` |
+| `disabled_until_timestep` | finite, `[0, 1]` | Keep early high-noise steps dense until normalized `t <= D` |
+
+```bash
+vllm-omni serve Wan-AI/Wan2.2-T2V-A14B-Diffusers \
+  --diffusion-attention-config '{"default":{"backend":"TRTLLM_ATTN",\
+    "skip_softmax":{"target_sparsity":0.65,"disabled_until_timestep":0.86}}}'
+```
+
+```python
+from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, SkipSoftmaxSpec
+
+config = AttentionConfig(
+    default=AttentionSpec(
+        backend="TRTLLM_ATTN",
+        skip_softmax=SkipSoftmaxSpec(
+            target_sparsity=0.65,
+            disabled_until_timestep=0.86,
+        ),
+    ),
+)
+```
+
+Start with `target_sparsity=0.65` and `disabled_until_timestep=0.86`. Increase
+sparsity only after comparing quality against dense output at the same seed.
+The benefit grows with sequence length because only attention is accelerated.
+
+## SAGE quantization
+
+TRTLLM SAGE quantizes Q/K per block and V per channel so both attention
+matrix multiplications use low precision.
+
+| Key | Valid values | Meaning |
+| --- | --- | --- |
+| `dtype_qk` | `int8`, `fp8_e4m3` | Q/K quantization dtype; absent means dense |
+| `q_block_size` | `1`, `4`, `16` | Q scale block size; default `1` |
+| `k_block_size` | `1`, `4`, `16` | K scale block size; default `16` |
+
+V is always quantized per channel to FP8 E4M3 and K-smoothing is internal.
+Only block sizes with compiled kernels are accepted.
+
+```bash
+vllm-omni serve Wan-AI/Wan2.2-T2V-A14B-Diffusers \
+  --diffusion-attention-config '{"default":{"backend":"TRTLLM_ATTN",\
+    "quant":{"dtype_qk":"fp8_e4m3","q_block_size":1,"k_block_size":16}}}'
+```
+
+```python
+from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, AttnQuantSpec
+
+config = AttentionConfig(
+    default=AttentionSpec(
+        backend="TRTLLM_ATTN",
+        quant=AttnQuantSpec(
+            dtype_qk="fp8_e4m3",
+            q_block_size=1,
+            k_block_size=16,
+        ),
+    ),
+)
+```
+
+This path requires FlashInfer 0.6.16rc1 or newer. FP8 Q/K kernels are
+available on sm_100 and sm_103; INT8 Q/K kernels are compiled for sm_100 only.
+The shared `AttnQuantSpec` is backend typed: values intended for
+`FLASHINFER_ATTN` are rejected here and vice versa.
@@ -1,389 +1,75 @@
 # CPU Offloading for Diffusion Models

-## Overview
+vLLM-Omni provides three CPU-offload strategies for diffusion models. Choose
+the coarsest strategy that meets your memory target: finer-grained streaming
+saves more device memory but adds more transfer and setup overhead.

-vLLM-Omni provides two offloading strategies to reduce GPU memory usage for diffusion models:
+For the shared factory, discovery, and lifecycle contract, see the
+[CPU Offloading design](../../design/feature/offloader/README.md).

-1. **Model-level (Sequential) Offloading**: Mutual exclusion between DiT model and encoder - only one is on GPU at a time.
-2. **Layerwise (Blockwise) Offloading**: Keeps only one transformer block on GPU at a time with compute-memory overlap.
+## Choose a strategy

-Both strategies use pinned memory for faster CPU-GPU transfers. The strategies are **mutually exclusive** for now - if both are enabled, layerwise takes priority.
+| Strategy | Device residency | Parallel scope | Guide |
+| --- | --- | --- | --- |
+| Model-level (sequential) | One pipeline component group at a time | Single device | [Model-Level Offloading](offloader/cpu_offload.md) |
+| Layerwise (blockwise) | One transformer block, with next-block prefetch | Single device | [Layerwise Offloading](offloader/layerwise_offload.md) |
+| Distributed layerwise | Fixed two-block device buffer; optional host-weight sharding and AllGather | Multiple GPU/NPU ranks | [Distributed Layerwise Offloading](offloader/distributed_layerwise_offload.md) |

+All strategies use pinned host memory for faster transfers where applicable.
+Configuration priority is:
+
+1. Distributed layerwise offloading.
+2. Layerwise offloading.
+3. Model-level offloading.
+
+Treat the flags as mutually exclusive. If more than one is enabled, the
+higher-priority strategy is selected.
+
+## Quick selection
+
+- Use [model-level offloading](offloader/cpu_offload.md) when swapping whole
+  encoders and DiTs is enough to fit and phase-boundary transfers are
+  acceptable.
+- Use [layerwise offloading](offloader/layerwise_offload.md) for compute-heavy video
+  DiTs where block transfers can overlap computation.
+- Use [distributed layerwise offloading](offloader/distributed_layerwise_offload.md)
+  when a multi-rank deployment also needs bounded device residency and,
+  optionally, sharded host weights.
+
+## Supported models
+
+| Architecture | Example models | DiT class | Model-level | Layerwise | Distributed layerwise | Layerwise block attributes |
+| --- | --- | --- | :---: | :---: | :---: | --- |
+| Flux2Pipeline | `black-forest-labs/FLUX.2-dev` | `Flux2Transformer2DModel` | yes | yes | — | `transformer_blocks`, `single_transformer_blocks` |
+| LongCatImagePipeline | `meituan-longcat/LongCat-Image` | `LongCatImageTransformer2DModel` | — | yes | — | `transformer_blocks`, `single_transformer_blocks` |
+| NextStep11Pipeline | `stepfun-ai/NextStep-1.1` | `NextStepModel` | — | yes | — | `layers` |
+| OvisImagePipeline | `AIDC-AI/Ovis-Image-7B` | `OvisImageTransformer2DModel` | — | yes | — | `transformer` |
+| QwenImagePipeline | `Qwen/Qwen-Image` | `QwenImageTransformer2DModel` | yes | yes | — | `transformer_blocks` |
+| StableDiffusionXLPipeline | `stabilityai/stable-diffusion-xl-base-1.0` | `SDXLUNet2DConditionModel` | yes | yes | — | `down_blocks`, `up_blocks` |
+| StableDiffusion3Pipeline | `stabilityai/stable-diffusion-3.5-medium` | `SD3Transformer2DModel` | — | yes | — | `transformer_blocks` |
+| Wan22I2VPipeline | `Wan-AI/Wan2.2-I2V-A14B-Diffusers` | `WanTransformer3DModel` | yes | yes | — | `blocks` |
+| Wan22Pipeline | `Wan-AI/Wan2.2-T2V-A14B-Diffusers` | `WanTransformer3DModel` | yes | yes | — | `blocks` |
+| SoulXSingerPipeline / SoulXSingerSVCPipeline | `Soul-AILab/SoulX-Singer` | `DiffLlama` | yes | yes | — | `layers` |
+| BagelPipeline | `ByteDance-Seed/BAGEL-7B-MoT` | `Qwen2MoTModel` | — | yes | — | `layers`, customized modules |
+| Cosmos3OmniDiffusersPipeline | `nvidia/Cosmos3-Nano`, `nvidia/Cosmos3-Super` | `Cosmos3VFMTransformer`, `Cosmos3LanguageModel` | yes | yes | yes | `layers`, `gen_layers` |
+
+Model-level support requires discoverable DiT and encoder components.
+Layerwise support requires transformer block topology. Distributed support
+reuses that topology but still requires validation for the model, checkpoint,
+and parallel configuration.
+
+## Compatibility anchors
+
+These headings preserve links to sections that moved into dedicated guides.

 ## Model-level (Sequential) Offloading

-### How It Works
-
-Model-level offloading implements mutual exclusion between DiT transformer and encoder modules using pre forward hooks:
-
- **When encoders run**: DiT transformer is offloaded to CPU
- **When DiT runs**: Encoders are offloaded to CPU, if more than one dit models, only one loaded on GPU, others get offloaded to CPU.
- **VAE**: Stays resident on GPU
-
-Before each module's forward pass, the hook automatically moves it to GPU while offloading the other module group to CPU. Transfers use pinned memory for speed.
-
-### Usage
-
-**Python API:**
-```python
-from vllm_omni import Omni
-
-m = Omni(model="Wan-AI/Wan2.2-T2V-A14B-Diffusers", enable_cpu_offload=True)
-```
-
-**CLI:**
-```bash
-vllm serve Wan-AI/Wan2.2-T2V-A14B-Diffusers --omni --enable-cpu-offload
-```
-
-### To Support a Model
-
-Implement the `SupportsComponentDiscovery` protocol to declare which
-submodules serve as pipeline components (used by offloading, HSDP
-sharding, and other framework features):
-
-```python
-from typing import ClassVar
-from vllm_omni.diffusion.models.interface import SupportsComponentDiscovery
-
-class MyPipeline(nn.Module, SupportsComponentDiscovery):
-    _dit_modules: ClassVar[list[str]] = ["transformer"]
-    _encoder_modules: ClassVar[list[str]] = ["text_encoder", "vision_model"]
-    _vae_modules: ClassVar[list[str]] = ["vae"]
-    _resident_modules: ClassVar[list[str]] = []  # optional
-
-    def __init__(self):
-        super().__init__()
-        self.transformer = ...     # DiT — stays on GPU during denoising
-        self.text_encoder = ...    # Encoder — offloaded to CPU during denoising
-        self.vision_model = ...    # Encoder — offloaded to CPU during denoising
-        self.vae = ...             # VAE — always on GPU
-```
-
- `_dit_modules`: attribute names of denoising submodules (kept on GPU
-  during the diffusion loop).
- `_encoder_modules`: attribute names of encoder/vision submodules
-  (offloaded to CPU during the diffusion loop).
- `_vae_modules`: attribute names of VAE(s) (always kept on GPU, not
-  part of the mutual exclusion hooks).
- `_resident_modules`: attribute names of small submodules that must
-  stay on GPU during layerwise offloading (e.g. embedders, connectors).
-  Optional — defaults to `[]`.
-
-All attribute names support dotted paths for nested submodules
-(e.g. `"pipe.transformer"`, `"bagel.time_embedder"`).
-
-Both DiT and encoder lists are needed because the offload hooks use
-mutual exclusion: when one group runs, the other moves to CPU.
-
-### Limitations
- Cold start latency increases
- Adds overhead from CPU-GPU transfers between encoder and denoising phases
- Support single GPU only for now
-
-
-### Component offloading for split models (e.g. Cosmos3)
-
-Some models split their transformer into mutually-exclusive *components* that run
-in different phases of a single forward pass rather than as separate pipeline
-components -- e.g. Cosmos3's understanding (reasoner) component runs once per
-generation while the generation (generator) component runs every denoising step.
-Such models have no separate text encoder to swap against, so the transformer
-owns a small model-local offload path and wraps each phase with
-`with self._offload_context(name):`
-
-```python
-class Cosmos3VFMTransformer(nn.Module):
-    def forward(self, ...):
-        with self._offload_context("reasoner"):
-            ...  # understanding pass, runs once
-        with self._offload_context("generator"):
-            ...  # denoising pass, runs every step
-```
-
-Model-level offloading then keeps exactly one component GPU-resident at a time
-(the other on CPU), reusing the same `SequentialOffloadHook` `.to()` movers. The
-pipeline opts in by exposing `enable_omni_model_cpu_offload` (which drives the
-transformer's `enable_model_cpu_offload` and pins the VAE). Layerwise offloading
-works for these models too -- each component declares its own block container via
-`_layerwise_offload_blocks_attrs`.
-
+See [Model-Level Offloading](offloader/cpu_offload.md).

 ## Layerwise (Blockwise) Offloading

-### How It Works
-
-Layerwise offloading keeps only one transformer block on GPU at a time.
-
-As each block completes, the next block is prefetched to GPU while the current block is freed. The pre and forward hooks utilized by layerwise offloading apply a separate CUDA stream (`copy_stream`) to overlap weight transfer with computation, and retain flattened tensors in pinned CPU memory for block parameters re-materialization. Encoders, VAE, and non-block DiT modules (embeddings, norms) always stay on GPU.
-
-**Execution Flow:**
-
-| Block | Pre-forward Hook | Forward | Post-forward Hook |
-|-------|------------------|---------|-------------------|
-| block-0 | Prefetch block-1 (async) | Compute block-0 | Free block-0 |
-| block-1 | Prefetch block-2 (async) | Compute block-1 | Free block-1 |
-| ... | ... | ... | ... |
-| block-(n-1) | **Prefetch block-0** (async) | Compute block-(n-1) | Free block-(n-1) |
-
-Each transformer block has a `LayerwiseOffloadHook` that prefetches the next block before forward and frees the current block after forward.
-
-Layerwise offloading is primarily recommended for large **video generation models** where the compute cost per block is high enough to effectively overlap with memory prefetch operations. For example, Wan2.2 T2V and I2V pipelines.
-
-### Usage
-
-**Python API:**
-```python
-from vllm_omni import Omni
-
-# Text-to-video
-m = Omni(model="Wan-AI/Wan2.2-T2V-A14B-Diffusers", enable_layerwise_offload=True)
-
-# Or image-to-video
-m = Omni(model="Wan-AI/Wan2.2-I2V-A14B-Diffusers", enable_layerwise_offload=True)
-```
-
-**CLI:**
-```bash
-# Text-to-video
-vllm serve Wan-AI/Wan2.2-T2V-A14B-Diffusers --omni --enable-layerwise-offload
-
-# Or image-to-video
-vllm serve Wan-AI/Wan2.2-I2V-A14B-Diffusers --omni --enable-layerwise-offload
-```
-
-### To Support a Model
-
-Models must define the blocks attribute name for layerwise offloading:
-
-```python
-class WanTransformer3DModel(nn.Module):
-    _layerwise_offload_blocks_attrs = ["blocks"]  # Attribute names containing transformer blocks
-
-    def __init__(self):
-        self.blocks = nn.ModuleList([...])  # Transformer blocks
-```
-
-For models with multiple block types:
-
-```python
-class Flux2Transformer2DModel(nn.Module):
-    _layerwise_offload_blocks_attrs = ["transformer_blocks", "single_transformer_blocks"]
-```
-
-### Limitations
- Cold start latency increases because offloaded components must be moved to CPU
-  during setup; layerwise offload may add extra weight consolidation and pinning
-  work.
- Performance depends on compute cost and H2D bandwidth as well
- Support single GPU only for now
-
+See [Layerwise Offloading](offloader/layerwise_offload.md).

 ## Distributed Layerwise Offloading

-### How It Works
-
-Distributed layerwise offloading extends single-GPU layerwise offloading to
-multi-device deployments. In the default mode, each rank stores only
-**1/dp_size** of the model weights in host memory; full layer weights are
-reconstructed at runtime via **AllGather** on a dedicated communication stream,
-overlapped with computation via a fixed double-buffer scheme. The
-`--dlo-no-use-allgather` mode instead streams standard-loader rank-local
-weights independently and does not shard weights across ranks.
-
-For the implementation design and compatibility matrix, see
-[Distributed Layerwise Offload](../../design/feature/distributed_layerwise_offload.md).
-
-**Key features:**
- **Weight sharding + AllGather**: each rank stores 1/dp_size of weights,
-  reconstructed per-layer via `all_gather_into_tensor`
- **Fixed double-buffer**: exactly 2 transformer blocks on each device at any
-  time, regardless of model size
- **DP multi-concurrency**: N concurrent requests processed in parallel
-  (AllGather only gathers weight shards, which are request-independent)
- **mmap weight loading**: weights loaded as mmap views pointing to shared OS
-  page cache, eliminating O(dp_size × model_size) RSS during model creation
- **Platform-agnostic**: works on NVIDIA GPU (CUDA/NCCL) and Ascend NPU
-  (CANN/HCCL) via vLLM-Omni's platform abstraction
-
-**Execution Flow (per device):**
-
-```
-Compute Stream:  [Layer N]          [Layer N+1]          [Layer N+2]
-H2D Stream:      [H2D Shard N+1]   [H2D Shard N+2]
-AllGather:       [AG N+1]          [AG N+2]
-Slot Usage:      Slot0: Layer N    Slot1: Layer N+1
-```
-
-### Usage
-
-**CLI:**
-```bash
-# 4× GPU/NPU with AllGather (recommended)
-vllm serve /path/to/model --omni \
-  --enable-distributed-layerwise-offload \
-  --data-parallel-size 4
-
-# Without DLO AllGather (each rank streams standard-loader rank-local weights)
-vllm serve /path/to/model --omni \
-  --enable-distributed-layerwise-offload \
-  --data-parallel-size 4 \
-  --dlo-no-use-allgather
-
-# With SP instead of DP (long sequences)
-vllm serve /path/to/model --omni \
-  --enable-distributed-layerwise-offload \
-  --usp 4
-```
-
-**Python API:**
-```python
-from vllm_omni import Omni
-
-m = Omni(
-    model="/path/to/model",
-    enable_distributed_layerwise_offload=True,
-    dlo_use_allgather=True,  # default
-)
-```
-
-### CLI Flags
-
-| Flag | Description | Default |
-|------|-------------|---------|
-| `--enable-distributed-layerwise-offload` | Enable DLO | `false` |
-| `--data-parallel-size N` | Number of DP ranks; DLO weight-sharding group when AllGather is enabled | `1` |
-| `--dlo-use-allgather` | Shard + AllGather (saves CPU, requires concurrent requests) | `true` |
-| `--dlo-no-use-allgather` | Use standard-loader rank-local weights; no DLO sharding or AllGather | `false` |
-
-### How model weights are loaded (mmap path)
-
-When DLO + AllGather is active, the offloader:
-
-1. **Saves non-persistent buffers** (e.g. RoPE `inv_freq`, timestep `freqs`)
-   from the normally-created transformer
-2. **Converts to meta device** via `to_empty(device="meta")`, releasing random
-   initialization weights
-3. **Loads checkpoint weights as mmap views** via `safe_open().get_tensor()`,
-   which return views into the OS page cache (shared across all ranks, 0 RSS)
-4. **Calls `post_load_weights()`** to apply model-specific dtype conversions
-   (e.g. Cosmos3's `time_embedder` → FP32)
-5. **Restores non-persistent buffers** from saved copies
-
-This approach requires **zero model-specific code changes** — no pipeline or
-transformer modifications are needed.
-
-### OffloadPlan (declarative topology metadata)
-
-Models can optionally declare an `OffloadPlan` class variable to provide
-topology metadata (block attribute names, submodules to offload) without
-any offload-specific logic:
-
-```python
-from vllm_omni.diffusion.offloader import OffloadPlan
-
-class MyPipeline(nn.Module):
-    _dit_modules = ["transformer"]
-    _offload_plan = OffloadPlan(
-        block_attrs={"transformer": ("blocks",)},
-        offload_submodules={"context_encoder": "layers"},
-    )
-```
-
-When not declared, the offloader falls back to `_layerwise_offload_blocks_attrs`
-and heuristic attribute search.  This is backward-compatible — existing models
-work without any changes.
-
-### DP Multi-concurrency
-
-When `--data-parallel-size > 1` and AllGather is enabled, the scheduler batches
-up to `dp_size` requests per denoise step.  Each DP rank processes a different
-request while AllGather synchronizes only weight shards (request-independent).
-
-**Requirements for concurrent requests:**
- `num_inference_steps` must be specified explicitly (None is not allowed)
- All concurrent requests must have the same `num_inference_steps` value
- These constraints exist because AllGather is a collective that requires
-  every rank to participate at each denoise step
-
-### Limitations
- Online quantization (FP8) is incompatible with mmap loading — falls back to
-  regular loading only when the model is outside the DLO+AllGather mmap path;
-  online quantization with DLO+AllGather is rejected. Use
-  `--dlo-no-use-allgather` or disable online quantization.
- Tensor Parallel > 1 is rejected in the DLO+AllGather mmap path because TP-aware
-  weight-loader callbacks are bypassed. The no-AllGather path retains the
-  standard TP loader and is experimental.
- HSDP + AllGather is rejected because it would double-shard parameters. HSDP
-  with no-AllGather is accepted at configuration level, but has limited
-  end-to-end validation.
- `num_inference_steps=None` is not allowed in DP multi-concurrency mode
-
-**Module Discovery**
-
-The offloader discovers pipeline components in two ways:
-
-1. **Protocol-based** (preferred): If the pipeline implements
-    `SupportsComponentDiscovery`, its `_dit_modules`, `_encoder_modules`,
-    `_vae_modules`, and `_resident_modules` class variables are used
-    directly.  All attribute names support dotted paths (e.g.
-    `"pipe.transformer"`, `"bagel.time_embedder"`) for nested submodules.
-
-2. **Fallback attribute scan**: Otherwise, the offloader scans for
-    well-known attribute names:
-    - **DiT modules**: `transformer`, `transformer_2`, `dit`, `sr_dit`, `language_model`, `transformer_blocks`, `model`
-    - **Encoders**: `text_encoder`, `text_encoder_2`, `text_encoder_3`, `image_encoder`
-    - **VAE**: `vae`, `audio_vae`
-
-**Hook System**
-
-Both strategies use vLLM-Omni's hook registry system (`HookRegistry` and `ModelHook`) to register pre/post forward callbacks on modules, enabling automatic swapping without modifying model code.
-
-**Backend Architecture**
-
-```
-OffloadBackend (base class)
-├── ModelLevelOffloadBackend → uses SequentialOffloadHook (.to() swap)
-│                              (delegates to a pipeline's enable_omni_model_cpu_offload
-│                               for split models like Cosmos3)
-├── LayerWiseOffloadBackend → uses LayerwiseOffloadHook
-│                          (single-GPU, full weights on host)
-└── DistributedLayerwiseOffloadBackend → uses DistributedLayerwiseOffloadHook
-                                         (multi-GPU, sharded AllGather or rank-local mode)
-```
-
-Factory function `get_offload_backend()` selects the appropriate backend based on
-configuration.
-
-For split models, `ModelLevelOffloadBackend.enable()` detects a pipeline's
-`enable_omni_model_cpu_offload` hook and delegates to it; Cosmos3 then swaps its
-reasoner/generator components inside the model forward pass.
-
-
-## Supported Models
-
-| Architecture | Example Models | DiT Class | Model-Level Offload | Layerwise Offload | Distributed Layerwise Offload | Blocks Attrs (Layerwise specific) |
-|--------------|----------------|-----------|---------------------|-------------------|-------------------------------|-----------------------------------|
-| Flux2Pipeline | `black-forest-labs/FLUX.2-dev` | `Flux2Transformer2DModel` | ✓ | ✓ | - | `"transformer_blocks"`, `"single_transformer_blocks"` |
-| LongCatImagePipeline | `meituan-longcat/LongCat-Image` | `LongCatImageTransformer2DModel` | - | ✓ | - | `"transformer_blocks"`, `"single_transformer_blocks"` |
-| NextStep11Pipeline | `stepfun-ai/NextStep-1.1` | `NextStepModel` | - | ✓ | - | `"layers"` |
-| OvisImagePipeline | `AIDC-AI/Ovis-Image-7B` | `OvisImageTransformer2DModel` | - | ✓ | - | `"transformer"` |
-| QwenImagePipeline | `Qwen/Qwen-Image` | `QwenImageTransformer2DModel` | ✓ | ✓ | - | `"transformer_blocks"` |
-| StableDiffusionXLPipeline | `stabilityai/stable-diffusion-xl-base-1.0` | `SDXLUNet2DConditionModel` | ✓ | ✓ | - | `"down_blocks"`, `"up_blocks"` |
-| StableDiffusion3Pipeline | `stabilityai/stable-diffusion-3.5-medium` | `SD3Transformer2DModel` | - | ✓ | - | `"transformer_blocks"` |
-| Wan22I2VPipeline | `Wan-AI/Wan2.2-I2V-A14B-Diffusers` | `WanTransformer3DModel` | ✓ | ✓ | - | `"blocks"` |
-| Wan22Pipeline | `Wan-AI/Wan2.2-T2V-A14B-Diffusers` | `WanTransformer3DModel` | ✓ | ✓ | - | `"blocks"` |
-| SoulXSingerPipeline / SoulXSingerSVCPipeline | `Soul-AILab/SoulX-Singer` | `DiffLlama` (`cfm_decoder.model.diff_estimator`) | ✓ | ✓ | - | `"layers"` |
-| BagelPipeline | `ByteDance-Seed/BAGEL-7B-MoT` | `Qwen2MoTModel` | - | ✓ | - | `"layers"`, `"customized modules"` |
-| Cosmos3OmniDiffusersPipeline | `nvidia/Cosmos3-Nano`, `nvidia/Cosmos3-Super` | `Cosmos3VFMTransformer`, `Cosmos3LanguageModel` | ✓ | ✓ | ✓ | `"layers"`, `"gen_layers"` |
-
-**Notes:**
- Model-Level Offloading is expected to be supported by all common diffusion models (DiT and encoders) naturally
- Layerwise Offloading requires DiT class to define `_layerwise_offload_blocks_attrs` pointing to transformer blocks
- Distributed Layerwise Offloading reuses the layerwise block topology, but the
-  model and parallelism combination still needs validation. AllGather mode
-  additionally requires a compatible checkpoint/mmap layout; no-AllGather mode
-  retains the standard loader's rank-local tensors. See the
-  [DLO design](../../design/feature/distributed_layerwise_offload.md) and
-  [Cosmos3 DistOffload recipe](https://github.com/vllm-project/vllm-omni/blob/main/recipes/cosmos3/Cosmos3-DistOffload.md)
-  for usage examples.
+See [Distributed Layerwise Offloading](offloader/distributed_layerwise_offload.md).
@@ -0,0 +1,81 @@
+# Model-Level Offloading
+
+Model-level, or sequential, offloading keeps only the pipeline component group
+currently executing on the accelerator. It is the simplest offload strategy
+and is selected with `--enable-cpu-offload`.
+
+## How it works
+
+Pre-forward hooks enforce mutual exclusion between DiT and encoder modules:
+
+- before an encoder runs, the DiT moves to CPU;
+- before a DiT runs, encoders and other DiTs move to CPU; and
+- VAE modules remain on the accelerator.
+
+Pinned host memory reduces transfer overhead. Transfers occur at phase
+boundaries, so cold-start and encoder-to-denoiser transitions become slower.
+
+## Usage
+
+```python
+from vllm_omni import Omni
+
+omni = Omni(
+    model="Wan-AI/Wan2.2-T2V-A14B-Diffusers",
+    enable_cpu_offload=True,
+)
+```
+
+```bash
+vllm serve Wan-AI/Wan2.2-T2V-A14B-Diffusers \
+  --omni --enable-cpu-offload
+```
+
+## Model integration
+
+Pipelines should implement `SupportsComponentDiscovery`:
+
+```python
+from typing import ClassVar
+
+from vllm_omni.diffusion.models.interface import SupportsComponentDiscovery
+
+
+class MyPipeline(nn.Module, SupportsComponentDiscovery):
+    _dit_modules: ClassVar[list[str]] = ["transformer"]
+    _encoder_modules: ClassVar[list[str]] = ["text_encoder", "vision_model"]
+    _vae_modules: ClassVar[list[str]] = ["vae"]
+    _resident_modules: ClassVar[list[str]] = []
+```
+
+All entries may be dotted paths. DiT and encoder lists are both required for
+mutual exclusion. VAE modules are pinned but not swapped; resident modules are
+small modules that must stay on the accelerator for layerwise paths.
+
+## Split-model components
+
+Some models, such as Cosmos3, split one transformer into mutually exclusive
+components that run in different phases. The pipeline exposes
+`enable_omni_model_cpu_offload`, and the backend delegates to the model-local
+contexts:
+
+```python
+class Cosmos3VFMTransformer(nn.Module):
+    def forward(self, ...):
+        with self._offload_context("reasoner"):
+            ...
+        with self._offload_context("generator"):
+            ...
+```
+
+This preserves the same invariant—exactly one component is device resident—
+while reusing sequential `.to()` movers.
+
+## Limitations
+
+- Single device only.
+- Higher cold-start latency.
+- Transfers between encoder and denoising phases add latency.
+
+See the [model-level design](../../../design/feature/offloader/cpu_offload.md)
+for lifecycle and extension invariants.
@@ -0,0 +1,118 @@
+# Distributed Layerwise Offloading
+
+Distributed layerwise offloading (DLO) extends block streaming to multi-device
+deployments. With AllGather enabled, each rank stores roughly `1 / dp_size` of
+the host weights and reconstructs each layer at runtime. Without AllGather,
+each rank streams its standard-loader rank-local weights independently.
+
+See the [DLO feature design](../../../design/feature/offloader/distributed_layerwise_offload.md)
+for the implementation contract and compatibility matrix.
+
+## Execution model
+
+DLO overlaps three operations with a fixed two-block device buffer:
+
+```text
+Compute stream:  [Layer N]          [Layer N+1]        [Layer N+2]
+H2D stream:      [H2D shard N+1]    [H2D shard N+2]
+AllGather:       [AG N+1]           [AG N+2]
+Slots:           slot 0: Layer N    slot 1: Layer N+1
+```
+
+AllGather communicates only request-independent weight shards, so data-
+parallel ranks may process different requests concurrently.
+
+## Usage
+
+```bash
+# Four ranks with sharded host weights and AllGather
+vllm serve /path/to/model --omni \
+  --enable-distributed-layerwise-offload \
+  --data-parallel-size 4
+
+# Standard-loader rank-local weights, without DLO AllGather
+vllm serve /path/to/model --omni \
+  --enable-distributed-layerwise-offload \
+  --data-parallel-size 4 \
+  --dlo-no-use-allgather
+
+# Sequence parallel deployment
+vllm serve /path/to/model --omni \
+  --enable-distributed-layerwise-offload \
+  --usp 4
+```
+
+```python
+from vllm_omni import Omni
+
+omni = Omni(
+    model="/path/to/model",
+    enable_distributed_layerwise_offload=True,
+    dlo_use_allgather=True,
+)
+```
+
+## Flags
+
+| Flag | Meaning | Default |
+| --- | --- | --- |
+| `--enable-distributed-layerwise-offload` | Enable DLO | `false` |
+| `--data-parallel-size N` | DP ranks and AllGather weight-sharding group | `1` |
+| `--dlo-use-allgather` | Shard host weights and reconstruct with AllGather | `true` |
+| `--dlo-no-use-allgather` | Keep standard-loader rank-local weights | `false` |
+| `--dlo-resident-layers N` | Keep N leading main-DiT blocks on device; requires no-AllGather and model-declared resident paths | `0` |
+
+## mmap weight loading
+
+The DLO plus AllGather path:
+
+1. saves non-persistent buffers such as RoPE frequencies;
+2. moves the normally created transformer to the meta device;
+3. loads checkpoint tensors as mmap views backed by the shared OS page cache;
+4. calls model-specific `post_load_weights()` conversions; and
+5. restores the saved non-persistent buffers.
+
+This avoids one full checkpoint RSS copy per rank and does not require
+model-specific loading code.
+
+## Declarative topology
+
+Models may declare an `OffloadPlan` instead of embedding offload logic:
+
+```python
+from vllm_omni.diffusion.offloader import OffloadPlan
+
+
+class MyPipeline(nn.Module):
+    _dit_modules = ["transformer"]
+    _offload_plan = OffloadPlan(
+        block_attrs={"transformer": ("blocks",)},
+        offload_submodules={"context_encoder": "layers"},
+    )
+```
+
+When no plan exists, discovery falls back to
+`_layerwise_offload_blocks_attrs` and then heuristic attribute lookup.
+
+## Data-parallel concurrency
+
+With `data_parallel_size > 1` and AllGather enabled, the scheduler can process
+up to `dp_size` requests per denoising step. Every concurrent request must set
+the same explicit `num_inference_steps`; `None` is rejected because every rank
+must enter each collective.
+
+## Limitations
+
+- Online FP8 quantization is rejected with the DLO plus AllGather mmap path.
+  Use `--dlo-no-use-allgather` or disable online quantization.
+- Tensor parallel size greater than one is rejected in the mmap path because
+  it bypasses TP-aware loader callbacks. The no-AllGather path retains the
+  standard loader and remains experimental with TP.
+- HSDP plus AllGather is rejected to avoid double sharding. HSDP without
+  AllGather has limited end-to-end validation.
+- Resident leading layers require `--dlo-no-use-allgather` and a model
+  `OffloadPlan` that declares eligible `resident_dit_paths`.
+- DP concurrency requires an explicit, identical inference-step count.
+
+See the [Cosmos3 DistOffload recipe](https://github.com/vllm-project/vllm-omni/blob/main/recipes/cosmos3/Cosmos3-DistOffload.md)
+for an end-to-end example.
@@ -0,0 +1,65 @@
+# Layerwise Offloading
+
+Layerwise, or blockwise, offloading keeps one transformer block on the
+accelerator and prefetches the next block while the current block computes.
+It is best suited to compute-heavy video DiTs whose block execution can hide
+host-to-device transfers.
+
+## Execution flow
+
+Each block has a pre-forward and post-forward hook. Parameters are consolidated
+in pinned host tensors and rematerialized for execution on a dedicated copy
+stream.
+
+| Block | Pre-forward hook | Forward | Post-forward hook |
+| --- | --- | --- | --- |
+| block 0 | Prefetch block 1 | Compute block 0 | Free block 0 |
+| block 1 | Prefetch block 2 | Compute block 1 | Free block 1 |
+| ... | ... | ... | ... |
+| last block | Prefetch block 0 | Compute last block | Free last block |
+
+Encoders, VAE modules, and non-block DiT modules such as embeddings and norms
+remain device resident.
+
+## Usage
+
+```python
+from vllm_omni import Omni
+
+omni = Omni(
+    model="Wan-AI/Wan2.2-T2V-A14B-Diffusers",
+    enable_layerwise_offload=True,
+)
+```
+
+```bash
+vllm serve Wan-AI/Wan2.2-T2V-A14B-Diffusers \
+  --omni --enable-layerwise-offload
+```
+
+## Model integration
+
+Transformer classes declare containers of executable blocks:
+
+```python
+class WanTransformer3DModel(nn.Module):
+    _layerwise_offload_blocks_attrs = ["blocks"]
+
+
+class Flux2Transformer2DModel(nn.Module):
+    _layerwise_offload_blocks_attrs = [
+        "transformer_blocks",
+        "single_transformer_blocks",
+    ]
+```
+
+See the [layerwise design](../../../design/feature/offloader/layerwise_offload.md) for
+the discovery and hook invariants. `OffloadPlan` is a separate declarative
+topology path for distributed layerwise offload.
+
+## Limitations
+
+- Single device only.
+- Setup consolidates and pins block parameters, increasing cold-start time.
+- Performance depends on block compute time and host-to-device bandwidth;
+  lightweight blocks may not hide transfers.
@@ -64,6 +64,41 @@ omni = Omni(
 )
 ```

+### Combined with Online FP8 Quantization
+
+HSDP can be combined with online FP8 quantization for diffusion models that
+support both features. HSDP shards the quantized weights across GPUs to reduce
+the per-GPU weight memory footprint.
+
+```python
+from vllm_omni import Omni
+from vllm_omni.diffusion.data import DiffusionParallelConfig
+
+omni = Omni(
+    model="<your-model>",
+    quantization="fp8",
+    parallel_config=DiffusionParallelConfig(
+        use_hsdp=True,
+        hsdp_shard_size=2,
+    ),
+)
+```
+
+For online serving:
+
+```bash
+vllm serve <your-model> \
+    --omni \
+    --quantization fp8 \
+    --use-hsdp \
+    --hsdp-shard-size 2
+```
+
+The model must support both HSDP and online FP8 quantization. HSDP remains
+incompatible with Tensor Parallelism. See the
+[FP8 Quantization Guide](../../quantization/fp8.md) for supported models and
+hardware.
+
 ---

 ## Example Script
@@ -147,3 +182,4 @@ For detailed instructions on adding HSDP support to new models, see the [HSDP Co
 1. ✅ **Enable HSDP** - Set `use_hsdp=True` and `hsdp_shard_size` to reduce per-GPU memory for large models
 2. ✅ **Combine with SP** - Use together with `ulysses_degree` for video models requiring both memory reduction and sequence parallelism
 3. ⚠️ **Incompatible with TP** - `tensor_parallel_size` must be 1 when HSDP is enabled
+4. ✅ **Combine with online FP8** - Use `quantization="fp8"` or `--quantization fp8` to reduce the sharded weight memory footprint
@@ -10,7 +10,7 @@ This guide covers the parallelism methods in vLLM-Omni for speeding up diffusion
 | **[Sequence Parallelism](sequence_parallel.md)**   | Splits sequence dimension across GPUs (Ulysses-SP, Ring-Attention, or hybrid) for high-resolution images and videos |
 | **[CFG-Parallel](cfg_parallel.md)**                | Runs CFG positive/negative branches on separate GPUs for ~1.8x speedup on guided generation                         |
 | **[Pipeline Parallelism](pipeline_parallel.md)**   | Splits the denoising transformer block-wise across sequential GPU stages to reduce per-GPU model memory             |
-| **[VAE Parallel](vae_parallel.md)** | Distributes VAE decode spatially across GPUs to reduce peak VAE memory                                              |
+| **[VAE Parallelism](vae_parallelism.md)** | Distributes VAE decode spatially across GPUs to reduce peak VAE memory                                              |
 | **[HSDP](hsdp.md)**                                | Shards full model weights via PyTorch FSDP2 to enable large-model inference on memory-constrained GPUs              |
 | **[Expert Parallelism](expert_parallel.md)**       | Shards MoE expert blocks across GPUs for MoE models (e.g. HunyuanImage3.0)                                          |

@@ -1,4 +1,4 @@
-# VAE Parallel Guide
+# VAE Parallelism Guide


 ## Table of Content
@@ -113,7 +113,7 @@ In `DiffusionParallelConfig`:
 | Parameter | Type | Default | Description |
 |-----------|------|---------|-------------|
 | `vae_patch_parallel_size` | int | 1 | Number of GPUs for VAE patch/tile parallelism. Set to 2 or higher to enable. Should typically match `tensor_parallel_size` as they share the same process group. |
-| `vae_parallel_mode` | str | `"tile"` | VAE parallel decode strategy: `"tile"`, `"auto"` (Wan request-shape selection), `"spatial_shard_height"`, or `"spatial_shard_width"`. Spatial sharding is currently Wan-only. See [Spatially-Sharded Decode](#spatially-sharded-decode-wan). |
+| `vae_parallel_mode` | str | `"tile"` | VAE parallel decode strategy: `"tile"` (default tile/patch parallel decode), `"spatial_shard_height"`, or `"spatial_shard_width"` (spatially-sharded decode, Wan only). See [Spatially-Sharded Decode](#spatially-sharded-decode-wan). |

 Additional requirements:

@@ -128,7 +128,7 @@ Additional requirements:

 ## Spatially-Sharded Decode (Wan)

-The default `vae_parallel_mode="tile"` keeps patch/tile parallel decode unchanged. For Wan, set `vae_parallel_mode="auto"` to select spatial-shard decode only after a request is already large enough to enter Diffusers' tiled-decode path and the VAE uses the full DiT decode group. Auto mode splits the longer eligible spatial axis to reduce halo traffic. Direct-decode requests and partial groups fall back to tile decode. Use `"spatial_shard_height"` or `"spatial_shard_width"` to force an axis.
+The default `vae_parallel_mode="tile"` distributes whole tiles across ranks. For the **Wan** VAE there is an alternative decode strategy, **spatially-sharded decode**, selected via `vae_parallel_mode="spatial_shard_height"` or `vae_parallel_mode="spatial_shard_width"`.

 Instead of assigning independent tiles to ranks, spatial-shard decode shards the decoder feature maps along the height (`spatial_shard_height`) or width (`spatial_shard_width`) dimension and exchanges halo rows/columns between neighboring ranks around the spatial convolutions. This keeps the receptive field correct across shard boundaries, so the result matches the single-GPU decode within numerical tolerance.

@@ -141,7 +141,7 @@ omni = Omni(
    parallel_config=DiffusionParallelConfig(
        tensor_parallel_size=2,
        vae_patch_parallel_size=2,               # must match the DiT group size
-        vae_parallel_mode="auto",                # opt in to Wan request-shape selection
+        vae_parallel_mode="spatial_shard_width", # or "spatial_shard_height"
    ),
 )
 ```
@@ -152,14 +152,14 @@ Or from the CLI / serving entrypoint:
 vllm serve Wan-AI/Wan2.1-T2V-1.3B-Diffusers --omni \
    --tensor-parallel-size 2 \
    --vae-patch-parallel-size 2 \
-    --vae-parallel-mode auto
+    --vae-parallel-mode spatial_shard_width
 ```

 **Constraints and behavior:**

- Spatial-shard decode is **decode-only** and currently implemented for the **Wan** VAE. Other VAE families retain their existing mode handling and may reject explicit Wan spatial modes.
- It requires `vae_patch_parallel_size` to **match the DiT process group size**. `auto` quietly falls back to tile decode for a partial group; an explicitly forced spatial mode logs a warning.
- The decoder's context-aware wrappers retain the direct path, so one process can alternate small/large and portrait/landscape requests safely.
+- Spatial-shard decode is **decode-only** and currently implemented for the **Wan** VAE. Other models ignore `spatial_shard_*` modes.
+- It requires `vae_patch_parallel_size` to **match the DiT process group size**. If it does not, the VAE logs a warning and **falls back to tile-parallel decode** at runtime.
+- `spatial_shard_height` and `spatial_shard_width` are mutually exclusive for a given VAE instance (the decoder is patched in place for a single split dimension).

 For end-to-end latency/throughput, launch serving with the desired `vae_parallel_mode` and use the existing diffusion serving benchmark:

@@ -61,7 +61,7 @@ Memory optimization methods help reduce GPU memory usage, enabling inference on
 |--------|-------------|----------|
 | **[CPU Offload](diffusion/cpu_offload.md)** | Offloads model components to CPU memory | Limited VRAM, large models on consumer GPUs |
 | **[Quantization](quantization/overview.md)** | Reduces transformer stages from BF16 to FP8/INT8/etc. | Limited VRAM, minimal accuracy loss    |
-| **[VAE Parallel](diffusion/parallelism/vae_parallel.md)** | Distributes VAE decode work across GPUs | High-resolution generation with reduced VAE memory peak |
+| **[VAE Parallelism](diffusion/parallelism/vae_parallelism.md)** | Distributes VAE decode work across GPUs | High-resolution generation with reduced VAE memory peak |

 ### Extensions

@@ -106,7 +106,10 @@ The following tables show which models support each feature:

 > Notes:

-> 1. CPU Offload has two methods: Module-wise (default for models with DiT + text encoder) and Layerwise. The tables below show **Layerwise support** only. Split models like Cosmos3 (no separate text encoder) swap their reasoner/generator components for module-wise offload; see the [CPU Offload Guide](diffusion/cpu_offload.md).
+> 1. CPU Offload has three strategies: model-level, layerwise, and distributed
+>    layerwise. The tables below show **layerwise support** only. Split models
+>    like Cosmos3 swap their reasoner/generator components for model-level
+>    offload; see the [CPU Offload Guide](diffusion/cpu_offload.md).
 > 2. The **💾Quantization** column is collapsed for readability. See [Quantization](quantization/overview.md) for per-method and per-model support details.

 ### ImageGen
@@ -214,7 +217,7 @@ The following tables show which models support each feature:
    3. CPU Offloading (Layerwise) and CPU Offloading (Module-wise) are not compatible.
    4. The CPU Offloading (Layerwise) row describes local layerwise offload.
       Multi-device Distributed Layerwise Offload has a separate topology and
-       compatibility matrix in the [CPU Offloading guide](diffusion/cpu_offload.md#distributed-layerwise-offloading).
+       compatibility matrix in the [Distributed Layerwise Offloading guide](diffusion/offloader/distributed_layerwise_offload.md).
    5. The compatibility matrix uses FP8 as the representative quantization method.
    6. Step Execution is not compatible with any diffusion cache backend. LoRA is supported, but each scheduled batch must use a single adapter (requests with different `lora_request` or `lora_scale` are kept in separate batches).

@@ -165,7 +165,7 @@ vllm serve Qwen/Qwen-Image --omni --port 8091 \

 - **Tensor Parallelism — Text Encoder Not Sharded**: TP currently only shards the DiT blocks. Each TP rank retains a **full copy of the text encoder weights**, leading to significant GPU memory overhead proportional to TP degree. Tracked in [Issue #771](https://github.com/vllm-project/vllm-omni/issues/771).

- **CPU Offloading — Two Modes Are Mutually Exclusive**: Model-level offload (`enable_cpu_offload`) and layerwise offload (`enable_layerwise_offload`) cannot be used simultaneously. If both are set, layerwise takes priority and model-level is silently ignored.
+- **CPU Offloading — Strategies Are Mutually Exclusive**: Distributed layerwise, layerwise, and model-level offload should not be enabled together. Selection priority is distributed layerwise, then layerwise, then model-level.

 - **CPU Offloading — VAE stays on GPU**: Both offloading strategies keep the VAE on GPU at all times. For high-resolution generation, VAE decode can still cause OOM. Mitigate by combining with `vae_use_tiling=True` or VAE Patch Parallelism.

@@ -80,12 +80,14 @@ warmup_quack_fp8([(14040, 2048, 6144), (14040, 2048, 2048)])

 ## Model Type Support

-### Diffusion Model (Qwen-Image, Wan2.2)
+### Diffusion Models

 | Model | HF models | Online | Pre-calibrated | Recommendation | `ignored_layers` | Text-Encoder quantization |
 |-------|-----------|:-------:|:------:|----------------|------------------|------------------|
 | Qwen-Image | `Qwen/Qwen-Image`, `Qwen/Qwen-Image-2512` | Yes | Yes | Skip sensitive image-stream MLPs when quality regresses | `img_mlp` | |
 | Wan2.2 | Wan2.2 diffusion pipelines | Not validated | Not validated | Validate against BF16 before documenting as supported | TBD | |
+| LTX-2 | `Lightricks/LTX-2`, `rootonchair/LTX-2-19b-distilled` | Yes | Not validated | Transformer only; use dynamic phase LoRA for ordinary two-stage | None | |
+| LTX-2.3 | `diffusers/LTX-2.3-Diffusers`, `diffusers/LTX-2.3-Distilled-Diffusers` | Yes | Not validated | Transformer only; use dynamic phase LoRA for ordinary two-stage | None | |
 | Z-Image | `Tongyi-MAI/Z-Image-Turbo` | Yes | Yes | All layers | None | ✅︎ |
 | FLUX.1 | `black-forest-labs/FLUX.1-dev`, `black-forest-labs/FLUX.1-schnell` | Yes | Yes | All layers | None | |
 | FLUX.2-dev | `black-forest-labs/FLUX.2-dev` | Yes | Not validated | All layers | None | ✅︎ |
@@ -23,6 +23,7 @@ list of supported architectures across all modalities, see
 | OmniVoice | `k2-fsa/OmniVoice` | 2 (gen + dec) | ✓ | — | voice design, language hint | 24 kHz |
 | Qwen3-TTS | `Qwen/Qwen3-TTS-12Hz-1.7B-{CustomVoice,VoiceDesign,Base}` | 2 (talker + code2wav) | ✓ (Base) | ✓ | 3 task variants | 24 kHz |
 | VoxCPM2 | `openbmb/VoxCPM2` | single (native AR) | ✓ | ✓ (online) | continuation | 48 kHz |
+| dots.tts | `rednote-hilab/dots.tts-soar` | single (native AR) | — (not wired yet) | — | — | 48 kHz |
 | IndexTTS-2 | `IndexTeam/IndexTTS-2` | 2 (AR talker + S2Mel DiT + BigVGAN) | ✓ (required) | — | emotion control (`--emo-audio`, `--emo-text`, `--emo-vector`) | 22.05 kHz |
 | IndexTTS-2.5 | native `checkpoints/` bundle | 2 (AR talker + EnhancedCodec + S2Mel DiT + BigVGAN) | ✓ (required) | — | multilingual (`--lang`) + emotion control | 22.05 kHz |
 | Voxtral TTS | `mistralai/Voxtral-4B-TTS-2603` | varies | ✓ | ✓ | voice presets | 24 kHz |
@@ -430,6 +431,31 @@ Streaming is exposed through the online OpenAI Speech API (`stream=true`). See [

 ---

+## dots.tts
+
+Single-stage native AR TTS at 48 kHz (rednote-hilab). Pipeline: `Qwen2.5-1.5B base LM → DiT (10-step Euler flow matching) → patch_encoder AR loopback → AudioVAE (streaming decode)`. Same "vLLM-native base LM + side-path computation" pattern as VoxCPM2, with a plain Qwen2 backbone instead of MiniCPM4 and no FSQ / residual LM stage.
+
+### Quick start
+```bash
+python examples/offline_inference/text_to_speech/dots_tts/end2end.py \
+    --model rednote-hilab/dots.tts-soar \
+    --text "Hello, this is a test of dots TTS running on vLLM Omni."
+```
+
+### Voice cloning
+Not wired in this release — generation is zero-shot only. The CAM++ x-vector speaker encoder weights load, but `end2end.py` has no `--ref-audio`/`--ref-text` flags yet.
+
+### Streaming
+The AudioVAE decoder has an internal streaming path (`init_stream_state` / `stream_step` / `stream_flush`) used to avoid boundary artifacts between 160 ms patches, but it is not yet exposed through an online serving endpoint or example.
+
+### Notes
+- Output: 48 kHz mono WAV.
+- Deploy config: `vllm_omni/deploy/dots_tts.yaml` (auto-loaded by HF `model_type`).
+- Checkpoints: `rednote-hilab/dots.tts-soar` is the validated default. `dots.tts-base` shares the same architecture but is unvalidated in this repo. `dots.tts-mf` (MeanFlow, 2-4 step) is not supported yet.
+- Known limitation: no CUDA graph capture and no batched side-path computation yet, so concurrent requests do not currently scale (each request's DiT Euler steps run serially). See `recipes/rednote-hilab/dots.tts.md` for details and the roadmap.
+
+---
+
 ## IndexTTS-2 and IndexTTS-2.5

 Both versions use a 2-stage TTS pipeline and produce 22.05 kHz mono audio. Stage 0 is a GPT AR talker. IndexTTS-2 Stage 1 uses RepCodec semantic embeddings, GPT latent, S2Mel CFM/DiT, and BigVGAN. IndexTTS-2.5 replaces RepCodec with EnhancedCodec and uses the official code-only path (`use_gpt_latent=false`). Every request requires reference audio for zero-shot voice cloning.
@@ -0,0 +1,125 @@
+"""Offline dots.tts inference example (single-stage native AR pipeline).
+
+Uses the single-stage native AR config (vllm_omni/deploy/dots_tts.yaml).
+Reference checkpoint: rednote-hilab/dots.tts-soar.
+"""
+
+from __future__ import annotations
+
+import os
+import time
+from pathlib import Path
+
+import soundfile as sf
+import torch
+
+from vllm_omni import Omni
+from vllm_omni.utils.tracking_parser import TrackingArgumentParser
+
+REPO_ROOT = Path(__file__).resolve().parents[4]
+SAMPLE_RATE = 48_000
+
+
+def parse_args():
+    parser = TrackingArgumentParser(description="Offline dots.tts native AR inference")
+    parser.add_argument(
+        "--model",
+        type=str,
+        default="rednote-hilab/dots.tts-soar",
+        help="dots.tts model path or HuggingFace repo ID.",
+    )
+    parser.add_argument(
+        "--text",
+        type=str,
+        default="Hello, this is a test of dots TTS running on vLLM Omni.",
+        help="Text to synthesize.",
+    )
+    parser.add_argument(
+        "--output-dir",
+        type=str,
+        default="output_audio",
+        help="Directory for output WAV files.",
+    )
+    parser.add_argument(
+        "--deploy-config",
+        type=str,
+        default=None,
+        help="Override the deploy config path. If unset, auto-loads "
+        "vllm_omni/deploy/dots_tts.yaml based on the HF model_type.",
+    )
+    return parser.parse_args()
+
+
+def extract_audio(multimodal_output: dict) -> torch.Tensor:
+    """Extract the final complete audio tensor from multimodal output.
+
+    The output processor concatenates per-step delta tensors under
+    ``model_outputs``.  Falls back to ``audio`` for backwards compat.
+    """
+    audio = multimodal_output.get("model_outputs")
+    if audio is None:
+        audio = multimodal_output.get("audio")
+    if audio is None:
+        raise ValueError(f"No audio key in multimodal_output: {list(multimodal_output.keys())}")
+
+    if isinstance(audio, list):
+        valid = [torch.as_tensor(a).float().cpu().reshape(-1) for a in audio if a is not None]
+        if not valid:
+            raise ValueError("Audio list is empty or all elements are None.")
+        return torch.cat(valid, dim=0) if len(valid) > 1 else valid[0]
+
+    return torch.as_tensor(audio).float().cpu().reshape(-1)
+
+
+def main():
+    args = parse_args()
+
+    output_dir = Path(args.output_dir)
+    output_dir.mkdir(parents=True, exist_ok=True)
+
+    engine = Omni(
+        model=args.model,
+        deploy_config=args.deploy_config,
+    )
+
+    from transformers import AutoTokenizer
+
+    from vllm_omni.model_executor.models.dots_tts.dots_tts_prompt import (
+        build_dots_tts_prompt,
+    )
+
+    tokenizer = AutoTokenizer.from_pretrained(args.model, trust_remote_code=True)
+
+    prompt = build_dots_tts_prompt(
+        tokenizer=tokenizer,
+        text=args.text,
+    )
+
+    print(f"Model       : {args.model}")
+    print(f"Text        : {args.text}")
+    print(f"Output dir  : {output_dir}")
+    print(f"Prompt len  : {len(prompt['prompt_token_ids'])} tokens")
+
+    t_start = time.perf_counter()
+    outputs = engine.generate([prompt])
+    elapsed = time.perf_counter() - t_start
+
+    request_output = outputs[0]
+    mm = request_output.outputs[0].multimodal_output
+    audio = extract_audio(mm)
+
+    duration = audio.numel() / SAMPLE_RATE
+    rtf = elapsed / duration if duration > 0 else float("inf")
+
+    output_path = output_dir / "output.wav"
+    sf.write(str(output_path), audio.numpy(), SAMPLE_RATE, format="WAV")
+
+    print(f"Saved       : {output_path}")
+    print(f"Duration    : {duration:.2f}s")
+    print(f"Inference   : {elapsed:.2f}s")
+    print(f"RTF         : {rtf:.3f}")
+
+
+if __name__ == "__main__":
+    os.environ["VLLM_WORKER_MULTIPROC_METHOD"] = "spawn"
+    main()
@@ -81,6 +81,10 @@ plugins:
        user_guide/quantization/index.md: user_guide/quantization/overview.md
        design/feature/diffusion_request_level_batching.md: design/feature/diffusion_continuous_batching.md
        design/feature/diffusion_step_execution.md: design/feature/diffusion_continuous_batching.md
+        design/feature/cpu_offloading.md: design/feature/offloader/README.md
+        design/feature/model_level_offload.md: design/feature/offloader/cpu_offload.md
+        design/feature/layerwise_offload.md: design/feature/offloader/layerwise_offload.md
+        design/feature/distributed_layerwise_offload.md: design/feature/offloader/distributed_layerwise_offload.md
  - glightbox
  - git-revision-date-localized:
      # exclude files
@@ -40,6 +40,7 @@ dev = [
    "pytest-asyncio==1.4.0",
    "pytest-cov==7.1.0",
    "pytest-mock==3.15.1",
+    "pytest-shard==0.1.2",
    "websockets==16.1",
    "pytest-subtests==0.15.0",
    "datasets==3.6.0",
@@ -301,6 +302,21 @@ filterwarnings = [
    "ignore:.*does not have '__bases__' attribute.*:UserWarning",
 ]

+[tool.coverage.run]
+core = "sysmon"
+source = ["vllm_omni"]
+# Online-serving tests launch the server via subprocess.Popen and stop it with
+# SIGTERM. pytest-cov no longer instruments subprocesses implicitly, and a
+# SIGTERM-ed process discards its data, so without patch/sigterm a coverage run
+# over an online test measures only the pytest parent.
+patch = ["subprocess"]
+concurrency = ["multiprocessing", "thread"]
+sigterm = true
+
+[tool.coverage.report]
+show_missing = true
+skip_covered = true
+
 [tool.typos.default]
 extend-ignore-identifiers-re = [
    ".*_thw",
@@ -4,11 +4,20 @@

 ## Pipelines

-| Model | Checkpoint | Task | `--model-class-name` | Batching |
-|---|---|---|---|---|
-| LTX-2 | `Lightricks/LTX-2` | One-stage T2V/I2V | `LTX2Pipeline` | Yes |
-| LTX-2 distilled | `rootonchair/LTX-2-19b-distilled` | Two-stage T2V/I2V | `LTX2DistilledPipeline` | No |
-| LTX-2.3 | `diffusers/LTX-2.3-Diffusers` | One-stage T2V/I2V | `LTX2Pipeline` | Yes |
+| `--model-class-name` | Task | Required checkpoint repositories |
+|---|---|---|
+| `LTX2Pipeline` | LTX-2 one-stage T2V/I2V | `Lightricks/LTX-2` |
+| `LTX2TwoStagePipeline` | LTX-2 ordinary two-stage T2V/I2V | `Lightricks/LTX-2` |
+| `LTX2DistilledPipeline` | LTX-2 full-distilled two-stage T2V/I2V | `rootonchair/LTX-2-19b-distilled` |
+| `LTX2Pipeline` | LTX-2.3 one-stage T2V/I2V | `diffusers/LTX-2.3-Diffusers` |
+| `LTX2TwoStagePipeline` | LTX-2.3 ordinary two-stage T2V/I2V | `diffusers/LTX-2.3-Diffusers`<br>`Lightricks/LTX-2.3` |
+| `LTX2DistilledPipeline` | LTX-2.3 full-distilled two-stage T2V/I2V | `diffusers/LTX-2.3-Distilled-Diffusers`<br>`Lightricks/LTX-2.3` |
+
+Repositories in the table are download units. A full pipeline repository
+contains the Transformer, text encoder, connectors, VAEs, vocoder, scheduler,
+and tokenizer; an additional repository supplies LoRA or upsampler sidecars.
+The runtime first searches the model root for sidecars under their official
+filenames, then downloads them from the matching Lightricks Hub repository.

 `LTX2Pipeline` is the unified one-stage entry. Checkpoint metadata selects the
 LTX-2 or LTX-2.3 profile; omitting an image selects T2V, while one initial
@@ -16,10 +25,10 @@ image selects I2V. Both one-stage repositories declare this class, so
 `--model-class-name` is optional. LTX-2.3 requires the Diffusers checkpoint;
 the raw `Lightricks/LTX-2.3` safetensors repository is not directly loadable.

-`LTX2DistilledPipeline` is the unified distilled entry: omit `image` for T2V
-or provide one initial image for I2V. As in the official LTX pipelines,
-distilled mode is selected explicitly rather than inferred from the checkpoint
-name or metadata; always pass `--model-class-name LTX2DistilledPipeline`.
+`LTX2TwoStagePipeline` samples the regular model at half resolution, upsamples,
+then refines with the distilled LoRA. `LTX2DistilledPipeline` uses a fully
+merged distilled Transformer in both stages. Both entries support T2V and I2V;
+select their class explicitly.

 ## API Migration

@@ -68,24 +77,28 @@ The model recipe also supplies the default negative prompt. Top-level
 API defaults `num_frames` to `1`, so set it explicitly; the offline LTX scripts
 default to `121`.

-## Distilled Two-Stage Defaults
+## Two-Stage Defaults

-| Parameter | Value |
-|---|---:|
-| Final width × height | 1536 × 1024 |
-| Stage 1 width × height | 768 × 512 |
-| Frames / frame rate | 121 / 24 |
-| Stage 1 / Stage 2 denoise steps | 8 / 3 |
-| Guidance | Fixed positive-only |
+| Parameter | Ordinary | Full-distilled |
+|---|---:|---:|
+| Final width × height | 1536 × 1024 | 1536 × 1024 |
+| Stage 1 width × height | 768 × 512 | 768 × 512 |
+| Frames / frame rate | 121 / 24 | 121 / 24 |
+| Stage 1 / Stage 2 steps | 40 (LTX-2) or 30 (LTX-2.3) / 3 | 8 / 3 |
+| Guidance | Stage 1 guided; Stage 2 positive-only | Fixed positive-only |

-The dimensions passed to the API are the final output dimensions and must be
-divisible by 64. Both one- and two-stage requests require `num_frames = 8k+1`.
-The distilled recipe fixes both sigma schedules; `num_inference_steps` must be
-`8` when supplied.
+API dimensions are final dimensions and must be divisible by 64. All LTX
+requests require `num_frames = 8k+1`. Ordinary Stage 1 uses the LTX-2 or
+LTX-2.3 one-stage defaults shown above. Distilled schedules are fixed, so
+`num_inference_steps`, when supplied, must be `8`. Both entries reject custom
+sigmas and input latents.
+
+Ordinary two-stage uses layer-fused LoRA for an unquantized BF16 Transformer
+and automatically switches to dynamic LoRA when quantization is enabled.

 ## Serving

-Start either checkpoint:
+Start either one-stage checkpoint:

 ```bash
 vllm serve Lightricks/LTX-2 --omni --stage-init-timeout 600
@@ -95,11 +108,19 @@ vllm serve Lightricks/LTX-2 --omni --stage-init-timeout 600
 vllm serve diffusers/LTX-2.3-Diffusers --omni --stage-init-timeout 600
 ```

-Start the distilled checkpoint with its two-stage entry:
+Start a two-stage checkpoint with its explicit class:

 ```bash
 vllm serve rootonchair/LTX-2-19b-distilled --omni \
  --model-class-name LTX2DistilledPipeline --stage-init-timeout 600
+# LTX-2.3 full-distilled; the v1.1 x2 upsampler is resolved when absent
+vllm serve diffusers/LTX-2.3-Distilled-Diffusers --omni \
+  --model-class-name LTX2DistilledPipeline --stage-init-timeout 600
+# Ordinary LTX-2.3 two-stage
+vllm serve diffusers/LTX-2.3-Diffusers --omni \
+  --model-class-name LTX2TwoStagePipeline \
+  --enable-layerwise-offload \
+  --stage-init-timeout 600
 ```

 The same server handles T2V and I2V. A T2V request using the selected recipe's
@@ -135,8 +156,9 @@ together with `input_reference`.

 ## Guidance

-LTX supports independent video/audio CFG, spatio-temporal guidance (STG),
-cross-modality guidance, and rescaling:
+One-stage and ordinary Stage 1 support independent video/audio CFG,
+spatio-temporal guidance (STG), cross-modality guidance, and rescaling.
+Distilled stages and ordinary Stage 2 are fixed positive-only.

 | Parameter | Default | Effect | Alias |
 |---|---:|---|---|
@@ -172,9 +194,9 @@ complete guidance plan in parallel. Despite the legacy `cfg` name, it covers
 text CFG, STG, and cross-modality guidance passes; guidance rescaling is
 applied after all pass predictions are gathered.

-The default one-stage recipe has four Transformer passes per denoise step:
-`cond`, `uncond`, `ptb` (STG), and `mod` (cross-modality). The useful balanced
-configurations are therefore:
+The default one-stage recipe and ordinary Stage 1 have four Transformer passes
+per denoise step: `cond`, `uncond`, `ptb` (STG), and `mod`
+(cross-modality). The useful balanced configurations are therefore:

 | `--cfg-parallel-size` | Passes per rank | Guidance-slot utilization | Notes |
 |---:|---:|---:|---|
@@ -195,9 +217,8 @@ vllm serve diffusers/LTX-2.3-Diffusers --omni \
 ```

 The total device count is the product of `cfg_parallel_size` and the other
-configured parallel dimensions. Distilled pipelines use positive-only
-guidance, so `cfg_parallel_size > 1` does not provide guidance speedup for
-them.
+configured parallel dimensions. Positive-only distilled phases do not benefit
+from `cfg_parallel_size > 1`.

 ### Python API

@@ -242,22 +263,22 @@ noted below.
 | `image` | image or batch, `None` | Direct value wins over request images; no image selects T2V. I2V accepts one image per prompt, and a batch cannot mix T2V/I2V. |
 | `prompt` | string or list, `None` | Positive-text fallback; request prompts win. Mutually exclusive with `prompt_embeds`. |
 | `negative_prompt` | string or list, `None` | Fallback after request values and before the recipe default; mutually exclusive with negative embeddings. |
-| `height` | `int`, `None` | Request → direct value → recipe default; divisible by 32 one-stage or 64 distilled. |
+| `height` | `int`, `None` | Request → direct value → recipe default; divisible by 32 one-stage or 64 two-stage. |
 | `width` | `int`, `None` | Same precedence and alignment as `height`. |
 | `num_frames` | `int`, `None` | Request → direct value → recipe default; must be `8k+1` and also determines audio duration with `frame_rate`. |
 | `frame_rate` | `float`, `None` | Request `frame_rate` → request `fps` → direct value → recipe default. |
-| `num_inference_steps` | `int`, `None` | Request → direct value → recipe default; minimum 2 one-stage, fixed at 8 for distilled Stage 1. Custom one-stage `sigmas` determine actual steps. |
+| `num_inference_steps` | `int`, `None` | Request → direct value → recipe default; minimum 2 one-stage and ordinary Stage 1, fixed at 8 for distilled Stage 1. Custom one-stage `sigmas` determine actual steps. |
 | `sigmas` | list of float, `None` | One-stage only. Request values win; every request in a fused batch must use the same schedule. |
 | `timesteps` | list of int, `None` | Compatibility slot; LTX accepts only `None`. Use `sigmas`. |
-| `guidance_scale` | `float`, `None` | One-stage common video/audio CFG fallback; an explicit request value wins. Distilled two-stage uses fixed positive-only guidance. |
+| `guidance_scale` | `float`, `None` | One-stage and ordinary Stage 1 common video/audio CFG fallback; an explicit request value wins. Distilled phases use fixed positive-only guidance. |
 | `guidance_rescale` | `float`, `None` | Accepts only `None` or `0.0`; use the modality rescale fields. |
 | `noise_scale` | `float`, `0.0` | Compatibility slot; LTX accepts only `0.0`. |
 | `num_videos_per_prompt` | `int`, `1` | Output-count fallback; positive request `num_outputs_per_prompt` wins. |
 | `generator` | generator or list, `None` | Explicit RNG; otherwise request generators/seeds are collated. Lists must match the effective output batch. |
-| `latents` | tensor, `None` | Request tensors win. One-stage accepts packed `[B, S, C]` and validated unpacked video layouts. |
-| `audio_latents` | tensor, `None` | One-stage initial audio latents; request tensors win and are collated. |
+| `latents` | tensor, `None` | One-stage only. Request tensors win; packed `[B, S, C]` and validated unpacked video layouts are accepted. |
+| `audio_latents` | tensor, `None` | One-stage only. Request tensors win and are collated. |
 | `prompt_embeds` | tensor, `None` | Precomputed positive conditioning; requires `prompt_attention_mask` and cannot accompany `prompt`. |
-| `negative_prompt_embeds` | tensor, `None` | One-stage precomputed negative conditioning; requires its mask and cannot accompany `negative_prompt`. |
+| `negative_prompt_embeds` | tensor, `None` | One-stage and ordinary Stage 1 conditioning; requires its mask and cannot accompany `negative_prompt`. |
 | `prompt_attention_mask` | tensor, `None` | Mask for positive embeddings; request `prompt_attention_mask` or `attention_mask` wins. |
 | `negative_prompt_attention_mask` | tensor, `None` | Mask for negative embeddings; request negative mask fields win. |
 | `decode_timestep` | float or list, `0.0` | Video-VAE decode timestep; request value wins. Lists may match 1, prompt batch, or output batch. |
@@ -273,13 +294,13 @@ request prompt payload; LTX guidance fields live in sampling `extra_args`.

 ### Recipe-Specific Request Capabilities

-| Override | One-stage | Distilled two-stage |
-|---|---|---|
-| Guidance | Supported | Fixed positive-only |
-| Negative prompt/embeddings | Supported | Rejected |
-| `num_inference_steps` | Supported | Fixed at 8 for Stage 1; Stage 2 uses 3 |
-| Custom `sigmas` | Supported | Rejected; both phases use fixed schedules |
-| Video/audio latents | Supported | Rejected |
+| Override | One-stage | Ordinary two-stage | Distilled two-stage |
+|---|---|---|---|
+| Guidance | Supported | Stage 1 only; Stage 2 is positive-only | Fixed positive-only |
+| Negative prompt/embeddings | Supported | Supported by Stage 1 | Rejected |
+| `num_inference_steps` | Supported | Controls Stage 1; Stage 2 uses 3 | Fixed at 8 for Stage 1; Stage 2 uses 3 |
+| Custom `sigmas` | Supported | Rejected | Rejected; both phases use fixed schedules |
+| Video/audio latents | Supported | Rejected | Rejected |

 These capability checks apply equally to direct `forward` keywords and
 values in `OmniDiffusionSamplingParams`; unsupported values fail instead of
@@ -309,6 +330,10 @@ bundled offline CLI do not currently expose `sigmas`.
  `--max-num-seqs 1`.
 - `--cfg-parallel-size` shards the complete LTX guidance plan, including STG,
  modality guidance, and rescale-compatible prediction gathering.
+- Sequence parallelism may pad audio latents. Pure Ulysses masks the padding;
+  Ring cannot, so audio length must be SP-divisible. Use `ring_degree=1` or a
+  divisible request shape.
+- Cache-DiT is one-stage only; multi-stage configurations are rejected.

 ## Operational Notes

@@ -323,14 +348,15 @@ bundled offline CLI do not currently expose `sigmas`.
  request parameter.
 - For benchmarks, use `tests/dfx/perf/tests/test_ltx2_vllm_omni.json` with
  `tests/dfx/perf/scripts/run_diffusion_benchmark.py`.
- Current two-stage support is limited to the distilled checkpoint; official
-  LTX two-stage/HQ execution remains out of scope.
+- Ordinary and full-distilled two-stage T2V/I2V are supported; HQ execution
+  remains out of scope.

 ## References

 - <https://huggingface.co/Lightricks/LTX-2>
 - <https://huggingface.co/Lightricks/LTX-2.3>
 - <https://huggingface.co/diffusers/LTX-2.3-Diffusers>
+- <https://huggingface.co/diffusers/LTX-2.3-Distilled-Diffusers>
 - [Online video generation](../../docs/user_guide/examples/online_serving/text_to_video.md)
 - [Diffusion execution modes](../../docs/user_guide/diffusion/execution_modes.md)
 - [T2V offline example](../../examples/offline_inference/text_to_video/text_to_video.md)
@@ -0,0 +1,152 @@
+# dots.tts
+
+> Offline continuous-AR TTS at 48 kHz (rednote-hilab)
+
+## Summary
+
+- Vendor: rednote-hilab
+- Model: `rednote-hilab/dots.tts-soar`
+- Task: Text-to-speech, zero-shot synthesis only
+- Mode: Offline end-to-end example only (not yet wired for online serving —
+  see [Known limitations](#known-limitations))
+- Maintainer: Community
+
+## When to use this recipe
+
+Use this recipe as a known-good starting point for running
+`rednote-hilab/dots.tts-soar` offline on vLLM-Omni on consumer-class GPUs.
+dots.tts is a ~1.7B-parameter continuous-AR TTS model (Qwen2.5-1.5B base LM
+ 344M DiT flow-matching head + 180M AudioVAE) that emits 48 kHz mono audio.
+It follows the same "vLLM-native base LM + side-path computation" pattern as
+VoxCPM2 — single-stage pipeline
+`Qwen2.5-1.5B base LM → DiT (10-step Euler flow matching) → patch_encoder AR
+loopback → AudioVAE (streaming decode)` — with a plain Qwen2 backbone
+instead of MiniCPM4, and no FSQ / residual-LM stage.
+
+This is an early integration (see [Known limitations](#known-limitations)
+below before depending on it for anything beyond offline experimentation).
+
+## References
+
+- Offline end-to-end script:
+  [`examples/offline_inference/text_to_speech/dots_tts/end2end.py`](../../examples/offline_inference/text_to_speech/dots_tts/end2end.py)
+- Example guide:
+  [`examples/offline_inference/text_to_speech/README.md`](../../examples/offline_inference/text_to_speech/README.md#dotstts)
+- Default deploy config:
+  [`vllm_omni/deploy/dots_tts.yaml`](../../vllm_omni/deploy/dots_tts.yaml)
+- Talker / pipeline source:
+  [`vllm_omni/model_executor/models/dots_tts/`](../../vllm_omni/model_executor/models/dots_tts/)
+- Upstream: [rednote-hilab/dots.tts](https://github.com/rednote-hilab/dots.tts)
+
+## Hardware Support
+
+This recipe documents one tested 16 GB consumer-GPU configuration. Other
+vendor sections (ROCm, NPU) and larger-VRAM configurations are welcome as
+community validation lands.
+
+## GPU
+
+### 1 x RTX 5080 16GB (Single GPU, Minimum Recommended)
+
+dots.tts (~1.7B params across the base LM + DiT + AudioVAE + CAM++
+speaker encoder, bfloat16) fits comfortably on a single 16 GB GPU. The
+bundled default config at
+[`vllm_omni/deploy/dots_tts.yaml`](../../vllm_omni/deploy/dots_tts.yaml)
+(`gpu_memory_utilization: 0.8`, `max_num_seqs: 4`, `enforce_eager: true`,
+`enable_prefix_caching: false`) loads cleanly with ~5.1 GiB for model
+weights and ~0.3 GiB peak activation; the remainder of the configured
+budget is available for KV cache. Total resident footprint at idle is
+roughly **7-8 GiB / 16 GB** — the only tight spot in the full CUDA-Graph
+roadmap would be step 8's graph capture (not implemented yet; this
+release runs `enforce_eager: true`, so it doesn't apply today).
+
+#### Environment
+
+- OS: Linux (WSL2)
+- Python: 3.12
+- Driver / runtime: NVIDIA driver 595.95
+- torch: 2.11.0+cu130
+- vLLM: 0.26.0
+- vLLM-Omni: 0.22.1.dev (current `main`)
+
+#### Command
+
+```bash
+python examples/offline_inference/text_to_speech/dots_tts/end2end.py \
+    --model rednote-hilab/dots.tts-soar \
+    --text "Hello, this is a test of dots TTS running on vLLM Omni."
+```
+
+The deploy config at
+[`vllm_omni/deploy/dots_tts.yaml`](../../vllm_omni/deploy/dots_tts.yaml)
+is loaded automatically by the model registry (HF `model_type=dots_tts`).
+Pass `--deploy-config <path>` to override.
+
+#### Verification
+
+**T1 — offline zero-shot synthesis**:
+
+```bash
+python examples/offline_inference/text_to_speech/dots_tts/end2end.py \
+    --model rednote-hilab/dots.tts-soar \
+    --text "Hello, this is a test of dots TTS running on vLLM Omni."
+```
+
+Observed: `output_audio/output.wav`, 3.52 s @ 48 kHz mono. Single
+one-shot process (init → one `generate()` → exit), so the reported
+numbers include engine init and first-request warmup, not just steady-state
+per-step throughput — same caveat as VoxCPM2's recipe. `Inference: 5.52s`,
+`RTF: 1.569`.
+
+Weight-loading breakdown from the same run (all tensors matched, no
+missing/extra keys): 951/951 AudioVAE, 244/244 DiT, 270/270
+patch_encoder, 198 Qwen2, 938/938 CAM++ speaker encoder.
+
+Whisper transcription of the output matched the input text with no
+dropped leading word (confirms the streaming-vocoder patch-boundary fix
+described in [Known limitations](#known-limitations)).
+
+#### Notes
+
+- Output: 48 kHz mono WAV.
+- Checkpoints: `rednote-hilab/dots.tts-soar` is the validated default
+  used throughout this recipe. `rednote-hilab/dots.tts-base` shares the
+  same architecture but is unvalidated in this repo. `rednote-hilab/dots.tts-mf`
+  (MeanFlow, 2-4 step) is not supported — see below.
+- `enforce_eager: true` and `enable_prefix_caching: false` in the deploy
+  config are load-bearing, not just conservative defaults: with prefix
+  caching enabled, vLLM-Omni's prefix-cache multimodal-output merge path
+  does not preserve the `sparse_audio` marker this model relies on to
+  route audio output correctly, and generation silently truncates to a
+  single ~160 ms patch. Do not override `enable_prefix_caching` for this
+  model until that framework-level gap is fixed.
+
+## Known limitations
+
+- **No online serving yet.** dots.tts is not registered in
+  `vllm_omni/entrypoints/openai/tts_adapters/` (the framework other TTS
+  models in this repo use for `/v1/audio/speech`), so `vllm serve
+  rednote-hilab/dots.tts-soar --omni` will start but the OpenAI-compatible
+  speech endpoint does not know how to build request params for it yet.
+  Offline `Omni()` (as used in `end2end.py`) is the only supported path
+  today.
+- **Voice cloning is not wired.** The CAM++ x-vector speaker encoder
+  weights load and are exercised by `load_weights()`, but nothing in the
+  prompt builder consumes a reference audio yet — generation is zero-shot
+  only. `end2end.py` has no `--ref-audio`/`--ref-text` flags.
+- **`dots.tts-mf` (MeanFlow, 2-4 step) checkpoint is not supported.** Only
+  the fixed 10-step Euler DiT sampler used by `dots.tts-soar` /
+  `dots.tts-base` is implemented.
+- **No CUDA graph capture.** The talker runs fully eager. voxcpm2's three
+  captured graphs (base LM decode, CFM solver, VAE decode) have no
+  dots.tts equivalent yet.
+- **Concurrent requests do not scale.** Each request's 10-step DiT Euler
+  integration runs serially in the side path (no cross-request batching,
+  unlike voxcpm2's `enable_batched_cfm`). A community review of this
+  integration measured no throughput gain at `c=4` concurrent requests
+  versus `c=1`.
+- **`SamplingParams.seed` does not control audio-generation randomness.**
+  The DiT's flow-matching noise is deterministically derived per-request
+  from a fixed internal seed (reproducible run-to-run, matching
+  voxcpm2's `deterministic_cfm_seed` convention) rather than from the
+  caller-supplied `seed` field — the same limitation voxcpm2 has today.
@@ -5,6 +5,7 @@ import sys
 from pathlib import Path

 import pytest
+import yaml

 sys.path.insert(0, str(Path(__file__).resolve().parents[2] / ".buildkite" / "common" / "scripts"))

@@ -139,3 +140,48 @@ def test_mirror_hardwares_a2b3_npu_4_expands_agents_image_and_plugins() -> None:
    assert step["plugins"][0]["kubernetes"]["podSpecPatch"]["imagePullSecrets"] == [
        {"name": "swr-secret"},
    ]
+
+
+COVERAGE_PILOT_LABELS = ("Diffusion · Bagel Test", "TTS · Qwen3-TTS Base Test")
+MERGE_PIPELINE = Path(".buildkite/cuda/test-merge.yml")
+
+
+def _surviving_labels(changed_files: list[str]) -> set[str]:
+    doc = yaml.safe_load(MERGE_PIPELINE.read_text(encoding="utf-8"))
+    rendered = _render_test_pipeline(doc, changed_files=changed_files)
+    labels: set[str] = set()
+
+    def walk(steps: list | None) -> None:
+        for step in steps or []:
+            if not isinstance(step, dict):
+                continue
+            if "label" in step:
+                labels.add(step["label"])
+            walk(step.get("steps"))
+
+    walk(rendered.get("steps"))
+    return labels
+
+
+@pytest.mark.parametrize(
+    "changed_file",
+    [
+        ".buildkite/common/scripts/run_cov_split.sh",
+        "pyproject.toml",
+    ],
+)
+def test_coverage_pilots_selected_by_their_shared_inputs(changed_file: str) -> None:
+    """The split helper and the coverage config both change what these jobs measure.
+
+    Without them in source_file_dependencies the jobs are filtered out of normal
+    PR builds, so a regression would only surface in a later nightly.
+    """
+    labels = _surviving_labels([changed_file])
+    for pilot in COVERAGE_PILOT_LABELS:
+        assert pilot in labels, f"{pilot} dropped when {changed_file} changed"
+
+
+def test_coverage_pilots_not_selected_by_unrelated_change() -> None:
+    labels = _surviving_labels(["vllm_omni/entrypoints/openai/serving_chat.py"])
+    for pilot in COVERAGE_PILOT_LABELS:
+        assert pilot not in labels
@@ -471,6 +471,7 @@ def test_sub_config_fields_match_structured_scopes():
        "trust_remote_code",
        "dtype",
        "attention_backend",
+        "attention_config",
        "moe_backend",
        "hf_overrides",
        "limit_mm_per_prompt",
@@ -582,12 +583,8 @@ def test_diffusion_parallel_config_keeps_current_diffusion_parallel_surface():
    assert cfg.data_parallel_size == 3
    assert cfg.cfg_parallel_size == 3
    assert cfg.mask_sp_padding is True
-    assert cfg.vae_parallel_mode == "tile"
    assert cfg.world_size == 72

-    auto_cfg = OmniStageDiffusionParallelConfig(vae_parallel_mode="auto")
-    assert auto_cfg.vae_parallel_mode == "auto"
-

 def test_parallel_config_derived_fields_are_not_init_inputs():
    with pytest.raises(ValidationError):
@@ -60,9 +60,18 @@ def _dense_attention(query: torch.Tensor, key: torch.Tensor, value: torch.Tensor
    return torch.einsum("bhqk,bkhd->bqhd", probs, value)


-def _cuda_flash_attn_usable() -> bool:
+def _gpu_flash_attn_usable() -> bool:
    if not torch.cuda.is_available():
        return False
+    if torch.version.hip is not None:
+        for module in ("aiter", "flash_attn"):
+            try:
+                imported = __import__(module, fromlist=["flash_attn_varlen_func"])
+                if getattr(imported, "flash_attn_varlen_func", None) is not None:
+                    return True
+            except ImportError:
+                pass
+        return False
    try:
        spec = find_spec("vllm.vllm_flash_attn")
        if spec is None or spec.origin is None:
@@ -227,12 +236,11 @@ def test_paged_attention_matches_dense_reference_cpu(history_chunks, action_len,
    assert st.adapter(POS).completed_chunks == before + (1 if commit_current else 0)


-@pytest.mark.skipif(not _cuda_flash_attn_usable(), reason="usable CUDA FlashAttention is required")
+@pytest.mark.skipif(not _gpu_flash_attn_usable(), reason="usable GPU FlashAttention is required")
@pytest.mark.parametrize("history_chunks", [1, 3])
@pytest.mark.parametrize("action_len", [0, 3])
@pytest.mark.parametrize("commit_current", [False, True])
 def test_paged_attention_matches_dense_reference_gpu(history_chunks, action_len, commit_current):
-    pytest.importorskip("vllm.vllm_flash_attn")
    torch.manual_seed(0)
    device = torch.device("cuda")
    dtype = torch.float16
@@ -265,10 +265,15 @@ def test_fa_vs_sdpa():
    print(f"Max absolute difference (valid region): {max_diff:.6f}")
    print(f"Mean absolute difference (valid region): {mean_diff:.6f}")

-    # Assert that outputs are close
-    # Using higher tolerance for bfloat16 and different implementations
-    assert max_diff < 0.01, f"Max difference {max_diff} exceeds threshold 0.01"
-    assert mean_diff < 0.001, f"Mean difference {mean_diff} exceeds threshold 0.001"
+    # AITER's ROCm BF16 kernel is deterministic, but differs from ROCm SDPA by
+    # about 1.3% relative error. An absolute-only bound misclassifies that
+    # expected rounding difference for larger output values (for example,
+    # -2.65625 vs -2.6875). Keep the tighter historical bounds elsewhere.
+    if current_omni_platform.is_rocm():
+        torch.testing.assert_close(output_fa_valid, output_sdpa_valid, rtol=1e-2, atol=1.5e-2)
+    else:
+        assert max_diff < 0.01, f"Max difference {max_diff} exceeds threshold 0.01"
+        assert mean_diff < 0.001, f"Mean difference {mean_diff} exceeds threshold 0.001"

    print("✓ Case 2 PASSED: FA and SDPA outputs are very close!")

@@ -416,8 +416,8 @@ def test_sp_correctness_advanced(model_name: str):


@pytest.mark.skipif(
-    not (current_omni_platform.is_cuda() or current_omni_platform.is_xpu()),
-    reason="Only tested on CUDA and XPU",
+    not (current_omni_platform.is_cuda() or current_omni_platform.is_rocm() or current_omni_platform.is_xpu()),
+    reason="Only tested on CUDA, ROCM and XPU",
 )
@pytest.mark.diffusion
@pytest.mark.parallel
@@ -447,8 +447,8 @@ def test_local_sp_padding_mask(monkeypatch: pytest.MonkeyPatch) -> None:


@pytest.mark.skipif(
-    not (current_omni_platform.is_cuda() or current_omni_platform.is_xpu()),
-    reason="Only tested on CUDA and XPU",
+    not (current_omni_platform.is_cuda() or current_omni_platform.is_rocm() or current_omni_platform.is_xpu()),
+    reason="Only tested on CUDA, ROCM and XPU",
 )
@pytest.mark.diffusion
@pytest.mark.parallel
@@ -468,8 +468,8 @@ def test_local_sp_padding_mask_no_padding(monkeypatch: pytest.MonkeyPatch) -> No


@pytest.mark.skipif(
-    not (current_omni_platform.is_cuda() or current_omni_platform.is_xpu()),
-    reason="Only tested on CUDA and XPU",
+    not (current_omni_platform.is_cuda() or current_omni_platform.is_rocm() or current_omni_platform.is_xpu()),
+    reason="Only tested on CUDA, ROCM and XPU",
 )
@pytest.mark.diffusion
@pytest.mark.parallel
@@ -1,5 +1,4 @@
 import os
-from contextlib import contextmanager
 from types import SimpleNamespace

 import pytest
@@ -11,7 +10,6 @@ from tests.helpers.mark import hardware_test
 from vllm_omni.diffusion.distributed.autoencoders import wan_spatial_shard
 from vllm_omni.diffusion.distributed.autoencoders.autoencoder_kl_wan import (
    DistributedAutoencoderKLWan,
-    select_auto_spatial_shard_split_dim,
 )
 from vllm_omni.platforms import current_omni_platform

@@ -19,23 +17,6 @@ from vllm_omni.platforms import current_omni_platform
 # the end uses hardware_test (H100 x2) plus full_model / diffusion / parallel.


-@contextmanager
-def _active_spatial_shard(split_dim: str):
-    token = wan_spatial_shard._SPATIAL_SHARD_CONTEXT.set(
-        wan_spatial_shard.SpatialShardContext(
-            input_extent=2,
-            local_input_extent=2,
-            split_dim=split_dim,
-            rank=0,
-            world_size=3,
-        )
-    )
-    try:
-        yield
-    finally:
-        wan_spatial_shard._SPATIAL_SHARD_CONTEXT.reset(token)
-
-
@pytest.mark.core_model
@pytest.mark.cpu
 def test_split_for_parallel_decode_pads_uneven_height():
@@ -351,7 +332,7 @@ def test_attention_wrapper_rejects_spatial_extent_change(monkeypatch: pytest.Mon
        lambda x, *, group, dim, dst=None: torch.cat([x, x], dim=dim),
    )
    module = ShrinkingWanAttentionBlock()
-    wan_spatial_shard._patch_attention_block(module, group=object())
+    wan_spatial_shard._patch_attention_block(module, group=object(), split_dim="width")
    x = torch.zeros((1, 1, 1, 1, 2), dtype=torch.float32)
    token = wan_spatial_shard._SPATIAL_SHARD_CONTEXT.set(
        wan_spatial_shard.SpatialShardContext(
@@ -396,14 +377,12 @@ def test_dist_zero_pad_only_applies_global_height_edges(monkeypatch: pytest.Monk

    monkeypatch.setattr(wan_spatial_shard, "_rank_world", lambda group: (1, 3))
    mid_rank_pad = wan_spatial_shard.WanDistZeroPad2d((0, 1, 1, 1), group=object())
-    with _active_spatial_shard("height"):
-        mid = mid_rank_pad(x)
+    mid = mid_rank_pad(x)
    assert mid.shape == (1, 1, 2, 3)

    monkeypatch.setattr(wan_spatial_shard, "_rank_world", lambda group: (2, 3))
    last_rank_pad = wan_spatial_shard.WanDistZeroPad2d((0, 1, 1, 1), group=object())
-    with _active_spatial_shard("height"):
-        last = last_rank_pad(x)
+    last = last_rank_pad(x)
    assert last.shape == (1, 1, 3, 3)


@@ -413,15 +392,13 @@ def test_dist_zero_pad_only_applies_global_width_edges(monkeypatch: pytest.Monke
    x = torch.ones((1, 1, 2, 2))

    monkeypatch.setattr(wan_spatial_shard, "_rank_world", lambda group: (1, 3))
-    mid_rank_pad = wan_spatial_shard.WanDistZeroPad2d((1, 1, 0, 0), group=object())
-    with _active_spatial_shard("width"):
-        mid = mid_rank_pad(x)
+    mid_rank_pad = wan_spatial_shard.WanDistZeroPad2d((1, 1, 0, 0), group=object(), split_dim="width")
+    mid = mid_rank_pad(x)
    assert mid.shape == (1, 1, 2, 2)

    monkeypatch.setattr(wan_spatial_shard, "_rank_world", lambda group: (2, 3))
-    last_rank_pad = wan_spatial_shard.WanDistZeroPad2d((1, 1, 0, 0), group=object())
-    with _active_spatial_shard("width"):
-        last = last_rank_pad(x)
+    last_rank_pad = wan_spatial_shard.WanDistZeroPad2d((1, 1, 0, 0), group=object(), split_dim="width")
+    last = last_rank_pad(x)
    assert last.shape == (1, 1, 2, 3)


@@ -444,111 +421,6 @@ def test_spatial_shard_height_gate_falls_back_for_partial_group(monkeypatch: pyt
    assert vae._spatial_shard_decode_enabled(z) is False


-@pytest.mark.core_model
-@pytest.mark.cpu
-@pytest.mark.parametrize(
-    "z_shape,world_size,expected",
-    [
-        ((1, 16, 1, 60, 104), 4, "width"),
-        ((1, 16, 5, 60, 104), 4, "width"),
-        ((1, 16, 5, 104, 60), 4, "height"),
-        ((1, 16, 5, 2, 2), 4, None),
-        ((1, 16, 5, 64, 64), 4, "width"),
-        ((1, 16, 64, 64), 4, None),
-        ((1, 16, 5, 64, 64), 1, None),
-    ],
-)
-def test_auto_spatial_shard_selector_is_shape_aware(z_shape, world_size, expected):
-    assert select_auto_spatial_shard_split_dim(z_shape, world_size) == expected
-
-
-@pytest.mark.core_model
-@pytest.mark.cpu
-def test_auto_spatial_shard_gate_requires_full_group(monkeypatch: pytest.MonkeyPatch):
-    monkeypatch.setattr(
-        "vllm_omni.diffusion.distributed.autoencoders.autoencoder_kl_wan.dist.get_world_size",
-        lambda group=None: 4,
-    )
-
-    vae = DistributedAutoencoderKLWan.__new__(DistributedAutoencoderKLWan)
-    vae.use_tiling = True
-    vae.distributed_executor = SimpleNamespace(group=object(), parallel_size=4, parallel_mode="auto")
-    vae.is_distributed_enabled = lambda: True
-    large = torch.zeros((1, 16, 5, 60, 104))
-    unshardable = torch.zeros((1, 16, 1, 2, 2))
-    unsupported = torch.zeros((1, 16, 60, 104))
-
-    assert vae._spatial_shard_decode_split_dim(large, 4) == "width"
-    assert vae._spatial_shard_decode_enabled(large) is True
-    assert vae._spatial_shard_decode_enabled(unshardable) is False
-    assert vae._spatial_shard_decode_enabled(unsupported) is False
-
-    vae.distributed_executor.parallel_size = 2
-    assert vae._spatial_shard_decode_enabled(large) is False
-
-
-@pytest.mark.core_model
-@pytest.mark.cpu
-def test_set_parallel_size_prepares_spatial_wrappers_only_for_a_full_group(monkeypatch: pytest.MonkeyPatch):
-    vae = DistributedAutoencoderKLWan.__new__(DistributedAutoencoderKLWan)
-    executor = SimpleNamespace(group=object(), parallel_size=1, parallel_mode="tile")
-
-    def set_parallel_size(parallel_size: int, mode: str = "tile") -> None:
-        executor.parallel_size = parallel_size
-        executor.parallel_mode = mode
-
-    executor.set_parallel_size = set_parallel_size
-    object.__setattr__(vae, "distributed_executor", executor)
-    installed = []
-
-    monkeypatch.setattr(
-        "vllm_omni.diffusion.distributed.autoencoders.autoencoder_kl_wan.dist.get_world_size",
-        lambda group=None: 4,
-    )
-    monkeypatch.setattr(
-        wan_spatial_shard,
-        "install_wan_spatial_shard_decode",
-        lambda candidate, group, split_dim: installed.append((candidate, group, split_dim)),
-    )
-
-    vae.set_parallel_size(2, mode="auto")
-    assert installed == []
-
-    vae.set_parallel_size(4, mode="auto")
-    assert installed == [(vae, vae.distributed_executor.group, "height")]
-
-
-@pytest.mark.core_model
-@pytest.mark.cpu
-@torch.no_grad()
-def test_installed_spatial_decoder_keeps_direct_fallback_and_can_change_axis():
-    vae = DistributedAutoencoderKLWan(
-        base_dim=8,
-        decoder_base_dim=8,
-        z_dim=4,
-        dim_mult=[1, 2],
-        num_res_blocks=1,
-        attn_scales=[],
-        temperal_downsample=[False],
-        latents_mean=[0.0] * 4,
-        latents_std=[1.0] * 4,
-        scale_factor_temporal=1,
-        scale_factor_spatial=2,
-    ).eval()
-    latents = torch.randn(1, 4, 1, 4, 5)
-    parameter_names = {name for name, _ in vae.named_parameters()}
-    state_dict = {name: value.clone() for name, value in vae.state_dict().items()}
-    reference = vae.decode(latents, return_dict=False)[0]
-
-    wan_spatial_shard.install_wan_spatial_shard_decode(vae, group=object(), split_dim="height")
-    assert any(getattr(module, "_vllm_omni_dynamic_spatial_shard_conv", False) for module in vae.decoder.modules())
-    wan_spatial_shard.install_wan_spatial_shard_decode(vae, group=object(), split_dim="width")
-
-    assert {name for name, _ in vae.named_parameters()} == parameter_names
-    vae.load_state_dict(state_dict, strict=True)
-    assert torch.equal(vae.decode(latents, return_dict=False)[0], reference)
-
-
@pytest.mark.core_model
@pytest.mark.cpu
 def test_spatial_shard_width_gate_selects_width(monkeypatch: pytest.MonkeyPatch):
@@ -694,99 +566,3 @@ def test_spatial_shard_decode_matches_reference(split_dim: str):
    assert mean_abs_diff <= _SPATIAL_SHARD_TOLERANCE, (
        f"spatial_shard_{split_dim} mean_abs_diff {mean_abs_diff} exceeds {_SPATIAL_SHARD_TOLERANCE}"
    )
-
-
-def _auto_spatial_shard_decode_worker(rank: int, return_dict, master_port: str) -> None:
-    from vllm_omni.diffusion.distributed.parallel_state import (
-        destroy_model_parallel,
-        init_distributed_environment,
-        initialize_model_parallel,
-    )
-
-    os.environ["MASTER_ADDR"] = "127.0.0.1"
-    os.environ["MASTER_PORT"] = master_port
-    device = current_omni_platform.get_torch_device(rank)
-    current_omni_platform.set_device(device)
-    backend = current_omni_platform.dist_backend
-    init_distributed_environment(world_size=2, rank=rank, local_rank=rank, backend=backend)
-    initialize_model_parallel(sequence_parallel_size=2, ulysses_degree=2, backend=backend)
-
-    try:
-        vae = DistributedAutoencoderKLWan.from_pretrained(
-            _SPATIAL_SHARD_MODEL,
-            subfolder=_SPATIAL_SHARD_SUBFOLDER,
-            torch_dtype=torch.float32,
-        ).to(device=device, dtype=torch.float32)
-        vae.eval()
-        generator = torch.Generator(device=device).manual_seed(1)
-        small = torch.randn((1, vae.config.z_dim, 1, 60, 104), generator=generator, device=device)
-        landscape = torch.randn((1, vae.config.z_dim, 2, 60, 104), generator=generator, device=device)
-        portrait = torch.randn((1, vae.config.z_dim, 2, 104, 60), generator=generator, device=device)
-
-        with torch.inference_mode():
-            vae.use_tiling = False
-            vae.set_parallel_size(1, mode="tile")
-            small_reference = vae.decode(small, return_dict=False)[0]
-
-            vae.use_tiling = True
-            vae.set_parallel_size(2, mode="tile")
-            tile_before_install = vae.decode(small, return_dict=False)[0]
-
-            vae.set_parallel_size(2, mode="auto")
-            assert vae._spatial_shard_decode_split_dim(small, 2) == "width"
-            small_sharded = vae.decode(small, return_dict=False)[0]
-            assert vae._spatial_shard_decode_split_dim(landscape, 2) == "width"
-            vae.decode(landscape, return_dict=False)
-
-            vae.set_parallel_size(2, mode="tile")
-            tile_after_install = vae.decode(small, return_dict=False)[0]
-
-            vae.use_tiling = False
-            vae.set_parallel_size(1, mode="tile")
-            portrait_reference = vae.decode(portrait, return_dict=False)[0]
-
-            vae.use_tiling = True
-            vae.set_parallel_size(2, mode="auto")
-            assert vae._spatial_shard_decode_split_dim(portrait, 2) == "height"
-            portrait_sharded = vae.decode(portrait, return_dict=False)[0]
-
-        assert all(
-            buffer is None
-            for module in vae.decoder.modules()
-            for name, buffer in module._buffers.items()
-            if name in {"_halo_recv_top_buf", "_halo_recv_bottom_buf"}
-        )
-
-        if rank == 0:
-            small_diff = (small_sharded.float() - small_reference.float()).abs()
-            portrait_diff = (portrait_sharded.float() - portrait_reference.float()).abs()
-            return_dict["tile_fallback_exact"] = torch.equal(tile_after_install, tile_before_install)
-            return_dict["small_max_abs_diff"] = small_diff.max().item()
-            return_dict["small_mean_abs_diff"] = small_diff.mean().item()
-            return_dict["portrait_max_abs_diff"] = portrait_diff.max().item()
-            return_dict["portrait_mean_abs_diff"] = portrait_diff.mean().item()
-    finally:
-        destroy_model_parallel()
-        if dist.is_initialized():
-            dist.destroy_process_group()
-
-
-@pytest.mark.full_model
-@pytest.mark.diffusion
-@pytest.mark.parallel
-@hardware_test(res={"cuda": "H100"}, num_cards=2)
-def test_auto_spatial_shard_switches_shape_paths_safely():
-    manager = mp.get_context("spawn").Manager()
-    return_dict = manager.dict()
-    mp.spawn(
-        _auto_spatial_shard_decode_worker,
-        args=(return_dict, "29502"),
-        nprocs=2,
-        join=True,
-    )
-
-    assert return_dict.get("tile_fallback_exact") is True
-    assert return_dict["small_max_abs_diff"] <= _SPATIAL_SHARD_TOLERANCE
-    assert return_dict["small_mean_abs_diff"] <= _SPATIAL_SHARD_TOLERANCE
-    assert return_dict["portrait_max_abs_diff"] <= _SPATIAL_SHARD_TOLERANCE
-    assert return_dict["portrait_mean_abs_diff"] <= _SPATIAL_SHARD_TOLERANCE
@@ -0,0 +1,82 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import torch
+import torch.nn.functional as F
+from vllm.triton_utils import HAS_TRITON
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cuda, pytest.mark.diffusion]
+
+_HEAD_DIM = 128
+_ROTARY_DIM = 96
+_EPS = 1e-5
+
+
+def _reference(q, k, q_weight, k_weight, rope_table):
+    q = F.rms_norm(q, (_HEAD_DIM,), q_weight, _EPS)
+    k = F.rms_norm(k, (_HEAD_DIM,), k_weight, _EPS)
+    half = _ROTARY_DIM // 2
+    cos = rope_table[..., :half].unsqueeze(1)
+    sin = rope_table[..., half:].unsqueeze(1)
+
+    def apply(x):
+        first = x[..., :half]
+        second = x[..., half:_ROTARY_DIM]
+        return torch.cat(
+            (
+                first * cos - second * sin,
+                second * cos + first * sin,
+                x[..., _ROTARY_DIM:],
+            ),
+            dim=-1,
+        )
+
+    return apply(q), apply(k)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+@pytest.mark.skipif(not HAS_TRITON, reason="Triton required")
+@pytest.mark.parametrize("seq_len", [1, 257, 1024])
+def test_fused_qk_norm_rope_matches_bf16_reference(seq_len):
+    from vllm_omni.diffusion.layers.fused_qk_norm_rope import (
+        fused_qk_norm_rope,
+    )
+
+    torch.manual_seed(17)
+    heads = 14
+    qkv = torch.randn(
+        seq_len,
+        heads * _HEAD_DIM * 3,
+        device="cuda",
+        dtype=torch.bfloat16,
+    )
+    q = qkv[:, : heads * _HEAD_DIM].view(seq_len, heads, _HEAD_DIM)
+    k = qkv[:, heads * _HEAD_DIM : 2 * heads * _HEAD_DIM].view(
+        seq_len,
+        heads,
+        _HEAD_DIM,
+    )
+    q_weight = torch.randn(_HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    k_weight = torch.randn(_HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    freqs = torch.randn(seq_len, _ROTARY_DIM // 2, device="cuda")
+    rope_table = torch.cat((torch.cos(freqs), torch.sin(freqs)), dim=-1).to(torch.bfloat16)
+
+    expected_q, expected_k = _reference(
+        q,
+        k,
+        q_weight,
+        k_weight,
+        rope_table,
+    )
+    actual_q, actual_k = fused_qk_norm_rope(
+        q,
+        k,
+        q_weight,
+        k_weight,
+        rope_table,
+        _EPS,
+    )
+
+    torch.testing.assert_close(actual_q, expected_q, atol=0.0625, rtol=0.02)
+    torch.testing.assert_close(actual_k, expected_k, atol=0.0625, rtol=0.02)
@@ -24,11 +24,20 @@ from safetensors.torch import save_file

 from tests.helpers.mark import hardware_test
 from tests.helpers.runtime import OmniServer, OmniServerParams
+from vllm_omni.platforms import current_omni_platform

 os.environ["VLLM_WORKER_MULTIPROC_METHOD"] = "spawn"

 MODEL = "Tongyi-MAI/Z-Image-Turbo"

+# ROCm's default FLASH_ATTN diffusion backend resolves to AITER. On a cold CI
+# worker, JIT-building its MHA kernel can take longer than the server readiness
+# timeout. This test validates LoRA switching rather than attention-backend or
+# compiled-execution performance, so use the eager SDPA path on ROCm only.
+ROCM_SERVER_ARGS = (
+    ["--enforce-eager", "--diffusion-attention-backend", "TORCH_SDPA"] if current_omni_platform.is_rocm() else []
+)
+

 PROMPT = "a photo of a cat sitting on a laptop keyboard"
 SIZE = "256x256"
@@ -42,7 +51,14 @@ server_params = [
            server_args=[
                "--num-gpus",
                "1",
+                *ROCM_SERVER_ARGS,
            ],
+            # A cold-cache Z-Image load can take slightly more than the
+            # orchestrator's 900s default on CI (weights alone have been
+            # observed at ~690s and full orchestrator readiness at ~906s).
+            # Match the large-model offline runner's startup allowance without
+            # weakening timeout detection for every online-serving test.
+            init_timeout=1800,
        ),
        id="zimage_turbo_lora",
    ),
@@ -23,7 +23,7 @@ def _make_pipeline(pipeline_cls, sequence_parallel_size: int = 1):
    return pipeline


-def test_prepare_video_latents_samples_directly_in_packed_token_space():
+def test_prepare_video_latents_matches_official_values_and_token_major_layout():
    pipeline = _make_pipeline(LTX2Pipeline)
    pipeline.vae_spatial_compression_ratio = 8
    pipeline.vae_temporal_compression_ratio = 8
@@ -44,6 +44,7 @@ def test_prepare_video_latents_samples_directly_in_packed_token_space():
    )

    torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+    assert actual.stride()[1:] == (1, actual.shape[1])


 def test_prepare_audio_latents_samples_directly_in_packed_token_space():
@@ -81,6 +82,42 @@ def test_prepare_audio_latents_pads_generated_dummy_length_for_sp():
    assert original_num_frames == 1
    assert padded_num_frames == 2
    assert latents.shape == (1, 2, 128)
+    torch.testing.assert_close(latents[:, 1:], torch.zeros_like(latents[:, 1:]))
+
+
+def test_prepare_audio_latents_request_rng_is_invariant_to_sp_padding():
+    pipeline_sp1 = _make_pipeline(LTX2Pipeline, sequence_parallel_size=1)
+    pipeline_sp4 = _make_pipeline(LTX2Pipeline, sequence_parallel_size=4)
+    generator_sp1 = torch.Generator().manual_seed(42)
+    generator_sp4 = torch.Generator().manual_seed(42)
+
+    latents_sp1, _, _ = pipeline_sp1.prepare_audio_latents(
+        batch_size=1,
+        num_channels_latents=2,
+        num_mel_bins=8,
+        audio_latent_length=3,
+        dtype=torch.float32,
+        device=torch.device("cpu"),
+        generator=generator_sp1,
+    )
+    latents_sp4, _, _ = pipeline_sp4.prepare_audio_latents(
+        batch_size=1,
+        num_channels_latents=2,
+        num_mel_bins=8,
+        audio_latent_length=3,
+        dtype=torch.float32,
+        device=torch.device("cpu"),
+        generator=generator_sp4,
+    )
+
+    torch.testing.assert_close(latents_sp4[:, :3], latents_sp1, rtol=0, atol=0)
+    torch.testing.assert_close(latents_sp4[:, 3:], torch.zeros_like(latents_sp4[:, 3:]))
+    torch.testing.assert_close(
+        torch.randn(8, generator=generator_sp4),
+        torch.randn(8, generator=generator_sp1),
+        rtol=0,
+        atol=0,
+    )


 def test_prepare_audio_latents_pads_packed_sequence_dim_for_provided_latents():
@@ -137,7 +174,9 @@ def test_prepare_audio_latents_accepts_already_padded_4d_latents_for_sp():
    assert original_num_frames == 10
    assert padded_num_frames == 12
    assert padded.shape == (1, 12, 8)
-    torch.testing.assert_close(padded, latent_ops.pack_audio_latents(latents))
+    packed = latent_ops.pack_audio_latents(latents)
+    torch.testing.assert_close(padded[:, :10], packed[:, :10])
+    torch.testing.assert_close(padded[:, 10:], torch.zeros_like(padded[:, 10:]))


 def test_resolve_audio_latent_length_preserves_legacy_4d_shape_inference():
@@ -12,6 +12,19 @@ import torch
 pytestmark = [pytest.mark.core_model, pytest.mark.cpu]


+def test_velocity_from_x0_uses_official_host_scalar_sigma(mocker):
+    from vllm_omni.diffusion.models.ltx2.ltx2_guidance import velocity_from_x0
+
+    sigma = mocker.MagicMock()
+    sigma.to.return_value.item.return_value = 0.5
+
+    actual = velocity_from_x0(torch.tensor([1.0]), torch.tensor([0.0]), sigma)
+
+    sigma.to.assert_called_once_with(torch.float32)
+    sigma.to.return_value.item.assert_called_once_with()
+    torch.testing.assert_close(actual, torch.tensor([2.0]))
+
+
 class TestCFGParallelHelpers:
    """Test LTX-2.3 CFG helper math without loading model weights."""

@@ -199,6 +212,7 @@ class TestCFGParallelHelpers:
            ),
            attention_kwargs=None,
            audio_scheduler=SimpleNamespace(sigmas=torch.stack([audio_sigma])),
+            original_audio_num_frames=state.audio.shape[1],
        )

        actual_video, actual_audio = executor.predict_parallel_guidance(
@@ -0,0 +1,483 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""Unit tests for the fixed LTX refinement-phase adapter."""
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+from safetensors.torch import save_file
+
+from vllm_omni.diffusion.models.ltx2 import ltx2_phase_adapter
+from vllm_omni.diffusion.models.ltx2.ltx2_adapter_parser import (
+    AdapterTarget,
+    parse_ltx_adapter,
+)
+from vllm_omni.diffusion.models.ltx2.ltx2_components import LTX23_TWO_STAGE_COMPONENT_PROFILE
+from vllm_omni.diffusion.models.ltx2.ltx2_phase_adapter import LTXPhaseAdapterRuntime
+from vllm_omni.diffusion.models.ltx2.ltx2_recipes import (
+    LTX2_TWO_STAGE_RECIPE,
+    LTX_DISTILLED_ADAPTER_SLOT,
+)
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+
+
+def _target(
+    lora_a,
+    lora_b,
+    *,
+    source_module="proj",
+    module="proj",
+    target_slice=None,
+):
+    return AdapterTarget(
+        source_module=source_module,
+        module=module,
+        slice=target_slice,
+        a_key="a",
+        b_key="b",
+        a_shape=tuple(lora_a.shape),
+        b_shape=tuple(lora_b.shape),
+    )
+
+
+def _wrap(
+    layer,
+    lora_a,
+    lora_b,
+    *,
+    source_module="proj",
+    module="proj",
+    target_slice=None,
+    layer_fused=False,
+):
+    target = _target(
+        lora_a,
+        lora_b,
+        source_module=source_module,
+        module=module,
+        target_slice=target_slice,
+    )
+    wrapper = ltx2_phase_adapter._PhaseAdapterLinear(
+        layer,
+        [target],
+        adapter_dtype=lora_a.dtype,
+        layer_fused=layer_fused,
+    )
+    wrapper.load_target(target, lora_a, lora_b)
+    wrapper.set_enabled(True)
+    return wrapper
+
+
+class _LoRAAttention(torch.nn.Module):
+    def __init__(self) -> None:
+        super().__init__()
+        self.to_qkv = torch.nn.Identity()
+
+
+class _LoRATransformer(torch.nn.Module):
+    stacked_params_mapping = ((".attn1.to_qkv", ".attn1.to_q", "q"),)
+
+    def __init__(self) -> None:
+        super().__init__()
+        self.proj_in = torch.nn.Linear(4, 6, bias=False)
+        block = torch.nn.Module()
+        block.attn1 = _LoRAAttention()
+        self.transformer_blocks = torch.nn.ModuleList([block])
+
+
+def test_ltx_adapter_parser_maps_official_and_packed_module_names(tmp_path):
+    adapter_path = tmp_path / "distilled.safetensors"
+    save_file(
+        {
+            "diffusion_model.text_embedding_projection.aggregate_embed.lora_A.weight": torch.randn(2, 4),
+            "diffusion_model.patchify_proj.lora_A.weight": torch.randn(2, 4),
+            "diffusion_model.patchify_proj.lora_B.weight": torch.randn(6, 2),
+            "diffusion_model.transformer_blocks.0.attn1.to_q.lora_A.weight": torch.randn(2, 4),
+            "diffusion_model.transformer_blocks.0.attn1.to_q.lora_B.weight": torch.randn(4, 2),
+        },
+        adapter_path,
+    )
+
+    manifest = parse_ltx_adapter(_LoRATransformer(), str(adapter_path))
+    assert {(target.source_module, target.module, target.slice) for target in manifest.targets} == {
+        ("proj_in", "proj_in", None),
+        ("transformer_blocks.0.attn1.to_q", "transformer_blocks.0.attn1.to_qkv", "q"),
+    }
+
+
+def test_ltx_adapter_parser_rejects_unpaired_b_tensor(tmp_path):
+    adapter_path = tmp_path / "distilled.safetensors"
+    save_file(
+        {"diffusion_model.patchify_proj.lora_B.weight": torch.randn(6, 2)},
+        adapter_path,
+    )
+
+    with pytest.raises(ValueError, match="Missing paired LoRA A tensors"):
+        parse_ltx_adapter(_LoRATransformer(), str(adapter_path))
+
+
+def _pipeline(*, dtype, quantization_config):
+    return SimpleNamespace(
+        pipeline_recipe=LTX2_TWO_STAGE_RECIPE,
+        component_profile=LTX23_TWO_STAGE_COMPONENT_PROFILE,
+        transformer=object(),
+        od_config=SimpleNamespace(
+            model="unused",
+            lora_path=None,
+            dtype=dtype,
+            quantization_config=quantization_config,
+        ),
+    )
+
+
+@pytest.mark.parametrize(
+    ("dtype", "quantization_config", "layer_fused"),
+    [
+        (torch.float32, object(), False),
+        (torch.bfloat16, None, True),
+    ],
+)
+def test_ltx_phase_adapter_selects_execution_from_quantization(
+    monkeypatch,
+    dtype,
+    quantization_config,
+    layer_fused,
+):
+    pipeline = _pipeline(dtype=dtype, quantization_config=quantization_config)
+    manifest = object()
+    captured = {}
+
+    def resolve_artifact(*_args, **kwargs):
+        assert not kwargs
+        return "adapter"
+
+    def parse(transformer, path):
+        assert transformer is pipeline.transformer
+        assert path == "adapter"
+        return manifest
+
+    class Runtime:
+        def __init__(self, transformer, parsed_manifest, **kwargs):
+            captured.update(transformer=transformer, manifest=parsed_manifest, **kwargs)
+
+    monkeypatch.setattr(ltx2_phase_adapter, "resolve_ltx_artifact", resolve_artifact)
+    monkeypatch.setattr(ltx2_phase_adapter, "parse_ltx_adapter", parse)
+    monkeypatch.setattr(ltx2_phase_adapter, "LTXPhaseAdapterRuntime", Runtime)
+
+    runtime = ltx2_phase_adapter.build_ltx_phase_adapter(pipeline)
+
+    assert isinstance(runtime, Runtime)
+    assert captured == {
+        "transformer": pipeline.transformer,
+        "manifest": manifest,
+        "dtype": dtype,
+        "layer_fused": layer_fused,
+    }
+
+
+def test_ltx_phase_adapter_rejects_non_bfloat16_layer_fused(monkeypatch):
+    pipeline = _pipeline(dtype=torch.float16, quantization_config=None)
+    monkeypatch.setattr(ltx2_phase_adapter, "resolve_ltx_artifact", lambda *_args, **_kwargs: "adapter")
+
+    with pytest.raises(ValueError, match="bfloat16"):
+        ltx2_phase_adapter.build_ltx_phase_adapter(pipeline)
+
+
+def test_ltx_phase_adapter_keeps_one_transformer_and_switches_a_fixed_slot(tmp_path):
+    class TinyTransformer(torch.nn.Module):
+        stacked_params_mapping = ()
+
+        def __init__(self) -> None:
+            super().__init__()
+            self.proj = torch.nn.Linear(2, 2, bias=False)
+
+    adapter_path = tmp_path / "distilled.safetensors"
+    save_file(
+        {
+            "diffusion_model.proj.lora_A.weight": torch.tensor([[1.0, 2.0]]),
+            "diffusion_model.proj.lora_B.weight": torch.tensor([[3.0], [4.0]]),
+        },
+        adapter_path,
+    )
+    transformer = TinyTransformer()
+    runtime = LTXPhaseAdapterRuntime(
+        transformer,
+        parse_ltx_adapter(transformer, str(adapter_path)),
+        dtype=torch.float32,
+    )
+
+    assert runtime.base_parameter_name("proj.weight") == "proj.base_layer.weight"
+    assert {name for name, _ in transformer.named_parameters()} == {"proj.base_layer.weight"}
+    with torch.no_grad():
+        transformer.proj.base_layer.weight.copy_(torch.eye(2))
+    runtime.finalize()
+
+    x = torch.tensor([[1.0, 1.0]])
+    runtime.activate(None)
+    torch.testing.assert_close(transformer.proj(x), torch.tensor([[1.0, 1.0]]))
+    runtime.activate(LTX_DISTILLED_ADAPTER_SLOT)
+    torch.testing.assert_close(transformer.proj(x), torch.tensor([[10.0, 13.0]]))
+
+
+def test_ltx_layer_fused_adapter_matches_official_bf16_weight_fusion():
+    layer = torch.nn.Linear(4, 3, bias=True, dtype=torch.bfloat16)
+    base_weight = torch.tensor(
+        [
+            [0.5, -0.25, 0.125, 0.75],
+            [-0.375, 0.625, 0.25, -0.5],
+            [0.875, -0.125, -0.75, 0.375],
+        ],
+        dtype=torch.bfloat16,
+    )
+    lora_a = torch.tensor(
+        [
+            [0.1015625, -0.203125, 0.3046875, -0.40625],
+            [0.5078125, -0.609375, 0.7109375, -0.8125],
+        ],
+        dtype=torch.bfloat16,
+    )
+    lora_b = torch.tensor(
+        [[0.9140625, -0.1171875], [-0.21875, 0.3203125], [0.421875, -0.5234375]],
+        dtype=torch.bfloat16,
+    )
+    with torch.no_grad():
+        layer.weight.copy_(base_weight)
+        layer.bias.copy_(torch.tensor([0.125, -0.25, 0.5], dtype=torch.bfloat16))
+    wrapper = _wrap(layer, lora_a, lora_b, layer_fused=True)
+
+    expected_weight = torch.matmul(lora_b, lora_a)
+    expected_weight.add_(base_weight)
+    x = torch.tensor([[0.6015625, -0.703125, 0.8046875, -0.90625]], dtype=torch.bfloat16)
+    expected = torch.nn.functional.linear(x, expected_weight, layer.bias)
+
+    assert torch.equal(wrapper(x), expected)
+    assert torch.equal(layer.weight, base_weight)
+
+
+def test_ltx_layer_fused_row_parallel_uses_one_collective(monkeypatch):
+    from vllm.model_executor.layers.linear import UnquantizedLinearMethod
+
+    class FakeRow(torch.nn.Module):
+        input_size = 4
+        output_size = 2
+        input_size_per_partition = 2
+        input_is_parallel = True
+        tp_rank = 0
+        tp_size = 2
+        reduce_results = True
+        skip_bias_add = False
+        return_bias = False
+
+        def __init__(self):
+            super().__init__()
+            self.weight = torch.nn.Parameter(torch.eye(2, dtype=torch.bfloat16))
+            self.bias = torch.nn.Parameter(torch.tensor([0.25, -0.5], dtype=torch.bfloat16))
+            self.quant_method = UnquantizedLinearMethod()
+
+    monkeypatch.setattr(ltx2_phase_adapter, "RowParallelLinear", FakeRow)
+    reductions = []
+
+    def all_reduce(value):
+        reductions.append(value.clone())
+        return value + 1
+
+    monkeypatch.setattr(ltx2_phase_adapter, "tensor_model_parallel_all_reduce", all_reduce)
+    layer = FakeRow()
+    lora_a = torch.tensor([[0.5, -0.25, 0.75, -0.125]], dtype=torch.bfloat16)
+    lora_b = torch.tensor([[0.25], [-0.5]], dtype=torch.bfloat16)
+    wrapper = _wrap(layer, lora_a, lora_b, layer_fused=True)
+
+    expected_weight = torch.matmul(lora_b, lora_a[:, :2])
+    expected_weight.add_(layer.weight)
+    x = torch.tensor([[1.0, 2.0]], dtype=torch.bfloat16)
+    expected_local = torch.nn.functional.linear(x, expected_weight, layer.bias)
+
+    actual = wrapper(x)
+    assert len(reductions) == 1
+    assert torch.equal(reductions[0], expected_local)
+    assert torch.equal(actual, expected_local + 1)
+
+
+def test_ltx_layer_fused_column_parallel_uses_local_weight_shard(monkeypatch):
+    from vllm.model_executor.layers.linear import UnquantizedLinearMethod
+
+    class FakeColumn(torch.nn.Module):
+        input_size = 3
+        output_size = 4
+        output_size_per_partition = 2
+        tp_rank = 1
+        tp_size = 2
+        gather_output = False
+        skip_bias_add = False
+        return_bias = False
+
+        def __init__(self):
+            super().__init__()
+            self.weight = torch.nn.Parameter(
+                torch.tensor([[0.5, -0.25, 0.125], [-0.375, 0.625, 0.25]], dtype=torch.bfloat16)
+            )
+            self.bias = torch.nn.Parameter(torch.tensor([0.25, -0.5], dtype=torch.bfloat16))
+            self.quant_method = UnquantizedLinearMethod()
+
+    monkeypatch.setattr(ltx2_phase_adapter, "ColumnParallelLinear", FakeColumn)
+    layer = FakeColumn()
+    base_weight = layer.weight.detach().clone()
+    lora_a = torch.tensor([[0.5, -0.25, 0.75]], dtype=torch.bfloat16)
+    lora_b = torch.tensor([[0.25], [-0.5], [0.75], [-1.0]], dtype=torch.bfloat16)
+    wrapper = _wrap(layer, lora_a, lora_b, layer_fused=True)
+
+    expected_weight = torch.matmul(lora_b[2:], lora_a)
+    expected_weight.add_(base_weight)
+    x = torch.tensor([[1.0, -0.5, 0.25]], dtype=torch.bfloat16)
+    expected = torch.nn.functional.linear(x, expected_weight, layer.bias)
+
+    assert torch.equal(wrapper(x), expected)
+    assert torch.equal(layer.weight, base_weight)
+
+
+def test_ltx_layer_fused_packed_qkv_updates_only_the_target_slice(monkeypatch):
+    from vllm.model_executor.layers.linear import UnquantizedLinearMethod
+
+    class FakeColumn(torch.nn.Module):
+        pass
+
+    class FakeQKV(FakeColumn):
+        input_size = 4
+        head_size = 2
+        v_head_size = 2
+        total_num_heads = 2
+        total_num_kv_heads = 2
+        num_heads = 1
+        num_kv_heads = 1
+        num_kv_head_replicas = 1
+        tp_rank = 1
+        tp_size = 2
+        gather_output = False
+        skip_bias_add = True
+        return_bias = True
+
+        def __init__(self):
+            super().__init__()
+            self.weight = torch.nn.Parameter(
+                torch.arange(24, dtype=torch.float32).reshape(6, 4).to(torch.bfloat16) / 16
+            )
+            self.bias = torch.nn.Parameter(torch.arange(6, dtype=torch.float32).to(torch.bfloat16) / 8)
+            self.quant_method = UnquantizedLinearMethod()
+
+    monkeypatch.setattr(ltx2_phase_adapter, "ColumnParallelLinear", FakeColumn)
+    monkeypatch.setattr(ltx2_phase_adapter, "QKVParallelLinear", FakeQKV)
+    layer = FakeQKV()
+    base_weight = layer.weight.detach().clone()
+    lora_a = torch.tensor([[0.5, -0.25, 0.75, -0.125]], dtype=torch.bfloat16)
+    lora_b = torch.tensor([[0.25], [-0.5], [0.75], [-1.0]], dtype=torch.bfloat16)
+    wrapper = _wrap(
+        layer,
+        lora_a,
+        lora_b,
+        source_module="to_k",
+        module="to_qkv",
+        target_slice="k",
+        layer_fused=True,
+    )
+
+    expected_weight = base_weight.clone()
+    expected_k = torch.matmul(lora_b[2:], lora_a)
+    expected_k.add_(base_weight[2:4])
+    expected_weight[2:4].copy_(expected_k)
+    x = torch.tensor([[1.0, -0.5, 0.25, -0.125]], dtype=torch.bfloat16)
+    expected_output = torch.nn.functional.linear(x, expected_weight)
+
+    actual_output, returned_bias = wrapper(x)
+    assert torch.equal(actual_output, expected_output)
+    assert returned_bias is layer.bias
+    assert torch.equal(layer.weight, base_weight)
+
+
+@pytest.mark.parametrize(
+    ("layer_factory", "target_slice", "expected_a", "expected_b", "expected_output"),
+    [
+        (
+            lambda: type(
+                "FakeColumn",
+                (torch.nn.Module,),
+                {
+                    "input_size": 4,
+                    "output_size": 4,
+                    "output_size_per_partition": 2,
+                    "tp_rank": 1,
+                    "gather_output": False,
+                    "forward": lambda self, x: x.new_zeros((*x.shape[:-1], 2)),
+                },
+            )(),
+            None,
+            torch.tensor([[1.0, 10.0, 100.0, 1000.0]]),
+            torch.tensor([[1.0], [2.0], [3.0], [4.0]]),
+            torch.tensor([[12963.0, 17284.0]]),
+        ),
+        (
+            lambda: type(
+                "FakeRow",
+                (torch.nn.Module,),
+                {
+                    "input_size": 4,
+                    "output_size": 2,
+                    "input_size_per_partition": 2,
+                    "tp_rank": 1,
+                    "reduce_results": False,
+                    "forward": lambda self, x: x.new_zeros((*x.shape[:-1], 2)),
+                },
+            )(),
+            None,
+            torch.tensor([[1.0, 10.0, 100.0, 1000.0]]),
+            torch.tensor([[2.0], [3.0]]),
+            torch.tensor([[8600.0, 12900.0]]),
+        ),
+        (
+            lambda: type(
+                "FakeQKV",
+                (torch.nn.Module,),
+                {
+                    "input_size": 4,
+                    "head_size": 2,
+                    "v_head_size": 2,
+                    "total_num_heads": 2,
+                    "total_num_kv_heads": 2,
+                    "num_heads": 1,
+                    "num_kv_heads": 1,
+                    "num_kv_head_replicas": 1,
+                    "tp_rank": 1,
+                    "forward": lambda self, x: x.new_zeros((*x.shape[:-1], 6)),
+                },
+            )(),
+            "k",
+            torch.tensor([[1.0, 10.0, 100.0, 1000.0]]),
+            torch.tensor([[1.0], [2.0], [3.0], [4.0]]),
+            torch.tensor([[0.0, 0.0, 12963.0, 17284.0, 0.0, 0.0]]),
+        ),
+    ],
+)
+def test_ltx_phase_adapter_uses_local_tp_lora_slices(
+    monkeypatch,
+    layer_factory,
+    target_slice,
+    expected_a,
+    expected_b,
+    expected_output,
+):
+    layer = layer_factory()
+    layer.register_buffer("device_anchor", torch.empty(0))
+    layer_type = type(layer)
+    if target_slice == "k":
+        monkeypatch.setattr(ltx2_phase_adapter, "QKVParallelLinear", layer_type)
+    elif "Column" in layer_type.__name__:
+        monkeypatch.setattr(ltx2_phase_adapter, "ColumnParallelLinear", layer_type)
+    else:
+        monkeypatch.setattr(ltx2_phase_adapter, "RowParallelLinear", layer_type)
+
+    wrapper = _wrap(layer, expected_a, expected_b, target_slice=target_slice)
+    torch.testing.assert_close(wrapper(torch.tensor([[1.0, 2.0, 3.0, 4.0]])), expected_output)
+    expected_local_a_width = getattr(layer, "input_size_per_partition", layer.input_size)
+    assert wrapper.pieces[0].lora_a.shape[-1] == expected_local_a_width
@@ -13,13 +13,21 @@ import pytest
 import torch

 from vllm_omni.diffusion.data import DiffusionOutput
-from vllm_omni.diffusion.models.ltx2 import ltx2_components
+from vllm_omni.diffusion.models.ltx2 import (
+    ltx2_components,
+    ltx2_transformer,
+)
 from vllm_omni.diffusion.models.ltx2.ltx2_components import (
    LTX2_COMPONENT_PROFILE,
    LTX2_DISTILLED_COMPONENT_PROFILE,
+    LTX2_TWO_STAGE_COMPONENT_PROFILE,
    LTX23_COMPONENT_PROFILE,
+    LTX23_DISTILLED_COMPONENT_PROFILE,
+    LTX23_TWO_STAGE_COMPONENT_PROFILE,
    _install_connector_attention,
    detect_ltx_model_version,
+    resolve_ltx_artifact,
+    resolve_ltx_component_profile,
 )
 from vllm_omni.diffusion.models.ltx2.ltx2_conditioning import LTXI2VConditioningMixin
 from vllm_omni.diffusion.models.ltx2.ltx2_denoise import (
@@ -33,6 +41,7 @@ from vllm_omni.diffusion.models.ltx2.ltx2_guidance import (
    LTXGuidancePlan,
    LTXGuidanceSpec,
    LTXModalityGuidance,
+    _repeat_batch,
    build_perturbation_kwargs,
    combine_guided_x0,
 )
@@ -40,18 +49,29 @@ from vllm_omni.diffusion.models.ltx2.ltx2_latents import LTXAVState
 from vllm_omni.diffusion.models.ltx2.ltx2_recipes import (
    LTX2_DISTILLED_TWO_STAGE_RECIPE,
    LTX2_ONE_STAGE_RECIPE,
+    LTX2_TWO_STAGE_RECIPE,
+    LTX23_DISTILLED_TWO_STAGE_RECIPE,
    LTX23_ONE_STAGE_RECIPE,
+    LTX23_TWO_STAGE_RECIPE,
    LTX_DEFAULT_NEGATIVE_PROMPT,
    LTX_POSITIVE_ONLY_RECIPE,
+    resolve_ltx_pipeline_recipe,
+)
+from vllm_omni.diffusion.models.ltx2.ltx2_request import (
+    LTXRequestInputs,
+    validate_ltx_checkpoint,
+    validate_pipeline_request,
 )
-from vllm_omni.diffusion.models.ltx2.ltx2_request import LTXRequestInputs, validate_pipeline_request
 from vllm_omni.diffusion.models.ltx2.ltx2_runtime import LTXRuntime
 from vllm_omni.diffusion.models.ltx2.pipeline_ltx2 import (
    LTX2I2VDMD2Pipeline,
    LTX2Pipeline,
    LTX2T2VDMD2Pipeline,
 )
-from vllm_omni.diffusion.models.ltx2.pipeline_ltx2_two_stage import LTX2DistilledPipeline
+from vllm_omni.diffusion.models.ltx2.pipeline_ltx2_two_stage import (
+    LTX2DistilledPipeline,
+    LTX2TwoStagePipeline,
+)

 pytestmark = [pytest.mark.core_model, pytest.mark.cpu]

@@ -104,10 +124,22 @@ def test_ltx_public_entries_share_runtime_and_keep_recipe_boundaries():
    )

    assert issubclass(LTX2Pipeline, LTXRuntime)
+    assert issubclass(LTX2TwoStagePipeline, LTXRuntime)
    assert issubclass(LTX2DistilledPipeline, LTXRuntime)
    assert LTX2Pipeline.component_profile is LTX2_COMPONENT_PROFILE
    assert LTX2Pipeline.pipeline_recipe is LTX2_ONE_STAGE_RECIPE
+    assert LTX2TwoStagePipeline.component_profile is LTX2_TWO_STAGE_COMPONENT_PROFILE
+    assert LTX2TwoStagePipeline.pipeline_recipe is LTX2_TWO_STAGE_RECIPE
    assert LTX2DistilledPipeline.component_profile is LTX2_DISTILLED_COMPONENT_PROFILE
+    for pipeline_cls, profile in (
+        (LTX2Pipeline, LTX2_COMPONENT_PROFILE),
+        (LTX2TwoStagePipeline, LTX2_TWO_STAGE_COMPONENT_PROFILE),
+        (LTX2DistilledPipeline, LTX2_DISTILLED_COMPONENT_PROFILE),
+    ):
+        assert pipeline_cls._dit_modules == list(profile.dit_modules)
+        assert pipeline_cls._encoder_modules == list(profile.encoder_modules)
+        assert pipeline_cls._vae_modules == list(profile.vae_modules)
+        assert pipeline_cls._resident_modules == list(profile.resident_modules)
    assert LTX2_COMPONENT_PROFILE.video_vae_cls is DistributedAutoencoderKLLTX2Video
    assert LTX23_COMPONENT_PROFILE.video_vae_cls is DistributedAutoencoderKLLTX2Video

@@ -123,6 +155,72 @@ def test_ltx_checkpoint_version_detection_uses_metadata(tmp_path):
    assert detect_ltx_model_version(str(tmp_path)) == "2"


+def test_ltx_artifact_falls_back_to_unpinned_hub_download(tmp_path, monkeypatch):
+    filename = "ltx-sidecar.safetensors"
+    calls = []
+
+    def fake_download(**kwargs):
+        calls.append(kwargs)
+        return "/cache/ltx-sidecar.safetensors"
+
+    monkeypatch.setattr(ltx2_components, "hf_hub_download", fake_download)
+
+    assert (
+        resolve_ltx_artifact(str(tmp_path / "model"), "Lightricks/test", filename) == "/cache/ltx-sidecar.safetensors"
+    )
+    assert calls == [{"repo_id": "Lightricks/test", "filename": filename}]
+
+
+def test_ltx_artifact_prefers_model_root(tmp_path, monkeypatch):
+    filename = "ltx-sidecar.safetensors"
+    expected = tmp_path / filename
+    expected.write_bytes(b"sidecar")
+    monkeypatch.setattr(ltx2_components, "hf_hub_download", lambda **_kwargs: pytest.fail("unexpected Hub lookup"))
+
+    assert resolve_ltx_artifact(str(tmp_path), "Lightricks/test", filename) == str(expected)
+
+
+@pytest.mark.parametrize(
+    ("expected_kind", "scheduler_config"),
+    [
+        ("regular", {"use_dynamic_shifting": True, "shift_terminal": 0.1}),
+        ("distilled", {"use_dynamic_shifting": False, "shift_terminal": None}),
+    ],
+)
+def test_ltx_checkpoint_validation_accepts_matching_scheduler_metadata(expected_kind, scheduler_config):
+    validate_ltx_checkpoint(
+        scheduler_config,
+        expected_kind=expected_kind,
+        pipeline_name="TestLTXPipeline",
+    )
+
+
+@pytest.mark.parametrize(
+    ("expected_kind", "scheduler_config", "error"),
+    [
+        ("regular", {"use_dynamic_shifting": False, "shift_terminal": None}, "regular non-distilled"),
+        ("distilled", {"use_dynamic_shifting": True, "shift_terminal": 0.1}, "merged distilled"),
+    ],
+)
+def test_ltx_checkpoint_validation_rejects_mismatched_scheduler_metadata(expected_kind, scheduler_config, error):
+    with pytest.raises(ValueError, match=error):
+        validate_ltx_checkpoint(
+            scheduler_config,
+            expected_kind=expected_kind,
+            pipeline_name="TestLTXPipeline",
+        )
+
+
+def test_ltx_rejects_advanced_uaa_before_component_initialization():
+    od_config = SimpleNamespace(
+        model="unused",
+        parallel_config=SimpleNamespace(ulysses_mode="advanced_uaa"),
+    )
+
+    with pytest.raises(ValueError, match="does not support ulysses_mode='advanced_uaa'"):
+        LTX2Pipeline(od_config=od_config)
+
+
 def test_ltx23_checkpoint_selects_version_specific_one_stage_profile(tmp_path, monkeypatch):
    from vllm_omni.diffusion.models.ltx2 import ltx2_runtime

@@ -144,16 +242,104 @@ def test_ltx23_checkpoint_selects_version_specific_one_stage_profile(tmp_path, m
    assert pipe.reports_stage_durations


-def test_ltx_one_stage_entry_exposes_batch_image_and_distributed_decode_contracts():
+def test_ltx_multistage_rejects_cache_dit_before_component_initialization(tmp_path, monkeypatch):
+    from vllm_omni.diffusion.models.ltx2 import ltx2_runtime
+
+    (tmp_path / "model_index.json").write_text(json.dumps({"vocoder": ["ltx2", "LTX2Vocoder"]}))
+    initialized = False
+
+    def stub_components(*_args, **_kwargs):
+        nonlocal initialized
+        initialized = True
+
+    monkeypatch.setattr(ltx2_runtime, "initialize_pipeline_components", stub_components)
+
+    with pytest.raises(ValueError, match="only one-stage LTX recipes"):
+        LTX2TwoStagePipeline(
+            od_config=SimpleNamespace(
+                model=str(tmp_path),
+                cache_backend="cache_dit",
+            )
+        )
+
+    assert not initialized
+
+
+def test_ltx23_checkpoint_selects_version_specific_two_stage_profiles(tmp_path, monkeypatch):
+    from vllm_omni.diffusion.models.ltx2 import ltx2_runtime
+
+    (tmp_path / "model_index.json").write_text(json.dumps({"vocoder": ["ltx2", "LTX2VocoderWithBWE"]}))
+
+    def stub_components(pipe, od_config):
+        pipe.od_config = od_config
+        pipe.vae_spatial_compression_ratio = 32
+
+    monkeypatch.setattr(ltx2_runtime, "initialize_pipeline_components", stub_components)
+    monkeypatch.setattr(LTXRuntime, "setup_diffusion_pipeline_profiler", lambda *_args, **_kwargs: None)
+    monkeypatch.setattr(ltx2_runtime, "build_ltx_phase_adapter", lambda *_args: SimpleNamespace())
+    pipe = LTX2TwoStagePipeline(
+        od_config=SimpleNamespace(
+            model=str(tmp_path),
+            enable_diffusion_pipeline_profiler=False,
+            lora_path=None,
+            model_config={},
+        )
+    )
+
+    assert pipe.model_version == "2.3"
+    assert pipe.component_profile is LTX23_TWO_STAGE_COMPONENT_PROFILE
+    assert pipe.pipeline_recipe is LTX23_TWO_STAGE_RECIPE
+
+
+def test_ltx23_checkpoint_selects_full_distilled_profile_and_recipe(tmp_path, monkeypatch):
+    from vllm_omni.diffusion.models.ltx2 import ltx2_runtime
+
+    (tmp_path / "model_index.json").write_text(json.dumps({"vocoder": ["ltx2", "LTX2VocoderWithBWE"]}))
+
+    def stub_components(pipe, od_config):
+        pipe.od_config = od_config
+        pipe.vae_spatial_compression_ratio = 32
+
+    monkeypatch.setattr(ltx2_runtime, "initialize_pipeline_components", stub_components)
+    monkeypatch.setattr(LTXRuntime, "setup_diffusion_pipeline_profiler", lambda *_args, **_kwargs: None)
+
+    pipe = LTX2DistilledPipeline(
+        od_config=SimpleNamespace(
+            model=str(tmp_path),
+            enable_diffusion_pipeline_profiler=False,
+        )
+    )
+
+    assert pipe.model_version == "2.3"
+    assert pipe.component_profile is LTX23_DISTILLED_COMPONENT_PROFILE
+    assert pipe.pipeline_recipe is LTX23_DISTILLED_TWO_STAGE_RECIPE
+    assert pipe._phase_adapter is None
+    assert pipe.preserve_sp_padded_audio_duration
+    assert pipe.reports_stage_durations
+
+
+def test_ltx_entries_expose_batch_image_and_distributed_decode_contracts():
    assert LTX2Pipeline.supports_request_batch
    assert LTX2Pipeline.support_image_input
    assert LTX2Pipeline.unified_text_image_entry
    assert LTX2Pipeline.distributed_video_decode
+    assert not LTX2TwoStagePipeline.supports_request_batch
+    assert LTX2TwoStagePipeline.support_image_input
+    assert LTX2TwoStagePipeline.unified_text_image_entry
    assert not LTX2DistilledPipeline.supports_request_batch
    assert LTX2DistilledPipeline.support_image_input
    assert LTX2DistilledPipeline.unified_text_image_entry


+def test_ltx_two_stage_resolves_request_embedded_images():
+    pipe = object.__new__(LTX2TwoStagePipeline)
+    request_inputs = SimpleNamespace(latents=None)
+    image = torch.zeros(3, 8, 8)
+    req = SimpleNamespace(prompts=[{"prompt": "image conditioned", "multi_modal_data": {"image": image}}])
+
+    assert pipe._resolve_request_image(req, None, request_inputs) is image
+
+
 def test_ltx_dmd2_entries_retain_their_task_boundaries():
    assert not LTX2T2VDMD2Pipeline.support_image_input
    assert LTX2I2VDMD2Pipeline.support_image_input
@@ -330,6 +516,15 @@ def test_ltx_guidance_uses_official_close_to_disabled_semantics():
    assert LTXGuidancePlan.build(LTXGuidanceSpec(video=guidance)).names == ("cond",)


+def test_ltx_positive_only_guidance_preserves_token_major_layout():
+    tensor = torch.arange(24).view(1, 3, 8).transpose(1, 2).contiguous().transpose(1, 2)
+
+    repeated = _repeat_batch(tensor, 1)
+
+    assert repeated is tensor
+    assert repeated.stride() == tensor.stride()
+
+
 def test_ltx_guidance_rescale_is_invariant_to_request_batching():
    cond = torch.tensor(
        [
@@ -380,6 +575,32 @@ def test_ltx_guidance_rescale_handles_zero_variance_prediction():
    torch.testing.assert_close(actual, torch.zeros_like(cond))


+def test_ltx_audio_guidance_rescale_ignores_sp_padding_tokens():
+    logical_cond = torch.tensor([[[1.0, 2.0], [3.0, 5.0]]])
+    logical_uncond = torch.tensor([[[0.5, -1.0], [2.0, 1.0]]])
+    guidance = LTXModalityGuidance(cfg_scale=3.0, rescale_scale=1.0)
+    expected = combine_guided_x0(
+        cond=logical_cond,
+        uncond_text=logical_uncond,
+        uncond_perturbed=0.0,
+        uncond_modality=0.0,
+        guidance=guidance,
+    )
+    padded_cond = torch.cat([logical_cond, torch.tensor([[[1000.0, -1000.0]]])], dim=1)
+    padded_uncond = torch.cat([logical_uncond, torch.tensor([[[-500.0, 500.0]]])], dim=1)
+
+    actual = combine_guided_x0(
+        cond=padded_cond,
+        uncond_text=padded_uncond,
+        uncond_perturbed=0.0,
+        uncond_modality=0.0,
+        guidance=guidance,
+        rescale_token_count=2,
+    )
+
+    torch.testing.assert_close(actual[:, :2], expected)
+
+
 def test_ltx_guidance_builds_per_sample_stg_and_modality_masks():
    plan = LTXGuidancePlan.build(LTX23_ONE_STAGE_RECIPE.request_guidance)
    perturbations = build_perturbation_kwargs(plan, batch_size=2, reference=torch.ones(8, 1, 1))
@@ -547,11 +768,19 @@ def test_ltx_connector_attention_dispatches_through_omni_kernel(has_learned_regi

 def test_ltx2_distilled_pipeline_shares_recipe_runtime():
    assert issubclass(LTX2DistilledPipeline, LTXI2VConditioningMixin)
-    assert issubclass(LTX2DistilledPipeline, LTX2Pipeline)
+    assert not issubclass(LTX2DistilledPipeline, LTX2Pipeline)
    assert LTX2DistilledPipeline._run_recipe is LTXRuntime._run_recipe
    assert "_forward_request" not in LTX2DistilledPipeline.__dict__


+def test_ltx_two_stage_entry_is_a_thin_pipeline_variant():
+    assert "_run_recipe" not in LTX2TwoStagePipeline.__dict__
+    assert "_forward_request" not in LTX2TwoStagePipeline.__dict__
+    assert "load_weights" not in LTX2TwoStagePipeline.__dict__
+    assert issubclass(LTX2TwoStagePipeline, LTXI2VConditioningMixin)
+    assert not issubclass(LTX2TwoStagePipeline, LTX2Pipeline)
+
+
 def test_ltx_variants_share_denoise_loop_and_i2v_conditioning():
    assert "_denoise_loop" not in LTX2Pipeline.__dict__
    assert issubclass(LTX2Pipeline, LTXI2VConditioningMixin)
@@ -594,7 +823,8 @@ def test_denoise_executor_owns_progress_and_interrupt():
    torch.testing.assert_close(state.audio, torch.tensor(11.0))


-def test_ltx2_distilled_two_stage_executes_declarative_phase_plan():
+@pytest.mark.parametrize("pipeline_cls", [LTX2TwoStagePipeline, LTX2DistilledPipeline])
+def test_ltx_two_stage_executes_declarative_i2v_phase_plan(pipeline_cls):
    request_inputs = LTXRequestInputs(
        prompt="prompt",
        negative_prompt="",
@@ -671,12 +901,15 @@ def test_ltx2_distilled_two_stage_executes_declarative_phase_plan():
            return latents.repeat_interleave(2, dim=-2).repeat_interleave(2, dim=-1)

    prompt_context_sentinel = prompt_context
-    pipeline = object.__new__(LTX2DistilledPipeline)
+    pipeline = object.__new__(pipeline_cls)
    torch.nn.Module.__init__(pipeline)
    pipeline.device = torch.device("cpu")
    pipeline.vae_spatial_compression_ratio = 32
    pipeline.vae_temporal_compression_ratio = 8
    pipeline.latent_upsampler = FakeUpsampler()
+    activated_slots = []
+    if pipeline_cls is LTX2TwoStagePipeline:
+        pipeline._phase_adapter = SimpleNamespace(activate=activated_slots.append)
    object.__setattr__(pipeline, "_resolve_request_inputs", resolve_request_inputs)
    object.__setattr__(pipeline, "run_phase", run_phase)
    object.__setattr__(pipeline, "decode_phase", decode_phase)
@@ -686,14 +919,18 @@ def test_ltx2_distilled_two_stage_executes_declarative_phase_plan():

    assert len(phase_calls) == 2
    assert phase_calls[1][2] is prompt_context_sentinel
+    if pipeline_cls is LTX2TwoStagePipeline:
+        assert activated_slots == [None, "ltx_distilled"]
    torch.testing.assert_close(output.output[0], torch.full((1, 128, 1, 2, 2), 3.0))
-    torch.testing.assert_close(output.output[1], torch.full((1, 8, 1, 2), 4.0))
+    expected_audio = 2.0 if pipeline_cls is LTX2TwoStagePipeline else 4.0
+    torch.testing.assert_close(output.output[1], torch.full((1, 8, 1, 2), expected_audio))


-def test_ltx2_distilled_stage2_reapplies_final_resolution_i2v_conditioning():
+@pytest.mark.parametrize("pipeline_cls", [LTX2TwoStagePipeline, LTX2DistilledPipeline])
+def test_ltx_two_stage_stage2_reapplies_final_resolution_i2v_conditioning(pipeline_cls):
    from vllm_omni.diffusion.models.ltx2 import ltx2_latents

-    pipeline = object.__new__(LTX2DistilledPipeline)
+    pipeline = object.__new__(pipeline_cls)
    torch.nn.Module.__init__(pipeline)
    pipeline.vae_spatial_compression_ratio = 32
    pipeline.vae_temporal_compression_ratio = 8
@@ -750,6 +987,18 @@ def test_ltx2_distilled_two_stage_recipe_is_fixed_positive_only():
    assert (LTX2_DISTILLED_TWO_STAGE_RECIPE.height, LTX2_DISTILLED_TWO_STAGE_RECIPE.width) == (1024, 1536)


+def test_ltx23_full_distilled_route_uses_merged_weights_without_lora():
+    assert resolve_ltx_component_profile("distilled_two_stage", "2.3") is LTX23_DISTILLED_COMPONENT_PROFILE
+    assert resolve_ltx_pipeline_recipe("distilled_two_stage", "2.3") is LTX23_DISTILLED_TWO_STAGE_RECIPE
+    assert LTX23_DISTILLED_TWO_STAGE_RECIPE is LTX2_DISTILLED_TWO_STAGE_RECIPE
+    assert all(phase.adapter_slot is None for phase in LTX23_DISTILLED_TWO_STAGE_RECIPE.phases)
+    assert LTX23_DISTILLED_COMPONENT_PROFILE.distilled_lora_filename is None
+    assert LTX23_DISTILLED_COMPONENT_PROFILE.vocoder_cls is LTX23_COMPONENT_PROFILE.vocoder_cls
+    assert LTX23_DISTILLED_COMPONENT_PROFILE.latent_upsampler_filename == (
+        "ltx-2.3-spatial-upscaler-x2-1.1.safetensors"
+    )
+
+
@pytest.mark.parametrize(
    ("request_inputs", "error"),
    [
@@ -958,6 +1207,98 @@ def test_ltx_custom_sigmas_bypass_scheduler_shifting():
    torch.testing.assert_close(timesteps, expected[:-1] * scheduler.config.num_train_timesteps)


+@pytest.mark.parametrize(
+    ("recipe", "steps", "stg_block"),
+    [
+        (LTX2_TWO_STAGE_RECIPE, 40, 29),
+        (LTX23_TWO_STAGE_RECIPE, 30, 28),
+    ],
+)
+def test_ltx_official_two_stage_recipes_only_vary_by_model_defaults(recipe, steps, stg_block):
+    stage1, stage2 = recipe.phases
+    assert (recipe.width, recipe.height) == (1536, 1024)
+    assert recipe.num_inference_steps == steps
+    assert stage1.spatial_downscale == 2
+    assert stage1.adapter_slot is None
+    assert stage1.guidance.video.stg_blocks == (stg_block,)
+    assert stage1.guidance.audio.stg_blocks == (stg_block,)
+    assert stage2.input_transform == "spatial_upsample"
+    assert stage2.adapter_slot == "ltx_distilled"
+    assert stage2.guidance == LTXGuidanceSpec.positive_only()
+    assert stage2.num_inference_steps == 3
+    assert stage2.sigmas == (0.909375, 0.725, 0.421875, 0.0)
+    # Official Stage 2 refines video only; its audio result is discarded.
+    assert (recipe.video_output_phase, recipe.audio_output_phase) == (1, 0)
+
+
+def test_ltx_two_stage_entries_select_the_official_adapter_slots():
+    phase_switches = []
+    ordinary_pipeline = object.__new__(LTX2TwoStagePipeline)
+    ordinary_pipeline._phase_adapter = SimpleNamespace(activate=phase_switches.append)
+    for phase in LTX2_TWO_STAGE_RECIPE.phases:
+        ordinary_pipeline._enter_phase(phase)
+    assert phase_switches == [None, "ltx_distilled"]
+
+    distilled_pipeline = object.__new__(LTX2DistilledPipeline)
+    for phase in LTX2_DISTILLED_TWO_STAGE_RECIPE.phases:
+        distilled_pipeline._enter_phase(phase)
+
+
+def test_ltx23_gate_projection_declares_its_global_column_parallel_layout(monkeypatch):
+    class FakeColumn(torch.nn.Module):
+        def __init__(
+            self,
+            input_size,
+            output_size,
+            *,
+            bias,
+            gather_output,
+            return_bias,
+            quant_config,
+            prefix,
+        ):
+            super().__init__()
+            self.input_size = input_size
+            self.output_size = output_size
+            self.bias_enabled = bias
+            self.gather_output = gather_output
+            self.return_bias = return_bias
+            self.quant_config = quant_config
+            self.prefix = prefix
+
+    class FakeRow(torch.nn.Module):
+        def __init__(self, *args, **kwargs):
+            super().__init__()
+
+    class FakeAttention(torch.nn.Module):
+        def __init__(self, **kwargs):
+            super().__init__()
+
+    monkeypatch.setattr(ltx2_transformer, "ColumnParallelLinear", FakeColumn)
+    monkeypatch.setattr(ltx2_transformer, "RowParallelLinear", FakeRow)
+    monkeypatch.setattr(ltx2_transformer, "_LTX2ParallelAttention", FakeAttention)
+    monkeypatch.setattr(ltx2_transformer, "get_tensor_model_parallel_world_size", lambda: 2)
+
+    attention = ltx2_transformer.LTX2Attention(
+        query_dim=16,
+        heads=4,
+        dim_head=4,
+        qk_norm="rms_norm",
+        apply_gated_attention=True,
+        pack_qkv=False,
+        prefix="transformer_blocks.0.attn1",
+    )
+
+    gate = attention.to_gate_logits
+    assert isinstance(gate, FakeColumn)
+    assert gate.input_size == 16
+    assert gate.output_size == 4
+    assert gate.gather_output is False
+    assert gate.return_bias is False
+    assert gate.quant_config is None
+    assert gate.prefix == "transformer_blocks.0.attn1.to_gate_logits"
+
+
 class TestLTXRequestParsing:
    def test_request_resolves_official_and_explicit_negative_prompts_per_sample(self):
        from vllm_omni.diffusion.request import OmniDiffusionRequest
@@ -1432,6 +1773,7 @@ class TestPipelineComponents:

    def test_distilled_profile_keeps_upsampler_in_managed_resident_components(self):
        assert LTX2_DISTILLED_COMPONENT_PROFILE.resident_modules == ("vocoder", "latent_upsampler")
+        assert LTX23_DISTILLED_COMPONENT_PROFILE.resident_modules == ("vocoder", "latent_upsampler")


 class TestLTX23DecodeConditioning:
@@ -1477,6 +1819,11 @@ class TestRegistryIntegration:
        from vllm_omni.diffusion.registry import _DIFFUSION_MODELS

        assert _DIFFUSION_MODELS["LTX2Pipeline"] == ("ltx2", "pipeline_ltx2", "LTX2Pipeline")
+        assert _DIFFUSION_MODELS["LTX2TwoStagePipeline"] == (
+            "ltx2",
+            "pipeline_ltx2_two_stage",
+            "LTX2TwoStagePipeline",
+        )
        assert _DIFFUSION_MODELS["LTX2DistilledPipeline"] == (
            "ltx2",
            "pipeline_ltx2_two_stage",
@@ -1488,6 +1835,9 @@ class TestRegistryIntegration:
            "LTX23ImageToVideoPipeline",
            "LTX2TwoStagesPipeline",
            "LTX2ImageToVideoTwoStagesPipeline",
+            "LTX2ImageToVideoTwoStagePipeline",
+            "LTX23TwoStagePipeline",
+            "LTX23ImageToVideoTwoStagePipeline",
        }
        assert removed_entries.isdisjoint(_DIFFUSION_MODELS)

@@ -1497,6 +1847,7 @@ class TestRegistryIntegration:

        expected = [
            "LTX2Pipeline",
+            "LTX2TwoStagePipeline",
            "LTX2DistilledPipeline",
            "LTX2T2VDMD2Pipeline",
            "LTX2I2VDMD2Pipeline",
@@ -1509,6 +1860,7 @@ class TestRegistryIntegration:
        "model_class_name",
        [
            "LTX2Pipeline",
+            "LTX2TwoStagePipeline",
            "LTX2DistilledPipeline",
            "LTX2T2VDMD2Pipeline",
            "LTX2I2VDMD2Pipeline",
@@ -1,13 +1,20 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project

+from types import SimpleNamespace
+
 import pytest
+import torch
 from torch import nn

+from vllm_omni.diffusion.attention.backends.abstract import AttentionMetadata
 from vllm_omni.diffusion.cache.cachedit import CacheDiTAdapterConfig
-from vllm_omni.diffusion.models.ltx2.ltx2_transformer import LTX2VideoTransformer3DModel, _make_rms_norm
-
-pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+from vllm_omni.diffusion.models.ltx2.ltx2_transformer import (
+    LTX2AudioVideoAttnProcessor,
+    LTX2VideoTransformer3DModel,
+    _LTX2ParallelAttention,
+    _make_rms_norm,
+)

 pytestmark = [pytest.mark.core_model, pytest.mark.diffusion, pytest.mark.cpu]

@@ -66,3 +73,83 @@ def test_ltx_sp_plan_shards_video_and_audio_timesteps_together(rope_type):
    assert video_timestep.expected_dims == audio_timestep.expected_dims == 2
    assert not video_timestep.split_output
    assert not audio_timestep.split_output
+
+
+def test_ltx_sp_keeps_cross_attention_key_padding_mask(monkeypatch):
+    monkeypatch.setattr(LTX2AudioVideoAttnProcessor, "_is_sp_enabled", staticmethod(lambda: True))
+    processor = LTX2AudioVideoAttnProcessor()
+    hidden_states = torch.zeros(1, 3, 4)
+    encoder_hidden_states = torch.zeros(1, 4, 4)
+    key_padding_mask = torch.tensor([[True, True, True, False]])
+
+    actual = processor._prepare_attention_mask(
+        SimpleNamespace(),
+        hidden_states,
+        encoder_hidden_states,
+        key_padding_mask,
+        batch_size=1,
+        sequence_length=4,
+    )
+
+    torch.testing.assert_close(actual, key_padding_mask[:, None, None, :])
+
+
+def _fake_ltx_parallel_attention():
+    calls = []
+
+    class Backend:
+        @staticmethod
+        def get_name():
+            return "FLASH_ATTN"
+
+        @staticmethod
+        def supports_attention_mask():
+            return True
+
+    def run(name):
+        def forward(query, _key, _value, _metadata):
+            calls.append(name)
+            return torch.zeros_like(query)
+
+        return forward
+
+    layer = object.__new__(_LTX2ParallelAttention)
+    nn.Module.__init__(layer)
+    layer.attn_backend = Backend
+    layer.attention = SimpleNamespace(forward=run("native"))
+    layer.sdpa_fallback = SimpleNamespace(forward=run("sdpa"))
+    layer.backend_pref = "FLASH_ATTN"
+    return layer, calls
+
+
+@pytest.mark.parametrize(("query_length", "key_length"), [(4, 4), (3, 4)])
+def test_ltx_masked_cross_attention_uses_sdpa(query_length, key_length):
+    layer, calls = _fake_ltx_parallel_attention()
+    query = torch.zeros(1, query_length, 2, 4, dtype=torch.bfloat16)
+    key = value = torch.zeros(1, key_length, 2, 4, dtype=torch.bfloat16)
+    key_padding_mask = torch.tensor([[True] * (key_length - 1) + [False]])
+    metadata = AttentionMetadata(attn_mask=key_padding_mask[:, None, None, :])
+
+    layer._run_local_attention(query, key, value, metadata)
+
+    assert calls == ["sdpa"]
+
+
+def test_ltx_masked_self_attention_uses_sdpa():
+    layer, calls = _fake_ltx_parallel_attention()
+    query = key = value = torch.zeros(1, 4, 2, 4, dtype=torch.bfloat16)
+    key_padding_mask = torch.tensor([[True, True, True, False]])
+    metadata = AttentionMetadata(attn_mask=key_padding_mask[:, None, None, :])
+
+    layer._run_local_attention(query, key, value, metadata)
+
+    assert calls == ["sdpa"]
+
+
+def test_ltx_unmasked_cross_attention_keeps_native_backend():
+    layer, calls = _fake_ltx_parallel_attention()
+    query = key = value = torch.zeros(1, 4, 2, 4, dtype=torch.bfloat16)
+
+    layer._run_local_attention(query, key, value, None)
+
+    assert calls == ["native"]
@@ -2,6 +2,7 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project

 import json
+import sys
 from multiprocessing.reduction import ForkingPickler
 from types import SimpleNamespace
 from unittest.mock import Mock
@@ -1428,34 +1429,219 @@ def test_r7_r8_ref2va_video_segment_matrix(monkeypatch, tmp_path, case, start_ti
    assert transcode_calls[0][1]["target_frame_count"] == 209


-def test_ref2va_qwen_sampling_decodes_prepared_video_once(monkeypatch):
+def test_ref2va_qwen_sampling_uses_one_selective_decode(monkeypatch):
    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module

-    decoded = np.arange(25 * 2 * 2 * 3, dtype=np.uint8).reshape(25, 2, 2, 3)
-    load_calls = []
+    decoded = np.arange(3 * 2 * 2 * 3, dtype=np.uint8).reshape(3, 2, 2, 3)
+    decode_calls = []
    monkeypatch.setattr(
        reference_video_module,
        "_probe_video",
-        lambda _path: {"frame_count": len(decoded)},
+        lambda _path: {
+            "frame_count": 25,
+            "width": 2,
+            "height": 2,
+        },
    )
    monkeypatch.setattr(
        reference_video_module,
-        "load_video_frames",
-        lambda path: load_calls.append(path) or decoded,
+        "_decode_video_frames_ffmpeg",
+        lambda path, **kwargs: decode_calls.append((path, kwargs)) or decoded,
    )

    sampled = reference_video_module.sample_reference_video_frames("prepared.mp4")

-    assert load_calls == ["prepared.mp4"]
-    for frame, expected in zip(
-        sampled["frames"],
-        (decoded[0], decoded[12], decoded[24]),
-        strict=True,
-    ):
+    assert decode_calls == [
+        (
+            "prepared.mp4",
+            {
+                "frame_count": 25,
+                "indices": [0, 12, 24],
+                "width": 2,
+                "height": 2,
+            },
+        )
+    ]
+    for frame, expected in zip(sampled["frames"], decoded, strict=True):
        np.testing.assert_array_equal(frame, expected)
    assert sampled["block_timestamps"] == [0.25, 1.0]


+@pytest.mark.parametrize(
+    ("indices", "frame_count"),
+    [([0, 12], 25), (None, 2)],
+)
+def test_ref2va_ffmpeg_decode_reads_exact_rgb_frames(monkeypatch, indices, frame_count):
+    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module
+
+    expected = np.arange(2 * 2 * 3 * 3, dtype=np.uint8).reshape(2, 2, 3, 3)
+    observed = []
+
+    def fake_run(command, **kwargs):
+        observed.append((command, kwargs))
+        return SimpleNamespace(stdout=expected.tobytes())
+
+    monkeypatch.setattr(reference_video_module.subprocess, "run", fake_run)
+
+    actual = reference_video_module._decode_video_frames_ffmpeg(
+        "prepared.mp4",
+        frame_count=frame_count,
+        indices=indices,
+        width=3,
+        height=2,
+    )
+
+    np.testing.assert_array_equal(actual, expected)
+    command, kwargs = observed[0]
+    assert command[0] == "ffmpeg"
+    assert command[command.index("-frames:v") + 1] == "2"
+    if indices is None:
+        assert "-vf" not in command
+    else:
+        assert command[command.index("-vf") + 1] == "select=eq(n\\,0)+eq(n\\,12)"
+    assert command[-1] == "pipe:1"
+    assert kwargs == {"check": True, "capture_output": True}
+
+
+def test_ref2va_load_video_frames_uses_ffmpeg_without_decord(monkeypatch):
+    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module
+
+    expected = np.zeros((2, 2, 3, 3), dtype=np.uint8)
+    decode_calls = []
+    monkeypatch.setitem(sys.modules, "decord", None)
+    monkeypatch.setattr(
+        reference_video_module,
+        "_probe_video",
+        lambda _path: {
+            "frame_count": 2,
+            "width": 3,
+            "height": 2,
+        },
+    )
+    monkeypatch.setattr(
+        reference_video_module,
+        "_decode_video_frames_ffmpeg",
+        lambda path, **kwargs: decode_calls.append((path, kwargs)) or expected,
+    )
+
+    actual = reference_video_module.load_video_frames("prepared.mp4")
+
+    assert actual is expected
+    assert decode_calls == [
+        (
+            "prepared.mp4",
+            {
+                "frame_count": 2,
+                "width": 3,
+                "height": 2,
+            },
+        )
+    ]
+
+
+@pytest.mark.parametrize(
+    ("kwargs", "message"),
+    [
+        ({"frame_count": 0, "width": 3, "height": 2}, "video has no frames"),
+        ({"frame_count": 1, "width": 0, "height": 2}, "invalid dimensions"),
+        (
+            {"frame_count": 1, "indices": [], "width": 3, "height": 2},
+            "invalid frame indices",
+        ),
+    ],
+)
+def test_ref2va_ffmpeg_decode_rejects_invalid_metadata(kwargs, message):
+    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module
+    from vllm_omni.errors import OmniClientError
+
+    with pytest.raises(OmniClientError, match=message):
+        reference_video_module._decode_video_frames_ffmpeg("prepared.mp4", **kwargs)
+
+
+def test_ref2va_probe_uses_container_frame_count_without_scan(monkeypatch):
+    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module
+
+    calls = []
+    probe = {
+        "streams": [
+            {
+                "codec_type": "video",
+                "width": 4,
+                "height": 2,
+                "r_frame_rate": "24/1",
+                "nb_frames": "3",
+                "duration": "0.125",
+                "sample_aspect_ratio": "1:1",
+            }
+        ],
+        "format": {"duration": "0.125", "size": "123"},
+    }
+    monkeypatch.setattr(
+        reference_video_module.subprocess,
+        "run",
+        lambda command, **kwargs: calls.append((command, kwargs)) or SimpleNamespace(stdout=json.dumps(probe)),
+    )
+
+    metadata = reference_video_module._probe_video("prepared.mp4")
+
+    assert metadata["frame_count"] == 3
+    assert len(calls) == 1
+    assert "-count_frames" not in calls[0][0]
+
+
+def test_ref2va_probe_counts_frames_only_when_metadata_is_missing(monkeypatch):
+    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module
+
+    calls = []
+    stream = {
+        "codec_type": "video",
+        "width": 4,
+        "height": 2,
+        "r_frame_rate": "24/1",
+        "duration": "0.125",
+        "sample_aspect_ratio": "1:1",
+    }
+    first_probe = {
+        "streams": [stream],
+        "format": {"duration": "0.125", "size": "123"},
+    }
+    second_probe = {
+        "streams": [{**stream, "nb_read_frames": "3"}],
+        "format": {"duration": "0.125", "size": "123"},
+    }
+    monkeypatch.setattr(
+        reference_video_module.subprocess,
+        "run",
+        lambda command, **kwargs: calls.append((command, kwargs))
+        or SimpleNamespace(stdout=json.dumps(first_probe if len(calls) == 1 else second_probe)),
+    )
+
+    metadata = reference_video_module._probe_video("prepared.mp4")
+
+    assert metadata["frame_count"] == 3
+    assert len(calls) == 2
+    assert "-count_frames" in calls[1][0]
+
+
+def test_ref2va_ffmpeg_decode_rejects_incomplete_output(monkeypatch):
+    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module
+    from vllm_omni.errors import OmniClientError
+
+    monkeypatch.setattr(
+        reference_video_module.subprocess,
+        "run",
+        lambda *_args, **_kwargs: SimpleNamespace(stdout=b"short"),
+    )
+
+    with pytest.raises(OmniClientError, match="decoded 5 bytes.*expected 18"):
+        reference_video_module._decode_video_frames_ffmpeg(
+            "prepared.mp4",
+            frame_count=1,
+            width=3,
+            height=2,
+        )
+
+
 def test_ref2va_two_video_recipe_tolerates_container_rounding(monkeypatch, tmp_path):
    from vllm_omni.diffusion.models.minimax_h3 import reference_video as reference_video_module

@@ -0,0 +1,104 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import torch
+
+from vllm_omni.platforms import current_omni_platform
+
+pytestmark = [pytest.mark.core_model, pytest.mark.diffusion]
+
+
+@pytest.mark.skipif(not current_omni_platform.is_npu(), reason="requires Ascend NPU")
+@pytest.mark.parametrize("num_q_heads,num_kv_heads,seq_len", [(8, 2, 16), (64, 8, 1)])
+def test_qwen3_vl_npu_rope_matches_bnsd_reference(
+    num_q_heads: int,
+    num_kv_heads: int,
+    seq_len: int,
+) -> None:
+    from vllm_omni.platforms.npu.models.minimax_h3 import _apply_rotary_pos_emb_npu
+
+    device = torch.device("npu")
+    dtype = torch.bfloat16
+    q = torch.randn(1, num_q_heads, seq_len, 128, device=device, dtype=dtype)
+    k = torch.randn(1, num_kv_heads, seq_len, 128, device=device, dtype=dtype)
+    cos = torch.randn(1, seq_len, 128, device=device, dtype=dtype)
+    sin = torch.randn(1, seq_len, 128, device=device, dtype=dtype)
+
+    q_out, k_out = _apply_rotary_pos_emb_npu(q, k, cos, sin)
+    cos_bnsd = cos.unsqueeze(1)
+    sin_bnsd = sin.unsqueeze(1)
+
+    def rotate_half(x: torch.Tensor) -> torch.Tensor:
+        x1, x2 = x.chunk(2, dim=-1)
+        return torch.cat((-x2, x1), dim=-1)
+
+    torch.testing.assert_close(
+        q_out,
+        q * cos_bnsd + rotate_half(q) * sin_bnsd,
+        atol=2e-2,
+        rtol=2e-2,
+    )
+    torch.testing.assert_close(
+        k_out,
+        k * cos_bnsd + rotate_half(k) * sin_bnsd,
+        atol=2e-2,
+        rtol=2e-2,
+    )
+
+
+@pytest.mark.skipif(not current_omni_platform.is_npu(), reason="requires Ascend NPU")
+def test_qwen3_vl_npu_rope_preserves_batched_mrope_positions() -> None:
+    """Keep per-sample Qwen3-VL M-RoPE frequencies on the BNSD NPU path."""
+    from types import SimpleNamespace
+
+    from vllm_omni.diffusion.models.minimax_h3.encoder import (
+        MiniMaxH3Qwen3VLTextRotaryEmbedding,
+    )
+    from vllm_omni.platforms.npu.models.minimax_h3 import _apply_rotary_pos_emb_npu
+
+    device = torch.device("npu")
+    dtype = torch.bfloat16
+    batch_size, seq_len, head_dim = 2, 16, 128
+    config = SimpleNamespace(
+        hidden_size=head_dim,
+        num_attention_heads=1,
+        rope_parameters={
+            "rope_theta": 1_000_000.0,
+            "mrope_section": [24, 20, 20],
+        },
+    )
+    rope = MiniMaxH3Qwen3VLTextRotaryEmbedding(config).to(device)
+    positions = torch.arange(seq_len, device=device)
+    position_ids = torch.stack(
+        (
+            torch.stack((positions, positions + 7)),
+            torch.stack((positions * 2, positions * 2 + 3)),
+            torch.stack((positions * 3, positions * 3 + 5)),
+        )
+    )
+    q = torch.randn(batch_size, 8, seq_len, head_dim, device=device, dtype=dtype)
+    k = torch.randn(batch_size, 2, seq_len, head_dim, device=device, dtype=dtype)
+    cos, sin = rope(q, position_ids)
+
+    q_out, k_out = _apply_rotary_pos_emb_npu(q, k, cos, sin)
+    cos_bnsd = cos.unsqueeze(1)
+    sin_bnsd = sin.unsqueeze(1)
+
+    def rotate_half(x: torch.Tensor) -> torch.Tensor:
+        x1, x2 = x.chunk(2, dim=-1)
+        return torch.cat((-x2, x1), dim=-1)
+
+    assert not torch.equal(cos[0], cos[1])
+    torch.testing.assert_close(
+        q_out,
+        q * cos_bnsd + rotate_half(q) * sin_bnsd,
+        atol=2e-2,
+        rtol=2e-2,
+    )
+    torch.testing.assert_close(
+        k_out,
+        k * cos_bnsd + rotate_half(k) * sin_bnsd,
+        atol=2e-2,
+        rtol=2e-2,
+    )
@@ -0,0 +1,46 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import torch
+import torch.nn.functional as F
+
+from vllm_omni.platforms import current_omni_platform
+
+pytestmark = [pytest.mark.core_model, pytest.mark.diffusion]
+
+
+@pytest.mark.skipif(not current_omni_platform.is_npu(), reason="requires Ascend NPU")
+@pytest.mark.parametrize("shape", [(1, 16, 256), (2, 1, 256)])
+def test_qwen3_vl_npu_swiglu_matches_split_reference(
+    shape: tuple[int, int, int],
+) -> None:
+    from vllm_omni.platforms.npu.models.minimax_h3 import npu_swiglu_from_packed
+
+    packed = torch.randn(*shape, device="npu", dtype=torch.bfloat16)
+    gate, up = packed.chunk(2, dim=-1)
+
+    torch.testing.assert_close(
+        npu_swiglu_from_packed(packed),
+        F.silu(gate) * up,
+        atol=2e-2,
+        rtol=2e-2,
+    )
+
+
+@pytest.mark.skipif(not current_omni_platform.is_npu(), reason="requires Ascend NPU")
+def test_qwen3_vl_npu_swiglu_matches_two_projection_reference() -> None:
+    from vllm_omni.platforms.npu.models.minimax_h3 import npu_swiglu_from_packed
+
+    x = torch.randn(1, 16, 128, device="npu", dtype=torch.bfloat16)
+    weight = torch.randn(256, 128, device="npu", dtype=torch.bfloat16)
+    packed = F.linear(x, weight)
+    gate = F.linear(x, weight[:128])
+    up = F.linear(x, weight[128:])
+
+    torch.testing.assert_close(
+        npu_swiglu_from_packed(packed),
+        F.silu(gate) * up,
+        atol=2e-2,
+        rtol=2e-2,
+    )
@@ -36,15 +36,23 @@ class _ToyKernel:
 # --- helper to build a toy module simulating online-FP8 post-load state ---


-def _make_fp8_toy_module(out_features: int = 16, in_features: int = 32):
+def _make_fp8_toy_module(
+    out_features: int = 16,
+    in_features: int = 32,
+    *,
+    quant_method_cls=None,
+):
    """Return a module whose weight is ``qweight.t()`` (non-contiguous) and whose
    ``quant_method`` is an ``Fp8LinearMethod`` with a ``_ToyKernel``.
    """
    from vllm.model_executor.layers.quantization.fp8 import Fp8LinearMethod

+    if quant_method_cls is None:
+        quant_method_cls = Fp8LinearMethod
+
    # Create the quant-method instance without calling __init__ (avoids
    # heavyweight upstream config dependencies).  isinstance() still works.
-    qm = object.__new__(Fp8LinearMethod)
+    qm = object.__new__(quant_method_cls)
    qm.fp8_linear = _ToyKernel()

    # randn does not support float8 on CPU; zeros does and the test only
@@ -104,6 +112,28 @@ def test_rewrites_non_contiguous_weight():
    assert module.weight.shape == (16, 32)


+def test_rewrites_per_tensor_online_fp8_weight():
+    """Regression test for vLLM's new online FP8 frontend (#45463)."""
+    from vllm.model_executor.layers.quantization.online.fp8 import (
+        Fp8PerTensorOnlineLinearMethod,
+    )
+
+    module = _make_fp8_toy_module(
+        16,
+        32,
+        quant_method_cls=Fp8PerTensorOnlineLinearMethod,
+    )
+
+    n_rewritten_layers = prepare_fp8_layers_for_fsdp(module)
+
+    assert n_rewritten_layers == 1
+    assert module.weight.is_contiguous()
+    assert module.weight.shape == (16, 32)
+    w, *_ = module.quant_method.fp8_linear._get_layer_params(module)
+    assert w.shape == (32, 16)
+    assert w.stride() == (1, 32)
+
+
 def test_patched_get_layer_params_returns_transposed_view():
    module = _make_fp8_toy_module(16, 32)
    prepare_fp8_layers_for_fsdp(module)
@@ -20,6 +20,11 @@ from PIL import Image
 def _parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser()
    parser.add_argument("--backend", choices=("official", "omni"), required=True)
+    parser.add_argument(
+        "--pipeline-kind",
+        choices=("one_stage", "distilled", "two_stage"),
+        default="one_stage",
+    )
    parser.add_argument("--request", type=Path, required=True)
    parser.add_argument("--output-dir", type=Path, required=True)
    parser.add_argument("--model")
@@ -27,6 +32,8 @@ def _parse_args() -> argparse.Namespace:
    parser.add_argument("--official-root", type=Path)
    parser.add_argument("--checkpoint", type=Path)
    parser.add_argument("--gemma-root", type=Path)
+    parser.add_argument("--spatial-upsampler", type=Path)
+    parser.add_argument("--distilled-lora", type=Path)
    parser.add_argument("--enable-layerwise-offload", action="store_true")
    return parser.parse_args()

@@ -91,9 +98,13 @@ def _configure_official_sdpa(pipeline: Any) -> None:
        attention=attention,
        masked_attention=attention,
    )
-    owners_and_attributes = (
-        (pipeline.stage, "_transformer_builder"),
+    owners_and_attributes = [
        (pipeline.prompt_encoder, "_embeddings_processor_builder"),
+    ]
+    owners_and_attributes.extend(
+        (stage, "_transformer_builder")
+        for name in ("stage", "stage_1", "stage_2")
+        if (stage := getattr(pipeline, name, None)) is not None
    )
    for owner, attribute in owners_and_attributes:
        builder = getattr(owner, attribute)
@@ -104,6 +115,12 @@ def _configure_official_sdpa(pipeline: Any) -> None:
        )


+def _require_path(path: Path | None, description: str) -> str:
+    if path is None or not path.is_file():
+        raise ValueError(f"Official {description} is required and must exist: {path}")
+    return str(path)
+
+
@torch.inference_mode()
 def _run_official(args: argparse.Namespace, request: dict[str, Any]) -> None:
    if args.official_root is None or args.checkpoint is None or args.gemma_root is None:
@@ -113,16 +130,54 @@ def _run_official(args: argparse.Namespace, request: dict[str, Any]) -> None:
    torch.backends.cuda.enable_cudnn_sdp(False)

    from ltx_core.components.guiders import MultiModalGuiderParams
+    from ltx_core.loader import LTXV_LORA_COMFY_RENAMING_MAP, LoraPathStrengthAndSDOps
+    from ltx_pipelines.distilled import DistilledPipeline
    from ltx_pipelines.ti2vid_one_stage import TI2VidOneStagePipeline
+    from ltx_pipelines.ti2vid_two_stages import TI2VidTwoStagesPipeline
    from ltx_pipelines.utils.args import ImageConditioningInput
    from ltx_pipelines.utils.types import OffloadMode

-    pipeline = TI2VidOneStagePipeline(
-        checkpoint_path=str(args.checkpoint),
-        gemma_root=str(args.gemma_root),
-        loras=(),
-        offload_mode=OffloadMode.CPU if args.enable_layerwise_offload else OffloadMode.NONE,
-    )
+    if request.get("pipeline_kind", "one_stage") != args.pipeline_kind:
+        raise ValueError(
+            f"Request pipeline kind {request.get('pipeline_kind')!r} does not match {args.pipeline_kind!r}"
+        )
+    checkpoint = _require_path(args.checkpoint, "checkpoint")
+    gemma_root = str(args.gemma_root)
+    if not Path(gemma_root).is_dir():
+        raise ValueError(f"Official Gemma root is required and must be a directory: {gemma_root}")
+    offload_mode = OffloadMode.CPU if args.enable_layerwise_offload else OffloadMode.NONE
+    pipeline: Any
+    if args.pipeline_kind == "one_stage":
+        pipeline = TI2VidOneStagePipeline(
+            checkpoint_path=checkpoint,
+            gemma_root=gemma_root,
+            loras=(),
+            offload_mode=offload_mode,
+        )
+    elif args.pipeline_kind == "distilled":
+        pipeline = DistilledPipeline(
+            distilled_checkpoint_path=checkpoint,
+            spatial_upsampler_path=_require_path(args.spatial_upsampler, "spatial upsampler"),
+            gemma_root=gemma_root,
+            loras=(),
+            offload_mode=offload_mode,
+        )
+    else:
+        distilled_lora = [
+            LoraPathStrengthAndSDOps(
+                path=_require_path(args.distilled_lora, "distilled LoRA"),
+                strength=1.0,
+                sd_ops=LTXV_LORA_COMFY_RENAMING_MAP,
+            )
+        ]
+        pipeline = TI2VidTwoStagesPipeline(
+            checkpoint_path=checkpoint,
+            distilled_lora=distilled_lora,
+            spatial_upsampler_path=_require_path(args.spatial_upsampler, "spatial upsampler"),
+            gemma_root=gemma_root,
+            loras=(),
+            offload_mode=offload_mode,
+        )
    _configure_official_sdpa(pipeline)
    image_path = request.get("image")
    images = (
@@ -140,34 +195,45 @@ def _run_official(args: argparse.Namespace, request: dict[str, Any]) -> None:
            )
        ]
    )
-    video, audio = pipeline(
-        prompt=request["prompt"],
-        negative_prompt=request["negative_prompt"],
-        seed=request["seed"],
-        height=request["height"],
-        width=request["width"],
-        num_frames=request["num_frames"],
-        frame_rate=request["fps"],
-        num_inference_steps=request["num_inference_steps"],
-        video_guider_params=MultiModalGuiderParams(
-            cfg_scale=request["video_cfg_scale"],
-            stg_scale=request["video_stg_scale"],
-            rescale_scale=request["video_rescale_scale"],
-            modality_scale=request["video_modality_scale"],
-            skip_step=0,
-            stg_blocks=request["video_stg_blocks"],
-        ),
-        audio_guider_params=MultiModalGuiderParams(
-            cfg_scale=request["audio_cfg_scale"],
-            stg_scale=request["audio_stg_scale"],
-            rescale_scale=request["audio_rescale_scale"],
-            modality_scale=request["audio_modality_scale"],
-            skip_step=0,
-            stg_blocks=request["audio_stg_blocks"],
-        ),
-        images=images,
-        max_batch_size=4,
-    )
+    if args.pipeline_kind == "distilled":
+        video, audio = pipeline(
+            prompt=request["prompt"],
+            seed=request["seed"],
+            height=request["height"],
+            width=request["width"],
+            num_frames=request["num_frames"],
+            frame_rate=request["fps"],
+            images=images,
+        )
+    else:
+        video, audio = pipeline(
+            prompt=request["prompt"],
+            negative_prompt=request["negative_prompt"],
+            seed=request["seed"],
+            height=request["height"],
+            width=request["width"],
+            num_frames=request["num_frames"],
+            frame_rate=request["fps"],
+            num_inference_steps=request["num_inference_steps"],
+            video_guider_params=MultiModalGuiderParams(
+                cfg_scale=request["video_cfg_scale"],
+                stg_scale=request["video_stg_scale"],
+                rescale_scale=request["video_rescale_scale"],
+                modality_scale=request["video_modality_scale"],
+                skip_step=0,
+                stg_blocks=request["video_stg_blocks"],
+            ),
+            audio_guider_params=MultiModalGuiderParams(
+                cfg_scale=request["audio_cfg_scale"],
+                stg_scale=request["audio_stg_scale"],
+                rescale_scale=request["audio_rescale_scale"],
+                modality_scale=request["audio_modality_scale"],
+                skip_step=0,
+                stg_blocks=request["audio_stg_blocks"],
+            ),
+            images=images,
+            max_batch_size=4,
+        )
    video_tensor = torch.cat([chunk.detach().cpu() for chunk in video], dim=0)
    _save_outputs(
        args.output_dir,
@@ -179,6 +245,7 @@ def _run_official(args: argparse.Namespace, request: dict[str, Any]) -> None:
            "attention_backend": "torch_sdpa",
            "official_revision": os.environ.get("VLLM_TEST_LTX_OFFICIAL_REVISION"),
            "checkpoint": str(args.checkpoint),
+            "pipeline_kind": args.pipeline_kind,
        },
    )

@@ -193,15 +260,9 @@ def _unwrap_omni_output(output: Any) -> tuple[Any, Any, int]:
        multimodal_output = frames.multimodal_output or {}
        audio = multimodal_output.get("audio")
        audio_sample_rate = multimodal_output.get("audio_sample_rate")
-        if frames.is_pipeline_output and isinstance(frames.request_output, OmniRequestOutput):
-            frames = frames.request_output
-            multimodal_output = frames.multimodal_output or {}
-            audio = multimodal_output.get("audio", audio)
-            audio_sample_rate = multimodal_output.get("audio_sample_rate", audio_sample_rate)
-        if isinstance(frames, OmniRequestOutput):
-            if not frames.images:
-                raise ValueError("No video frames found in OmniRequestOutput")
-            frames = frames.images
+        if not frames.images:
+            raise ValueError("No video frames found in OmniRequestOutput")
+        frames = frames.images

    if isinstance(frames, list) and len(frames) == 1:
        frames = frames[0]
@@ -253,6 +314,10 @@ def _canonical_video(video: Any) -> torch.Tensor:
 def _run_omni(args: argparse.Namespace, request: dict[str, Any]) -> None:
    if args.model is None or args.model_class_name is None:
        raise ValueError("Omni backend requires --model and --model-class-name")
+    if request.get("pipeline_kind", "one_stage") != args.pipeline_kind:
+        raise ValueError(
+            f"Request pipeline kind {request.get('pipeline_kind')!r} does not match {args.pipeline_kind!r}"
+        )

    attention_config = {"default": {"backend": "TORCH_SDPA"}}

@@ -273,7 +338,8 @@ def _run_omni(args: argparse.Namespace, request: dict[str, Any]) -> None:
        parallel_config=DiffusionParallelConfig(),
    )
    try:
-        model_class_name = get_model_class_name(omni)
+        detected_model_class_name = get_model_class_name(omni)
+        model_class_name = args.model_class_name
        sampling_params = OmniDiffusionSamplingParams(
            height=request["height"],
            width=request["width"],
@@ -289,10 +355,9 @@ def _run_omni(args: argparse.Namespace, request: dict[str, Any]) -> None:
            key: value for key, value in request.items() if key.startswith("video_") or key.startswith("audio_")
        }
        apply_declared_extra_args(sampling_params, get_extra_body_params(model_class_name), guidance)
-        prompt: dict[str, Any] = {
-            "prompt": request["prompt"],
-            "negative_prompt": request["negative_prompt"],
-        }
+        prompt: dict[str, Any] = {"prompt": request["prompt"]}
+        if request.get("negative_prompt"):
+            prompt["negative_prompt"] = request["negative_prompt"]
        image_path = request.get("image")
        if image_path is not None:
            with Image.open(str(image_path)) as source_image:
@@ -312,6 +377,8 @@ def _run_omni(args: argparse.Namespace, request: dict[str, Any]) -> None:
                "attention_backend": "torch_sdpa",
                "model": args.model,
                "model_class_name": model_class_name,
+                "model_config_class_name": detected_model_class_name,
+                "pipeline_kind": args.pipeline_kind,
            },
        )
    finally:
@@ -3,10 +3,9 @@

 """E2E accuracy guard against a pinned Lightricks LTX pipeline revision.

-The comparison runs both runtimes through PyTorch SDPA and uses
-``max_batch_size=4`` in the official reference to match Omni's fused guidance
-batch. Video and audio guidance use the official non-HQ one-stage defaults;
-only the generation shape and step count are reduced for CI runtime.
+The original reduced one-stage guards remain unchanged. Four complementary
+default-shape cases cover both model versions, both two-stage weight families,
+and T2V/I2V without expanding every combination into a duplicate golden.
 """

 from __future__ import annotations
@@ -17,8 +16,9 @@ import shutil
 import subprocess
 import sys
 import time
-from dataclasses import dataclass
+from dataclasses import asdict, dataclass, replace
 from pathlib import Path
+from typing import Literal

 import numpy as np
 import pytest
@@ -53,74 +53,260 @@ NEGATIVE_PROMPT = (

 # Both runtimes use PyTorch SDPA with the current Torch dispatch defaults.
 ATTENTION_BACKEND = "torch_sdpa"
-VIDEO_SSIM_MEAN_THRESHOLD = 0.95
-VIDEO_SSIM_MIN_THRESHOLD = 0.90
-VIDEO_PSNR_MEAN_THRESHOLD = 30.0
-AUDIO_RELATIVE_L2_THRESHOLD = 0.2
-AUDIO_COSINE_THRESHOLD = 0.95
+
+
+@dataclass(frozen=True)
+class LTXArtifact:
+    repo_id: str
+    filename: str
+    revision: str
+    env: str
+    repo_type: Literal["model", "dataset"] = "model"
+
+
+@dataclass(frozen=True)
+class LTXAccuracyThresholds:
+    video_ssim_mean: float
+    video_ssim_min: float
+    video_psnr_mean_db: float
+    # Zero disables gating for waveform-sensitive audio that is not expected
+    # to align with the official output.
+    audio_relative_l2: float
+    audio_cosine_similarity: float
+
+
+STRICT_THRESHOLDS = LTXAccuracyThresholds(
+    video_ssim_mean=0.95,
+    video_ssim_min=0.90,
+    video_psnr_mean_db=30.0,
+    audio_relative_l2=0.20,
+    audio_cosine_similarity=0.95,
+)


@dataclass(frozen=True)
 class LTXAccuracyCase:
    name: str
+    pipeline_kind: Literal["one_stage", "distilled", "two_stage"]
    model_id: str
    model_revision: str
    model_env: str
    model_class_name: str
-    checkpoint_repo: str
-    checkpoint_filename: str
-    checkpoint_revision: str
-    checkpoint_env: str
-    stg_block: int
+    checkpoint: LTXArtifact
+    width: int
+    height: int
+    num_frames: int
+    num_inference_steps: int
+    seed: int
+    stg_block: int | None
+    thresholds: LTXAccuracyThresholds
    prompt: str = PROMPT
-    image_repo: str | None = None
-    image_filename: str | None = None
-    image_revision: str | None = None
+    gemma_model_id: str | None = None
+    gemma_model_revision: str | None = None
+    gemma_model_env: str | None = None
+    image: LTXArtifact | None = None
+    spatial_upsampler: LTXArtifact | None = None
+    distilled_lora: LTXArtifact | None = None
+    # Official uses block streaming; Omni uses layerwise offload. Both preserve
+    # bf16 arithmetic while keeping full-resolution two-stage cases on one H100.
+    enable_layerwise_offload: bool = False


-CASES = (
+LTX2_REVISION = "47da56e2ad66ce4125a9922b4a8826bf407f9d0a"
+LTX23_REVISION = "4229404625088d21c4f112eb640fb04a0900ee25"
+LTX2_CHECKPOINT = LTXArtifact(
+    repo_id="Lightricks/LTX-2",
+    filename="ltx-2-19b-dev.safetensors",
+    revision=LTX2_REVISION,
+    env="VLLM_TEST_LTX2_OFFICIAL_CHECKPOINT",
+)
+LTX2_DISTILLED_CHECKPOINT = replace(
+    LTX2_CHECKPOINT,
+    filename="ltx-2-19b-distilled.safetensors",
+    env="VLLM_TEST_LTX2_DISTILLED_OFFICIAL_CHECKPOINT",
+)
+LTX2_UPSAMPLER = replace(
+    LTX2_CHECKPOINT,
+    filename="ltx-2-spatial-upscaler-x2-1.0.safetensors",
+    env="VLLM_TEST_LTX2_UPSAMPLER",
+)
+LTX2_DISTILLED_LORA = replace(
+    LTX2_CHECKPOINT,
+    filename="ltx-2-19b-distilled-lora-384.safetensors",
+    env="VLLM_TEST_LTX2_DISTILLED_LORA",
+)
+LTX23_CHECKPOINT = LTXArtifact(
+    repo_id="Lightricks/LTX-2.3",
+    filename="ltx-2.3-22b-dev.safetensors",
+    revision=LTX23_REVISION,
+    env="VLLM_TEST_LTX23_OFFICIAL_CHECKPOINT",
+)
+LTX23_DISTILLED_CHECKPOINT = replace(
+    LTX23_CHECKPOINT,
+    # The pinned Diffusers checkpoint contains the original merged distilled
+    # Transformer. Keep the reference Transformer on the same version; the
+    # spatial upsampler is versioned independently and remains on 1.1 below.
+    filename="ltx-2.3-22b-distilled.safetensors",
+    env="VLLM_TEST_LTX23_DISTILLED_OFFICIAL_CHECKPOINT",
+)
+LTX23_UPSAMPLER = replace(
+    LTX23_CHECKPOINT,
+    filename="ltx-2.3-spatial-upscaler-x2-1.1.safetensors",
+    env="VLLM_TEST_LTX23_UPSAMPLER",
+)
+LTX23_DISTILLED_LORA = replace(
+    LTX23_CHECKPOINT,
+    filename="ltx-2.3-22b-distilled-lora-384-1.1.safetensors",
+    env="VLLM_TEST_LTX23_DISTILLED_LORA",
+)
+I2V_IMAGE = LTXArtifact(
+    repo_id="huggingface/documentation-images",
+    filename="diffusers/svd/rocket.png",
+    revision="645d8364f0c7a101180b364811b5a11a362e4010",
+    env="VLLM_TEST_LTX_I2V_IMAGE",
+    repo_type="dataset",
+)
+
+
+LEGACY_CASES = (
    LTXAccuracyCase(
        name="ltx2",
+        pipeline_kind="one_stage",
        model_id="Lightricks/LTX-2",
-        model_revision="47da56e2ad66ce4125a9922b4a8826bf407f9d0a",
+        model_revision=LTX2_REVISION,
        model_env="VLLM_TEST_LTX2_MODEL",
        model_class_name="LTX2Pipeline",
-        checkpoint_repo="Lightricks/LTX-2",
-        checkpoint_filename="ltx-2-19b-dev.safetensors",
-        checkpoint_revision="47da56e2ad66ce4125a9922b4a8826bf407f9d0a",
-        checkpoint_env="VLLM_TEST_LTX2_OFFICIAL_CHECKPOINT",
+        checkpoint=LTX2_CHECKPOINT,
+        width=512,
+        height=384,
+        num_frames=25,
+        num_inference_steps=20,
+        seed=42,
        stg_block=29,
+        thresholds=STRICT_THRESHOLDS,
    ),
    LTXAccuracyCase(
        name="ltx2_3",
+        pipeline_kind="one_stage",
        model_id="diffusers/LTX-2.3-Diffusers",
        model_revision="8eee8edcf067e838b843f926ec4d4cc9b2be1aaf",
        model_env="VLLM_TEST_LTX23_MODEL",
        model_class_name="LTX2Pipeline",
-        checkpoint_repo="Lightricks/LTX-2.3",
-        checkpoint_filename="ltx-2.3-22b-dev.safetensors",
-        checkpoint_revision="4229404625088d21c4f112eb640fb04a0900ee25",
-        checkpoint_env="VLLM_TEST_LTX23_OFFICIAL_CHECKPOINT",
+        checkpoint=LTX23_CHECKPOINT,
+        width=512,
+        height=384,
+        num_frames=25,
+        num_inference_steps=20,
+        seed=42,
        stg_block=28,
+        thresholds=STRICT_THRESHOLDS,
    ),
    LTXAccuracyCase(
        name="ltx2_3_i2v",
+        pipeline_kind="one_stage",
        model_id="diffusers/LTX-2.3-Diffusers",
        model_revision="8eee8edcf067e838b843f926ec4d4cc9b2be1aaf",
        model_env="VLLM_TEST_LTX23_MODEL",
        model_class_name="LTX2Pipeline",
-        checkpoint_repo="Lightricks/LTX-2.3",
-        checkpoint_filename="ltx-2.3-22b-dev.safetensors",
-        checkpoint_revision="4229404625088d21c4f112eb640fb04a0900ee25",
-        checkpoint_env="VLLM_TEST_LTX23_OFFICIAL_CHECKPOINT",
+        checkpoint=LTX23_CHECKPOINT,
+        width=512,
+        height=384,
+        num_frames=25,
+        num_inference_steps=20,
+        seed=42,
        stg_block=28,
-        image_repo="huggingface/documentation-images",
-        image_filename="diffusers/svd/rocket.png",
-        image_revision="645d8364f0c7a101180b364811b5a11a362e4010",
+        thresholds=STRICT_THRESHOLDS,
+        image=I2V_IMAGE,
    ),
 )


+TWO_STAGE_CASES = (
+    LTXAccuracyCase(
+        name="ltx2_distilled_t2v",
+        pipeline_kind="distilled",
+        model_id="rootonchair/LTX-2-19b-distilled",
+        model_revision="388e2846f54aae51687498ffb6b27c7c2c9ce9e5",
+        model_env="VLLM_TEST_LTX2_DISTILLED_MODEL",
+        model_class_name="LTX2DistilledPipeline",
+        checkpoint=LTX2_DISTILLED_CHECKPOINT,
+        spatial_upsampler=LTX2_UPSAMPLER,
+        width=1536,
+        height=1024,
+        num_frames=121,
+        num_inference_steps=8,
+        seed=10,
+        stg_block=None,
+        gemma_model_id="Lightricks/LTX-2",
+        gemma_model_revision=LTX2_REVISION,
+        gemma_model_env="VLLM_TEST_LTX2_MODEL",
+        thresholds=STRICT_THRESHOLDS,
+    ),
+    LTXAccuracyCase(
+        name="ltx23_distilled_i2v",
+        pipeline_kind="distilled",
+        model_id="diffusers/LTX-2.3-Distilled-Diffusers",
+        model_revision="432e0d3c2d1769aaa4d295f9243f7062bf6b47ee",
+        model_env="VLLM_TEST_LTX23_DISTILLED_MODEL",
+        model_class_name="LTX2DistilledPipeline",
+        checkpoint=LTX23_DISTILLED_CHECKPOINT,
+        spatial_upsampler=LTX23_UPSAMPLER,
+        width=1536,
+        height=1024,
+        num_frames=121,
+        num_inference_steps=8,
+        seed=10,
+        stg_block=None,
+        gemma_model_id="diffusers/LTX-2.3-Diffusers",
+        gemma_model_revision="8eee8edcf067e838b843f926ec4d4cc9b2be1aaf",
+        gemma_model_env="VLLM_TEST_LTX23_MODEL",
+        thresholds=STRICT_THRESHOLDS,
+        image=I2V_IMAGE,
+    ),
+    LTXAccuracyCase(
+        name="ltx2_two_stage_layer_fused_t2v",
+        pipeline_kind="two_stage",
+        model_id="Lightricks/LTX-2",
+        model_revision=LTX2_REVISION,
+        model_env="VLLM_TEST_LTX2_MODEL",
+        model_class_name="LTX2TwoStagePipeline",
+        checkpoint=LTX2_CHECKPOINT,
+        spatial_upsampler=LTX2_UPSAMPLER,
+        distilled_lora=LTX2_DISTILLED_LORA,
+        width=1536,
+        height=1024,
+        num_frames=121,
+        num_inference_steps=40,
+        seed=10,
+        stg_block=29,
+        enable_layerwise_offload=True,
+        thresholds=STRICT_THRESHOLDS,
+    ),
+    LTXAccuracyCase(
+        name="ltx23_two_stage_layer_fused_i2v",
+        pipeline_kind="two_stage",
+        model_id="diffusers/LTX-2.3-Diffusers",
+        model_revision="8eee8edcf067e838b843f926ec4d4cc9b2be1aaf",
+        model_env="VLLM_TEST_LTX23_MODEL",
+        model_class_name="LTX2TwoStagePipeline",
+        checkpoint=LTX23_CHECKPOINT,
+        spatial_upsampler=LTX23_UPSAMPLER,
+        distilled_lora=LTX23_DISTILLED_LORA,
+        width=1536,
+        height=1024,
+        num_frames=121,
+        num_inference_steps=30,
+        seed=10,
+        stg_block=28,
+        enable_layerwise_offload=True,
+        thresholds=STRICT_THRESHOLDS,
+        image=I2V_IMAGE,
+    ),
+)
+
+CASES = (*LEGACY_CASES, *TWO_STAGE_CASES)
+
+
 def _run(command: list[str], *, env: dict[str, str], timeout: int = 1800) -> None:
    start = time.perf_counter()
    subprocess.run(command, env=env, timeout=timeout, check=True)
@@ -209,6 +395,7 @@ def _resolve_model(case: LTXAccuracyCase) -> Path:
                "model_index.json",
                "audio_vae/*",
                "connectors/*",
+                "latent_upsampler/*",
                "processor/*",
                "scheduler/*",
                "text_encoder/config.json",
@@ -223,74 +410,82 @@ def _resolve_model(case: LTXAccuracyCase) -> Path:
    )


-def _resolve_gemma_root(model: Path) -> Path:
+def _resolve_gemma_root(case: LTXAccuracyCase, model: Path) -> Path:
    configured_root = os.environ.get("VLLM_TEST_LTX_GEMMA_ROOT")
    if configured_root:
        root = Path(configured_root)
        assert root.is_dir(), f"Gemma root not found: {root}"
        return root
+    if case.gemma_model_id is not None:
+        configured_model = os.environ.get(case.gemma_model_env or "")
+        if configured_model and Path(configured_model).is_dir():
+            return Path(configured_model)
+        return Path(
+            snapshot_download(
+                repo_id=case.gemma_model_id,
+                revision=case.gemma_model_revision,
+                allow_patterns=[
+                    "processor/*",
+                    "text_encoder/*",
+                    "tokenizer/*",
+                ],
+            )
+        )
    return model


-def _resolve_checkpoint(case: LTXAccuracyCase, model: Path) -> Path:
-    configured_checkpoint = os.environ.get(case.checkpoint_env)
-    if configured_checkpoint:
-        checkpoint = Path(configured_checkpoint)
-        assert checkpoint.is_file(), f"Official checkpoint not found: {checkpoint}"
-        return checkpoint
-    model_checkpoint = model / case.checkpoint_filename
-    if model_checkpoint.is_file():
-        return model_checkpoint
+def _resolve_artifact(artifact: LTXArtifact, model: Path | None = None) -> Path:
+    configured_path = os.environ.get(artifact.env)
+    if configured_path:
+        path = Path(configured_path)
+        assert path.is_file(), f"Configured LTX artifact not found: {path}"
+        return path
+    if model is not None:
+        model_path = model / artifact.filename
+        if model_path.is_file():
+            return model_path
    return Path(
        hf_hub_download(
-            repo_id=case.checkpoint_repo,
-            filename=case.checkpoint_filename,
-            revision=case.checkpoint_revision,
+            repo_id=artifact.repo_id,
+            repo_type=None if artifact.repo_type == "model" else artifact.repo_type,
+            filename=artifact.filename,
+            revision=artifact.revision,
        )
    )


 def _resolve_image(case: LTXAccuracyCase) -> Path | None:
-    if case.image_filename is None:
+    if case.image is None:
        return None
-    if case.image_repo is None or case.image_revision is None:
-        raise ValueError(f"Incomplete image source for LTX accuracy case {case.name!r}.")
-    configured_image = os.environ.get("VLLM_TEST_LTX_I2V_IMAGE")
-    if configured_image:
-        image = Path(configured_image)
-        assert image.is_file(), f"LTX I2V conditioning image not found: {image}"
-        return image
-    return Path(
-        hf_hub_download(
-            repo_id=case.image_repo,
-            repo_type="dataset",
-            filename=case.image_filename,
-            revision=case.image_revision,
-        )
-    )
+    return _resolve_artifact(case.image)


 def _request(case: LTXAccuracyCase, image: Path | None) -> dict[str, object]:
    request: dict[str, object] = {
+        "pipeline_kind": case.pipeline_kind,
        "prompt": case.prompt,
-        "negative_prompt": NEGATIVE_PROMPT,
-        "width": 512,
-        "height": 384,
-        "num_frames": 25,
+        "negative_prompt": "" if case.pipeline_kind == "distilled" else NEGATIVE_PROMPT,
+        "width": case.width,
+        "height": case.height,
+        "num_frames": case.num_frames,
        "fps": 24,
-        "num_inference_steps": 20,
-        "seed": 42,
-        "video_cfg_scale": 3.0,
-        "audio_cfg_scale": 7.0,
-        "video_stg_scale": 1.0,
-        "audio_stg_scale": 1.0,
-        "video_modality_scale": 3.0,
-        "audio_modality_scale": 3.0,
-        "video_rescale_scale": 0.7,
-        "audio_rescale_scale": 0.7,
-        "video_stg_blocks": [case.stg_block],
-        "audio_stg_blocks": [case.stg_block],
+        "num_inference_steps": case.num_inference_steps,
+        "seed": case.seed,
    }
+    if case.pipeline_kind != "distilled":
+        assert case.stg_block is not None
+        request.update(
+            video_cfg_scale=3.0,
+            audio_cfg_scale=7.0,
+            video_stg_scale=1.0,
+            audio_stg_scale=1.0,
+            video_modality_scale=3.0,
+            audio_modality_scale=3.0,
+            video_rescale_scale=0.7,
+            audio_rescale_scale=0.7,
+            video_stg_blocks=[case.stg_block],
+            audio_stg_blocks=[case.stg_block],
+        )
    if image is not None:
        request["image"] = str(image.resolve())
    return request
@@ -299,20 +494,30 @@ def _request(case: LTXAccuracyCase, image: Path | None) -> dict[str, object]:
 def _video_metrics(reference: np.ndarray, prediction: np.ndarray) -> dict[str, float]:
    assert reference.shape == prediction.shape
    assert reference.ndim == 4 and reference.shape[-1] == 3
+    device = torch.device("cpu")
+    ssim = StructuralSimilarityIndexMeasure(data_range=1.0).to(device)
+    psnr = PeakSignalNoiseRatio(data_range=1.0).to(device)
    ssim_scores: list[float] = []
    psnr_scores: list[float] = []
-    for reference_frame, prediction_frame in zip(reference, prediction, strict=True):
-        reference_tensor = torch.from_numpy(reference_frame).permute(2, 0, 1).unsqueeze(0)
-        prediction_tensor = torch.from_numpy(prediction_frame).permute(2, 0, 1).unsqueeze(0)
-        ssim_scores.append(float(StructuralSimilarityIndexMeasure(data_range=1.0)(prediction_tensor, reference_tensor)))
-        psnr_scores.append(float(PeakSignalNoiseRatio(data_range=1.0)(prediction_tensor, reference_tensor)))
-    difference = np.abs(reference.astype(np.float64) - prediction.astype(np.float64))
+    max_abs = 0.0
+    absolute_error_sum = 0.0
+    with torch.inference_mode():
+        for reference_frame, prediction_frame in zip(reference, prediction, strict=True):
+            reference_tensor = torch.from_numpy(reference_frame).permute(2, 0, 1).unsqueeze(0).to(device)
+            prediction_tensor = torch.from_numpy(prediction_frame).permute(2, 0, 1).unsqueeze(0).to(device)
+            ssim_scores.append(float(ssim(prediction_tensor, reference_tensor)))
+            psnr_scores.append(float(psnr(prediction_tensor, reference_tensor)))
+            ssim.reset()
+            psnr.reset()
+            difference = np.abs(reference_frame.astype(np.float64) - prediction_frame.astype(np.float64))
+            max_abs = max(max_abs, float(difference.max()))
+            absolute_error_sum += float(difference.sum())
    return {
        "ssim_mean": float(np.mean(ssim_scores)),
        "ssim_min": float(np.min(ssim_scores)),
        "psnr_mean_db": float(np.mean(psnr_scores)),
-        "max_abs": float(difference.max()),
-        "mean_abs": float(difference.mean()),
+        "max_abs": max_abs,
+        "mean_abs": absolute_error_sum / reference.size,
    }


@@ -343,13 +548,19 @@ def _audio_metrics(reference: np.ndarray, prediction: np.ndarray) -> dict[str, f
@pytest.mark.benchmark
@pytest.mark.diffusion
@hardware_test(res={"cuda": "H100"}, num_cards=1)
-def test_ltx_one_stage_matches_official(case: LTXAccuracyCase, accuracy_artifact_root: Path) -> None:
-    """Compare official and Omni raw AV outputs from the same E2E request."""
+def test_ltx_matches_official(case: LTXAccuracyCase, accuracy_artifact_root: Path) -> None:
+    """Compare official and Omni raw AV outputs from the same single-H100 request."""
+    configured_artifact_root = os.environ.get("VLLM_TEST_LTX_ARTIFACT_ROOT")
+    if configured_artifact_root:
+        accuracy_artifact_root = Path(configured_artifact_root)
+        accuracy_artifact_root.mkdir(parents=True, exist_ok=True)
    output_root = reset_artifact_dir(accuracy_artifact_root / "ltx_official" / case.name)
    official_root, official_revision = _official_source(accuracy_artifact_root / "ltx_official")
    model = _resolve_model(case)
-    gemma_root = _resolve_gemma_root(model)
-    checkpoint = _resolve_checkpoint(case, model)
+    gemma_root = _resolve_gemma_root(case, model)
+    checkpoint = _resolve_artifact(case.checkpoint, model)
+    spatial_upsampler = _resolve_artifact(case.spatial_upsampler, model) if case.spatial_upsampler is not None else None
+    distilled_lora = _resolve_artifact(case.distilled_lora, model) if case.distilled_lora is not None else None
    image = _resolve_image(case)
    request_path = output_root / "request.json"
    request_path.write_text(json.dumps(_request(case, image), indent=2) + "\n")
@@ -359,8 +570,13 @@ def test_ltx_one_stage_matches_official(case: LTXAccuracyCase, accuracy_artifact
        str(runner),
        "--request",
        str(request_path),
+        "--pipeline-kind",
+        case.pipeline_kind,
    ]
-    if os.environ.get("VLLM_TEST_LTX_ENABLE_LAYERWISE_OFFLOAD", "").lower() in {"1", "true", "yes", "on"}:
+    enable_layerwise_offload = case.enable_layerwise_offload or os.environ.get(
+        "VLLM_TEST_LTX_ENABLE_LAYERWISE_OFFLOAD", ""
+    ).lower() in {"1", "true", "yes", "on"}
+    if enable_layerwise_offload:
        runner_args.append("--enable-layerwise-offload")
    env = os.environ.copy()
    env["VLLM_TEST_LTX_OFFICIAL_REVISION"] = official_revision
@@ -372,10 +588,9 @@ def test_ltx_one_stage_matches_official(case: LTXAccuracyCase, accuracy_artifact
    )

    official_output = output_root / "official"
-    _run(
-        _official_runner_prefix()
-        + runner_args
-        + [
+    official_args = _official_runner_prefix() + runner_args
+    official_args.extend(
+        [
            "--backend",
            "official",
            "--output-dir",
@@ -386,12 +601,16 @@ def test_ltx_one_stage_matches_official(case: LTXAccuracyCase, accuracy_artifact
            str(checkpoint),
            "--gemma-root",
            str(gemma_root),
-        ],
-        env=env,
+        ]
    )
+    if spatial_upsampler is not None:
+        official_args.extend(["--spatial-upsampler", str(spatial_upsampler)])
+    if distilled_lora is not None:
+        official_args.extend(["--distilled-lora", str(distilled_lora)])
+    _run(official_args, env=env)

    omni_output = output_root / "omni"
-    _run(
+    omni_args = (
        [sys.executable]
        + runner_args
        + [
@@ -403,9 +622,9 @@ def test_ltx_one_stage_matches_official(case: LTXAccuracyCase, accuracy_artifact
            str(model),
            "--model-class-name",
            case.model_class_name,
-        ],
-        env=env,
+        ]
    )
+    _run(omni_args, env=env)

    official_metadata = json.loads((official_output / "metadata.json").read_text())
    omni_metadata = json.loads((omni_output / "metadata.json").read_text())
@@ -423,18 +642,27 @@ def test_ltx_one_stage_matches_official(case: LTXAccuracyCase, accuracy_artifact
    result = {
        "case": case.name,
        "task": "i2v" if image is not None else "t2v",
+        "pipeline_kind": case.pipeline_kind,
+        "model_class_name": case.model_class_name,
        "attention_backend": ATTENTION_BACKEND,
        "official_revision": official_revision,
        "model_revision": case.model_revision,
-        "checkpoint_revision": case.checkpoint_revision,
+        "checkpoint_revision": case.checkpoint.revision,
+        "spatial_upsampler_revision": (case.spatial_upsampler.revision if case.spatial_upsampler is not None else None),
+        "distilled_lora_revision": case.distilled_lora.revision if case.distilled_lora is not None else None,
+        "enable_layerwise_offload": enable_layerwise_offload,
+        "thresholds": asdict(case.thresholds),
+        "request": _request(case, image),
        "video": video_metrics,
        "audio": audio_metrics,
    }
    (output_root / "metrics.json").write_text(json.dumps(result, indent=2) + "\n")
    print(json.dumps(result, indent=2))

-    assert video_metrics["ssim_mean"] >= VIDEO_SSIM_MEAN_THRESHOLD
-    assert video_metrics["ssim_min"] >= VIDEO_SSIM_MIN_THRESHOLD
-    assert video_metrics["psnr_mean_db"] >= VIDEO_PSNR_MEAN_THRESHOLD
-    assert audio_metrics["relative_l2"] <= AUDIO_RELATIVE_L2_THRESHOLD
-    assert audio_metrics["cosine_similarity"] >= AUDIO_COSINE_THRESHOLD
+    assert video_metrics["ssim_mean"] >= case.thresholds.video_ssim_mean
+    assert video_metrics["ssim_min"] >= case.thresholds.video_ssim_min
+    assert video_metrics["psnr_mean_db"] >= case.thresholds.video_psnr_mean_db
+    if case.thresholds.audio_relative_l2 > 0:
+        assert audio_metrics["relative_l2"] <= case.thresholds.audio_relative_l2
+    if case.thresholds.audio_cosine_similarity > 0:
+        assert audio_metrics["cosine_similarity"] >= case.thresholds.audio_cosine_similarity
@@ -0,0 +1,109 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Online-serving smoke tests for the public LTX-2 pipeline entries."""
+
+import os
+
+import pytest
+
+from tests.helpers.mark import hardware_marks
+from tests.helpers.media import generate_synthetic_image
+from tests.helpers.runtime import OmniServer, OmniServerParams, OpenAIClientHandler
+
+pytestmark = [pytest.mark.diffusion, pytest.mark.full_model]
+
+DISTILLED_MODEL = os.getenv("VLLM_TEST_LTX23_DISTILLED_MODEL", "diffusers/LTX-2.3-Distilled-Diffusers")
+BASE_MODEL = os.getenv("VLLM_TEST_LTX23_MODEL", "diffusers/LTX-2.3-Diffusers")
+PROMPT = "A serene lake at sunset with mountains in the background."
+SINGLE_CARD_MARKS = hardware_marks(res={"cuda": "H100"})
+HSDP_MARKS = hardware_marks(res={"cuda": "H100"}, num_cards=2)
+
+
+def _server(
+    model: str,
+    model_class_name: str,
+    *extra_args: str,
+) -> OmniServerParams:
+    return OmniServerParams(
+        model=model,
+        server_args=[
+            "--enforce-eager",
+            *extra_args,
+            "--model-class-name",
+            model_class_name,
+        ],
+    )
+
+
+def _cases():
+    return [
+        pytest.param(
+            _server(BASE_MODEL, "LTX2Pipeline"),
+            3,
+            id="one_stage",
+            marks=SINGLE_CARD_MARKS,
+        ),
+        pytest.param(
+            _server(DISTILLED_MODEL, "LTX2DistilledPipeline"),
+            8,
+            id="distilled",
+            marks=SINGLE_CARD_MARKS,
+        ),
+        pytest.param(
+            _server(
+                DISTILLED_MODEL,
+                "LTX2DistilledPipeline",
+                "--use-hsdp",
+                "--hsdp-shard-size",
+                "2",
+            ),
+            8,
+            id="distilled_hsdp2",
+            marks=HSDP_MARKS,
+        ),
+        pytest.param(
+            _server(
+                BASE_MODEL,
+                "LTX2TwoStagePipeline",
+                "--enable-layerwise-offload",
+            ),
+            3,
+            id="two_stage",
+            marks=SINGLE_CARD_MARKS,
+        ),
+    ]
+
+
+@pytest.mark.parametrize(
+    ("omni_server", "num_inference_steps"),
+    _cases(),
+    indirect=["omni_server"],
+)
+def test_ltx_pipeline_entries(
+    omni_server: OmniServer,
+    num_inference_steps: int,
+    openai_client: OpenAIClientHandler,
+    subtests,
+):
+    for task in ("t2v", "i2v"):
+        with subtests.test(task=task):
+            form_data = {
+                "prompt": PROMPT,
+                "model": omni_server.model,
+                "height": 128,
+                "width": 128,
+                "num_frames": 9,
+                "fps": 8,
+                "num_inference_steps": num_inference_steps,
+                "seed": 42,
+            }
+            request_config = {
+                "model": omni_server.model,
+                "form_data": form_data,
+            }
+            if task == "i2v":
+                request_config["image_reference"] = (
+                    f"data:image/jpeg;base64,{generate_synthetic_image(512, 512)['base64']}"
+                )
+
+            openai_client.send_video_diffusion_request(request_config)
@@ -1,91 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Expansion tests for LTX-2 two-stage pipelines.
-
-Coverage:
- HSDP with LTX2DistilledPipeline (text-to-video and image-to-video)
-"""
-
-import os
-
-import pytest
-
-from tests.helpers.mark import hardware_marks
-from tests.helpers.media import generate_synthetic_image
-from tests.helpers.runtime import OmniServer, OmniServerParams, OpenAIClientHandler
-
-pytestmark = [pytest.mark.diffusion, pytest.mark.full_model]
-
-MODEL = os.getenv("VLLM_TEST_LTX2_MODEL", "rootonchair/LTX-2-19b-distilled")
-PROMPT = "A serene lake at sunset with mountains in the background."
-PARALLEL_MARKS = hardware_marks(res={"cuda": "H100"}, num_cards=2)
-HSDP_ARGS = ["--use-hsdp", "--hsdp-shard-size", "2"]
-
-
-def _cases():
-    cases = []
-
-    cases.append(
-        pytest.param(
-            OmniServerParams(
-                model=MODEL,
-                server_args=[
-                    *HSDP_ARGS,
-                    "--model-class-name",
-                    "LTX2DistilledPipeline",
-                ],
-            ),
-            False,
-            id="t2v_hsdp",
-            marks=PARALLEL_MARKS,
-        )
-    )
-
-    cases.append(
-        pytest.param(
-            OmniServerParams(
-                model=MODEL,
-                server_args=[
-                    *HSDP_ARGS,
-                    "--model-class-name",
-                    "LTX2DistilledPipeline",
-                ],
-            ),
-            True,
-            id="i2v_hsdp",
-            marks=PARALLEL_MARKS,
-        )
-    )
-
-    return cases
-
-
-@pytest.mark.parametrize(("omni_server", "is_i2v"), _cases(), indirect=["omni_server"])
-def test_ltx2_two_stage_hsdp(
-    omni_server: OmniServer,
-    is_i2v: bool,
-    openai_client: OpenAIClientHandler,
-):
-    # Keep CI small while exercising the fixed 8+3-step distilled recipe.
-    # Height and width describe the final output; Stage 1 runs at half size.
-    form_data = {
-        "prompt": PROMPT,
-        "model": omni_server.model,
-        "height": 128,
-        "width": 128,
-        "num_frames": 9,
-        "fps": 8,
-        "num_inference_steps": 8,
-        "seed": 42,
-    }
-
-    request_config = {
-        "model": omni_server.model,
-        "form_data": form_data,
-    }
-
-    if is_i2v:
-        request_config["image_reference"] = f"data:image/jpeg;base64,{generate_synthetic_image(512, 512)['base64']}"
-
-    openai_client.send_video_diffusion_request(request_config)
@@ -3,6 +3,7 @@

 from __future__ import annotations

+import os
 import subprocess
 import sys
 from pathlib import Path
@@ -14,9 +15,18 @@ REPO_ROOT = Path(__file__).resolve().parents[2]


 def _assert_isolated_import_succeeds(script: str) -> None:
+    # These tests validate Python import/package boundaries and do not use a
+    # GPU. Other tests in the combined entrypoints/engine suite may temporarily
+    # configure only one of the CUDA/ROCm visibility variables. Do not let an
+    # unrelated, inconsistent device-selection environment prevent the child
+    # interpreter from reaching the boundary assertion.
+    env = os.environ.copy()
+    env.pop("CUDA_VISIBLE_DEVICES", None)
+    env.pop("HIP_VISIBLE_DEVICES", None)
    result = subprocess.run(
        [sys.executable, "-c", script],
        cwd=REPO_ROOT,
+        env=env,
        capture_output=True,
        text=True,
        timeout=60,
@@ -25,6 +35,14 @@ def _assert_isolated_import_succeeds(script: str) -> None:
    assert result.returncode == 0, result.stderr


+def test_isolated_import_ignores_conflicting_gpu_visibility(monkeypatch: pytest.MonkeyPatch) -> None:
+    """Import-boundary checks must not depend on GPU state leaked by prior tests."""
+    monkeypatch.setenv("HIP_VISIBLE_DEVICES", "0")
+    monkeypatch.setenv("CUDA_VISIBLE_DEVICES", "0,1")
+
+    _assert_isolated_import_succeeds("import vllm_omni.outputs")
+
+
 def test_stable_engine_imports_do_not_load_experimental_duplex() -> None:
    _assert_isolated_import_succeeds("""
 import sys
@@ -139,6 +139,7 @@ def _engine_arg_inputs(tmp_path: Path) -> tuple[PipelineConfig, DeployConfig, st
                engine_extras={
                    "model": str(stage_model),
                    "attention_backend": "FLASHINFER",
+                    "attention_config": {"use_trtllm_attention": False},
                    "hf_overrides": {"rope_scaling": {"factor": 2}},
                    "limit_mm_per_prompt": {"audio": 1},
                    "logits_processors": ["test.custom.LogitsProcessor"],
@@ -229,6 +230,7 @@ def test_typed_llm_engine_args_preserve_legacy_adapter_behavior(tmp_path):
    assert thinker_args["omni_kv_config"] == {"need_send_cache": True}
    assert thinker_args["worker_cls"] == "test.custom.Worker"
    assert thinker_args["attention_backend"] == "FLASHINFER"
+    assert thinker_args["attention_config"] == {"use_trtllm_attention": False}
    assert thinker_args["trust_remote_code"] is True
    assert thinker_args["dtype"] == "float16"
    assert thinker_args["distributed_executor_backend"] == "mp"
@@ -5,6 +5,7 @@
 import asyncio
 import base64
 import os
+import tempfile

 import httpx
 import pytest
@@ -101,7 +102,7 @@ class TestDecodeAudioUrlSuffixSanitization:
        try:
            # ".." contains non-alnum dots, so falls back to .wav
            assert path.endswith(".wav")
-            assert os.path.dirname(path) == "/tmp"
+            assert os.path.dirname(path) == tempfile.gettempdir()
        finally:
            _cleanup(path)

@@ -76,18 +76,6 @@ def test_serve_parser_accepts_four_way_cfg_parallelism() -> None:
    assert args.cfg_parallel_size == 4


-def test_serve_parser_defaults_vae_parallel_mode_to_tile_and_accepts_auto() -> None:
-    parser = TrackingArgumentParser()
-    subparsers = parser.add_subparsers(dest="subcommand")
-    OmniServeCommand().subparser_init(subparsers)
-
-    default_args = parser.parse_args(["serve", "fake-model", "--omni"])
-    auto_args = parser.parse_args(["serve", "fake-model", "--omni", "--vae-parallel-mode", "auto"])
-
-    assert default_args.vae_parallel_mode == "tile"
-    assert auto_args.vae_parallel_mode == "auto"
-
-
 def _make_headless_args(**kwargs) -> TrackingNamespace:
    defaults = {
        "model": "fake-model",
@@ -0,0 +1,2 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
@@ -0,0 +1,137 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Regression tests for the dots.tts talker's stop-signal pairing and
+per-request state lifecycle.
+
+Pure-CPU tests on a bare talker (``__new__`` + the handful of attributes
+the methods under test read) — no weights, no GPU, no engine.  Mirrors
+tests/model_executor/models/voxcpm2/test_talker_state_eviction.py.
+"""
+
+from __future__ import annotations
+
+import functools
+from types import SimpleNamespace
+
+import pytest
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+
+
+torch = pytest.importorskip("torch")
+
+_VOCAB_SIZE = 151672
+
+
+@functools.lru_cache(maxsize=1)
+def _dots_tts_talker_mod():
+    """Defer talker import (pulls vLLM model_executor) until first use."""
+    from vllm_omni.model_executor.models.dots_tts.dots_tts_talker import (
+        DotsTTSForConditionalGeneration,
+        _RequestState,
+    )
+
+    return DotsTTSForConditionalGeneration, _RequestState
+
+
+def _make_bare_talker():
+    DotsTTSForConditionalGeneration, _ = _dots_tts_talker_mod()
+    talker = DotsTTSForConditionalGeneration.__new__(DotsTTSForConditionalGeneration)
+    talker.config = SimpleNamespace(vocab_size=_VOCAB_SIZE)
+    talker._active_states = {}
+    talker._pending_requests = []
+    talker._results_queue = []
+    talker._audio_queue = []
+    talker._deferred_cleanup_ids = set()
+    talker._beta_trace = False
+    return talker
+
+
+def _add_state(talker, req_id: str, *, is_stopping: bool = False):
+    _, _RequestState = _dots_tts_talker_mod()
+    state = _RequestState(request_id=req_id)
+    state.prefill_completed = True
+    state.is_stopping = is_stopping
+    talker._active_states[req_id] = state
+    return state
+
+
+def _hidden(bsz: int) -> torch.Tensor:
+    return torch.zeros(bsz, 1536, dtype=torch.float32)
+
+
+class TestComputeLogitsPairing:
+    """compute_logits pairs _results_queue entries with batch rows
+    positionally, so every request must push exactly one entry per step —
+    including prefills, which push a (req_id, None) placeholder."""
+
+    def test_mixed_batch_prefill_placeholder_keeps_rows_aligned(self) -> None:
+        talker = _make_bare_talker()
+        _add_state(talker, "prefill-req")
+        decode_state = _add_state(talker, "decode-req", is_stopping=True)
+        decode_state.precomputed_stop_logits = torch.tensor([[0.1, 0.9]])
+
+        # Row 0 is the prefill (None placeholder), row 1 the stopping decode.
+        talker._results_queue = [
+            ("prefill-req", None),
+            ("decode-req", torch.tensor([[0.1, 0.9]])),
+        ]
+        logits = talker.compute_logits(_hidden(2))
+
+        assert logits.shape == (2, _VOCAB_SIZE)
+        # Prefill row: forced continue.
+        assert logits[0, 0].item() == 1.0
+        assert logits[0, 1].item() == float("-inf")
+        # Stopping decode row: stop slot wins.
+        assert logits[1, 1].item() == 1.0
+        assert logits[1, 0].item() == 0.0
+        # Everything else stays masked.
+        assert logits[0, 2].item() == float("-inf")
+        assert logits[1, 2].item() == float("-inf")
+        # Queue drained.
+        assert talker._results_queue == []
+
+    def test_below_threshold_stop_logits_force_continue(self) -> None:
+        """Raw (continue, stop) softmax must not reach the sampler: a
+        prob_stop in (0.5, 0.8] would flip greedy argmax even though
+        is_stopping stayed False (the 0.8-threshold contract)."""
+        talker = _make_bare_talker()
+        _add_state(talker, "req", is_stopping=False)
+        talker._results_queue = [("req", torch.tensor([[0.3, 0.7]]))]
+
+        logits = talker.compute_logits(_hidden(1))
+
+        assert logits[0, 0].item() == 1.0
+        assert logits[0, 1].item() == float("-inf")
+
+    def test_empty_queue_defaults_to_all_continue(self) -> None:
+        talker = _make_bare_talker()
+        logits = talker.compute_logits(_hidden(3))
+
+        assert torch.all(logits[:, 0] == 1.0)
+        assert torch.all(logits[:, 1] == float("-inf"))
+
+
+class TestStateEviction:
+    """on_requests_finished fires BEFORE forward(), so eviction must be
+    deferred to _flush_deferred_cleanup at the end of the step — the
+    finishing request's audio still drains through the current forward."""
+
+    def test_finished_request_survives_until_flush(self) -> None:
+        talker = _make_bare_talker()
+        _add_state(talker, "req-a")
+        _add_state(talker, "req-b")
+
+        talker.on_requests_finished(["req-a"])
+        assert "req-a" in talker._active_states  # not evicted yet
+
+        talker._flush_deferred_cleanup()
+        assert "req-a" not in talker._active_states
+        assert "req-b" in talker._active_states
+        assert talker._deferred_cleanup_ids == set()
+
+    def test_unknown_request_id_is_ignored(self) -> None:
+        talker = _make_bare_talker()
+        talker.on_requests_finished(["never-seen"])
+        talker._flush_deferred_cleanup()
+        assert talker._active_states == {}
@@ -1326,7 +1326,9 @@ class TestRegistry:

        assert "higgs_multimodal_qwen3" in OMNI_PIPELINES

-    def test_deploy_yaml_exists(self):
+    def test_deploy_yaml_uses_native_flashinfer(self):
+        from vllm_omni.config.stage_config import load_deploy_config
+
        for name in (
            "higgs_multimodal_qwen3.yaml",
            "higgs_multimodal_qwen3_high_throughput.yaml",
@@ -1334,6 +1336,10 @@ class TestRegistry:
        ):
            yaml_path = get_deploy_config_path(name)
            assert os.path.isfile(yaml_path), f"Deploy YAML not found at {yaml_path}"
+            deploy_config = load_deploy_config(yaml_path)
+            stage_0 = next(stage for stage in deploy_config.stages if stage.stage_id == 0)
+            assert stage_0.engine_extras["attention_backend"] == "FLASHINFER"
+            assert stage_0.engine_extras["attention_config"] == {"use_trtllm_attention": False}


 # ---- AC-3: Prompt Builder ----
@@ -233,10 +233,13 @@ def test_ltx_extra_registry_declares_official_guidance_params() -> None:
        }
    )

-    assert get_extra_body_params("LTX2Pipeline") == expected
-    assert get_extra_output_params("LTX2Pipeline") == frozenset()
-    assert get_extra_body_params("LTX2DistilledPipeline") == expected
-    assert get_extra_output_params("LTX2DistilledPipeline") == frozenset()
+    for pipeline_name in (
+        "LTX2Pipeline",
+        "LTX2TwoStagePipeline",
+        "LTX2DistilledPipeline",
+    ):
+        assert get_extra_body_params(pipeline_name) == expected
+        assert get_extra_output_params(pipeline_name) == frozenset()


@pytest.mark.core_model
@@ -27,6 +27,7 @@ EXCLUDED_MODELS = [
    "OvisImagePipeline",
    "WanPipeline",
    "WanVACEPipeline",
+    "LTX2TwoStagePipeline",
    "LTX2DistilledPipeline",
    "LTX2T2VDMD2Pipeline",
    "LTX2I2VDMD2Pipeline",
@@ -90,6 +90,7 @@ class _ModelEngineOverrides(TypedDict, total=False):
    trust_remote_code: bool
    dtype: Any
    attention_backend: Any
+    attention_config: Any
    moe_backend: str
    hf_overrides: Any
    limit_mm_per_prompt: dict[str, Any]
@@ -321,6 +322,7 @@ class OmniStageModelConfig:
    trust_remote_code: bool = False
    dtype: Any = "auto"
    attention_backend: Any = None
+    attention_config: Any = None
    moe_backend: str = "auto"
    hf_overrides: Any = None
    limit_mm_per_prompt: dict[str, Any] | None = None
@@ -473,10 +475,9 @@ class OmniStageDiffusionParallelConfig(OmniStageParallelConfig):
            raise ValueError("allgather_degree > 1 is mutually exclusive with ulysses_degree/ring_degree > 1")
        if self.ulysses_mode not in {"strict", "advanced_uaa"}:
            raise ValueError("ulysses_mode must be 'strict' or 'advanced_uaa'")
-        if self.vae_parallel_mode not in {"tile", "auto", "spatial_shard_height", "spatial_shard_width"}:
+        if self.vae_parallel_mode not in {"tile", "spatial_shard_height", "spatial_shard_width"}:
            raise ValueError(
-                "vae_parallel_mode must be one of "
-                "{'tile', 'auto', 'spatial_shard_height', 'spatial_shard_width'}, "
+                "vae_parallel_mode must be one of {'tile', 'spatial_shard_height', 'spatial_shard_width'}, "
                f"but got {self.vae_parallel_mode!r}."
            )

@@ -47,6 +47,7 @@ from vllm_omni.model_executor.models.bagel.pipeline import (
 )
 from vllm_omni.model_executor.models.cosyvoice3.pipeline import COSYVOICE3_PIPELINE
 from vllm_omni.model_executor.models.covo_audio.pipeline import COVO_AUDIO_PIPELINE
+from vllm_omni.model_executor.models.dots_tts.pipeline import DOTS_TTS_PIPELINE
 from vllm_omni.model_executor.models.dreamzero.pipeline import DREAMZERO_PIPELINE
 from vllm_omni.model_executor.models.dynin_omni.pipeline import DYNIN_OMNI_PIPELINE
 from vllm_omni.model_executor.models.fish_speech.pipeline import FISH_SPEECH_PIPELINE
@@ -144,6 +145,7 @@ OMNI_PIPELINES: dict[str, PipelineConfig | PipelineResolverFunc] = {
    "hunyuan_video_15": HUNYUAN_VIDEO_15_PIPELINE,
    "wan2_2_ti2v": WAN2_2_TI2V_PIPELINE,
    "voxcpm2": VOXCPM2_PIPELINE,
+    "dots_tts": DOTS_TTS_PIPELINE,
    "cosyvoice3": COSYVOICE3_PIPELINE,
    "audex_tts": AUDEX_TTS_PIPELINE,
    "audex_tta": AUDEX_TTA_PIPELINE,
@@ -44,6 +44,65 @@ vllm-omni serve bosonai/higgs-audio-v3-tts-4b \
    --deploy-config vllm_omni/deploy/higgs_multimodal_qwen3_low_latency.yaml
 ```

+## Why Stage 0 Pins Native FlashInfer
+
+The three Stage 0 profiles set `attention_backend: FLASHINFER` and also set:
+
+```yaml
+attention_config:
+  use_trtllm_attention: false
+```
+
+`FLASHINFER` names the overall vLLM attention backend. Starting with vLLM
+0.25, that backend can internally route Hopper decode to the TRT-LLM XQA
+kernel. The XQA eligibility check answers whether the kernel can run, not
+whether it is the fastest kernel for the current workload. In particular, the
+decode builder can select XQA statically when all of the following hold:
+
+- the explicit `use_trtllm_attention` override is not `false`;
+- the device is Hopper/SM90 and the XQA cubin is available; and
+- the query-head count is divisible by the KV-head count.
+
+Higgs has 32 query heads and 8 KV heads, so it passes the head-layout check.
+The static gate does not include a model dtype, KV-cache dtype, context-length,
+batch-shape, or measured-performance comparison. This matters because kernel
+support is not a performance guarantee.
+
+Higgs also has an unusual output shape that must not be confused with its
+attention shape. One autoregressive audio step advances one token per request.
+The resulting hidden row is then projected into eight codebook distributions
+with shape `[1, 8, codebook_size]`. The eight codebooks therefore do not turn
+the attention query into `q_len=8`; at concurrency one, XQA still receives a
+single BF16 query row. At this small per-step shape, any XQA advantage for
+other KV dtypes, batch sizes, or context lengths does not necessarily apply.
+The difference is paid in every transformer layer and every generated audio
+step, so even a small per-call disadvantage accumulates in TTFP and RTF.
+
+The upstream XQA integration also showed that the result is dtype-sensitive:
+its FP8 benchmark favored XQA, while its BF16 benchmark reported 21.16 ms mean
+TPOT for XQA versus 19.47 ms for FA3. That is not a direct native-FlashInfer
+comparison, but it demonstrates why SM90 compatibility alone is insufficient
+to choose the fastest BF16 kernel. See the
+[vLLM XQA benchmark](https://github.com/vllm-project/vllm/pull/43232#issuecomment-4501176335).
+
+The end-to-end H800 A/B in
+[issue #5584](https://github.com/vllm-project/vllm-omni/issues/5584#issuecomment-5188700218)
+isolated the routing decision on the same vLLM 0.25 stack:
+
+| Decode path | TTFP (ms) | RTF |
+|---|---:|---:|
+| XQA (automatic default) | 175.76 | 0.3405 |
+| Native FlashInfer (`use_trtllm_attention: false`) | 166.39 | 0.3233 |
+| vLLM 0.24 reference | 166.68 | 0.3209 |
+
+This proves that XQA selection causes the observed local regression and that
+native FlashInfer recovers the pre-v0.25 range. It does not by itself identify
+one defective XQA instruction or memory transaction. Establishing that deeper
+kernel-level cause requires an attention-only microbenchmark and GPU profiler
+trace for the exact Higgs batch/context distribution. Until vLLM has a
+workload-aware selector or XQA is faster for this BF16 shape, the model-level
+pin is the deterministic choice.
+
 ## Notes

 - Stage 1 remains `enforce_eager: true` in both profiles.
@@ -0,0 +1,28 @@
+# dots.tts deployment config (single-stage AR talker).
+#
+# Conservative first-integration settings: eager execution, no CUDA
+# Graphs, no batched side-path optimizations yet — see voxcpm2.yaml for
+# the runtime knobs a mature deployment of this model family grows.
+#
+# Verified locally on RTX 5080 16GB.
+
+async_chunk: false
+dtype: bfloat16
+
+stages:
+  - stage_id: 0
+    max_num_seqs: 4
+    gpu_memory_utilization: 0.8
+    enforce_eager: true
+    trust_remote_code: true
+    enable_prefix_caching: false
+    async_scheduling: true
+    max_num_batched_tokens: 4096
+    max_model_len: 4096
+    devices: "0"
+    default_sampling_params:
+      temperature: 0.0
+      top_p: 1.0
+      top_k: -1
+      max_tokens: 4096
+      repetition_penalty: 1.0
@@ -38,6 +38,11 @@ stages:
    gpu_memory_utilization: 0.6
    enforce_eager: true
    attention_backend: FLASHINFER
+    # Eight codebooks share one attention row, so decode remains BF16 q_len=1
+    # per request. vLLM's SM90 XQA capability gate does not benchmark this
+    # shape; keep the faster native path (README_higgs_audio_v3.md, #5584).
+    attention_config:
+      use_trtllm_attention: false
    trust_remote_code: true
    # Stage-0 prefix caching is on by default. The talker's HS opt-out and
    # deferred ``codes.audio`` mm-output write (see higgs_audio_v3_talker.py)
@@ -37,6 +37,11 @@ stages:
    gpu_memory_utilization: 0.6
    enforce_eager: true
    attention_backend: FLASHINFER
+    # Eight codebooks share one attention row, so decode remains BF16 q_len=1
+    # per request. vLLM's SM90 XQA capability gate does not benchmark this
+    # shape; keep the faster native path (README_higgs_audio_v3.md, #5584).
+    attention_config:
+      use_trtllm_attention: false
    trust_remote_code: true
    # Stage-0 prefix caching is on by default. The talker's HS opt-out and
    # deferred ``codes.audio`` mm-output write (see higgs_audio_v3_talker.py)
@@ -43,6 +43,11 @@ stages:
    gpu_memory_utilization: 0.6
    enforce_eager: false
    attention_backend: FLASHINFER
+    # Eight codebooks share one attention row, so decode remains BF16 q_len=1
+    # per request. vLLM's SM90 XQA capability gate does not benchmark this
+    # shape; keep the faster native path (README_higgs_audio_v3.md, #5584).
+    attention_config:
+      use_trtllm_attention: false
    trust_remote_code: true
    enable_prefix_caching: true
    async_scheduling: false
@@ -1263,6 +1263,7 @@ def extract_minimax_h3_context(
        _BF16_DTYPE,
        _FORWARD_SUPPORTED_KWARGS,
        MINIMAX_H3_ADALN_MODALITY_NUM,
+        _build_rope_table,
        _modulate_scale_shift,
        _required_kwarg,
    )
@@ -1318,7 +1319,7 @@ def extract_minimax_h3_context(
        raise ValueError(f"inverse_indices must be [{seq_len}], got {list(inverse_indices.shape)}")
    device = x.device

-    rope_freqs = module.rope(img_position_ids).to(device)
+    rope_table = _build_rope_table(module.rope(img_position_ids).to(device))

    decoder_input, t_emb = module._embed(
        x=x,
@@ -1351,7 +1352,7 @@ def extract_minimax_h3_context(
    def run_transformer_blocks() -> tuple[torch.Tensor, ...]:
        hidden, block_rope, block_combined = module.sp_prepare(
            decoder_input,
-            rope_freqs,
+            rope_table,
            combined_indices,
        )
        for block in module.blocks:
@@ -1359,7 +1360,7 @@ def extract_minimax_h3_context(
                hidden,
                t_emb=t_emb,
                combined_indices=block_combined,
-                rope_freqs=block_rope,
+                rope_table=block_rope,
                cu_seqlens=cu_seqlens,
                max_seqlen=max_seqlen,
                packed_total=seq_len,
@@ -212,9 +212,6 @@ class DiffusionParallelConfig:

    - "tile": Patch/tile parallel decode (default). Each rank decodes a subset
      of spatial tiles and the results are stitched on rank 0.
-    - "auto": Select the longer spatial-shard axis for validated Wan tiled
-      decodes on a full DiT-group topology; otherwise use patch/tile parallel
-      decode.
    - "spatial_shard_height": Spatially-sharded decode that splits decoder
      feature maps along height and exchanges halo rows around spatial
      convolutions.
@@ -254,9 +251,8 @@ class DiffusionParallelConfig:
        assert self.allgather_degree > 0, "AllGather degree must be > 0"
        assert self.cfg_parallel_size > 0, "CFG parallel size must be > 0"
        assert self.vae_patch_parallel_size > 0, "VAE patch parallel size must be > 0"
-        assert self.vae_parallel_mode in {"tile", "auto", "spatial_shard_height", "spatial_shard_width"}, (
-            "vae_parallel_mode must be one of "
-            "{'tile', 'auto', 'spatial_shard_height', 'spatial_shard_width'}, "
+        assert self.vae_parallel_mode in {"tile", "spatial_shard_height", "spatial_shard_width"}, (
+            "vae_parallel_mode must be one of {'tile', 'spatial_shard_height', 'spatial_shard_width'}, "
            f"but got {self.vae_parallel_mode!r}."
        )
        if self.allgather_degree > 1:
@@ -23,30 +23,6 @@ from vllm_omni.platforms import current_omni_platform
 logger = init_logger(__name__)


-def select_auto_spatial_shard_split_dim(
-    z_shape: tuple[int, ...],
-    world_size: int,
-) -> str | None:
-    """Select a profitable Wan spatial-shard axis for one latent shape.
-
-    This selector runs only after Diffusers has already chosen ``tiled_decode``
-    for a spatially large request. Splitting the longer eligible axis reduces
-    halo traffic for rectangular requests and improves per-rank balance.
-    """
-    if len(z_shape) != 5 or world_size <= 1:
-        return None
-
-    _, _, _, height, width = z_shape
-    candidates = [
-        (extent, split_dim) for extent, split_dim in ((height, "height"), (width, "width")) if extent >= world_size
-    ]
-    if not candidates:
-        return None
-    # Prefer width for square latents: it was modestly faster in the existing
-    # four-GPU Wan benchmark and keeps the choice deterministic.
-    return max(candidates, key=lambda candidate: (candidate[0], candidate[1] == "width"))[1]
-
-
 class OmniAutoencoderKLWan(AutoencoderKLWan):
    def _execution_context(self):
        try:
@@ -80,19 +56,6 @@ class DistributedAutoencoderKLWan(OmniAutoencoderKLWan, DistributedVaeMixin):
        model.init_distributed()
        return model

-    def set_parallel_size(self, parallel_size: int, mode: str = "tile") -> None:
-        super().set_parallel_size(parallel_size, mode=mode)
-        if mode not in {"auto", "spatial_shard_height", "spatial_shard_width"}:
-            return
-
-        executor = self.distributed_executor
-        world_size = dist.get_world_size(group=executor.group)
-        if parallel_size <= 1 or parallel_size != world_size:
-            return
-
-        split_dim = "width" if mode == "spatial_shard_width" else "height"
-        wan_spatial_shard.install_wan_spatial_shard_decode(self, executor.group, split_dim=split_dim)
-
    def tile_split(self, z: torch.Tensor) -> tuple[list[TileTask], GridSpec]:
        _, _, num_frames, height, width = z.shape
        sample_height = height * self.spatial_compression_ratio
@@ -307,26 +270,20 @@ class DistributedAutoencoderKLWan(OmniAutoencoderKLWan, DistributedVaeMixin):
        dec = torch.clamp(dec, min=-1.0, max=1.0)
        return dec

-    def _spatial_shard_decode_split_dim(
-        self,
-        z: torch.Tensor | None = None,
-        world_size: int | None = None,
-    ) -> str | None:
+    def _spatial_shard_decode_split_dim(self) -> str | None:
        mode = self.distributed_executor.parallel_mode
        if mode == "spatial_shard_width":
            return "width"
        if mode == "spatial_shard_height":
            return "height"
-        if mode == "auto" and z is not None and world_size is not None:
-            return select_auto_spatial_shard_split_dim(tuple(z.shape), world_size)
        return None

    def _spatial_shard_decode_enabled(self, z: torch.Tensor) -> bool:
-        if self.distributed_executor.parallel_mode == "tile":
+        split_dim = self._spatial_shard_decode_split_dim()
+        if split_dim is None:
            return False
        if z.ndim != 5:
-            if self.distributed_executor.parallel_mode != "tile":
-                logger.warning("Wan VAE spatial sharded decode expects 5D latent input; falling back to tiled decode.")
+            logger.warning("Wan VAE spatial sharded decode expects 5D latent input; falling back to tiled decode.")
            return False
        if not self.is_distributed_enabled():
            return False
@@ -335,32 +292,23 @@ class DistributedAutoencoderKLWan(OmniAutoencoderKLWan, DistributedVaeMixin):
        world_size = dist.get_world_size(group=group)
        requested_size = int(self.distributed_executor.parallel_size)
        if requested_size != world_size:
-            if self.distributed_executor.parallel_mode == "auto":
-                logger.debug(
-                    "Wan VAE auto decode selected tile mode for a partial DiT group (requested=%s dit_group=%s)",
-                    requested_size,
-                    world_size,
-                )
-            elif self._spatial_shard_decode_split_dim() is not None:
-                logger.warning(
-                    "Wan VAE spatial sharded decode currently requires vae_patch_parallel_size "
-                    "to match the DIT group size; falling back to tiled decode. "
-                    "requested=%s dit_group=%s split_dim=%s",
-                    requested_size,
-                    world_size,
-                    self._spatial_shard_decode_split_dim(),
-                )
+            logger.warning(
+                "Wan VAE spatial sharded decode currently requires vae_patch_parallel_size "
+                "to match the DIT group size; falling back to tiled decode. "
+                "requested=%s dit_group=%s split_dim=%s",
+                requested_size,
+                world_size,
+                split_dim,
+            )
            return False
-        return self._spatial_shard_decode_split_dim(z, world_size) is not None
+        return True

    def tiled_decode(self, z: torch.Tensor, return_dict: bool = True):
        if not self.is_distributed_enabled():
            return super().tiled_decode(z, return_dict=return_dict)

        if self._spatial_shard_decode_enabled(z):
-            world_size = dist.get_world_size(group=self.distributed_executor.group)
-            split_dim = self._spatial_shard_decode_split_dim(z, world_size)
-            assert split_dim is not None
+            split_dim = self._spatial_shard_decode_split_dim()
            logger.debug("Decode running with Wan VAE spatial_shard_%s mode", split_dim)
            return wan_spatial_shard.spatial_shard_decode(
                self,
@@ -52,11 +52,6 @@ _SPATIAL_SHARD_CONTEXT: ContextVar[SpatialShardContext | None] = ContextVar(
 )


-def _active_spatial_shard_split_dim() -> str | None:
-    context = _SPATIAL_SHARD_CONTEXT.get()
-    return context.split_dim if context is not None else None
-
-
 def _spatial_dim(split_dim: str) -> int:
    if split_dim == "height":
        return -2
@@ -337,27 +332,21 @@ class WanDistZeroPad2d(nn.Module):
        padding: tuple[int, int, int, int],
        group: dist.ProcessGroup,
        *,
-        defer_split_after_padding: bool = False,
+        split_dim: str = "height",
+        split_padding: tuple[int, int] | None = None,
    ) -> None:
        super().__init__()
        self.padding = padding
-        self.defer_split_after_padding = defer_split_after_padding
+        self.split_dim = split_dim
+        default_split_padding = (padding[2], padding[3]) if split_dim == "height" else (padding[0], padding[1])
+        self.split_padding = split_padding or default_split_padding
        self.group = group

    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        split_dim = _active_spatial_shard_split_dim()
-        if split_dim is None:
-            return F.pad(x, self.padding)
-
        rank, world_size = _rank_world(self.group)
        left, right, top, bottom = self.padding
-        if self.defer_split_after_padding:
-            if split_dim == "height":
-                bottom = 0
-            else:
-                right = 0
        if world_size > 1:
-            if split_dim == "height":
+            if self.split_dim == "height":
                top = top if rank == 0 else 0
                bottom = bottom if rank == world_size - 1 else 0
            else:
@@ -371,7 +360,8 @@ class WanDistConv2d(nn.Conv2d):
        self,
        source: nn.Conv2d,
        group: dist.ProcessGroup,
-        deferred_padding: tuple[int, int, int, int] | None = None,
+        split_dim: str = "height",
+        split_padding: tuple[int, int] | None = None,
    ):
        super().__init__(
            source.in_channels,
@@ -388,117 +378,75 @@ class WanDistConv2d(nn.Conv2d):
        )
        self.load_state_dict(source.state_dict())
        self.group = group
-        self._direct_padding = source.padding
-        self._direct_padding_mode = source.padding_mode
-        self._direct_reversed_padding = source._reversed_padding_repeated_twice
-        self._deferred_padding = deferred_padding
-        source_padding = source.padding if isinstance(source.padding, tuple) else (source.padding, source.padding)
-        if len(source_padding) != 2 or not all(isinstance(value, int) for value in source_padding):
-            raise ValueError(f"Wan spatial-shard Conv2d requires integer padding, got {source.padding!r}.")
-        self._spatial_pad_h = int(source_padding[-2])
-        self._spatial_pad_w = int(source_padding[-1])
-        self.register_buffer("_halo_recv_top_buf", None, persistent=False)
-        self.register_buffer("_halo_recv_bottom_buf", None, persistent=False)
-        self._trim_cache: dict[tuple[str, int], tuple[int, int, int]] = {}
+        self.split_dim = split_dim
+        self.split_tensor_dim = _spatial_dim(split_dim)
+        kernel_extent = self.kernel_size[-2] if split_dim == "height" else self.kernel_size[-1]
+        self.halo_size = (kernel_extent - 1) // 2
+        pad_h = source.padding[-2] if isinstance(source.padding, tuple) else source.padding
+        pad_w = source.padding[-1] if isinstance(source.padding, tuple) else source.padding
+        if split_dim == "height":
+            if split_padding is None:
+                split_padding = (pad_h, pad_h)
+            self._non_split_padding = (pad_w, pad_w, 0, 0)
+            self.kernel_extent = self.kernel_size[-2]
+            self.stride_extent = self.stride[-2]
+        else:
+            if split_padding is None:
+                split_padding = (pad_w, pad_w)
+            self._non_split_padding = (0, 0, pad_h, pad_h)
+            self.kernel_extent = self.kernel_size[-1]
+            self.stride_extent = self.stride[-1]
+        self.split_pad_left, self.split_pad_right = split_padding
+        self._halo_recv_top_buf: torch.Tensor | None = None
+        self._halo_recv_bottom_buf: torch.Tensor | None = None
+        self._trim_cache: dict[int, tuple[int, int, int]] = {}

    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        split_dim = _active_spatial_shard_split_dim()
-        if split_dim is None:
-            if self._direct_padding_mode != "zeros":
-                x = F.pad(x, self._direct_reversed_padding, mode=self._direct_padding_mode)
-                padding: str | tuple[int, int] = (0, 0)
-            else:
-                padding = self._direct_padding
-            return F.conv2d(x, self.weight, self.bias, self.stride, padding, self.dilation, self.groups)
-
-        split_tensor_dim = _spatial_dim(split_dim)
-        if split_dim == "height":
-            kernel_extent = self.kernel_size[-2]
-            stride_extent = self.stride[-2]
-            non_split_padding = (self._spatial_pad_w, self._spatial_pad_w, 0, 0)
-            default_split_padding = (self._spatial_pad_h, self._spatial_pad_h)
-            deferred_split_padding = (
-                (self._deferred_padding[2], self._deferred_padding[3]) if self._deferred_padding is not None else None
-            )
-        else:
-            kernel_extent = self.kernel_size[-1]
-            stride_extent = self.stride[-1]
-            non_split_padding = (0, 0, self._spatial_pad_h, self._spatial_pad_h)
-            default_split_padding = (self._spatial_pad_w, self._spatial_pad_w)
-            deferred_split_padding = (
-                (self._deferred_padding[0], self._deferred_padding[1]) if self._deferred_padding is not None else None
-            )
-        split_pad_left, split_pad_right = deferred_split_padding or default_split_padding
-        halo_size = (kernel_extent - 1) // 2
-
-        x = F.pad(x, non_split_padding)
+        x = F.pad(x, self._non_split_padding)
        x_padded, self._halo_recv_top_buf, self._halo_recv_bottom_buf = halo_exchange(
            x,
            group=self.group,
-            halo_size=halo_size,
-            split_dim=split_dim,
+            halo_size=self.halo_size,
+            split_dim=self.split_dim,
            recv_top_buf=self._halo_recv_top_buf,
            recv_bottom_buf=self._halo_recv_bottom_buf,
        )
-        shift, start, upper_bound = self._get_trim_params(
-            x.shape[split_tensor_dim],
-            split_dim=split_dim,
-            halo_size=halo_size,
-            split_pad_left=split_pad_left,
-            split_pad_right=split_pad_right,
-            kernel_extent=kernel_extent,
-            stride_extent=stride_extent,
-        )
+        shift, start, upper_bound = self._get_trim_params(x.shape[self.split_tensor_dim])
        if shift:
            x_padded = _narrow_along_dim(
                x_padded,
-                split_tensor_dim,
+                self.split_tensor_dim,
                shift,
-                x_padded.shape[split_tensor_dim] - shift,
+                x_padded.shape[self.split_tensor_dim] - shift,
            )
        out = super().forward(x_padded)
-        out = _trim_local_conv_output(out, halo_size, start, upper_bound, split_dim=split_dim)
-        return _zero_invalid_extent(out, split_dim=split_dim)
+        out = _trim_local_conv_output(out, self.halo_size, start, upper_bound, split_dim=self.split_dim)
+        return _zero_invalid_extent(out, split_dim=self.split_dim)

-    def _get_trim_params(
-        self,
-        local_extent: int,
-        *,
-        split_dim: str,
-        halo_size: int,
-        split_pad_left: int,
-        split_pad_right: int,
-        kernel_extent: int,
-        stride_extent: int,
-    ) -> tuple[int, int, int]:
-        cache_key = (split_dim, local_extent)
-        trim_params = self._trim_cache.get(cache_key)
+    def _get_trim_params(self, local_extent: int) -> tuple[int, int, int]:
+        trim_params = self._trim_cache.get(local_extent)
        if trim_params is None:
            rank, world_size = _rank_world(self.group)
            trim_params = _compute_conv_trim_params(
                local_extent=local_extent,
                rank=rank,
                world_size=world_size,
-                halo_size=halo_size,
-                pad_before=split_pad_left,
-                pad_after=split_pad_right,
-                kernel_extent=kernel_extent,
-                stride_extent=stride_extent,
+                halo_size=self.halo_size,
+                pad_before=self.split_pad_left,
+                pad_after=self.split_pad_right,
+                kernel_extent=self.kernel_extent,
+                stride_extent=self.stride_extent,
            )
-            self._trim_cache[cache_key] = trim_params
+            self._trim_cache[local_extent] = trim_params
        return trim_params


 class WanDistCausalConv3d(nn.Conv3d):
-    # Cross-PR compatibility contract for the lossless Wan data-movement
-    # installer: it may fuse the no-context direct path, but must leave this
-    # wrapper in control whenever a spatial-shard context is active.
-    _vllm_omni_dynamic_spatial_shard_conv = True
-
    def __init__(
        self,
        source: nn.Conv3d,
        group: dist.ProcessGroup,
+        split_dim: str = "height",
    ):
        super().__init__(
            source.in_channels,
@@ -515,98 +463,88 @@ class WanDistCausalConv3d(nn.Conv3d):
        )
        self.load_state_dict(source.state_dict())
        self.group = group
+        self.split_dim = split_dim
+        self.split_tensor_dim = _spatial_dim(split_dim)
        source_padding = getattr(source, "_padding", None)
        if source_padding is None:
            p_t, p_h, p_w = source.padding
            source_padding = (p_w, p_w, p_h, p_h, 2 * p_t, 0)
        self._source_padding = tuple(source_padding)
-        self.register_buffer("_halo_recv_top_buf", None, persistent=False)
-        self.register_buffer("_halo_recv_bottom_buf", None, persistent=False)
-        self._trim_cache: dict[tuple[str, int], tuple[int, int, int]] = {}
+        if split_dim == "height":
+            self.split_pad_left = int(self._source_padding[2])
+            self.split_pad_right = int(self._source_padding[3])
+            self.kernel_extent = self.kernel_size[-2]
+            self.stride_extent = self.stride[-2]
+            padding = (
+                self._source_padding[0],
+                self._source_padding[1],
+                0,
+                0,
+                self._source_padding[4],
+                self._source_padding[5],
+            )
+        else:
+            self.split_pad_left = int(self._source_padding[0])
+            self.split_pad_right = int(self._source_padding[1])
+            self.kernel_extent = self.kernel_size[-1]
+            self.stride_extent = self.stride[-1]
+            padding = (
+                0,
+                0,
+                self._source_padding[2],
+                self._source_padding[3],
+                self._source_padding[4],
+                self._source_padding[5],
+            )
+        self.halo_size = (self.kernel_extent - 1) // 2
+        self._padding = padding if self.halo_size > 0 else self._source_padding
+        self._halo_recv_top_buf: torch.Tensor | None = None
+        self._halo_recv_bottom_buf: torch.Tensor | None = None
+        self._trim_cache: dict[int, tuple[int, int, int]] = {}

    def forward(self, x: torch.Tensor, cache_x: torch.Tensor | None = None) -> torch.Tensor:
-        split_dim = _active_spatial_shard_split_dim()
-        padding = list(self._source_padding)
+        padding = list(self._padding)
        if cache_x is not None and padding[4] > 0:
            cache_x = cache_x.to(x.device)
            x = torch.cat([cache_x, x], dim=2)
            padding[4] -= cache_x.shape[2]

-        if split_dim is None:
-            return super().forward(F.pad(x, padding))
-
-        split_tensor_dim = _spatial_dim(split_dim)
-        if split_dim == "height":
-            split_pad_left = int(padding[2])
-            split_pad_right = int(padding[3])
-            kernel_extent = self.kernel_size[-2]
-            stride_extent = self.stride[-2]
-            padding[2] = 0
-            padding[3] = 0
-        else:
-            split_pad_left = int(padding[0])
-            split_pad_right = int(padding[1])
-            kernel_extent = self.kernel_size[-1]
-            stride_extent = self.stride[-1]
-            padding[0] = 0
-            padding[1] = 0
-        halo_size = (kernel_extent - 1) // 2
-
        x = F.pad(x, padding)
        x_padded, self._halo_recv_top_buf, self._halo_recv_bottom_buf = halo_exchange(
            x,
            group=self.group,
-            halo_size=halo_size,
-            split_dim=split_dim,
+            halo_size=self.halo_size,
+            split_dim=self.split_dim,
            recv_top_buf=self._halo_recv_top_buf,
            recv_bottom_buf=self._halo_recv_bottom_buf,
        )
-        shift, start, upper_bound = self._get_trim_params(
-            x.shape[split_tensor_dim],
-            split_dim=split_dim,
-            halo_size=halo_size,
-            split_pad_left=split_pad_left,
-            split_pad_right=split_pad_right,
-            kernel_extent=kernel_extent,
-            stride_extent=stride_extent,
-        )
+        shift, start, upper_bound = self._get_trim_params(x.shape[self.split_tensor_dim])
        if shift:
            x_padded = _narrow_along_dim(
                x_padded,
-                split_tensor_dim,
+                self.split_tensor_dim,
                shift,
-                x_padded.shape[split_tensor_dim] - shift,
+                x_padded.shape[self.split_tensor_dim] - shift,
            )
        out = super().forward(x_padded)
-        out = _trim_local_conv_output(out, halo_size, start, upper_bound, split_dim=split_dim)
-        return _zero_invalid_extent(out, split_dim=split_dim)
+        out = _trim_local_conv_output(out, self.halo_size, start, upper_bound, split_dim=self.split_dim)
+        return _zero_invalid_extent(out, split_dim=self.split_dim)

-    def _get_trim_params(
-        self,
-        local_extent: int,
-        *,
-        split_dim: str,
-        halo_size: int,
-        split_pad_left: int,
-        split_pad_right: int,
-        kernel_extent: int,
-        stride_extent: int,
-    ) -> tuple[int, int, int]:
-        cache_key = (split_dim, local_extent)
-        trim_params = self._trim_cache.get(cache_key)
+    def _get_trim_params(self, local_extent: int) -> tuple[int, int, int]:
+        trim_params = self._trim_cache.get(local_extent)
        if trim_params is None:
            rank, world_size = _rank_world(self.group)
            trim_params = _compute_conv_trim_params(
                local_extent=local_extent,
                rank=rank,
                world_size=world_size,
-                halo_size=halo_size,
-                pad_before=split_pad_left,
-                pad_after=split_pad_right,
-                kernel_extent=kernel_extent,
-                stride_extent=stride_extent,
+                halo_size=self.halo_size,
+                pad_before=self.split_pad_left,
+                pad_after=self.split_pad_right,
+                kernel_extent=self.kernel_extent,
+                stride_extent=self.stride_extent,
            )
-            self._trim_cache[cache_key] = trim_params
+            self._trim_cache[local_extent] = trim_params
        return trim_params


@@ -653,15 +591,12 @@ def _trim_local_conv_output(
    return out


-def _patch_attention_block(module: nn.Module, group: dist.ProcessGroup) -> None:
+def _patch_attention_block(module: nn.Module, group: dist.ProcessGroup, split_dim: str) -> None:
    if getattr(module, "_vllm_omni_spatial_shard_attention", False):
        return
    orig_forward = module.forward

    def _forward(self: nn.Module, x: torch.Tensor, *args: Any, **kwargs: Any) -> torch.Tensor:
-        split_dim = _active_spatial_shard_split_dim()
-        if split_dim is None:
-            return orig_forward(x, *args, **kwargs)
        _, world_size = _rank_world(group)
        if world_size <= 1:
            return orig_forward(x, *args, **kwargs)
@@ -695,6 +630,7 @@ def _replace_child(
    name: str,
    child: nn.Module,
    group: dist.ProcessGroup,
+    split_dim: str,
 ) -> None:
    if child.__class__.__name__ == "WanCausalConv3d":
        setattr(
@@ -703,39 +639,51 @@ def _replace_child(
            WanDistCausalConv3d(
                child,
                group,
+                split_dim=split_dim,
            ),
        )
        return
    if isinstance(child, nn.ZeroPad2d):
        padding = tuple(int(p) for p in child.padding)
+        module_padding = padding
+        if parent.__class__.__name__ == "Sequential":
+            # Let the following WanDistConv2d account for global after-edge
+            # padding; this module handles the non-split dimension and before edge.
+            if split_dim == "height":
+                module_padding = (padding[0], padding[1], padding[2], 0)
+            else:
+                module_padding = (padding[0], 0, padding[2], padding[3])
        setattr(
            parent,
            name,
            WanDistZeroPad2d(
-                padding,
+                module_padding,
                group,
-                defer_split_after_padding=parent.__class__.__name__ == "Sequential",
+                split_dim=split_dim,
+                split_padding=(padding[2], padding[3]) if split_dim == "height" else (padding[0], padding[1]),
            ),
        )
        return
    if isinstance(child, nn.Conv2d):
-        deferred_padding = None
+        split_padding = None
        if name == "1" and parent.__class__.__name__ == "Sequential":
            # WanResample downsample uses ZeroPad2d((0, 1, 0, 1)) before a
            # stride-2 conv with padding=0.  Only the last rank should see the
            # bottom/right global padding, which is approximated by split_padding.
            prev = getattr(parent, "0", None)
            if isinstance(prev, WanDistZeroPad2d):
-                deferred_padding = prev.padding
+                split_padding = prev.split_padding
            elif isinstance(prev, nn.ZeroPad2d):
-                deferred_padding = tuple(int(value) for value in prev.padding)
+                pad = prev.padding
+                split_padding = (int(pad[2]), int(pad[3])) if split_dim == "height" else (int(pad[0]), int(pad[1]))
        setattr(
            parent,
            name,
            WanDistConv2d(
                child,
                group,
-                deferred_padding=deferred_padding,
+                split_dim=split_dim,
+                split_padding=split_padding,
            ),
        )

@@ -743,20 +691,21 @@ def _replace_child(
 def _patch_decoder_modules(
    module: nn.Module,
    group: dist.ProcessGroup,
+    split_dim: str,
    inside_attention: bool = False,
 ) -> None:
    if module.__class__.__name__ == "WanAttentionBlock":
-        _patch_attention_block(module, group)
+        _patch_attention_block(module, group, split_dim)
        inside_attention = True

    for name, child in list(module.named_children()):
        if child.__class__.__name__ == "WanCausalConv3d":
-            _replace_child(module, name, child, group)
+            _replace_child(module, name, child, group, split_dim)
            continue
        if isinstance(child, nn.Conv2d) and not inside_attention:
-            _replace_child(module, name, child, group)
+            _replace_child(module, name, child, group, split_dim)
            continue
-        _patch_decoder_modules(child, group, inside_attention=inside_attention)
+        _patch_decoder_modules(child, group, split_dim, inside_attention=inside_attention)


 def _decoder_upsample_count(decoder: nn.Module) -> int:
@@ -767,25 +716,16 @@ def _decoder_upsample_count(decoder: nn.Module) -> int:
    return count


-def _clear_wan_spatial_shard_runtime_buffers(vae: Any) -> None:
-    decoder = getattr(vae, "decoder", None)
-    if decoder is None:
-        return
-    for module in decoder.modules():
-        if "_halo_recv_top_buf" in module._buffers:
-            module._halo_recv_top_buf = None
-        if "_halo_recv_bottom_buf" in module._buffers:
-            module._halo_recv_bottom_buf = None
-
-
 def install_wan_spatial_shard_decode(vae: Any, group: dist.ProcessGroup, split_dim: str = "height") -> None:
    """Patch ``vae.decoder`` once for spatially-sharded decode.

    This mutates the already-loaded decoder in place by swapping its spatial
-    convolutions/padding for context-aware variants and attaching a dedicated
-    spatial forward. Outside that forward the replacements execute their exact
-    local PyTorch behavior, so a later auto-selected tile request remains safe.
-    The same installed decoder can alternate height and width sharding.
+    convolutions/padding for halo-exchanging variants and wrapping
+    ``decoder.forward``. The patch is permanent for the lifetime of the VAE
+    instance and is applied only once (subsequent calls are no-ops). A given
+    instance is bound to a single ``split_dim``; switching between
+    ``"height"`` and ``"width"`` requires a fresh VAE instance and raises here
+    otherwise.

    Only group-relative rank 0 assembles the final decoded frame, mirroring the
    distributed tiled-decode ``broadcast_result=False`` contract; the other ranks
@@ -793,12 +733,18 @@ def install_wan_spatial_shard_decode(vae: Any, group: dist.ProcessGroup, split_d
    """
    _spatial_dim(split_dim)
    if getattr(vae, "_vllm_omni_wan_spatial_shard_installed", False):
+        installed_split_dim = getattr(vae, "_vllm_omni_wan_spatial_shard_split_dim", "height")
+        if installed_split_dim != split_dim:
+            raise ValueError(
+                "Wan spatial-shard VAE decoder was already patched for "
+                f"{installed_split_dim!r} split; create a fresh VAE instance to use {split_dim!r} split."
+            )
        return
    decoder = getattr(vae, "decoder", None)
    if decoder is None:
        raise ValueError("Wan spatial-shard VAE decode requires a decoder module.")

-    _patch_decoder_modules(decoder, group)
+    _patch_decoder_modules(decoder, group, split_dim)
    upsample_count = _decoder_upsample_count(decoder)
    orig_forward = decoder.forward

@@ -808,8 +754,6 @@ def install_wan_spatial_shard_decode(vae: Any, group: dist.ProcessGroup, split_d
        feat_cache: list[torch.Tensor] | None = None,
        feat_idx: list[int] | None = None,
        first_chunk: bool = False,
-        *,
-        split_dim: str,
    ) -> torch.Tensor:
        if feat_idx is None:
            feat_idx = [0]
@@ -837,9 +781,10 @@ def install_wan_spatial_shard_decode(vae: Any, group: dist.ProcessGroup, split_d
            _SPATIAL_SHARD_CONTEXT.reset(token)
        return gather_and_trim_extent(out, expected_extent=expected_extent, split_dim=split_dim, group=group, dst=0)

-    decoder._vllm_omni_spatial_shard_forward = MethodType(_forward, decoder)
+    decoder.forward = MethodType(_forward, decoder)
    vae._vllm_omni_wan_spatial_shard_installed = True
-    logger.info("Installed Wan VAE dynamic-axis spatial-shard decode.")
+    vae._vllm_omni_wan_spatial_shard_split_dim = split_dim
+    logger.info("Installed Wan VAE %s-sharded decode.", split_dim)


 def spatial_shard_decode(
@@ -868,12 +813,11 @@ def spatial_shard_decode(
            decoded_chunks = []
            for i in range(z.shape[2]):
                vae._conv_idx = [0]
-                chunk = vae.decoder._vllm_omni_spatial_shard_forward(
+                chunk = vae.decoder(
                    x[:, :, i : i + 1, :, :],
                    feat_cache=vae._feat_map,
                    feat_idx=vae._conv_idx,
                    first_chunk=(i == 0),
-                    split_dim=split_dim,
                )
                if produce_output:
                    decoded_chunks.append(chunk)
@@ -887,10 +831,6 @@ def spatial_shard_decode(
                out = z.new_zeros(0)
    finally:
        vae.clear_cache()
-        # Halo buffers are request-shape scratch space allocated outside the
-        # tagged weight pool. Drop live references so sleep/offload never has
-        # to preserve or discard them as model state.
-        _clear_wan_spatial_shard_runtime_buffers(vae)

    if not return_dict:
        return (out,)
@@ -18,6 +18,7 @@ import zmq
 from vllm.distributed.device_communicators.shm_broadcast import Handle, MessageQueue
 from vllm.logger import init_logger
 from vllm.v1.engine.exceptions import EngineDeadError
+from vllm.v1.executor.multiproc_executor import set_multiprocessing_worker_envs

 from vllm_omni.diffusion.data import SHUTDOWN_MESSAGE, AsyncDiffusionOutput, AsyncOutputKind, DiffusionOutput
 from vllm_omni.diffusion.executor.abstract import DiffusionExecutor
@@ -294,6 +295,10 @@ class MultiprocDiffusionExecutor(DiffusionExecutor):
        logger.info("Starting server...")

        num_gpus = cast(int, od_config.num_gpus)
+        # Without this, every worker inherits one Torch thread per core, so an
+        # N-GPU run oversubscribes the host by N x core_count. Honours a
+        # user-provided OMP_NUM_THREADS.
+        set_multiprocessing_worker_envs()
        mp.set_start_method("spawn", force=True)
        processes = []

@@ -0,0 +1,299 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""Q/K RMSNorm followed by packed non-interleaved RoPE.
+
+The public contract is shared by diffusion attention implementations:
+
+* ``q`` and ``k`` are ``[tokens, heads, head_dim]``;
+* norm weights are one-dimensional ``[head_dim]`` tensors;
+* ``rope_table`` is ``[tokens, rotary_dim]`` and stores
+  ``[cos(theta), sin(theta)]`` with ``theta`` of width ``rotary_dim // 2``.
+
+The CUDA fast path fuses RMSNorm and RoPE without materializing normalized Q/K
+or rotary-product intermediates. Unsupported inputs use the eager reference.
+"""
+
+from __future__ import annotations
+
+import torch
+import torch.nn.functional as F
+from torch.library import Library
+from vllm.platforms import current_platform
+from vllm.triton_utils import HAS_TRITON, tl, triton
+from vllm.utils.torch_utils import direct_register_custom_op
+
+_FUSED_HEAD_DIM = 128
+_FUSED_ROTARY_DIM = 96
+_HEADS_PER_PROGRAM = 8
+
+
+def _apply_rope_table(
+    x: torch.Tensor,
+    rope_table: torch.Tensor,
+    rotary_dim: int,
+) -> torch.Tensor:
+    half = rotary_dim // 2
+    cos = rope_table[..., :half].to(x.dtype).unsqueeze(1)
+    sin = rope_table[..., half:].to(x.dtype).unsqueeze(1)
+    first = x[..., :half]
+    second = x[..., half:rotary_dim]
+    return torch.cat(
+        (
+            first * cos - second * sin,
+            second * cos + first * sin,
+            x[..., rotary_dim:],
+        ),
+        dim=-1,
+    )
+
+
+if HAS_TRITON:
+
+    @triton.jit
+    def _rms_norm_rope_kernel(
+        x_ptr,
+        weight_ptr,
+        rope_table_ptr,
+        out_ptr,
+        x_stride_t,
+        x_stride_h,
+        x_stride_d,
+        rope_stride_t,
+        out_stride_t,
+        out_stride_h,
+        out_stride_d,
+        num_heads: tl.constexpr,
+        head_dim: tl.constexpr,
+        rotary_half: tl.constexpr,
+        eps: tl.constexpr,
+        heads_per_program: tl.constexpr,
+    ):
+        token = tl.program_id(0)
+        head_group = tl.program_id(1)
+        heads = head_group * heads_per_program + tl.arange(0, heads_per_program)
+        dims = tl.arange(0, head_dim)
+        mask = heads[:, None] < num_heads
+        offsets = token * x_stride_t + heads[:, None] * x_stride_h + dims[None, :] * x_stride_d
+
+        x = tl.load(x_ptr + offsets, mask=mask, other=0.0).to(tl.float32)
+        weight = tl.load(weight_ptr + dims).to(tl.float32)
+        inv_rms = tl.rsqrt(tl.sum(x * x, axis=1) / head_dim + eps)
+        normalized = (x * inv_rms[:, None] * weight[None, :]).to(tl.bfloat16)
+
+        rotary_dim = rotary_half * 2
+        pair_dims = tl.where(
+            dims < rotary_half,
+            dims + rotary_half,
+            tl.where(dims < rotary_dim, dims - rotary_half, dims),
+        )
+        pair_offsets = token * x_stride_t + heads[:, None] * x_stride_h + pair_dims[None, :] * x_stride_d
+        pair_x = tl.load(x_ptr + pair_offsets, mask=mask, other=0.0).to(tl.float32)
+        pair_weight = tl.load(weight_ptr + pair_dims).to(tl.float32)
+        pair_normalized = (pair_x * inv_rms[:, None] * pair_weight[None, :]).to(tl.bfloat16)
+
+        freq_dims = tl.where(
+            dims < rotary_half,
+            dims,
+            tl.where(dims < rotary_dim, dims - rotary_half, 0),
+        )
+        table_offsets = token * rope_stride_t + freq_dims
+        cos = tl.load(rope_table_ptr + table_offsets).to(tl.float32)
+        sin = tl.load(rope_table_ptr + table_offsets + rotary_half).to(tl.float32)
+        first = normalized.to(tl.float32) * cos - pair_normalized.to(tl.float32) * sin
+        second = normalized.to(tl.float32) * cos + pair_normalized.to(tl.float32) * sin
+        output = tl.where(
+            dims < rotary_dim,
+            tl.where(dims < rotary_half, first, second),
+            normalized.to(tl.float32),
+        )
+
+        out_offsets = token * out_stride_t + heads[:, None] * out_stride_h + dims[None, :] * out_stride_d
+        tl.store(out_ptr + out_offsets, output, mask=mask)
+
+
+def _eager_qk_norm_rope(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    q_weight: torch.Tensor,
+    k_weight: torch.Tensor,
+    rope_table: torch.Tensor,
+    eps: float,
+    head_dim: int,
+    rotary_dim: int,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    q_norm = F.rms_norm(q, (head_dim,), q_weight, eps)
+    k_norm = F.rms_norm(k, (head_dim,), k_weight, eps)
+    return (
+        _apply_rope_table(q_norm, rope_table, rotary_dim),
+        _apply_rope_table(k_norm, rope_table, rotary_dim),
+    )
+
+
+def _fused_cuda_supported(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    head_dim: int,
+    rotary_dim: int,
+) -> bool:
+    return (
+        HAS_TRITON
+        and current_platform.is_cuda()
+        and q.is_cuda
+        and k.is_cuda
+        and q.dtype == torch.bfloat16
+        and k.dtype == torch.bfloat16
+        and head_dim == _FUSED_HEAD_DIM
+        and rotary_dim == _FUSED_ROTARY_DIM
+    )
+
+
+def _launch_fused_rms_norm_rope(
+    x: torch.Tensor,
+    weight: torch.Tensor,
+    rope_table: torch.Tensor,
+    eps: float,
+) -> torch.Tensor:
+    tokens, heads, head_dim = x.shape
+    rotary_half = rope_table.shape[-1] // 2
+    out = torch.empty(x.shape, dtype=x.dtype, device=x.device)
+    if tokens == 0:
+        return out
+    grid = (tokens, triton.cdiv(heads, _HEADS_PER_PROGRAM))
+    _rms_norm_rope_kernel[grid](
+        x,
+        weight,
+        rope_table,
+        out,
+        x.stride(0),
+        x.stride(1),
+        x.stride(2),
+        rope_table.stride(0),
+        out.stride(0),
+        out.stride(1),
+        out.stride(2),
+        num_heads=heads,
+        head_dim=head_dim,
+        rotary_half=rotary_half,
+        eps=eps,
+        heads_per_program=_HEADS_PER_PROGRAM,
+        num_warps=8,
+    )
+    return out
+
+
+def _fused_qk_norm_rope_impl(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    q_weight: torch.Tensor,
+    k_weight: torch.Tensor,
+    rope_table: torch.Tensor,
+    eps: float,
+    head_dim: int,
+    rotary_dim: int,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    if not _fused_cuda_supported(q, k, head_dim, rotary_dim):
+        return _eager_qk_norm_rope(
+            q,
+            k,
+            q_weight,
+            k_weight,
+            rope_table,
+            eps,
+            head_dim,
+            rotary_dim,
+        )
+    return (
+        _launch_fused_rms_norm_rope(q, q_weight, rope_table, eps),
+        _launch_fused_rms_norm_rope(k, k_weight, rope_table, eps),
+    )
+
+
+def _fused_qk_norm_rope_fake(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    q_weight: torch.Tensor,
+    k_weight: torch.Tensor,
+    rope_table: torch.Tensor,
+    eps: float,
+    head_dim: int,
+    rotary_dim: int,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    del q_weight, k_weight, rope_table, eps, head_dim, rotary_dim
+    return torch.empty_like(q), torch.empty_like(k)
+
+
+_OMNI_OP_LIB = Library("vllm_omni", "FRAGMENT")
+if not hasattr(torch.ops.vllm_omni, "fused_qk_norm_rope"):
+    direct_register_custom_op(
+        op_name="fused_qk_norm_rope",
+        op_func=_fused_qk_norm_rope_impl,
+        fake_impl=_fused_qk_norm_rope_fake,
+        mutates_args=[],
+        target_lib=_OMNI_OP_LIB,
+    )
+
+
+def fused_qk_norm_rope(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    q_weight: torch.Tensor,
+    k_weight: torch.Tensor,
+    rope_table: torch.Tensor,
+    eps: float,
+    *,
+    head_dim: int | None = None,
+    rotary_dim: int | None = None,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Apply Q/K RMSNorm and packed non-interleaved RoPE."""
+    if q.ndim != 3 or k.ndim != 3:
+        raise ValueError(f"q and k must be [tokens, heads, head_dim], got {q.shape} and {k.shape}")
+    if q.shape[0] != k.shape[0] or q.shape[2] != k.shape[2]:
+        raise ValueError(f"q and k shapes are incompatible: {q.shape} and {k.shape}")
+    if q.dtype != k.dtype or q.device != k.device:
+        raise ValueError("q and k must have the same dtype and device")
+    if q.dtype not in (torch.bfloat16, torch.float16, torch.float32):
+        raise TypeError(f"Fused QK RMSNorm/RoPE requires floating inputs, got {q.dtype}")
+
+    head_dim = q.shape[-1] if head_dim is None else head_dim
+    rotary_dim = rope_table.shape[-1] if rotary_dim is None else rotary_dim
+    if q.shape[-1] != head_dim:
+        raise ValueError(f"Expected q/k head_dim={head_dim}, got {q.shape[-1]}")
+    if rotary_dim <= 0 or rotary_dim > head_dim or rotary_dim % 2:
+        raise ValueError(f"rotary_dim must be even and in [2, {head_dim}], got {rotary_dim}")
+    if q_weight.shape != (head_dim,) or k_weight.shape != (head_dim,):
+        raise ValueError(f"Expected norm weights [{head_dim}], got {tuple(q_weight.shape)} and {tuple(k_weight.shape)}")
+    if q_weight.device != q.device or k_weight.device != q.device:
+        raise ValueError("Q/K norm weights must be on the activation device")
+    if rope_table.device != q.device or rope_table.dtype != q.dtype:
+        raise ValueError("rope_table must have the same dtype and device as q/k")
+    if rope_table.shape != (q.shape[0], rotary_dim):
+        raise ValueError(f"Expected rope_table [{q.shape[0]}, {rotary_dim}], got {tuple(rope_table.shape)}")
+
+    q_weight = q_weight.contiguous()
+    k_weight = k_weight.contiguous()
+    rope_table = rope_table.contiguous()
+    if not _fused_cuda_supported(q, k, head_dim, rotary_dim):
+        return _fused_qk_norm_rope_impl(
+            q,
+            k,
+            q_weight,
+            k_weight,
+            rope_table,
+            eps,
+            head_dim,
+            rotary_dim,
+        )
+    return torch.ops.vllm_omni.fused_qk_norm_rope(
+        q,
+        k,
+        q_weight,
+        k_weight,
+        rope_table,
+        eps,
+        head_dim,
+        rotary_dim,
+    )
+
+
+__all__ = ["fused_qk_norm_rope"]
@@ -12,12 +12,16 @@ from vllm_omni.diffusion.models.ltx2.pipeline_ltx2 import (
    LTX2Pipeline,
    LTX2T2VDMD2Pipeline,
 )
-from vllm_omni.diffusion.models.ltx2.pipeline_ltx2_two_stage import LTX2DistilledPipeline
+from vllm_omni.diffusion.models.ltx2.pipeline_ltx2_two_stage import (
+    LTX2DistilledPipeline,
+    LTX2TwoStagePipeline,
+)

 __all__ = [
    "LTX2Pipeline",
    "LTX2T2VDMD2Pipeline",
    "LTX2I2VDMD2Pipeline",
+    "LTX2TwoStagePipeline",
    "LTX2DistilledPipeline",
    "get_ltx2_post_process_func",
    "load_transformer_config",
@@ -0,0 +1,156 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""Parse official LTX distilled adapters without materializing their tensors.
+
+The official LTX adapters are raw safetensors files rather than PEFT
+checkpoints.  This module is deliberately LTX-specific: it translates the
+official names into the locally constructed Diffusers transformer names and
+describes the resulting adapter targets.  Applying those targets is the job
+of :mod:`ltx2_phase_adapter`.
+"""
+
+from __future__ import annotations
+
+from collections.abc import Iterator
+from dataclasses import dataclass
+from pathlib import Path
+
+import torch
+import torch.nn as nn
+from safetensors import safe_open
+
+_IGNORED_OFFICIAL_LORA_MODULES = {"text_embedding_projection.aggregate_embed"}
+_OFFICIAL_TO_DIFFUSERS_PREFIXES = (
+    ("audio_prompt_adaln_single.", "audio_prompt_adaln."),
+    ("prompt_adaln_single.", "prompt_adaln."),
+    ("audio_adaln_single.", "audio_time_embed."),
+    ("adaln_single.", "time_embed."),
+    ("audio_patchify_proj", "audio_proj_in"),
+    ("patchify_proj", "proj_in"),
+    ("av_ca_video_scale_shift_adaln_single.", "av_cross_attn_video_scale_shift."),
+    ("av_ca_audio_scale_shift_adaln_single.", "av_cross_attn_audio_scale_shift."),
+    ("av_ca_a2v_gate_adaln_single.", "av_cross_attn_video_a2v_gate."),
+    ("av_ca_v2a_gate_adaln_single.", "av_cross_attn_audio_v2a_gate."),
+)
+
+
+@dataclass(frozen=True)
+class AdapterTarget:
+    """One A/B pair mapped onto a local transformer linear layer."""
+
+    # Keep the original Diffusers module name for diagnostics after LTX packs
+    # Q/K/V into the runtime module path.
+    source_module: str
+    module: str
+    slice: str | None
+    a_key: str
+    b_key: str
+    a_shape: tuple[int, ...]
+    b_shape: tuple[int, ...]
+
+
+@dataclass(frozen=True)
+class AdapterManifest:
+    """Tensor-free description of the fixed refinement adapter."""
+
+    path: str
+    targets: tuple[AdapterTarget, ...]
+
+
+def _to_diffusers_module_name(name: str) -> str:
+    for official_prefix, diffusers_prefix in _OFFICIAL_TO_DIFFUSERS_PREFIXES:
+        if name.startswith(official_prefix):
+            return diffusers_prefix + name[len(official_prefix) :]
+    return name
+
+
+def _resolve_lora_target(transformer: nn.Module, module_name: str) -> tuple[str, str | int | None] | None:
+    modules = dict(transformer.named_modules())
+    if module_name in modules:
+        return module_name, None
+    for packed_suffix, sub_suffix, shard_id in getattr(transformer, "stacked_params_mapping", ()):
+        if sub_suffix in module_name:
+            packed_name = module_name.replace(sub_suffix, packed_suffix)
+            if packed_name in modules:
+                return packed_name, shard_id
+    return None
+
+
+def parse_ltx_adapter(transformer: nn.Module, path: str) -> AdapterManifest:
+    """Translate one official raw LTX adapter into local Transformer targets."""
+    path = str(Path(path))
+    targets: list[AdapterTarget] = []
+    unresolved: list[str] = []
+    seen: set[tuple[str, str | int | None]] = set()
+
+    with safe_open(path, framework="pt", device="cpu") as handle:
+        tensor_names = set(handle.keys())
+        official_name_set = {
+            tensor_name.removeprefix("diffusion_model.").rsplit(".lora_", 1)[0]
+            for tensor_name in tensor_names
+            if tensor_name.endswith(".lora_A.weight")
+        }
+        official_b_names = {
+            tensor_name.removeprefix("diffusion_model.").rsplit(".lora_", 1)[0]
+            for tensor_name in tensor_names
+            if tensor_name.endswith(".lora_B.weight")
+        }
+        missing_a = sorted(official_b_names - official_name_set - _IGNORED_OFFICIAL_LORA_MODULES)
+        if missing_a:
+            raise ValueError(f"Missing paired LoRA A tensors for {missing_a} in {path}.")
+
+        official_names = sorted(official_name_set)
+        for official_name in official_names:
+            if official_name in _IGNORED_OFFICIAL_LORA_MODULES:
+                continue
+            key_prefix = f"diffusion_model.{official_name}"
+            key_a = f"{key_prefix}.lora_A.weight"
+            key_b = f"{key_prefix}.lora_B.weight"
+            if key_b not in tensor_names:
+                raise ValueError(f"Missing paired LoRA tensor {key_b!r} in {path}.")
+
+            a_shape = tuple(handle.get_slice(key_a).get_shape())
+            b_shape = tuple(handle.get_slice(key_b).get_shape())
+            if len(a_shape) != 2 or len(b_shape) != 2 or b_shape[1] != a_shape[0]:
+                raise ValueError(f"Invalid LoRA pair for {official_name}: A={a_shape}, B={b_shape}.")
+
+            source_module = _to_diffusers_module_name(official_name)
+            resolved = _resolve_lora_target(transformer, source_module)
+            if resolved is None:
+                unresolved.append(official_name)
+                continue
+            module, shard_id = resolved
+            target_id = (module, shard_id)
+            if target_id in seen:
+                raise ValueError(f"Official LTX LoRA contains duplicate target {module!r}, slice {shard_id!r}.")
+            seen.add(target_id)
+            targets.append(
+                AdapterTarget(
+                    source_module=source_module,
+                    module=module,
+                    slice=shard_id,
+                    a_key=key_a,
+                    b_key=key_b,
+                    a_shape=a_shape,
+                    b_shape=b_shape,
+                )
+            )
+
+    if unresolved:
+        raise ValueError(f"Official LTX LoRA contains unmapped modules: {unresolved}.")
+    if not targets:
+        raise ValueError(f"No Transformer LoRA tensors were found in {path}.")
+    return AdapterManifest(path=path, targets=tuple(targets))
+
+
+def iter_adapter_tensors(manifest: AdapterManifest) -> Iterator[tuple[AdapterTarget, torch.Tensor, torch.Tensor]]:
+    """Stream adapter pairs in manifest order without a full checkpoint copy."""
+
+    with safe_open(manifest.path, framework="pt", device="cpu") as handle:
+        for target in manifest.targets:
+            yield (
+                target,
+                handle.get_tensor(target.a_key),
+                handle.get_tensor(target.b_key),
+            )
@@ -9,7 +9,8 @@ import inspect
 import json
 import logging
 import os
-from dataclasses import dataclass
+from dataclasses import dataclass, replace
+from pathlib import Path
 from typing import TYPE_CHECKING, Any

 import torch
@@ -19,6 +20,8 @@ from diffusers.pipelines.ltx2.latent_upsampler import LTX2LatentUpsamplerModel
 from diffusers.pipelines.ltx2.vocoder import LTX2Vocoder
 from diffusers.video_processor import VideoProcessor
 from huggingface_hub import hf_hub_download
+from safetensors import safe_open
+from safetensors.torch import load_file
 from transformers import AutoTokenizer, Gemma3ForConditionalGeneration

 from vllm_omni.diffusion.attention.backends.abstract import AttentionMetadata
@@ -32,6 +35,7 @@ from vllm_omni.diffusion.offloader.module_collector import ModuleDiscovery
 if TYPE_CHECKING:
    from vllm.model_executor.layers.quantization.base_config import QuantizationConfig

+from .ltx2_request import LTXCheckpointKind, validate_ltx_checkpoint
 from .ltx2_transformer import (
    LTX2VideoTransformer3DModel,
    apply_interleaved_rotary_emb,
@@ -70,6 +74,9 @@ class LTXComponentProfile:
    video_vae_cls: type = AutoencoderKLLTX2Video
    vocoder_cls: type = LTX2Vocoder
    vocoder_fallback_cls: type | None = None
+    artifact_repo_id: str | None = None
+    latent_upsampler_filename: str | None = None
+    distilled_lora_filename: str | None = None


 LTX2_COMPONENT_PROFILE = LTXComponentProfile(
@@ -101,16 +108,96 @@ LTX2_DISTILLED_COMPONENT_PROFILE = LTXComponentProfile(
    video_vae_cls=DistributedAutoencoderKLLTX2Video,
 )

+LTX23_DISTILLED_COMPONENT_PROFILE = replace(
+    LTX23_COMPONENT_PROFILE,
+    name="ltx2_3_distilled",
+    resident_modules=(*LTX23_COMPONENT_PROFILE.resident_modules, "latent_upsampler"),
+    artifact_repo_id="Lightricks/LTX-2.3",
+    latent_upsampler_filename="ltx-2.3-spatial-upscaler-x2-1.1.safetensors",
+)
+
+LTX2_TWO_STAGE_COMPONENT_PROFILE = replace(
+    LTX2_COMPONENT_PROFILE,
+    name="ltx2_two_stage",
+    resident_modules=(*LTX2_COMPONENT_PROFILE.resident_modules, "latent_upsampler"),
+    artifact_repo_id="Lightricks/LTX-2",
+    latent_upsampler_filename="ltx-2-spatial-upscaler-x2-1.0.safetensors",
+    distilled_lora_filename="ltx-2-19b-distilled-lora-384.safetensors",
+)
+
+LTX23_TWO_STAGE_COMPONENT_PROFILE = replace(
+    LTX23_COMPONENT_PROFILE,
+    name="ltx2_3_two_stage",
+    resident_modules=(*LTX23_COMPONENT_PROFILE.resident_modules, "latent_upsampler"),
+    artifact_repo_id="Lightricks/LTX-2.3",
+    latent_upsampler_filename="ltx-2.3-spatial-upscaler-x2-1.1.safetensors",
+    distilled_lora_filename="ltx-2.3-22b-distilled-lora-384-1.1.safetensors",
+)
+

 _COMPONENT_PROFILES: dict[tuple[str, str], LTXComponentProfile] = {
    ("one_stage", "2"): LTX2_COMPONENT_PROFILE,
    ("one_stage", "2.3"): LTX23_COMPONENT_PROFILE,
+    ("two_stage", "2"): LTX2_TWO_STAGE_COMPONENT_PROFILE,
+    ("two_stage", "2.3"): LTX23_TWO_STAGE_COMPONENT_PROFILE,
    ("distilled_two_stage", "2"): LTX2_DISTILLED_COMPONENT_PROFILE,
+    ("distilled_two_stage", "2.3"): LTX23_DISTILLED_COMPONENT_PROFILE,
    ("dmd2", "2"): LTX2_COMPONENT_PROFILE,
    ("dmd2", "2.3"): LTX23_COMPONENT_PROFILE,
 }


+def resolve_ltx_artifact(
+    model: str,
+    repo_id: str,
+    filename: str,
+) -> str:
+    """Resolve an official LTX sidecar from the model root or its Hub repository."""
+    candidate = Path(model) / filename
+    if candidate.is_file():
+        return str(candidate)
+
+    try:
+        return hf_hub_download(repo_id=repo_id, filename=filename)
+    except Exception as exc:
+        raise FileNotFoundError(
+            f"Unable to resolve LTX artifact {filename!r}. Searched {candidate}; "
+            f"place the file in the model root or make {repo_id} available."
+        ) from exc
+
+
+def _load_ltx_latent_upsampler_single_file(path: str, dtype: torch.dtype) -> LTX2LatentUpsamplerModel:
+    """Load an official single-file upsampler into the Diffusers module."""
+    with safe_open(path, framework="pt", device="cpu") as handle:
+        metadata = handle.metadata() or {}
+    raw_config = json.loads(metadata.get("config", "{}"))
+    config = {
+        "in_channels": raw_config.get("in_channels", 128),
+        "mid_channels": raw_config.get("mid_channels", 1024),
+        "num_blocks_per_stage": raw_config.get("num_blocks_per_stage", 4),
+        "dims": raw_config.get("dims", 3),
+        "spatial_upsample": raw_config.get("spatial_upsample", True),
+        "temporal_upsample": raw_config.get("temporal_upsample", False),
+        "rational_spatial_scale": raw_config.get("spatial_scale", 2.0),
+        "use_rational_resampler": raw_config.get("rational_resampler", False),
+    }
+    with torch.device("cpu"):
+        upsampler = LTX2LatentUpsamplerModel(**config).to(dtype=dtype)
+
+    state_dict = load_file(path, device="cpu")
+    if "upsampler.0.weight" in state_dict and hasattr(upsampler.upsampler, "conv"):
+        state_dict["upsampler.conv.weight"] = state_dict.pop("upsampler.0.weight")
+        state_dict["upsampler.conv.bias"] = state_dict.pop("upsampler.0.bias")
+    missing, unexpected = upsampler.load_state_dict(state_dict, strict=False)
+    unresolved_missing = set(missing) - {"upsampler.blur_down.kernel"}
+    if unresolved_missing or unexpected:
+        raise ValueError(
+            f"Invalid LTX latent upsampler {path}: missing={sorted(unresolved_missing)}, "
+            f"unexpected={sorted(unexpected)}."
+        )
+    return upsampler
+
+
 def resolve_ltx_component_profile(pipeline_kind: str, model_version: str) -> LTXComponentProfile:
    """Resolve component construction independently from execution recipes."""
    try:
@@ -392,16 +479,34 @@ def initialize_pipeline_components(pipeline: Any, od_config: Any) -> None:
        )

    if "latent_upsampler" in profile.resident_modules:
-        # BlurDownsample constructs an integer kernel that must be initialized
-        # on CPU; component placement is handled uniformly after construction.
-        with torch.device("cpu"):
-            pipeline.latent_upsampler = _load_component(
-                LTX2LatentUpsamplerModel,
+        upsampler_config = os.path.join(model, "latent_upsampler", "config.json")
+        if os.path.isfile(upsampler_config) or not local_files_only:
+            try:
+                pipeline.latent_upsampler = _load_component(
+                    LTX2LatentUpsamplerModel,
+                    model,
+                    "latent_upsampler",
+                    local_files_only=local_files_only,
+                    dtype=dtype,
+                )
+            except (OSError, ValueError):
+                if profile.latent_upsampler_filename is None or profile.artifact_repo_id is None:
+                    raise
+                upsampler_path = resolve_ltx_artifact(
+                    model,
+                    profile.artifact_repo_id,
+                    profile.latent_upsampler_filename,
+                )
+                pipeline.latent_upsampler = _load_ltx_latent_upsampler_single_file(upsampler_path, dtype)
+        else:
+            if profile.latent_upsampler_filename is None or profile.artifact_repo_id is None:
+                raise FileNotFoundError(f"LTX latent upsampler component not found under {model}.")
+            upsampler_path = resolve_ltx_artifact(
                model,
-                "latent_upsampler",
-                local_files_only=local_files_only,
-                dtype=dtype,
+                profile.artifact_repo_id,
+                profile.latent_upsampler_filename,
            )
+            pipeline.latent_upsampler = _load_ltx_latent_upsampler_single_file(upsampler_path, dtype)

    transformer_config = load_transformer_config(model, "transformer", local_files_only)
    quant_config = getattr(od_config, "quantization_config", None)
@@ -412,6 +517,16 @@ def initialize_pipeline_components(pipeline: Any, od_config: Any) -> None:
        subfolder="scheduler",
        local_files_only=local_files_only,
    )
+    expected_checkpoint_kind: LTXCheckpointKind | None = None
+    if pipeline.pipeline_kind == "two_stage":
+        expected_checkpoint_kind = "regular"
+    elif pipeline.pipeline_kind == "distilled_two_stage":
+        expected_checkpoint_kind = "distilled"
+    validate_ltx_checkpoint(
+        pipeline.scheduler.config,
+        expected_kind=expected_checkpoint_kind,
+        pipeline_name=type(pipeline).__name__,
+    )

    pipeline.vae_spatial_compression_ratio = pipeline.vae.spatial_compression_ratio
    pipeline.vae_temporal_compression_ratio = pipeline.vae.temporal_compression_ratio
@@ -62,13 +62,14 @@ class LTXDenoiseContext:
    audio_latents: torch.Tensor
    video_coords: torch.Tensor
    audio_coords: torch.Tensor
+    audio_attention_mask: torch.Tensor | None = None
    conditioning_mask: torch.Tensor | None = None
    conditioning_mask_for_model: torch.Tensor | None = None


@dataclass
 class LTXPhaseResult:
-    """Denoised, unpacked AV latents and the context used to produce them."""
+    """Denoised AV latents and the context used to produce them."""

    forward_context: LTXForwardContext
    video: torch.Tensor
@@ -284,7 +285,8 @@ def prepare_rope_coords_stage(
    latents: torch.Tensor,
    audio_latents: torch.Tensor,
 ) -> tuple[torch.Tensor, torch.Tensor]:
-    video_coords = pipeline.transformer.rope.prepare_video_coords(
+    transformer = pipeline.transformer
+    video_coords = transformer.rope.prepare_video_coords(
        latents.shape[0],
        forward_ctx.latent_num_frames,
        forward_ctx.latent_height,
@@ -292,7 +294,7 @@ def prepare_rope_coords_stage(
        latents.device,
        fps=forward_ctx.request_inputs.frame_rate,
    )
-    audio_coords = pipeline.transformer.audio_rope.prepare_audio_coords(
+    audio_coords = transformer.audio_rope.prepare_audio_coords(
        audio_latents.shape[0],
        forward_ctx.padded_audio_num_frames,
        audio_latents.device,
@@ -331,6 +333,7 @@ def build_transformer_kwargs(
        # with a learned register, making all output context tokens valid.
        "encoder_attention_mask": None,
        "audio_encoder_attention_mask": None,
+        "audio_attention_mask": denoise_ctx.audio_attention_mask,
        "num_frames": forward_ctx.latent_num_frames,
        "height": forward_ctx.latent_height,
        "width": forward_ctx.latent_width,
@@ -442,11 +445,25 @@ class LTXPhaseExecutor:
            video_audio_step_adapter=video_audio_step_adapter,
        )
        video_coords, audio_coords = prepare_rope_coords_stage(pipeline, forward_ctx, latents, audio_latents)
+        ring_degree = getattr(pipeline.od_config.parallel_config, "ring_degree", 1) or 1
+        if padded_audio_num_frames > original_audio_num_frames and ring_degree > 1:
+            raise ValueError(
+                "LTX audio padding requires an attention mask, which Ring sequence parallelism does not support. "
+                "Use Ulysses-only SP or choose a request whose audio latent length is divisible by the SP size."
+            )
        denoise_ctx = LTXDenoiseContext(
            latents=latents,
            audio_latents=audio_latents,
            video_coords=video_coords,
            audio_coords=audio_coords,
+            audio_attention_mask=(
+                torch.arange(padded_audio_num_frames, device=audio_latents.device)
+                .lt(original_audio_num_frames)
+                .unsqueeze(0)
+                .expand(audio_latents.shape[0], -1)
+                if padded_audio_num_frames > original_audio_num_frames
+                else None
+            ),
            conditioning_mask=conditioning_mask,
        )
        denoise_ctx = pipeline._prepare_denoise_context_for_guidance(forward_ctx, denoise_ctx)
@@ -126,7 +126,10 @@ def x0_from_velocity(sample: torch.Tensor, velocity: torch.Tensor, sigma: torch.

 def velocity_from_x0(sample: torch.Tensor, x0: torch.Tensor, sigma: torch.Tensor) -> torch.Tensor:
    """Convert x0 back to velocity using the official fp32 arithmetic order."""
-    sigma = sigma.to(torch.float32)
+    # Official LTX converts the scalar scheduler sigma to a host value before
+    # division. A CUDA scalar tensor selects different kernel arithmetic and
+    # can cross bf16 rounding boundaries late in the denoise trajectory.
+    sigma = sigma.to(torch.float32).item()
    return ((sample.to(torch.float32) - x0.to(torch.float32)) / sigma).to(sample.dtype)


@@ -148,6 +151,7 @@ def combine_guided_x0(
    uncond_perturbed: torch.Tensor | float,
    uncond_modality: torch.Tensor | float,
    guidance: LTXModalityGuidance,
+    rescale_token_count: int | None = None,
 ) -> torch.Tensor:
    dtype = cond.dtype
    cond = cond.float()
@@ -163,9 +167,18 @@ def combine_guided_x0(
        # uses its full-tensor standard deviation. In Omni, dimension 0 may
        # instead contain independently scheduled requests. Reduce each item
        # separately so its result cannot depend on unrelated co-batched work.
-        reduce_dims = tuple(range(1, pred.ndim))
-        cond_std = cond.std(dim=reduce_dims, keepdim=True)
-        pred_std = pred.std(dim=reduce_dims, keepdim=True)
+        stats_cond = cond
+        stats_pred = pred
+        if rescale_token_count is not None:
+            if cond.ndim < 2 or not 0 < rescale_token_count <= cond.shape[1]:
+                raise ValueError(
+                    f"Guidance rescale token count must be in [1, {cond.shape[1]}], got {rescale_token_count}."
+                )
+            stats_cond = cond[:, :rescale_token_count]
+            stats_pred = pred[:, :rescale_token_count]
+        reduce_dims = tuple(range(1, stats_pred.ndim))
+        cond_std = stats_cond.std(dim=reduce_dims, keepdim=True)
+        pred_std = stats_pred.std(dim=reduce_dims, keepdim=True)
        eps = torch.finfo(pred.dtype).eps
        factor = cond_std / pred_std.clamp_min(eps)
        factor = torch.where(pred_std > eps, factor, torch.ones_like(factor))
@@ -175,6 +188,8 @@ def combine_guided_x0(


 def _repeat_batch(tensor: torch.Tensor, repeats: int) -> torch.Tensor:
+    if repeats == 1:
+        return tensor
    return tensor.repeat((repeats,) + (1,) * (tensor.ndim - 1))


@@ -278,6 +293,11 @@ class LTXGuidanceExecutor:
        if model_pass_count > 1:
            denoise_ctx.video_coords = _repeat_batch(denoise_ctx.video_coords, model_pass_count)
            denoise_ctx.audio_coords = _repeat_batch(denoise_ctx.audio_coords, model_pass_count)
+            if denoise_ctx.audio_attention_mask is not None:
+                denoise_ctx.audio_attention_mask = _repeat_batch(
+                    denoise_ctx.audio_attention_mask,
+                    model_pass_count,
+                )
        return denoise_ctx

    @staticmethod
@@ -322,6 +342,7 @@ class LTXGuidanceExecutor:
        guidance: LTXModalityGuidance,
        *,
        model_sigma: torch.Tensor | None = None,
+        rescale_token_count: int | None = None,
    ) -> torch.Tensor:
        model_sigma = sigma if model_sigma is None else model_sigma
        # Official guidance reduces contiguous BSC tensors. The fused
@@ -334,6 +355,7 @@ class LTXGuidanceExecutor:
            uncond_perturbed=x0.get("ptb", 0.0),
            uncond_modality=x0.get("mod", 0.0),
            guidance=guidance,
+            rescale_token_count=rescale_token_count,
        )
        return velocity_from_x0(sample, guided, sigma)

@@ -434,6 +456,7 @@ class LTXGuidanceExecutor:
                audio_splits,
                forward_ctx.audio_scheduler.sigmas[index],
                plan.spec.audio,
+                rescale_token_count=forward_ctx.original_audio_num_frames,
            ),
        )

@@ -509,6 +532,7 @@ class LTXGuidanceExecutor:
                audio_splits,
                forward_ctx.audio_scheduler.sigmas[index],
                plan.spec.audio,
+                rescale_token_count=forward_ctx.original_audio_num_frames,
            ),
        )

@@ -117,7 +117,9 @@ def create_noised_state(
        device=latents.device,
        dtype=latents.dtype,
    )
-    return noise_scale * noise + (1 - noise_scale) * latents
+    if isinstance(noise_scale, torch.Tensor):
+        noise_scale = noise_scale.float()
+    return torch.lerp(latents.float(), noise.float(), noise_scale).to(latents.dtype)


 def pack_audio_latents(latents: torch.Tensor) -> torch.Tensor:
@@ -135,6 +137,16 @@ def unpad_audio_latents(latents: torch.Tensor, num_frames: int) -> torch.Tensor:
    return latents[:, :num_frames]


+def clear_audio_padding(latents: torch.Tensor, num_frames: int) -> torch.Tensor:
+    """Keep Omni's SP-only audio padding outside the sampler state."""
+    if not 0 < num_frames <= latents.shape[1]:
+        raise ValueError(f"Audio frame count must be in [1, {latents.shape[1]}], got {num_frames}.")
+    if num_frames == latents.shape[1]:
+        return latents
+    padding = latents.new_zeros(latents.shape[0], latents.shape[1] - num_frames, latents.shape[2])
+    return torch.cat([latents[:, :num_frames], padding], dim=1)
+
+
 def get_sp_padded_audio_latent_length(audio_latent_length: int, sp_size: int) -> int:
    if sp_size > 1:
        audio_latent_length += (sp_size - (audio_latent_length % sp_size)) % sp_size
@@ -205,7 +217,12 @@ def prepare_video_latents(
            f"You have passed a list of generators of length {len(generator)}, but requested an effective batch"
            f" size of {batch_size}. Make sure the batch size matches the length of the generators."
        )
-    return randn_tensor(shape, generator=generator, device=device, dtype=dtype)
+    latents = randn_tensor(shape, generator=generator, device=device, dtype=dtype)
+    # Official LTX samples the 5D video latent first and then patchifies it,
+    # leaving token-major storage behind the [B, tokens, channels] view.  The
+    # logical values are unchanged, but preserving that layout is important:
+    # CUDA selects a different bf16 GEMM path for a contiguous token tensor.
+    return latents.transpose(1, 2).contiguous().transpose(1, 2)


 def prepare_audio_latents(
@@ -226,6 +243,16 @@ def prepare_audio_latents(
    sp_size = getattr(pipeline.od_config.parallel_config, "sequence_parallel_size", 1) or 1
    padded_latent_length = get_sp_padded_audio_latent_length(original_latent_length, int(sp_size))

+    def pad_logical_latents(logical_latents: torch.Tensor) -> torch.Tensor:
+        if padded_latent_length == original_latent_length:
+            return logical_latents
+        padding = logical_latents.new_zeros(
+            logical_latents.shape[0],
+            padded_latent_length - original_latent_length,
+            logical_latents.shape[2],
+        )
+        return torch.cat([logical_latents, padding], dim=1)
+
    if latents is not None:
        if latents.ndim == 4:
            latents = pack_audio_latents(latents)
@@ -237,29 +264,24 @@ def prepare_audio_latents(
                pipeline.audio_vae.latents_mean,
                pipeline.audio_vae.latents_std,
            )
-        latents = create_noised_state(latents, noise_scale, generator)

        if latents.shape[1] not in {original_latent_length, padded_latent_length}:
            raise ValueError(
                "Provided `audio_latents` has incompatible audio frame count "
                f"{latents.shape[1]}; expected {original_latent_length} or {padded_latent_length}."
            )
-        if latents.shape[1] == original_latent_length and padded_latent_length > original_latent_length:
-            padding = torch.zeros(
-                latents.shape[0],
-                padded_latent_length - original_latent_length,
-                latents.shape[2],
-                dtype=latents.dtype,
-                device=latents.device,
-            )
-            latents = torch.cat([latents, padding], dim=1)
+        # Padding is an Omni implementation detail, not part of the official
+        # audio state. Draw request RNG only for logical tokens so later phases
+        # remain seed-invariant across SP degrees, then append deterministic 0s.
+        latents = create_noised_state(latents[:, :original_latent_length], noise_scale, generator)
+        latents = pad_logical_latents(latents)
        return latents.to(device=device, dtype=dtype), original_latent_length, padded_latent_length

-    shape = (batch_size, padded_latent_length, num_channels_latents * latent_mel_bins)
+    shape = (batch_size, original_latent_length, num_channels_latents * latent_mel_bins)
    if isinstance(generator, list) and len(generator) != batch_size:
        raise ValueError(
            f"You have passed a list of generators of length {len(generator)}, but requested an effective batch"
            f" size of {batch_size}. Make sure the batch size matches the length of the generators."
        )
-    latents = randn_tensor(shape, generator=generator, device=device, dtype=dtype)
+    latents = pad_logical_latents(randn_tensor(shape, generator=generator, device=device, dtype=dtype))
    return latents, original_latent_length, padded_latent_length
@@ -0,0 +1,430 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""Construction and execution of a fixed LTX refinement-phase adapter.
+
+Checkpoint key parsing remains in :mod:`ltx2_adapter_parser`; this module
+resolves the declared artifact and makes the existing Transformer linears
+adapter-capable while retaining their quantization and tensor-parallel paths.
+"""
+
+from __future__ import annotations
+
+from collections import defaultdict
+from collections.abc import Iterable
+from dataclasses import dataclass
+from typing import Any
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+import vllm.envs as envs
+from vllm.distributed import (
+    split_tensor_along_last_dim,
+    tensor_model_parallel_all_reduce,
+)
+from vllm.logger import init_logger
+from vllm.model_executor.layers.batch_invariant import linear_batch_invariant
+from vllm.model_executor.layers.linear import (
+    ColumnParallelLinear,
+    QKVParallelLinear,
+    RowParallelLinear,
+    UnquantizedLinearMethod,
+)
+from vllm.model_executor.layers.utils import dispatch_unquantized_gemm
+from vllm.platforms import current_platform
+
+from .ltx2_adapter_parser import AdapterManifest, AdapterTarget, iter_adapter_tensors, parse_ltx_adapter
+from .ltx2_components import resolve_ltx_artifact
+from .ltx2_recipes import LTX_DISTILLED_ADAPTER_SLOT
+
+logger = init_logger(__name__)
+
+
+@dataclass(frozen=True)
+class _AdapterLayout:
+    """Rank-local storage and execution layout for one LoRA pair."""
+
+    kind: str
+    input_size: int
+    output_size: int
+    a_input_start: int
+    a_input_size: int
+    b_output_start: int
+    b_output_size: int
+    output_start: int
+    reduce_results: bool = False
+
+
+def _module_device(module: nn.Module) -> torch.device:
+    for tensor in (*module.parameters(), *module.buffers()):
+        if tensor.device.type != "meta":
+            return tensor.device
+    raise RuntimeError(f"Cannot materialize an LTX phase adapter for meta module {module._get_name()}.")
+
+
+def _build_layout(layer: nn.Module, target: AdapterTarget) -> _AdapterLayout:
+    if isinstance(layer, QKVParallelLinear):
+        if target.slice not in {"q", "k", "v"}:
+            raise ValueError(
+                f"LTX packed QKV target {target.module!r} requires a q/k/v adapter slice, got {target.slice!r}."
+            )
+        q_local = layer.num_heads * layer.head_size
+        k_local = layer.num_kv_heads * layer.head_size
+        v_local = layer.num_kv_heads * layer.v_head_size
+        if target.slice == "q":
+            output_size = layer.total_num_heads * layer.head_size
+            b_output_start = layer.tp_rank * q_local
+            b_output_size = q_local
+            output_start = 0
+        elif target.slice == "k":
+            output_size = layer.total_num_kv_heads * layer.head_size
+            shard_rank = layer.tp_rank // layer.num_kv_head_replicas
+            b_output_start = shard_rank * k_local
+            b_output_size = k_local
+            output_start = q_local
+        else:
+            output_size = layer.total_num_kv_heads * layer.v_head_size
+            shard_rank = layer.tp_rank // layer.num_kv_head_replicas
+            b_output_start = shard_rank * v_local
+            b_output_size = v_local
+            output_start = q_local + k_local
+        return _AdapterLayout(
+            kind="qkv",
+            input_size=layer.input_size,
+            output_size=output_size,
+            a_input_start=0,
+            a_input_size=layer.input_size,
+            b_output_start=b_output_start,
+            b_output_size=b_output_size,
+            output_start=output_start,
+        )
+
+    if isinstance(layer, ColumnParallelLinear):
+        if layer.gather_output:
+            raise ValueError(
+                f"LTX phase adapter does not support gathered ColumnParallelLinear target {target.module!r}."
+            )
+        return _AdapterLayout(
+            kind="column",
+            input_size=layer.input_size,
+            output_size=layer.output_size,
+            a_input_start=0,
+            a_input_size=layer.input_size,
+            b_output_start=layer.tp_rank * layer.output_size_per_partition,
+            b_output_size=layer.output_size_per_partition,
+            output_start=0,
+        )
+
+    if isinstance(layer, RowParallelLinear):
+        return _AdapterLayout(
+            kind="row",
+            input_size=layer.input_size,
+            output_size=layer.output_size,
+            a_input_start=layer.tp_rank * layer.input_size_per_partition,
+            a_input_size=layer.input_size_per_partition,
+            b_output_start=0,
+            b_output_size=layer.output_size,
+            output_start=0,
+            reduce_results=layer.reduce_results,
+        )
+
+    if isinstance(layer, nn.Linear):
+        return _AdapterLayout(
+            kind="replicated",
+            input_size=layer.in_features,
+            output_size=layer.out_features,
+            a_input_start=0,
+            a_input_size=layer.in_features,
+            b_output_start=0,
+            b_output_size=layer.out_features,
+            output_start=0,
+        )
+
+    raise TypeError(f"LTX phase adapter target {target.module!r} has unsupported layer type {type(layer).__name__}.")
+
+
+class _AdapterPiece(nn.Module):
+    """One rank-local A/B pair stored as non-persistent buffers."""
+
+    def __init__(self, target: AdapterTarget, layout: _AdapterLayout) -> None:
+        super().__init__()
+        self.target = target
+        self.layout = layout
+        # The fixed graph is installed before base checkpoint loading.  The
+        # empty buffers are replaced only after the base layer's quant/TP layout
+        # has settled and before offload/compile snapshots are taken.
+        self.register_buffer("lora_a", torch.empty(0), persistent=False)
+        self.register_buffer("lora_b", torch.empty(0), persistent=False)
+
+    def load(self, lora_a: torch.Tensor, lora_b: torch.Tensor, *, device: torch.device, dtype: torch.dtype) -> None:
+        if tuple(lora_a.shape) != self.target.a_shape or tuple(lora_b.shape) != self.target.b_shape:
+            raise ValueError(
+                f"LTX adapter source shape changed for {self.target.source_module!r}: "
+                f"A={tuple(lora_a.shape)}, B={tuple(lora_b.shape)}."
+            )
+        if lora_a.ndim != 2 or lora_b.ndim != 2:
+            raise ValueError(f"LTX adapter target {self.target.source_module!r} must use rank-2 A/B tensors.")
+        rank = self.target.a_shape[0]
+        if lora_a.shape[1] != self.layout.input_size or lora_b.shape[0] != self.layout.output_size:
+            raise ValueError(
+                f"LTX adapter shape does not match {self.target.module!r}: "
+                f"expected A=({rank}, {self.layout.input_size}), "
+                f"B=({self.layout.output_size}, {rank}); "
+                f"got A={tuple(lora_a.shape)}, B={tuple(lora_b.shape)}."
+            )
+
+        local_a = lora_a.narrow(1, self.layout.a_input_start, self.layout.a_input_size)
+        local_b = lora_b.narrow(0, self.layout.b_output_start, self.layout.b_output_size)
+        self.lora_a = local_a.to(device=device, dtype=dtype)
+        self.lora_b = local_b.to(device=device, dtype=dtype)
+
+    def delta(self, input_: torch.Tensor) -> torch.Tensor:
+        if self.lora_a.numel() == 0 or self.lora_b.numel() == 0:
+            raise RuntimeError(f"LTX phase adapter target {self.target.module!r} was not materialized.")
+        if self.layout.kind == "row" and input_.shape[-1] != self.layout.a_input_size:
+            input_ = input_.narrow(-1, self.layout.a_input_start, self.layout.a_input_size)
+        input_ = input_.to(dtype=self.lora_a.dtype)
+        delta = F.linear(F.linear(input_, self.lora_a), self.lora_b)
+        if self.layout.kind == "row" and self.layout.reduce_results:
+            delta = tensor_model_parallel_all_reduce(delta)
+        return delta
+
+    def weight_delta(self) -> torch.Tensor:
+        """Materialize the rank-local BF16 weight delta used by official fusion."""
+        if self.lora_a.numel() == 0 or self.lora_b.numel() == 0:
+            raise RuntimeError(f"LTX phase adapter target {self.target.module!r} was not materialized.")
+        return torch.matmul(self.lora_b, self.lora_a)
+
+
+def _apply_unquantized_gemm_with_weight(
+    layer: nn.Module,
+    input_: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor | None,
+) -> torch.Tensor:
+    if envs.VLLM_BATCH_INVARIANT and current_platform.is_cuda_alike():
+        return linear_batch_invariant(input_, weight, bias)
+    return dispatch_unquantized_gemm()(layer, input_, weight, bias)
+
+
+def _forward_with_weight(
+    layer: nn.Module,
+    input_: torch.Tensor,
+    weight: torch.Tensor,
+) -> torch.Tensor | tuple[torch.Tensor, torch.Tensor | None]:
+    """Run an unquantized base layer with an ephemeral fused weight."""
+    if isinstance(layer, nn.Linear):
+        return F.linear(input_, weight, layer.bias)
+
+    quant_method = getattr(layer, "quant_method", None)
+    if not isinstance(quant_method, UnquantizedLinearMethod):
+        raise ValueError("LTX layer-fused phase LoRA requires unquantized Transformer linears.")
+
+    if isinstance(layer, RowParallelLinear):
+        if layer.input_is_parallel:
+            input_parallel = input_
+        else:
+            input_parallel = split_tensor_along_last_dim(input_, num_partitions=layer.tp_size)[
+                layer.tp_rank
+            ].contiguous()
+        bias = None if layer.tp_rank > 0 or layer.skip_bias_add else layer.bias
+        output_parallel = _apply_unquantized_gemm_with_weight(layer, input_parallel, weight, bias)
+        if layer.reduce_results and layer.tp_size > 1:
+            output = tensor_model_parallel_all_reduce(output_parallel)
+        else:
+            output = output_parallel
+    elif isinstance(layer, ColumnParallelLinear):
+        bias = layer.bias if not layer.skip_bias_add else None
+        output = _apply_unquantized_gemm_with_weight(layer, input_, weight, bias)
+    else:
+        raise TypeError(f"Unsupported LTX layer-fused linear type {type(layer).__name__}.")
+
+    if not layer.return_bias:
+        return output
+    return output, layer.bias if layer.skip_bias_add else None
+
+
+class _PhaseAdapterLinear(nn.Module):
+    """Delegate the base linear unchanged and add an enabled adapter slot."""
+
+    def __init__(
+        self,
+        base_layer: nn.Module,
+        targets: Iterable[AdapterTarget],
+        *,
+        adapter_dtype: torch.dtype,
+        layer_fused: bool = False,
+    ) -> None:
+        super().__init__()
+        self.base_layer = base_layer
+        self.adapter_dtype = adapter_dtype
+        if layer_fused and not (
+            isinstance(base_layer, nn.Linear)
+            or isinstance(getattr(base_layer, "quant_method", None), UnquantizedLinearMethod)
+        ):
+            raise ValueError("LTX layer-fused phase LoRA requires unquantized Transformer linears.")
+        self.layer_fused = layer_fused
+        pieces = [_AdapterPiece(target, _build_layout(base_layer, target)) for target in targets]
+        self.pieces = nn.ModuleList(pieces)
+        self._pieces_by_target = {piece.target: piece for piece in pieces}
+        self._enabled = False
+
+    def load_target(self, target: AdapterTarget, lora_a: torch.Tensor, lora_b: torch.Tensor) -> None:
+        try:
+            piece = self._pieces_by_target[target]
+        except KeyError as exc:
+            raise ValueError(f"Adapter target {target.module!r} is not installed on this layer.") from exc
+        piece.load(lora_a, lora_b, device=_module_device(self.base_layer), dtype=self.adapter_dtype)
+
+    def set_enabled(self, enabled: bool) -> None:
+        self._enabled = enabled
+
+    def _add_adapter_delta(self, input_: torch.Tensor, output: torch.Tensor) -> torch.Tensor:
+        if not self._enabled:
+            return output
+        if output.requires_grad:
+            output = output.clone()
+        for piece in self.pieces:
+            delta = piece.delta(input_).to(dtype=output.dtype)
+            start = piece.layout.output_start
+            output[..., start : start + piece.layout.b_output_size].add_(delta)
+        return output
+
+    @torch.no_grad()
+    def _fused_weight(self) -> torch.Tensor:
+        base_weight = self.base_layer.weight
+        fused_weight = base_weight.clone()
+        for piece in self.pieces:
+            start = piece.layout.output_start
+            base_slice = base_weight.narrow(0, start, piece.layout.b_output_size)
+            delta = piece.weight_delta()
+            if delta.shape != base_slice.shape:
+                raise ValueError(
+                    f"LTX layer-fused delta shape {tuple(delta.shape)} does not match base weight slice "
+                    f"{tuple(base_slice.shape)} for {piece.target.module!r}."
+                )
+            # Match official fusion exactly: materialize B@A in the adapter
+            # dtype, then add the base tensor in-place before the main GEMM.
+            delta.add_(base_slice)
+            fused_weight.narrow(0, start, piece.layout.b_output_size).copy_(delta)
+        return fused_weight
+
+    def forward(self, input_: torch.Tensor):
+        if not self._enabled:
+            return self.base_layer(input_)
+        if self.layer_fused:
+            return _forward_with_weight(self.base_layer, input_, self._fused_weight())
+        output = self.base_layer(input_)
+        if isinstance(output, tuple):
+            return (self._add_adapter_delta(input_, output[0]), *output[1:])
+        return self._add_adapter_delta(input_, output)
+
+
+class LTXPhaseAdapterRuntime:
+    """Install and operate one fixed phase adapter on an existing transformer."""
+
+    def __init__(
+        self,
+        transformer: nn.Module,
+        manifest: AdapterManifest,
+        *,
+        dtype: torch.dtype,
+        layer_fused: bool = False,
+    ) -> None:
+        self.transformer = transformer
+        self.manifest = manifest
+        self._wrappers: dict[str, _PhaseAdapterLinear] = {}
+        self._target_names: tuple[str, ...] = ()
+        self._materialized = False
+        grouped: dict[str, list[AdapterTarget]] = defaultdict(list)
+        for target in manifest.targets:
+            grouped[target.module].append(target)
+
+        original_layers = {name: self.transformer.get_submodule(name) for name in grouped}
+        for name, targets in grouped.items():
+            wrapper = _PhaseAdapterLinear(
+                original_layers[name],
+                targets,
+                adapter_dtype=dtype,
+                layer_fused=layer_fused,
+            )
+            parent_name, _, child_name = name.rpartition(".")
+            parent = self.transformer.get_submodule(parent_name) if parent_name else self.transformer
+            if child_name not in parent._modules:
+                raise ValueError(f"LTX phase adapter target {name!r} is no longer a direct module child.")
+            parent._modules[child_name] = wrapper
+            self._wrappers[name] = wrapper
+
+        self._target_names = tuple(sorted(self._wrappers, key=len, reverse=True))
+        # LTX's custom loader receives the original checkpoint names.  Keep the
+        # adaptation local to that loader rather than changing AutoWeightsLoader.
+        self.transformer._phase_adapter_parameter_name = self.base_parameter_name
+        logger.info(
+            "Installed %d LTX %s phase-adapter linears",
+            len(self._wrappers),
+            "layer-fused" if layer_fused else "dynamic",
+        )
+
+    def base_parameter_name(self, name: str) -> str:
+        for target_name in self._target_names:
+            prefix = f"{target_name}."
+            if name.startswith(prefix):
+                return f"{target_name}.base_layer.{name[len(prefix) :]}"
+        return name
+
+    def finalize(self) -> None:
+        """Stream rank-local A/B buffers after base weight processing completes."""
+        if self._materialized:
+            return
+        for target, lora_a, lora_b in iter_adapter_tensors(self.manifest):
+            self._wrappers[target.module].load_target(target, lora_a, lora_b)
+        self._materialized = True
+        logger.info("Materialized %d LTX phase-adapter tensor pairs", len(self.manifest.targets))
+
+    def activate(self, adapter_slot: str | None) -> None:
+        if adapter_slot is not None:
+            if adapter_slot != LTX_DISTILLED_ADAPTER_SLOT:
+                raise ValueError(f"Unknown LTX phase adapter slot {adapter_slot!r}.")
+            if not self._materialized:
+                raise RuntimeError("LTX phase adapter data must be materialized before activation.")
+        for wrapper in self._wrappers.values():
+            wrapper.set_enabled(adapter_slot is not None)
+
+
+def build_ltx_phase_adapter(pipeline: Any) -> LTXPhaseAdapterRuntime | None:
+    """Build the fixed adapter required by a multi-phase LTX recipe."""
+    adapter_slots = {phase.adapter_slot for phase in pipeline.pipeline_recipe.phases if phase.adapter_slot is not None}
+    if not adapter_slots:
+        return None
+    if adapter_slots != {LTX_DISTILLED_ADAPTER_SLOT}:
+        raise ValueError(f"Unsupported LTX phase adapter slots: {sorted(adapter_slots)}.")
+
+    profile = pipeline.component_profile
+    if profile.distilled_lora_filename is None or profile.artifact_repo_id is None:
+        raise ValueError(f"{profile.name} does not declare the required distilled adapter artifact.")
+    if getattr(pipeline.od_config, "lora_path", None) is not None:
+        raise ValueError(
+            f"{pipeline.__class__.__name__} reserves LoRA execution for its phase adapter; "
+            "request or static LoRA composition is not supported yet."
+        )
+
+    quantization_config = getattr(pipeline.od_config, "quantization_config", None)
+    layer_fused = quantization_config is None
+    adapter_path = resolve_ltx_artifact(
+        pipeline.od_config.model,
+        profile.artifact_repo_id,
+        profile.distilled_lora_filename,
+    )
+
+    if layer_fused and pipeline.od_config.dtype is not torch.bfloat16:
+        raise ValueError("LTX layer-fused phase LoRA requires a bfloat16 pipeline dtype.")
+
+    manifest = parse_ltx_adapter(pipeline.transformer, adapter_path)
+    runtime = LTXPhaseAdapterRuntime(
+        pipeline.transformer,
+        manifest,
+        dtype=pipeline.od_config.dtype,
+        layer_fused=layer_fused,
+    )
+    return runtime
@@ -24,6 +24,7 @@ LTX_DEFAULT_NEGATIVE_PROMPT = (

 LTX_DISTILLED_SIGMAS = (1.0, 0.99375, 0.9875, 0.98125, 0.975, 0.909375, 0.725, 0.421875, 0.0)
 LTX_STAGE_2_DISTILLED_SIGMAS = (0.909375, 0.725, 0.421875, 0.0)
+LTX_DISTILLED_ADAPTER_SLOT = "ltx_distilled"


@dataclass(frozen=True)
@@ -36,6 +37,7 @@ class LTXPhaseRecipe:
    sigmas: tuple[float, ...] | None = None
    noise_scale: float = 0.0
    input_transform: Literal["initial", "spatial_upsample"] = "initial"
+    adapter_slot: str | None = None
    allow_guidance_override: bool = True
    use_official_sigma_schedule: bool = True

@@ -157,11 +159,59 @@ LTX2_DISTILLED_TWO_STAGE_RECIPE = LTXPipelineRecipe(
    fixed_num_inference_steps=True,
 )

+# LTX-2.3 full-distilled uses the same official fixed 8 + 3 sigma schedules
+# and request contract. Its distinct component profile selects the 22B merged
+# distilled Transformer, BWE vocoder, and matching spatial upsampler; no
+# distilled LoRA is loaded for either phase.
+LTX23_DISTILLED_TWO_STAGE_RECIPE = LTX2_DISTILLED_TWO_STAGE_RECIPE
+
+
+def _official_two_stage_recipe(one_stage_recipe: LTXPipelineRecipe) -> LTXPipelineRecipe:
+    return LTXPipelineRecipe(
+        height=one_stage_recipe.height * 2,
+        width=one_stage_recipe.width * 2,
+        num_frames=one_stage_recipe.num_frames,
+        frame_rate=one_stage_recipe.frame_rate,
+        num_inference_steps=one_stage_recipe.num_inference_steps,
+        negative_prompt=one_stage_recipe.negative_prompt,
+        phases=(
+            LTXPhaseRecipe(
+                name="generate_lowres",
+                guidance=one_stage_recipe.request_guidance,
+                spatial_downscale=2,
+                noise_scale=1.0,
+            ),
+            LTXPhaseRecipe(
+                name="refine",
+                guidance=LTXGuidanceSpec.positive_only(),
+                sigmas=LTX_STAGE_2_DISTILLED_SIGMAS,
+                noise_scale=LTX_STAGE_2_DISTILLED_SIGMAS[0],
+                input_transform="spatial_upsample",
+                adapter_slot=LTX_DISTILLED_ADAPTER_SLOT,
+                allow_guidance_override=False,
+                use_official_sigma_schedule=False,
+            ),
+        ),
+        video_output_phase=1,
+        # The official second stage refines video only and deliberately
+        # discards its audio result. Decode the full-context Stage-1 audio.
+        audio_output_phase=0,
+        allow_request_sigmas=False,
+        allow_request_latents=False,
+    )
+
+
+LTX2_TWO_STAGE_RECIPE = _official_two_stage_recipe(LTX2_ONE_STAGE_RECIPE)
+LTX23_TWO_STAGE_RECIPE = _official_two_stage_recipe(LTX23_ONE_STAGE_RECIPE)
+

 _PIPELINE_RECIPES: dict[tuple[str, str], LTXPipelineRecipe] = {
    ("one_stage", "2"): LTX2_ONE_STAGE_RECIPE,
    ("one_stage", "2.3"): LTX23_ONE_STAGE_RECIPE,
+    ("two_stage", "2"): LTX2_TWO_STAGE_RECIPE,
+    ("two_stage", "2.3"): LTX23_TWO_STAGE_RECIPE,
    ("distilled_two_stage", "2"): LTX2_DISTILLED_TWO_STAGE_RECIPE,
+    ("distilled_two_stage", "2.3"): LTX23_DISTILLED_TWO_STAGE_RECIPE,
    ("dmd2", "2"): LTX_POSITIVE_ONLY_RECIPE,
    ("dmd2", "2.3"): LTX_POSITIVE_ONLY_RECIPE,
 }
@@ -5,8 +5,9 @@

 from __future__ import annotations

+from collections.abc import Mapping
 from dataclasses import dataclass, replace
-from typing import Any
+from typing import Any, Literal

 import torch

@@ -49,6 +50,34 @@ class LTXRequestInputs:
        return self.guidance.video.cfg_scale


+LTXCheckpointKind = Literal["distilled", "regular"]
+
+
+def validate_ltx_checkpoint(
+    scheduler_config: Mapping[str, Any],
+    *,
+    expected_kind: LTXCheckpointKind | None,
+    pipeline_name: str,
+) -> None:
+    """Ensure scheduler metadata matches the checkpoint required by a profile."""
+    if expected_kind is None:
+        return
+
+    if expected_kind == "distilled":
+        if scheduler_config.get("use_dynamic_shifting", True) or scheduler_config.get("shift_terminal") is not None:
+            raise ValueError(
+                f"{pipeline_name} requires a merged distilled LTX2 checkpoint; "
+                "the loaded scheduler metadata describes a non-distilled checkpoint."
+            )
+        return
+
+    if not scheduler_config.get("use_dynamic_shifting", False) or scheduler_config.get("shift_terminal") is None:
+        raise ValueError(
+            f"{pipeline_name} requires a regular non-distilled LTX checkpoint; "
+            "the loaded scheduler metadata describes a distilled checkpoint."
+        )
+
+
 def validate_pipeline_request(
    request_inputs: LTXRequestInputs,
    *,
Author	SHA1	Message	Date
Hongsheng Liu	7ec78360bb	Merge branch 'main' into codex-split-attention-offload-guides	2026-08-15 07:00:22 +08:00
Joshna-Medisetty Cursor Chendi.Xue	4b6d6d5cf7	[Bugfix][XPU][Tests]Scope XPU pytest to explicit paths to fix collection crash (#6175 ) Signed-off-by: Joshna Medisetty <joshna.medisetty@intel.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Chendi.Xue <chendi.xue@intel.com>	2026-08-14 15:23:41 -05:00
Alex Brooks	f7a5146093	[Feat/Bugfix] Fix API Server model tag <-> model sync & flag normalization (#3805 ) Signed-off-by: Alex Brooks <albrooks@redhat.com>	2026-08-14 19:13:00 +00:00
Wenjia Li	8ecd1f6d5c	[Perf][NPU] Fuse MiniMax H3 Qwen3-VL SwiGLU (#6167 ) Signed-off-by: Wenjia Li <wjialish@gmail.com>	2026-08-14 11:47:17 +00:00
Wenjia Li	ffd11c181e	[Perf][NPU] Fuse MiniMax H3 Qwen3-VL RoPE (#6061 ) Signed-off-by: Wenjia Li <wjialish@gmail.com>	2026-08-14 09:49:26 +00:00
R0CKSTAR	728640f4d9	[Bugfix][MUSA][Qwen Image] Avoid complex RoPE alias guards (#6110 ) Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>	2026-08-14 15:51:23 +08:00
R0CKSTAR	022d7a7695	[Perf][Diffusion] Avoid redundant MiniMax-H3 reference video scans (#6064 ) Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>	2026-08-14 15:49:26 +08:00
baonudesifeizhai	b300f66656	[Bugfix] Fix HSDP compatibility with the new online FP8 linear method (#5677 ) Signed-off-by: roG0d <baonudesifeizhai@gmail.com>	2026-08-14 15:45:55 +08:00
TJianandHongsheng Liu	bc73546c06	[ROCm] [CI] Migrate CI to mi300x for v0.27.x (#5886 ) Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com> Signed-off-by: TJian <tunjian.tan@embeddedllm.com> Co-authored-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-08-14 15:38:47 +08:00
Shenglei Fu Claude Sonnet 5 Alicia wangyu Hongsheng Liu	2b08419d8b	[CI] Per-model, per-entry-mode code coverage collection (#5593 ) Signed-off-by: Shenglei Fu <sfu@confluent.io> Signed-off-by: Shenglei Fu <117230642+ShengleiFu@users.noreply.github.com> Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com> Co-authored-by: Alicia <115451386+congw729@users.noreply.github.com> Co-authored-by: wangyu <53896905+yenuo26@users.noreply.github.com> Co-authored-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-08-14 15:36:56 +08:00
zhengjiaandCursor	5d50bdc2ab	[Perf] Cap diffusion worker thread count before spawn (#6165 ) Signed-off-by: zjli2013 <leezhengjiang@126.com> Co-authored-by: Cursor <cursoragent@cursor.com>	2026-08-14 12:14:55 +08:00
Samit	8a5bf2fb3f	[doc]: Update latest news section for verl-omni release (#6187 ) Signed-off-by: Samit <285365963@qq.com>	2026-08-14 12:04:15 +08:00
Bo Li lishunyang12 Claude Opus 5	596c16a550	[Kernel] Fuse Q/K RMSNorm and RoPE (#5990 ) Signed-off-by: Bo Li <22713281+bobboli@users.noreply.github.com> Signed-off-by: lishunyang12 <lishunyang12@163.com> Signed-off-by: Mengjie Zhao <zmj0129@gmail.com> Co-authored-by: lishunyang12 <lishunyang12@163.com> Co-authored-by: Claude Opus 5 <noreply@anthropic.com>	2026-08-14 00:46:59 +00:00
Lucas(Ruijie)	90a08c0eda	[Model][TTS][WIP] Add dots.tts (rednote-hilab) — continuous-AR 48kHz TTS (#4765 ) Signed-off-by: Moore-Z <zhumoore9661996@gmail.com>	2026-08-13 17:58:33 +00:00
Sy03	dbc0dd6d85	[Bugfix][Higgs-Audio-V3] Disable XQA decode (#6068 ) Signed-off-by: Sy03 <1370724210@qq.com> Signed-off-by: Sysy <1370724210@qq.com>	2026-08-13 10:00:50 -07:00
Mu GuanLin	83d1f11449	[Feat] LTX Standard Two-Stage Pipeline (#5500 ) Signed-off-by: mglyn <1203789601@qq.com>	2026-08-14 00:13:14 +08:00
Hongsheng Liu	5569abf0a2	Update docs/design/feature/offloader/README.md Signed-off-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-08-13 23:37:08 +08:00
Hongsheng Liu	10853fb4f3	Update docs/design/feature/offloader/README.md Signed-off-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-08-13 23:37:01 +08:00
Hongsheng Liu	a52c19ea8b	docs: organize offloader feature designs Signed-off-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-08-12 11:35:07 +08:00
Hongsheng Liu	8b61499662	docs: rename offloader guides Signed-off-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-08-12 09:30:34 +08:00
Hongsheng Liu	7872663bf0	docs: split diffusion feature guides Signed-off-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-08-12 08:55:41 +08:00