Merge branch 'main' into codex-diffusion-shutdown-escalation

[Rebase] Rebase to vLLM 0.29.0 (#7230 )
Signed-off-by: tzhouam <tzhouam@connect.ust.hk> Signed-off-by: Zhou Taichang <tzhouam@connect.ust.hk> Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
2026-09-10 13:27:51 +08:00 · 2026-09-10 12:08:28 +08:00 · 2026-09-10 11:38:11 +08:00 · 2026-09-10 03:07:35 +00:00 · 2026-09-10 02:16:14 +00:00 · 2026-09-10 10:09:39 +08:00
222 changed files with 11676 additions and 5113 deletions
@@ -14,6 +14,8 @@ set -euo pipefail
 # shellcheck disable=SC1091
 source .buildkite/common/scripts/resolve_skip_ci.sh

+AMD_SUITE_SELECTOR=".buildkite/amd/scripts/select_test_suites.py"
+
 if [[ -z "${RUN_ALL:-}" ]]; then
    RUN_ALL=0
 fi
@@ -30,31 +32,20 @@ if [[ -z "${AMD_MIRROR_HW:-}" ]]; then
    AMD_MIRROR_HW="amdproduction"
 fi

-fail_fast() {
-    DISABLE_LABEL="ci-no-fail-fast"
-    # If BUILDKITE_PULL_REQUEST != "false", then we check the PR labels using curl and jq
-    if [ "$BUILDKITE_PULL_REQUEST" != "false" ]; then
-        PR_LABELS=$(curl -s "https://api.github.com/repos/vllm-project/vllm-omni/pulls/$BUILDKITE_PULL_REQUEST" | jq -r '.labels[].name')
-        if [[ $PR_LABELS == *"$DISABLE_LABEL"* ]]; then
-            echo false
-        else
-            echo true
-        fi
-    else
-        echo false  # not a PR or BUILDKITE_PULL_REQUEST not set
+PR_LABELS=""
+# Explicit debug suites do not depend on PR labels.
+if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" && -z "${DEBUG_TEST_YAML:-}" ]]; then
+    if ! PR_LABELS=$(curl -fsSL \
+        "https://api.github.com/repos/vllm-project/vllm-omni/pulls/$BUILDKITE_PULL_REQUEST" \
+        | jq -r '.labels[].name'); then
+        echo "ERROR: Could not read PR labels; refusing to select an incorrect AMD test tier." >&2
+        exit 1
    fi
-}
+fi

-check_run_all_label() {
-    RUN_ALL_LABEL="ready-run-all-tests"
-    # If BUILDKITE_PULL_REQUEST != "false", then we check the PR labels using curl and jq
-    if [ "$BUILDKITE_PULL_REQUEST" != "false" ]; then
-        PR_LABELS=$(curl -s "https://api.github.com/repos/vllm-project/vllm-omni/pulls/$BUILDKITE_PULL_REQUEST" | jq -r '.labels[].name')
-        if [[ $PR_LABELS == *"$RUN_ALL_LABEL"* ]]; then
-            echo true
-        else
-            echo false
-        fi
+fail_fast() {
+    if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" ]]; then
+        echo true
    else
        echo false  # not a PR or BUILDKITE_PULL_REQUEST not set
    fi
@@ -64,6 +55,57 @@ if [[ -z "${COV_ENABLED:-}" ]]; then
    COV_ENABLED=0
 fi

+resolve_test_specs() {
+    local selected_specs
+    if ! selected_specs=$(python3 "$AMD_SUITE_SELECTOR" \
+        --branch "$BUILDKITE_BRANCH" \
+        --labels "$PR_LABELS" \
+        --debug-test-yaml "${DEBUG_TEST_YAML:-}"); then
+        exit 1
+    fi
+
+    TEST_SPECS=()
+    while IFS= read -r suite_spec; do
+        [[ -z "$suite_spec" ]] && continue
+        TEST_SPECS+=("$suite_spec")
+    done <<< "$selected_specs"
+}
+
+filter_test_specs_by_skip_ci() {
+    local suite_spec level decision
+    local skip_all=0
+    local -a runnable_specs=()
+
+    for suite_spec in "${TEST_SPECS[@]}"; do
+        case "$suite_spec" in
+            READY_TESTS:*) level="l2" ;;
+            MERGE_TESTS:*) level="l3" ;;
+            *)
+                echo "ERROR: unknown AMD test suite spec '$suite_spec'" >&2
+                exit 1
+                ;;
+        esac
+
+        if decision=$(python3 "$SKIP_CI_PY" gate amd "$level"); then
+            echo "Skipping AMD $level suite due to $decision decision."
+            [[ "$decision" == "skip-all" ]] && skip_all=1
+        else
+            runnable_specs+=("$suite_spec")
+        fi
+    done
+
+    TEST_SPECS=("${runnable_specs[@]}")
+    if [[ ${#TEST_SPECS[@]} -eq 0 ]]; then
+        if [[ $skip_all -eq 1 ]]; then
+            buildkite-agent annotate \
+                ":memo: CI skipped — docs or pytest skip-mark changes only" \
+                --style "info" 2>/dev/null || true
+        fi
+        echo "No AMD L2/L3 suites remain after skip-ci filtering."
+        exit 0
+    fi
+}
+
 upload_pipeline() {
    echo "Uploading pipeline..."
    # Install minijinja
@@ -84,7 +126,7 @@ upload_pipeline() {
    # (WIP) Use pipeline generator instead of jinja template
    if [ -e ".buildkite/amd/pipeline_generator/pipeline_generator.py" ]; then
        python -m pip install click pydantic
-        python .buildkite/amd/pipeline_generator/pipeline_generator.py --run_all=$RUN_ALL --list_file_diff="$LIST_FILE_DIFF" --nightly="$NIGHTLY" --mirror_hw="$AMD_MIRROR_HW"
+        python .buildkite/amd/pipeline_generator/pipeline_generator.py --run_all="$RUN_ALL" --list_file_diff="$LIST_FILE_DIFF" --nightly="$NIGHTLY" --mirror_hw="$AMD_MIRROR_HW"
        buildkite-agent pipeline upload .buildkite/amd/pipeline.yaml
        exit 0
    fi
@@ -97,49 +139,12 @@ upload_pipeline() {

    cd .buildkite/amd

-    # Select test definition file: merge suite for main, ready suite for PRs.
-    # For debugging, DEBUG_TEST_YAML accepts a comma-separated list containing
-    # "merge" and/or "ready" (case-insensitive). Multiple suites are combined
-    # before rendering so they share one amd-build step.
-    if [[ -n "${DEBUG_TEST_YAML:-}" ]]; then
-        declare -a DEBUG_TEST_SPECS=()
-        declare -A SEEN_DEBUG_TESTS=()
-        IFS=',' read -ra REQUESTED_DEBUG_TESTS <<< "${DEBUG_TEST_YAML,,}"
-
-        for requested_test in "${REQUESTED_DEBUG_TESTS[@]}"; do
-            # Permit readable values such as "ready, merge," and ignore the
-            # empty item produced by a trailing comma.
-            requested_test="${requested_test//[[:space:]]/}"
-            [[ -z "$requested_test" ]] && continue
-
-            if [[ -n "${SEEN_DEBUG_TESTS[$requested_test]:-}" ]]; then
-                echo "ERROR: duplicate DEBUG_TEST_YAML suite '$requested_test'" >&2
-                exit 1
-            fi
-            SEEN_DEBUG_TESTS[$requested_test]=1
-
-            case "$requested_test" in
-                ready)
-                    DEBUG_TEST_SPECS+=("READY_TESTS:test-amd-ready.yml")
-                    ;;
-                merge)
-                    DEBUG_TEST_SPECS+=("MERGE_TESTS:test-amd-merge.yml")
-                    ;;
-                *)
-                    echo "ERROR: DEBUG_TEST_YAML entries must be 'merge' or 'ready', got '$requested_test'" >&2
-                    exit 1
-                    ;;
-            esac
-        done
-
-        if [[ ${#DEBUG_TEST_SPECS[@]} -eq 0 ]]; then
-            echo "ERROR: DEBUG_TEST_YAML did not contain a test suite" >&2
-            exit 1
-        elif [[ ${#DEBUG_TEST_SPECS[@]} -eq 1 ]]; then
-            TEST_YAML="${DEBUG_TEST_SPECS[0]#*:}"
-        else
-            TEST_YAML=$(mktemp "${TMPDIR:-/tmp}/amd-debug-tests.XXXXXX.yml")
-            python - "$TEST_YAML" "${DEBUG_TEST_SPECS[@]}" <<'PY'
+    # Multiple label-selected or debug-selected suites share one image build.
+    if [[ ${#TEST_SPECS[@]} -eq 1 ]]; then
+        TEST_YAML="${TEST_SPECS[0]#*:}"
+    else
+        TEST_YAML=$(mktemp "${TMPDIR:-/tmp}/amd-selected-tests.XXXXXX.yml")
+        python - "$TEST_YAML" "${TEST_SPECS[@]}" <<'PY'
 import sys

 import yaml
@@ -172,13 +177,8 @@ for suite_spec in suite_specs:
 with open(output_path, "w", encoding="utf-8") as output_file:
    yaml.safe_dump(combined, output_file, sort_keys=False)
 PY
-        fi
-        echo "DEBUG_TEST_YAML override: using ${DEBUG_TEST_SPECS[*]}"
-    elif [[ $BUILDKITE_BRANCH == "main" ]]; then
-        TEST_YAML="test-amd-merge.yml"
-    else
-        TEST_YAML="test-amd-ready.yml"
    fi
+    echo "AMD test suites: ${TEST_SPECS[*]}"

    (
        set -x
@@ -198,7 +198,7 @@ PY
            > pipeline.yaml
    )
    cat pipeline.yaml
-    if [[ "$TEST_YAML" == "${TMPDIR:-/tmp}/amd-debug-tests."*.yml ]]; then
+    if [[ "$TEST_YAML" == "${TMPDIR:-/tmp}/amd-selected-tests."*.yml ]]; then
        rm -f -- "$TEST_YAML"
    fi
    buildkite-agent artifact upload pipeline.yaml
@@ -219,12 +219,10 @@ if [[ $BUILDKITE_BRANCH == "main" ]]; then
    file_diff=$(get_diff_main)
 fi

-# Early exit: unified skip-ci (docs / skip marks) and CI-yaml-only level targeting.
-if [[ $BUILDKITE_BRANCH == "main" ]]; then
-    gate_bootstrap_ci amd l3
-else
-    gate_bootstrap_ci amd l2
-fi
+# Resolve PR tier labels before skip-ci so ready and merge suites can be
+# filtered independently for CI-YAML-only changes.
+resolve_test_specs
+filter_test_specs_by_skip_ci

 patterns=(
    "docker/Dockerfile"
@@ -274,14 +272,6 @@ for file in $file_diff; do
    fi
 done

-# Check for ready-run-all-tests label
-LABEL_RUN_ALL=$(check_run_all_label)
-if [[ $LABEL_RUN_ALL == true ]]; then
-    RUN_ALL=1
-    NIGHTLY=1
-    echo "Found 'ready-run-all-tests' label. Running all tests including optional tests."
-fi
-
 # Decide whether to use precompiled wheels
 # Relies on existing patterns array as a basis.
 if [[ -n "${VLLM_USE_PRECOMPILED:-}" ]]; then
@@ -0,0 +1,93 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+"""Select AMD L2/L3 suites after Buildkite has created an AMD build.
+
+This repository-side selector does not decide which GitHub label events start
+the external ``vllm-omni-amd-ci`` pipeline. Its Buildkite PR build condition
+must admit the relevant label event before this code can inspect the labels.
+"""
+
+from __future__ import annotations
+
+import argparse
+import sys
+from collections.abc import Iterable
+
+SUITE_SPECS = {
+    "ready": "READY_TESTS:test-amd-ready.yml",
+    "merge": "MERGE_TESTS:test-amd-merge.yml",
+}
+
+
+def _parse_debug_suites(value: str) -> tuple[str, ...]:
+    suites = tuple(item.strip().lower() for item in value.split(",") if item.strip())
+    if not suites:
+        raise ValueError("DEBUG_TEST_YAML did not contain a test suite")
+
+    seen: set[str] = set()
+    for suite in suites:
+        if suite not in SUITE_SPECS:
+            raise ValueError(f"DEBUG_TEST_YAML entries must be 'merge' or 'ready', got {suite!r}")
+        if suite in seen:
+            raise ValueError(f"duplicate DEBUG_TEST_YAML suite {suite!r}")
+        seen.add(suite)
+    return suites
+
+
+def select_amd_test_suites(
+    *,
+    branch: str,
+    labels: Iterable[str],
+    debug_test_yaml: str = "",
+) -> tuple[str, ...]:
+    """Return ordered suite names while preserving legacy unlabeled PRs."""
+
+    if debug_test_yaml:
+        return _parse_debug_suites(debug_test_yaml)
+    if branch == "main":
+        return ("merge",)
+
+    label_set = {label.strip() for label in labels if label.strip()}
+    selected = tuple(suite for label, suite in (("ready", "ready"), ("merge-test", "merge")) if label in label_set)
+    # AMD historically uploaded L2 for every PR build. Keep that behavior for
+    # triggers such as amd-test while adding tier selection for ready/merge-test.
+    return selected or ("ready",)
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--branch", required=True)
+    parser.add_argument("--labels", default="")
+    parser.add_argument("--debug-test-yaml", default="")
+    args = parser.parse_args()
+
+    labels = tuple(args.labels.splitlines())
+    try:
+        suites = select_amd_test_suites(
+            branch=args.branch,
+            labels=labels,
+            debug_test_yaml=args.debug_test_yaml,
+        )
+    except ValueError as exc:
+        parser.error(str(exc))
+
+    label_set = {label.strip() for label in labels if label.strip()}
+    if not args.debug_test_yaml and args.branch != "main" and not ({"ready", "merge-test"} & label_set):
+        print(
+            "No AMD L2/L3 tier label found; preserving the legacy ready-suite fallback.",
+            file=sys.stderr,
+        )
+    if "nightly-test" in label_set:
+        print(
+            "AMD nightly-test has no suite yet and does not affect L2/L3 selection.",
+            file=sys.stderr,
+        )
+
+    for suite in suites:
+        print(SUITE_SPECS[suite])
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
@@ -17,6 +17,8 @@
 #   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 (A3 CI image tag …:$BUILDKITE_COMMIT-a3)
+#   a5_npu_2    — ascend-a5, resource_class npu-2
+#   310p_npu_1  — ascend-310p, resource_class npu-1
 #
 # B200 (b200-k8s): matches vLLM device ``b200-k8s`` from ci-infra
 # (AgentQueue.B200_K8S + b200_plugin_template). Queue routes the K8s pool;
@@ -26,6 +28,9 @@
 # packed onto g6.12xlarge L4 nodes (``vllm.ci/gpu-pool=l4x4``). GPU / CPU /
 # memory / shm requests scale with card count (1 GPU = 10 CPU / 40Gi / 24Gi shm).
 # Retry covers lost agents / killed pods, not pytest or OOM (exit 1 / 137).
+# The Kubernetes HF secret deliberately uses VLLM_CI_HF_TOKEN. Scheduled
+# build-level HF_TOKEN values can override pod env entries with the same name;
+# upload_pipeline.py restores the secret-backed value inside the job shell.

 x-npu-k8s-plugins: &npu_k8s_plugins
  - kubernetes:
@@ -84,7 +89,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /fsx/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -135,7 +140,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /fsx/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -186,7 +191,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /fsx/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -237,7 +242,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /fsx/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -271,7 +276,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /root/.cache/huggingface
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -306,7 +311,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /root/.cache/huggingface
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -341,7 +346,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /root/.cache/huggingface
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -376,7 +381,7 @@ mirror_hardwares:
                env:
                  - name: HF_HOME
                    value: /root/.cache/huggingface
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -426,7 +431,7 @@ mirror_hardwares:
                    value: "0"
                  - name: HF_HOME
                    value: /raid/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -477,7 +482,7 @@ mirror_hardwares:
                    value: "0"
                  - name: HF_HOME
                    value: /raid/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -528,7 +533,7 @@ mirror_hardwares:
                    value: "0"
                  - name: HF_HOME
                    value: /raid/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -579,7 +584,7 @@ mirror_hardwares:
                    value: "0"
                  - name: HF_HOME
                    value: /raid/hf_cache
-                  - name: HF_TOKEN
+                  - name: VLLM_CI_HF_TOKEN
                    valueFrom:
                      secretKeyRef:
                        name: hf-token-secret
@@ -646,3 +651,52 @@ mirror_hardwares:
      resource_class: "npu-16"
    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-a3"
    plugins: *npu_k8s_plugins
+
+  a5_npu_2:
+    agents:
+      queue: "ascend-a5"
+      resource_class: "npu-2"
+    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-a5"
+    plugins: *npu_k8s_plugins
+
+  a5_npu_4:
+    agents:
+      queue: "ascend-a5"
+      resource_class: "npu-4"
+    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-a5"
+    plugins: *npu_k8s_plugins
+
+  a5_npu_8:
+    agents:
+      queue: "ascend-a5"
+      resource_class: "npu-8"
+    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-a5"
+    plugins: *npu_k8s_plugins
+
+  a5_npu_16:
+    agents:
+      queue: "ascend-a5"
+      resource_class: "npu-16"
+    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-a5"
+    plugins: *npu_k8s_plugins
+
+  310p_npu_1:
+    agents:
+      queue: "ascend-310p"
+      resource_class: "npu-1"
+    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-310p"
+    plugins: *npu_k8s_plugins
+
+  310p_npu_2:
+    agents:
+      queue: "ascend-310p"
+      resource_class: "npu-2"
+    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-310p"
+    plugins: *npu_k8s_plugins
+
+  310p_npu_4:
+    agents:
+      queue: "ascend-310p"
+      resource_class: "npu-4"
+    image: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry/vllm-omni-ci-npu:${BUILDKITE_COMMIT}-310p"
+    plugins: *npu_k8s_plugins
@@ -70,7 +70,15 @@ from tests.helpers.mark import (  # noqa: E402

 LOG = "upload_pipeline"
 BOOTSTRAP_STEPS_FILENAME = "bootstrap-upload-steps.yml"
-BOOTSTRAP_IMAGE_BUILD_KEYS = frozenset({"image-build", "image-build-a2", "image-build-a3"})
+BOOTSTRAP_IMAGE_BUILD_KEYS = frozenset(
+    {
+        "image-build",
+        "image-build-a2",
+        "image-build-a3",
+        "image-build-a5",
+        "image-build-310p",
+    }
+)
 BOOTSTRAP_UPLOAD_IF_KEYS = {
    "upload-ready-pipeline": "ready",
    "upload-merge-pipeline": "merge",
@@ -79,6 +87,8 @@ BOOTSTRAP_UPLOAD_IF_KEYS = {
 }
 E2E_GROUP_MARKER = "E2E Test"
 CI_MIRROR_HARDWARES_PATH = ROOT / ".buildkite/common/ci_mirror_hardwares.yml"
+CUDA_HF_TOKEN_ENV = "VLLM_CI_HF_TOKEN"
+CUDA_HF_TOKEN_EXPORT = f'if [ -n "$${{{CUDA_HF_TOKEN_ENV}:-}}" ]; then export HF_TOKEN="$${{{CUDA_HF_TOKEN_ENV}}}"; fi'

 # Bootstrap Buildkite ``if`` expressions.
 # ``*_MAIN_IF``: main + env schedule. ``*_LABEL_IF``: PR label (and/or composed with MAIN).
@@ -501,6 +511,17 @@ def _expand_mirror_hardwares(step: dict[str, Any]) -> dict[str, Any] | None:

    expanded = copy.deepcopy(preset)
    merged = {key: value for key, value in step.items() if key != "mirror_hardwares"} | expanded
+    # A scheduled build's HF_TOKEN may override a pod env entry with the same
+    # name. The preset injects the Kubernetes secret under a collision-proof
+    # alias; restore the conventional name after Buildkite has merged env.
+    if any(preset_name == chip or preset_name.startswith(f"{chip}_") for chip in _cuda_mirror_chips()):
+        commands = merged.get("commands")
+        if isinstance(commands, list):
+            merged["commands"] = [CUDA_HF_TOKEN_EXPORT, *commands]
+        elif commands is None:
+            merged["commands"] = [CUDA_HF_TOKEN_EXPORT]
+        else:
+            merged["commands"] = [CUDA_HF_TOKEN_EXPORT, commands]
    # Preset retry (K8S_RETRY on l4_*) must not clobber a step that opted out.
    if "retry" in step:
        merged["retry"] = step["retry"]
@@ -155,6 +155,7 @@ steps:
        source_file_dependencies:
          - tests/e2e/offline_inference/test_minicpmo_4_5.py
          - tests/e2e/online_serving/test_minicpmo_4_5.py
+          - vllm_omni/transformers_utils/chat_templates/minicpmo45_native.jinja
          - vllm_omni/model_executor/models/minicpmo_4_5/
          - vllm_omni/model_executor/models/step_audio2/
          - vllm_omni/model_executor/models/cosyvoice3/code2wav_core/
@@ -188,6 +188,7 @@ steps:
      - label: "Omni · MiniCPM-o 4.5 Online Test"
        source_file_dependencies:
          - tests/e2e/online_serving/test_minicpmo_4_5.py
+          - vllm_omni/transformers_utils/chat_templates/minicpmo45_native.jinja
          - vllm_omni/model_executor/models/minicpmo_4_5/
          - vllm_omni/model_executor/models/step_audio2/
          - vllm_omni/model_executor/models/cosyvoice3/code2wav_core/
@@ -3,10 +3,10 @@
 steps:
  - label: ":buildkit: Build and Push NPU Test Image - A2"
    key: image-build-a2
-    agents:
+    agents: &npu-image-build-agents
      queue: "ascend-a2b3"
      resource_class: "npu-1"
-    plugins:
+    plugins: &npu-image-build-plugins
      - kubernetes:
          metadata:
            annotations:
@@ -46,123 +46,47 @@ steps:
      IMAGE_NAME: "vllm-omni-ci-npu"
      IMAGE_REGISTRY: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry"
      BUILDKITD_ADDR: "tcp://buildkitd-service.buildkitd:1234"
-    command: |
-      set -ex
-
-      echo "--- Building and pushing NPU Test Image"
-      echo "Image: $${IMAGE_REGISTRY}/$${IMAGE_NAME}:$${VLLM_IMAGE_TAG}"
-      echo "buildkitd address: $${BUILDKITD_ADDR}"
-
-      if ! command -v buildctl &> /dev/null; then
-        echo "Installing buildctl..."
-        mkdir -p /tmp/buildkit
-        BUILDKIT_VERSION="v0.29.0"
-        wget -q "https://github.com/moby/buildkit/releases/download/$${BUILDKIT_VERSION}/buildkit-$${BUILDKIT_VERSION}.linux-arm64.tar.gz" -O /tmp/buildkit.tar.gz
-        tar -xzf /tmp/buildkit.tar.gz -C /tmp/buildkit
-        cp /tmp/buildkit/bin/buildctl /usr/local/bin/
-      fi
-
-      export DOCKER_CONFIG=/home/user/.docker
-      buildctl \
-        --addr="$${BUILDKITD_ADDR}" \
-        --tlscacert=/home/user/.docker/ca.pem \
-        --tlscert=/home/user/.docker/cert.pem \
-        --tlskey=/home/user/.docker/key.pem \
-        build \
-        --frontend dockerfile.v0 \
-        --local context=. \
-        --local dockerfile=./docker \
-        --opt filename=Dockerfile.npu.ci \
-        --secret id=dockerconfig,src=/home/user/.docker/config.json \
-        --output type=image,name=$${IMAGE_REGISTRY}/$${IMAGE_NAME}:$${VLLM_IMAGE_TAG},push=true \
-        --progress=plain
-
-      echo "--- Image pushed successfully"
-      echo "$${IMAGE_REGISTRY}/$${IMAGE_NAME}:$${VLLM_IMAGE_TAG}"
+    command: bash .buildkite/npu/scripts/build-npu-image.sh Dockerfile.npu.ci "$${VLLM_IMAGE_TAG}"

  - label: ":buildkit: Build and Push NPU Test Image - A3"
    key: image-build-a3
-    agents:
-      queue: "ascend-a2b3"
-      resource_class: "npu-1"
-    plugins:
-      - kubernetes:
-          metadata:
-            annotations:
-              vault.hashicorp.com/agent-init-first: "true"
-              vault.hashicorp.com/agent-inject: "true"
-              vault.hashicorp.com/agent-inject-perms-ca.pem: "0400"
-              vault.hashicorp.com/agent-inject-perms-cert.pem: "0400"
-              vault.hashicorp.com/agent-inject-perms-config.json: "0400"
-              vault.hashicorp.com/agent-inject-perms-key.pem: "0400"
-              vault.hashicorp.com/agent-inject-secret-ca.pem: internal/data/ascend/buildkitd
-              vault.hashicorp.com/agent-inject-secret-cert.pem: internal/data/ascend/buildkitd
-              vault.hashicorp.com/agent-inject-secret-config.json: internal/data/ascend/buildkitd
-              vault.hashicorp.com/agent-inject-secret-key.pem: internal/data/ascend/buildkitd
-              vault.hashicorp.com/agent-inject-template-ca.pem: "{{- with secret \"internal/data/ascend/buildkitd\" -}}\n{{ .Data.data.RootCA }}\n{{- end }}"
-              vault.hashicorp.com/agent-inject-template-cert.pem: "{{- with secret \"internal/data/ascend/buildkitd\" -}}\n{{ .Data.data.ClientCaCert }}\n{{- end }}"
-              vault.hashicorp.com/agent-inject-template-config.json: "{{- with secret \"internal/data/ascend/buildkitd\" -}}\n{{ .Data.data.dockerConfig }}\n{{- end }}"
-              vault.hashicorp.com/agent-inject-template-key.pem: "{{- with secret \"internal/data/ascend/buildkitd\" -}}\n{{ .Data.data.ClientCaKey }}\n{{- end }}"
-              vault.hashicorp.com/agent-pre-populate-only: "true"
-              vault.hashicorp.com/agent-run-as-group: "1000"
-              vault.hashicorp.com/agent-run-as-user: "1000"
-              vault.hashicorp.com/agent-service-account-token-volume-name: token-vol
-              vault.hashicorp.com/role: ascend-gha-runners
-              vault.hashicorp.com/secret-volume-path: /home/user/.docker/
-              vault.hashicorp.com/tls-skip-verify: "true"
-          podSpecPatch:
-            volumes:
-            - name: token-vol
-              projected:
-                defaultMode: 420
-                sources:
-                - serviceAccountToken:
-                    audience: api
-                    expirationSeconds: 600
-                    path: token
+    agents: *npu-image-build-agents
+    plugins: *npu-image-build-plugins
    env:
      VLLM_IMAGE_TAG: "${BUILDKITE_COMMIT}-a3"
      IMAGE_NAME: "vllm-omni-ci-npu"
      IMAGE_REGISTRY: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry"
      BUILDKITD_ADDR: "tcp://buildkitd-service.buildkitd:1234"
-    command: |
-      set -ex
+    command: bash .buildkite/npu/scripts/build-npu-image.sh Dockerfile.npu.ci "$${VLLM_IMAGE_TAG}" -a3

-      echo "--- Building and pushing NPU Test Image"
-      echo "Image: $${IMAGE_REGISTRY}/$${IMAGE_NAME}:$${VLLM_IMAGE_TAG}"
-      echo "buildkitd address: $${BUILDKITD_ADDR}"
+  - label: ":buildkit: Build and Push NPU Test Image - A5"
+    key: image-build-a5
+    agents: *npu-image-build-agents
+    plugins: *npu-image-build-plugins
+    env:
+      VLLM_IMAGE_TAG: "${BUILDKITE_COMMIT}-a5"
+      IMAGE_NAME: "vllm-omni-ci-npu"
+      IMAGE_REGISTRY: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry"
+      BUILDKITD_ADDR: "tcp://buildkitd-service.buildkitd:1234"
+    command: bash .buildkite/npu/scripts/build-npu-image.sh Dockerfile.npu.ci "$${VLLM_IMAGE_TAG}" -a5

-      if ! command -v buildctl &> /dev/null; then
-        echo "Installing buildctl..."
-        mkdir -p /tmp/buildkit
-        BUILDKIT_VERSION="v0.29.0"
-        wget -q "https://github.com/moby/buildkit/releases/download/$${BUILDKIT_VERSION}/buildkit-$${BUILDKIT_VERSION}.linux-arm64.tar.gz" -O /tmp/buildkit.tar.gz
-        tar -xzf /tmp/buildkit.tar.gz -C /tmp/buildkit
-        cp /tmp/buildkit/bin/buildctl /usr/local/bin/
-      fi
-
-      export DOCKER_CONFIG=/home/user/.docker
-      buildctl \
-        --addr="$${BUILDKITD_ADDR}" \
-        --tlscacert=/home/user/.docker/ca.pem \
-        --tlscert=/home/user/.docker/cert.pem \
-        --tlskey=/home/user/.docker/key.pem \
-        build \
-        --frontend dockerfile.v0 \
-        --local context=. \
-        --local dockerfile=./docker \
-        --opt filename=Dockerfile.npu.ci.a3 \
-        --secret id=dockerconfig,src=/home/user/.docker/config.json \
-        --output type=image,name=$${IMAGE_REGISTRY}/$${IMAGE_NAME}:$${VLLM_IMAGE_TAG},push=true \
-        --progress=plain
-
-      echo "--- Image pushed successfully"
-      echo "$${IMAGE_REGISTRY}/$${IMAGE_NAME}:$${VLLM_IMAGE_TAG}"
+  - label: ":buildkit: Build and Push NPU Test Image - 310P"
+    key: image-build-310p
+    agents: *npu-image-build-agents
+    plugins: *npu-image-build-plugins
+    env:
+      VLLM_IMAGE_TAG: "${BUILDKITE_COMMIT}-310p"
+      IMAGE_NAME: "vllm-omni-ci-npu"
+      IMAGE_REGISTRY: "swr.cn-southwest-2.myhuaweicloud.com/modelfoundry"
+      BUILDKITD_ADDR: "tcp://buildkitd-service.buildkitd:1234"
+    command: bash .buildkite/npu/scripts/build-npu-image.sh Dockerfile.npu.ci "$${VLLM_IMAGE_TAG}" -310p

  - label: "Upload Ready Pipeline"
    depends_on:
      - image-build-a2
      - image-build-a3
+      - image-build-a5
+      - image-build-310p
    key: upload-ready-pipeline
    agents:
      queue: "cpu_queue_premerge"
@@ -174,6 +98,8 @@ steps:
    depends_on:
      - image-build-a2
      - image-build-a3
+      - image-build-a5
+      - image-build-310p
    key: upload-nightly-pipeline
    commands:
      - python3 .buildkite/common/scripts/upload_pipeline.py --upload .buildkite/npu/test-npu-nightly.yml
@@ -0,0 +1,48 @@
+#!/usr/bin/env bash
+set -ex
+
+dockerfile=${1:?Dockerfile is required}
+image_tag=${2:?Image tag is required}
+device_suffix=${3:-}
+
+image_name=${IMAGE_NAME:-vllm-omni-ci-npu}
+image_registry=${IMAGE_REGISTRY:-swr.cn-southwest-2.myhuaweicloud.com/modelfoundry}
+buildkitd_addr=${BUILDKITD_ADDR:-tcp://buildkitd-service.buildkitd:1234}
+
+echo "--- Building and pushing NPU Test Image"
+echo "Image: ${image_registry}/${image_name}:${image_tag}"
+echo "Dockerfile: ${dockerfile}"
+
+if ! command -v buildctl &> /dev/null; then
+  mkdir -p /tmp/buildkit
+  buildkit_version=v0.29.0
+  wget -q \
+    "https://github.com/moby/buildkit/releases/download/${buildkit_version}/buildkit-${buildkit_version}.linux-arm64.tar.gz" \
+    -O /tmp/buildkit.tar.gz
+  tar -xzf /tmp/buildkit.tar.gz -C /tmp/buildkit
+  cp /tmp/buildkit/bin/buildctl /usr/local/bin/
+fi
+
+build_args=()
+if [[ -n "${device_suffix}" ]]; then
+  build_args+=(--opt "build-arg:VLLM_ASCEND_DEVICE_SUFFIX=${device_suffix}")
+fi
+
+export DOCKER_CONFIG=/home/user/.docker
+buildctl \
+  --addr="${buildkitd_addr}" \
+  --tlscacert=/home/user/.docker/ca.pem \
+  --tlscert=/home/user/.docker/cert.pem \
+  --tlskey=/home/user/.docker/key.pem \
+  build \
+  --frontend dockerfile.v0 \
+  --local context=. \
+  --local dockerfile=./docker \
+  --opt "filename=${dockerfile}" \
+  "${build_args[@]}" \
+  --secret id=dockerconfig,src=/home/user/.docker/config.json \
+  --output "type=image,name=${image_registry}/${image_name}:${image_tag},push=true" \
+  --progress=plain
+
+echo "--- Image pushed successfully"
+echo "${image_registry}/${image_name}:${image_tag}"
@@ -1,28 +1,59 @@
 steps:
-  - label: "Platforms Test"
-    key: ready-platforms-test
+  - group: ":card_index_dividers: Simple Test"
    depends_on: upload-ready-pipeline
-    mirror_hardwares: a3_npu_2
-    commands:
-      - >-
-        pytest -s -v
-        tests/platforms/npu/test_minicpmo_code2wav_npugraph.py
-        -m 'core_model and npu and A3 and omni'
-        --run-level 'core_model'
+    steps:
+      - label: "Platforms Test"
+        key: ready-platforms-test
+        mirror_hardwares: a3_npu_2
+        commands:
+          - >-
+            pytest -s -v
+            tests/platforms/npu/test_minicpmo_code2wav_npugraph.py
+            -m 'core_model and npu and A3 and omni'
+            --run-level 'core_model'

-  - label: "Diffusion NPU layer tests"
-    key: ready-diffusion-npu-layer-tests
-    depends_on: upload-ready-pipeline
-    mirror_hardwares: a3_npu_2
-    commands:
-      - >-
-        pytest -s -v
-        tests/diffusion/layers/test_activation.py
-        tests/diffusion/layers/test_fused_qk_norm_rope_npu.py
-        tests/diffusion/models/minimax_h3/test_minimax_h3_qwen3vl_swiglu.py
-        tests/diffusion/models/minimax_h3/test_minimax_h3_modulation.py
-        -m 'core_model and npu and A3'
-        --run-level 'core_model'
+      - label: "Diffusion NPU layer tests - A2"
+        key: ready-diffusion-npu-layer-tests-a2
+        mirror_hardwares: a2b3_npu_1
+        commands:
+          - >-
+            pytest -s -v
+            tests/diffusion/layers/test_activation.py
+            -m 'core_model and npu and A2'
+            --run-level 'core_model'
+
+      - label: "Diffusion NPU layer tests - A3"
+        key: ready-diffusion-npu-layer-tests
+        mirror_hardwares: a3_npu_2
+        commands:
+          - >-
+            pytest -s -v
+            tests/diffusion/layers/test_activation.py
+            tests/diffusion/layers/test_fused_qk_norm_rope_npu.py
+            tests/diffusion/models/minimax_h3/test_minimax_h3_qwen3vl_swiglu.py
+            tests/diffusion/models/minimax_h3/test_minimax_h3_modulation.py
+            -m 'core_model and npu and A3'
+            --run-level 'core_model'
+
+      - label: "Diffusion NPU layer tests - A5"
+        key: ready-diffusion-npu-layer-tests-a5
+        mirror_hardwares: a5_npu_2
+        commands:
+          - >-
+            pytest -s -v
+            tests/diffusion/layers/test_activation.py
+            -m 'core_model and npu and A5'
+            --run-level 'core_model'
+
+      - label: "Diffusion NPU layer tests - 310P"
+        key: ready-diffusion-npu-layer-tests-310p
+        mirror_hardwares: 310p_npu_1
+        commands:
+          - >-
+            pytest -s -v
+            tests/diffusion/layers/test_activation.py
+            -m 'core_model and npu and 310P'
+            --run-level 'core_model'

  - group: ":card_index_dividers: Omni Model Test"
    key: ready-omni-group
@@ -123,10 +123,11 @@ it changes again, report the churn and wait for a stable target.

 For a trusted PR head, materialize the pinned head in an isolated detached
 worktree. A worktree freezes identity but is not a security sandbox. Treat fork
-heads as untrusted unless the user and environment policy explicitly establish
-otherwise: execute them only in a disposable, secret-free sandbox with restricted
-filesystem, network, and resources; without one, use static SHA-addressed reads
-and CI evidence only. For a local review, freeze the committed, index, worktree,
+heads as untrusted until the user and environment policy explicitly establish
+trust: execute their code only after that trust is recorded for this host,
+with the trusted SHA noted in the review state and secrets kept out of the
+execution scope; without recorded trust, use static SHA-addressed reads and CI
+evidence only. For a local review, freeze the committed, index, worktree,
 and NUL-safe in-scope untracked contents. Follow
 [review-execution.md](references/process/review-execution.md) for trust gates,
 state fingerprints, and byte-for-byte staleness checks.
@@ -196,7 +197,7 @@ style preferences as simplification findings.

 Before each validation group, verify the frozen SHA plus the tracked, index,
 untracked, and ignored-file fingerprint, or recreate a pristine snapshot. On a
-trusted head or inside the required sandbox, run an import/version preflight,
+head the user explicitly trusted for this host, run an import/version preflight,
 then the narrowest relevant tests and low-cost static checks. Bind every result
 to the head SHA, snapshot fingerprint, and environment fingerprint. Never run
 imports, tests, builds, hooks, or repo-configurable tooling from an untrusted
@@ -22,6 +22,19 @@ Keep exact commands and report variability rather than a cherry-picked run.
 Use the repository's benchmark when it covers the claim; otherwise use the
 smallest reproducible workload that reaches the changed production path.

+## Isolation per claim
+
+Isolate one claim per comparison. When a PR stacks several changes (a fused or
+restructured load path, a distilled step schedule, a new cache), request one
+column per claim at equal steps/workload instead of one blended base-vs-head
+delta, and name the unchanged stage or metric that pins equivalence.
+
+When the changed path is runtime-switchable (feature flag or environment
+variable), run every column on the same head and toggle the switch rather than
+diffing against a base commit that cannot separate the stacked changes:
+base, head with the switch off, head with the switch on, then head with the
+switch on plus each further claim.
+
 ## Evidence table

 | Dimension | Typical evidence |
@@ -34,6 +47,100 @@ smallest reproducible workload that reaches the changed production path.
 Every optimization needs correctness or quality evidence. A faster result on a
 different workload, precision, seed, or topology is not a valid comparison.

+## E2e stage-attribution table
+
+For an end-to-end latency claim, request the stage-attribution table the
+repository's perf PRs use:
+
+- Rows carry the benchmark's own field names: total e2e, per-stage execution,
+  postprocess, response encoding, and a request-to-output roll-up. Add a stage
+  handoff/idle row when the pipeline spans stages — connector transfer and
+  scheduling gaps are invisible in per-stage latency deltas.
+- Sanity rows pin comparability: inference steps, output frames/pixels/samples,
+  and generated tokens held identical across columns.
+- Report P50 and P100 from the stated warmup and measured counts. One
+  population per number: never mix cold-start and warm samples into one value.
+- Name the baseline and candidate commit SHAs, and close with an isolation
+  statement naming the metric that stayed unchanged ("X remained effectively
+  unchanged, isolating the gain to Y").
+
+When the serving path caches, captures, or regenerates at runtime (CUDA-graph
+bucket growth, cache rebuilds, recompilation), medians hide the resulting tail.
+Also request the event count and the maximum observed stall.
+
+Component-level tables pair with, not replace, the e2e table: per-op reference
+vs candidate timings (with the numerical-equality result per op) plus one
+complete-pipeline check under the production shapes.
+
+## Proportionality
+
+When a PR adds a model and explicitly labels its numbers as smoke observations
+rather than claims, do not demand the attribution machinery. Ask only for
+evidence protecting the model's stated operating contract — for a
+realtime/duplex model that is sustained cadence over a realistic session
+(enough ticks to leave warmup), with any mid-session capture/regeneration
+event and its stall reported, since a cadence that collapses mid-session is a
+functional failure, not a perf claim.
+
+## Realtime/duplex world models
+
+Request the operating-contract evidence as serving metrics, not claims:
+
+- Sustained per-tick latency over one realistic session: median, max, and the
+  argmax tick.
+- Stall events with their causes.
+- Peak memory at first versus last tick — session-owned state must not grow
+  per tick.
+- Control latency from a submitted action to the tick that applies it.
+- Realtime-versus-offline parity under the same controls, using the
+  repository's video-similarity helpers. Parity plotted over tick index is the
+  drift-over-horizon curve and the honest measure of session lifetime;
+  distribution-level suites (FVD, VBench) belong to model evaluation, not the
+  serving integration.
+
+Watch the contract boundaries the model documents: if a control input is
+tick-API-only, the parity pair must use a horizon both paths can run.
+
+## Speech-generation (TTS) models
+
+The streaming contract outranks the average — a synthesis that pauses
+mid-stream is broken for realtime use regardless of its averages. Request:
+
+- The benchmark's own speech metrics over one long utterance (enough audio to
+  leave warmup, tens of seconds): `audio_ttfp` (time to first packet),
+  `audio_rtf` (real-time factor; above 1.0 cannot keep up), `e2el`, and
+  `ttft`/`tpot` where tokens precede audio.
+- Mid-utterance stalls as the benchmark's continuity fields, not prose:
+  `audio_underrun` percentiles and the `audio_continuity_ok` rate
+  (`vllm_omni/benchmarks/metrics/metrics.py`), computed by
+  `vllm_omni/benchmarks/audio_continuity.py` — the companion to `audio_ttfp`
+  and `audio_rtf`. A synthesis that bursts chunks fast enough on average but
+  underruns mid-stream fails the streaming contract even when the averages
+  pass.
+- Quality as a paired comparison against the reference implementation under
+  the same text, speaker prompt, and seed; report sample rate and frame
+  count, and assert output completeness — truncated or dropped trailing audio
+  is a known failure class, not a quality nuance.
+- For realtime or duplex TTS: interruption (barge-in) latency and the
+  per-session memory trend across turns.
+
+## Omni models
+
+Apply the e2e stage-attribution table directly to the voice-to-voice path of
+omni (speech-to-speech and multimodal AR) models:
+
+- TTFT for text output, `audio_ttfp` for audio output, and per-stage
+  ITL/TPOT cadence across the perception-understanding-generation chain.
+- Cross-modal parity — the omni analogue of realtime-versus-offline parity:
+  the same semantic request through text and audio inputs, compared on output
+  content or timing.
+- For duplex omni sessions: turn-taking latency, interruption handling, the
+  concurrent session count actually validated, and a per-turn memory trend —
+  session-owned KV state must not grow per turn. When either turn produces
+  audio, continuity is judged with the same `audio_underrun` and
+  `audio_continuity_ok` fields as TTS, since turn-taking latency alone does
+  not expose mid-utterance underruns.
+
 ## Reviewer verification ladder

 1. Run base/head A/B on suitable hardware when affordable.
@@ -121,6 +121,19 @@ Search bounded callers and sibling implementations before reporting dead code,
 duplication, or a missed consumer. Keep new abstractions only when they have a
 distinct live caller, invariant, or compatibility need.

+Verify absence and equality claims with whole-file evidence. A line window, a
+grep for one symbol, or a commit message proves presence, never absence: to
+assert that a file matches another version or lacks a change, diff or hash the
+full files (for example `git diff <sha1>:<path> <sha2>:<path>` or the contents
+API's size and SHA fields) before anchoring the finding. A verification
+shortcut that would itself be wrong to run is not evidence.
+
+For a value written or finalized in more than one place, establish which site
+executes first on each entry path — and whether the earlier site makes the
+later ones unreachable — before naming where a fix belongs. Enumerating all
+call sites correctly still misses the bug when a path the enumeration ranked
+last actually finalizes first and clears the state the proposed fix checks.
+
 ## Finding bar

 Anchor each finding to a changed `path:line`; name the trigger or call path,
@@ -39,13 +39,13 @@ isolated detached worktree. Run every source read, `rg` search, import, and test
 from that snapshot, not the caller's checkout. A detached worktree provides
 snapshot isolation, not a security boundary.

-Treat a fork head as untrusted unless the user and environment policy explicitly
-establish trust. Never run its imports, tests, builds, hooks, package setup, or
-repo-configurable linters/plugins on the reviewer host. Execute them only in a
-disposable sandbox or VM with no credentials or inherited secrets, no host agent
-or service sockets, minimal read-only host mounts, disabled or explicitly
-allowlisted network, resource/time limits, and destruction after the run. If
-that boundary is unavailable, limit the review to the remote diff, SHA-addressed
+Treat a fork head as untrusted until the user and environment policy
+explicitly establish trust. Without recorded trust, never run its imports,
+tests, builds, hooks, package setup, or repo-configurable linters/plugins on
+the reviewer host. Once the user explicitly trusts this pinned commit for
+this host, record the trusted SHA with the review state and execute with
+credentials, agent sockets, and other secrets out of scope. Without recorded
+trust, limit the review to the remote diff, SHA-addressed
 reads such as `git show <head_sha>:<path>`, and existing CI evidence; report all
 executable validation as a gap.

@@ -1,27 +1,10 @@
-# Target vllm version: v0.28.0 (tag v0.28.0 @ 2cf0a6915ce5).
-# Keep VLLM_BASE_TAG aligned with the target vLLM release on every rebase.
-#
-# Release-boundary form: there is no precompiled-wheel pin here, because the
-# base image already ships exactly the target vLLM. Everything the runtime
-# needs — torch, flashinfer, nvJitLink, numpy — comes from that image as one
-# self-consistent set.
-#
-# This block reappears only when the rebase targets a vLLM commit that is NOT a
-# released tag. During the 0.27 rebase, targeting v0.27.0rc1 on top of the
-# v0.26.0 base image required reinstalling vLLM from wheels.vllm.ai and then
-# hand-repairing the fallout: pinning torch ahead of flashinfer (installing it
-# after bound flashinfer's TVM FFI runtime to the old torch and aborted engine
-# core init), registering libnvJitLink.so.13 with ldconfig, and re-pinning numpy
-# under numba's ceiling. All of it was compensation for a base image and a vLLM
-# that disagreed. At a real release boundary they agree, so none of it is
-# needed. The 0.28 rebase repeated the cycle: the v0.28.0rc / final-tag wheels
-# rode on the v0.27.0 base image with the full reinstall block until the
-# official v0.28.0 image shipped, at which point this file dropped back to the
-# form you are reading. If you pin an untagged commit again, restore the
-# reinstall block — agent/tasks/20_update_wheel_url.sh fails the run rather
-# than letting the Dockerfile silently ship the wrong vLLM.
+# Target vLLM version: v0.29.0 (tag @ 98dff2a81d74).
+# Use the matching release image and its coherent CUDA/PyTorch/FlashInfer
+# runtime. Do not reinstall a commit wheel or repair dependencies on top of
+# an older base image at a release boundary. For a future unreleased target,
+# restore an explicit wheel pin and validate its complete runtime stack.
 ARG VLLM_BASE_IMAGE=vllm/vllm-openai
-ARG VLLM_BASE_TAG=v0.28.0
+ARG VLLM_BASE_TAG=v0.29.0

 # We use registry-based layer caching for the dependencies,
 # where the tag is determined based on the hash of the contents
@@ -1,20 +1,7 @@
 ARG VLLM_ASCEND_IMAGE=quay.io/atlas-ci/vllm-ascend
-ARG VLLM_ASCEND_TAG=v0.28.0
-FROM ${VLLM_ASCEND_IMAGE}:${VLLM_ASCEND_TAG}
-
-# WORKDIR /vllm-workspace/vllm
-# RUN git fetch origin --tags && git checkout v0.18.0
-
-# WORKDIR /vllm-workspace/vllm-ascend
-# RUN git fetch origin releases/v0.18.0 && git checkout d781902ce9dbda8ab1e11bb0f2f0c1bc508fee7a
-# # Install vllm-ascend
-# # Append `libascend_hal.so` path (devlib) to LD_LIBRARY_PATH
-# RUN export PIP_EXTRA_INDEX_URL=https://mirrors.huaweicloud.com/ascend/repos/pypi && \
-#     source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
-#     source /usr/local/Ascend/nnal/atb/set_env.sh && \
-#     export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/Ascend/ascend-toolkit/latest/`uname -i`-linux/devlib && \
-#     python3 -m pip install -v -e /vllm-workspace/vllm-ascend/ --extra-index https://download.pytorch.org/whl/cpu/ && \
-#     python3 -m pip cache purge
+ARG VLLM_ASCEND_VERSION=v0.28.0
+ARG VLLM_ASCEND_DEVICE_SUFFIX=
+FROM ${VLLM_ASCEND_IMAGE}:${VLLM_ASCEND_VERSION}${VLLM_ASCEND_DEVICE_SUFFIX}

 ARG APP_DIR=/vllm-workspace/vllm-omni
 WORKDIR ${APP_DIR}
@@ -24,6 +11,7 @@ RUN export PIP_EXTRA_INDEX_URL=https://mirrors.huaweicloud.com/ascend/repos/pypi
    source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
    source /usr/local/Ascend/nnal/atb/set_env.sh && \
    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/Ascend/ascend-toolkit/latest/`uname -i`-linux/devlib && \
+    python3 -m pip install -v "mindiesd==3.1.0" && \
    python3 -m pip install -v -e /vllm-workspace/vllm-omni/ --no-build-isolation

 RUN python3 -m pip install -v --no-deps \
@@ -1,36 +0,0 @@
-ARG VLLM_ASCEND_IMAGE=quay.io/atlas-ci/vllm-ascend
-ARG VLLM_ASCEND_TAG=v0.28.0-a3
-FROM ${VLLM_ASCEND_IMAGE}:${VLLM_ASCEND_TAG}
-
-# WORKDIR /vllm-workspace/vllm
-# RUN git fetch origin --tags && git checkout v0.18.0
-
-# WORKDIR /vllm-workspace/vllm-ascend
-# RUN git fetch origin releases/v0.18.0 && git checkout d781902ce9dbda8ab1e11bb0f2f0c1bc508fee7a
-# # Install vllm-ascend
-# # Append `libascend_hal.so` path (devlib) to LD_LIBRARY_PATH
-# RUN export PIP_EXTRA_INDEX_URL=https://mirrors.huaweicloud.com/ascend/repos/pypi && \
-#     source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
-#     source /usr/local/Ascend/nnal/atb/set_env.sh && \
-#     export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/Ascend/ascend-toolkit/latest/`uname -i`-linux/devlib && \
-#     python3 -m pip install -v -e /vllm-workspace/vllm-ascend/ --extra-index https://download.pytorch.org/whl/cpu/ && \
-#     python3 -m pip cache purge
-
-ARG APP_DIR=/vllm-workspace/vllm-omni
-WORKDIR ${APP_DIR}
-COPY . .
-
-RUN export PIP_EXTRA_INDEX_URL=https://mirrors.huaweicloud.com/ascend/repos/pypi && \
-    source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
-    source /usr/local/Ascend/nnal/atb/set_env.sh && \
-    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/Ascend/ascend-toolkit/latest/`uname -i`-linux/devlib && \
-    python3 -m pip install -v -e /vllm-workspace/vllm-omni/ --no-build-isolation
-
-RUN python3 -m pip install -v --no-deps \
-      "stepaudio2-minicpmo>=0.1.1" \
-      "step-audio2>=1.0.0"
-
-ENV VLLM_TARGET_DEVICE=npu
-ENV VLLM_WORKER_MULTIPROC_METHOD=spawn
-
-ENTRYPOINT []
@@ -1,10 +1,10 @@
 ARG VLLM_ASCEND_IMAGE=quay.io/atlas-ci/vllm-ascend
-ARG VLLM_ASCEND_TAG=v0.28.0
-FROM ${VLLM_ASCEND_IMAGE}:${VLLM_ASCEND_TAG}
+ARG VLLM_ASCEND_VERSION=v0.28.0
+ARG VLLM_ASCEND_DEVICE_SUFFIX=
+FROM ${VLLM_ASCEND_IMAGE}:${VLLM_ASCEND_VERSION}${VLLM_ASCEND_DEVICE_SUFFIX}

 ARG APP_DIR=/vllm-workspace/vllm-omni
 WORKDIR ${APP_DIR}
-COPY . .

 RUN sed -i "s@^\(deb.*\)https\?://[a-z0-9.-]*\.ubuntu\.com@\1http://cache-service.nginx-pypi-cache.svc.cluster.local:8081@g" /etc/apt/sources.list

@@ -22,11 +22,9 @@ RUN apt-get update && \
    rm -rf /var/cache/apt/* && \
    rm -rf /var/lib/apt/lists/*

-RUN source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
-    source /usr/local/Ascend/nnal/atb/set_env.sh && \
-    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/Ascend/ascend-toolkit/latest/`uname -i`-linux/devlib && \
-    python3 -m pip install -U pip && \
+RUN python3 -m pip install -U pip && \
    python3 -m pip install -v \
+      "mindiesd==3.1.0" \
      "pytest>=7.0.0" \
      "pytest-asyncio>=0.21.0" \
      "pytest-cov>=4.0.0" \
@@ -48,7 +46,13 @@ RUN source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
      "datasets>=3.0.0" \
      --index-url http://cache-service.nginx-pypi-cache.svc.cluster.local/pypi/simple \
      --trusted-host cache-service.nginx-pypi-cache.svc.cluster.local \
-      --default-timeout=60 && \
+      --default-timeout=60
+
+COPY . .
+
+RUN source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
+    source /usr/local/Ascend/nnal/atb/set_env.sh && \
+    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/Ascend/ascend-toolkit/latest/`uname -i`-linux/devlib && \
    python3 -m pip install -v -e /vllm-workspace/vllm-omni/ --no-build-isolation \
      --index-url http://cache-service.nginx-pypi-cache.svc.cluster.local/pypi/simple \
      --trusted-host cache-service.nginx-pypi-cache.svc.cluster.local \
@@ -1,72 +0,0 @@
-ARG VLLM_ASCEND_IMAGE=quay.io/atlas-ci/vllm-ascend
-ARG VLLM_ASCEND_TAG=v0.28.0-a3
-FROM ${VLLM_ASCEND_IMAGE}:${VLLM_ASCEND_TAG}
-
-ARG APP_DIR=/vllm-workspace/vllm-omni
-WORKDIR ${APP_DIR}
-COPY . .
-
-RUN sed -i "s@^\(deb.*\)https\?://[a-z0-9.-]*\.ubuntu\.com@\1http://cache-service.nginx-pypi-cache.svc.cluster.local:8081@g" /etc/apt/sources.list
-
-RUN apt-get update && \
-    apt-get install -y --no-install-recommends \
-        ffmpeg \
-        libgl1 \
-        libglib2.0-0 \
-        libgomp1 \
-        libsm6 \
-        libxext6 \
-        libxrender1 \
-        libxcb1 \
-        espeak-ng && \
-    rm -rf /var/cache/apt/* && \
-    rm -rf /var/lib/apt/lists/*
-
-RUN source /usr/local/Ascend/ascend-toolkit/set_env.sh && \
-    source /usr/local/Ascend/nnal/atb/set_env.sh && \
-    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/Ascend/ascend-toolkit/latest/`uname -i`-linux/devlib && \
-    python3 -m pip install -U pip && \
-    python3 -m pip install -v \
-      "pytest>=7.0.0" \
-      "pytest-asyncio>=0.21.0" \
-      "pytest-cov>=4.0.0" \
-      "pytest-mock>=3.10.0" \
-      "openpyxl>=3.0.0" \
-      "psutil>=7.2.0" \
-      "soundfile>=0.13.1" \
-      "pyttsx3>=2.99" \
-      "opencc==1.4.1" \
-      "opencv-python>=4.12.0.88" \
-      "imageio[ffmpeg]>=0.6.0" \
-      "jiwer>=3.0.0" \
-      "zhon>=2.0.0" \
-      "zhconv>=1.4.2" \
-      "funasr>=1.0.0" \
-      "hyperpyyaml>=1.2.0" \
-      "onnx>=1.16.0" \
-      "librosa>=0.10.0" \
-      "torchmetrics>=1.9.0" \
-      "datasets>=3.0.0" \
-      --index-url http://cache-service.nginx-pypi-cache.svc.cluster.local/pypi/simple \
-      --trusted-host cache-service.nginx-pypi-cache.svc.cluster.local \
-      --default-timeout=60 && \
-    python3 -m pip install -v -e /vllm-workspace/vllm-omni/ --no-build-isolation \
-      --index-url http://cache-service.nginx-pypi-cache.svc.cluster.local/pypi/simple \
-      --trusted-host cache-service.nginx-pypi-cache.svc.cluster.local \
-      --default-timeout=60
-
-# MiniCPM-o 4.5 Stage-2 vocoder assets: s3tokenizer plus flashcosyvoice, which
-# is shipped by step-audio2. --no-deps keeps pip from replacing the Ascend torch
-# build and from shadowing onnxruntime-cann with the vanilla onnxruntime wheel;
-# the deps these two actually need are installed above.
-RUN python3 -m pip install -v --no-deps \
-      "s3tokenizer>=0.3.0" \
-      "step-audio2>=1.0.0" \
-      --index-url http://cache-service.nginx-pypi-cache.svc.cluster.local/pypi/simple \
-      --trusted-host cache-service.nginx-pypi-cache.svc.cluster.local \
-      --default-timeout=60
-
-ENV VLLM_TARGET_DEVICE=npu
-ENV VLLM_WORKER_MULTIPROC_METHOD=spawn
-
-ENTRYPOINT []
@@ -2,7 +2,7 @@

 This section lists the most common options for running vLLM-Omni.

-For options within a vLLM Engine, please refer to the [vLLM 0.28 configuration guide](https://docs.vllm.ai/en/v0.28.0/configuration/index.html).
+For options within a vLLM Engine, please refer to the [vLLM 0.29 configuration guide](https://docs.vllm.ai/en/v0.29.0/configuration/index.html).

 Each model defines fixed topology in a registered `PipelineConfig` and runtime
 overrides in a deploy YAML.
@@ -151,7 +151,7 @@ their keys only.
 ## Inherited vLLM variables

 vLLM-Omni also reads variables through its aligned vLLM dependency. Refer to
-the [vLLM 0.28 environment-variable reference](https://docs.vllm.ai/en/v0.28.0/configuration/env_vars.html)
+the [vLLM 0.29 environment-variable reference](https://docs.vllm.ai/en/v0.29.0/configuration/env_vars.html)
 for their definitions. This includes vLLM launch, cache, logging, plugin, ROCm,
 XPU, ModelScope, and FlashInfer workspace settings.

@@ -129,7 +129,7 @@ Canonical layout (prefer these paths for new changes):

 === "NPU"

-    **Bootstrap:** [`npu/pipeline-npu.yml`](https://github.com/vllm-project/vllm-omni/blob/main/.buildkite/npu/pipeline-npu.yml) + [`npu/bootstrap-upload-steps.yml`](https://github.com/vllm-project/vllm-omni/blob/main/.buildkite/npu/bootstrap-upload-steps.yml)—same split as CUDA; builds A2/B3 and A3 CI images, then uploads child test pipelines.
+    **Bootstrap:** [`npu/pipeline-npu.yml`](https://github.com/vllm-project/vllm-omni/blob/main/.buildkite/npu/pipeline-npu.yml) + [`npu/bootstrap-upload-steps.yml`](https://github.com/vllm-project/vllm-omni/blob/main/.buildkite/npu/bootstrap-upload-steps.yml)—same split as CUDA; builds A2/B3, A3, A5, and 310P CI images, then uploads child test pipelines.

    **Test YAML:** `npu/test-npu-ready.yml` (L2), `test-npu-nightly.yml` (L4).

@@ -137,7 +137,12 @@ Canonical layout (prefer these paths for new changes):

    **Upload:** same as CUDA—`upload_pipeline.py --upload`.

-    **Hardware in YAML:** `mirror_hardwares` preset (string), expanded to `agents`, top-level `image`, and `plugins`. Presets: `a2b3_npu_1`, `a2b3_npu_4`, `a3_npu_2` in `common/ci_mirror_hardwares.yml`.
+    **Hardware in YAML:** `mirror_hardwares` preset (string), expanded to `agents`, top-level `image`, and `plugins`. NPU presets in `common/ci_mirror_hardwares.yml` are:
+
+    - A2: `a2b3_npu_1`, `a2b3_npu_4`, `a2b3_npu_8`
+    - A3: `a3_npu_2`, `a3_npu_4`, `a3_npu_8`, `a3_npu_16`
+    - A5: `a5_npu_2`, `a5_npu_4`, `a5_npu_8`, `a5_npu_16`
+    - 310P: `310p_npu_1`, `310p_npu_2`, `310p_npu_4`

    **Conventions**

@@ -161,7 +166,9 @@ Canonical layout (prefer these paths for new changes):

    **Bootstrap:** [`amd/scripts/bootstrap-amd-omni.sh`](https://github.com/vllm-project/vllm-omni/blob/main/.buildkite/amd/scripts/bootstrap-amd-omni.sh)—skip-ci, diff filtering, Jinja render, then `buildkite-agent pipeline upload`.

-    **Test YAML (data):** `amd/test-amd-ready.yml` (L2 / PR), `test-amd-merge.yml` (L3 / main).
+    **Test YAML (data):** `amd/test-amd-ready.yml` (L2 / `ready`), `test-amd-merge.yml` (L3 / `merge-test` and `main`). On a PR carrying both labels, AMD combines both suites behind one image build. PR builds without either tier label retain the legacy L2 fallback for compatibility with the `amd-test` trigger. `DEBUG_TEST_YAML` remains an explicit override. AMD has no L4 file yet, so `nightly-test` does not select an additional suite.
+
+    **Trigger boundary:** AMD label handling has two layers. First, the external `vllm-omni-amd-ci` Buildkite pipeline condition decides whether a GitHub label event creates a build. Only after that build starts does this repository's bootstrap inspect all current PR labels and select L2, L3, or both. Repository-side selection therefore cannot make a `merge-test` event start AMD CI by itself. The external condition must admit both `ready` and `merge-test`, while preserving the legacy `amd-test` trigger and non-PR `main` / scheduled builds. For PR builds without a `DEBUG_TEST_YAML` override, an unreadable label set fails the bootstrap instead of silently selecting the wrong tier. Debug overrides skip label lookup and are validated directly; invalid values, including whitespace-only values, fail the bootstrap. Keep `nightly-test` out of the AMD condition until an AMD L4 suite exists; otherwise it would start a build without providing nightly coverage.

    **Rendering:** [`test-template-amd-omni.j2`](https://github.com/vllm-project/vllm-omni/blob/main/.buildkite/amd/test-template-amd-omni.j2) wraps data steps with `amd-build` image build and `amd_<agent_pool>` queues. Do **not** hand-edit generated `pipeline.yaml`.

@@ -48,6 +48,8 @@ Defined in `pyproject.toml`:
 | `MI325`            | Tests that require MI325 GPU (AMD/ROCm) *                                                               |
 | `A2`               | Tests that require A2 NPU *                                                                             |
 | `A3`               | Tests that require A3 NPU *                                                                             |
+| `A5`               | Tests that require A5 NPU *                                                                             |
+| `310P`             | Tests that require 310P NPU *                                                                           |
 | `cards_{n}`        | Tests that require *n* accelerator cards (`cards_1` … `cards_8`); auto-added from `num_cards` *         |
 | `slow`             | Slow tests (may skip in quick CI)                                                                       |
 | `benchmark`        | Benchmark tests (decorator on runner test functions; perf JSON uses `full_model` + type marker instead) |
@@ -76,7 +78,7 @@ def test_video_to_audio()
 This decorator is intended to make hardware-aware, cross-platform test authoring easier and more robust for CI/CD environments. The `hardware_test` decorator in `vllm-omni/tests/helpers/mark.py` performs the following actions:

 1. **Applies platform and resource markers**  
-   Adds the appropriate pytest markers for each specified hardware platform (e.g., `cuda`, `rocm`, `xpu`, `npu`) and resource type (e.g., `L4`, `H100`, `H200`, `B200`, `MI325`, `B60`, `A2`, `A3`).
+   Adds the appropriate pytest markers for each specified hardware platform (e.g., `cuda`, `rocm`, `xpu`, `npu`) and resource type (e.g., `L4`, `H100`, `H200`, `B200`, `MI325`, `B60`, `A2`, `A3`, `A5`, `310P`).
   ```python
   @pytest.mark.cuda
   @pytest.mark.L4
@@ -1,7 +1,7 @@
 # Installation

 !!! important
-    vLLM-Omni is released against the matching upstream vLLM major/minor version. vLLM-Omni 0.28.x requires vLLM 0.28.x; the stable v0.28.0 instructions pin vLLM 0.28.0.
+    vLLM-Omni is released against the matching upstream vLLM major/minor version. The 0.29 development line requires vLLM 0.29.x; the source-install instructions pin vLLM 0.29.0. Published vLLM-Omni 0.28.0 wheels and images still require vLLM 0.28.x.

 vLLM-Omni supports the following hardware platforms:

@@ -34,6 +34,8 @@ vLLM-Omni is a Python library that supports the following GPU variants. The libr

 ### Pre-built wheels

+The source-install instructions target vLLM 0.29.0. Until vLLM-Omni 0.29.0 artifacts are published, the CUDA/ROCm pre-built examples below remain on the matching 0.28.0 release pair.
+
 Note: Pre-built wheels are currently available for vLLM-Omni 0.11.0rc1, 0.12.0rc1, 0.14.0rc1, 0.14.0, 0.16.0, 0.18.0, 0.20.0, 0.21.0, 0.22.0, 0.23.0, 0.24.0, 0.25.0, 0.26.0, and 0.28.0. If you need a newer unreleased revision, please [build from source](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/gpu/#build-wheel-from-source).

 === "NVIDIA CUDA"
@@ -5,7 +5,7 @@
 # --8<-- [end:requirements]
 # --8<-- [start:set-up-using-python]

-vLLM-Omni depends on the matching major/minor release of vLLM. The vLLM-Omni 0.28.x release line uses vLLM 0.28.x.
+vLLM-Omni depends on the matching major/minor release of vLLM. The 0.29 development line uses vLLM 0.29.x. Published 0.28.0 wheels and images use vLLM 0.28.x.

 !!! note
    PyTorch installed via `conda` will statically link `NCCL` library, which can cause issues when vLLM tries to use `NCCL`. See <gh-issue:8420> for more details.
@@ -18,6 +18,8 @@ Therefore, it is recommended to install vLLM and vLLM-Omni with a **fresh new**

 #### Installation of vLLM

+These pre-built wheel instructions install the published vLLM-Omni 0.28.0 release. For the 0.29 development line, use the source-install instructions below.
+
 vLLM-Omni is built based on vLLM. Please install it with command below.
 ```bash
 uv pip install vllm==0.28.0 --torch-backend=auto
@@ -26,12 +28,12 @@ uv pip install vllm==0.28.0 --torch-backend=auto
 #### Installation of vLLM-Omni

 ```bash
-uv pip install vllm-omni
+uv pip install vllm-omni==0.28.0
 ```

 To run Gradio demos, also install the optional extras:
 ```bash
-uv pip install 'vllm-omni[demo]'
+uv pip install 'vllm-omni[demo]==0.28.0'
 ```

 # --8<-- [end:pre-built-wheels]
@@ -39,13 +41,13 @@ uv pip install 'vllm-omni[demo]'
 # --8<-- [start:build-wheel-from-source]

 #### Installation of vLLM
-If you do not need to modify source code of vLLM, you can directly install the stable 0.28.0 release version of the library
+If you do not need to modify source code of vLLM, you can directly install the stable 0.29.0 release version of the library

 ```bash
-uv pip install vllm==0.28.0 --torch-backend=auto
+uv pip install vllm==0.29.0 --torch-backend=auto
 ```

-The 0.28.0 release of vLLM ships CUDA 13.0-compatible binaries by default. If you need a different CUDA variant or want to reuse an existing PyTorch installation, build vLLM from source instead.
+The 0.29.0 release of vLLM ships CUDA 13.0-compatible binaries by default. If you need a different CUDA variant or want to reuse an existing PyTorch installation, build vLLM from source instead.

 #### Installation of vLLM-Omni
 Since vllm-omni is rapidly evolving, it's recommended to install it from source
@@ -66,12 +68,12 @@ If you want to check, modify or debug with source code of vLLM, install the libr
 ```bash
 git clone https://github.com/vllm-project/vllm.git
 cd vllm
-git checkout v0.28.0
+git checkout v0.29.0
 ```
 Set up environment variables to get pre-built wheels. If there are internet problems, just download the whl file manually. And set `VLLM_PRECOMPILED_WHEEL_LOCATION` as your local absolute path of whl file.
 ```bash
-#For CUDA 13.0 (the default for v0.28.0; the wheel filename has no `+cu130` suffix)
-export VLLM_PRECOMPILED_WHEEL_LOCATION=https://github.com/vllm-project/vllm/releases/download/v0.28.0/vllm-0.28.0-cp38-abi3-manylinux_2_28_x86_64.whl
+#For CUDA 13.0 (the default for v0.29.0; the wheel filename has no `+cu130` suffix)
+export VLLM_PRECOMPILED_WHEEL_LOCATION=https://github.com/vllm-project/vllm/releases/download/v0.29.0/vllm-0.29.0-cp38-abi3-manylinux_2_28_x86_64.whl
 ```
 Install vllm with command below (If you have no existing PyTorch).
 ```bash
@@ -131,7 +133,7 @@ If you want to specify the base vLLM version:
 ```bash
 DOCKER_BUILDKIT=1 docker build \
  -f docker/Dockerfile.cuda \
-  --build-arg BASE_IMAGE=vllm/vllm-openai:v0.28.0 \
+  --build-arg BASE_IMAGE=vllm/vllm-openai:v0.29.0 \
  -t vllm-omni-cuda .
 ```

@@ -11,6 +11,8 @@ vLLM-Omni current recommends the steps in under setup through Docker Images.

 #### Installation of vLLM

+These pre-built wheel instructions install the published vLLM-Omni 0.28.0 release. For the 0.29 development line, use the source-install instructions below.
+
 vLLM-Omni is built based on vLLM. Please install it with command below.
 ```bash
 uv pip install vllm==0.28.0+rocm723 --extra-index-url https://wheels.vllm.ai/rocm/0.28.0/rocm723
@@ -22,7 +24,7 @@ uv pip install vllm==0.28.0+rocm723 --extra-index-url https://wheels.vllm.ai/roc
 # we need to add --no-build-isolation as the torch
 # is not obtained from pypi, we have to install using the
 # torch installed in our environment
-uv pip install vllm-omni
+uv pip install vllm-omni==0.28.0

 # Optional if want to run Qwen3 TTS
 uv pip uninstall onnxruntime # should be removed before we can install onnxruntime-rocm
@@ -34,13 +36,13 @@ uv pip install onnxruntime-rocm
 # --8<-- [start:build-wheel-from-source]

 #### Installation of vLLM
-If you do not need to modify source code of vLLM, you can directly install the stable 0.28.0 release version of the library
+If you do not need to modify source code of vLLM, you can directly install the stable 0.29.0 release version of the library

 ```bash
-uv pip install vllm==0.28.0+rocm723 --extra-index-url https://wheels.vllm.ai/rocm/0.28.0/rocm723
+uv pip install vllm==0.29.0+rocm723 --extra-index-url https://wheels.vllm.ai/rocm/0.29.0/rocm723
 ```

-The pre-built 0.28.0 vLLM wheel targets ROCm 7.2.3. If you need a different ROCm stack or want to reuse an existing PyTorch installation, build vLLM from source instead.
+The pre-built 0.29.0 vLLM wheel targets ROCm 7.2.3. If you need a different ROCm stack or want to reuse an existing PyTorch installation, build vLLM from source instead.

 #### Installation of vLLM-Omni
 Since vllm-omni is rapidly evolving, it's recommended to install it from source
@@ -58,7 +60,7 @@ If you want to check, modify or debug with source code of vLLM, install the libr
 ```bash
 git clone https://github.com/vllm-project/vllm.git
 cd vllm
-git checkout v0.28.0
+git checkout v0.29.0
 python3 -m pip install -r requirements/rocm.txt
 python3 setup.py develop
 ```
@@ -7,13 +7,34 @@ For detailed hardware and software requirements, please refer to the [vLLM-Ascen

 The recommended way to use vLLM-Omni on NPU is through the vLLM-Ascend pre-built Docker images:

+=== "A2"
+
+    ```bash
+    export IMAGE=quay.io/atlas-ci/vllm-ascend:v0.28.0
+    ```
+
+=== "A3"
+
+    ```bash
+    export IMAGE=quay.io/atlas-ci/vllm-ascend:v0.28.0-a3
+    ```
+
+=== "A5"
+
+    ```bash
+    export IMAGE=quay.io/atlas-ci/vllm-ascend:v0.28.0-a5
+    ```
+
+=== "310P (Experimental)"
+
+    ```bash
+    export IMAGE=quay.io/atlas-ci/vllm-ascend:v0.28.0-310p
+    ```
+
 ```bash
-# vLLM-Ascend image
-# Atlas A2: quay.io/atlas-ci/vllm-ascend:v0.28.0
-# Atlas A3: quay.io/atlas-ci/vllm-ascend:v0.28.0-a3
 docker run --rm \
    --name vllm-omni-npu \
-    --shm-size=1g \
+    --shm-size=64g \
    --device /dev/davinci0 \
    --device /dev/davinci1 \
    --device /dev/davinci2 \
@@ -27,8 +48,8 @@ docker run --rm \
    -v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
    -v /etc/ascend_install.info:/etc/ascend_install.info \
    -v /root/.cache:/root/.cache \
-    -p 8000:8000 \
-    -it quay.io/atlas-ci/vllm-ascend:v0.28.0-a3 bash
+    -p 8091:8091 \
+    -it "$IMAGE" bash

 # Inside the container, install vLLM-Omni from source
 cd /vllm-workspace
@@ -72,53 +93,146 @@ Supported images as following.

 | image name | Hardware | OS |
 |-|-|-|
-| image-tag | Atlas A2 | Ubuntu |
-| image-tag-a3 | Atlas A3 | Ubuntu |
+| image-tag | A2 | Ubuntu |
+| image-tag-a3 | A3 | Ubuntu |
+| image-tag-a5 | A5 | Ubuntu |
+| image-tag-310p | 310P (experimental) | Ubuntu |

 Here's an example deployment command that has been verified on 4 x NPUs:

-```bash
-# Atlas A2: quay.io/ascend/vllm-omni:v0.28.0
-# Atlas A3: quay.io/ascend/vllm-omni:v0.28.0-a3
-docker run --rm \
-    --name vllm-omni-a3 \
-    --shm-size=4g \
-    --device /dev/davinci0 \
-    --device /dev/davinci1 \
-    --device /dev/davinci2 \
-    --device /dev/davinci3 \
-    --device /dev/davinci_manager \
-    --device /dev/devmm_svm \
-    --device /dev/hisi_hdc \
-    -v /usr/local/dcmi:/usr/local/dcmi \
-    -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
-    -v /usr/local/Ascend/driver/lib64/:/usr/local/Ascend/driver/lib64/ \
-    -v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
-    --device /dev/davinci_manager \
-    --device /dev/devmm_svm \
-    --device /dev/hisi_hdc \
-    -v /usr/local/dcmi:/usr/local/dcmi \
-    -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
-    -v /usr/local/Ascend/driver/lib64/:/usr/local/Ascend/driver/lib64/ \
-    -v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
-    -v /etc/ascend_install.info:/etc/ascend_install.info \
-    -v ~/.cache:/root/.cache \
-    -p 8091:8091 \
-    -it quay.io/ascend/vllm-omni:v0.28.0-a3 bash
-```
+=== "A2"
+
+    ```bash
+    export IMAGE=quay.io/ascend/vllm-omni:v0.28.0
+    docker run --rm \
+        --name vllm-omni-a2 \
+        --shm-size=64g \
+        --device /dev/davinci0 \
+        --device /dev/davinci1 \
+        --device /dev/davinci2 \
+        --device /dev/davinci3 \
+        --device /dev/davinci_manager \
+        --device /dev/devmm_svm \
+        --device /dev/hisi_hdc \
+        -v /usr/local/dcmi:/usr/local/dcmi \
+        -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
+        -v /usr/local/Ascend/driver/lib64/:/usr/local/Ascend/driver/lib64/ \
+        -v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
+        -v /etc/ascend_install.info:/etc/ascend_install.info \
+        -v ~/.cache:/root/.cache \
+        -p 8091:8091 \
+        -it "$IMAGE" bash
+    ```
+
+=== "A3"
+
+    ```bash
+    export IMAGE=quay.io/ascend/vllm-omni:v0.28.0-a3
+    docker run --rm \
+        --name vllm-omni-a3 \
+        --shm-size=64g \
+        --device /dev/davinci0 \
+        --device /dev/davinci1 \
+        --device /dev/davinci2 \
+        --device /dev/davinci3 \
+        --device /dev/davinci_manager \
+        --device /dev/devmm_svm \
+        --device /dev/hisi_hdc \
+        -v /usr/local/dcmi:/usr/local/dcmi \
+        -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
+        -v /usr/local/Ascend/driver/lib64/:/usr/local/Ascend/driver/lib64/ \
+        -v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
+        -v /etc/ascend_install.info:/etc/ascend_install.info \
+        -v ~/.cache:/root/.cache \
+        -p 8091:8091 \
+        -it "$IMAGE" bash
+    ```
+
+=== "A5"
+
+    ```bash
+    export IMAGE=quay.io/ascend/vllm-omni:v0.28.0-a5
+    docker run --rm \
+        --name vllm-omni-a5 \
+        --shm-size=64g \
+        --device /dev/davinci0 \
+        --device /dev/davinci1 \
+        --device /dev/davinci2 \
+        --device /dev/davinci3 \
+        --device /dev/davinci_manager \
+        --device /dev/devmm_svm \
+        --device /dev/hisi_hdc \
+        -v /usr/local/dcmi:/usr/local/dcmi \
+        -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
+        -v /usr/local/Ascend/driver/lib64/:/usr/local/Ascend/driver/lib64/ \
+        -v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
+        -v /etc/ascend_install.info:/etc/ascend_install.info \
+        -v ~/.cache:/root/.cache \
+        -p 8091:8091 \
+        -it "$IMAGE" bash
+    ```
+
+=== "310P (Experimental)"
+
+    ```bash
+    export IMAGE=quay.io/ascend/vllm-omni:v0.28.0-310p
+    docker run --rm \
+        --name vllm-omni-310p \
+        --shm-size=64g \
+        --device /dev/davinci0 \
+        --device /dev/davinci1 \
+        --device /dev/davinci2 \
+        --device /dev/davinci3 \
+        --device /dev/davinci_manager \
+        --device /dev/devmm_svm \
+        --device /dev/hisi_hdc \
+        -v /usr/local/dcmi:/usr/local/dcmi \
+        -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
+        -v /usr/local/Ascend/driver/lib64/:/usr/local/Ascend/driver/lib64/ \
+        -v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
+        -v /etc/ascend_install.info:/etc/ascend_install.info \
+        -v ~/.cache:/root/.cache \
+        -p 8091:8091 \
+        -it "$IMAGE" bash
+    ```

 !!! tip
    You can use this docker image to serve models the same way you would with in vLLM! To do so, make sure you overwrite the default entrypoint (`vllm serve --omni`) which works only for models supported in the vLLM-Omni project.

-Or build IMAGE from **source code**:
+Or build the image from **source code**:

 ```bash
 git clone https://github.com/vllm-project/vllm-omni.git
 cd vllm-omni
-# A2
-# docker build -t vllm-omni-dev-image:latest -f ./docker/Dockerfile.npu .
-# A3
-docker build -t vllm-omni-dev-image:latest -f ./docker/Dockerfile.npu.a3 .
 ```

+All hardware variants use `docker/Dockerfile.npu`. Choose the device suffix for your hardware:
+
+=== "A2"
+
+    ```bash
+    docker build -t vllm-omni-dev-image:latest -f docker/Dockerfile.npu .
+    ```
+
+=== "A3"
+
+    ```bash
+    docker build -t vllm-omni-dev-image:latest -f docker/Dockerfile.npu \
+        --build-arg VLLM_ASCEND_DEVICE_SUFFIX=-a3 .
+    ```
+
+=== "A5"
+
+    ```bash
+    docker build -t vllm-omni-dev-image:latest -f docker/Dockerfile.npu \
+        --build-arg VLLM_ASCEND_DEVICE_SUFFIX=-a5 .
+    ```
+
+=== "310P (Experimental)"
+
+    ```bash
+    docker build -t vllm-omni-dev-image:latest -f docker/Dockerfile.npu \
+        --build-arg VLLM_ASCEND_DEVICE_SUFFIX=-310p .
+    ```
+
 # --8<-- [end:pre-built-images]
@@ -19,10 +19,10 @@ uv venv --python 3.12 --seed
 source .venv/bin/activate

 # On CUDA
-uv pip install vllm==0.28.0 --torch-backend=auto
+uv pip install vllm==0.29.0 --torch-backend=auto

 # On ROCm
-uv pip install vllm==0.28.0+rocm723 --extra-index-url https://wheels.vllm.ai/rocm/0.28.0/rocm723
+uv pip install vllm==0.29.0+rocm723 --extra-index-url https://wheels.vllm.ai/rocm/0.29.0/rocm723

 git clone https://github.com/vllm-project/vllm-omni.git
 cd vllm-omni
@@ -34,7 +34,7 @@ For additional installation methods — please see the [installation guide](inst
 !!! note
    It is important to install the same major & minor version of vLLM and vLLM Omni, otherwise things may not work as expected. If the versions are misaligned, you will see a warning when you import vLLM Omni.

-    If you are seeing strange behavior with the `vllm` command not handling the `--omni` flag correctly, you most likely have a version mismatch with vLLM < `0.28.0` and vLLM Omni `0.28.0`, as vLLM Omni no longer hijacks the vLLM entrypoint. Updating vLLM should resolve this issue.
+    If you are seeing strange behavior with the `vllm` command not handling the `--omni` flag correctly, you most likely have a version mismatch with vLLM < `0.29.0` and vLLM Omni `0.29.0`, as vLLM Omni no longer hijacks the vLLM entrypoint. Updating vLLM should resolve this issue.

 ## Offline Inference

@@ -12,7 +12,7 @@ implementations to reduce numerical drift and sampling divergence.
 ## Installation

 ```bash
-pip install kernels==0.15.2
+pip install kernels==0.16.1
 ```

 If `kernels` is unavailable, an explicit `FLASH_ATTN_HUB` or
@@ -199,12 +199,21 @@ curl -X POST http://localhost:8091/v1/videos \
 | `boundary_ratio`      | float  | None    | Boundary split ratio for low/high DiT (Wan2.2)   |
 | `flow_shift`          | float  | None    | Scheduler flow shift (Wan2.2)                    |
 | `seed`                | int    | None    | Random seed (reproducible)                       |
+| `num_outputs_per_prompt` | int  | 1       | Number of videos to generate (1-10, MiniMax H3); only the first is returned, see below |
 | `lora`                | object | None    | LoRA configuration                               |
 | `enable_frame_interpolation` | bool | false | Enable RIFE frame interpolation before MP4 encoding |
 | `frame_interpolation_exp` | int | 1 | Interpolation exponent; 1=2x temporal resolution, 2=4x |
 | `frame_interpolation_scale` | float | 1.0 | RIFE inference scale; use 0.5 for high-resolution inputs |
 | `frame_interpolation_model_path` | str | None | Local directory or Hugging Face repo ID with `flownet.pkl`; defaults to `elfgum/RIFE-4.22.lite` |

+!!! note "Only the first output is returned"
+    `num_outputs_per_prompt` is forwarded to the pipeline, so the model does
+    generate that many videos and you pay the generation cost for all of them.
+    `/v1/videos` and `/v1/videos/sync` then return only the **first** one,
+    because an async video job stores a single file per video id. The server
+    logs a warning when it discards the extra outputs. Keep this at `1`
+    unless you have a specific reason to generate videos you will not receive.
+
 ## Frame Interpolation

 Frame interpolation is an optional post-processing step for `/v1/videos` and
@@ -161,7 +161,7 @@ where = ["."]
 include = ["vllm_omni*"]

 [tool.setuptools.package-data]
-"vllm_omni" = ["_version.py", "py.typed"]
+"vllm_omni" = ["_version.py", "py.typed", "transformers_utils/chat_templates/*.jinja"]

 [tool.setuptools_scm]
 # no extra settings needed, presence enables setuptools-scm
@@ -283,6 +283,8 @@ markers = [
    "S5000: [hardware-resource] [musa] Tests that require S5000 GPU (Moore Threads/MUSA)",
    "A2: [hardware-resource] [npu] Tests that require A2 NPU",
    "A3: [hardware-resource] [npu] Tests that require A3 NPU",
+    "A5: [hardware-resource] [npu] Tests that require A5 NPU",
+    "310P: [hardware-resource] [npu] Tests that require 310P NPU",
    # Tag with [hardware-cards] for tests.helpers.mark.get_supported_card_counts().
    "cards_1: [hardware-cards] Tests that require 1 accelerator card",
    "cards_2: [hardware-cards] Tests that require 2 accelerator cards",
@@ -152,7 +152,7 @@ The following configuration was validated:
 - vLLM-Omni version or commit: Use a commit that contains TorchAO FP8
  checkpoint loading for diffusion models
 - TorchAO version: 0.17.0
- `kernels` / `kernels-data`: 0.15.2 / 0.16.1
+- `kernels` / `kernels-data`: 0.16.1 / matching kernels-data

 #### Command

@@ -7,13 +7,14 @@ recipes it uses **no offload of any kind** — see the capacity note below.
 Validated on:

 - Host: DGX Spark (GB10), aarch64
- GPUs: 1 (unified memory)
- Driver: <FILL_DRIVER_VERSION>
+- GPUs: 1, or 2 hosts with one unified-memory GPU each
+- Driver: 580.173.02 for the two-host qualification
 - vLLM: 0.26.0
 - vLLM-Omni: `main` at `e1aa6eae75c460cd1893bc320546e81e66973831`
- Workloads: T2VA and Ref2VA, 960x576, `duration=8.0`, 50 steps, `flow_shift=12`,
-  `seed=1101`
-
+- One-host workloads: T2VA and Ref2VA, 960x576, `duration=8.0`, 50 steps
+- Two-host workload: T2VA, 1344x768, `duration=5.0` and `duration=10.0`,
+  50 steps, text-encoder TP2, and DiT USP2
+- Common sampling settings: `flow_shift=12`, `seed=1101`

 ## Capacity requirements

@@ -24,6 +25,7 @@ Validated on:
 | Checkpoint storage | 135 GiB per partition |
 | Measured peak allocator high-water, T2VA (FP8) | 97.7 GiB |
 | Measured peak allocator high-water, Ref2VA (FP8) | 102.8 GiB |
+| Measured peak per rank, two-host T2VA 5 s / 10 s (FP8) | 75.45 / 87.17 GiB |

 `FL2VA` and `Ref2VA` are separate 135 GiB checkpoint partitions. Start one server
 at a time.
@@ -164,6 +166,41 @@ curl --fail-with-body -sS -X POST http://127.0.0.1:8000/v1/videos/sync \
 The response is `200 OK` with an MP4 body: H.264 at the requested geometry plus a
 32 kHz stereo AAC track. `ffprobe` reports 8.032 s for `duration=8.0`.

+## Two GB10 hosts: T2VA at 1344x768
+
+The two-host qualification uses one GB10 GPU per DGX Spark, connected through a
+single ConnectX-7 RoCE interface. The FL2VA checkpoint and software revision are
+identical on both hosts. The distributed configuration is:
+
+- world size 2;
+- text-encoder tensor parallel size 2;
+- DiT Ulysses sequence parallel size 2 and ring degree 1;
+- VAE patch parallel size 1;
+- online FP8, eager execution, cuDNN attention, and tiled VAE decode.
+
+The current diffusion multiprocess executor starts all ranks on one host and sets
+`MASTER_ADDR=localhost`. These measurements therefore used an external `env://`
+launcher with one rank on each host and a benchmark-only adaptation that maps
+both hosts' local worker to `cuda:0`. This qualifies the two-rank MiniMax-H3
+pipeline and its output, but it is not yet a supported cross-host `vllm serve`
+command. Do not infer that the single-host launch command above works unchanged
+across two machines.
+
+One 5-second, 2-step request warmed both workers before the measured runs. The
+reported pipeline elapsed time covers distributed model execution through GPU
+completion; MP4 encoding happened afterward and is not included.
+
+| Duration | Steps | Text encode | Denoise | Per step | VAE decode | Pipeline elapsed | Max peak per rank |
+| --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: |
+| 5 s | 50 | 0.21 s | 1570.19 s | 31.40 s | 79.45 s | 1649.87 s (27 min 30 s) | 75.45 GiB |
+| 10 s | 50 | 0.15 s | 4448.12 s | 88.96 s | 160.71 s | 4609.00 s (76 min 49 s) | 87.17 GiB |
+
+Both ranks reported stage times and allocator peaks within measurement noise.
+The generated outputs were H.264 at 1344x768 and 24 FPS with 32 kHz stereo AAC:
+124 frames (5.167 s) and 243 frames (10.125 s). Both files decoded successfully,
+and neither run encountered an OOM or NCCL error. These are single qualification
+runs (`n=1`), not throughput guarantees.
+
 ## Measured latency

 Single GB10, FP8, eager, cuDNN attention, tiled VAE, one request at a time:
@@ -238,9 +275,11 @@ shorter than a 50-step Ref2VA run.

 ## Known limitations

- One GPU means no Ulysses, ring, TP, or VAE patch parallelism. Pass
+- With one GPU there is no Ulysses, ring, TP, or VAE patch parallelism. Pass
  `--usp 1 --ring 1 --vae-patch-parallel-size 1` explicitly if a profile or
  script would otherwise supply higher degrees.
+- The two-host results require the external launch adaptation described above;
+  cross-host diffusion worker launch is not exposed by the documented server CLI.
 - `--enforce-eager` is used because regional compile adds startup time without a
  parallelism win on a single rank.
 - `t2va` rejects requests without an explicit `aspect_ratio`. When `width` and
@@ -255,8 +294,7 @@ shorter than a 50-step Ref2VA run.
  adding references or raising the output shape.
 - Ref2VA requires a vLLM-Omni build newer than `v0.26.0`. See the version note
  under **Validated on**: FP8 weight loading and image-only Ref2VA are both
-  broken on `release/v0.26.0`.suggestion is  `v0.26.1`.
+  broken on `release/v0.26.0`. The suggested version is `v0.26.1`.
 - Online FP8 is incompatible with layerwise offload — the offload path produces a
  weight stride the Cutlass FP8 kernel rejects. This is not a practical
  restriction here since offload is unusable on GB10 anyway.
-```
@@ -28,13 +28,13 @@ also provided.

 - Default deploy configs (auto-loaded by HF `model_type=minicpmo` +
  `hf_config.version="4.5"`):
-  - Default single-GPU compatibility layout (auto-loaded):
+    - Default single-GPU compatibility layout (auto-loaded):
    [`vllm_omni/deploy/minicpmo_4_5.yaml`](../../vllm_omni/deploy/minicpmo_4_5.yaml)
-  - Recommended 2-GPU continuous-batching layout:
+    - Recommended 2-GPU continuous-batching layout:
    [`vllm_omni/deploy/minicpmo_4_5_2gpu.yaml`](../../vllm_omni/deploy/minicpmo_4_5_2gpu.yaml),
-  - 3-GPU layout:
+    - 3-GPU layout:
    [`vllm_omni/deploy/minicpmo_4_5_3gpu.yaml`](../../vllm_omni/deploy/minicpmo_4_5_3gpu.yaml)
-  - 8x RTX 4090 layout:
+    - 8x RTX 4090 layout:
    [`vllm_omni/deploy/minicpmo_4_5_8x4090.yaml`](../../vllm_omni/deploy/minicpmo_4_5_8x4090.yaml)
 - Online example + Gradio demo:
  [`examples/online_serving/minicpmo/`](../../examples/online_serving/minicpmo/)
@@ -126,9 +126,20 @@ video. Use the two-GPU profile for production throughput.
 ```bash
 vllm serve openbmb/MiniCPM-o-4_5 --omni \
    --trust-remote-code \
+    --chat-template vllm_omni/transformers_utils/chat_templates/minicpmo45_native.jinja \
+    --chat-template-content-format openai \
    --host 0.0.0.0 --port 8099
 ```

+Run these commands from the repository root. The bundled chat template
+matches the Hugging Face `chat(omni_mode=True)` content layout: media and text
+parts are concatenated without adding separators, while whitespace explicitly
+included in a text part is preserved. This matters for speech generation because
+an inserted newline changes the Thinker condition passed to Talker. For image
+questions, place the image before the question, as in the native image-chat
+layout. Continue to request `use_tts_template: true` and
+`enable_thinking: false` for spoken answers.
+
 The deploy config is auto-loaded by the model registry — no
 `--deploy-config` flag needed for this default single-GPU layout.

@@ -271,6 +282,8 @@ uses GPU 5. GPUs 6–7 are left free.
 vllm serve openbmb/MiniCPM-o-4_5 --omni \
    --deploy-config vllm_omni/deploy/minicpmo_4_5_8x4090.yaml \
    --trust-remote-code \
+    --chat-template vllm_omni/transformers_utils/chat_templates/minicpmo45_native.jinja \
+    --chat-template-content-format openai \
    --host 0.0.0.0 --port 8099
 ```

@@ -263,6 +263,94 @@ curl -X POST http://localhost:8091/v1/audio/speech \
      gpu_memory_utilization: 0.15
  ```

+### 1x A100 40GB (0.6B CustomVoice)
+
+#### Environment
+
+- OS: Linux 4.18.0-553 (RHEL 8.10), x86_64
+- Python: 3.12.14
+- PyTorch: 2.13.0+cu129
+- Driver / runtime: NVIDIA 570.124.06 / CUDA 12.8 (`nvidia-smi`); PyTorch CUDA 12.9
+- GPU: NVIDIA A100-SXM4-40GB, 40960 MiB
+- vLLM version: 0.28.0+cu129
+- vLLM-Omni version or commit: `3204a0b2ade0f05424b7f11e3bcc03cb3542e0c6`
+
+#### Command
+
+```bash
+vllm serve Qwen/Qwen3-TTS-12Hz-0.6B-CustomVoice \
+    --deploy-config vllm_omni/deploy/qwen3_tts.yaml \
+    --omni --port 8091
+```
+
+The default deploy config (`qwen3_tts.yaml`) works without modification on this
+A100 40GB. Both stages (talker + code2wav) share GPU 0 with
+`gpu_memory_utilization: 0.3` each.
+
+#### Verification
+
+**English synthesis:**
+
+```bash
+curl -X POST http://localhost:8091/v1/audio/speech \
+    -H "Content-Type: application/json" \
+    -d '{
+        "input": "Hello, this is Qwen3-TTS running on A100 40GB.",
+        "voice": "vivian",
+        "language": "English"
+    }' --output test_english.wav
+```
+
+**Chinese synthesis:**
+
+```bash
+curl -X POST http://localhost:8091/v1/audio/speech \
+    -H "Content-Type: application/json" \
+    -d '{
+        "input": "你好，这是在A100 40GB上运行的语音合成测试。",
+        "voice": "vivian",
+        "language": "Chinese"
+    }' --output test_chinese.wav
+```
+
+**With emotion instruction:**
+
+```bash
+curl -X POST http://localhost:8091/v1/audio/speech \
+    -H "Content-Type: application/json" \
+    -d '{
+        "input": "I am so excited about this!",
+        "voice": "vivian",
+        "language": "English",
+        "instructions": "Speak with great enthusiasm"
+    }' --output test_emotion.wav
+```
+
+#### Notes
+
+- Memory usage: **16283 MiB / 40960 MiB** peak (`nvidia-smi`) during the three
+  verification requests with the default deploy config
+  (`gpu_memory_utilization: 0.3` per stage). Stage 0 used 12404 MiB and stage 1
+  used 3862 MiB. The 0.6B weights occupy only ~2.4 GiB (Stage 0: 1.91 GiB,
+  Stage 1: 0.45 GiB); the rest is KV cache reserved as a fraction of device
+  memory, so the same `0.3` setting reads higher in GiB on 40GB than the
+  ~13.5 GiB idle figure on the 24GB 4090 profile above. To shrink the
+  reservation, use the `gpu_memory_utilization: 0.15` override documented in
+  that 4090 section.
+- vLLM 0.28.0's default wheel targets CUDA 13.0. Driver 570 cannot load that
+  variant. This qualification used the official CUDA 12.9 wheel (same pin as
+  the A100 Wan2.2 recipe):
+
+  ```bash
+  uv pip install \
+    'https://github.com/vllm-project/vllm/releases/download/v0.28.0/vllm-0.28.0%2Bcu129-cp38-abi3-manylinux_2_28_x86_64.whl' \
+    --torch-backend=cu129
+  ```
+
+  Confirm `torch.__version__` contains `cu129`. If a CPU build is already
+  installed, `uv pip install --torch-backend=cu129` is a no-op; use
+  `--reinstall`.
+
 ### 1x AMD MI300X, 1.7B checkpoints

 #### Environment
@@ -42,6 +42,7 @@ recipes/
 | -------- | ------ | ---------- |
 | [`Baidu/ERNIE-Image.md`](./Baidu/ERNIE-Image.md) | Text-to-image online serving (ERNIE-Image 8B) | 1x or 2x RTX 4090 24GB |
 | [`Bagel/BAGEL-7B-MoT.md`](./Bagel/BAGEL-7B-MoT.md) | Text-to-image with shared online/offline examples | 1x A100 80GB / 2x CUDA GPUs |
+| [`black-forest-labs/FLUX.1-schnell.md`](./black-forest-labs/FLUX.1-schnell.md) | Offline text-to-image generation with CPU offload | 1x NVIDIA GeForce RTX 5090 32GB |
 | [`Boogu/Boogu-Image.md`](./Boogu/Boogu-Image.md) | Text-to-image online serving (Boogu-Image-0.1-Base) | 1x A100/H100 40GB+ |
 | [`BosonAI/Higgs-Audio-V3-TTS.md`](./BosonAI/Higgs-Audio-V3-TTS.md) | Online + offline multilingual TTS with voice cloning | 1x H100 80GB |
 | [`ByteDance/Lance.md`](./ByteDance/Lance.md) | Unified AR+diffusion: text/img/video gen + understanding (Lance 3B) | 1x B300 / A100 80GB |
@@ -54,6 +55,7 @@ recipes/
 | [`IndexTeam/IndexTTS-2.md`](./IndexTeam/IndexTTS-2.md) | Online serving for voice-cloned TTS with optional emotion control | 1x L4 24GB or larger CUDA GPU |
 | [`IndexTeam/IndexTTS-2_5.md`](./IndexTeam/IndexTTS-2_5.md) | Online + offline multilingual voice-cloned TTS with native speed and emotion control | 1x NVIDIA H20 96GB |
 | [`MiniMaxAI/MiniMax-H3.md`](./MiniMaxAI/MiniMax-H3.md) | T2VA, FL2VA, image+audio Ref2VA, and multi-video Ref2VA serving | 1x GPU with CPU offload / 4x B300 |
+| [`MiniMaxAI/MiniMax-H3-Spark-GB10.md`](./MiniMaxAI/MiniMax-H3-Spark-GB10.md) | MiniMax H3 T2VA and Ref2VA on DGX Spark | 1x or 2x DGX Spark (GB10) |
 | [`MiniMaxAI/MiniMax-H3-MUSA.md`](./MiniMaxAI/MiniMax-H3-MUSA.md) | MiniMax H3 T2VA, FL2VA, and Ref2VA on MUSA | 4x MTT S5000 (TP4 Ref2VA profile) |
 | [`krea/Krea-2.md`](./krea/Krea-2.md) | Text-to-image (Turbo + Raw), offline + online, with LoRA | 1x H100 80GB |
 | [`LTX/LTX-2.md`](./LTX/LTX-2.md) | LTX-2/LTX-2.3 text-to-video and image-to-video with synchronized audio | H200 141GB / 96GB-class GPU |
@@ -0,0 +1,160 @@
+# FLUX.1-schnell
+
+> Offline text-to-image generation on one NVIDIA GeForce RTX 5090 32 GB
+
+## Summary
+
+- Vendor: Black Forest Labs
+- Model: `black-forest-labs/FLUX.1-schnell`
+- Task: Text-to-image generation
+- Mode: Offline inference with the shared text-to-image runner
+- Maintainer: Community
+
+## When to use this recipe
+
+Use this recipe as a personally validated starting point for generating a
+1024x1024 image with FLUX.1-schnell on one NVIDIA GeForce RTX 5090. The tested
+configuration uses BF16, four denoising steps and model-level CPU offload.
+
+This recipe covers single-image generation. It does not attempt to optimize
+throughput with quantization or cache acceleration.
+
+## References
+
+- Model card: <https://huggingface.co/black-forest-labs/FLUX.1-schnell>
+- Canonical offline text-to-image guide:
+  [`docs/user_guide/examples/offline_inference/text_to_image.md`](../../docs/user_guide/examples/offline_inference/text_to_image.md)
+- Shared runnable example:
+  [`examples/offline_inference/text_to_image`](../../examples/offline_inference/text_to_image)
+- CPU offload guide:
+  [`docs/user_guide/diffusion/cpu_offload.md`](../../docs/user_guide/diffusion/cpu_offload.md)
+- Community recipe tracker:
+  [#2645](https://github.com/vllm-project/vllm-omni/issues/2645)
+
+## Hardware Support
+
+## GPU
+
+### 1x NVIDIA GeForce RTX 5090 32 GB
+
+This configuration was personally validated end to end on September 9, 2026.
+
+#### Environment
+
+- OS: Ubuntu 22.04.3 x86_64
+- Python: 3.12.7
+- GPU: 1x NVIDIA GeForce RTX 5090, 32,607 MiB reported by `nvidia-smi`
+- Host RAM: approximately 62 GiB (64 GB class)
+- NVIDIA driver: 595.71.05; CUDA 13.2 reported by `nvidia-smi`
+- PyTorch: 2.13.0+cu130; PyTorch CUDA runtime: 13.0
+- vLLM: 0.28.0
+- vLLM-Omni: `0.28.1.dev112+ga31690041`
+- Diffusers: 0.40.0
+- Transformers: 5.14.1
+
+#### Command
+
+Run the shared example from the vLLM-Omni repository root in an environment
+with vLLM-Omni installed:
+
+```bash
+python examples/offline_inference/text_to_image/text_to_image.py \
+  --model black-forest-labs/FLUX.1-schnell \
+  --prompt 'A red fox standing in fresh snow, cinematic photography' \
+  --height 1024 --width 1024 \
+  --num-inference-steps 4 --guidance-scale 0 --seed 42 \
+  --enable-cpu-offload \
+  --output flux-schnell-rtx5090.png
+```
+
+The command uses the public model ID for portability. The measured runs used
+already downloaded FLUX.1-schnell weights; initial Hub download and
+authentication were not part of the inference validation.
+
+#### Verification
+
+Verify that the runner succeeded and inspect the saved image:
+
+```bash
+python - <<'PY'
+import hashlib
+from pathlib import Path
+
+from PIL import Image, ImageStat
+
+path = Path("flux-schnell-rtx5090.png")
+with Image.open(path) as image:
+    image.load()
+    assert image.format == "PNG"
+    assert image.size == (1024, 1024)
+    rgb = image.convert("RGB")
+    stats = ImageStat.Stat(rgb)
+    assert max(stats.stddev) > 1, "Image is nearly uniform"
+    print(f"format={image.format}")
+    print(f"size={image.width}x{image.height}")
+    print(f"mode={image.mode}")
+    print(f"bytes={path.stat().st_size}")
+    print("rgb_mean=" + ",".join(f"{value:.3f}" for value in stats.mean))
+    print(f"rgb_extrema={rgb.getextrema()}")
+
+print(f"sha256={hashlib.sha256(path.read_bytes()).hexdigest()}")
+PY
+```
+
+The repeated experiment completed 36 image requests (6 additional warmups and
+30 measured outputs) with successful process exits. Every PNG passed decoding,
+1024x1024 dimension and non-uniform pixel checks. SHA-256 hashes were recorded
+for the artifacts; matching those hashes is not required across environments.
+These checks establish file integrity, not a quantitative image-quality score.
+
+#### Performance
+
+The experiment used three independent processes. Each process initialized the
+model once, generated two warmup images, then generated ten measured images
+sequentially. All 30 measured images used BF16, CPU offload, 1024x1024, four
+steps, guidance scale 0 and TP=1, without concurrent requests.
+
+| Generation measurement | Result |
+| --- | ---: |
+| Measured images | 30 |
+| Mean latency | 25.8411 seconds |
+| P95 latency (nearest-rank) | 26.1446 seconds |
+
+| Startup / resource measurement | Maximum |
+| --- | ---: |
+| Model initialization (seconds) | 64.73 |
+| Full-lifetime RAM PSS peak (MiB) | 49306 |
+| Full-lifetime device GPU memory peak (MiB) | 28121 |
+| Measured PyTorch reserved GPU memory peak (MiB) | 27360 |
+
+Latency measures only `Omni.generate()`, excluding image saving and
+verification. The two warmup images per process are excluded. P95 uses the
+nearest-rank method over the 30 measured images. Model initialization includes
+the engine's built-in dummy warmup; weights were already downloaded and OS file
+caches were not cleared.
+
+RAM and whole-device GPU memory were sampled throughout each process lifetime
+at a target interval of 200 ms. RAM is the maximum simultaneous sum of PSS for
+the inference process and its workers, read from Linux `smaps_rollup`.
+Whole-device GPU memory is the maximum reported by `nvidia-smi`. Both are
+sampled peaks and may miss shorter spikes.
+
+The PyTorch value is the maximum worker-reported `peak_memory_mb` across the
+30 measured requests. Despite the field name, its unit is MiB; it measures the
+PyTorch allocator's reserved-memory high-water mark rather than whole-device
+usage.
+
+#### Notes
+
+- `--enable-cpu-offload` swaps the transformer and text encoders between host
+  and device memory at phase boundaries. It reduces GPU residency but adds
+  transfer overhead and requires sufficient host RAM. The minimum host RAM
+  requirement was not measured.
+- CPU offload does not imply CPU-only computation. The tested model-level
+  offload backend keeps the VAE on the GPU.
+- Only offline single-image generation on one GPU with BF16, 1024x1024,
+  four steps and CPU offload was validated. Online serving, FP8, other
+  resolutions, concurrency, execution without offload and smaller host RAM
+  configurations are outside this recipe's validated scope.
+- The upstream revision of the downloaded weights was not recorded, so the
+  measurements do not establish exact checkpoint reproducibility.
@@ -9,13 +9,15 @@
 # That breakage is now handled in code instead of by pinning:
 # MossTTSNanoForGeneration wraps weight loading in
 # transformers_keys_to_ignore_compat (vllm_omni/model_executor/models/utils.py).
-transformers >= 5.10.1, < 5.15
+# Floor is 5.13.0 so transformers.integrations.hub_kernels LayerRepository
+# entries always pass version=/revision= required by kernels>=0.16 (#6971).
+transformers >= 5.13.0, < 5.15
 av>=14.0.0
 omegaconf>=2.3.0
 diffusers==0.40.0
-# Hub kernels, including LTX-2.5 NATTEN. Keep this pin in sync with the
-# supported transformers range above.
-kernels==0.15.2
+# Hub kernels, including LTX-2.5 NATTEN. 0.16+ resolves torch-stable-abi*
+# FA3 builds on torch 2.13 / cu130; keep in sync with transformers above.
+kernels==0.16.1
 safetensors>=0.8.0
 accelerate==1.12.0
 soundfile>=0.13.1
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 import pytest
 import torch
@@ -214,11 +214,20 @@ def test_fish_kvcache_backend_wraps_only_model_instance(monkeypatch):
    assert fish_kvcache_backend.get_fish_kvcache_attn_stats()["small_hit_count"] == 1


-def _vllm_kv_cache_shape(num_blocks, block_size, num_kv_heads, head_size):
+def _vllm_kv_cache_shape(num_blocks, block_size, num_kv_heads, head_size, dtype=torch.float16):
    # Source of truth for the paged KV cache layout the backend must unpack.
-    from vllm.v1.attention.backends.flash_attn import FlashAttentionBackend
+    # vLLM 0.29 removed AttentionBackend.get_kv_cache_shape; the per-layer shape
+    # now comes from the spec, with the content dim expressed in bytes.
+    from vllm.v1.kv_cache_interface import FullAttentionSpec, compute_layer_kv_cache_shape_bytes

-    return FlashAttentionBackend.get_kv_cache_shape(num_blocks, block_size, num_kv_heads, head_size)
+    spec = FullAttentionSpec(
+        block_size=block_size,
+        num_kv_heads=num_kv_heads,
+        head_size=head_size,
+        dtype=dtype,
+    )
+    num_pages, num_heads, num_states, content_bytes = compute_layer_kv_cache_shape_bytes(spec, num_blocks)
+    return (num_pages, num_heads, num_states, content_bytes // dtype.itemsize)


 def test_fish_kvcache_backend_unbinds_kv_on_vllm_cache_layout(monkeypatch):
@@ -0,0 +1,156 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
+from __future__ import annotations
+
+import os
+import shutil
+import subprocess
+import sys
+from pathlib import Path
+
+import pytest
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+
+REPO_ROOT = Path(__file__).resolve().parents[2]
+BOOTSTRAP_PATH = REPO_ROOT / ".buildkite/amd/scripts/bootstrap-amd-omni.sh"
+
+
+def _run_bootstrap(
+    tmp_path: Path,
+    *,
+    debug_test_yaml: str | None = None,
+    curl_status: int = 0,
+    jq_status: int = 0,
+    labels: tuple[str, ...] = (),
+    pull_request: bool = True,
+) -> subprocess.CompletedProcess[str]:
+    for relative_path in (
+        ".buildkite/common/scripts/resolve_skip_ci.sh",
+        ".buildkite/amd/scripts/select_test_suites.py",
+    ):
+        destination = tmp_path / relative_path
+        destination.parent.mkdir(parents=True, exist_ok=True)
+        shutil.copyfile(REPO_ROOT / relative_path, destination)
+
+    # Stop after real suite selection, before rendering or uploading a pipeline.
+    (tmp_path / ".buildkite/common/scripts/skip_ci.py").write_text("print('test-skip')\n", encoding="utf-8")
+    env = {
+        "PATH": os.environ["PATH"],
+        "BUILDKITE_PULL_REQUEST": "123" if pull_request else "false",
+        "BUILDKITE_BRANCH": "feature" if pull_request else "main",
+        "TEST_PYTHON": sys.executable,
+        "TEST_CURL_STATUS": str(curl_status),
+        "TEST_JQ_STATUS": str(jq_status),
+        "TEST_PR_LABELS": "\n".join(labels),
+    }
+    if debug_test_yaml is not None:
+        env["DEBUG_TEST_YAML"] = debug_test_yaml
+
+    # Source the unmodified bootstrap with subprocess-local command doubles.
+    script = r"""
+curl() {
+    echo 'curl called' >&2
+    return "$TEST_CURL_STATUS"
+}
+jq() {
+    echo 'jq called' >&2
+    printf '%s\n' "$TEST_PR_LABELS"
+    return "$TEST_JQ_STATUS"
+}
+git() { :; }
+python3() { "$TEST_PYTHON" "$@"; }
+source "$1"
+"""
+    return subprocess.run(
+        ["bash", "-c", script, "test-amd-bootstrap", str(BOOTSTRAP_PATH)],
+        cwd=tmp_path,
+        env=env,
+        capture_output=True,
+        text=True,
+        timeout=10,
+        check=False,
+    )
+
+
+@pytest.mark.parametrize(("curl_status", "jq_status"), [(22, 0), (0, 4)], ids=["request-failure", "parse-failure"])
+@pytest.mark.parametrize(
+    ("debug_test_yaml", "expected_levels"),
+    [
+        ("ready", ["l2"]),
+        ("merge", ["l3"]),
+        ("ready,merge", ["l2", "l3"]),
+        (" MERGE, ready,", ["l3", "l2"]),
+    ],
+)
+def test_debug_override_skips_label_lookup(tmp_path, debug_test_yaml, expected_levels, curl_status, jq_status):
+    result = _run_bootstrap(
+        tmp_path,
+        debug_test_yaml=debug_test_yaml,
+        curl_status=curl_status,
+        jq_status=jq_status,
+    )
+
+    assert result.returncode == 0, result.stderr
+    assert "curl called" not in result.stderr
+    assert "jq called" not in result.stderr
+    assert "legacy ready-suite fallback" not in result.stderr
+    assert [
+        line.split()[2] for line in result.stdout.splitlines() if line.startswith("Skipping AMD ")
+    ] == expected_levels
+
+
+@pytest.mark.parametrize(("curl_status", "jq_status"), [(22, 0), (0, 4)], ids=["request-failure", "parse-failure"])
+@pytest.mark.parametrize("debug_test_yaml", [None, ""], ids=["unset", "empty"])
+def test_label_lookup_failure_without_override(tmp_path, debug_test_yaml, curl_status, jq_status):
+    result = _run_bootstrap(
+        tmp_path,
+        debug_test_yaml=debug_test_yaml,
+        curl_status=curl_status,
+        jq_status=jq_status,
+    )
+
+    assert result.returncode == 1
+    assert "Could not read PR labels" in result.stderr
+    assert "Skipping AMD " not in result.stdout
+
+
+@pytest.mark.parametrize(
+    ("labels", "expected_levels"),
+    [
+        (("ready",), ["l2"]),
+        (("merge-test",), ["l3"]),
+        (("ready", "merge-test"), ["l2", "l3"]),
+        ((), ["l2"]),
+    ],
+)
+def test_label_lookup_selects_tiers(tmp_path, labels, expected_levels):
+    result = _run_bootstrap(tmp_path, labels=labels)
+
+    assert result.returncode == 0, result.stderr
+    assert "curl called" in result.stderr
+    assert "jq called" in result.stderr
+    assert [
+        line.split()[2] for line in result.stdout.splitlines() if line.startswith("Skipping AMD ")
+    ] == expected_levels
+
+
+@pytest.mark.parametrize("debug_test_yaml", ["ready,ready", "nightly", ", ,", " \t"])
+def test_invalid_debug_override_fails_without_label_lookup(tmp_path, debug_test_yaml):
+    result = _run_bootstrap(tmp_path, debug_test_yaml=debug_test_yaml, curl_status=22)
+
+    assert result.returncode == 1
+    assert "DEBUG_TEST_YAML" in result.stderr
+    assert "curl called" not in result.stderr
+    assert "jq called" not in result.stderr
+    assert "Skipping AMD " not in result.stdout
+
+
+def test_main_does_not_lookup_labels(tmp_path):
+    result = _run_bootstrap(tmp_path, pull_request=False, curl_status=22, jq_status=4)
+
+    assert result.returncode == 0, result.stderr
+    assert "curl called" not in result.stderr
+    assert "jq called" not in result.stderr
+    assert "Skipping AMD l3 suite" in result.stdout
@@ -0,0 +1,61 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
+from __future__ import annotations
+
+import importlib.util
+from pathlib import Path
+
+import pytest
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+
+SELECTOR_PATH = Path(".buildkite/amd/scripts/select_test_suites.py")
+SPEC = importlib.util.spec_from_file_location("select_test_suites", SELECTOR_PATH)
+assert SPEC is not None and SPEC.loader is not None
+SELECTOR = importlib.util.module_from_spec(SPEC)
+SPEC.loader.exec_module(SELECTOR)
+
+
+@pytest.mark.parametrize(
+    ("branch", "labels", "expected"),
+    [
+        ("main", (), ("merge",)),
+        ("feature", ("ready",), ("ready",)),
+        ("feature", ("merge-test",), ("merge",)),
+        ("feature", ("merge-test", "ready"), ("ready", "merge")),
+        ("feature", ("amd-test",), ("ready",)),
+        ("feature", ("amd-test", "merge-test"), ("merge",)),
+        ("feature", ("amd-test", "ready", "merge-test"), ("ready", "merge")),
+        ("feature", ("nightly-test",), ("ready",)),
+        ("feature", ("not-ready", "merge-test-extra"), ("ready",)),
+    ],
+)
+def test_label_suite_selection(branch, labels, expected):
+    assert SELECTOR.select_amd_test_suites(branch=branch, labels=labels) == expected
+
+
+def test_debug_override_takes_precedence_and_normalizes_input():
+    assert SELECTOR.select_amd_test_suites(
+        branch="main",
+        labels=("ready",),
+        debug_test_yaml=" MERGE, ready,",
+    ) == ("merge", "ready")
+
+
+def test_empty_debug_override_uses_normal_selection():
+    assert SELECTOR.select_amd_test_suites(
+        branch="feature",
+        labels=("merge-test",),
+        debug_test_yaml="",
+    ) == ("merge",)
+
+
+@pytest.mark.parametrize("value", ["ready,ready", "nightly", ", ,", " \t"])
+def test_invalid_debug_override(value):
+    with pytest.raises(ValueError):
+        SELECTOR.select_amd_test_suites(
+            branch="feature",
+            labels=(),
+            debug_test_yaml=value,
+        )
@@ -12,6 +12,7 @@ sys.path.insert(0, str(Path(__file__).resolve().parents[2] / ".buildkite" / "com

 from skip_ci import resolve_ci_decision  # noqa: E402
 from upload_pipeline import (  # noqa: E402
+    CUDA_HF_TOKEN_EXPORT,
    _expand_mirror_hardwares,
    _get_mirror_hw_selector,
    _load_bootstrap_steps,
@@ -143,6 +144,10 @@ def test_mirror_hardwares_l4_1_expands_to_agents_and_plugins(monkeypatch: pytest
    container = step["plugins"][0]["kubernetes"]["podSpec"]["containers"][0]
    assert container["image"].endswith("$BUILDKITE_COMMIT")
    assert container["resources"]["limits"]["nvidia.com/gpu"] == 1
+    env_names = {item["name"] for item in container["env"]}
+    assert "VLLM_CI_HF_TOKEN" in env_names
+    assert "HF_TOKEN" not in env_names
+    assert step["commands"] == [CUDA_HF_TOKEN_EXPORT, "pytest -sv tests/example"]


 def test_mirror_hardwares_l4_preserves_explicit_retry() -> None:
@@ -182,6 +187,33 @@ def test_mirror_hardwares_a2b3_npu_4_expands_agents_image_and_plugins() -> None:
    assert step["plugins"][0]["kubernetes"]["podSpecPatch"]["imagePullSecrets"] == [
        {"name": "swr-secret"},
    ]
+    assert step["commands"] == ["pytest -sv tests/example"]
+
+
+def test_all_cuda_mirror_hardwares_restore_hf_token_at_runtime() -> None:
+    for name in (
+        "l4_1",
+        "l4_2",
+        "l4_3",
+        "l4_4",
+        "h100_1",
+        "h100_2",
+        "h100_3",
+        "h100_4",
+        "b200_1",
+        "b200_2",
+        "b200_3",
+        "b200_4",
+    ):
+        step = _expand_mirror_hardwares(
+            {"label": name, "mirror_hardwares": name, "commands": ["pytest -sv tests/example"]},
+        )
+        assert step is not None
+        container = step["plugins"][0]["kubernetes"]["podSpec"]["containers"][0]
+        env_names = {item["name"] for item in container["env"]}
+        assert "VLLM_CI_HF_TOKEN" in env_names
+        assert "HF_TOKEN" not in env_names
+        assert step["commands"][0] == CUDA_HF_TOKEN_EXPORT


 def _gpu_limit(step: dict) -> int:
@@ -215,6 +215,16 @@ def test_from_pipeline_config_normalizes_stage_engine_extras_without_expanding_s
    assert stage.diffusion_config.model_config["default_robot_embodiment"] == "roboarena"


+@pytest.mark.parametrize("disabled", [True, False])
+def test_frontend_log_stats_flag_is_not_an_unowned_stage_argument(disabled):
+    from vllm_omni.engine.stage_init_utils import build_engine_args_dict_from_omni_stage_config
+
+    config = _from_pipeline_key("dots_tts", cli_overrides={"disable_log_stats": disabled})
+    assert config.stage_configs
+    engine_args = build_engine_args_dict_from_omni_stage_config(config.stage_by_id(0), model="test-model")
+    assert "disable_log_stats" not in engine_args
+
+
 def test_from_pipeline_config_applies_cli_overrides_without_stage_config_runtime_bridge():
    omni_config = _from_pipeline_key(
        "qwen3_tts",
@@ -0,0 +1,59 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+"""The omni scheduler must stay on the v1 model runner, like every omni worker.
+
+vLLM 0.28 resolved this config to v1; 0.29 ends its fallback chain with v2. When
+only the workers forced v1, the scheduler took its v2 fast path -- which omits
+resumed-request token ids -- while the v1 runner still read them, so resuming a
+request raised ``KeyError: <req_id>`` in ``_update_states`` and killed the engine
+core (Multi-Replica Startup Test).
+
+The default is upstream's to change, so this pins Omni's side of the contract
+rather than the default itself.
+"""
+
+from __future__ import annotations
+
+from types import SimpleNamespace
+
+import pytest
+
+from vllm_omni.core.sched.omni_scheduler_mixin import OmniSchedulerMixin
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+
+
+class _Scheduler(OmniSchedulerMixin):
+    """Minimal carrier for the shared init hook both omni schedulers call."""
+
+    def __init__(self, *, use_v2_model_runner: bool) -> None:
+        self.use_v2_model_runner = use_v2_model_runner
+        self.vllm_config = SimpleNamespace(
+            model_config=SimpleNamespace(
+                async_chunk=False,
+                stage_id=0,
+                pooling_output_decoder=None,
+            )
+        )
+        self._init_omni_io_scheduling_state()
+
+
+@pytest.mark.parametrize("initial", [True, False])
+def test_scheduler_ends_on_v1_model_runner(initial: bool) -> None:
+    """Whatever upstream resolves, the omni scheduler runs v1."""
+    assert _Scheduler(use_v2_model_runner=initial).use_v2_model_runner is False
+
+
+def test_v2_default_does_not_reach_the_resumed_request_fast_path() -> None:
+    """The v2 fast path is what dropped resumed-request token ids.
+
+    ``OmniGenerationScheduler`` branches on ``use_v2_model_runner`` when it
+    assembles SchedulerOutput; forcing it false keeps the v1 assembly that
+    carries ``all_token_ids`` for resumed requests, which the v1 runner reads.
+    """
+    scheduler = _Scheduler(use_v2_model_runner=True)
+
+    assert not scheduler.use_v2_model_runner, (
+        "scheduler would take the v2 fast path while omni workers run v1; "
+        "resumed requests then KeyError in _update_states"
+    )
@@ -13,9 +13,6 @@ from tests.helpers.runtime import OmniServer, OmniServerParams, OnlineOmniClient

 pytestmark = [pytest.mark.slow, pytest.mark.diffusion]

-# https://github.com/vllm-project/vllm-omni/issues/3649 — invalid inputs still HTTP 200 / timeout / missing validation.
-_SKIP_ISSUE_3649 = pytest.mark.skip(reason="https://github.com/vllm-project/vllm-omni/issues/3649")
-
 _PARAMS = [
    pytest.param(
        OmniServerParams(
@@ -36,10 +33,8 @@ _PARAMS = [
@pytest.mark.parametrize(
    "overrides, err_message",
    [
-        pytest.param({"input": ""}, ("cannot be empty", "input"), id="input_empty", marks=_SKIP_ISSUE_3649),
-        pytest.param(
-            {"input": "   "}, ("cannot be empty", "input"), id="input_whitespace_only", marks=_SKIP_ISSUE_3649
-        ),
+        pytest.param({"input": ""}, ("cannot be empty", "input"), id="input_empty"),
+        pytest.param({"input": "   "}, ("cannot be empty", "input"), id="input_whitespace_only"),
        pytest.param({"input": 123}, ("input", "a valid string"), id="input_wrong_type"),
        pytest.param(
            {"response_format": "mpeg"},
@@ -54,10 +49,8 @@ _PARAMS = [
            ("sse", "stream_format", "value_error", "not a supported stream_format"),
            id="stream_format_sse_blocked",
        ),
-        pytest.param(
-            {"audio_length": -1.0}, ("audio_end_in_s", "higher than", "audio_start_in_s"), id="audio_length_negative"
-        ),
-        pytest.param({"audio_length": 0}, "audio_length", id="audio_length_zero", marks=_SKIP_ISSUE_3649),
+        pytest.param({"audio_length": -1.0}, ("audio_length", "greater_than", "0"), id="audio_length_negative"),
+        pytest.param({"audio_length": 0}, ("audio_length", "greater_than", "0"), id="audio_length_zero"),
        # ``OmniOpenAIServingAudioGenerate`` only forwards ``audio_start`` when ``audio_length`` is set.
        pytest.param(
            {"audio_start": -0.5, "audio_length": 5.0},
@@ -67,26 +60,30 @@ _PARAMS = [
        pytest.param(
            {"negative_prompt": ["noise"]}, ("negative_prompt", "a valid string"), id="negative_prompt_wrong_type"
        ),
-        pytest.param({"guidance_scale": -1.0}, "guidance_scale", id="guidance_scale_negative", marks=_SKIP_ISSUE_3649),
-        pytest.param({"guidance_scale": 0}, "guidance_scale", id="guidance_scale_zero", marks=_SKIP_ISSUE_3649),
+        pytest.param(
+            {"guidance_scale": -1.0},
+            ("guidance_scale", "greater_than_equal", "0"),
+            id="guidance_scale_negative",
+        ),
        pytest.param(
            {"num_inference_steps": 0},
-            ("num_inference_steps", "greater than or equal to 1"),
+            ("num_inference_steps", "greater_than_equal", "1"),
            id="num_inference_steps_zero",
        ),
        pytest.param(
            {"num_inference_steps": -1},
-            ("num_inference_steps", "greater than or equal to 1"),
+            ("num_inference_steps", "greater_than_equal", "1"),
            id="num_inference_steps_negative",
        ),
        pytest.param(
            {"num_inference_steps": 6000},
-            "num_inference_steps",
+            ("num_inference_steps", "less_than_equal", "1000"),
            id="num_inference_steps_above_max",
-            marks=_SKIP_ISSUE_3649,
        ),
        pytest.param(
-            {"guidance_scale": 1001.0}, "guidance_scale", id="guidance_scale_above_max", marks=_SKIP_ISSUE_3649
+            {"guidance_scale": 1001.0},
+            ("guidance_scale", "less_than_equal", "1000"),
+            id="guidance_scale_above_max",
        ),
        pytest.param(
            {"audio_length": 86401.0}, ("Requested audio length", "exceeds maximum"), id="audio_length_above_max"
@@ -13,6 +13,12 @@ from vllm_omni.diffusion.attention.backends.sdpa import SDPAImpl
 from vllm_omni.platforms import current_omni_platform


+def _clear_hub_module_cache(monkeypatch: pytest.MonkeyPatch) -> None:
+    from vllm_omni.diffusion.attention.backends import flash_attn_hub
+
+    monkeypatch.setattr(flash_attn_hub, "_hub_modules", {})
+
+
@pytest.mark.core_model
@pytest.mark.cpu
 def test_explicit_kernels_hub_selection_does_not_fallback(monkeypatch: pytest.MonkeyPatch):
@@ -54,6 +60,74 @@ def test_explicit_kernels_hub_selection_does_not_fallback(monkeypatch: pytest.Mo
        CudaOmniPlatform.get_diffusion_attn_backend_cls("FLASH_ATTN_3_HUB", head_size=64)


+@pytest.mark.core_model
+@pytest.mark.cpu
+def test_explicit_kernels_hub_variant_miss_does_not_fallback(monkeypatch: pytest.MonkeyPatch):
+    """Explicit Hub backends must fail loudly when no build variant resolves (#6971)."""
+    from vllm.platforms.interface import DeviceCapability
+
+    from vllm_omni.diffusion.envs import PACKAGES_CHECKER
+    from vllm_omni.platforms.cuda.platform import CudaOmniPlatform
+
+    monkeypatch.setattr(
+        CudaOmniPlatform,
+        "get_device_capability",
+        classmethod(lambda cls, device_id=0: DeviceCapability(9, 0)),
+    )
+    monkeypatch.setattr(PACKAGES_CHECKER, "get_packages_info", lambda: {"has_flash_attn": True})
+    _clear_hub_module_cache(monkeypatch)
+
+    kernels_module = types.ModuleType("kernels")
+
+    def _missing_variant(repo_id, version=None, **kwargs):
+        raise FileNotFoundError(f"Cannot find a build variant for this system in {repo_id}")
+
+    kernels_module.get_kernel = _missing_variant  # type: ignore[attr-defined]
+    monkeypatch.setitem(sys.modules, "kernels", kernels_module)
+
+    with pytest.raises(RuntimeError, match="no compatible build"):
+        CudaOmniPlatform.get_diffusion_attn_backend_cls("FLASH_ATTN_3_HUB", head_size=64)
+    with pytest.raises(RuntimeError, match="no compatible build"):
+        CudaOmniPlatform.get_diffusion_attn_backend_cls("FLASH_ATTN_HUB", head_size=64)
+
+
+@pytest.mark.core_model
+@pytest.mark.cpu
+def test_explicit_kernels_hub_preflight_uses_v2_when_v1_misses(monkeypatch: pytest.MonkeyPatch):
+    """Preflight must share _load_hub_module's (1, 2) policy, not pin version=1."""
+    from vllm.platforms.interface import DeviceCapability
+
+    from vllm_omni.diffusion.attention.backends.registry import DiffusionAttentionBackendEnum
+    from vllm_omni.diffusion.envs import PACKAGES_CHECKER
+    from vllm_omni.platforms.cuda.platform import CudaOmniPlatform
+
+    monkeypatch.setattr(
+        CudaOmniPlatform,
+        "get_device_capability",
+        classmethod(lambda cls, device_id=0: DeviceCapability(9, 0)),
+    )
+    monkeypatch.setattr(PACKAGES_CHECKER, "get_packages_info", lambda: {"has_flash_attn": True})
+    _clear_hub_module_cache(monkeypatch)
+
+    kernels_module = types.ModuleType("kernels")
+    versions_tried: list[int | None] = []
+
+    def _v1_miss_v2_hit(repo_id, version=None, **kwargs):
+        versions_tried.append(version)
+        if version == 1:
+            raise FileNotFoundError(f"Cannot find a build variant for this system in {repo_id}")
+        return types.SimpleNamespace(name=f"{repo_id}@v{version}")
+
+    kernels_module.get_kernel = _v1_miss_v2_hit  # type: ignore[attr-defined]
+    monkeypatch.setitem(sys.modules, "kernels", kernels_module)
+
+    fa3_path = CudaOmniPlatform.get_diffusion_attn_backend_cls("FLASH_ATTN_3_HUB", head_size=64)
+    assert fa3_path == DiffusionAttentionBackendEnum.FLASH_ATTN_3_HUB.get_path()
+    fa2_path = CudaOmniPlatform.get_diffusion_attn_backend_cls("FLASH_ATTN_HUB", head_size=64)
+    assert fa2_path == DiffusionAttentionBackendEnum.FLASH_ATTN_HUB.get_path()
+    assert versions_tried == [1, 2, 1, 2]
+
+
@pytest.mark.core_model
@pytest.mark.cpu
 def test_explicit_flash_attention_unavailable_does_not_fallback(monkeypatch: pytest.MonkeyPatch):
@@ -279,17 +279,17 @@ def _install_native_modules(monkeypatch: pytest.MonkeyPatch, events: list[tuple]
        group = SimpleNamespace(get_metadata_builder=lambda _index: SimpleNamespace(reorder_batch_threshold=None))
        return [[group], [group]], "cg-support", [4, 8]

+    # vLLM 0.29 dropped attn_groups and cache_dtype from init_kv_cache: the
+    # layout now rides on the resolved CacheConfig.kv_cache_layout instead.
    def init_cache(
        runner_kv_caches,
        forward_context,
        kv_cache_config,
-        attn_groups,
        device,
-        cache_dtype,
        kernel_block_sizes,
        vllm_config,
    ):
-        events.append(("cache", kv_cache_config, attn_groups, kernel_block_sizes))
+        events.append(("cache", kv_cache_config, kernel_block_sizes))
        runner_kv_caches.extend(["cache-0", "cache-1"])
        return {"layer-0": "tensor-0", "layer-1": "tensor-1"}

@@ -299,8 +299,10 @@ def test_minus_one_max_model_len_is_auto_fitted_by_native_cache_sizing() -> None
    )

    assert vllm_config.model_config.original_max_model_len == -1
-    # Sixteen physical pages provide ``16 * block_size`` tokens.  The native
-    # backend chooses the block geometry, so keep this assertion backend
+    # Sixteen physical pages provide ``(16 - 1) * block_size`` tokens: vLLM 0.29
+    # subtracts the null block that BlockPool permanently holds back before
+    # sizing (check_enough_kv_cache_memory), so only 15 pages are usable.  The
+    # native backend chooses the block geometry, so keep this assertion backend
    # independent instead of assuming the generic 16-token block size.
-    expected_max_model_len = min(model_limit, spec.block_size * 16)
+    expected_max_model_len = min(model_limit, spec.block_size * (16 - 1))
    assert vllm_config.model_config.max_model_len == expected_max_model_len
@@ -0,0 +1,147 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+"""Physical KV cache layout selection and propagation for the Diffusion cache."""
+
+from __future__ import annotations
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+from vllm.config import CacheConfig
+from vllm.v1.kv_cache_interface import FullAttentionSpec
+from vllm.v1.kv_cache_layout import KVCacheLayout
+
+from vllm_omni.diffusion.diffusion_kv.layout import (
+    DEFAULT_DIFFUSION_KV_CACHE_LAYOUT,
+    adopt_kv_cache_layout,
+    assert_backend_layout_supported,
+    build_kv_cache_tensor,
+    resolve_diffusion_kv_cache_layout,
+)
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu, pytest.mark.diffusion]
+
+
+def _config():
+    return SimpleNamespace(cache_config=CacheConfig())
+
+
+class _DenseBackend:
+    @classmethod
+    def indexes_kv_by_block_stride(cls) -> bool:
+        return False
+
+
+class _BlockStrideBackend:
+    @classmethod
+    def indexes_kv_by_block_stride(cls) -> bool:
+        return True
+
+
+def test_default_layout_is_not_block_outermost() -> None:
+    """The default reproduces the pre-0.29 indexes_kv_by_block_stride=False."""
+    vllm_config = _config()
+    layout = resolve_diffusion_kv_cache_layout(vllm_config)
+
+    assert layout is DEFAULT_DIFFUSION_KV_CACHE_LAYOUT
+    assert layout.is_block_outermost is False
+    assert vllm_config.cache_config.get_resolved_kv_cache_layout() is layout
+
+
+def test_block_stride_backend_selects_block_outermost_layout() -> None:
+    layout = resolve_diffusion_kv_cache_layout(_config(), indexes_kv_by_block_stride=True)
+
+    assert layout.is_block_outermost is True
+
+
+def test_resolution_is_idempotent() -> None:
+    vllm_config = _config()
+
+    assert resolve_diffusion_kv_cache_layout(vllm_config) is resolve_diffusion_kv_cache_layout(vllm_config)
+
+
+def test_contradicting_requirement_raises_instead_of_misaddressing() -> None:
+    vllm_config = _config()
+    resolve_diffusion_kv_cache_layout(vllm_config)
+
+    with pytest.raises(ValueError, match="contradicts the attention"):
+        resolve_diffusion_kv_cache_layout(vllm_config, indexes_kv_by_block_stride=True)
+
+
+def test_preset_layout_is_honoured() -> None:
+    vllm_config = _config()
+    vllm_config.cache_config.kv_cache_layout = KVCacheLayout.BLNHC.name
+
+    layout = resolve_diffusion_kv_cache_layout(vllm_config, indexes_kv_by_block_stride=True)
+
+    assert layout is KVCacheLayout.BLNHC
+
+
+def test_backend_assertion_is_quiet_until_the_layout_is_resolved() -> None:
+    """Workers collect specs before the control plane hands them the layout."""
+    vllm_config = _config()
+
+    assert_backend_layout_supported(vllm_config, _BlockStrideBackend)  # unresolved: nothing to check
+
+    resolve_diffusion_kv_cache_layout(vllm_config)
+    assert_backend_layout_supported(vllm_config, _DenseBackend)
+    assert_backend_layout_supported(vllm_config, None)
+    with pytest.raises(ValueError, match="contradicts the attention"):
+        assert_backend_layout_supported(vllm_config, _BlockStrideBackend)
+
+
+def test_worker_adopts_the_layout_carried_on_the_kv_cache_config() -> None:
+    """Diffusion workers receive only a KVCacheConfig, so the name rides on it."""
+    worker = _config()
+    shipped = SimpleNamespace(kv_cache_layout=KVCacheLayout.BLNHC.name)
+
+    assert adopt_kv_cache_layout(worker, shipped) is KVCacheLayout.BLNHC
+    assert worker.cache_config.get_resolved_kv_cache_layout() is KVCacheLayout.BLNHC
+
+
+def test_worker_without_a_carried_layout_falls_back_to_the_same_default() -> None:
+    worker = _config()
+
+    assert adopt_kv_cache_layout(worker, SimpleNamespace(kv_cache_layout=None)) is DEFAULT_DIFFUSION_KV_CACHE_LAYOUT
+
+
+def test_worker_disagreeing_with_the_control_plane_raises() -> None:
+    worker = _config()
+    resolve_diffusion_kv_cache_layout(worker)
+
+    with pytest.raises(ValueError, match="disagrees with the layout"):
+        adopt_kv_cache_layout(worker, SimpleNamespace(kv_cache_layout=KVCacheLayout.BLNHC.name))
+
+
+def test_block_stride_requirement_propagates_control_plane_to_worker() -> None:
+    """A block-stride backend must reach the worker as a block-outermost layout.
+
+    Covers the whole seam: the control plane resolves with the requirement, the
+    name rides on each rank-local KVCacheConfig, the worker adopts it, and the
+    backend contract check then passes instead of raising.
+    """
+    control_plane = _config()
+    layout = resolve_diffusion_kv_cache_layout(control_plane, indexes_kv_by_block_stride=True)
+    assert layout.is_block_outermost is True
+
+    shipped = SimpleNamespace(kv_cache_layout=layout.name)
+
+    worker = _config()
+    assert adopt_kv_cache_layout(worker, shipped) is layout
+    # The worker's own backend requirement now agrees with what it was handed.
+    assert_backend_layout_supported(worker, _BlockStrideBackend)
+    with pytest.raises(ValueError, match="contradicts the attention"):
+        assert_backend_layout_supported(worker, _DenseBackend)
+
+
+def test_tensor_strides_come_from_the_upstream_derivation() -> None:
+    spec = FullAttentionSpec(block_size=16, num_kv_heads=2, head_size=8, dtype=torch.bfloat16)
+
+    tensor = build_kv_cache_tensor(spec, 4, ["layer0", "layer1"])
+
+    assert tensor.layers == ["layer0", "layer1"]
+    assert tensor.size == spec.page_size_bytes * 2 * 4
+    # Layer-outermost: each layer owns a contiguous run of its own blocks.
+    assert tensor.layer_stride == spec.page_size_bytes * 4
+    assert tensor.block_stride == spec.page_size_bytes
@@ -1,11 +1,12 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 import pytest
 import torch
 import vllm.v1.core.single_type_kv_cache_manager as native_kv_managers
-from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec, KVCacheTensor
+from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec

+from tests.helpers.kv_layout import build_kv_cache_tensor
 from vllm_omni.diffusion.diffusion_kv.manager import DiffusionKVAdmissionError, DiffusionKVCacheManager
 from vllm_omni.diffusion.diffusion_kv.request import DiffusionKVContext, DiffusionKVRequest

@@ -24,7 +25,7 @@ def _config(num_blocks: int) -> KVCacheConfig:
    )
    return KVCacheConfig(
        num_blocks=num_blocks,
-        kv_cache_tensors=[KVCacheTensor(size=spec.page_size_bytes * num_blocks, shared_by=["layer0"])],
+        kv_cache_tensors=[build_kv_cache_tensor(spec, num_blocks, ["layer0"])],
        kv_cache_groups=[KVCacheGroupSpec(layer_names=["layer0"], kv_cache_spec=spec)],
    )

@@ -15,6 +15,7 @@ import torch
 from torch import nn
 from vllm.v1.kv_cache_interface import FullAttentionSpec

+from tests.helpers.kv_layout import layout_for_backend
 from vllm_omni.diffusion.attention.backends.flash_attn import FlashAttentionImpl
 from vllm_omni.diffusion.attention.layer import Attention
 from vllm_omni.diffusion.diffusion_kv import paged_attention_adapter as adapter_module
@@ -1239,7 +1240,8 @@ def test_layer_adapter_accepts_platform_native_backend_and_uses_rank_local_heads
    assert native_layer.num_heads == 4
    assert native_layer.num_kv_heads == 2
    assert native_layer.spec.num_kv_heads == 2
-    assert native_layer.spec.indexes_kv_by_block_stride is True
+    # 0.29: the flag no longer rides on the spec, it selects the layout.
+    assert layout_for_backend(native_layer.attn_backend).is_block_outermost is True
    assert native_layer._q_scale_float == 1.0
    assert native_layer._k_scale_float == 1.0
    assert native_layer._v_scale_float == 1.0
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 from __future__ import annotations

@@ -10,12 +10,14 @@ import torch
 from torch import nn
 from torch.nn import functional as F
 from vllm.config import AttentionConfig, CacheConfig, VllmConfig, set_current_vllm_config
-from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec, KVCacheTensor
+from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec
 from vllm.v1.worker.gpu.attn_utils import init_attn_backend, init_kv_cache
 from vllm.v1.worker.gpu.block_table import BlockTables

+from tests.helpers.kv_layout import build_kv_cache_tensor
 from vllm_omni.diffusion.attention.backends.abstract import AttentionMetadata
 from vllm_omni.diffusion.attention.backends.flash_attn import FlashAttentionImpl
+from vllm_omni.diffusion.diffusion_kv.layout import resolve_diffusion_kv_cache_layout
 from vllm_omni.diffusion.diffusion_kv.paged_attention_adapter import (
    DiffusionPagedAttentionAdapter,
    DiffusionPagedAttentionLayerAdapter,
@@ -88,12 +90,7 @@ def test_adapter_executes_native_paged_attention_on_non_contiguous_blocks() -> N
    canonical_spec = native_layer.spec
    kv_cache_config = KVCacheConfig(
        num_blocks=4,
-        kv_cache_tensors=[
-            KVCacheTensor(
-                size=canonical_spec.page_size_bytes * 4,
-                shared_by=[_LAYER_NAME],
-            )
-        ],
+        kv_cache_tensors=[build_kv_cache_tensor(canonical_spec, 4, [_LAYER_NAME])],
        kv_cache_groups=[
            KVCacheGroupSpec(
                layer_names=[_LAYER_NAME],
@@ -101,6 +98,9 @@ def test_adapter_executes_native_paged_attention_on_non_contiguous_blocks() -> N
            )
        ],
    )
+    # 0.29 reads the resolved physical layout when allocating; the diffusion
+    # control plane pins it, so this standalone harness must too.
+    resolve_diffusion_kv_cache_layout(vllm_config)
    with set_current_vllm_config(vllm_config):
        attn_groups, _, kernel_block_sizes = init_attn_backend(
            kv_cache_config,
@@ -108,13 +108,13 @@ def test_adapter_executes_native_paged_attention_on_non_contiguous_blocks() -> N
            device,
        )
        runner_kv_caches: list[torch.Tensor | list[torch.Tensor]] = []
+        # vLLM 0.29 dropped attn_groups and cache_dtype here: the layout now
+        # comes from the resolved CacheConfig.kv_cache_layout inside the call.
        init_kv_cache(
            runner_kv_caches,
            vllm_config.compilation_config.static_forward_context,
            kv_cache_config,
-            attn_groups,
            device,
-            vllm_config.cache_config.cache_dtype,
            kernel_block_sizes,
            vllm_config,
        )
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 import pytest
 import torch
@@ -10,10 +10,10 @@ from vllm.v1.kv_cache_interface import (
    FullAttentionSpec,
    KVCacheConfig,
    KVCacheGroupSpec,
-    KVCacheTensor,
 )
 from vllm.v1.request import RequestStatus

+from tests.helpers.kv_layout import build_kv_cache_tensor
 from vllm_omni.diffusion.diffusion_kv.request import DiffusionKVContext, DiffusionKVRequest

 pytestmark = [pytest.mark.core_model, pytest.mark.cpu, pytest.mark.diffusion]
@@ -119,7 +119,7 @@ def _manager(*, enable_caching: bool) -> KVCacheManager:
    group = KVCacheGroupSpec(layer_names=["layer0"], kv_cache_spec=spec)
    config = KVCacheConfig(
        num_blocks=8,
-        kv_cache_tensors=[KVCacheTensor(size=spec.page_size_bytes * 8, shared_by=["layer0"])],
+        kv_cache_tensors=[build_kv_cache_tensor(spec, 8, ["layer0"])],
        kv_cache_groups=[group],
    )
    return KVCacheManager(
@@ -8,9 +8,10 @@ from unittest.mock import Mock
 import pytest
 import torch
 from torch import nn
-from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec, KVCacheTensor
+from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec

 import vllm_omni.diffusion.worker.diffusion_worker as diffusion_worker_module
+from tests.helpers.kv_layout import build_kv_cache_tensor, layout_for_backend
 from vllm_omni.diffusion.attention.layer import Attention
 from vllm_omni.diffusion.diffusion_kv.config import DiffusionKVCacheMode
 from vllm_omni.diffusion.worker.diffusion_model_runner import DiffusionModelRunner
@@ -79,7 +80,9 @@ def test_runner_discovers_native_spec_from_loaded_attention() -> None:
    assert spec.num_kv_heads == 2
    assert spec.head_size == 8
    assert spec.dtype is torch.bfloat16
-    assert spec.indexes_kv_by_block_stride is True
+    # vLLM 0.29 moved the block-stride decision off the spec onto the resolved
+    # physical layout; the backend's declaration must still select one.
+    assert layout_for_backend(_Backend).is_block_outermost is True
    assert spec.non_causal is True


@@ -112,7 +115,7 @@ def test_runner_retains_matching_rank_local_config() -> None:
    spec = runner.get_kv_cache_spec()["image_attention"]
    config = KVCacheConfig(
        num_blocks=8,
-        kv_cache_tensors=[KVCacheTensor(size=spec.page_size_bytes * 8, shared_by=["image_attention"])],
+        kv_cache_tensors=[build_kv_cache_tensor(spec, 8, ["image_attention"])],
        kv_cache_groups=[KVCacheGroupSpec(layer_names=["image_attention"], kv_cache_spec=spec)],
    )

@@ -126,7 +129,7 @@ def test_runner_rejects_rank_local_config_for_different_layers() -> None:
    spec = runner.get_kv_cache_spec()["image_attention"]
    config = KVCacheConfig(
        num_blocks=8,
-        kv_cache_tensors=[KVCacheTensor(size=spec.page_size_bytes * 8, shared_by=["other_attention"])],
+        kv_cache_tensors=[build_kv_cache_tensor(spec, 8, ["other_attention"])],
        kv_cache_groups=[KVCacheGroupSpec(layer_names=["other_attention"], kv_cache_spec=spec)],
    )

@@ -5,7 +5,7 @@ import pytest
 import torch
 import torch.nn.functional as F

-from tests.helpers.mark import hardware_test
+from tests.helpers.mark import hardware_marks
 from vllm_omni.diffusion.layers.activation import SiluAndMul

 pytestmark = [pytest.mark.core_model, pytest.mark.diffusion]
@@ -19,9 +19,17 @@ def test_silu_and_mul_native_matches_packed_reference() -> None:
    torch.testing.assert_close(SiluAndMul().forward_native(packed), F.silu(gate) * up)


-@hardware_test(res={"npu": "A3"}, num_cards=1)
-def test_silu_and_mul_npu_matches_packed_reference() -> None:
-    packed = torch.randn(257, 256, device="npu", dtype=torch.bfloat16)
+@pytest.mark.parametrize(
+    "dtype",
+    [
+        pytest.param(torch.bfloat16, marks=hardware_marks(res={"npu": "A2"}, num_cards=1), id="A2-bf16"),
+        pytest.param(torch.bfloat16, marks=hardware_marks(res={"npu": "A3"}, num_cards=1), id="A3-bf16"),
+        pytest.param(torch.bfloat16, marks=hardware_marks(res={"npu": "A5"}, num_cards=1), id="A5-bf16"),
+        pytest.param(torch.float16, marks=hardware_marks(res={"npu": "310P"}, num_cards=1), id="310P-fp16"),
+    ],
+)
+def test_silu_and_mul_npu_matches_packed_reference(dtype: torch.dtype) -> None:
+    packed = torch.randn(257, 256, device="npu", dtype=dtype)
    gate, up = packed.chunk(2, dim=-1)

    torch.testing.assert_close(
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Unit tests for shared AdaLayerNorm layers used by FLUX and other models."""

 import os
@@ -46,7 +46,7 @@ def _force_default_gemm(monkeypatch):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )


@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """
 Shape broadcasting regression tests for RotaryEmbedding.forward_native.

@@ -16,7 +16,11 @@ import types
 import pytest
 import torch

-from vllm_omni.diffusion.layers.rope import RotaryEmbedding, apply_rotary_emb_mindiesd
+from vllm_omni.diffusion.layers.rope import (
+    RotaryEmbedding,
+    RotaryEmbeddingWan,
+    apply_rotary_emb_mindiesd,
+)

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

@@ -239,6 +243,25 @@ class TestRotaryEmbeddingNativeShapeRegression:
        assert not torch.isnan(output).any()


+def test_cuda_paths_fall_back_without_vendored_rotary_kernel() -> None:
+    """Source/CPU installs need not contain the wheel-vendored CUDA kernel."""
+    x = torch.randn(1, 5, 2, 8)
+
+    rope = RotaryEmbedding(is_neox_style=True)
+    rope.apply_rotary_emb_vllm_flash_attn = None
+    output = rope.forward_cuda(x, torch.randn(5, 4), torch.randn(5, 4))
+    assert output.shape == x.shape
+
+    wan_rope = RotaryEmbeddingWan(is_neox_style=False, half_head_dim=True)
+    wan_rope.apply_rotary_emb_vllm_flash_attn = None
+    output = wan_rope.forward_cuda(
+        x,
+        torch.randn(1, 5, 1, 4),
+        torch.randn(1, 5, 1, 4),
+    )
+    assert output.shape == x.shape
+
+
 def test_mindie_full_dim_rope_keeps_h3_cos_sin_layout(monkeypatch: pytest.MonkeyPatch) -> None:
    """H3 passes complete rotary dims to MindIE instead of the CUDA half-dim layout."""
    captured: dict[str, object] = {}
@@ -59,7 +59,7 @@ def _force_default_gemm(monkeypatch):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )


@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 """Equivalence test for DreamZero self-attn QKV fusion.

@@ -52,7 +52,7 @@ def _force_default_gemm(monkeypatch):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )


@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Unit tests for GLM-Image quantization support (W4A16/AutoRound).

 These tests verify that the GLM-Image DiT transformer correctly accepts and uses
@@ -69,7 +69,7 @@ def _force_default_gemm(monkeypatch):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )


@@ -186,7 +186,7 @@ def test_ltx25_missing_gemma4_recommends_supported_transformers_range(monkeypatc
        ltx2_components.initialize_pipeline_components(pipe, od_config)

    assert str(exc_info.value) == (
-        "LTX-2.5 requires Gemma4UnifiedForConditionalGeneration; install transformers>=5.10.1,<5.15."
+        "LTX-2.5 requires Gemma4UnifiedForConditionalGeneration; install transformers>=5.13.0,<5.15."
    )


@@ -387,7 +387,7 @@ def test_ltx25_natten_failure_has_actionable_omni_remedy(monkeypatch):
        MissingNattenProcessor,
    )

-    with pytest.raises(RuntimeError, match=r"kernels==0\.15\.2.*supported GPU.*allow Hub access"):
+    with pytest.raises(RuntimeError, match=r"kernels==0\.16\.1.*supported GPU.*allow Hub access"):
        ltx2_components._create_ltx25_natten_processor()


@@ -0,0 +1,89 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
+"""Native FNA against an independent, test-only explicit neighborhood mask."""
+
+import pytest
+import torch
+from torch.nn.attention import SDPBackend, sdpa_kernel
+
+from vllm_omni.diffusion.models.ltx2.ops import resolve_ltx2_vae_operators
+from vllm_omni.platforms import current_omni_platform
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cuda, pytest.mark.diffusion]
+
+
+def _selected_fna():
+    if not current_omni_platform.is_available():
+        pytest.skip("No accelerator is available")
+    operators = resolve_ltx2_vae_operators(current_omni_platform.get_torch_device())
+    if operators is None or operators.fna is None:
+        pytest.skip("LTX DiffVAE FNA is unavailable for this device or installation")
+    return operators.fna
+
+
+def _load_fna_ops():
+    return pytest.importorskip("vllm_omni.diffusion.models.ltx2.ops.fna", exc_type=ImportError)
+
+
+def _pack(x, tile):
+    b, t, h, w, heads, dim = x.shape
+    tt, th, tw = tile
+    nt, nh, nw = (t + tt - 1) // tt, (h + th - 1) // th, (w + tw - 1) // tw
+    padded = x.new_zeros((b, nt * tt, nh * th, nw * tw, heads, dim))
+    padded[:, :t, :h, :w] = x
+    # Outer and inner spatial tiles both flatten in W,H,T order.
+    return (
+        padded.view(b, nt, tt, nh, th, nw, tw, heads, dim)
+        .permute(0, 5, 3, 1, 6, 4, 2, 7, 8)
+        .reshape(b, nt * nh * nw * tt * th * tw, heads, dim)
+    )
+
+
+def _unpack(x, shape):
+    t, h, w = shape
+    b, _, heads, dim = x.shape
+    nt, nh, nw = ((n + 3) // 4 for n in shape)
+    return (
+        x.view(b, nw, nh, nt, 4, 4, 4, heads, dim)
+        .permute(0, 3, 6, 2, 5, 1, 4, 7, 8)
+        .reshape(b, nt * 4, nh * 4, nw * 4, heads, dim)[:, :t, :h, :w]
+    )
+
+
+@pytest.mark.parametrize("batch,shape", [(1, (11, 11, 11)), (2, (13, 12, 17))])
+@torch.inference_mode()
+def test_fna_matches_explicit_inward_shift_mask(batch, shape):
+    fna = _selected_fna()
+    generator = torch.Generator(device="cuda").manual_seed(42)
+    q, k, v = [
+        torch.randn(batch, *shape, 4, 64, device="cuda", dtype=torch.bfloat16, generator=generator) for _ in range(3)
+    ]
+    q = q * 0.125
+    coordinates = torch.stack(
+        torch.meshgrid(*(torch.arange(n, device="cuda") for n in shape), indexing="ij"), dim=-1
+    ).reshape(-1, 3)
+    start = (coordinates - 5).clamp_min(0).minimum(torch.tensor(shape, device="cuda") - 11)
+    mask = ((coordinates[None] >= start[:, None]) & (coordinates[None] < start[:, None] + 11)).all(-1)
+    with sdpa_kernel(SDPBackend.MATH):
+        expected = (
+            torch.nn.functional.scaled_dot_product_attention(
+                *(x.flatten(1, 3).transpose(1, 2).float() for x in (q, k, v)), attn_mask=mask, scale=1.0
+            )
+            .transpose(1, 2)
+            .reshape(q.shape)
+        )
+    actual = _unpack(
+        fna(_pack(q, (4, 4, 4)), _pack(k, (4, 4, 8)), _pack(v, (4, 4, 8)), shape=shape, window=(11, 11, 11)),
+        shape,
+    ).float()
+    assert actual.isfinite().all()
+    assert float((actual - expected).norm() / expected.norm()) < 0.005
+    torch.testing.assert_close(actual, expected, atol=0.008, rtol=0.025)
+
+
+def test_fna_rejects_cpu_without_fallback():
+    fna_ops = _load_fna_ops()
+    x = torch.zeros(1, 64, 4, 64, dtype=torch.bfloat16)
+    with pytest.raises(ValueError, match="CUDA BF16"):
+        fna_ops.fna3d_tilelang(x, x, x, shape=(11, 11, 11), window=(11, 11, 11))
@@ -0,0 +1,409 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
+from __future__ import annotations
+
+import subprocess
+import sys
+
+import pytest
+import torch
+
+from vllm_omni.diffusion.models.ltx2.ops import qk_rms_norm as qk_rms_norm_ops
+from vllm_omni.diffusion.models.ltx2.ops import residual_adaln as residual_adaln_ops
+from vllm_omni.diffusion.models.ltx2.ops import resolve_ltx2_vae_operators
+from vllm_omni.diffusion.models.ltx2.ops import swiglu as swiglu_ops
+from vllm_omni.platforms import current_omni_platform
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cuda, pytest.mark.diffusion]
+
+_DIM_SPLIT = (16, 24, 24)
+
+
+def _selected_operators():
+    if not current_omni_platform.is_available():
+        pytest.skip("No accelerator is available")
+    operators = resolve_ltx2_vae_operators(current_omni_platform.get_torch_device())
+    if operators is None:
+        pytest.skip("No LTX DiffVAE operators are available for this device")
+    return operators
+
+
+def _selected_fna():
+    operators = _selected_operators()
+    if operators.fna is None:
+        pytest.skip("LTX DiffVAE FNA is unavailable for this device or installation")
+    return operators.fna
+
+
+def _load_fna_ops():
+    return pytest.importorskip("vllm_omni.diffusion.models.ltx2.ops.fna", exc_type=ImportError)
+
+
+def test_diffvae_model_and_tests_do_not_import_tilelang_backend() -> None:
+    # Use a fresh interpreter: reloading the model in this process would leave
+    # already-imported distributed subclasses pointing at the old base class.
+    result = subprocess.run(
+        [
+            sys.executable,
+            "-c",
+            "import sys; "
+            "sys.modules['vllm_omni.diffusion.models.ltx2.ops.fna'] = None; "
+            "sys.modules['vllm.tilelang_utils'] = None; "
+            "from vllm_omni.diffusion.models.ltx2.vae.decoder import LTX2VideoDiffusionDecoderModel; "
+            "import tests.diffusion.models.ltx2.test_ltx2_diffvae_attention; "
+            "import tests.diffusion.models.ltx2.test_ltx2_diffvae_ops",
+        ],
+        capture_output=True,
+        text=True,
+        timeout=60,
+    )
+    assert result.returncode == 0, result.stdout + result.stderr
+
+
+def test_diffvae_fna_tile_map_covers_exact_neighborhood() -> None:
+    fna_ops = _load_fna_ops()
+    problem = fna_ops.Problem(13, 14, 17, (11, 11, 11))
+    indices, counts = fna_ops.build_kv_tile_map(problem)
+    q_grid_t, q_grid_h, _ = problem.q_grid
+    kv_grid_t, kv_grid_h, _ = problem.kv_grid
+
+    for q_block in range(indices.shape[0]):
+        selected = set(indices[q_block, : counts[q_block]].tolist())
+        q_outer_t = q_block % q_grid_t
+        q_outer_h = (q_block // q_grid_t) % q_grid_h
+        q_outer_w = q_block // (q_grid_t * q_grid_h)
+        for q_inner in range(fna_ops.Q_TOKENS):
+            query = (
+                q_outer_t * 4 + q_inner % 4,
+                q_outer_h * 4 + (q_inner // 4) % 4,
+                q_outer_w * 4 + q_inner // 16,
+            )
+            if any(index >= length for index, length in zip(query, problem.shape, strict=True)):
+                continue
+            starts = tuple(
+                fna_ops._window_start(index, window, length)
+                for index, window, length in zip(query, problem.window, problem.shape, strict=True)
+            )
+            for key_t in range(starts[0], starts[0] + problem.window[0]):
+                for key_h in range(starts[1], starts[1] + problem.window[1]):
+                    for key_w in range(starts[2], starts[2] + problem.window[2]):
+                        key_block = ((key_w // 8) * kv_grid_h + key_h // 4) * kv_grid_t + key_t // 4
+                        assert key_block in selected
+
+
+def test_diffvae_fna_metadata_uses_static_kernel_shapes() -> None:
+    fna_ops = _load_fna_ops()
+    for shape in ((11, 11, 11), (13, 14, 17), (66, 96, 128), (79, 192, 192)):
+        problem = fna_ops.Problem(*shape, (11, 11, 11))
+        _, t_patterns, h_patterns, w_patterns = fna_ops.build_kv_tile_metadata(problem)
+        assert {table.numel() for table in (t_patterns, h_patterns, w_patterns)} == {fna_ops.PATTERN_TABLE_SIZE}
+        assert {limit for _, _, limit in fna_ops.build_kv_tile_buckets(problem)} <= set(fna_ops.KV_BUCKET_LIMITS)
+        for ids, metadata, limit in fna_ops.build_kv_tile_buckets(problem):
+            assert metadata.shape == (ids.numel(), limit)
+
+
+def test_diffvae_fna_rejects_oversized_packed_tile_grid() -> None:
+    fna_ops = _load_fna_ops()
+    with pytest.raises(ValueError, match="15-bit metadata"):
+        fna_ops.Problem(128, 256, 256, (11, 11, 11))
+
+
+def test_diffvae_fna_propagates_launch_failures(
+    monkeypatch: pytest.MonkeyPatch,
+) -> None:
+    fna_ops = _load_fna_ops()
+    fna = _selected_fna()
+    problem = fna_ops.Problem(11, 11, 11, (11, 11, 11))
+    query = torch.randn(1, problem.q_length, fna_ops.HEADS, 64, device="cuda", dtype=torch.bfloat16)
+    key = torch.randn(1, problem.kv_length, fna_ops.HEADS, 64, device="cuda", dtype=torch.bfloat16)
+    value = torch.randn_like(key)
+    calls = 0
+
+    def fail_launch(*args, **kwargs):
+        nonlocal calls
+        calls += 1
+        raise RuntimeError("injected launch failure")
+
+    monkeypatch.setattr(fna_ops, "_fna3d_tilelang", fail_launch)
+    with torch.inference_mode():
+        for _ in range(2):
+            with pytest.raises(RuntimeError, match="injected launch failure"):
+                fna(
+                    query,
+                    key,
+                    value,
+                    shape=problem.shape,
+                    window=problem.window,
+                )
+    assert calls == 2
+
+
+@pytest.fixture
+def isolated_runtime_state():
+    targets = (
+        (qk_rms_norm_ops, "_FAILED_ROPE_KEYS"),
+        (qk_rms_norm_ops, "_VERIFIED_ROPE_KEYS"),
+        (swiglu_ops, "_FAILED_KEYS"),
+        (swiglu_ops, "_VERIFIED_KEYS"),
+        (residual_adaln_ops, "_FAILED_ADALN_KEYS"),
+        (residual_adaln_ops, "_VERIFIED_ADALN_KEYS"),
+        (residual_adaln_ops, "_FAILED_ADD_DEVICES"),
+        (residual_adaln_ops, "_VERIFIED_ADD_DEVICES"),
+    )
+    saved = [set(getattr(module, name)) for module, name in targets]
+    for module, name in targets:
+        getattr(module, name).clear()
+    try:
+        yield
+    finally:
+        for (module, name), state in zip(targets, saved, strict=True):
+            target = getattr(module, name)
+            target.clear()
+            target.update(state)
+
+
+@pytest.mark.parametrize(
+    ("dim", "rows"),
+    [
+        (256, 257),
+        (256, 16385),
+        (256, 131075),
+        (512, 65539),
+        (1024, 32771),
+        (2048, 16387),
+    ],
+)
+def test_diffvae_swiglu_tiled_is_bit_exact(dim: int, rows: int, isolated_runtime_state) -> None:
+    from diffusers.models.autoencoders.ltx2_diffusion_decoder import LTX2VideoVaeSwiGLU
+
+    operators = _selected_operators()
+    generator = torch.Generator(device="cuda").manual_seed(20260829)
+    hidden_states = torch.randn((1, 1, 1, rows, dim), device="cuda", dtype=torch.bfloat16, generator=generator)
+    gate_weight = torch.randn((4 * dim, dim), device="cuda", dtype=torch.bfloat16, generator=generator)
+    up_weight = torch.randn_like(gate_weight)
+    down_weight = torch.randn((dim, 4 * dim), device="cuda", dtype=torch.bfloat16, generator=generator)
+    with torch.device("meta"):
+        reference = LTX2VideoVaeSwiGLU(dim, 4 * dim)
+    reference.w_gate.weight = torch.nn.Parameter(gate_weight, requires_grad=False)
+    reference.w_up.weight = torch.nn.Parameter(up_weight, requires_grad=False)
+    reference.w_down.weight = torch.nn.Parameter(down_weight, requires_grad=False)
+
+    with torch.inference_mode():
+        expected = reference(hidden_states)
+        actual = operators.swiglu(hidden_states, gate_weight, up_weight, down_weight)
+
+    assert actual is not None
+    assert torch.equal(actual, expected)
+
+
+@pytest.mark.parametrize(
+    ("dim", "expected_rows"),
+    [(256, 131072), (512, 65536), (1024, 32768), (2048, 16384)],
+)
+def test_diffvae_swiglu_tile_rows_keep_a_constant_hidden_budget(dim: int, expected_rows: int) -> None:
+    assert swiglu_ops._tile_rows(1_000_000, 4 * dim) == expected_rows
+
+
+def test_diffvae_swiglu_tiled_permanently_falls_back_after_failure(
+    monkeypatch: pytest.MonkeyPatch,
+    isolated_runtime_state,
+) -> None:
+    operators = _selected_operators()
+    hidden_states = torch.randn(1, 1, 1, 17, 256, device="cuda", dtype=torch.bfloat16)
+    gate_weight = torch.randn(1024, 256, device="cuda", dtype=torch.bfloat16)
+    up_weight = torch.randn_like(gate_weight)
+    down_weight = torch.randn(256, 1024, device="cuda", dtype=torch.bfloat16)
+    calls = 0
+
+    def fail_launch(*args, **kwargs):
+        nonlocal calls
+        calls += 1
+        raise RuntimeError("injected launch failure")
+
+    monkeypatch.setattr(swiglu_ops, "_launch", fail_launch)
+    with torch.inference_mode():
+        assert operators.swiglu(hidden_states, gate_weight, up_weight, down_weight) is None
+        assert operators.swiglu(hidden_states, gate_weight, up_weight, down_weight) is None
+
+    assert calls == 1
+
+
+@pytest.mark.parametrize("residual_count", [1, 2])
+def test_diffvae_residual_rms_norm_modulate_is_bit_exact(
+    residual_count: int,
+    isolated_runtime_state,
+) -> None:
+    operators = _selected_operators()
+    generator = torch.Generator(device="cuda").manual_seed(20260829)
+    shape = (2, 3, 7, 7, 256)
+    x = torch.randn(shape, device="cuda", dtype=torch.bfloat16, generator=generator)
+    residual_a = torch.randn_like(x)
+    residual_b = torch.randn_like(x) if residual_count == 2 else None
+    norm_weight = torch.randn((256,), device="cuda", dtype=torch.bfloat16, generator=generator)
+    scale = torch.randn((2, 1, 1, 1, 256), device="cuda", dtype=torch.bfloat16, generator=generator)
+    shift = torch.randn_like(scale)
+    hidden_states = x + residual_a
+    if residual_b is not None:
+        hidden_states = hidden_states + residual_b
+    expected = torch.nn.functional.rms_norm(hidden_states, (256,), norm_weight, eps=1e-6) * (1 + scale) + shift
+
+    with torch.inference_mode():
+        actual = operators.residual_norm(
+            x,
+            residual_a,
+            residual_b,
+            norm_weight,
+            scale,
+            shift,
+            1e-6,
+        )
+
+    assert actual is not None
+    assert torch.equal(actual, expected)
+
+
+def test_diffvae_residual_add3_is_bit_exact(isolated_runtime_state) -> None:
+    operators = _selected_operators()
+    generator = torch.Generator(device="cuda").manual_seed(20260829)
+    tensors = [
+        torch.randn((2, 3, 7, 7, 256), device="cuda", dtype=torch.bfloat16, generator=generator) for _ in range(4)
+    ]
+    expected = tensors[0] + tensors[1]
+    expected = expected + tensors[2]
+    expected = expected + tensors[3]
+
+    with torch.inference_mode():
+        actual = operators.residual_add(*tensors)
+
+    assert actual is not None
+    assert torch.equal(actual, expected)
+
+
+def test_diffvae_residual_adaln_permanently_falls_back_after_failure(
+    monkeypatch: pytest.MonkeyPatch,
+    isolated_runtime_state,
+) -> None:
+    operators = _selected_operators()
+    shape = (1, 3, 7, 7, 256)
+    x = torch.randn(shape, device="cuda", dtype=torch.bfloat16)
+    residual = torch.randn_like(x)
+    weight = torch.randn(256, device="cuda", dtype=torch.bfloat16)
+    scale = torch.randn(1, 1, 1, 1, 256, device="cuda", dtype=torch.bfloat16)
+    shift = torch.randn_like(scale)
+    calls = 0
+
+    def fail_launch(*args, **kwargs):
+        nonlocal calls
+        calls += 1
+        raise RuntimeError("injected launch failure")
+
+    monkeypatch.setattr(residual_adaln_ops, "_launch_adaln", fail_launch)
+    with torch.inference_mode():
+        for _ in range(2):
+            assert operators.residual_norm(x, residual, None, weight, scale, shift, 1e-6) is None
+
+    assert calls == 1
+
+
+def _tables(
+    frames: int,
+    height: int,
+    width: int,
+    dim_split: tuple[int, int, int] = _DIM_SPLIT,
+    base: float = 10000.0,
+) -> tuple[tuple[torch.Tensor, torch.Tensor], ...]:
+    tables = []
+    for length, dim in zip((frames, height, width), dim_split, strict=True):
+        exponents = torch.arange(0, dim, 2, dtype=torch.float64, device="cuda") / dim
+        inv_freqs = (1.0 / base**exponents).to(torch.float32)
+        positions = torch.arange(length, dtype=torch.float32, device="cuda")
+        angles = positions[:, None] * inv_freqs[None, :]
+        tables.append((angles.cos(), angles.sin()))
+    return tuple(tables)
+
+
+def _apply_rope_reference(
+    hidden_states: torch.Tensor,
+    tables: tuple[tuple[torch.Tensor, torch.Tensor], ...],
+    dim_split: tuple[int, int, int] = _DIM_SPLIT,
+) -> torch.Tensor:
+    outputs = []
+    offset = 0
+    for axis, (dim, (cos, sin)) in enumerate(zip(dim_split, tables, strict=True), 1):
+        chunk = hidden_states[..., offset : offset + dim]
+        pairs = chunk.reshape(*chunk.shape[:-1], dim // 2, 2)
+        even = pairs[..., 0].float()
+        odd = pairs[..., 1].float()
+        shape = [1, 1, 1, 1, 1, dim // 2]
+        shape[axis] = cos.shape[0]
+        cos = cos.reshape(shape)
+        sin = sin.reshape(shape)
+        rotated = torch.stack([even * cos - odd * sin, even * sin + odd * cos], dim=-1)
+        outputs.append(rotated.reshape(chunk.shape).to(hidden_states.dtype))
+        offset += dim
+    return torch.cat(outputs, dim=-1)
+
+
+@pytest.mark.parametrize(
+    ("frames", "height", "width", "heads"),
+    [(3, 7, 7, 1), (31, 136, 192, 4)],
+)
+def test_diffvae_qk_rms_norm_scale_rope_3d_is_bit_exact(
+    frames: int,
+    height: int,
+    width: int,
+    heads: int,
+    isolated_runtime_state,
+) -> None:
+    operators = _selected_operators()
+    generator = torch.Generator(device="cuda").manual_seed(20260829)
+    shape = (1, frames, height, width, heads, 64)
+    query = torch.randn(shape, device="cuda", dtype=torch.bfloat16, generator=generator)
+    key = torch.randn_like(query)
+    query_weight = torch.randn((64,), device="cuda", dtype=torch.bfloat16, generator=generator)
+    key_weight = torch.randn((64,), device="cuda", dtype=torch.bfloat16, generator=generator)
+    tables = _tables(frames, height, width)
+    expected = (
+        _apply_rope_reference(torch.nn.functional.rms_norm(query, (64,), query_weight, 1e-6) * 0.125, tables),
+        _apply_rope_reference(torch.nn.functional.rms_norm(key, (64,), key_weight, 1e-6), tables),
+    )
+
+    with torch.inference_mode():
+        actual = operators.qk_norm_rope(
+            query,
+            key,
+            query_weight,
+            key_weight,
+            1e-6,
+            0.125,
+            _DIM_SPLIT,
+            10000.0,
+        )
+
+    assert actual is not None
+    assert torch.equal(actual[0], expected[0])
+    assert torch.equal(actual[1], expected[1])
+
+
+def test_diffvae_qk_rms_norm_scale_rope_3d_permanently_falls_back_after_failure(
+    monkeypatch: pytest.MonkeyPatch,
+    isolated_runtime_state,
+) -> None:
+    operators = _selected_operators()
+    query = torch.randn(1, 3, 7, 7, 1, 64, device="cuda", dtype=torch.bfloat16)
+    key = torch.randn_like(query)
+    weight = torch.ones(64, device="cuda", dtype=torch.bfloat16)
+    calls = 0
+
+    def fail_launch(*args, **kwargs):
+        nonlocal calls
+        calls += 1
+        raise RuntimeError("injected launch failure")
+
+    monkeypatch.setattr(qk_rms_norm_ops, "_launch_combined", fail_launch)
+    with torch.inference_mode():
+        for _ in range(2):
+            assert operators.qk_norm_rope(query, key, weight, weight, 1e-6, 0.125, _DIM_SPLIT, 10000.0) is None
+
+    assert calls == 1
@@ -0,0 +1,125 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
+from __future__ import annotations
+
+import sys
+from dataclasses import dataclass
+
+import pytest
+import torch
+from vllm.platforms.interface import DeviceCapability
+
+from vllm_omni.diffusion.models.ltx2 import ops as ltx2_platform
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+
+
+@dataclass(frozen=True)
+class _TensorProbe:
+    is_cuda: bool
+    device: torch.device
+
+
+@pytest.fixture(autouse=True)
+def clear_platform_cache():
+    ltx2_platform._cuda_compute_capability.cache_clear()
+    try:
+        yield
+    finally:
+        ltx2_platform._cuda_compute_capability.cache_clear()
+
+
+@pytest.mark.parametrize(
+    ("capability", "triton", "tilelang", "fusion_expected", "fna_expected"),
+    [
+        (80, True, True, False, False),
+        (90, True, True, True, True),
+        (90, True, False, True, False),
+        (90, False, True, False, False),
+        (100, True, True, True, False),
+        (100, True, False, True, False),
+        (103, True, True, True, False),
+        (103, True, False, True, False),
+        (110, True, True, False, False),
+        (120, True, True, False, False),
+        (121, True, True, False, False),
+    ],
+)
+def test_production_resolver_and_tensor_eligibility_agree(
+    monkeypatch, capability, triton, tilelang, fusion_expected, fna_expected
+):
+    monkeypatch.setattr(ltx2_platform, "HAS_TRITON", triton)
+    monkeypatch.setattr(ltx2_platform, "has_tilelang", lambda: tilelang)
+    monkeypatch.setattr(ltx2_platform.current_omni_platform, "is_cuda", lambda: True)
+    monkeypatch.setattr(ltx2_platform.current_omni_platform, "is_available", lambda: True)
+    monkeypatch.setattr(
+        ltx2_platform.current_omni_platform,
+        "get_device_capability",
+        lambda device_id=0: DeviceCapability(major=capability // 10, minor=capability % 10),
+    )
+    # Resolution must be lazy even when FNA is selected.
+    monkeypatch.setitem(sys.modules, "vllm_omni.diffusion.models.ltx2.ops.fna", None)
+    monkeypatch.setattr(torch.compiler, "is_compiling", lambda: False)
+    device = torch.device("cuda:0")
+    operators = ltx2_platform.resolve_ltx2_vae_operators(device)
+    assert (operators is not None) is fusion_expected
+    assert (operators is not None and operators.fna is not None) is fna_expected
+    if operators is not None:
+        from vllm_omni.diffusion.models.ltx2.ops import qk_rms_norm, residual_adaln, swiglu
+
+        assert operators.qk_norm_rope is qk_rms_norm.try_qk_rms_norm_scale_rope_3d_exact
+        assert operators.residual_norm is residual_adaln.try_residual_rms_norm_modulate_exact
+        assert operators.residual_add is residual_adaln.try_residual_add3_exact
+        assert operators.swiglu is swiglu.try_swiglu_tiled_exact
+
+    tensor = _TensorProbe(is_cuda=True, device=device)
+    with torch.inference_mode():
+        assert ltx2_platform.is_ltx2_fusion_eligible(tensor) is fusion_expected
+        assert ltx2_platform.is_ltx2_fna_eligible(tensor) is fna_expected
+
+
+@pytest.mark.parametrize(
+    ("device", "is_cuda", "available"),
+    [("cpu", True, True), ("cuda:0", False, True), ("cuda:0", True, False)],
+)
+def test_resolver_rejects_unsupported_platform_before_capability_lookup(monkeypatch, device, is_cuda, available):
+    monkeypatch.setattr(ltx2_platform.current_omni_platform, "is_cuda", lambda: is_cuda)
+    monkeypatch.setattr(ltx2_platform.current_omni_platform, "is_available", lambda: available)
+
+    def unexpected_lookup(_index):
+        raise AssertionError("Unsupported platforms must not query CUDA capability")
+
+    monkeypatch.setattr(ltx2_platform, "_cuda_compute_capability", unexpected_lookup)
+    assert ltx2_platform.resolve_ltx2_vae_operators(torch.device(device)) is None
+
+
+@pytest.mark.parametrize(("compiling", "grad_enabled"), [(True, False), (False, True)])
+def test_runtime_guards_bypass_operator_resolution(monkeypatch, compiling, grad_enabled):
+    monkeypatch.setattr(torch.compiler, "is_compiling", lambda: compiling)
+    monkeypatch.setattr(torch, "is_grad_enabled", lambda: grad_enabled)
+
+    def unexpected_resolve(_device):
+        raise AssertionError("Ineligible runtime must not select eager operators")
+
+    monkeypatch.setattr(ltx2_platform, "resolve_ltx2_vae_operators", unexpected_resolve)
+    tensor = _TensorProbe(is_cuda=True, device=torch.device("cuda:0"))
+    assert not ltx2_platform.is_ltx2_fusion_eligible(tensor)
+    assert not ltx2_platform.is_ltx2_fna_eligible(tensor)
+
+
+def test_device_capability_is_cached(monkeypatch):
+    calls = 0
+    monkeypatch.setattr(ltx2_platform, "HAS_TRITON", True)
+
+    def get_device_capability(device_id=0):
+        nonlocal calls
+        calls += 1
+        return DeviceCapability(major=9, minor=0)
+
+    monkeypatch.setattr(ltx2_platform.current_omni_platform, "is_cuda", lambda: True)
+    monkeypatch.setattr(ltx2_platform.current_omni_platform, "is_available", lambda: True)
+    monkeypatch.setattr(ltx2_platform.current_omni_platform, "get_device_capability", get_device_capability)
+    assert ltx2_platform._cuda_compute_capability(0) == 90
+    assert ltx2_platform._cuda_compute_capability(0) == 90
+    assert calls == 1
@@ -3,6 +3,8 @@

 """Unit tests for LTX video VAE tiling and distributed decode behavior."""

+from collections.abc import Callable
+from dataclasses import dataclass
 from types import SimpleNamespace
 from typing import Any

@@ -12,6 +14,45 @@ import torch
 pytestmark = [pytest.mark.core_model, pytest.mark.cpu]


+@pytest.mark.parametrize("initial_deterministic", [False, True])
+@pytest.mark.parametrize("fail", [False, True])
+def test_ltx_vocoder_deterministic_context(monkeypatch, initial_deterministic, fail):
+    from vllm_omni.diffusion.models.ltx2.ltx2_runtime import _deterministic_ltx_vocoder
+
+    monkeypatch.setattr(torch.backends.cudnn, "deterministic", initial_deterministic)
+    if torch.backends.cudnn.is_available():
+        monkeypatch.setattr(torch.backends.cudnn, "benchmark_limit", 17)
+    settings = {
+        name: getattr(torch.backends.cudnn, name) for name in ("enabled", "benchmark", "benchmark_limit", "allow_tf32")
+    }
+    try:
+        with _deterministic_ltx_vocoder():
+            assert torch.backends.cudnn.deterministic
+            assert {name: getattr(torch.backends.cudnn, name) for name in settings} == settings
+            if fail:
+                raise RuntimeError("injected failure")
+    except RuntimeError as exc:
+        assert fail and str(exc) == "injected failure"
+    assert torch.backends.cudnn.deterministic == initial_deterministic
+    assert {name: getattr(torch.backends.cudnn, name) for name in settings} == settings
+
+
+@dataclass(frozen=True)
+class _OperatorSetProbe:
+    fna: Callable[..., torch.Tensor] | None
+
+
+@dataclass
+class _AttentionProbe:
+    kernel_size: tuple[int, int, int]
+    head_dim: int
+    to_out: tuple[torch.nn.Module, ...]
+    projected: torch.Tensor
+
+    def project_qkv(self, _hidden_states: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+        return self.projected, self.projected, self.projected
+
+
 def test_ltx_base_vocoder_keeps_native_dtype(monkeypatch):
    import vllm_omni.diffusion.models.ltx2.ltx2_runtime as ltx_runtime

@@ -38,9 +79,164 @@ def test_ltx_base_vocoder_keeps_native_dtype(monkeypatch):


 class TestLTXDiffusionDecoder:
+    def test_diffusion_decoder_reuses_diffusers_with_scoped_overrides(self, monkeypatch):
+        from diffusers.models.autoencoders.ltx2_diffusion_decoder import (
+            LTX2VideoDiffusionDecoderModel as DiffusersModel,
+        )
+
+        from vllm_omni.diffusion.models.ltx2.vae.decoder import (
+            LTX2VideoDiffusionDecoderModel,
+            LTX2VideoVaeDiffusionNABlock,
+            LTX2VideoVaeNeighborhoodAttention,
+            LTX2VideoVaeSwiGLU,
+        )
+
+        original_init = DiffusersModel.__init__
+        original_parameters = {}
+
+        def capture_parameters(self, *args, **kwargs):
+            original_init(self, *args, **kwargs)
+            original_parameters.update(self.named_parameters())
+
+        with torch.device("meta"):
+            reference = DiffusersModel()
+            monkeypatch.setattr(DiffusersModel, "__init__", capture_parameters)
+            model = LTX2VideoDiffusionDecoderModel()
+        assert isinstance(model, DiffusersModel)
+        assert all(parameter is original_parameters[name] for name, parameter in model.named_parameters())
+        assert {name: tensor.shape for name, tensor in model.state_dict().items()} == {
+            name: tensor.shape for name, tensor in reference.state_dict().items()
+        }
+
+        modules = tuple(model.decoder.modules())
+        expected_blocks = sum(model.config.decoder_stage_depths)
+        assert sum(isinstance(module, LTX2VideoVaeNeighborhoodAttention) for module in modules) == expected_blocks
+        assert sum(isinstance(module, LTX2VideoVaeSwiGLU) for module in modules) == expected_blocks
+        assert (
+            sum(isinstance(module, LTX2VideoVaeDiffusionNABlock) for module in modules)
+            == model.config.decoder_stage_depths[-1]
+        )
+
+        assert type(model.decoder).__module__.startswith("vllm_omni.")
+        assert model.tiled_decode.__func__ is DiffusersModel.tiled_decode
+        assert type(reference.decoder) is type(model.decoder).__bases__[0]
+
+    def test_natten_processor_routes_stage5_through_tilelang_fna(self, monkeypatch):
+        from vllm_omni.diffusion.models.ltx2.vae import decoder as diffvae_modules
+
+        processor = object.__new__(diffvae_modules.LTX2VideoVaeNeighborhoodNattenProcessor)
+        processor.backend = None
+        fallback_calls = []
+        processor._na3d = lambda *args, **kwargs: fallback_calls.append((args, kwargs))
+
+        hidden_states = torch.randn(1, 11, 11, 11, 256, dtype=torch.bfloat16)
+        projected = hidden_states.view(1, 11, 11, 11, 4, 64)
+
+        class FakeAttention:
+            kernel_size = (11, 11, 11)
+            head_dim = 64
+            to_out = (torch.nn.Identity(),)
+
+            @staticmethod
+            def project_qkv(_hidden_states):
+                return projected, projected, projected
+
+        def token_permute(value, _tile, *, flip_tiled_dims):
+            assert flip_tiled_dims
+            return value, value.shape, None
+
+        def token_unpermute(value, *, token_layout_shape, tile_shape, flip_tiled_dims):
+            assert token_layout_shape == projected.shape
+            assert tile_shape == (4, 4, 4)
+            assert flip_tiled_dims
+            return value
+
+        fast_path_calls = []
+
+        def fast_path(query, key, value, *, shape, window):
+            fast_path_calls.append((query, key, value, shape, window))
+            return query
+
+        monkeypatch.setattr(diffvae_modules, "is_ltx2_fna_eligible", lambda _tensor: True)
+        monkeypatch.setattr(
+            diffvae_modules, "resolve_ltx2_vae_operators", lambda _device: _OperatorSetProbe(fna=fast_path)
+        )
+        monkeypatch.setattr(
+            diffvae_modules,
+            "_load_natten_token_permutation",
+            lambda: (token_permute, token_unpermute),
+        )
+
+        output = processor(FakeAttention(), hidden_states)
+
+        assert torch.equal(output, hidden_states)
+        assert len(fast_path_calls) == 1
+        assert fast_path_calls[0][3:] == ((11, 11, 11), (11, 11, 11))
+        assert not fallback_calls
+
+    @pytest.mark.parametrize("reason", ["device", "head_dim", "channels", "dtype", "tilelang_missing", "kernel"])
+    def test_natten_processor_keeps_natten_when_fast_path_does_not_apply(self, monkeypatch, reason):
+        from vllm_omni.diffusion.models.ltx2.vae import decoder as ops
+
+        processor = object.__new__(ops.LTX2VideoVaeNeighborhoodNattenProcessor)
+        processor.backend = None
+        channels = 128 if reason == "channels" else 256
+        head_dim = 32 if reason == "head_dim" else 64
+        x = torch.zeros(1, 11, 11, 11, channels, dtype=torch.float32 if reason == "dtype" else torch.bfloat16)
+        projected = x.view(1, 11, 11, 11, channels // head_dim, head_dim)
+        attn = _AttentionProbe(
+            kernel_size=(3, 5, 5) if reason == "kernel" else (11, 11, 11),
+            head_dim=head_dim,
+            to_out=(torch.nn.Identity(),),
+            projected=projected,
+        )
+        calls = []
+
+        def natten(q, k, v, **kwargs):
+            calls.append(kwargs)
+            return q
+
+        def unexpected_native_load(*_args, **_kwargs):
+            raise AssertionError("Native FNA must not be loaded for an unsupported configuration")
+
+        processor._na3d = natten
+        monkeypatch.setattr(ops, "is_ltx2_fna_eligible", lambda _tensor: reason != "device")
+        monkeypatch.setattr(
+            ops,
+            "resolve_ltx2_vae_operators",
+            lambda _device: _OperatorSetProbe(fna=None if reason == "tilelang_missing" else unexpected_native_load),
+        )
+        assert torch.equal(processor(attn, x), x)
+        assert calls == [{"kernel_size": attn.kernel_size, "scale": 1.0, "backend": None}]
+
+    def test_diffusion_decoder_noneligible_block_uses_diffusers_forward(self, monkeypatch):
+        from vllm_omni.diffusion.models.ltx2.vae import decoder as diffvae_modules
+
+        block = object.__new__(diffvae_modules.LTX2VideoVaeDiffusionNABlock)
+        torch.nn.Module.__init__(block)
+        expected = torch.randn(1, 1, 1, 1, 1)
+        calls = []
+
+        def upstream_forward(_self, hidden_states, latent_context, modulation, block_mask):
+            assert block_mask is None
+            calls.append((hidden_states, latent_context, modulation))
+            return expected
+
+        monkeypatch.setattr(
+            diffvae_modules.diffusers_decoder.LTX2VideoVaeDiffusionNABlock,
+            "forward",
+            upstream_forward,
+        )
+        hidden_states = torch.randn(1, 1, 1, 1, 1)
+        latent_context = torch.randn_like(hidden_states)
+        modulation = (torch.randn_like(hidden_states),)
+
+        assert block.forward(hidden_states, latent_context, modulation) is expected
+        assert calls == [(hidden_states, latent_context, modulation)]
+
    @pytest.mark.parametrize("mode", ["spatial_shard_height", "spatial_shard_width"])
    def test_distributed_diffusion_decoder_rejects_non_tile_parallel_modes(self, mode):
-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder_distributed import (
+        from vllm_omni.diffusion.models.ltx2.vae.distributed import (
            DistributedLTX2VideoDiffusionDecoderModel,
        )

@@ -52,7 +248,7 @@ class TestLTXDiffusionDecoder:
            model.set_parallel_size(2, mode=mode)

    def test_distributed_diffusion_decoder_accepts_tile_parallel_mode(self):
-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder_distributed import (
+        from vllm_omni.diffusion.models.ltx2.vae.distributed import (
            DistributedLTX2VideoDiffusionDecoderModel,
        )

@@ -68,7 +264,7 @@ class TestLTXDiffusionDecoder:
        assert calls == [(4, "tile")]

    def test_short_clip_keeps_stage5_temporal_context_then_crops_output(self):
-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder import (
+        from vllm_omni.diffusion.models.ltx2.vae.decoder import (
            LTX2VideoDiffusionDecoder3d,
            LTX2VideoDiffusionDecoderModel,
        )
@@ -143,7 +339,7 @@ class TestLTXDiffusionDecoder:
            _ltx2_use_diffusion_decoder(SimpleNamespace(extras={"ltx2_use_conv_vae": "true"}), "2.5")

    def test_native_diffusion_decoder_conversion_splits_qkv_and_folds_gates(self):
-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder import (
+        from vllm_omni.diffusion.models.ltx2.vae.decoder import (
            convert_ltx25_native_diffusion_decoder_state_dict,
        )

@@ -190,7 +386,7 @@ class TestLTXDiffusionDecoder:
        assert not any("type_emb" in key or "coarse" in key or "gate_msa" in key for key in converted)

    def test_native_diffusion_decoder_conversion_rejects_invalid_fused_qkv(self):
-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder import (
+        from vllm_omni.diffusion.models.ltx2.vae.decoder import (
            convert_ltx25_native_diffusion_decoder_state_dict,
        )

@@ -388,8 +584,8 @@ class TestLTXDiffusionDecoder:
        assert output.output[1].numel() == 0

    def test_diffusion_decoder_patch_parallel_size_one_uses_native_tiling(self, monkeypatch):
-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder import LTX2VideoDiffusionDecoderModel
-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder_distributed import (
+        from vllm_omni.diffusion.models.ltx2.vae.decoder import LTX2VideoDiffusionDecoderModel
+        from vllm_omni.diffusion.models.ltx2.vae.distributed import (
            DistributedLTX2VideoDiffusionDecoderModel,
        )

@@ -424,7 +620,7 @@ class TestLTXDiffusionDecoder:
    def test_distributed_diffusion_tiles_preserve_serial_noise_order(self, monkeypatch):
        from diffusers.utils.torch_utils import randn_tensor

-        from vllm_omni.diffusion.models.ltx2.ltx2_diffusion_decoder_distributed import (
+        from vllm_omni.diffusion.models.ltx2.vae.distributed import (
            DistributedLTX2VideoDiffusionDecoderModel,
        )

@@ -1,7 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

-"""CUDA regression tests for LTX-2 BWE vocoder precision."""
+"""GPU regression tests for LTX-2 BWE vocoder precision."""

 import pytest
 import torch
@@ -14,10 +14,13 @@ pytestmark = [pytest.mark.core_model, pytest.mark.diffusion]
 class _BWEConvVocoder(torch.nn.Module):
    def __init__(self):
        super().__init__()
+        # Match the production LTX vocoder's conv_pre geometry. The previous
+        # single-channel 1x1 probe intermittently had no MIOpen algorithm.
        self.conv = torch.nn.Conv1d(
-            1,
-            1,
-            kernel_size=1,
+            128,
+            1024,
+            kernel_size=7,
+            padding=3,
            bias=False,
            device="cuda",
            dtype=torch.bfloat16,
@@ -33,12 +36,12 @@ class _BWEConvVocoder(torch.nn.Module):
        return output


-@hardware_test(res={"cuda": "L4"}, num_cards=1)
+@hardware_test(res={"cuda": "L4", "rocm": "MI325"}, num_cards=1)
 def test_ltx_bwe_vocoder_runs_real_cuda_autocast_in_fp32():
    from vllm_omni.diffusion.models.ltx2.ltx2_runtime import _run_ltx_vocoder

    vocoder = _BWEConvVocoder()
-    generated_mel = torch.ones((1, 1, 4), device="cuda", dtype=torch.bfloat16)
+    generated_mel = torch.ones((1, 128, 64), device="cuda", dtype=torch.bfloat16)

    output = _run_ltx_vocoder(vocoder, generated_mel)

@@ -13,12 +13,14 @@ from PIL import Image
 pytestmark = [pytest.mark.core_model, pytest.mark.cpu, pytest.mark.diffusion]


-def test_h3_profiler_targets_cover_aggregate_reference_encoders():
+def test_h3_profiler_targets_cover_reference_encoders_and_vae_decoders():
    from vllm_omni.diffusion.models.minimax_h3 import MiniMaxH3Pipeline

    targets = MiniMaxH3Pipeline._PROFILER_TARGETS
    assert "_encode_visual_conditions" in targets
    assert "_encode_reference_audio_conditions" in targets
+    assert "video_vae.decode_latent" in targets
+    assert "audio_vae.decode_latent" in targets
    assert "_encode_video_conditions" not in targets
    assert "_encode_video_audio_conditions" not in targets
    assert "_encode_audio_conditions" not in targets
@@ -0,0 +1,357 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
+import gc
+import tempfile
+from types import SimpleNamespace
+
+import pytest
+import torch
+import torch.distributed as dist
+
+from vllm_omni.diffusion.models.minimax_h3.chunked_decode import decode_h3_chunks
+from vllm_omni.diffusion.models.minimax_h3.temporal_chunks import decode_temporal_chunks
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu, pytest.mark.diffusion]
+
+
+@pytest.fixture(scope="module")
+def cpu_process_group():
+    if dist.is_initialized():
+        yield dist.group.WORLD
+        return
+
+    with tempfile.NamedTemporaryFile(prefix="h3_chunk_dist_") as rendezvous:
+        init_method = f"file://{rendezvous.name}"
+    dist.init_process_group("gloo", rank=0, world_size=1, init_method=init_method)
+    try:
+        yield dist.group.WORLD
+    finally:
+        dist.destroy_process_group()
+        gc.collect()
+
+
+class _FakeTemporalModel:
+    use_3d_conv = True
+    token_drop = 3
+    tokens_chunk_size = 5
+    token_overlap = 2
+    vae_ratio_t = 4
+    frame_pre_padding = 3
+    frame_overlap = 5
+    isolated_first_frame = False
+    isolated_last_frame = False
+
+    def _decode_temporal_output_frame_plan(self, z, z_head, z_tail, num_chunks, pad_tokens):
+        del z, z_head, z_tail, num_chunks, pad_tokens
+        return 35, 0, 35
+
+    def _adaptive_decode(self, clip):
+        value = float(clip[:, :, 0].mean())
+        return torch.full((1, 1, 24, 2, 2), value)
+
+    @staticmethod
+    def blend(overlap, part, frame_overlap, dim):
+        del overlap, frame_overlap, dim
+        return part
+
+
+class _FakeHost:
+    def __init__(self):
+        self.model = _FakeTemporalModel()
+
+    @staticmethod
+    def _denormalize_latent(latent):
+        return latent
+
+    @staticmethod
+    def _normalize_decoded_frames(frames):
+        return frames.float()
+
+
+def test_temporal_chunks_emit_ordered_frames_and_collect_when_unconsumed():
+    model = _FakeTemporalModel()
+    latent = torch.arange(8, dtype=torch.float32).view(1, 1, 8, 1, 1)
+    chunks = []
+    marker = decode_temporal_chunks(model, latent, chunks.append)
+
+    assert marker.shape == (0,)
+    assert [chunk.shape[2] for chunk in chunks] == [17, 17, 1]
+    assert torch.equal(torch.cat(chunks, dim=2), decode_temporal_chunks(model, latent, None))
+
+
+def test_h3_callback_failure_is_deferred_until_temporal_decode_finishes():
+    host = _FakeHost()
+    latent = torch.zeros(1, 1, 8, 1, 1)
+    seen = []
+
+    def fail_once(frames):
+        seen.append(frames.shape[2])
+        raise RuntimeError("sink failed")
+
+    with pytest.raises(RuntimeError, match="sink failed"):
+        decode_h3_chunks(host, latent, fail_once, group=None)
+    assert seen == [17]
+
+
+def test_h3_without_callback_is_a_plain_full_decode_on_a_vae_group(cpu_process_group):
+    host = _FakeHost()
+    latent = torch.arange(8, dtype=torch.float32).view(1, 1, 8, 1, 1)
+
+    full = decode_h3_chunks(host, latent, None, group=cpu_process_group)
+
+    chunks = []
+    decode_h3_chunks(host, latent, chunks.append, group=cpu_process_group)
+    assert torch.equal(torch.cat(chunks, dim=2), full)
+
+
+class _SimulatedGroup:
+    """Stand in for a VAE group so one process can check every rank's verdict.
+
+    ``all_reduce`` is replaced with the sum the real collective would produce
+    for a group of ``world_size`` ranks, ``supplied`` of which handed in a
+    callback. That sum is what makes a rank-local verdict observable from a
+    single process.
+    """
+
+    def __init__(self, world_size: int, supplied: int):
+        self.world_size = world_size
+        self.supplied = supplied
+
+
+def _patch_simulated_ranks(monkeypatch, simulated: _SimulatedGroup, rank: int) -> None:
+    from vllm_omni.diffusion.models.minimax_h3 import chunked_decode as mod
+
+    def fake_all_reduce(tensor, group=None):
+        del group  # the simulated census below stands in for the real reduction
+        tensor.copy_(torch.tensor([simulated.supplied], dtype=tensor.dtype))
+
+    monkeypatch.setattr(mod.dist, "get_rank", lambda _group: rank)
+    monkeypatch.setattr(mod.dist, "get_world_size", lambda _group: simulated.world_size)
+    monkeypatch.setattr(mod.dist, "all_reduce", fake_all_reduce)
+    monkeypatch.setattr(mod.dist, "broadcast", lambda tensor, src=0, group=None: None)
+    monkeypatch.setattr(mod.dist, "get_global_rank", lambda _group, _rank: 0)
+
+
+@pytest.mark.parametrize("rank", [0, 1, 2])
+def test_h3_every_rank_rejects_a_partially_supplied_callback(monkeypatch, rank):
+    """Ranks that disagree must all refuse, not let one decode alone.
+
+    A rank-local verdict let a peer fall through into the decoder collectives
+    while its peers raised, hanging the stage instead of failing it.
+    """
+    host = _FakeHost()
+    latent = torch.zeros(1, 1, 8, 1, 1)
+    # Two ranks of three handed in a callback.
+    group = _SimulatedGroup(world_size=3, supplied=2)
+    _patch_simulated_ranks(monkeypatch, group, rank)
+
+    with pytest.raises(ValueError, match="every rank"):
+        decode_h3_chunks(host, latent, lambda _frames: None, group=object())
+
+
+@pytest.mark.parametrize("rank", [0, 1, 2])
+def test_h3_only_rank_zero_publishes_when_every_rank_supplies(monkeypatch, rank):
+    """Ownership is positional: all ranks stream, rank 0 publishes."""
+    host = _FakeHost()
+    latent = torch.zeros(1, 1, 8, 1, 1)
+    group = _SimulatedGroup(world_size=3, supplied=3)
+    _patch_simulated_ranks(monkeypatch, group, rank)
+
+    seen: list[int] = []
+    decode_h3_chunks(host, latent, lambda f: seen.append(int(f.shape[2])), group=object())
+
+    # Peers run the loop to stay in the collectives but publish nothing.
+    assert (seen != []) is (rank == 0)
+
+
+def _stub_video_vae(*, parallel_size: int, tile_count: int):
+    """A MiniMaxH3VideoVAE with only the pieces the decode entry points read."""
+    from contextlib import contextmanager
+
+    from vllm_omni.diffusion.models.minimax_h3.vae import MiniMaxH3VideoVAE
+
+    vae = object.__new__(MiniMaxH3VideoVAE)
+    vae.parallel_size = parallel_size
+    vae.model = SimpleNamespace(_adaptive_decode=lambda clip: clip)
+    vae.entered_rank_local = False
+
+    @contextmanager
+    def rank_local():
+        vae.entered_rank_local = True
+        yield
+
+    vae._decoder_tile_count = lambda latent: tile_count
+    vae._rank_local_tiling = rank_local
+    vae.is_distributed_enabled = lambda: False
+    return vae
+
+
+def test_chunked_decode_falls_back_when_tiles_are_fewer_than_ranks(monkeypatch):
+    """Without the fallback, tileless ranks raise while their peers block."""
+    from vllm_omni.diffusion.models.minimax_h3 import vae as vae_mod
+
+    vae = _stub_video_vae(parallel_size=4, tile_count=1)
+    monkeypatch.setattr(vae_mod, "decode_h3_chunks", lambda *args, **kwargs: torch.zeros(1))
+
+    vae.decode_with_chunks(torch.zeros(1, 1, 8, 1, 1), on_chunk=lambda _frames: None)
+
+    assert vae.entered_rank_local, "chunked decode must share decode_latent's too-few-tiles fallback"
+
+
+def test_chunked_decode_keeps_shared_tiling_when_every_rank_has_a_tile(monkeypatch):
+    from vllm_omni.diffusion.models.minimax_h3 import vae as vae_mod
+
+    vae = _stub_video_vae(parallel_size=4, tile_count=8)
+    monkeypatch.setattr(vae_mod, "decode_h3_chunks", lambda *args, **kwargs: torch.zeros(1))
+
+    vae.decode_with_chunks(torch.zeros(1, 1, 8, 1, 1), on_chunk=lambda _frames: None)
+
+    assert not vae.entered_rank_local
+
+
+def test_h3_vae_declares_the_chunked_decode_capability():
+    """A consumer must be able to detect the capability and the pixel range."""
+    from vllm_omni.diffusion.models.interface import supports_chunked_vae_decode
+    from vllm_omni.diffusion.models.minimax_h3.vae import MiniMaxH3VideoVAE
+
+    vae = object.__new__(MiniMaxH3VideoVAE)
+
+    assert supports_chunked_vae_decode(vae)
+    # H3 reverts through the checkpoint's processor, which lands in [0, 1];
+    # a Wan VAE publishes [-1, 1], so the range cannot be assumed.
+    assert vae.chunk_value_range == (0.0, 1.0)
+
+
+def _native_decode_temporal_streaming(model, z, z_head, z_tail, num_chunks, pad_tokens):
+    """The released ``_decode_temporal_streaming`` loop, as the parity oracle.
+
+    Transcribed from the checkpoint's ``AutoencoderKLLegacy`` so the fork can be
+    compared against the control flow it mirrors, including the isolated
+    head/tail extraction.
+    """
+    total_frames, pad_frames, output_frames = model._decode_temporal_output_frame_plan(
+        z, z_head, z_tail, num_chunks, pad_tokens
+    )
+    chunk_dec = model.tokens_chunk_size * model.vae_ratio_t
+    split_count = int(model.token_drop > 0) + 1
+    dec = None
+    dec_overlap = None
+    write_pos = 0
+
+    def write_part(part):
+        nonlocal dec, write_pos
+        part_frames = int(part.shape[2])
+        if part_frames <= 0:
+            return
+        if dec is None:
+            out_shape = list(part.shape)
+            out_shape[2] = output_frames
+            dec = torch.empty(out_shape, dtype=part.dtype, device=part.device)
+        remaining = int(dec.shape[2]) - write_pos
+        copy_frames = min(part_frames, max(0, remaining))
+        if copy_frames > 0:
+            dec[:, :, write_pos : write_pos + copy_frames].copy_(part[:, :, :copy_frames])
+            write_pos += copy_frames
+
+    for i in range(num_chunks):
+        t_start = i * model.tokens_chunk_size
+        clip_z = z[:, :, t_start : t_start + model.tokens_chunk_size + model.token_overlap]
+        if i == 0 and z_head is not None:
+            clip_z = torch.cat([z_head, clip_z], dim=2)
+        if i == num_chunks - 1 and z_tail is not None:
+            clip_z = torch.cat([clip_z, z_tail], dim=2)
+
+        clip_dec = model._adaptive_decode(clip_z)
+
+        dec_tail = None
+        if i == 0 and z_head is not None:
+            write_part(clip_dec[:, :, model.vae_ratio_t - 1 : model.vae_ratio_t])
+            clip_dec = clip_dec[:, :, model.vae_ratio_t :]
+        if i == num_chunks - 1 and z_tail is not None:
+            dec_tail = clip_dec[:, :, -1:]
+            clip_dec = clip_dec[:, :, : -model.vae_ratio_t]
+
+        for j in range(split_count):
+            f_start = j * chunk_dec
+            f_end = min(f_start + chunk_dec, clip_dec.shape[2])
+            chunk = clip_dec[:, :, f_start:f_end][:, :, model.frame_pre_padding :]
+            if j == 0:
+                if dec_overlap is not None:
+                    chunk = model.blend(dec_overlap, chunk, model.frame_overlap, dim=-3)
+                    dec_overlap = None
+                write_part(chunk)
+            else:
+                dec_overlap = chunk.contiguous()
+
+        if i == num_chunks - 1:
+            if dec_overlap is not None:
+                write_part(dec_overlap)
+                dec_overlap = None
+            if dec_tail is not None:
+                write_part(dec_tail)
+
+    return dec
+
+
+class _IsolatedFrameModel(_FakeTemporalModel):
+    """Fake decoder with the isolated boundaries the released config can set."""
+
+    frame_pre_padding = 0
+
+    def __init__(self, *, first: bool, last: bool):
+        self.isolated_first_frame = first
+        self.isolated_last_frame = last
+
+    def _decode_temporal_output_frame_plan(self, z, z_head, z_tail, num_chunks, pad_tokens):
+        del pad_tokens
+        frames = num_chunks * self.tokens_chunk_size * self.vae_ratio_t
+        frames += int(z_head is not None) + int(z_tail is not None)
+        del z
+        return frames, 0, frames
+
+    def _adaptive_decode(self, clip):
+        # Every latent token decodes to a distinct, ordered block, so a shifted
+        # or reordered emission is visible in the values themselves.
+        base = float(clip[:, :, 0].mean())
+        frames = int(clip.shape[2]) * self.vae_ratio_t
+        return torch.arange(frames, dtype=torch.float32).view(1, 1, frames, 1, 1) + base
+
+
+@pytest.mark.parametrize(
+    ("first", "last"),
+    [(True, False), (False, True), (True, True)],
+)
+def test_temporal_chunks_match_the_native_loop_at_isolated_boundaries(first, last):
+    """The fork must extract the isolated head/tail the way the native loop does."""
+    model = _IsolatedFrameModel(first=first, last=last)
+    latent = torch.arange(12, dtype=torch.float32).view(1, 1, 12, 1, 1)
+
+    chunks: list[torch.Tensor] = []
+    decode_temporal_chunks(model, latent, chunks.append)
+    streamed = torch.cat(chunks, dim=2)
+
+    # Feed the oracle exactly what the released ``decode_temporal`` would:
+    # split the isolated tokens off, then pad the remainder to a whole chunk.
+    isolated = int(first) + int(last)
+    pseudo_tokens = int(latent.shape[2]) - isolated + model.token_drop
+    remainder = pseudo_tokens % model.tokens_chunk_size
+    pad_tokens = (model.tokens_chunk_size - remainder) if remainder else 0
+    num_chunks = (pseudo_tokens + pad_tokens) // model.tokens_chunk_size - int(model.token_drop > 0)
+
+    body = latent
+    z_head = None
+    if first:
+        z_head = body[:, :, :1]
+        body = body[:, :, 1:]
+    z_tail = None
+    if last:
+        z_tail = body[:, :, -1:]
+        body = body[:, :, :-1]
+    if pad_tokens:
+        body = torch.cat([body, body[:, :, -1:].repeat(1, 1, pad_tokens, 1, 1)], dim=2)
+
+    expected = _native_decode_temporal_streaming(model, body, z_head, z_tail, num_chunks, pad_tokens)
+
+    assert streamed.shape == expected.shape
+    assert torch.equal(streamed, expected)
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 from types import SimpleNamespace

 import pytest
@@ -64,6 +64,13 @@ def setup_tp_group(monkeypatch, mocker):
        lambda: 0,
    )

+    # vLLM 0.29 added a CUDA-only fused embedding kernel that engages whenever
+    # tp_size > 1; these tests drive CPU tensors, so keep the portable masked path.
+    monkeypatch.setattr(
+        "vllm.model_executor.layers.vocab_parallel_embedding.current_platform.is_cuda",
+        lambda: False,
+    )
+
    mock_tp_group = mocker.MagicMock()
    mock_tp_group.world_size = 2
    mocker.patch("vllm.distributed.parallel_state.get_tp_group", return_value=mock_tp_group)
@@ -94,7 +101,7 @@ def setup_tp_group(monkeypatch, mocker):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )

    with set_current_vllm_config(VllmConfig(device_config=device_config)):
@@ -1,12 +1,39 @@
 import pytest
 import torch

-from vllm_omni.diffusion.models.qwen_image.qwen_image_transformer import QwenEmbedRope
+from vllm_omni.diffusion.models.qwen_image.qwen_image_transformer import (
+    QwenEmbedRope,
+    _apply_qwen_image_rotary_emb,
+)
 from vllm_omni.diffusion.models.qwen_image.rope_utils import txt_seq_lens_from_embeds

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


+@pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16])
+@pytest.mark.parametrize("noncontiguous", [False, True])
+def test_qwen_rotary_preserves_frequency_precision(dtype, noncontiguous):
+    from diffusers.models.transformers.transformer_qwenimage import apply_rotary_emb_qwen
+
+    generator = torch.Generator().manual_seed(42)
+    x = torch.randn(2, 17, 3, 128, generator=generator).to(dtype)
+    if noncontiguous:
+        x = x.transpose(1, 2).contiguous().transpose(1, 2)
+    angles = torch.randn(17, 64, generator=generator)
+    freqs = torch.polar(torch.ones_like(angles), angles)
+    expected = apply_rotary_emb_qwen(x, freqs, use_real=False)
+    actual = _apply_qwen_image_rotary_emb(x, freqs)
+    assert actual.dtype == dtype
+    assert actual.shape == x.shape
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+
+    if dtype == torch.bfloat16:
+        # A regression to early BF16 frequency rounding must be observable.
+        rounded_freqs = torch.complex(freqs.real.to(dtype).float(), freqs.imag.to(dtype).float())
+        rounded = _apply_qwen_image_rotary_emb(x, rounded_freqs)
+        assert not torch.equal(actual, rounded)
+
+
 def test_txt_seq_lens_from_embeds_uses_padded_width_not_valid_token_count():
    prompt_embeds = torch.zeros(2, 64, 8)
    prompt_embeds_mask = torch.zeros(2, 64, dtype=torch.bool)
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 import os

@@ -22,7 +22,7 @@ def _init_distributed(monkeypatch):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )
    os.environ.setdefault("MASTER_ADDR", "localhost")
    os.environ.setdefault("MASTER_PORT", "29501")
@@ -41,8 +41,8 @@ def tp1_group():

    os.environ.setdefault("MASTER_ADDR", "localhost")
    os.environ.setdefault("MASTER_PORT", "29501")
-    init_distributed_environment(world_size=1, rank=0, local_rank=0, distributed_init_method="env://")
-    initialize_model_parallel()
+    init_distributed_environment(world_size=1, rank=0, local_rank=0, distributed_init_method="env://", backend="gloo")
+    initialize_model_parallel(backend="gloo")
    yield
    cleanup_dist_env_and_memory()

@@ -54,7 +54,7 @@ def force_default_gemm(monkeypatch):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )


@@ -482,8 +482,12 @@ class _LockstepP2PDist:
 def test_tiny_transformer_sp2_lockstep_matches_dense(tp1_group, force_default_gemm, mocker, task) -> None:
    """Two lockstep rank threads over uneven frame shards must reproduce the
    dense tiny-transformer output, and must actually have sharded the work."""
+    from vllm_omni.diffusion.attention.backends.sdpa import SDPABackend
    from vllm_omni.diffusion.models.sana_video import SanaVideoTransformer3DModel

+    # These CPU tensor tests must not probe a CUDA device during construction.
+    mocker.patch("vllm_omni.diffusion.attention.selector._cached_get_backend_cls", return_value=SDPABackend)
+
    torch.manual_seed(3)
    model = SanaVideoTransformer3DModel(**_TINY_SP_CONFIG).eval()
    for _, param in sorted(model.named_parameters()):
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 """TP unit tests for SANA-Video: distributed RMSNorm, sharded attention and
 sharded GLUMB temporal conv.
@@ -51,7 +51,7 @@ def force_default_gemm(monkeypatch):

    monkeypatch.setattr(
        "vllm.model_executor.layers.linear.dispatch_unquantized_gemm",
-        lambda: default_unquantized_gemm,
+        lambda *_args, **_kwargs: default_unquantized_gemm,
    )


@@ -17,6 +17,7 @@ import torch
 from vllm.lora.layers import BaseLayerWithLoRA

 import vllm_omni.diffusion.models.sensenova_u1.paged_decode as paged_decode
+from tests.helpers.mark import hardware_test
 from vllm_omni.diffusion.models.sensenova_u1.paged_decode import (
    BLOCK_SIZE,
    BUCKETS,
@@ -169,7 +170,8 @@ def test_length_lives_in_device_tensors():


 # ---------------------------------------------------------------------------
-# The capture contract. These need CUDA and the bundled flash-attention kernel.
+# The capture contract. Graph tests need a torch.cuda-compatible accelerator;
+# paged-attention equivalence also needs the bundled kernel's paged API.
 # ---------------------------------------------------------------------------

 cuda_only = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
@@ -181,8 +183,7 @@ vllm_flash_attn_only = pytest.mark.skipif(

@cuda_only
@vllm_flash_attn_only
-@pytest.mark.cuda
-@pytest.mark.L4
+@hardware_test(res={"cuda": "L4", "rocm": "MI325"})
 def test_paged_attention_matches_sdpa_over_the_same_buffer():
    import torch.nn.functional as F

@@ -223,8 +224,7 @@ def test_paged_attention_matches_sdpa_over_the_same_buffer():

@cuda_only
@vllm_flash_attn_only
-@pytest.mark.cuda
-@pytest.mark.L4
+@hardware_test(res={"cuda": "L4", "rocm": "MI325"})
 def test_a_captured_graph_follows_a_later_length():
    """One capture has to serve every length in the bucket.

@@ -285,8 +285,7 @@ def test_a_captured_graph_follows_a_later_length():

@cuda_only
@vllm_flash_attn_only
-@pytest.mark.cuda
-@pytest.mark.L4
+@hardware_test(res={"cuda": "L4", "rocm": "MI325"})
 def test_a_captured_graph_writes_each_step_to_its_own_slot():
    """The write slot has to be a device tensor, not a Python index.

@@ -347,8 +346,7 @@ def test_a_captured_graph_writes_each_step_to_its_own_slot():

@cuda_only
@vllm_flash_attn_only
-@pytest.mark.cuda
-@pytest.mark.L4
+@hardware_test(res={"cuda": "L4", "rocm": "MI325"})
 def test_a_whole_decode_loop_matches_the_unpaged_path():
    """Per-step kernel equivalence is not loop equivalence.

@@ -413,8 +411,7 @@ class _StubLM(torch.nn.Module):


@cuda_only
-@pytest.mark.cuda
-@pytest.mark.L4
+@hardware_test(res={"cuda": "L4", "rocm": "MI325"})
 def test_the_decode_graph_captures_into_the_shared_platform_pool(monkeypatch):
    """``_decode_context`` builds a runner per request, so a private capture pool
    hands every request its own arena and never gives it back. Measured on a
@@ -556,10 +553,12 @@ def test_a_kernel_missing_any_kwarg_we_pass_is_not_supported(monkeypatch):
 def test_only_q_k_and_v_reach_the_kernel_positionally(monkeypatch):
    """A positional standard argument binds to whatever a future wheel puts in
    that slot, and the probe above only checks that the names exist."""
-    seen = {}
+    seen_args: tuple[object, ...] = ()
+    seen_kwargs: dict[str, object] = {}

-    def recorder(*args, **kwargs):
-        seen["args"], seen["kwargs"] = args, kwargs
+    def recorder(*args: object, **kwargs: object):
+        nonlocal seen_args, seen_kwargs
+        seen_args, seen_kwargs = args, kwargs
        return torch.zeros(1, N_HEADS, HEAD_DIM)

    monkeypatch.setattr(paged_decode, "_flash_varlen", lambda: recorder)
@@ -572,8 +571,8 @@ def test_only_q_k_and_v_reach_the_kernel_positionally(monkeypatch):
        torch.zeros(1, KV_HEADS, 1, HEAD_DIM),
        1.0,
    )
-    assert len(seen["args"]) == 3, "only q, k and v may be positional"
-    assert set(seen["kwargs"]) == {
+    assert len(seen_args) == 3, "only q, k and v may be positional"
+    assert set(seen_kwargs) == {
        "max_seqlen_q",
        "cu_seqlens_q",
        "max_seqlen_k",
@@ -694,8 +693,7 @@ def test_no_capture_crosses_a_request_once_lora_wrappers_are_installed(monkeypat


@cuda_only
-@pytest.mark.cuda
-@pytest.mark.L4
+@hardware_test(res={"cuda": "L4", "rocm": "MI325"})
 def test_a_second_request_replays_the_first_capture(monkeypatch):
    """Reuse is only worth anything if the capture survives it."""
    dev = torch.device("cuda")
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 import pytest
 import torch
 from transformers import T5Config
@@ -54,6 +54,13 @@ def setup_tp_group(monkeypatch, mocker):
        lambda: 0,
    )

+    # vLLM 0.29 added a CUDA-only fused embedding kernel that engages whenever
+    # tp_size > 1; these tests drive CPU tensors, so keep the portable masked path.
+    monkeypatch.setattr(
+        "vllm.model_executor.layers.vocab_parallel_embedding.current_platform.is_cuda",
+        lambda: False,
+    )
+
    # TP group
    mock_tp_group = mocker.MagicMock()
    mock_tp_group.world_size = 2
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Tests for component routing for quantization."""

 from unittest.mock import MagicMock
@@ -372,7 +372,8 @@ class TestComponentApplyVllmMapper:
                "thinker.model": "thinker.language_model.model",
                "talker.model": "talker.language_model.model",
            }
-        ).get_unstacked_mapper()
+            # vLLM 0.29 renamed get_unstacked_mapper to get_rename_mapper.
+        ).get_rename_mapper()

        config.apply_vllm_mapper(mapper)

@@ -300,7 +300,8 @@ def test_single_stage_zimage_fp8_uses_less_memory():
    assert mem_fp8 < mem_bf16, f"FP8 ({mem_fp8:.2f} GiB) should use less memory than BF16 ({mem_bf16:.2f} GiB)"


-@hardware_test(res={"cuda": "L4"})
+@hardware_test(res={"cuda": "L4", "xpu": "B60"})
+@pytest.mark.advanced_model
 def test_single_stage_qwen_image_fp8():
    """Qwen-Image (random weights) with FP8 generates valid images."""
    model = "riverclouds/qwen_image_random"
@@ -140,6 +140,7 @@ class TestMetricKeys:
                raise AssertionError("diffusion_engine_total_time_ms should be attached in step_streaming()")

        assert '"diffusion_engine_exec_time_ms": exec_total_time * 1000' in step_streaming_source
+        assert '"output_ready_wait_time_ms": output_ready_wait_time * 1000' in step_streaming_source
        # step_total_ms is no longer emitted as a metric (no family consumes
        # it); it lives only in the debug log breakdown below.
        assert '"diffusion_engine_total_time_ms"' not in step_streaming_source
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 import asyncio
 import queue
@@ -10,8 +10,9 @@ import pytest
 import torch
 import vllm.v1.core.single_type_kv_cache_manager as native_kv_managers
 from pytest_mock import MockerFixture
-from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec, KVCacheTensor
+from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec

+from tests.helpers.kv_layout import build_kv_cache_tensor
 from vllm_omni.diffusion.data import DiffusionOutput, DiffusionRequestAbortedError
 from vllm_omni.diffusion.diffusion_engine import DiffusionEngine, DiffusionExecutionMode
 from vllm_omni.diffusion.diffusion_kv.config import DiffusionKVCacheMode
@@ -104,7 +105,7 @@ def _initialize_paged_scheduler(
    )
    config = KVCacheConfig(
        num_blocks=num_blocks,
-        kv_cache_tensors=[KVCacheTensor(size=spec.page_size_bytes * num_blocks, shared_by=["layer0"])],
+        kv_cache_tensors=[build_kv_cache_tensor(spec, num_blocks, ["layer0"])],
        kv_cache_groups=[KVCacheGroupSpec(layer_names=["layer0"], kv_cache_spec=spec)],
    )
    scheduler.initialize(
@@ -4,8 +4,10 @@
 import asyncio
 import multiprocessing as mp
 import queue
+import signal
 import threading
 import time
+import weakref
 from types import SimpleNamespace
 from unittest.mock import MagicMock, Mock

@@ -1330,6 +1332,7 @@ class TestExecutorShutdownCleaner:
        class FakeProcess:
            def __init__(self, name):
                self.name = name
+                self.pid = 123
                self.alive = True
                self.terminated = False
                self.join_timeouts = []
@@ -1358,6 +1361,132 @@ class TestExecutorShutdownCleaner:
        assert first.terminated and second.terminated
        assert not first.is_alive() and not second.is_alive()

+    @pytest.mark.parametrize("cooperative", [False, True])
+    def test_real_workers_exit_and_are_reaped(self, monkeypatch, cooperative):
+        from vllm_omni.diffusion.executor import multiproc_executor as executor_module
+
+        ctx = mp.get_context("fork")
+        ready, stop = ctx.Event(), ctx.Event()
+
+        def run_worker():
+            signal.signal(signal.SIGTERM, signal.SIG_IGN)
+            ready.set()
+            stop.wait()
+
+        dead = ctx.Process(target=lambda: None)
+        worker = ctx.Process(target=run_worker)
+        monkeypatch.setattr(executor_module, "_WORKER_SHUTDOWN_GRACE_S", 0.2)
+        monkeypatch.setattr(executor_module, "_WORKER_TERMINATE_GRACE_S", 0.1)
+        monkeypatch.setattr(executor_module, "_WORKER_KILL_GRACE_S", 5.0, raising=False)
+        dead.start()
+        worker.start()
+        try:
+            dead.join(5)
+            assert not dead.is_alive()
+            assert ready.wait(5), "worker did not install its signal handler"
+            mq = Mock()
+            if cooperative:
+                mq.enqueue.side_effect = lambda *args, **kwargs: stop.set()
+            else:
+                mq.enqueue.side_effect = OSError("shutdown queue unavailable")
+            cleaner = executor_module._ExecutorShutdownCleaner(mq, 2, [dead, worker])
+
+            cleaner()
+
+            assert not worker.is_alive()
+            assert worker.exitcode == (0 if cooperative else -signal.SIGKILL)
+            assert cleaner.processes == []
+            assert cleaner.broadcast_mq is None
+            cleaner()
+        finally:
+            for proc in (dead, worker):
+                if proc.is_alive():
+                    proc.kill()
+                proc.join(5)
+                proc.close()
+
+    @pytest.mark.parametrize("failed_action", ["terminate", "kill"])
+    def test_kill_phase_shares_deadline_and_continues_after_os_errors(self, monkeypatch, failed_action):
+        from vllm_omni.diffusion.executor import multiproc_executor as executor_module
+
+        first, second = Mock(pid=101), Mock(pid=102)
+        first.name, second.name = "first", "second"
+        for proc in (first, second):
+            proc.is_alive.side_effect = [True, True, True, False]
+        getattr(first, failed_action).side_effect = OSError("signal failed")
+        first.join.side_effect = [OSError("join failed"), None, None]
+        monotonic = Mock(side_effect=[100, 100, 110, 120, 120, 124, 130, 130, 134])
+        monkeypatch.setattr(executor_module, "time", SimpleNamespace(monotonic=monotonic))
+        cleaner = executor_module._ExecutorShutdownCleaner(processes=[first, second])
+
+        cleaner()
+
+        first.kill.assert_called_once()
+        second.kill.assert_called_once()
+        assert [call.args[0] for call in first.join.call_args_list] == [15, 5, 5]
+        assert [call.args[0] for call in second.join.call_args_list] == [5, 1, 1]
+        assert cleaner.processes == []
+
+    def test_executor_retains_survivor_for_explicit_shutdown_retry(self, monkeypatch):
+        from vllm_omni.diffusion.executor import multiproc_executor as executor_module
+
+        survivor = Mock(pid=123)
+        survivor.name = "surviving-worker"
+        survivor.is_alive.return_value = True
+        cleaner = executor_module._ExecutorShutdownCleaner(processes=[survivor])
+        executor, _, _ = _make_executor()
+        executor._shutdown_cleaner = cleaner
+        executor._processes = [survivor]
+        executor._finalizer = weakref.finalize(executor, cleaner)
+        executor._pump_stop = threading.Event()
+        executor._futures_lock = threading.RLock()
+        executor._rpc_futures = {}
+        executor._output_futures = {}
+        executor._batch_split_map = {}
+        log_error = Mock()
+        monkeypatch.setattr(executor_module.logger, "error", log_error)
+
+        executor.shutdown()
+
+        assert not executor._finalizer.alive
+        assert executor._shutdown_cleaner is cleaner
+        assert executor._processes == [survivor]
+        log_error.assert_called_once()
+        assert log_error.call_args.args[1] == [("surviving-worker", 123)]
+        survivor.is_alive.return_value = False
+
+        executor.shutdown()
+        executor.shutdown()
+
+        assert executor._shutdown_cleaner is None
+        assert executor._processes == []
+
+    def test_concurrent_cleanup_does_not_repeat_process_operations(self):
+        from vllm_omni.diffusion.executor import multiproc_executor as executor_module
+
+        joining, release = threading.Event(), threading.Event()
+        proc = Mock(pid=123)
+        proc.is_alive.return_value = True
+
+        def join(timeout):
+            joining.set()
+            assert release.wait(5)
+            proc.is_alive.return_value = False
+
+        proc.join.side_effect = join
+        cleaner = executor_module._ExecutorShutdownCleaner(processes=[proc])
+        thread = threading.Thread(target=cleaner)
+        thread.start()
+        try:
+            assert joining.wait(5)
+            cleaner()
+            proc.join.assert_called_once()
+        finally:
+            release.set()
+            thread.join(5)
+        assert not thread.is_alive()
+        assert cleaner.processes == []
+

 # ───────── monitor thread & death sentinel integration tests ─────────

@@ -1406,7 +1535,7 @@ class TestMultiprocExecutorWorkerMonitor:
        executor._result_mq = None
        executor._shutdown_cleaner = None
        # Use a no-op so shutdown() doesn't crash on None resources.
-        executor._finalizer = lambda: None
+        executor._finalizer = weakref.finalize(executor, lambda: None)
        # ------------------------------------------------------------------
        # Attributes added by remove_bubble_v2 (async D2H); shutdown() iterates
        # over them, so they need to exist even when constructed via __new__.
@@ -1441,7 +1570,7 @@ class TestMultiprocExecutorWorkerMonitor:
        executor._broadcast_mq = None
        executor._result_mq = None
        executor._shutdown_cleaner = None
-        executor._finalizer = lambda: None
+        executor._finalizer = weakref.finalize(executor, lambda: None)

        proc = _make_short_lived_process()
        executor._processes = [proc]
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project

 """Generate raw LTX video and audio outputs with the official or Omni runtime."""

@@ -114,6 +114,29 @@ def _insert_official_paths(official_root: Path) -> None:
            sys.path.insert(0, path)


+def _configure_official_vocoder_determinism() -> None:
+    """Make the pinned official LTX-2.5 vocoder a stable reference."""
+    from ltx_core.model.audio_vae.vocoder import VocoderWithBWE
+
+    original_forward = VocoderWithBWE.forward
+    if getattr(original_forward, "_vllm_omni_deterministic", False):
+        return
+
+    def deterministic_forward(self: Any, mel_spec: torch.Tensor) -> torch.Tensor:
+        device_type = mel_spec.device.type
+        if device_type != "cuda":
+            return original_forward(self, mel_spec)
+        previous = torch.backends.cudnn.deterministic
+        try:
+            torch.backends.cudnn.deterministic = True
+            return original_forward(self, mel_spec)
+        finally:
+            torch.backends.cudnn.deterministic = previous
+
+    setattr(deterministic_forward, "_vllm_omni_deterministic", True)
+    setattr(VocoderWithBWE, "forward", deterministic_forward)
+
+
 def _configure_official_sdpa(pipeline: Any) -> None:
    """Pin official connector and denoiser attention to cuDNN, with MATH fallback."""
    from ltx_core.loader.attention_ops import set_attention_module_op
@@ -187,6 +210,7 @@ def _run_official(args: argparse.Namespace, request: dict[str, Any]) -> None:
    if args.official_root is None:
        raise ValueError("Official backend requires --official-root")
    _insert_official_paths(args.official_root)
+    _configure_official_vocoder_determinism()
    # Keep Gemma on the same SDPA path in both subprocesses. DiT attention is
    # still pinned independently to the explicit cuDNN backend below.
    torch.backends.cuda.enable_cudnn_sdp(False)
@@ -164,6 +164,7 @@ def _run_diffusers_wan22_i2v(*, model: str, output_path: Path, conditioning_imag
        )
        pipe.scheduler = UniPCMultistepScheduler.from_config(pipe.scheduler.config, flow_shift=FLOW_SHIFT)
        pipe.to("cuda")
+        _set_matched_attention_backend(pipe)

        _diffusers_dummy_run(pipe)

@@ -34,11 +34,7 @@ TRUE_CFG_SCALE = 4.0
 SEED = 42
 SSIM_THRESHOLD = 0.94
 # With --fa-deterministic, H100 nightlies stably land near SSIM 0.958 / PSNR 27.8
-# vs Diffusers (see #5734 / #5963), so the primary gates are 0.94 / 27.0.
-# The SDPA-matched fallback (hub kernel unavailable for the image's torch)
-# landed at SSIM 0.9687 in build 2954 — keep a slightly tighter SSIM gate for
-# that path only.
-SSIM_THRESHOLD_SDPA_FALLBACK = 0.96
+# vs Diffusers on the FA3 hub path (see #5734 / #5963 / #6971).
 PSNR_THRESHOLD = 27.0

 MODEL_2512_ID = "Qwen/Qwen-Image-2512"
@@ -65,44 +61,6 @@ MEAN_ABS_DIFF_THRESHOLD_2512 = 3e-2
 P99_ABS_DIFF_THRESHOLD_2512 = 4e-1


-_FLASH_ATTN3_HUB_AVAILABLE: bool | None = None
-
-
-def _flash_attn3_hub_available() -> bool:
-    """Whether kernels-community/flash-attn3 has a build for this torch/CUDA.
-
-    The hub repo only ships prebuilt variants for specific torch versions;
-    a base-image torch bump can outrun it (build 2952 failed with "Cannot
-    find a build variant for this system"). When it is unavailable, both
-    sides of the comparison fall back to torch SDPA together — the test
-    compares implementations, so the attention math must stay matched.
-    """
-    global _FLASH_ATTN3_HUB_AVAILABLE
-    if _FLASH_ATTN3_HUB_AVAILABLE is None:
-        try:
-            from kernels import get_kernel
-
-            get_kernel("kernels-community/flash-attn3")
-            _FLASH_ATTN3_HUB_AVAILABLE = True
-        except Exception as exc:
-            print(f"kernels-community/flash-attn3 unavailable ({exc}); using matched torch SDPA on both sides.")
-            _FLASH_ATTN3_HUB_AVAILABLE = False
-    return _FLASH_ATTN3_HUB_AVAILABLE
-
-
-def _omni_server_env() -> dict[str, str] | None:
-    if _flash_attn3_hub_available():
-        return None
-    return {"DIFFUSION_ATTENTION_BACKEND": "TORCH_SDPA"}
-
-
-def _set_reference_attention_backend(pipe: DiffusionPipeline) -> None:
-    if _flash_attn3_hub_available():
-        pipe.transformer.set_attention_backend("_flash_3_hub")
-    else:
-        pipe.transformer.set_attention_backend("native")
-
-
 def _model_name() -> str:
    return os.environ.get(MODEL_ENV_VAR, MODEL_ID)

@@ -115,7 +73,22 @@ def _local_files_only(model: str) -> bool:
    return Path(model).exists()


+_OMNI_FA3_HUB_BACKEND = "FLASH_ATTN_3_HUB"
+_OMNI_FA3_HUB_ENV = {"DIFFUSION_ATTENTION_BACKEND": _OMNI_FA3_HUB_BACKEND}
+
+
+def _assert_omni_fa3_hub_resolves() -> None:
+    from vllm_omni.diffusion.attention.backends.registry import DiffusionAttentionBackendEnum
+    from vllm_omni.platforms.cuda.platform import CudaOmniPlatform
+
+    path = CudaOmniPlatform.get_diffusion_attn_backend_cls(_OMNI_FA3_HUB_BACKEND, head_size=128)
+    expected = DiffusionAttentionBackendEnum.FLASH_ATTN_3_HUB.get_path()
+    assert path == expected, path
+    print(f"Resolved Omni attention backend: {_OMNI_FA3_HUB_BACKEND} -> {path}")
+
+
 def _run_vllm_omni_qwen_image(*, model: str, output_path: Path) -> Image.Image:
+    _assert_omni_fa3_hub_resolves()
    server_args = [
        "--num-gpus",
        "1",
@@ -124,8 +97,10 @@ def _run_vllm_omni_qwen_image(*, model: str, output_path: Path) -> Image.Image:
        "--init-timeout",
        "900",
        "--fa-deterministic",
+        "--diffusion-attention-backend",
+        _OMNI_FA3_HUB_BACKEND,
    ]
-    with OmniServer(model, server_args, use_omni=True, env_dict=_omni_server_env()) as omni_server:
+    with OmniServer(model, server_args, use_omni=True, env_dict=_OMNI_FA3_HUB_ENV) as omni_server:
        response = requests.post(
            f"http://{omni_server.host}:{omni_server.port}/v1/images/generations",
            json={
@@ -161,7 +136,7 @@ def _run_diffusers_qwen_image(*, model: str, output_path: Path) -> Image.Image:
            trust_remote_code=True,
            local_files_only=_local_files_only(model),
        ).to("cuda")
-        _set_reference_attention_backend(pipe)
+        pipe.transformer.set_attention_backend("_flash_3_hub")
        generator = torch.Generator(device="cuda").manual_seed(SEED)
        result = pipe(  # pyright: ignore[reportCallIssue]
            prompt=PROMPT,
@@ -186,8 +161,18 @@ def _run_diffusers_qwen_image(*, model: str, output_path: Path) -> Image.Image:


 def _run_vllm_omni_qwen_image_2512(*, model: str, output_path: Path) -> Image.Image:
-    server_args = ["--num-gpus", "1", "--stage-init-timeout", "300", "--init-timeout", "900"]
-    with OmniServer(model, server_args, use_omni=True, env_dict=_omni_server_env()) as omni_server:
+    _assert_omni_fa3_hub_resolves()
+    server_args = [
+        "--num-gpus",
+        "1",
+        "--stage-init-timeout",
+        "300",
+        "--init-timeout",
+        "900",
+        "--diffusion-attention-backend",
+        _OMNI_FA3_HUB_BACKEND,
+    ]
+    with OmniServer(model, server_args, use_omni=True, env_dict=_OMNI_FA3_HUB_ENV) as omni_server:
        response = requests.post(
            f"http://{omni_server.host}:{omni_server.port}/v1/images/generations",
            json={
@@ -223,7 +208,7 @@ def _run_diffusers_qwen_image_2512(*, model: str, output_path: Path) -> Image.Im
            trust_remote_code=True,
            local_files_only=_local_files_only(model),
        ).to("cuda")
-        _set_reference_attention_backend(pipe)
+        pipe.transformer.set_attention_backend("_flash_3_hub")
        generator = torch.Generator(device="cuda").manual_seed(SEED_2512)
        result = pipe(  # pyright: ignore[reportCallIssue]
            prompt=PROMPT_2512,
@@ -262,7 +247,7 @@ def test_qwen_image_matches_diffusers(accuracy_artifact_root: Path) -> None:
        diffusers_image=diffusers_output,
        width=WIDTH,
        height=HEIGHT,
-        ssim_threshold=SSIM_THRESHOLD if _flash_attn3_hub_available() else SSIM_THRESHOLD_SDPA_FALLBACK,
+        ssim_threshold=SSIM_THRESHOLD,
        psnr_threshold=PSNR_THRESHOLD,
    )

@@ -65,7 +65,12 @@ def _stage_config():
    # gains enough VRAM headroom for a no-eager run, add a companion test.
    return modify_stage_config(
        _CI_DEPLOY,
-        updates={"stages": {0: {"enforce_eager": True}, 1: {"enforce_eager": True}}},
+        updates={
+            "stages": {
+                0: {"enforce_eager": True, "moe_backend": "auto"},
+                1: {"enforce_eager": True},
+            }
+        },
    )


@@ -85,5 +85,7 @@ def test_text_only_nonstreaming_pcm_001(omni_server, openai_client) -> None:
        "voice": "default",
        "min_audio_bytes": _MIN_AUDIO_BYTES,
        "min_hnr_db": -2.0,
+        # dots.tts-soar emits native 48 kHz PCM, which has no rate header.
+        "expected_sample_rate": 48_000,
    }
    openai_client.send_audio_speech_request(request_config)
@@ -24,6 +24,9 @@ from tests.helpers.stage_config import get_deploy_config_path, modify_stage_conf
 os.environ["VLLM_WORKER_MULTIPROC_METHOD"] = "spawn"

 _MODEL = "openbmb/MiniCPM-o-4_5"
+_NATIVE_CHAT_TEMPLATE = (
+    Path(__file__).resolve().parents[3] / "vllm_omni/transformers_utils/chat_templates/minicpmo45_native.jinja"
+)
 _CI_DEPLOY = modify_stage_config(
    get_deploy_config_path("minicpmo_4_5.yaml"),
    updates={
@@ -63,7 +66,14 @@ test_params = [
            model=_MODEL,
            stage_config_path=_CI_DEPLOY,
            use_stage_cli=False,
-            server_args=["--trust-remote-code", "--async-chunk"],
+            server_args=[
+                "--trust-remote-code",
+                "--async-chunk",
+                "--chat-template",
+                str(_NATIVE_CHAT_TEMPLATE),
+                "--chat-template-content-format",
+                "openai",
+            ],
            env_dict={
                "VLLM_CONFIGURE_LOGGING": "1",
                "VLLM_LOGGING_CONFIG_PATH": str(_PROMPT_LOGGING_CONFIG_PATH),
@@ -265,11 +275,18 @@ def test_image_to_text_audio_001(omni_server, openai_client) -> None:
    Input Setting: stream=True
    """
    image_data_url = f"data:image/jpeg;base64,{generate_synthetic_image(24, 24)['base64']}"
-    messages = dummy_messages_from_mix_data(
-        system_prompt=get_system_prompt(),
-        image_data_url=image_data_url,
-        content_text=get_prompt("text_image"),
-    )
+    # Match the native image-chat layout explicitly; do not rely on the
+    # renderer moving media placeholders ahead of text.
+    messages = [
+        get_system_prompt(),
+        {
+            "role": "user",
+            "content": [
+                {"type": "image_url", "image_url": {"url": image_data_url}},
+                {"type": "text", "text": get_prompt("text_image")},
+            ],
+        },
+    ]

    request_config = {
        "model": omni_server.model,
@@ -926,10 +926,7 @@ def test_remote_replicas_use_distinct_init_group_keys():
        replica.launch_mode = "remote"
        replica.metadata.runtime_cfg = None

-    assert [runtime._replica_init_group_key(replica) for replica in plan.replicas] == [
-        "remote:1:0",
-        "remote:1:1",
-    ]
+    assert runtime._init_group_keys(plan.replicas) == ["remote:1:0", "remote:1:1"]


 def test_initialize_stages_cleans_up_successful_replicas_after_partial_multi_replica_failure(monkeypatch):
@@ -1596,73 +1593,58 @@ def test_model_path_resolver_is_generic_and_model_owned(tmp_path):
    assert "model_path_resolver" not in engine_args


-def test_build_stage0_input_processor_uses_omni_input_preprocessor(monkeypatch):
+def test_build_stage0_input_processor_uses_omni_renderer_subclass(monkeypatch):
+    from vllm.renderers import BaseRenderer
+
    import vllm_omni.engine.stage_init_utils as init_mod
+    from vllm_omni.inputs.preprocess import OmniRenderer, omni_renderer_cls

-    class DummyInputProcessor:
-        def __init__(self, vllm_config, renderer=None):
-            self.vllm_config = vllm_config
-            self.renderer = renderer or object()
-            self.input_preprocessor = None
+    class _Base(BaseRenderer):
+        def __init__(self, config, tokenizer):
+            self.config, self.tokenizer = config, tokenizer

-    class DummyOmniInputPreprocessor:
-        def __init__(self, vllm_config, renderer=None):
-            self.vllm_config = vllm_config
-            self.renderer = renderer
+        def render_messages(self, messages, params):  # pragma: no cover - abstract stub
+            raise NotImplementedError

-    monkeypatch.setattr(init_mod, "InputProcessor", DummyInputProcessor)
-    monkeypatch.setattr(init_mod, "OmniInputPreprocessor", DummyOmniInputPreprocessor)
-
-    input_processor = build_stage0_input_processor(
-        types.SimpleNamespace(model_config=types.SimpleNamespace(try_get_generation_config=lambda: {}))
-    )
-
-    assert isinstance(input_processor.input_preprocessor, DummyOmniInputPreprocessor)
-    assert input_processor.input_preprocessor.renderer is input_processor.renderer
-
-
-def test_build_stage0_input_processor_does_not_resolve_tokenizer_when_skipped(
-    monkeypatch,
-):
-    import vllm_omni.engine.stage_init_utils as init_mod
-
-    token_only_renderer = object()
+    config = types.SimpleNamespace(model_config=types.SimpleNamespace(try_get_generation_config=lambda: {}))
+    built = omni_renderer_cls(_Base)(config, "tok")
    seen = {}

    class DummyInputProcessor:
        def __init__(self, vllm_config, renderer=None):
-            if renderer is None:
-                raise AssertionError("skip_tokenizer_init must supply a renderer")
            seen["renderer"] = renderer
            self.renderer = renderer
-            self.input_preprocessor = None
-
-    class DummyOmniInputPreprocessor:
-        def __init__(self, vllm_config, renderer=None):
-            seen["preprocessor_renderer"] = renderer

    monkeypatch.setattr(init_mod, "InputProcessor", DummyInputProcessor)
-    monkeypatch.setattr(
-        init_mod,
-        "_build_token_only_renderer",
-        lambda _config: token_only_renderer,
-    )
-    monkeypatch.setattr(
-        init_mod,
-        "OmniInputPreprocessor",
-        DummyOmniInputPreprocessor,
-    )
+    monkeypatch.setattr(init_mod, "build_omni_renderer", lambda cfg: built if cfg is config else None)
+    processor = build_stage0_input_processor(config)
+    assert seen["renderer"] is built
+    assert isinstance(processor.renderer, OmniRenderer)
+    assert isinstance(processor.renderer, _Base)
+    assert not hasattr(processor, "input_preprocessor")

+
+def test_build_stage0_input_processor_does_not_resolve_tokenizer_when_skipped(monkeypatch):
+    import vllm_omni.engine.stage_init_utils as init_mod
+
+    original = object()
+
+    class DummyInputProcessor:
+        def __init__(self, vllm_config, renderer=None):
+            assert renderer is original
+            self.renderer = renderer
+
+    def _must_not_resolve(_cfg):
+        raise AssertionError("tokenizer must not be resolved when skip_tokenizer_init=True")
+
+    monkeypatch.setattr(init_mod, "InputProcessor", DummyInputProcessor)
+    monkeypatch.setattr(init_mod, "_build_token_only_renderer", lambda _: original)
+    monkeypatch.setattr(init_mod, "build_omni_renderer", _must_not_resolve)
    config = types.SimpleNamespace(
-        model_config=types.SimpleNamespace(
-            skip_tokenizer_init=True,
-            try_get_generation_config=lambda: {},
-        )
+        model_config=types.SimpleNamespace(skip_tokenizer_init=True, try_get_generation_config=lambda: {})
    )
-    build_stage0_input_processor(config)
-
-    assert seen["renderer"] is token_only_renderer
-    assert seen["preprocessor_renderer"] is token_only_renderer
+    processor = build_stage0_input_processor(config)
+    assert processor.renderer is original


 def test_inject_kv_stage_info_infers_sender_tp_topology():
@@ -1,7 +1,10 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
 """Unit tests for parallel stage initialization (CPU-only).

 Covers the pure/lock logic added for ``parallel_stage_init``:
-  * init-group device keying (resolved physical set; parallel unique keys),
+  * init-group device keying (overlap components; parallel unique keys),
  * admission control arithmetic + graph reserve,
  * SH/EX device phase locks (real flock on a temp dir),
  * the phase-locked executor wrapper ordering,
@@ -18,6 +21,7 @@ import inspect
 import os
 import subprocess
 import threading
+import time
 import types
 from pathlib import Path

@@ -37,7 +41,9 @@ from vllm_omni.engine.stage_init_utils import (
    LogicalStageInitPlan,
    ReplicaInitPlan,
    device_init_lock_path,
+    device_overlap_group_keys,
    open_device_lock_file,
+    parse_physical_device_ids,
 )
 from vllm_omni.engine.stage_phase_lock import (
    DeviceLockTimeoutError,
@@ -46,6 +52,7 @@ from vllm_omni.engine.stage_phase_lock import (
    wrap_executor_with_phase_locks,
 )
 from vllm_omni.engine.stage_runtime import StageRuntime
+from vllm_omni.platforms import current_omni_platform

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

@@ -104,49 +111,114 @@ def _fake_vllm_config(util: float, *, model="m", dtype="bf16", cudagraph_mode="F
 # --------------------------------------------------------------------------- #
 # A1: init-group device key
 # --------------------------------------------------------------------------- #
-def test_init_group_key_resolves_physical_devices_serial():
+@pytest.mark.parametrize(
+    ("devices", "expected"),
+    [
+        ("0", {0}),
+        ("1,0,1", {0, 1}),
+        (" 2, 3 ", {2, 3}),
+        (None, None),
+        ("", None),
+        ("GPU-uuid", None),
+    ],
+)
+def test_parse_physical_device_ids(devices, expected):
+    assert parse_physical_device_ids(devices) == expected
+
+
+@pytest.mark.parametrize(
+    ("device_sets", "expected"),
+    [
+        # equal / disjoint
+        ([{0}, {1}, {0}], ["device-group:0", "device-group:1", "device-group:0"]),
+        # unequal overlap, canonical order
+        ([{0, 1}, {0}, {2}], ["device-group:0,1", "device-group:0,1", "device-group:2"]),
+        # transitive: {0,1} ~ {1,2} ~ {2}
+        ([{0, 1}, {1, 2}, {2}], ["device-group:0,1,2"] * 3),
+        # unresolved may touch any GPU -> everything serial
+        ([{0}, None, {1}], ["device-group:*"] * 3),
+        ([], []),
+    ],
+)
+def test_device_overlap_group_keys(device_sets, expected):
+    sets = [None if s is None else frozenset(s) for s in device_sets]
+    assert device_overlap_group_keys(sets) == expected
+
+
+def test_init_group_keys_serial_groups_by_device_overlap():
    runtime = _runtime(parallel_stage_init=False)
-    runtime._init_visible_devices_baseline = None  # no CUDA_VISIBLE_DEVICES mapping
-    r0 = _llm_replica(0, 0, "0")
-    r1 = _llm_replica(1, 0, "1")
-    r0b = _llm_replica(2, 0, "0")

-    k0 = runtime._replica_init_group_key(r0)
-    k1 = runtime._replica_init_group_key(r1)
-    k0b = runtime._replica_init_group_key(r0b)
+    keys = runtime._init_group_keys(
+        [_llm_replica(0, 0, "0,1"), _llm_replica(1, 0, "0"), _llm_replica(2, 0, "2"), _llm_replica(3, 0, "1,0")]
+    )

-    assert k0 == "device:0"
-    assert k1 == "device:1"
-    # Different physical devices -> different groups (parallel).
-    assert k0 != k1
-    # Same physical device -> same group (serialized under LOCK_EX).
-    assert k0 == k0b
+    assert keys == ["device-group:0,1", "device-group:0,1", "device-group:2", "device-group:0,1"]


-def test_init_group_key_parallel_is_unique_per_replica():
-    runtime = _runtime(parallel_stage_init=True)
-    runtime._init_visible_devices_baseline = None
-    # Both on the SAME device, but parallel mode must give distinct keys so they
-    # run concurrently, coordinated by the child SH/EX locks + admission.
-    keys = [
-        runtime._replica_init_group_key(_llm_replica(0, 0, "0")),
-        runtime._replica_init_group_key(_llm_replica(0, 1, "0")),
-        runtime._replica_init_group_key(_llm_replica(1, 0, "0")),
-    ]
-    assert keys == ["parallel:0:0", "parallel:0:1", "parallel:1:0"]
-    assert len(set(keys)) == 3
-
-
-def test_init_group_key_remote_and_diffusion_unchanged():
+def test_init_group_keys_fixed_for_parallel_remote_and_diffusion():
    runtime = _runtime(parallel_stage_init=True)
    remote = _llm_replica(1, 0, "0")
    remote.launch_mode = "remote"
    remote.metadata.runtime_cfg = None
-    assert runtime._replica_init_group_key(remote) == "remote:1:0"
-
    diffusion = _llm_replica(2, 0, "0")
    diffusion.metadata.stage_type = "diffusion"
-    assert runtime._replica_init_group_key(diffusion) == "inline:diffusion"
+
+    # Same device everywhere; parallel mode still keys per replica (child SH/EX locks coordinate).
+    keys = runtime._init_group_keys([_llm_replica(0, 0, "0"), _llm_replica(0, 1, "0"), remote, diffusion])
+
+    assert keys == ["parallel:0:0", "parallel:0:1", "remote:1:0", "inline:diffusion"]
+
+
+def test_unresolved_devices_collapse_serial_groups_and_warn(caplog):
+    runtime = _runtime(parallel_stage_init=False)
+    with caplog.at_level("WARNING"):
+        keys = runtime._init_group_keys([_llm_replica(0, 0, "0"), _llm_replica(1, 0, "GPU-uuid")])
+
+    assert keys == ["device-group:*", "device-group:*"]
+    assert "Stage-1 replica 0" in caplog.text
+    assert "GPU-uuid" in caplog.text
+
+
+def _stage_plans(*devices: str) -> list[LogicalStageInitPlan]:
+    return [
+        LogicalStageInitPlan(stage_idx=i, stage_id=i, replicas=[_llm_replica(i, 0, dev)])
+        for i, dev in enumerate(devices)
+    ]
+
+
+def test_overlapping_device_stages_initialize_sequentially(monkeypatch):
+    monkeypatch.delenv(current_omni_platform.device_control_env_var, raising=False)
+    runtime = _runtime(parallel_stage_init=False)
+    in_flight, peak = 0, 0
+    guard = threading.Lock()
+
+    def _initialize_replica(plan, _timeout):
+        nonlocal in_flight, peak
+        with guard:
+            in_flight += 1
+            peak = max(peak, in_flight)
+        time.sleep(0.05)
+        with guard:
+            in_flight -= 1
+        return object()
+
+    monkeypatch.setattr(runtime, "_initialize_replica", _initialize_replica)
+    runtime._initialize_stage_replicas(_stage_plans("0,1", "0"), stage_init_timeout=5)
+
+    assert peak == 1
+
+
+def test_disjoint_device_stages_initialize_in_parallel(monkeypatch):
+    monkeypatch.delenv(current_omni_platform.device_control_env_var, raising=False)
+    runtime = _runtime(parallel_stage_init=False)
+    barrier = threading.Barrier(2, timeout=2)  # breaks unless both inits are in flight together
+
+    def _initialize_replica(_plan, _timeout):
+        barrier.wait()
+        return object()
+
+    monkeypatch.setattr(runtime, "_initialize_replica", _initialize_replica)
+    runtime._initialize_stage_replicas(_stage_plans("0", "1"), stage_init_timeout=5)


 # --------------------------------------------------------------------------- #
@@ -11,6 +11,7 @@ from pathlib import Path

 import pytest
 from pydantic.fields import FieldInfo
+from transformers import Qwen3OmniMoeConfig
 from vllm.config import CacheConfig as VllmCacheConfig
 from vllm.config import CompilationConfig as VllmCompilationConfig
 from vllm.config import LoadConfig as VllmLoadConfig
@@ -19,6 +20,7 @@ from vllm.config import ProfilerConfig as VllmProfilerConfig
 from vllm.config import SchedulerConfig as VllmSchedulerConfig
 from vllm.engine.arg_utils import EngineArgs

+from tests.helpers.stage_config import get_deploy_config_path, modify_stage_config
 from vllm_omni.config.omni_config import (
    _LLM_STAGE_ENGINE_FIELDS,
    OmniStageCacheConfig,
@@ -155,6 +157,44 @@ def test_qwen3_omni_preserves_explicit_moe_backend():
    assert engine_args["moe_backend"] == "flashinfer"


+def test_qwen3_omni_preserves_explicit_auto_moe_backend():
+    engine_args = {
+        "model_arch": "Qwen3OmniMoeForConditionalGeneration",
+        "moe_backend": "auto",
+    }
+
+    stage_init_utils._maybe_set_qwen3_omni_moe_backend(engine_args)
+
+    assert engine_args["moe_backend"] == "auto"
+
+
+def test_qwen3_omni_nvfp4_ci_config_preserves_auto_moe_backend(tmp_path):
+    """The NVFP4 CI deploy opts stage 0 into vLLM's auto backend selection."""
+    hf_config = Qwen3OmniMoeConfig(enable_audio_output=True)
+    pipeline = resolve_pipeline_config("qwen3_omni_moe", hf_config)
+    assert pipeline is not None
+    nvfp4_deploy = modify_stage_config(
+        get_deploy_config_path("ci/qwen3_omni_moe.yaml"),
+        updates={"stages": {0: {"moe_backend": "auto"}}},
+    )
+    deploy = load_deploy_config(nvfp4_deploy)
+    legacy_stages = merge_pipeline_deploy(pipeline, copy.deepcopy(deploy))
+    omni_config = VllmOmniConfig.from_pipeline_config(
+        pipeline,
+        user_deploy_config=copy.deepcopy(deploy),
+        cli_overrides={"model": str(tmp_path)},
+    )
+
+    for stage_id, expected_backend in ((0, "auto"), (1, "triton")):
+        legacy_stage = next(stage for stage in legacy_stages if stage.stage_id == stage_id)
+        typed_stage = omni_config.stage_by_id(stage_id)
+        for engine_args in (
+            build_engine_args_dict(legacy_stage.to_omegaconf(), str(tmp_path)),
+            build_engine_args_dict_from_omni_stage_config(typed_stage, str(tmp_path)),
+        ):
+            assert engine_args["moe_backend"] == expected_backend
+
+
 def _engine_arg_inputs(tmp_path: Path) -> tuple[PipelineConfig, DeployConfig, str]:
    parent_model = tmp_path / "parent-model"
    stage_model = tmp_path / "stage-model"
@@ -0,0 +1,149 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+"""Check replica startup ordering without model-loading timing assumptions."""
+
+import contextlib
+import fcntl
+import os
+import threading
+from collections.abc import Iterator
+from pathlib import Path
+
+import pytest
+from pytest_mock import MockerFixture
+from vllm.config import VllmConfig
+from vllm.sampling_params import SamplingParams
+from vllm.v1.engine.utils import CoreEngineProcManager, EngineZmqAddresses
+
+from vllm_omni.config.stage_config import StageConfig
+from vllm_omni.engine import stage_init_utils, stage_runtime
+from vllm_omni.engine.stage_engine_startup import StageReplicaResources
+from vllm_omni.engine.stage_init_utils import LogicalStageInitPlan, ReplicaInitPlan, StageMetadata
+
+pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
+
+
+@pytest.fixture
+def launch_runtime(
+    monkeypatch: pytest.MonkeyPatch, tmp_path: Path, mocker: MockerFixture
+) -> Iterator[stage_runtime.StageRuntime]:
+    runtime = stage_runtime.StageRuntime(
+        stage_configs=[],
+        model="dummy-model",
+        config_path="dummy-config",
+        stage_init_timeout=5,
+        async_chunk=False,
+    )
+    # These are virtual device IDs: only the real per-device file locking and
+    # scheduling run. No model, CUDA context, or engine process is created.
+    monkeypatch.setattr(runtime, "_resolve_replica_physical_devices", lambda _stage, cfg: cfg["devices"])
+    monkeypatch.setattr(
+        stage_init_utils, "device_init_lock_path", lambda device: str(tmp_path / f"device-{device}.lock")
+    )
+    mocker.patch.object(
+        stage_runtime.StageEngineCoreClientBase,
+        "make_async_mp_client",
+        return_value=mocker.Mock(spec=stage_runtime.StageClient),
+    )
+    monkeypatch.setenv("VLLM_OMNI_TEST_LAUNCH_ENV", "parent")
+    yield runtime
+    if runtime._stage_init_executor is not None:
+        runtime._stage_init_executor.shutdown(wait=True)
+
+
+def _plan(stage_id: int, device: str) -> LogicalStageInitPlan:
+    metadata = StageMetadata(
+        stage_id=stage_id,
+        stage_type="llm",
+        engine_output_type="text",
+        is_comprehension=False,
+        requires_multimodal_data=False,
+        engine_input_source=[],
+        final_output=False,
+        final_output_type=None,
+        default_sampling_params=SamplingParams(),
+        custom_process_input_func=None,
+        model_stage="dummy",
+        runtime_cfg={"devices": device, "env": {"VLLM_OMNI_TEST_LAUNCH_ENV": str(stage_id)}},
+    )
+    replica = ReplicaInitPlan(
+        replica_id=0,
+        num_replicas=1,
+        launch_mode="local",
+        stage_cfg=StageConfig(stage_id=stage_id, model_stage="dummy", yaml_runtime=metadata.runtime_cfg),
+        metadata=metadata,
+        stage_connector_spec={},
+        omni_kv_connector=(None, None, None),
+        stage_vllm_config=VllmConfig(),
+        executor_class=object,
+        engine_args_dict={},
+    )
+    return LogicalStageInitPlan(stage_idx=stage_id, stage_id=stage_id, replicas=[replica])
+
+
+def _resources(mocker: MockerFixture) -> StageReplicaResources:
+    return StageReplicaResources(
+        manager=mocker.Mock(spec=CoreEngineProcManager),
+        addresses=EngineZmqAddresses(inputs=["unused-input"], outputs=["unused-output"]),
+    )
+
+
+def test_distinct_devices_spawn_before_either_is_ready(
+    launch_runtime: stage_runtime.StageRuntime, monkeypatch: pytest.MonkeyPatch, mocker: MockerFixture
+) -> None:
+    spawned = [threading.Event(), threading.Event()]
+
+    @contextlib.contextmanager
+    def launch(*, stage_id: int, stage_visible_devices: str, **_kwargs: object) -> Iterator[StageReplicaResources]:
+        assert os.environ["VLLM_OMNI_TEST_LAUNCH_ENV"] == str(stage_id)
+        spawned[stage_id].set()
+        yield _resources(mocker)
+        # The device remains exclusively locked throughout loading/profiling,
+        # even though another device's replica must be able to start.
+        with open(stage_init_utils.device_init_lock_path(int(stage_visible_devices)), "r+") as lock:
+            with pytest.raises(BlockingIOError):
+                fcntl.flock(lock.fileno(), fcntl.LOCK_EX | fcntl.LOCK_NB)
+        assert spawned[1 - stage_id].wait(5), "A different device's spawn was blocked by the READY wait"
+
+    monkeypatch.setattr(stage_runtime, "launch_stage_replica", launch)
+    clients = launch_runtime._initialize_stage_replicas([_plan(0, "0"), _plan(1, "1")], 5)
+    assert all(clients[stage][0] is not None for stage in (0, 1))
+    assert os.environ["VLLM_OMNI_TEST_LAUNCH_ENV"] == "parent"
+
+
+def test_same_device_waits_for_ready_before_next_spawn(
+    launch_runtime: stage_runtime.StageRuntime, monkeypatch: pytest.MonkeyPatch, mocker: MockerFixture
+) -> None:
+    events: list[tuple[int, str]] = []
+
+    @contextlib.contextmanager
+    def launch(*, stage_id: int, **_kwargs: object) -> Iterator[StageReplicaResources]:
+        events.append((stage_id, "spawn"))
+        yield _resources(mocker)
+        events.append((stage_id, "ready"))
+
+    monkeypatch.setattr(stage_runtime, "launch_stage_replica", launch)
+    launch_runtime._initialize_stage_replicas([_plan(0, "0"), _plan(1, "0")], 5)
+    assert events == [(0, "spawn"), (0, "ready"), (1, "spawn"), (1, "ready")]
+    assert os.environ["VLLM_OMNI_TEST_LAUNCH_ENV"] == "parent"
+
+
+def test_ready_failure_cleans_up_resources_and_releases_device(
+    launch_runtime: stage_runtime.StageRuntime, monkeypatch: pytest.MonkeyPatch, mocker: MockerFixture
+) -> None:
+    manager = mocker.Mock(spec=CoreEngineProcManager)
+    resources = StageReplicaResources(manager=manager, addresses=EngineZmqAddresses(inputs=[], outputs=[]))
+
+    @contextlib.contextmanager
+    def launch(**_kwargs: object) -> Iterator[StageReplicaResources]:
+        yield resources
+        raise RuntimeError("READY handshake failed")
+
+    monkeypatch.setattr(stage_runtime, "launch_stage_replica", launch)
+    with pytest.raises(RuntimeError, match="READY handshake failed"):
+        launch_runtime._initialize_local_llm_replica(_plan(0, "0").replicas[0], 5)
+
+    manager.shutdown.assert_called_once_with()
+    assert os.environ["VLLM_OMNI_TEST_LAUNCH_ENV"] == "parent"
+    with open(stage_init_utils.device_init_lock_path(0), "r+") as lock:
+        fcntl.flock(lock.fileno(), fcntl.LOCK_EX | fcntl.LOCK_NB)
@@ -1,3 +1,6 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+
 import asyncio
 from unittest.mock import AsyncMock

@@ -12,7 +15,7 @@ from vllm.entrypoints.openai.chat_completion.protocol import (
    ChatCompletionResponseChoice,
    ChatMessage,
 )
-from vllm.entrypoints.openai.engine.protocol import ErrorResponse, UsageInfo
+from vllm.entrypoints.serve.engine.protocol import ErrorResponse, UsageInfo

 from vllm_omni.entrypoints.openai.batch_serving import OmniOpenAIServingChatBatch

@@ -248,7 +248,7 @@ class TestResolveAudioFormat:
        assert result == "pcm"

    def test_invalid_format_returns_error(self, serving_chat):
-        from vllm.entrypoints.openai.engine.protocol import ErrorResponse
+        from vllm.entrypoints.serve.engine.protocol import ErrorResponse

        request = self._make_request({"format": "aac", "voice": "alloy"})
        result = serving_chat._resolve_audio_format(request)
@@ -256,7 +256,7 @@ class TestResolveAudioFormat:
        assert "aac" in result.error.message

    def test_all_supported_formats_accepted(self, serving_chat):
-        from vllm.entrypoints.openai.engine.protocol import ErrorResponse
+        from vllm.entrypoints.serve.engine.protocol import ErrorResponse

        for fmt in SUPPORTED_CHAT_AUDIO_FORMATS:
            request = self._make_request({"format": fmt, "voice": "alloy"})
@@ -1,4 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Regression tests for output modality validation in chat completions.

 Covers:
@@ -13,7 +14,7 @@ from unittest.mock import AsyncMock, MagicMock, patch

 import pytest
 from vllm.entrypoints.openai.chat_completion.protocol import ChatCompletionRequest
-from vllm.entrypoints.openai.engine.protocol import ErrorResponse
+from vllm.entrypoints.serve.engine.protocol import ErrorResponse

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

@@ -1,4 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Unit tests for OmniChatCompletionResponse/StreamResponse metrics field."""

 from types import SimpleNamespace
@@ -19,7 +20,7 @@ pytestmark = [pytest.mark.core_model, pytest.mark.cpu]

 def test_omni_chat_completion_response_metrics():
    """Test OmniChatCompletionResponse metrics field works correctly."""
-    from vllm.entrypoints.openai.engine.protocol import UsageInfo
+    from vllm.entrypoints.serve.engine.protocol import UsageInfo

    from vllm_omni.entrypoints.openai.protocol.chat_completion import (
        OmniChatCompletionResponse,
@@ -25,7 +25,7 @@ from fastapi.responses import JSONResponse
 from fastapi.testclient import TestClient
 from pydantic import ValidationError
 from pytest_mock import MockerFixture
-from vllm.entrypoints.openai.engine.protocol import ErrorInfo, ErrorResponse
+from vllm.entrypoints.serve.engine.protocol import ErrorInfo, ErrorResponse

 from vllm_omni.diffusion.request import OmniDiffusionRequest
 from vllm_omni.diffusion.sched.request_scheduler import RequestScheduler
@@ -1,5 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Runtime behavior tests for VideoStreamHandler.

 Every test creates a mock WebSocket, drives handle_session through a
@@ -281,7 +281,7 @@ class TestHandleQueryErrors:
    @pytest.mark.asyncio
    async def test_query_returns_error_response(self):
        """create_chat_completion returns ErrorResponse → error forwarded."""
-        from vllm.entrypoints.openai.engine.protocol import ErrorInfo, ErrorResponse
+        from vllm.entrypoints.serve.engine.protocol import ErrorInfo, ErrorResponse

        err = ErrorResponse(
            error=ErrorInfo(
@@ -1,4 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Tests for multi-modal streaming finish_reason behavior (commit 44c799bc).

 Verifies that the /v1/chat/completions streaming endpoint emits exactly one
@@ -411,7 +412,7 @@ async def test_audio_chunk_without_waveform_keeps_stream_alive():
@pytest.mark.asyncio
 async def test_audio_choice_error_response_is_not_iterated_as_choices():
    """Any ErrorResponse from the audio path is skipped, not iterated."""
-    from vllm.entrypoints.openai.engine.protocol import ErrorResponse
+    from vllm.entrypoints.serve.engine.protocol import ErrorResponse

    serving_chat = build_serving_chat()
    request = make_request(modalities=["text", "audio"])
@@ -0,0 +1,36 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
+"""Shared access to the production ``KVCacheTensor`` layout derivation.
+
+Tests deliberately go through the same code production uses
+(``vllm_omni.diffusion.diffusion_kv.layout``) rather than recomputing strides,
+so a regression in that derivation fails the suite instead of being masked by a
+parallel test-only copy.
+"""
+
+from types import SimpleNamespace
+
+from vllm.config import CacheConfig
+from vllm.v1.kv_cache_layout import KVCacheLayout
+
+from vllm_omni.diffusion.diffusion_kv.layout import (
+    build_kv_cache_tensor,
+    resolve_diffusion_kv_cache_layout,
+)
+
+__all__ = ["build_kv_cache_tensor", "layout_for_backend"]
+
+
+def layout_for_backend(attn_backend) -> KVCacheLayout:
+    """The physical layout ``attn_backend``'s block-stride declaration selects.
+
+    Before 0.29 a backend's ``indexes_kv_by_block_stride()`` was copied onto
+    ``AttentionSpec``; asserting on the resolved layout is the surviving way to
+    check that declaration still reaches the cache.
+    """
+
+    vllm_config = SimpleNamespace(cache_config=CacheConfig())
+    return resolve_diffusion_kv_cache_layout(
+        vllm_config,
+        indexes_kv_by_block_stride=attn_backend.indexes_kv_by_block_stride(),
+    )
@@ -968,7 +968,6 @@ def iter_omni_server(
                raise ValueError("omni_server with use_stage_cli=True requires use_omni=True")
            if stage_config_path is None:
                raise ValueError("omni_server with use_stage_cli=True requires a stage_config_path")
-            server_args += ["--deploy-config", stage_config_path]

            with OmniServerStageCli(
                model,
@@ -1,4 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
 """Shared scaffolding for OmniOpenAIServingChat unit tests.

 Provides factories for a minimal serving-chat instance, chat requests, and
@@ -8,11 +9,11 @@ OmniRequestOutput objects, plus SSE stream collection/parsing utilities.
 import json
 from unittest.mock import MagicMock

+from vllm.entrypoints.generate.base.protocol import DeltaMessage
 from vllm.entrypoints.openai.chat_completion.protocol import (
    ChatCompletionRequest,
    ChatCompletionResponseStreamChoice,
 )
-from vllm.entrypoints.openai.engine.protocol import DeltaMessage
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
 from vllm.outputs import CompletionOutput, RequestOutput
Author	SHA1	Message	Date
Hongsheng Liu	fd071ae081	Merge branch 'main' into codex-diffusion-shutdown-escalation	2026-09-10 13:27:51 +08:00
Zhou TaichangandClaude Opus 5	aff7d64948	[Rebase] Rebase to vLLM 0.29.0 (#7230 ) Signed-off-by: tzhouam <tzhouam@connect.ust.hk> Signed-off-by: Zhou Taichang <tzhouam@connect.ust.hk> Co-authored-by: Claude Opus 5 <noreply@anthropic.com>	2026-09-10 12:08:28 +08:00
Sy03	333f14e97e	[Bugfix][MiniCPM-o] Align structured chat content with native omni rendering (#7344 ) Signed-off-by: Sy03 <1370724210@qq.com>	2026-09-10 11:38:11 +08:00
Mu GuanLin	4b8dcac681	[Kernel] Enable LTX DiffVAE fusions on SM100 and SM103 (#7350 ) Signed-off-by: mglyn <1203789601@qq.com>	2026-09-10 03:07:35 +00:00
Weiming Liaoandwangyu	35a83a0885	[NPU][CI] Add A5 and 310P CI support (#6875 ) Signed-off-by: Weiming Liao <liaowm5@gmail.com> Co-authored-by: wangyu <53896905+yenuo26@users.noreply.github.com>	2026-09-10 02:16:14 +00:00
andyluo7	3d71a2b3ad	[CI/Build] Stabilize LTX2 vocoder autocast test on ROCm (#7336 ) Signed-off-by: andyluo7 <andy.luo@amd.com>	2026-09-10 10:09:39 +08:00
Gao Hanandamy-why-3459	6e63ea875e	[Bugfix] Allow LLM replicas on different GPUs to initialize concurrently (#7292 ) Signed-off-by: Gao Han <hgaoaf@connect.ust.hk> Co-authored-by: amy-why-3459 <wuhaiyan17@huawei.com>	2026-09-10 09:54:20 +08:00
Hongsheng Liuandhsliu_ustc	0637454301	[Skills] Add perf-evidence isolation, stage-attribution, and realtime-contract requirements (#6820 ) Signed-off-by: hsliu_ustc <hsliu_ustc@noreply.gitcode.com> Co-authored-by: hsliu_ustc <hsliu_ustc@noreply.gitcode.com>	2026-09-10 08:05:47 +08:00
Guangjian Dong	faad918588	[Doc] Document num_outputs_per_prompt for /v1/videos (#7341 ) Signed-off-by: Guangjian <hiro20833@gmail.com>	2026-09-10 08:00:13 +08:00
Joshna-Medisetty	51cf142bc1	[Bugfix][XPU] Restore N-D output shape for W8A16 FP8 linear (#7301 ) Signed-off-by: Joshna Medisetty <joshna.medisetty@intel.com> Signed-off-by: Joshna-Medisetty <joshna.medisetty@intel.com>	2026-09-09 17:11:31 -05:00
Sy03	b02862beee	[Model] Optimize MOSS-TTS Local batched execution and streaming codec (#7202 ) Signed-off-by: Sy03 <1370724210@qq.com>	2026-09-09 13:02:12 -07:00
bojiang-li	504e5a6a6e	[Doc] Add dual DGX Spark MiniMax-H3 results (#7343 ) Signed-off-by: bojiang-li <327132355+bojiang-li@users.noreply.github.com>	2026-09-09 18:48:41 +00:00
hyw	6fb7b0a05e	[Model] Skip unused frame packing in Wan2.2 S2V (#7155 ) Signed-off-by: hyw <yuweih205@gmail.com>	2026-09-09 14:34:46 +00:00
andyluo7	b9932aa17b	[CI/Build][ROCm] Normalize SenseNova paged-decode hardware markers (#6935 ) Signed-off-by: andyluo7 <andy.luo@amd.com>	2026-09-09 07:32:38 -07:00
shiyichuanandmershi	06f67695f1	[CI/Build] Avoid duplicate stage CLI deploy config (#7007 ) Signed-off-by: mershi <mershi@tencent.com> Co-authored-by: mershi <mershi@tencent.com>	2026-09-09 22:05:44 +08:00
andyluo7andHongsheng Liu	c94cf4ddde	[CI][ROCm] Match CUDA/NPU L2/L3 label routing (#6966 ) Signed-off-by: andyluo7 <andy.luo@amd.com> Co-authored-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-09-09 13:47:15 +00:00
Shaun Walsh Claude Opus 4.6 Nick Cao	0e602509aa	[BugFix] Add field validators for /v1/audio/generate request (#4741 ) Signed-off-by: Shaun Walsh <shaunwalsh24@gmail.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: Nick Cao <ncao@redhat.com>	2026-09-09 13:38:42 +00:00
kunkun	2f1845fa4e	fix: reserve Qwen3-Omni NVFP4 backend fix (#7200 ) Signed-off-by: kunkunblueberry <1833921874@qq.com>	2026-09-09 20:35:16 +08:00
Zheng Wengangandamy-why-3459	7c1b1e4189	[Bugfix][Engine] Group overlapping device stages into one sequential init component (#7328 ) Signed-off-by: ZhengWG <zwg0606@gmail.com> Co-authored-by: amy-why-3459 <wuhaiyan17@huawei.com>	2026-09-09 20:25:41 +08:00
zyz111222	6e79d582e3	[Bugfix][Ascend] fix npu 310/a5 bugs (#6685 ) Signed-off-by: zouyizhou <zouyizhou@huawei.com>	2026-09-09 19:03:09 +08:00
NumberWan	038c9b948f	[Bugfix] Resolve #6931 hub FA3 on torch 2.13 via kernels 0.16.1 (#7185 ) Signed-off-by: NumberWan <wantszkin2003@gmail.com>	2026-09-09 10:20:34 +00:00
Bo Li	cdd4638b17	[Feature][Diffusion] Expose detailed pipeline timings (#6822 ) Signed-off-by: Bo Li <22713281+bobboli@users.noreply.github.com>	2026-09-09 09:58:42 +00:00
Anjie HouandClaude Sonnet 5	ecdd762ff5	[2/N] Add a minimal temporal chunk callback for MiniMax-H3 (#7017 ) Signed-off-by: specture724 <specture724@gmail.com> Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com>	2026-09-09 16:40:32 +08:00
Mu GuanLin	ffa991b24f	[Perf][Model] Add optimized LTX-2.5 DiffVAE operators (#7308 ) Signed-off-by: mglyn <1203789601@qq.com>	2026-09-09 08:31:16 +00:00
chi030303	c66df06c0f	[Doc] Qwen3-TTS: add 0.6B on 1x A100 40GB (#7289 ) Signed-off-by: chi030303 <106855944+chi030303@users.noreply.github.com>	2026-09-09 07:56:14 +00:00
Sparks	7712701578	[Doc] [Recipe] Add FLUX.1-schnell recipe for RTX 5090 32GB (#7299 ) Signed-off-by: Sparks-M <41097544+Sparks-M@users.noreply.github.com>	2026-09-09 07:54:46 +00:00
Mu GuanLin	ab56148505	[Bugfix] Make LTX vocoder decoding deterministic (#7231 ) Signed-off-by: mglyn <1203789601@qq.com>	2026-09-09 06:37:01 +00:00
NATURE	5b002727b0	[Core] Split Omni connector model runner mixin (#6903 ) Signed-off-by: natureofnature <wzliu@connect.hku.hk>	2026-09-09 13:49:43 +08:00
andyluo7	e0d0879794	[Bugfix][Magi2] Allow import without an active Triton driver (#7239 ) Signed-off-by: andyluo7 <andy.luo@amd.com>	2026-09-08 22:48:08 -07:00
Hongsheng Liu	c79c841e00	[Bugfix] Escalate diffusion worker shutdown and retain survivors Signed-off-by: Hongsheng Liu <liuhongsheng4@huawei.com>	2026-09-06 15:56:17 +08:00