diff --git a/.buildkite/check-torch-abi.py b/.buildkite/check-torch-abi.py
new file mode 100644
index 00000000000..cb580fb56f1
--- /dev/null
+++ b/.buildkite/check-torch-abi.py
@@ -0,0 +1,103 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""Audit vLLM compiled libraries for PyTorch stable ABI compliance."""
+
+import fnmatch
+import sys
+from pathlib import Path
+
+from torch_abi_audit import inspect_package
+from torch_abi_audit.report import ExtensionReport, PackageReport
+
+# Temporary allowlist of extensions not yet on the stable ABI.
+# Shrink and remove over time.
+ALLOWED_UNSTABLE_LIBRARIES: tuple[str, ...] = (
+    "_flashkda_C.abi3.so",
+    "vllm_flash_attn/_vllm_fa2_C.abi3.so",
+    "vllm_flash_attn/_vllm_fa3_C.abi3.so",
+    "third_party/deep_gemm/_C*.so",
+)
+
+
+def _relative_path(lib: ExtensionReport, package_root: Path) -> str:
+    try:
+        return lib.path.relative_to(package_root).as_posix()
+    except ValueError:
+        return lib.path.name
+
+
+def _is_torch_unstable(lib: ExtensionReport) -> bool:
+    return lib.error is None and lib.torch.uses_torch and not lib.torch.stable
+
+
+def _matches_allowlist(rel_path: str, patterns: tuple[str, ...]) -> bool:
+    return any(fnmatch.fnmatch(rel_path, pattern) for pattern in patterns)
+
+
+def _iter_libs(report: PackageReport) -> tuple[ExtensionReport, ...]:
+    return (*report.extensions, *report.bundled_libs)
+
+
+def _collect_unstable(report: PackageReport) -> list[str]:
+    return sorted(
+        _relative_path(lib, report.root)
+        for lib in _iter_libs(report)
+        if _is_torch_unstable(lib)
+    )
+
+
+def _find_stale_allowlist_entries(
+    report: PackageReport, patterns: tuple[str, ...]
+) -> list[str]:
+    """Allowlist patterns that match a built library which is no longer unstable."""
+    stale: list[str] = []
+    for pattern in patterns:
+        for lib in _iter_libs(report):
+            if lib.error is not None:
+                continue
+            if not fnmatch.fnmatch(_relative_path(lib, report.root), pattern):
+                continue
+            if not _is_torch_unstable(lib):
+                stale.append(pattern)
+                break
+    return stale
+
+
+def check_torch_abi(
+    package: str = "vllm",
+    patterns: tuple[str, ...] = ALLOWED_UNSTABLE_LIBRARIES,
+) -> int:
+    report = inspect_package(package)
+    if report.error:
+        print(f"error: failed to inspect {package!r}: {report.error}", file=sys.stderr)
+        return 2
+
+    unstable = _collect_unstable(report)
+    unexpected = [
+        rel_path for rel_path in unstable if not _matches_allowlist(rel_path, patterns)
+    ]
+    stale = _find_stale_allowlist_entries(report, patterns)
+
+    if unexpected or stale:
+        if unexpected:
+            print(
+                "Not allowed: torch-unstable libraries outside "
+                f"ALLOWED_UNSTABLE_LIBRARIES: {', '.join(unexpected)}",
+                file=sys.stderr,
+            )
+        if stale:
+            print(
+                "Not allowed: stale ALLOWED_UNSTABLE_LIBRARIES entries: "
+                f"{', '.join(stale)}",
+                file=sys.stderr,
+            )
+        return 1
+
+    print("Torch stable ABI check passed.")
+    return 0
+
+
+if __name__ == "__main__":
+    print(">>> Auditing vLLM extension modules for PyTorch stable ABI compliance")
+    sys.exit(check_torch_abi())
diff --git a/.buildkite/ci_config.yaml b/.buildkite/ci_config.yaml
index 21ffa1b9b8d..9e1e46db67e 100644
--- a/.buildkite/ci_config.yaml
+++ b/.buildkite/ci_config.yaml
@@ -14,6 +14,7 @@ run_all_patterns:
   - "setup.py"
   - "csrc/"
   - "cmake/"
+  - ".buildkite/check-torch-abi.py"
 run_all_exclude_patterns:
   - "docker/Dockerfile."
   - "csrc/cpu/"
diff --git a/.buildkite/intel_jobs/test-intel.yaml b/.buildkite/intel_jobs/test-intel.yaml
index 3fadb07f391..ec5cb2fd9e7 100644
--- a/.buildkite/intel_jobs/test-intel.yaml
+++ b/.buildkite/intel_jobs/test-intel.yaml
@@ -147,6 +147,30 @@ steps:
         'cd tests &&
         pytest -v -s quantization/test_auto_round.py &&
         pytest -v -s quantization/test_online.py'
+  - label: "XPU GPQA Eval (GPT-OSS)"
+    depends_on:
+      - image-build-xpu
+    timeout_in_minutes: 60
+    device: intel_gpu
+    agent_tags:
+      label: production
+      gpu: 1+
+      mem: 24+
+    no_plugin: true
+    env:
+      REGISTRY: "public.ecr.aws/q9t5s3a7"
+      REPO: "vllm-ci-test-repo"
+      VLLM_TEST_DEVICE: "xpu"
+    source_file_dependencies:
+      - vllm/
+      - tests/evals/gpt_oss/
+      - .buildkite/intel_jobs/test-intel.yaml
+    commands:
+      - >-
+        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+        'pip install "gpt-oss[eval]==0.0.5" &&
+        cd tests &&
+        pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-xpu.txt'
   - label: "XPU compressed tensors FP8 test"
     depends_on:
       - image-build-xpu
diff --git a/.buildkite/test_areas/cuda.yaml b/.buildkite/test_areas/cuda.yaml
index 927b5bd27f2..431ce07af4d 100644
--- a/.buildkite/test_areas/cuda.yaml
+++ b/.buildkite/test_areas/cuda.yaml
@@ -16,6 +16,7 @@ steps:
   commands:
     - pytest -v -s cuda/test_cuda_context.py
     - pytest -v -s cuda/test_platform_no_cuda_init.py
+    - pytest -v -s cuda/test_cuda_compatibility_path.py
 
 - label: Cudagraph
   device: h200_35gb
diff --git a/.buildkite/test_areas/disaggregated.yaml b/.buildkite/test_areas/disaggregated.yaml
index a3342e362ed..f1a89b39682 100644
--- a/.buildkite/test_areas/disaggregated.yaml
+++ b/.buildkite/test_areas/disaggregated.yaml
@@ -131,6 +131,22 @@ steps:
         - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
         - HYBRID_SSM=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
+- label: NixlConnector PD edge case test (2 GPUs)
+  key: nixlconnector-pd-edge-cases-2-gpus
+  timeout_in_minutes: 40
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/v1/core/sched/
+    - tests/v1/kv_connector/nixl_integration/
+  env:
+    PREFILL_GPU_ID: "0"
+    DECODE_GPU_ID: "1"
+  commands:
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - bash v1/kv_connector/nixl_integration/run_edge_case_test.sh
+
 - label: Hybrid SSM NixlConnector PD prefix cache test (2 GPUs)
   key: hybrid-ssm-nixlconnector-pd-prefix-cache-2-gpus
   timeout_in_minutes: 25
diff --git a/.buildkite/test_areas/engine.yaml b/.buildkite/test_areas/engine.yaml
index ed593c4aba2..ce4fc590eec 100644
--- a/.buildkite/test_areas/engine.yaml
+++ b/.buildkite/test_areas/engine.yaml
@@ -40,9 +40,11 @@ steps:
   source_file_dependencies:
     - vllm/v1/engine/
     - tests/v1/engine/
+    - tests/v1/test_tensor_ipc_queue.py
   commands:
     - pytest -v -s v1/engine/test_preprocess_error_handling.py
     - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
+    - pytest -v -s v1/test_tensor_ipc_queue.py
   mirror:
     amd:
       device: mi250_1
diff --git a/.buildkite/test_areas/expert_parallelism.yaml b/.buildkite/test_areas/expert_parallelism.yaml
index a3b46b58285..1d1609d46b8 100644
--- a/.buildkite/test_areas/expert_parallelism.yaml
+++ b/.buildkite/test_areas/expert_parallelism.yaml
@@ -52,4 +52,5 @@ steps:
   - vllm/compilation/
   - tests/distributed/
   commands:
+  - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
   - pytest -v -s distributed/test_elastic_ep.py
diff --git a/.buildkite/test_areas/kernels.yaml b/.buildkite/test_areas/kernels.yaml
index e1951685f60..81fa9c5f1d8 100644
--- a/.buildkite/test_areas/kernels.yaml
+++ b/.buildkite/test_areas/kernels.yaml
@@ -61,9 +61,45 @@ steps:
   source_file_dependencies:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
   - vllm/models/deepseek_v4/common/ops/
+  - vllm/models/deepseek_v4/nvidia/
   - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
+  - tests/models/test_deepseek_v4_mega_moe.py
   commands:
     - pytest -v -s kernels/test_fused_deepseek_v4_*.py
+    - pytest -v -s models/test_deepseek_v4_mega_moe.py
+
+# Catch-all for test files at the tests/kernels root. This job collects
+# the whole root so new files are wired by default.
+# Files with dedicated jobs elsewhere in this file are excluded via --ignore
+# (test_kda, test_bf16x3_router_gemm_cutedsl and test_ll_bf16_gemm run in
+# their own jobs / Kernels (B200)).
+- label: Kernels Root Misc Test (B200)
+  key: kernels-root-misc-test-b200
+  timeout_in_minutes: 45
+  device: b200-k8s
+  source_file_dependencies:
+  - csrc/
+  - vllm/
+  - tests/kernels/
+  commands:
+    - pytest -v -s kernels/
+        --ignore=kernels/attention
+        --ignore=kernels/core
+        --ignore=kernels/helion
+        --ignore=kernels/ir
+        --ignore=kernels/mamba
+        --ignore=kernels/moe
+        --ignore=kernels/quantization
+        --ignore=kernels/test_concat_mla_q.py
+        --ignore=kernels/test_fused_qk_norm_rope_gate.py
+        --ignore=kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
+        --ignore=kernels/test_top_k_per_row.py
+        --ignore=kernels/test_kda.py
+        --ignore=kernels/test_bf16x3_router_gemm_cutedsl.py
+        --ignore=kernels/test_ll_bf16_gemm.py
+        --ignore=kernels/test_shuffle_rows.py
+    # BROKEN on main, pending kernel fixes (B200):
+    #   test_shuffle_rows.py (1: test_shuffle_rows_edge_cases)
 
 - label: Kernels Attention Test %N
   key: kernels-attention-test
@@ -178,17 +214,6 @@ steps:
   commands:
     - pytest -v -s kernels/mamba
 
-- label: Kernels KDA Test
-  timeout_in_minutes: 25
-  device: h200_18gb
-  source_file_dependencies:
-  - vllm/third_party/flash_linear_attention/ops/kda.py
-  - vllm/third_party/flash_linear_attention/ops/chunk_delta_h.py
-  - vllm/third_party/flash_linear_attention/ops/l2norm.py
-  - tests/kernels/test_kda.py
-  commands:
-    - pytest -v -s kernels/test_kda.py
-
 - label: Kernels DeepGEMM Test (H100)
   key: kernels-deepgemm-test-h100
   timeout_in_minutes: 35
diff --git a/.buildkite/test_areas/misc.yaml b/.buildkite/test_areas/misc.yaml
index 1a53a92961f..caa56c21b37 100644
--- a/.buildkite/test_areas/misc.yaml
+++ b/.buildkite/test_areas/misc.yaml
@@ -148,6 +148,7 @@ steps:
     - pytest -v -s -m 'cpu_test' v1/core
     - pytest -v -s v1/structured_output
     - pytest -v -s v1/test_serial_utils.py
+    - pytest -v -s v1/test_kv_cache_spec_registry.py
     - pytest -v -s v1/cudagraph/test_cudagraph_manager.py
     - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'cpu_test' v1/metrics
@@ -265,6 +266,7 @@ steps:
   - vllm/utils/
   - vllm/v1/
   - tests/v1/tracing
+  - tests/tracing/
   commands:
   - "pip install \
       'opentelemetry-sdk>=1.26.0' \
@@ -272,6 +274,7 @@ steps:
       'opentelemetry-exporter-otlp>=1.26.0' \
       'opentelemetry-semantic-conventions-ai>=0.4.1'"
   - pytest -v -s v1/tracing
+  - pytest -v -s tracing
   mirror:
     amd:
       dind: false
@@ -395,7 +398,7 @@ steps:
 
 - label: Batch Invariance (A100)
   key: batch-invariance-a100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 60
   device: a100
   source_file_dependencies:
     - vllm/v1/attention
@@ -405,11 +408,11 @@ steps:
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
     - pip install pytest-timeout pytest-forked
     - pytest -v -s v1/determinism/test_batch_invariance.py
-    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
+    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
 
 - label: Batch Invariance (H100)
   key: batch-invariance-h100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 60
   device: h100
   source_file_dependencies:
     - vllm/v1/attention
@@ -420,12 +423,12 @@ steps:
     - pip install pytest-timeout pytest-forked
     - pytest -v -s v1/determinism/test_batch_invariance.py
     - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
-    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
-    - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
+    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
+    - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k FLASH_ATTN
 
 - label: Batch Invariance (B200)
   key: batch-invariance-b200
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
   device: b200-k8s
   source_file_dependencies:
     - vllm/v1/attention
@@ -436,11 +439,14 @@ steps:
     - pip install pytest-timeout pytest-forked
     - pytest -v -s v1/determinism/test_batch_invariance.py
     - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
-    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
-    - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
+    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
+    - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k FLASH_ATTN
     - pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
     - pytest -v -s v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py
-  
+    - pytest -v -s v1/determinism/test_matmul_batch_invariant.py
+    - pytest -v -s v1/determinism/test_cutlass_batch_invariance.py
+    - pytest -v -s v1/determinism/test_online_batch_invariance.py
+
 - label: Acceptance Length Test (Large Models) # optional
   device: h200_35gb
   key: acceptance-length-test-large-models
diff --git a/.buildkite/test_areas/models_basic.yaml b/.buildkite/test_areas/models_basic.yaml
index af90308d6c3..e231e6d90f7 100644
--- a/.buildkite/test_areas/models_basic.yaml
+++ b/.buildkite/test_areas/models_basic.yaml
@@ -61,6 +61,20 @@ steps:
     # FA4 kernel tests require SM100; the suite skips them elsewhere.
     - pytest -v -s models/inkling
 
+- label: Kimi K3 Unit Tests (B200)
+  key: kimi-k3-unit-tests-b200
+  timeout_in_minutes: 40
+  device: b200-k8s
+  source_file_dependencies:
+  - vllm/models/kimi_k3/
+  - csrc/libtorch_stable/kimi_k3/
+  - tests/models/kimi_k3/
+  - tests/kernels/attention/test_kimi_k3_mla_fused_epilogue.py
+  - tests/kernels/test_bf16_skinny_gemm.py
+  commands:
+    # The native NVIDIA Kimi K3 kernels require the SM100 family.
+    - pytest -v -s models/kimi_k3 kernels/attention/test_kimi_k3_mla_fused_epilogue.py kernels/test_bf16_skinny_gemm.py
+
 - label: Basic Models Test (Other CPU) # 5min
   key: basic-models-test-other-cpu
   depends_on:
@@ -70,7 +84,8 @@ steps:
   - vllm/
   - tests/models/test_utils.py
   - tests/models/test_vision.py
+  - tests/models/test_adapters.py
   - tests/models/transformers/fusers/
   device: cpu-small
   commands:
-    - pytest -v -s models/test_utils.py models/test_vision.py models/transformers/fusers/
+    - pytest -v -s models/test_utils.py models/test_vision.py models/test_adapters.py models/transformers/fusers/
diff --git a/.buildkite/test_areas/spec_decode.yaml b/.buildkite/test_areas/spec_decode.yaml
index c63aaa18d8b..7cde7124cdc 100644
--- a/.buildkite/test_areas/spec_decode.yaml
+++ b/.buildkite/test_areas/spec_decode.yaml
@@ -90,8 +90,10 @@ steps:
     - vllm/v1/spec_decode/
     - vllm/v1/worker/gpu/spec_decode/
     - tests/v1/e2e/spec_decode/
+    - tests/spec_decode/
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
+    - python3 spec_decode/test_custom_proposer.py
   mirror:
     amd:
       dind: false
diff --git a/.buildkite/test_areas/torch_abi.yaml b/.buildkite/test_areas/torch_abi.yaml
new file mode 100644
index 00000000000..eaef3551664
--- /dev/null
+++ b/.buildkite/test_areas/torch_abi.yaml
@@ -0,0 +1,14 @@
+group: Torch ABI
+depends_on:
+  - image-build
+steps:
+- label: Torch Stable ABI Audit
+  key: torch-stable-abi-audit
+  timeout_in_minutes: 5
+  source_file_dependencies:
+  - .buildkite/check-torch-abi.py
+  - csrc/
+  - cmake/
+  - setup.py
+  commands:
+  - python3 /vllm-workspace/.buildkite/check-torch-abi.py
diff --git a/.github/workflows/new_pr_bot.yml b/.github/workflows/new_pr_bot.yml
index 4124583d96d..a2f09eb8a45 100644
--- a/.github/workflows/new_pr_bot.yml
+++ b/.github/workflows/new_pr_bot.yml
@@ -80,9 +80,9 @@ jobs:
                   '',
                   '\u{1f4ac} Join our developer Slack at https://slack.vllm.ai to discuss your PR in `#pr-reviews`, coordinate on features in `#feat-` channels, or join special interest groups in `#sig-` channels.',
                   '',
-                  'PRs do not trigger a full CI run by default. Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging.',
+                  'PRs do not trigger a full CI run by default. Reviewers with write access and configured trusted contributors can comment `/ci run` whenever CI signals are needed.',
                   '',
-                  'To run CI, PR reviewers can either: Add `ready` label to the PR or enable auto-merge.',
+                  'Once the PR is approved or has the `ready` label, the PR author can also use `/ci run` or `/ci retry`. New commits do not start CI automatically.',
                   '',
                   'If you have any questions, please reach out to us on Slack at https://slack.vllm.ai.',
                   '',
diff --git a/.github/workflows/pre-commit.yml b/.github/workflows/pre-commit.yml
index 143fc427a49..aa1f437ab6a 100644
--- a/.github/workflows/pre-commit.yml
+++ b/.github/workflows/pre-commit.yml
@@ -41,7 +41,7 @@ jobs:
           if (hasReadyLabel || hasVerifiedLabel || mergedCount >= 4) {
             core.info(`Check passed: verified label=${hasVerifiedLabel}, ready label=${hasReadyLabel}, 4+ merged PRs=${mergedCount >= 4}`);
           } else {
-            core.setFailed(`PR must have the 'verified', 'ready', or 'ready-run-all-tests' label (the ready labels also trigger tests) or the author must have at least 4 merged PRs (found ${mergedCount}).`);
+            core.setFailed(`PR must have the 'verified', 'ready', or 'ready-run-all-tests' label to run pre-commit, or the author must have at least 4 merged PRs (found ${mergedCount}).`);
           }
 
   pre-commit:
diff --git a/.github/workflows/run-ci-command.yml b/.github/workflows/run-ci-command.yml
new file mode 100644
index 00000000000..ff55c65dd75
--- /dev/null
+++ b/.github/workflows/run-ci-command.yml
@@ -0,0 +1,40 @@
+name: Run CI from PR comment
+
+on:
+  issue_comment:
+    types: [created]
+
+concurrency:
+  group: run-ci-comment-${{ github.event.issue.number }}
+  cancel-in-progress: false
+
+permissions:
+  contents: read
+  issues: write
+  pull-requests: read
+
+jobs:
+  run-ci-command:
+    if: >-
+      github.event.issue.pull_request &&
+      (github.event.comment.body == '/ci run' ||
+      github.event.comment.body == '/ci retry')
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@9c091bb21b7c1c1d1991bb908d89e4e9dddfe3e0 # v7.0.0
+        with:
+          ref: ${{ github.event.repository.default_branch }}
+          persist-credentials: false
+      - uses: astral-sh/setup-uv@37802adc94f370d6bfd71619e3f0bf239e1f3b78 # v7.6.0
+        with:
+          python-version: "3.12"
+      - name: Authorize and run CI command
+        run: >-
+          uv run --no-project --python 3.12
+          .github/workflows/scripts/run_ci_command.py
+        env:
+          BUILDKITE_API_TOKEN: ${{ secrets.BUILDKITE_API_TOKEN }}
+          BUILDKITE_ORGANIZATION: vllm
+          BUILDKITE_PIPELINE: ci
+          CI_TRUSTED_USERS: ${{ vars.CI_TRUSTED_USERS }}
+          GH_TOKEN: ${{ github.token }}
diff --git a/.github/workflows/scripts/run_ci_command.py b/.github/workflows/scripts/run_ci_command.py
new file mode 100644
index 00000000000..5bf196bb53f
--- /dev/null
+++ b/.github/workflows/scripts/run_ci_command.py
@@ -0,0 +1,581 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import json
+import os
+import sys
+import urllib.error
+import urllib.parse
+import urllib.request
+from collections.abc import Mapping, Sequence
+from typing import Any
+
+COMMAND_RUN_CI = "/ci run"
+COMMAND_RETRY_FAILED = "/ci retry"
+READY_LABELS = {"ready", "ready-run-all-tests"}
+TRUSTED_PERMISSIONS = {"admin", "maintain", "write"}
+ACTIVE_BUILD_STATES = {
+    "blocked",
+    "creating",
+    "scheduled",
+    "running",
+    "failing",
+    "canceling",
+    "waiting",
+    "waiting_failed",
+}
+RETRY_STATES = "failed,timed_out,expired"
+
+
+class ApiError(RuntimeError):
+    def __init__(self, status: int | None, message: str) -> None:
+        super().__init__(message)
+        self.status = status
+
+
+class HttpTransport:
+    def request(
+        self,
+        url: str,
+        *,
+        body: Mapping[str, Any] | None = None,
+        headers: Mapping[str, str] | None = None,
+        method: str = "GET",
+    ) -> Any:
+        data = None if body is None else json.dumps(body).encode()
+        request = urllib.request.Request(
+            url,
+            data=data,
+            headers=dict(headers or {}),
+            method=method,
+        )
+        try:
+            with urllib.request.urlopen(request, timeout=30) as response:
+                response_body = response.read().decode()
+        except urllib.error.HTTPError as error:
+            response_body = error.read().decode()
+            message = self._error_message(response_body, error.reason)
+            raise ApiError(
+                error.code,
+                f"API returned {error.code}: {message}",
+            ) from error
+        except urllib.error.URLError as error:
+            raise ApiError(None, f"API request failed: {error.reason}") from error
+
+        if not response_body:
+            return None
+        try:
+            return json.loads(response_body)
+        except json.JSONDecodeError as error:
+            raise ApiError(None, "API returned a non-JSON response.") from error
+
+    @staticmethod
+    def _error_message(response_body: str, fallback: str) -> str:
+        try:
+            parsed = json.loads(response_body)
+        except json.JSONDecodeError:
+            return fallback
+        return str(parsed.get("message", fallback))
+
+
+class GitHubClient:
+    def __init__(
+        self,
+        token: str,
+        repository: str,
+        transport: HttpTransport | None = None,
+    ) -> None:
+        if not token:
+            raise RuntimeError("GH_TOKEN is not set.")
+        self.owner, self.repo = repository.split("/", maxsplit=1)
+        self.transport = transport or HttpTransport()
+        self.headers = {
+            "Accept": "application/vnd.github+json",
+            "Authorization": f"Bearer {token}",
+            "Content-Type": "application/json",
+            "User-Agent": "vllm-ci-command",
+            "X-GitHub-Api-Version": "2022-11-28",
+        }
+
+    def _request(
+        self,
+        path: str,
+        *,
+        body: Mapping[str, Any] | None = None,
+        method: str = "GET",
+    ) -> Any:
+        return self.transport.request(
+            f"https://api.github.com{path}",
+            body=body,
+            headers=self.headers,
+            method=method,
+        )
+
+    def _repo_path(self, suffix: str) -> str:
+        owner = urllib.parse.quote(self.owner, safe="")
+        repo = urllib.parse.quote(self.repo, safe="")
+        return f"/repos/{owner}/{repo}{suffix}"
+
+    def _paginate(self, path: str) -> list[dict[str, Any]]:
+        results: list[dict[str, Any]] = []
+        separator = "&" if "?" in path else "?"
+        for page in range(1, 101):
+            response = self._request(f"{path}{separator}per_page=100&page={page}")
+            if not isinstance(response, list):
+                raise ApiError(None, "GitHub API returned an invalid list response.")
+            results.extend(response)
+            if len(response) < 100:
+                return results
+        raise ApiError(None, "GitHub API pagination exceeded 10,000 results.")
+
+    def get_pr(self, number: int) -> dict[str, Any]:
+        return self._request(self._repo_path(f"/pulls/{number}"))
+
+    def get_permission(self, actor: str) -> str:
+        username = urllib.parse.quote(actor, safe="")
+        try:
+            response = self._request(
+                self._repo_path(f"/collaborators/{username}/permission")
+            )
+        except ApiError as error:
+            if error.status == 404:
+                return "none"
+            raise
+        return str(response["permission"])
+
+    def get_review_decision(self, number: int) -> str | None:
+        query = """
+          query($owner: String!, $repo: String!, $number: Int!) {
+            repository(owner: $owner, name: $repo) {
+              pullRequest(number: $number) {
+                reviewDecision
+              }
+            }
+          }
+        """
+        response = self._request(
+            "/graphql",
+            body={
+                "query": query,
+                "variables": {
+                    "number": number,
+                    "owner": self.owner,
+                    "repo": self.repo,
+                },
+            },
+            method="POST",
+        )
+        return response["data"]["repository"]["pullRequest"]["reviewDecision"]
+
+    def list_reviews(self, number: int) -> list[dict[str, Any]]:
+        return self._paginate(self._repo_path(f"/pulls/{number}/reviews"))
+
+    def list_reactions(self, comment_id: int) -> list[dict[str, Any]]:
+        return self._paginate(
+            self._repo_path(f"/issues/comments/{comment_id}/reactions")
+        )
+
+    def add_reaction(self, comment_id: int, content: str) -> None:
+        self._request(
+            self._repo_path(f"/issues/comments/{comment_id}/reactions"),
+            body={"content": content},
+            method="POST",
+        )
+
+    def add_comment(self, issue_number: int, body: str) -> None:
+        self._request(
+            self._repo_path(f"/issues/{issue_number}/comments"),
+            body={"body": body},
+            method="POST",
+        )
+
+
+class BuildkiteClient:
+    def __init__(
+        self,
+        token: str,
+        organization: str,
+        pipeline: str,
+        transport: HttpTransport | None = None,
+    ) -> None:
+        self.token = token
+        self.transport = transport or HttpTransport()
+        organization = urllib.parse.quote(organization, safe="")
+        pipeline = urllib.parse.quote(pipeline, safe="")
+        self.base_url = (
+            "https://api.buildkite.com/v2/organizations/"
+            f"{organization}/pipelines/{pipeline}/builds"
+        )
+
+    def _request(
+        self,
+        *,
+        body: Mapping[str, Any] | None = None,
+        method: str = "GET",
+        path: str = "",
+        query: Sequence[tuple[str, str]] = (),
+    ) -> Any:
+        if not self.token:
+            raise RuntimeError("The BUILDKITE_API_TOKEN repository secret is not set.")
+        url = f"{self.base_url}{path}"
+        if query:
+            url = f"{url}?{urllib.parse.urlencode(query)}"
+        return self.transport.request(
+            url,
+            body=body,
+            headers={
+                "Authorization": f"Bearer {self.token}",
+                "Content-Type": "application/json",
+                "User-Agent": "vllm-ci-command",
+            },
+            method=method,
+        )
+
+    def list_builds(
+        self,
+        commit: str,
+        *,
+        metadata: tuple[str, str] | None = None,
+    ) -> list[dict[str, Any]]:
+        query = [
+            ("commit", commit),
+            ("exclude_jobs", "true"),
+            ("exclude_pipeline", "true"),
+            ("per_page", "100"),
+        ]
+        if metadata:
+            key, value = metadata
+            query.append((f"meta_data[{key}]", value))
+        response = self._request(query=query)
+        if not isinstance(response, list):
+            raise ApiError(None, "Buildkite API returned an invalid build list.")
+        return response
+
+    def create_build(self, body: Mapping[str, Any]) -> dict[str, Any]:
+        return self._request(body=body, method="POST")
+
+    def retry_failed_jobs(
+        self,
+        build_number: int,
+        states: str,
+    ) -> dict[str, Any]:
+        number = urllib.parse.quote(str(build_number), safe="")
+        return self._request(
+            body={"states": states},
+            method="PUT",
+            path=f"/{number}/retry_failed_jobs",
+        )
+
+
+def parse_command(body: str) -> str | None:
+    if body in {COMMAND_RUN_CI, COMMAND_RETRY_FAILED}:
+        return body
+    return None
+
+
+def parse_trusted_users(value: str = "") -> set[str]:
+    return {
+        user.casefold() for item in value.split(",") for user in item.split() if user
+    }
+
+
+def has_ready_label(pr: Mapping[str, Any]) -> bool:
+    return any(label["name"] in READY_LABELS for label in pr["labels"])
+
+
+def is_trusted_permission(permission: str) -> bool:
+    return permission in TRUSTED_PERMISSIONS
+
+
+def authorize(
+    *,
+    actor: str,
+    permission: str,
+    pr: Mapping[str, Any],
+    trusted_approval: bool = False,
+    trusted_users: set[str] | None = None,
+) -> tuple[bool, str]:
+    trusted_users = trusted_users or set()
+    if is_trusted_permission(permission):
+        return True, f"repository {permission} permission"
+    if actor.casefold() in trusted_users:
+        return True, "configured trusted contributor"
+    if actor.casefold() != pr["user"]["login"].casefold():
+        return (
+            False,
+            "Only reviewers with write access can run CI before it is "
+            "delegated to the PR author.",
+        )
+    if pr["draft"]:
+        return False, "PR authors cannot run CI while the PR is a draft."
+    if has_ready_label(pr):
+        return True, "ready label"
+    if trusted_approval:
+        return True, "approval from a trusted reviewer"
+    return (
+        False,
+        "A reviewer with write access must run `/ci run`, approve the PR, "
+        "or add the `ready` label first.",
+    )
+
+
+def has_trusted_approval(
+    github: GitHubClient,
+    number: int,
+    trusted_users: set[str],
+) -> bool:
+    if github.get_review_decision(number) != "APPROVED":
+        return False
+
+    latest_review_states: dict[str, tuple[str, str]] = {}
+    for review in github.list_reviews(number):
+        user = review.get("user") or {}
+        login = user.get("login")
+        state = review.get("state")
+        if login and state in {"APPROVED", "CHANGES_REQUESTED", "DISMISSED"}:
+            latest_review_states[login.casefold()] = (login, state)
+
+    for login, state in latest_review_states.values():
+        if state != "APPROVED":
+            continue
+        if login.casefold() in trusted_users:
+            return True
+        if is_trusted_permission(github.get_permission(login)):
+            return True
+    return False
+
+
+def is_build_for_pr(build: Mapping[str, Any], pr_number: int) -> bool:
+    pull_request = build.get("pull_request")
+    if isinstance(pull_request, Mapping):
+        build_pr_number = pull_request.get("id", pull_request.get("number"))
+        if build_pr_number is not None:
+            return str(build_pr_number) == str(pr_number)
+    metadata = build.get("meta_data") or {}
+    return str(metadata.get("github-pr-number")) == str(pr_number)
+
+
+def is_active_build(build: Mapping[str, Any]) -> bool:
+    return bool(build.get("blocked")) or build.get("state") in ACTIVE_BUILD_STATES
+
+
+def select_latest_build(
+    builds: Sequence[dict[str, Any]],
+    pr_number: int,
+) -> dict[str, Any] | None:
+    matching = [build for build in builds if is_build_for_pr(build, pr_number)]
+    return max(matching, key=lambda build: build.get("created_at", ""), default=None)
+
+
+def create_build_payload(
+    *,
+    actor: str,
+    comment_id: int,
+    pr: Mapping[str, Any],
+) -> dict[str, Any]:
+    return {
+        "commit": pr["head"]["sha"],
+        "branch": pr["head"]["ref"],
+        "message": f"PR #{pr['number']} {COMMAND_RUN_CI} by @{actor}",
+        "pull_request_id": pr["number"],
+        "pull_request_base_branch": pr["base"]["ref"],
+        "pull_request_repository": pr["head"]["repo"]["clone_url"],
+        "pull_request_labels": [label["name"] for label in pr["labels"]],
+        "ignore_pipeline_branch_filters": True,
+        "env": {
+            "VLLM_CI_GITHUB_COMMENT_ID": str(comment_id),
+            "VLLM_CI_TRIGGERED_BY": actor,
+        },
+        "meta_data": {
+            "github-comment-id": str(comment_id),
+            "github-pr-number": str(pr["number"]),
+            "github-triggered-by": actor,
+        },
+    }
+
+
+def add_reaction_safely(
+    github: GitHubClient,
+    comment_id: int,
+    content: str,
+) -> None:
+    try:
+        github.add_reaction(comment_id, content)
+    except Exception as error:
+        print(f"Could not add {content} reaction: {error}", file=sys.stderr)
+
+
+def is_already_handled(github: GitHubClient, comment_id: int) -> bool:
+    return any(
+        reaction.get("content") in {"rocket", "-1"}
+        and (reaction.get("user") or {}).get("login") == "github-actions[bot]"
+        for reaction in github.list_reactions(comment_id)
+    )
+
+
+def handle_run_ci(
+    *,
+    actor: str,
+    buildkite: BuildkiteClient,
+    comment_id: int,
+    github: GitHubClient,
+    pr: Mapping[str, Any],
+) -> str:
+    duplicate_builds = buildkite.list_builds(
+        pr["head"]["sha"],
+        metadata=("github-comment-id", str(comment_id)),
+    )
+    duplicate = select_latest_build(duplicate_builds, pr["number"])
+    if duplicate:
+        return f"CI was already requested by this comment: {duplicate['web_url']}"
+
+    current_builds = buildkite.list_builds(pr["head"]["sha"])
+    active_build = next(
+        (
+            build
+            for build in current_builds
+            if is_build_for_pr(build, pr["number"]) and is_active_build(build)
+        ),
+        None,
+    )
+    if active_build:
+        return f"CI is already running for this commit: {active_build['web_url']}"
+
+    current_pr = github.get_pr(pr["number"])
+    if current_pr["state"] != "open" or current_pr["head"]["sha"] != pr["head"]["sha"]:
+        return (
+            "The PR head changed while processing the command. Comment `/ci run` again."
+        )
+
+    build = buildkite.create_build(
+        create_build_payload(
+            actor=actor,
+            comment_id=comment_id,
+            pr=current_pr,
+        )
+    )
+    return (
+        f"Triggered [Buildkite CI #{build['number']}]({build['web_url']}) "
+        f"for commit `{current_pr['head']['sha'][:12]}`."
+    )
+
+
+def handle_retry_failed(
+    *,
+    buildkite: BuildkiteClient,
+    pr: Mapping[str, Any],
+) -> str:
+    builds = buildkite.list_builds(pr["head"]["sha"])
+    build = select_latest_build(builds, pr["number"])
+    if not build:
+        return "No CI build exists for the current PR commit. Use `/ci run` first."
+    if not build.get("finished_at") or is_active_build(build):
+        return f"CI is still running for this commit: {build['web_url']}"
+
+    retried = buildkite.retry_failed_jobs(build["number"], RETRY_STATES)
+    if retried["retried_jobs_count"] == 0:
+        return (
+            f"No failed, timed-out, or expired jobs need retrying: {build['web_url']}"
+        )
+    return (
+        f"Queued {retried['retried_jobs_count']} failed job(s) for retry in "
+        f"[Buildkite CI #{build['number']}]({build['web_url']})."
+    )
+
+
+def run(
+    event: Mapping[str, Any],
+    github: GitHubClient,
+    buildkite: BuildkiteClient,
+    trusted_users_value: str = "",
+) -> None:
+    command = parse_command(event["comment"]["body"])
+    if not command or "pull_request" not in event["issue"]:
+        return
+
+    issue_number = event["issue"]["number"]
+    comment_id = event["comment"]["id"]
+    actor = event["comment"]["user"]["login"]
+
+    if is_already_handled(github, comment_id):
+        print(f"Comment {comment_id} was already handled.")
+        return
+    add_reaction_safely(github, comment_id, "eyes")
+
+    try:
+        pr = github.get_pr(issue_number)
+        permission = github.get_permission(actor)
+        if pr["state"] != "open":
+            github.add_comment(issue_number, "CI commands require an open PR.")
+            return
+
+        trusted_users = parse_trusted_users(trusted_users_value)
+        should_check_approval = (
+            not is_trusted_permission(permission)
+            and actor.casefold() not in trusted_users
+            and actor.casefold() == pr["user"]["login"].casefold()
+            and not pr["draft"]
+            and not has_ready_label(pr)
+        )
+        trusted_approval = should_check_approval and has_trusted_approval(
+            github,
+            issue_number,
+            trusted_users,
+        )
+        allowed, reason = authorize(
+            actor=actor,
+            permission=permission,
+            pr=pr,
+            trusted_approval=trusted_approval,
+            trusted_users=trusted_users,
+        )
+        if not allowed:
+            add_reaction_safely(github, comment_id, "-1")
+            github.add_comment(issue_number, f"@{actor}, {reason}")
+            return
+
+        print(f"Authorized @{actor}: {reason}")
+        if command == COMMAND_RUN_CI:
+            message = handle_run_ci(
+                actor=actor,
+                buildkite=buildkite,
+                comment_id=comment_id,
+                github=github,
+                pr=pr,
+            )
+        else:
+            message = handle_retry_failed(buildkite=buildkite, pr=pr)
+        add_reaction_safely(github, comment_id, "rocket")
+        github.add_comment(issue_number, message)
+    except Exception:
+        add_reaction_safely(github, comment_id, "confused")
+        raise
+
+
+def main() -> None:
+    event_path = os.environ["GITHUB_EVENT_PATH"]
+    with open(event_path, encoding="utf-8") as event_file:
+        event = json.load(event_file)
+
+    if not parse_command(event["comment"]["body"]):
+        return
+
+    github = GitHubClient(
+        os.environ.get("GH_TOKEN", ""),
+        os.environ["GITHUB_REPOSITORY"],
+    )
+    buildkite = BuildkiteClient(
+        os.environ.get("BUILDKITE_API_TOKEN", ""),
+        os.environ.get("BUILDKITE_ORGANIZATION", "vllm"),
+        os.environ.get("BUILDKITE_PIPELINE", "ci"),
+    )
+    run(
+        event,
+        github,
+        buildkite,
+        os.environ.get("CI_TRUSTED_USERS", ""),
+    )
+
+
+if __name__ == "__main__":
+    main()
diff --git a/.github/workflows/scripts/test_run_ci_command.py b/.github/workflows/scripts/test_run_ci_command.py
new file mode 100644
index 00000000000..edbfff3f5f0
--- /dev/null
+++ b/.github/workflows/scripts/test_run_ci_command.py
@@ -0,0 +1,363 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import unittest
+from typing import Any
+
+from run_ci_command import (
+    COMMAND_RETRY_FAILED,
+    COMMAND_RUN_CI,
+    RETRY_STATES,
+    BuildkiteClient,
+    authorize,
+    create_build_payload,
+    has_trusted_approval,
+    is_active_build,
+    is_build_for_pr,
+    parse_command,
+    parse_trusted_users,
+    run,
+    select_latest_build,
+)
+
+
+def make_pr(**overrides: Any) -> dict[str, Any]:
+    pr = {
+        "base": {"ref": "main"},
+        "draft": False,
+        "head": {
+            "ref": "feature",
+            "repo": {"clone_url": "https://github.com/contributor/vllm.git"},
+            "sha": "0123456789abcdef",
+        },
+        "labels": [],
+        "number": 42,
+        "state": "open",
+        "user": {"login": "author"},
+    }
+    pr.update(overrides)
+    return pr
+
+
+def make_event(command: str, actor: str = "reviewer") -> dict[str, Any]:
+    return {
+        "comment": {
+            "body": command,
+            "id": 99,
+            "user": {"login": actor},
+        },
+        "issue": {
+            "number": 42,
+            "pull_request": {},
+        },
+    }
+
+
+class FakeGitHub:
+    def __init__(
+        self,
+        *,
+        permission: str = "write",
+        permissions: dict[str, str] | None = None,
+        pr: dict[str, Any] | None = None,
+        review_decision: str = "REVIEW_REQUIRED",
+        reviews: list[dict[str, Any]] | None = None,
+    ) -> None:
+        self.comments: list[str] = []
+        self.permission = permission
+        self.permissions = permissions or {}
+        self.pr = pr or make_pr()
+        self.reactions: list[str] = []
+        self.review_decision = review_decision
+        self.reviews = reviews or []
+
+    def get_pr(self, number: int) -> dict[str, Any]:
+        return self.pr
+
+    def get_permission(self, actor: str) -> str:
+        return self.permissions.get(actor, self.permission)
+
+    def get_review_decision(self, number: int) -> str:
+        return self.review_decision
+
+    def list_reviews(self, number: int) -> list[dict[str, Any]]:
+        return self.reviews
+
+    def list_reactions(self, comment_id: int) -> list[dict[str, Any]]:
+        return []
+
+    def add_reaction(self, comment_id: int, content: str) -> None:
+        self.reactions.append(content)
+
+    def add_comment(self, issue_number: int, body: str) -> None:
+        self.comments.append(body)
+
+
+class FakeBuildkite:
+    def __init__(
+        self,
+        build_lists: list[list[dict[str, Any]]] | None = None,
+    ) -> None:
+        self.build_lists = build_lists or []
+        self.created_builds: list[dict[str, Any]] = []
+        self.list_calls: list[tuple[str, tuple[str, str] | None]] = []
+        self.retry_calls: list[tuple[int, str]] = []
+
+    def list_builds(
+        self,
+        commit: str,
+        *,
+        metadata: tuple[str, str] | None = None,
+    ) -> list[dict[str, Any]]:
+        self.list_calls.append((commit, metadata))
+        return self.build_lists.pop(0)
+
+    def create_build(self, body: dict[str, Any]) -> dict[str, Any]:
+        self.created_builds.append(body)
+        return {
+            "number": 123,
+            "web_url": "https://buildkite.example/builds/123",
+        }
+
+    def retry_failed_jobs(
+        self,
+        build_number: int,
+        states: str,
+    ) -> dict[str, Any]:
+        self.retry_calls.append((build_number, states))
+        return {"retried_jobs_count": 3}
+
+
+class FakeTransport:
+    def __init__(self, response: Any) -> None:
+        self.calls: list[dict[str, Any]] = []
+        self.response = response
+
+    def request(self, url: str, **kwargs: Any) -> Any:
+        self.calls.append({"url": url, **kwargs})
+        return self.response
+
+
+class RunCiCommandTest(unittest.TestCase):
+    def test_only_exact_ci_commands_are_accepted(self) -> None:
+        self.assertEqual(parse_command(COMMAND_RUN_CI), COMMAND_RUN_CI)
+        self.assertEqual(
+            parse_command(COMMAND_RETRY_FAILED),
+            COMMAND_RETRY_FAILED,
+        )
+        self.assertIsNone(parse_command("/ci run please"))
+        self.assertIsNone(parse_command(" /ci run"))
+
+    def test_write_access_authorizes_reviewers_and_authors(self) -> None:
+        allowed, _ = authorize(
+            actor="reviewer",
+            permission="write",
+            pr=make_pr(),
+        )
+        self.assertTrue(allowed)
+
+    def test_configured_trusted_contributors_can_run_ci(self) -> None:
+        trusted_users = parse_trusted_users("trusted-one, TRUSTED-TWO")
+        allowed, _ = authorize(
+            actor="trusted-two",
+            permission="read",
+            pr=make_pr(),
+            trusted_users=trusted_users,
+        )
+        self.assertTrue(allowed)
+
+    def test_authors_need_an_approval_or_ready_label(self) -> None:
+        pending, _ = authorize(
+            actor="author",
+            permission="read",
+            pr=make_pr(),
+        )
+        approved, _ = authorize(
+            actor="author",
+            permission="read",
+            pr=make_pr(),
+            trusted_approval=True,
+        )
+        ready, _ = authorize(
+            actor="author",
+            permission="read",
+            pr=make_pr(labels=[{"name": "ready"}]),
+        )
+        self.assertFalse(pending)
+        self.assertTrue(approved)
+        self.assertTrue(ready)
+
+    def test_non_author_contributors_without_write_are_denied(self) -> None:
+        allowed, _ = authorize(
+            actor="contributor",
+            permission="read",
+            pr=make_pr(),
+            trusted_approval=True,
+        )
+        self.assertFalse(allowed)
+
+    def test_authors_cannot_use_ready_state_on_draft_prs(self) -> None:
+        allowed, _ = authorize(
+            actor="author",
+            permission="read",
+            pr=make_pr(draft=True, labels=[{"name": "ready"}]),
+            trusted_approval=True,
+        )
+        self.assertFalse(allowed)
+
+    def test_only_trusted_reviewers_can_delegate_through_approval(self) -> None:
+        approved_review = {
+            "state": "APPROVED",
+            "user": {"login": "reviewer"},
+        }
+        trusted = FakeGitHub(
+            permission="read",
+            permissions={"reviewer": "write"},
+            review_decision="APPROVED",
+            reviews=[approved_review],
+        )
+        untrusted = FakeGitHub(
+            permission="read",
+            review_decision="APPROVED",
+            reviews=[approved_review],
+        )
+        self.assertTrue(has_trusted_approval(trusted, 42, set()))
+        self.assertFalse(has_trusted_approval(untrusted, 42, set()))
+
+    def test_build_matching_is_scoped_to_the_pr(self) -> None:
+        self.assertTrue(is_build_for_pr({"pull_request": {"id": 42}}, 42))
+        self.assertFalse(is_build_for_pr({"pull_request": {"id": 43}}, 42))
+        self.assertTrue(
+            is_build_for_pr(
+                {"meta_data": {"github-pr-number": "42"}},
+                42,
+            )
+        )
+
+    def test_latest_build_selection_ignores_other_prs(self) -> None:
+        latest = select_latest_build(
+            [
+                {
+                    "created_at": "2026-07-28T02:00:00Z",
+                    "number": 3,
+                    "pull_request": {"id": 43},
+                },
+                {
+                    "created_at": "2026-07-28T01:00:00Z",
+                    "number": 2,
+                    "pull_request": {"id": 42},
+                },
+                {
+                    "created_at": "2026-07-28T00:00:00Z",
+                    "number": 1,
+                    "pull_request": {"id": 42},
+                },
+            ],
+            42,
+        )
+        self.assertEqual(latest["number"], 2)
+
+    def test_active_build_states_prevent_duplicate_runs(self) -> None:
+        self.assertTrue(is_active_build({"state": "scheduled"}))
+        self.assertTrue(is_active_build({"state": "running"}))
+        self.assertTrue(is_active_build({"state": "waiting"}))
+        self.assertTrue(is_active_build({"blocked": True, "state": "passed"}))
+        self.assertFalse(is_active_build({"state": "failed"}))
+
+    def test_build_payload_preserves_pr_context(self) -> None:
+        payload = create_build_payload(
+            actor="reviewer",
+            comment_id=99,
+            pr=make_pr(labels=[{"name": "ready"}, {"name": "v1"}]),
+        )
+        self.assertEqual(
+            payload,
+            {
+                "commit": "0123456789abcdef",
+                "branch": "feature",
+                "message": "PR #42 /ci run by @reviewer",
+                "pull_request_id": 42,
+                "pull_request_base_branch": "main",
+                "pull_request_repository": ("https://github.com/contributor/vllm.git"),
+                "pull_request_labels": ["ready", "v1"],
+                "ignore_pipeline_branch_filters": True,
+                "env": {
+                    "VLLM_CI_GITHUB_COMMENT_ID": "99",
+                    "VLLM_CI_TRIGGERED_BY": "reviewer",
+                },
+                "meta_data": {
+                    "github-comment-id": "99",
+                    "github-pr-number": "42",
+                    "github-triggered-by": "reviewer",
+                },
+            },
+        )
+
+    def test_ci_run_dispatches_build_with_current_pr_metadata(self) -> None:
+        github = FakeGitHub()
+        buildkite = FakeBuildkite([[], []])
+        run(make_event(COMMAND_RUN_CI), github, buildkite)
+
+        self.assertEqual(len(buildkite.created_builds), 1)
+        self.assertEqual(
+            buildkite.created_builds[0]["message"],
+            "PR #42 /ci run by @reviewer",
+        )
+        self.assertEqual(github.reactions, ["eyes", "rocket"])
+        self.assertIn("Buildkite CI #123", github.comments[0])
+
+    def test_unapproved_authors_are_denied_without_buildkite(self) -> None:
+        github = FakeGitHub(
+            permission="read",
+            pr=make_pr(),
+            review_decision="REVIEW_REQUIRED",
+        )
+        buildkite = FakeBuildkite()
+        run(make_event(COMMAND_RUN_CI, "author"), github, buildkite)
+
+        self.assertEqual(buildkite.list_calls, [])
+        self.assertEqual(github.reactions, ["eyes", "-1"])
+        self.assertIn("approve the PR", github.comments[0])
+
+    def test_ci_retry_uses_latest_current_sha_build(self) -> None:
+        github = FakeGitHub(
+            permission="read",
+            pr=make_pr(labels=[{"name": "ready"}]),
+        )
+        buildkite = FakeBuildkite(
+            [
+                [
+                    {
+                        "created_at": "2026-07-28T01:00:00Z",
+                        "finished_at": "2026-07-28T02:00:00Z",
+                        "number": 123,
+                        "pull_request": {"id": 42},
+                        "state": "failed",
+                        "web_url": "https://buildkite.example/builds/123",
+                    }
+                ]
+            ]
+        )
+        run(make_event(COMMAND_RETRY_FAILED, "author"), github, buildkite)
+
+        self.assertEqual(buildkite.retry_calls, [(123, RETRY_STATES)])
+        self.assertIn("Queued 3 failed job", github.comments[0])
+
+    def test_buildkite_retry_uses_retry_failed_jobs_endpoint(self) -> None:
+        transport = FakeTransport({"retried_jobs_count": 2})
+        client = BuildkiteClient(
+            "secret",
+            "vllm",
+            "ci",
+            transport=transport,
+        )
+        client.retry_failed_jobs(123, RETRY_STATES)
+
+        call = transport.calls[0]
+        self.assertEqual(call["method"], "PUT")
+        self.assertTrue(call["url"].endswith("/123/retry_failed_jobs"))
+        self.assertEqual(call["body"], {"states": RETRY_STATES})
+
+
+if __name__ == "__main__":
+    unittest.main()
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 29263a68491..3790db92631 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -219,10 +219,8 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   # the set of architectures we want to compile for and remove the from the
   # CMAKE_CUDA_FLAGS so that they are not applied globally.
   #
-  # `+PTX` in TORCH_CUDA_ARCH_LIST is not preserved here. It is emitted by torch
-  # as `code=compute_*`, while extract_unique_cuda_archs_ascending() records only
-  # `arch=compute_*`. If a kernel really needs PTX, add `+PTX` to that kernel's
-  # component-specific arch list below.
+  # `+PTX` in TORCH_CUDA_ARCH_LIST is not preserved here. If a kernel really
+  # needs PTX, add `+PTX` to that kernel's component-specific arch list below.
   #
   clear_cuda_arches(CUDA_ARCH_FLAGS)
   extract_unique_cuda_archs_ascending(CUDA_ARCHS "${CUDA_ARCH_FLAGS}")
@@ -232,6 +230,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   cuda_archs_loose_intersection(CUDA_ARCHS
     "${CUDA_SUPPORTED_ARCHS}" "${CUDA_ARCHS}")
   message(STATUS "CUDA supported target architectures: ${CUDA_ARCHS}")
+  if(NOT CUDA_ARCHS)
+    message(FATAL_ERROR
+      "No supported CUDA architectures; the build would produce a binary "
+      "with no usable kernels. Detected gencode flags: ${CUDA_ARCH_FLAGS}; "
+      "supported: ${CUDA_SUPPORTED_ARCHS}. "
+      "Set TORCH_CUDA_ARCH_LIST for your GPU (e.g. 12.0).")
+  endif()
 else()
   #
   # For other GPU targets override the GPU architectures detected by cmake/torch
diff --git a/cmake/utils.cmake b/cmake/utils.cmake
index 14a94eebb22..bbae89c1f57 100644
--- a/cmake/utils.cmake
+++ b/cmake/utils.cmake
@@ -241,14 +241,15 @@ endmacro()
 # `<major>.<minor>`, dedupes them and then sorts them in ascending order and 
 # stores them in `OUT_ARCHES`.
 #
-# Example:
-#   CUDA_ARCH_FLAGS="-gencode arch=compute_75,code=sm_75;...;-gencode arch=compute_90a,code=sm_90a" 
-#   extract_unique_cuda_archs_ascending(OUT_ARCHES CUDA_ARCH_FLAGS)
-#   OUT_ARCHES="7.5;...;9.0"
+# Prefer `code=sm_*`; fall back to `arch=compute_*` for PTX-only flags.
+# This handles mismatches such as `arch=compute_20,code=sm_121`.
 function(extract_unique_cuda_archs_ascending OUT_ARCHES CUDA_ARCH_FLAGS)
   set(_CUDA_ARCHES)
   foreach(_ARCH ${CUDA_ARCH_FLAGS})
-    string(REGEX MATCH "arch=compute_\([0-9]+[af]?\)" _COMPUTE ${_ARCH})
+    string(REGEX MATCH "code=sm_\([0-9]+[af]?\)" _COMPUTE ${_ARCH})
+    if (NOT _COMPUTE)
+      string(REGEX MATCH "arch=compute_\([0-9]+[af]?\)" _COMPUTE ${_ARCH})
+    endif()
     if (_COMPUTE)
       set(_COMPUTE ${CMAKE_MATCH_1})
     endif()
diff --git a/csrc/cpu/cpu_types_vxe.hpp b/csrc/cpu/cpu_types_vxe.hpp
index bf96554a8df..26a14e8ea6b 100644
--- a/csrc/cpu/cpu_types_vxe.hpp
+++ b/csrc/cpu/cpu_types_vxe.hpp
@@ -269,7 +269,7 @@ struct FP32Vec4 : public Vec<FP32Vec4> {
 
   explicit FP32Vec4(__vector float data) : reg(data) {}
 
-  explicit FP32Vec4(const FP32Vec4& data) : reg(data.reg) {}
+  FP32Vec4(const FP32Vec4& data) : reg(data.reg) {}
 };
 
 struct FP32Vec8 : public Vec<FP32Vec8> {
@@ -298,7 +298,7 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
 
   explicit FP32Vec8(f32x4x2_t data) : reg(data) {}
 
-  explicit FP32Vec8(const FP32Vec8& data) {
+  FP32Vec8(const FP32Vec8& data) {
     reg.val[0] = data.reg.val[0];
     reg.val[1] = data.reg.val[1];
   }
@@ -643,7 +643,7 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
 
   explicit FP32Vec16(f32x4x4_t data) : reg(data) {}
 
-  explicit FP32Vec16(const FP32Vec16& data) {
+  FP32Vec16(const FP32Vec16& data) {
     reg.val[0] = data.reg.val[0];
     reg.val[1] = data.reg.val[1];
     reg.val[2] = data.reg.val[2];
diff --git a/csrc/libtorch_stable/layernorm_kernels.cu b/csrc/libtorch_stable/layernorm_kernels.cu
index 878b44df936..7a1051d2c00 100644
--- a/csrc/libtorch_stable/layernorm_kernels.cu
+++ b/csrc/libtorch_stable/layernorm_kernels.cu
@@ -249,7 +249,9 @@ void rms_norm(torch::stable::Tensor& out,    // [..., hidden_size]
   int64_t input_shape_d3 = (num_dims >= 4) ? input.size(-3) : 0;
 
   // For large num_tokens, use smaller blocks to increase SM concurrency.
-  const int max_block_size = (num_tokens < 256) ? 1024 : 256;
+  const bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
+  const int max_block_size =
+      batch_invariant_launch ? 1024 : ((num_tokens < 256) ? 1024 : 256);
   dim3 grid(num_tokens);
   const torch::stable::accelerator::DeviceGuard device_guard(
       input.get_device_index());
@@ -325,8 +327,13 @@ void fused_add_rms_norm(torch::stable::Tensor& input,     // [..., hidden_size]
   /* This kernel is memory-latency bound in many scenarios.
      When num_tokens is large, a smaller block size allows
      for increased block occupancy on CUs and better latency
-     hiding on global mem ops. */
-  const int max_block_size = (num_tokens < 256) ? 1024 : 256;
+     hiding on global mem ops. In batch-invariant mode the block size must
+     not depend on num_tokens, otherwise the same token would use a different
+     reduction width (and thus a different floating-point summation order)
+     across batches; lock it to 1024 to keep results bit-exact. */
+  const bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
+  const int max_block_size =
+      batch_invariant_launch ? 1024 : ((num_tokens < 256) ? 1024 : 256);
   dim3 block(std::min(hidden_size, max_block_size));
   const torch::stable::accelerator::DeviceGuard device_guard(
       input.get_device_index());
@@ -337,7 +344,6 @@ void fused_add_rms_norm(torch::stable::Tensor& input,     // [..., hidden_size]
   auto res_ptr = reinterpret_cast<std::uintptr_t>(residual.data_ptr());
   bool offsets_are_multiple_of_vector_width =
       hidden_size % vector_width == 0 && input_stride % vector_width == 0;
-  bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
   const bool has_weight = weight.has_value();
   if (has_weight) {
     auto wt_ptr = reinterpret_cast<std::uintptr_t>(weight->data_ptr());
diff --git a/csrc/libtorch_stable/layernorm_quant_kernels.cu b/csrc/libtorch_stable/layernorm_quant_kernels.cu
index f3bf8882e77..f43be531de0 100644
--- a/csrc/libtorch_stable/layernorm_quant_kernels.cu
+++ b/csrc/libtorch_stable/layernorm_quant_kernels.cu
@@ -215,7 +215,9 @@ void rms_norm_static_fp8_quant(
   int num_tokens = input.numel() / hidden_size;
 
   // For large num_tokens, use smaller blocks to increase SM concurrency.
-  const int max_block_size = (num_tokens < 256) ? 1024 : 256;
+  const bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
+  const int max_block_size =
+      batch_invariant_launch ? 1024 : ((num_tokens < 256) ? 1024 : 256);
   dim3 grid(num_tokens);
   const torch::stable::accelerator::DeviceGuard device_guard(
       input.get_device_index());
@@ -279,7 +281,9 @@ void fused_add_rms_norm_static_fp8_quant(
      When num_tokens is large, a smaller block size allows
      for increased block occupancy on CUs and better latency
      hiding on global mem ops. */
-  const int max_block_size = (num_tokens < 256) ? 1024 : 256;
+  const bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
+  const int max_block_size =
+      batch_invariant_launch ? 1024 : ((num_tokens < 256) ? 1024 : 256);
   dim3 block(std::min(hidden_size, max_block_size));
   const torch::stable::accelerator::DeviceGuard device_guard(
       input.get_device_index());
@@ -296,7 +300,6 @@ void fused_add_rms_norm_static_fp8_quant(
   auto wt_ptr = reinterpret_cast<std::uintptr_t>(weight.data_ptr());
   bool ptrs_are_aligned =
       inp_ptr % 16 == 0 && res_ptr % 16 == 0 && wt_ptr % 16 == 0;
-  bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
   if (ptrs_are_aligned && hidden_size % 8 == 0 && input_stride % 8 == 0 &&
       !batch_invariant_launch) {
     LAUNCH_FUSED_ADD_RMS_NORM(8);
diff --git a/csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu b/csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu
index 2152e64dc96..56dd4703872 100644
--- a/csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu
+++ b/csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu
@@ -2,6 +2,7 @@
 #include "../../torch_utils.h"
 
 #include "../../dispatch_utils.h"
+#include "../../../core/batch_invariant.hpp"
 #include "layernorm_utils.cuh"
 #include "quant_conversions.cuh"
 
@@ -231,7 +232,9 @@ void rms_norm_per_block_quant_dispatch(
   auto num_tokens = input.numel() / hidden_size;
 
   dim3 grid(num_tokens);
-  const int max_block_size = (num_tokens <= 256) ? 512 : 256;
+  const bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
+  const int max_block_size =
+      batch_invariant_launch ? 512 : ((num_tokens <= 256) ? 512 : 256);
   dim3 block(std::min(hidden_size, max_block_size));
   const torch::stable::accelerator::DeviceGuard device_guard(
       input.get_device_index());
diff --git a/csrc/libtorch_stable/torch_bindings.cpp b/csrc/libtorch_stable/torch_bindings.cpp
index fc0322c67b4..fec337a8953 100644
--- a/csrc/libtorch_stable/torch_bindings.cpp
+++ b/csrc/libtorch_stable/torch_bindings.cpp
@@ -591,7 +591,7 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "limit=7.0) "
       "-> ()");
 
-  // Kimi SITU (SituGLU) gated activation. linear_beta<=0 means unset.
+  // SituGLU implementation used in Kimi models.
   ops.def(
       "situ_and_mul(Tensor! out, Tensor input, float beta=1.0, float "
       "linear_beta=-1.0) -> ()");
diff --git a/csrc/spinloop.cpp b/csrc/spinloop.cpp
index c29e48a5f0e..3285b9a3ded 100644
--- a/csrc/spinloop.cpp
+++ b/csrc/spinloop.cpp
@@ -7,7 +7,7 @@ extern "C" {
 
 #if defined(__i386__) || defined(__x86_64__)
   #include <cpuid.h>
-  #include <mwaitxintrin.h>
+  #include <x86intrin.h>
 #endif
 
 #if defined(CLOCK_MONOTONIC_RAW)
diff --git a/docker/Dockerfile.s390x b/docker/Dockerfile.s390x
index b71e035e152..d645e37e78c 100644
--- a/docker/Dockerfile.s390x
+++ b/docker/Dockerfile.s390x
@@ -61,13 +61,13 @@ ENV C_INCLUDE_PATH="/usr/local/include:$C_INCLUDE_PATH"
 
 FROM python-install AS torch-vision
 # Install torchvision
-ARG TORCH_VISION_VERSION=v0.26.0
+ARG TORCH_VISION_VERSION=v0.28.0
 WORKDIR /tmp
 RUN --mount=type=cache,target=/root/.cache/uv \
     git clone https://github.com/pytorch/vision.git && \
     cd vision && \
     git checkout $TORCH_VISION_VERSION && \
-    uv pip install torch==2.11.0 --index-url https://download.pytorch.org/whl/cpu && \
+    uv pip install torch==2.13.0 --index-url https://download.pytorch.org/whl/cpu && \
     python setup.py bdist_wheel
 
 FROM python-install AS hf-xet-builder
diff --git a/docs/contributing/README.md b/docs/contributing/README.md
index 34dc385db78..89acc6b7f5a 100644
--- a/docs/contributing/README.md
+++ b/docs/contributing/README.md
@@ -301,8 +301,10 @@ review process:
   isn't clear or you disagree with a suggestion, feel free to ask for
   clarification or discuss the suggestion.
 - Note that not all CI checks will be executed due to limited computational
-  resources. The reviewer will add `ready` label to the PR when the PR is
-  ready to merge or a full CI run is needed.
+  resources. Reviewers with write access and configured trusted contributors
+  can comment `/ci run` when CI signals are needed before a PR is ready. After
+  the PR is approved or has the `ready` label, the PR author can use `/ci run`
+  or `/ci retry`. New commits do not start CI automatically.
 
 ### Pull Request Limits and Escalation
 
diff --git a/docs/deployment/integrations/llm-d.md b/docs/deployment/integrations/llm-d.md
index 6060b98f642..7d261eb910b 100644
--- a/docs/deployment/integrations/llm-d.md
+++ b/docs/deployment/integrations/llm-d.md
@@ -1,5 +1,37 @@
 # llm-d
 
-vLLM can be deployed with [llm-d](https://github.com/llm-d/llm-d), a Kubernetes-native distributed inference serving stack providing well-lit paths for anyone to serve large generative AI models at scale. It helps achieve the fastest "time to state-of-the-art (SOTA) performance" for key OSS models across most hardware accelerators and infrastructure providers.
+[llm-d](https://llm-d.ai/) is a Kubernetes-native distributed inference framework for serving large language models at scale, with vLLM as its primary inference engine. llm-d coordinates a fleet of vLLM instances across a cluster so that performance holds up under real production traffic, achieving the fastest "time to state-of-the-art (SOTA) performance" for key OSS models across most hardware accelerators.
 
-You can use vLLM with llm-d directly by following [the official guides](https://llm-d.ai/docs/guides) or via [KServe's LLMInferenceService](https://kserve.github.io/website/docs/model-serving/generative-inference/llmisvc/llmisvc-overview).
+It is a [CNCF Sandbox project](https://www.cncf.io/blog/2026/03/24/welcome-llm-d-to-the-cncf-evolving-kubernetes-into-sota-ai-infrastructure/) founded by Red Hat, Google Cloud, IBM Research, CoreWeave, and NVIDIA.
+
+## What llm-d adds to vLLM
+
+A single vLLM server is fast, but at scale the picture changes: across many replicas, cache locality breaks under round-robin load balancing, long prompts inflate time-to-first-token, and accelerators sit underused. llm-d adds the cluster-level layer that vLLM does not aim to provide on its own:
+
+- **[Prefix-aware routing](https://llm-d.ai/docs/guides/precise-prefix-cache-aware).** Instead of round-robin, llm-d reads vLLM's KV-cache events and routes each request to the replica that already holds its prefix, reusing cache instead of recomputing it.
+- **[Distributed KV-cache management](https://llm-d.ai/docs/guides#advanced-kv-cache-management).** A global index tracks which token blocks live on which replica, and [tiered offloading](https://llm-d.ai/docs/guides/tiered-prefix-cache) spills cache to CPU memory or local SSD, extending the working set beyond accelerator HBM.
+- **[Prefill/decode disaggregation](https://llm-d.ai/docs/guides/pd-disaggregation).** Prompt processing and token generation run on separate vLLM workers, with KV-cache moved over the vLLM [NIXL connector](https://docs.vllm.ai/en/latest/features/nixl_connector_usage/), lowering TTFT and steadying per-token latency on long prompts.
+- **[Wide expert-parallelism](https://llm-d.ai/docs/guides/wide-expert-parallelism).** Serve large Mixture-of-Experts models such as DeepSeek-R1 and GPT-OSS across nodes with combined data and expert parallelism, for more KV-cache capacity and throughput.
+- **SLO-aware [autoscaling](https://llm-d.ai/docs/guides/workload-autoscaling) and [flow control](https://llm-d.ai/docs/guides/flow-control).** Scale vLLM pools on real inference signals (queue depth, true demand) rather than raw GPU utilization, with multi-tenant fairness and priority dispatch.
+
+These are composable. Most teams start by adding prefix-aware routing over an existing vLLM pool, then layer in the rest as specific bottlenecks appear.
+
+## Performance
+
+Representative benchmarked results across accelerators:
+
+- **3x higher output throughput** and **2x faster TTFT** from prefix-aware routing vs round-robin (Llama 3.1 70B, AMD MI300X)
+- **Up to 70% higher tokens/sec** from prefill/decode disaggregation (GPT-OSS, NVIDIA B200)
+- **13.9x throughput** from hierarchical KV offloading at high concurrency vs GPU-only (NVIDIA H100)
+
+See the [full list](https://github.com/llm-d/llm-d#performance-highlights) and reproducible benchmarks on [Prism](https://prism.llm-d.ai/).
+
+## Get started
+
+1. Deploy the [Optimized Baseline](https://llm-d.ai/docs/guides/optimized-baseline) with the [Quickstart](https://llm-d.ai/docs/getting-started/quickstart). It stands up an intelligent router over a vLLM pool on Kubernetes in a tested configuration.
+2. Browse the [well-lit path guides](https://llm-d.ai/docs/guides), each a tested recipe for one of the capabilities above, and add the optimization that fits your workload.
+3. Read the [Introduction](https://llm-d.ai/docs/getting-started) and [Architecture overview](https://llm-d.ai/docs/architecture) to see how the pieces wrap your vLLM deployment.
+
+You can also deploy vLLM with llm-d via [KServe's LLMInferenceService](https://kserve.github.io/website/docs/model-serving/generative-inference/llmisvc/llmisvc-overview).
+
+Questions and contributions are welcome on [GitHub](https://github.com/llm-d/llm-d) and [Slack](https://llm-d.ai/slack).
diff --git a/docs/features/kv_offloading_usage.md b/docs/features/kv_offloading_usage.md
index 13dbd5299d3..72f838c7d2c 100644
--- a/docs/features/kv_offloading_usage.md
+++ b/docs/features/kv_offloading_usage.md
@@ -70,7 +70,8 @@ vllm serve <model> \
 | `cpu_bytes_to_use` | yes | — | both | Total bytes of host memory reserved for the CPU tier across all workers (not per-worker). |
 | `block_size` | no | GPU block size | both | Offloaded block size in tokens; must be a multiple of the GPU block size. Mutually exclusive with `blocks_per_chunk`. |
 | `blocks_per_chunk` | no | `1` | both | Offloaded chunk size in GPU blocks; must be > 0. Alternative to `block_size` for models whose KV cache groups have different block sizes. |
-| `eviction_policy` | no | `lru` | both | Primary tier policy: `lru` or `arc`. |
+| `eviction_policy` | no | `lru` | both | Primary tier policy: built-in `lru`/`arc`, or a custom `CachePolicy` name (see [Custom Eviction Policies](#custom-eviction-policies)). |
+| `cache_policy_module_path` | no | — | both | Python import path for a custom `CachePolicy` not in the built-in registry. Required only when `eviction_policy` is not built-in and wasn't pre-registered via `CachePolicyFactory` (advanced). |
 | `store_threshold` | no | `0` | single-tier | Min lookups before a block is offloaded. Values ≥ 2 are rejected by `TieringOffloadingSpec`. |
 | `max_tracker_size` | no | `64000` | single-tier | Max entries in the lookup tracker. |
 | `secondary_tiers` | no | `[]` | multi-tier | List of secondary tier configs (see below). |
@@ -78,6 +79,36 @@ vllm serve <model> \
 | `self_describing_kv_events` | no | `false` | both | Opt-in. When `true` *and* KV cache events are enabled (`--kv-events-config` with `enable_kv_cache_events`), the connector emits self-describing block-granular `BlockStored`/`BlockRemoved` payloads (constituent block hashes, whole-chunk `token_ids`, per-block `block_size`, parent hash, LoRA + group/cache-spec metadata) instead of the placeholder fallback, so external KV-event consumers can index offloaded blocks. Inert unless events are enabled. With `TieringOffloadingSpec`, a CPU promotion is self-describing when a local request observes its primary-tier `HIT` before event translation; otherwise its stored event may retain the placeholder, while a later `HIT` can backfill metadata for removal. Pending-removal/re-promotion races and externally initiated promotions may also produce placeholders, and consumers must ignore removals for unknown hashes. Full-attention groups only; sliding-window/SSM groups keep the placeholder fallback. In chunk mode (`block_size` > GPU block size, or `blocks_per_chunk` > 1), overlapping chunks re-announce shared per-block hashes, so consumers must reference-count (deduplicate) repeated store/remove announcements. |
 | `spec_module_path` | no | — | both | Python import path for a custom `OffloadingSpec` not in the built-in registry. Required only when `spec_name` is not built-in (advanced). |
 
+## Custom Eviction Policies
+
+`eviction_policy` resolves through `CachePolicyFactory` (`vllm/v1/kv_offload/cpu/policies/factory.py`), which pre-registers the built-in `lru` and `arc` policies.
+
+### Out-of-tree (recommended)
+
+Implement `CachePolicy` (`vllm/v1/kv_offload/cpu/policies/base.py`) in your own package — no vLLM fork or patch required — and point `kv_connector_extra_config` at it directly:
+
+```json
+{
+  "cpu_bytes_to_use": 10737418240,
+  "eviction_policy": "MyCachePolicy",
+  "cache_policy_module_path": "my_package.my_module"
+}
+```
+
+`eviction_policy` is checked against the built-in registry first; if it isn't a registered name, vLLM imports `cache_policy_module_path` and looks up `eviction_policy` as a class name in that module — the same fallback `spec_module_path` provides for a custom `OffloadingSpec`. No import or registration call needs to run before the server starts.
+
+### Registering a friendly short name (in-process only)
+
+If you control the process that constructs the vLLM engine (e.g. an embedding application), you can register a short name once at startup instead of repeating the module path in every config:
+
+```python
+from vllm.v1.kv_offload.cpu.policies.factory import CachePolicyFactory
+
+CachePolicyFactory.register_cache_policy("my_policy", "my_package.my_module", "MyCachePolicy")
+```
+
+Then set `"eviction_policy": "my_policy"` in `kv_connector_extra_config`, the same as `"lru"`/`"arc"`. This only takes effect within the process that ran the `register_cache_policy` call — it does not help when the server is launched as a separate process (e.g. via the `vllm serve` CLI), where the out-of-tree `cache_policy_module_path` config above is the only option.
+
 ## Secondary Tiers
 
 Each entry in `secondary_tiers` is a dict with a required `type` field plus tier-specific fields.
diff --git a/docs/serving/offline_inference.md b/docs/serving/offline_inference.md
index 4512f4a0720..9a71612f262 100644
--- a/docs/serving/offline_inference.md
+++ b/docs/serving/offline_inference.md
@@ -65,6 +65,8 @@ For further details on Weight Transfer, please refer to [this page](../training/
 - `LLM.start_weight_update` - Starts a new weight update cycle.
 - `LLM.update_weights` - Updates the model weights.
 - `LLM.finish_weight_update` - Finishes the current weight update cycle.
+- `LLM.update_weight_version` - Sets the weight version without updating model weights.
+- `LLM.get_weight_version` - Returns the latest committed weight version.
 
 ## Additional APIs
 
diff --git a/docs/serving/online_serving/README.md b/docs/serving/online_serving/README.md
index f7914bc0582..90ff7a3e3d8 100644
--- a/docs/serving/online_serving/README.md
+++ b/docs/serving/online_serving/README.md
@@ -179,6 +179,8 @@ For further details on Weight Transfer, please refer to [this page](../../traini
 - `/start_weight_update` - Prepares the inference engine for a weight update.
 - `/update_weights` - Update model weights (can alter model behavior)
 - `/finish_weight_update` - Finalizes the weight update
+- `/update_weight_version` - Set the weight version without updating model weights
+- `/weight_info` - Get the latest committed weight version
 - `/get_world_size` - Get distributed world size
 
 ### Collective RPC
diff --git a/docs/training/async_rl.md b/docs/training/async_rl.md
index e655f9c39ff..9e75a24eaa1 100644
--- a/docs/training/async_rl.md
+++ b/docs/training/async_rl.md
@@ -38,11 +38,12 @@ Resumes the scheduler after a pause. Any requests frozen with `mode="keep"` will
 
 ### HTTP Endpoints
 
-When using the vLLM HTTP server, the same functionality is available via:
+With `VLLM_SERVER_DEV_MODE=1`, the vLLM HTTP server exposes the same functionality via:
 
 - `POST /pause?mode=keep` - Pause generation
 - `POST /resume` - Resume generation
 - `POST /abort_requests` - Abort in-flight requests without pausing the scheduler (send `{}` to abort all, or `{"request_ids": [...]}`)
+- `GET /weight_info` - Return the latest committed `weight_version`
 
 !!! note "Data Parallelism"
     When using data parallelism with vLLM's **internal load balancer** (i.e. `data_parallel_backend="ray"`), pause and resume are handled automatically across all DP ranks -- a single call is sufficient. When using an **external load balancer** (i.e. multiple independent vLLM instances behind a proxy), you must send pause and resume requests to **every** engine instance individually before and after the weight update.
diff --git a/docs/training/weight_transfer/README.md b/docs/training/weight_transfer/README.md
index 7579e5fd4d0..b8d39763181 100644
--- a/docs/training/weight_transfer/README.md
+++ b/docs/training/weight_transfer/README.md
@@ -53,7 +53,9 @@ When running vLLM as an HTTP server, the following endpoints are available for w
 | `/init_weight_transfer_engine` | POST | Initialize the weight transfer engine with backend-specific info |
 | `/start_weight_update` | POST | Start a weight update |
 | `/update_weights` | POST | Transfer a batch of weights with backend-specific metadata |
-| `/finish_weight_update` | POST | Finish the weight update and run post-processing |
+| `/finish_weight_update` | POST | Finish the update and optionally commit its `weight_version` |
+| `/update_weight_version` | POST | Update `weight_version` without changing model weights |
+| `/weight_info` | GET | Get the latest committed weight version |
 | `/pause` | POST | Pause generation before weight sync to handle inflight requests |
 | `/resume` | POST | Resume generation after weight sync |
 | `/get_world_size` | GET | Get the number of inference workers (useful for NCCL world size calculation) |
@@ -79,7 +81,7 @@ EngineClass.trainer_send_weights(
 )
 
 # 4. Finish weight update on inference side
-llm.finish_weight_update()
+llm.finish_weight_update(weight_version="step-42")
 ```
 
 See the [NCCL](nccl.md) and [IPC](ipc.md) pages for backend-specific trainer APIs and full examples.
diff --git a/pyproject.toml b/pyproject.toml
index 5ece92136a1..0766645fc74 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -129,10 +129,9 @@ extend-exclude = ["tests/models/fixtures/*", "tests/prompts/*", "tests/tokenizer
     "tests/entrypoints/speech_to_text/transcription/test_transcription_validation.py",
     "docs/governance/process.md", "docs/assets/contributing/vllm_bench_serve_timeline.html", 
     "tests/v1/engine/test_fast_incdec_prefix_err.py", ".git/*", "csrc/cpu/sgl-kernels/*",
-    "rust/src/chat/src/renderer/deepseek_v32/fixtures/*",
-    "rust/src/parser/src/tool/gemma4.rs", "rust/src/parser/src/unified/gemma4.rs",
+    "rust/src/chat/src/renderer/deepseek_v32/fixtures/*", "rust/src/parser/**",
     "rust/src/text/src/output/decoded.rs",
-    "rust/src/tokenizer/src/incremental.rs", "rust/src/parser/src/reasoning/tests.rs"]
+    "rust/src/tokenizer/src/incremental.rs"]
 ignore-hidden = false
 
 [tool.typos.default]
diff --git a/requirements/test/cpu.txt b/requirements/test/cpu.txt
index 3cb251a308b..923b54bb14e 100644
--- a/requirements/test/cpu.txt
+++ b/requirements/test/cpu.txt
@@ -1,5 +1,7 @@
 # This file was autogenerated by uv via the following command:
 #    uv pip compile requirements/test/cuda.in -o requirements/test/cpu.txt --index-strategy unsafe-best-match --torch-backend cpu --python-platform x86_64-manylinux_2_28 --python-version 3.12
+abi3info==2025.11.29
+    # via torch-abi-audit
 absl-py==2.1.0
     # via rouge-score
 accelerate==1.13.0
@@ -763,6 +765,8 @@ pycparser==2.22
     # via cffi
 pycryptodomex==3.22.0
     # via blobfile
+pycxxfilt==0.1.0
+    # via torch-abi-audit
 pydantic==2.12.0
     # via
     #   -r requirements/test/../common.txt
@@ -1127,6 +1131,8 @@ torch==2.13.0+cpu
     #   vector-quantize-pytorch
     #   vocos
     #   xgrammar
+torch-abi-audit==0.0.1
+    # via -r requirements/test/cuda.in
 torchaudio==2.11.0+cpu
     # via
     #   -r requirements/test/cuda.in
diff --git a/requirements/test/cuda.in b/requirements/test/cuda.in
index b33257250e7..377224eac36 100644
--- a/requirements/test/cuda.in
+++ b/requirements/test/cuda.in
@@ -45,7 +45,7 @@ schemathesis>=4.0.0 # Required for openai schema test.
 # quantization
 bitsandbytes==0.49.2
 buildkite-test-collector==0.1.9
-
+torch-abi-audit # CI check for PyTorch stable ABI compliance
 
 genai_perf>=0.0.8
 tritonclient>=2.51.0
diff --git a/requirements/test/cuda.txt b/requirements/test/cuda.txt
index 6490502cdda..6c155d89bd2 100644
--- a/requirements/test/cuda.txt
+++ b/requirements/test/cuda.txt
@@ -1,5 +1,7 @@
 # This file was autogenerated by uv via the following command:
 #    uv pip compile requirements/test/cuda.in -c requirements/cuda.txt -o requirements/test/cuda.txt --index-strategy unsafe-best-match --torch-backend cu130 --python-platform x86_64-manylinux_2_28 --python-version 3.12
+abi3info==2025.11.29
+    # via torch-abi-audit
 absl-py==2.1.0
     # via rouge-score
 accelerate==1.13.0
@@ -850,6 +852,8 @@ pycparser==2.22
     # via cffi
 pycryptodomex==3.22.0
     # via blobfile
+pycxxfilt==0.1.0
+    # via torch-abi-audit
 pydantic==2.12.0
     # via
     #   -c requirements/common.txt
@@ -1225,6 +1229,8 @@ torch==2.13.0+cu130
     #   vector-quantize-pytorch
     #   vocos
     #   xgrammar
+torch-abi-audit==0.0.1
+    # via -r requirements/test/cuda.in
 torchaudio==2.11.0+cu130
     # via
     #   -c requirements/cuda.txt
diff --git a/rust/Cargo.lock b/rust/Cargo.lock
index 00c134a33c8..38f1c257756 100644
--- a/rust/Cargo.lock
+++ b/rust/Cargo.lock
@@ -5503,9 +5503,9 @@ dependencies = [
  "tokio",
  "tokio-stream",
  "tracing",
- "tracing-subscriber",
  "url",
  "uuid",
+ "vllm-tracing",
 ]
 
 [[package]]
@@ -5570,17 +5570,16 @@ dependencies = [
  "serde_json",
  "serde_with",
  "thiserror-ext",
- "time",
  "tokio",
  "tokio-util",
  "tracing",
- "tracing-subscriber",
  "uuid",
  "vllm-bench",
  "vllm-chat",
  "vllm-engine-core-client",
  "vllm-managed-engine",
  "vllm-server",
+ "vllm-tracing",
 ]
 
 [[package]]
@@ -5827,6 +5826,15 @@ dependencies = [
  "vllm-parser",
 ]
 
+[[package]]
+name = "vllm-tracing"
+version = "0.1.0"
+dependencies = [
+ "time",
+ "tracing",
+ "tracing-subscriber",
+]
+
 [[package]]
 name = "walkdir"
 version = "2.5.0"
diff --git a/rust/Cargo.toml b/rust/Cargo.toml
index f0277eb79cc..09f55cf07cd 100644
--- a/rust/Cargo.toml
+++ b/rust/Cargo.toml
@@ -13,6 +13,7 @@ members = [
     "src/server",
     "src/text",
     "src/tokenizer",
+    "src/tracing",
 ]
 resolver = "3"
 
@@ -143,6 +144,7 @@ vllm-parser = { path = "src/parser" }
 vllm-server = { path = "src/server" }
 vllm-text = { path = "src/text" }
 vllm-tokenizer = { path = "src/tokenizer" }
+vllm-tracing = { path = "src/tracing" }
 winnow = { version = "1.0.2", features = ["simd"] }
 xgrammar-structural-tag = "0.2.0"
 zeromq = { version = "0.6.0", default-features = false, features = [
diff --git a/rust/src/bench/Cargo.toml b/rust/src/bench/Cargo.toml
index 2da0f13f1ce..960e7a62f7d 100644
--- a/rust/src/bench/Cargo.toml
+++ b/rust/src/bench/Cargo.toml
@@ -32,9 +32,9 @@ tokenizers.workspace = true
 tokio.workspace = true
 tokio-stream.workspace = true
 tracing.workspace = true
-tracing-subscriber.workspace = true
 url.workspace = true
 uuid.workspace = true
+vllm-tracing.workspace = true
 
 [lints]
 workspace = true
diff --git a/rust/src/bench/src/main.rs b/rust/src/bench/src/main.rs
index b9ebfc2a0ac..1764983e448 100644
--- a/rust/src/bench/src/main.rs
+++ b/rust/src/bench/src/main.rs
@@ -19,18 +19,8 @@ struct Cli {
     args: vllm_bench::BenchServeArgs,
 }
 
-// TODO: unify the tracing subscriber used by different binaries.
-fn init_tracing() {
-    let filter = tracing_subscriber::EnvFilter::try_from_default_env()
-        .unwrap_or_else(|_| tracing_subscriber::EnvFilter::new("info"));
-    let _ = tracing_subscriber::fmt()
-        .with_env_filter(filter)
-        .with_writer(std::io::stderr)
-        .try_init();
-}
-
 fn main() -> anyhow::Result<()> {
-    init_tracing();
+    vllm_tracing::init_tracing("Bench");
 
     let cli = Cli::parse();
     vllm_bench::prepare_process();
diff --git a/rust/src/chat/src/backend/hf.rs b/rust/src/chat/src/backend/hf.rs
index ee950d26b73..583a5fb936e 100644
--- a/rust/src/chat/src/backend/hf.rs
+++ b/rust/src/chat/src/backend/hf.rs
@@ -73,7 +73,7 @@ impl HfChatBackend {
             RendererSelection::DeepSeekV4 => Arc::new(DeepSeekV4ChatRenderer::new()),
             RendererSelection::Harmony => Arc::new(HarmonyChatRenderer::new()?),
             RendererSelection::Inkling => Arc::new(InklingChatRenderer::new(tokenizer.clone())?),
-            RendererSelection::KimiK3 => Arc::new(KimiK3ChatRenderer::new()),
+            RendererSelection::KimiK3 => Arc::new(KimiK3ChatRenderer::new(tokenizer.clone())),
         };
 
         info!(
diff --git a/rust/src/chat/src/renderer/inkling/tests.rs b/rust/src/chat/src/renderer/inkling/tests.rs
index 7267c25f031..1c4cd7236db 100644
--- a/rust/src/chat/src/renderer/inkling/tests.rs
+++ b/rust/src/chat/src/renderer/inkling/tests.rs
@@ -31,6 +31,10 @@ impl Tokenizer for FixtureTokenizer {
         Ok(text.bytes().map(u32::from).collect())
     }
 
+    fn encode_ordinary(&self, text: &str) -> vllm_tokenizer::Result<Vec<u32>> {
+        self.encode(text, false)
+    }
+
     fn decode(
         &self,
         token_ids: &[u32],
diff --git a/rust/src/chat/src/renderer/kimi_k3/encoding.rs b/rust/src/chat/src/renderer/kimi_k3/encoding.rs
index 05790514525..189d68c101c 100644
--- a/rust/src/chat/src/renderer/kimi_k3/encoding.rs
+++ b/rust/src/chat/src/renderer/kimi_k3/encoding.rs
@@ -6,9 +6,9 @@
 //! Port of Moonshot remote-code `encoding_k3.py::build_chat_segments()`.
 
 use std::collections::HashMap;
-use std::fmt::Write as _;
 
 use serde_json::{Map, Value, json};
+use vllm_tokenizer::Tokenizer;
 
 use crate::error::{Error, Result};
 use crate::request::{
@@ -25,12 +25,47 @@ pub(super) const IMAGE_PLACEHOLDER: &str = "<|media_pad|>";
 const DEFAULT_THINKING_EFFORT: &str = "max";
 const VALID_THINKING_EFFORTS: &[&str] = &["low", "high", "max"];
 
-/// Render one chat request into the K3 XTML prompt string.
-pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
+/// K3 prompt encoder preserving Python's per-segment tokenization boundaries.
+pub(super) struct K3TokenWriter<'a> {
+    tokenizer: &'a dyn Tokenizer,
+    token_ids: Vec<u32>,
+}
+
+impl<'a> K3TokenWriter<'a> {
+    pub(super) fn new(tokenizer: &'a dyn Tokenizer) -> Self {
+        Self {
+            tokenizer,
+            token_ids: Vec::new(),
+        }
+    }
+
+    /// Encode one trusted segment with normal added-token recognition.
+    pub(super) fn control(&mut self, text: &str) -> Result<()> {
+        if !text.is_empty() {
+            self.token_ids.extend(self.tokenizer.encode(text, false)?);
+        }
+        Ok(())
+    }
+
+    /// Encode one literal segment while bypassing every added-token matcher.
+    pub(super) fn ordinary(&mut self, text: &str) -> Result<()> {
+        if !text.is_empty() {
+            self.token_ids.extend(self.tokenizer.encode_ordinary(text)?);
+        }
+        Ok(())
+    }
+
+    pub(super) fn finish(self) -> Vec<u32> {
+        self.token_ids
+    }
+}
+
+/// Render and tokenize one chat request using K3's segment-aware contract.
+pub(super) fn render_request(request: &ChatRequest, tokenizer: &dyn Tokenizer) -> Result<Vec<u32>> {
     let thinking = thinking_enabled(request)?;
     let thinking_effort = thinking.then(|| thinking_effort(request)).transpose()?;
     let tools = request_tools(request);
-    let mut out = String::new();
+    let mut out = K3TokenWriter::new(tokenizer);
 
     if !tools.is_empty() {
         write_tool_declare(&mut out, tools, false)?;
@@ -48,14 +83,14 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
                  supported values include `low`, `medium`, `high`, and `max`.\n\
                  Now the system is invoked with `thinking_effort={effort}`."
             ),
-        );
+        )?;
     }
 
     // Track prior assistant tool-call ids for tool-result reordering / naming.
     let mut tool_call_id_index: HashMap<String, (usize, String)> = HashMap::new();
     let mut pending_tool_run: Vec<(usize, ChatMessage)> = Vec::new();
 
-    let flush_tool_run = |out: &mut String,
+    let flush_tool_run = |out: &mut K3TokenWriter<'_>,
                           run: &mut Vec<(usize, ChatMessage)>,
                           id_index: &HashMap<String, (usize, String)>|
      -> Result<()> {
@@ -163,7 +198,7 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
                 "tool-choice",
                 "The system is invoked with `tool_choice=required`.\n\
                  You MUST call tools in the next message.",
-            );
+            )?;
         }
         // Emit only when tools are present: Rust defaults tool_choice to None
         // for tool-free requests, which must not inject a tool-choice message.
@@ -173,7 +208,7 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
                 "tool-choice",
                 "The system is invoked with `tool_choice=none`.\n\
                  You MUST NOT call any tools in the next message.",
-            );
+            )?;
         }
         ChatToolChoice::None | ChatToolChoice::Auto | ChatToolChoice::Function { .. } => {}
     }
@@ -181,11 +216,11 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
     write_response_format(&mut out, request)?;
 
     if request.chat_options.add_generation_prompt() {
-        write_open_tag(&mut out, "message", &[("role", "assistant")]);
-        write_open_tag(&mut out, if thinking { "think" } else { "response" }, &[]);
+        write_open_tag(&mut out, "message", &[("role", "assistant")])?;
+        write_open_tag(&mut out, if thinking { "think" } else { "response" }, &[])?;
     }
 
-    Ok(out)
+    Ok(out.finish())
 }
 
 fn request_tools(request: &ChatRequest) -> &[ChatTool] {
@@ -248,7 +283,11 @@ fn content_is_empty(content: &ChatContent) -> bool {
     }
 }
 
-fn write_tool_declare(out: &mut String, tools: &[ChatTool], dynamic: bool) -> Result<()> {
+fn write_tool_declare(
+    out: &mut K3TokenWriter<'_>,
+    tools: &[ChatTool],
+    dynamic: bool,
+) -> Result<()> {
     let mut specs = Vec::with_capacity(tools.len());
     for tool in tools {
         let mut function = Map::new();
@@ -285,23 +324,26 @@ fn write_tool_declare(out: &mut String, tools: &[ChatTool], dynamic: bool) -> Re
         )
     };
 
-    write_internal_system(out, "tool-declare", &body);
-    Ok(())
+    write_internal_system(out, "tool-declare", &body)
 }
 
-fn write_internal_system(out: &mut String, message_type: &str, body: &str) {
+fn write_internal_system(
+    out: &mut K3TokenWriter<'_>,
+    message_type: &str,
+    body: &str,
+) -> Result<()> {
     write_open_tag(
         out,
         "message",
         &[("role", "system"), ("type", message_type)],
-    );
-    out.push_str(body.trim());
-    write_close_tag(out, "message");
-    out.push_str(END_OF_MSG);
+    )?;
+    out.ordinary(body.trim())?;
+    write_close_tag(out, "message")?;
+    out.control(END_OF_MSG)
 }
 
 fn write_role_message(
-    out: &mut String,
+    out: &mut K3TokenWriter<'_>,
     role: &str,
     name: Option<&str>,
     content: &ChatContent,
@@ -311,15 +353,14 @@ fn write_role_message(
         attrs.push(("name", name.to_string()));
     }
     let attr_refs: Vec<(&str, &str)> = attrs.iter().map(|(k, v)| (*k, v.as_str())).collect();
-    write_open_tag(out, "message", &attr_refs);
+    write_open_tag(out, "message", &attr_refs)?;
     write_content(out, content)?;
-    write_close_tag(out, "message");
-    out.push_str(END_OF_MSG);
-    Ok(())
+    write_close_tag(out, "message")?;
+    out.control(END_OF_MSG)
 }
 
 fn write_tool_message(
-    out: &mut String,
+    out: &mut K3TokenWriter<'_>,
     tool_name: &str,
     index: usize,
     content: &ChatContent,
@@ -329,19 +370,18 @@ fn write_tool_message(
         out,
         "message",
         &[("role", "tool"), ("tool", tool_name), ("index", &index_str)],
-    );
+    )?;
     write_content(out, content)?;
-    write_close_tag(out, "message");
-    out.push_str(END_OF_MSG);
-    Ok(())
+    write_close_tag(out, "message")?;
+    out.control(END_OF_MSG)
 }
 
 fn write_assistant_message(
-    out: &mut String,
+    out: &mut K3TokenWriter<'_>,
     content: &[AssistantContentBlock],
     thinking: bool,
 ) -> Result<()> {
-    write_open_tag(out, "message", &[("role", "assistant")]);
+    write_open_tag(out, "message", &[("role", "assistant")])?;
 
     let mut reasoning = String::new();
     let mut response = String::new();
@@ -358,32 +398,31 @@ fn write_assistant_message(
     // message carries open/close tags even when there is no reasoning content.
     // In non-thinking mode the channel is dropped entirely.
     if thinking {
-        write_open_tag(out, "think", &[]);
+        write_open_tag(out, "think", &[])?;
         if !reasoning.trim().is_empty() {
-            out.push_str(&reasoning);
+            out.ordinary(&reasoning)?;
         }
-        write_close_tag(out, "think");
+        write_close_tag(out, "think")?;
     }
 
-    write_open_tag(out, "response", &[]);
-    out.push_str(&response);
-    write_close_tag(out, "response");
+    write_open_tag(out, "response", &[])?;
+    out.ordinary(&response)?;
+    write_close_tag(out, "response")?;
 
     if !tool_calls.is_empty() {
-        write_open_tag(out, "tools", &[]);
+        write_open_tag(out, "tools", &[])?;
         for (index, tool_call) in tool_calls.into_iter().enumerate() {
             write_assistant_tool_call(out, tool_call, index + 1)?;
         }
-        write_close_tag(out, "tools");
+        write_close_tag(out, "tools")?;
     }
 
-    write_close_tag(out, "message");
-    out.push_str(END_OF_MSG);
-    Ok(())
+    write_close_tag(out, "message")?;
+    out.control(END_OF_MSG)
 }
 
 fn write_assistant_tool_call(
-    out: &mut String,
+    out: &mut K3TokenWriter<'_>,
     tool_call: &AssistantToolCall,
     index: usize,
 ) -> Result<()> {
@@ -395,34 +434,33 @@ fn write_assistant_tool_call(
             ("tool", tool_call.name.as_str()),
             ("index", index_str.as_str()),
         ],
-    );
+    )?;
 
     let (args, json_block) = normalize_tool_arguments(&tool_call.arguments)?;
     if let Some(raw) = json_block {
-        write_open_tag(out, "json", &[("type", "object")]);
-        out.push_str(&raw);
-        write_close_tag(out, "json");
+        write_open_tag(out, "json", &[("type", "object")])?;
+        out.ordinary(&raw)?;
+        write_close_tag(out, "json")?;
     } else {
         for (key, value) in args {
             let typ = xtml_type(&value);
-            write_open_tag(out, "argument", &[("key", key.as_str()), ("type", typ)]);
-            out.push_str(&xtml_value(&value));
-            write_close_tag(out, "argument");
+            write_open_tag(out, "argument", &[("key", key.as_str()), ("type", typ)])?;
+            out.ordinary(&xtml_value(&value))?;
+            write_close_tag(out, "argument")?;
         }
     }
 
-    write_close_tag(out, "call");
-    Ok(())
+    write_close_tag(out, "call")
 }
 
-fn write_content(out: &mut String, content: &ChatContent) -> Result<()> {
+fn write_content(out: &mut K3TokenWriter<'_>, content: &ChatContent) -> Result<()> {
     match content {
         ChatContent::Text(text) => write_text_with_images(out, text),
         ChatContent::Parts(parts) => {
             for part in parts {
                 match part {
                     ChatContentPart::Text { text } => write_text_with_images(out, text)?,
-                    ChatContentPart::ImageUrl { .. } => out.push_str(IMAGE_PLACEHOLDER),
+                    ChatContentPart::ImageUrl { .. } => out.control(IMAGE_PLACEHOLDER)?,
                     ChatContentPart::VideoUrl { .. } => {
                         return Err(Error::UnsupportedMultimodalContent("video_url"));
                     }
@@ -439,15 +477,14 @@ fn write_content(out: &mut String, content: &ChatContent) -> Result<()> {
     }
 }
 
-fn write_text_with_images(out: &mut String, text: &str) -> Result<()> {
+fn write_text_with_images(out: &mut K3TokenWriter<'_>, text: &str) -> Result<()> {
     // Placeholder expansion is left as the literal K3 image token; multimodal
     // preprocessing can replace it once image prompts are known.
-    out.push_str(text);
-    Ok(())
+    out.ordinary(text)
 }
 
-fn write_response_format(out: &mut String, request: &ChatRequest) -> Result<()> {
-    let Some(rf) = request.chat_options.template_kwargs.get("response_format") else {
+fn write_response_format(out: &mut K3TokenWriter<'_>, request: &ChatRequest) -> Result<()> {
+    let Some(rf) = request.chat_options.response_format.as_ref() else {
         return Ok(());
     };
 
@@ -461,7 +498,7 @@ fn write_response_format(out: &mut String, request: &ChatRequest) -> Result<()>
                 "The system is invoked with `response_format=json_object`.\n\
                  Your response must be raw JSON data without markdown code \
                  blocks (```json) or any additional formatting.",
-            );
+            )?;
         }
         "json_schema" => {
             let schema = extract_response_schema(rf);
@@ -478,7 +515,7 @@ fn write_response_format(out: &mut String, request: &ChatRequest) -> Result<()>
                      {schema_json}\n\
                      ```"
                 ),
-            );
+            )?;
         }
         _ => {}
     }
@@ -528,19 +565,22 @@ fn xtml_value(value: &Value) -> String {
     }
 }
 
-fn write_open_tag(out: &mut String, tag: &str, attrs: &[(&str, &str)]) {
-    out.push_str(OPEN);
-    out.push_str(tag);
+fn write_open_tag(out: &mut K3TokenWriter<'_>, tag: &str, attrs: &[(&str, &str)]) -> Result<()> {
+    out.control(OPEN)?;
+    out.ordinary(tag)?;
     for (key, value) in attrs {
-        let _ = write!(out, " {key}=\"{}\"", escape_attr_value(value));
+        out.ordinary(&format!(" {key}"))?;
+        out.ordinary("=\"")?;
+        out.ordinary(&escape_attr_value(value))?;
+        out.ordinary("\"")?;
     }
-    out.push_str(SEP);
+    out.control(SEP)
 }
 
-fn write_close_tag(out: &mut String, tag: &str) {
-    out.push_str(CLOSE);
-    out.push_str(tag);
-    out.push_str(SEP);
+fn write_close_tag(out: &mut K3TokenWriter<'_>, tag: &str) -> Result<()> {
+    out.control(CLOSE)?;
+    out.ordinary(tag)?;
+    out.control(SEP)
 }
 
 fn escape_attr_value(value: &str) -> String {
diff --git a/rust/src/chat/src/renderer/kimi_k3/fixtures/controls_thinking_off_input.json b/rust/src/chat/src/renderer/kimi_k3/fixtures/controls_thinking_off_input.json
index 90ab846e131..569201f3b97 100644
--- a/rust/src/chat/src/renderer/kimi_k3/fixtures/controls_thinking_off_input.json
+++ b/rust/src/chat/src/renderer/kimi_k3/fixtures/controls_thinking_off_input.json
@@ -6,10 +6,10 @@
     }
   ],
   "add_generation_prompt": true,
+  "response_format": {
+    "type": "json_object"
+  },
   "template_kwargs": {
-    "thinking": false,
-    "response_format": {
-      "type": "json_object"
-    }
+    "thinking": false
   }
 }
diff --git a/rust/src/chat/src/renderer/kimi_k3/mod.rs b/rust/src/chat/src/renderer/kimi_k3/mod.rs
index 1ff0e556dfa..09bf356f5d6 100644
--- a/rust/src/chat/src/renderer/kimi_k3/mod.rs
+++ b/rust/src/chat/src/renderer/kimi_k3/mod.rs
@@ -8,19 +8,22 @@ mod encoding;
 mod tests;
 
 use vllm_text::Prompt;
+use vllm_text::tokenizer::DynTokenizer;
 
 use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
 use crate::Result;
 use crate::request::ChatRequest;
 
 /// Dedicated Kimi K3 XTML renderer.
-#[derive(Debug, Clone, Copy, Default)]
-pub struct KimiK3ChatRenderer;
+#[derive(Clone)]
+pub struct KimiK3ChatRenderer {
+    tokenizer: DynTokenizer,
+}
 
 impl KimiK3ChatRenderer {
     /// Create a Kimi K3 renderer.
-    pub fn new() -> Self {
-        Self
+    pub fn new(tokenizer: DynTokenizer) -> Self {
+        Self { tokenizer }
     }
 }
 
@@ -29,7 +32,7 @@ impl ChatRenderer for KimiK3ChatRenderer {
         request.validate()?;
 
         Ok(RenderedPrompt {
-            prompt: Prompt::Text(encoding::render_request(request)?),
+            prompt: Prompt::TokenIds(encoding::render_request(request, self.tokenizer.as_ref())?),
             effective_template_kwargs: request_template_kwargs(request),
         })
     }
diff --git a/rust/src/chat/src/renderer/kimi_k3/tests.rs b/rust/src/chat/src/renderer/kimi_k3/tests.rs
index b96e04fad2a..7bab90ddcf7 100644
--- a/rust/src/chat/src/renderer/kimi_k3/tests.rs
+++ b/rust/src/chat/src/renderer/kimi_k3/tests.rs
@@ -4,23 +4,49 @@
 //! Golden fixtures generated from HF remote-code `encoding_k3.py`.
 
 use std::path::PathBuf;
+use std::sync::Arc;
 
 use expect_test::{expect, expect_file};
 use serde_json::json;
+use vllm_text::Prompt;
+use vllm_text::tokenizer::DynTokenizer;
+use vllm_tokenizer::Tokenizer;
+use vllm_tokenizer::test_utils::TestTokenizer;
 
 use super::KimiK3ChatRenderer;
 use crate::AssistantContentBlock;
 use crate::ChatRenderer;
+use crate::renderer::kimi_k3::encoding::{CLOSE, END_OF_MSG, IMAGE_PLACEHOLDER, OPEN, SEP};
 use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
-use crate::request::{ChatMessage, GenerationPromptMode, ReasoningEffort};
+use crate::request::{ChatContentPart, ChatMessage, GenerationPromptMode, ReasoningEffort};
+
+const OPEN_ID: u32 = 256;
+const CLOSE_ID: u32 = 257;
+const SEP_ID: u32 = 258;
+const END_OF_MSG_ID: u32 = 259;
+const MEDIA_ID: u32 = 260;
+
+fn test_tokenizer() -> TestTokenizer {
+    TestTokenizer::new()
+        .with_special_token(OPEN, OPEN_ID)
+        .with_special_token(CLOSE, CLOSE_ID)
+        .with_special_token(SEP, SEP_ID)
+        .with_special_token(END_OF_MSG, END_OF_MSG_ID)
+        .with_special_token(IMAGE_PLACEHOLDER, MEDIA_ID)
+}
+
+fn render_token_ids(request: &crate::request::ChatRequest, tokenizer: DynTokenizer) -> Vec<u32> {
+    let prompt = KimiK3ChatRenderer::new(tokenizer).render(request).unwrap().prompt;
+    let Prompt::TokenIds(token_ids) = prompt else {
+        panic!("kimi k3 renderer should return token IDs")
+    };
+    token_ids
+}
 
 fn render_request(request: &crate::request::ChatRequest) -> String {
-    KimiK3ChatRenderer::new()
-        .render(request)
-        .unwrap()
-        .prompt
-        .into_text()
-        .expect("kimi k3 renderer should return text prompt")
+    let tokenizer: DynTokenizer = Arc::new(test_tokenizer());
+    let token_ids = render_token_ids(request, tokenizer.clone());
+    tokenizer.decode(&token_ids, false).unwrap()
 }
 
 fn fixture_path(name: &str) -> PathBuf {
@@ -64,6 +90,44 @@ fn golden_dynamic_system_tool_declare() {
     assert_golden("dynamic_system_tool_declare");
 }
 
+#[test]
+fn token_writer_protects_literal_control_and_media_markers() {
+    let tokenizer = Arc::new(test_tokenizer());
+    let user_text = format!("literal {OPEN} and {}", super::encoding::IMAGE_PLACEHOLDER);
+    let mut request = crate::request::ChatRequest::for_test();
+    request.messages = vec![ChatMessage::user(vec![
+        ChatContentPart::text(user_text),
+        ChatContentPart::image_url("data:image/png;base64,test"),
+    ])];
+    request
+        .chat_options
+        .template_kwargs
+        .insert("thinking".to_string(), json!(false));
+    request.chat_options.generation_prompt_mode = GenerationPromptMode::NoGenerationPrompt;
+
+    let token_ids = render_token_ids(&request, tokenizer.clone());
+
+    assert_eq!(
+        token_ids.iter().filter(|&&token_id| token_id == OPEN_ID).count(),
+        1
+    );
+    assert_eq!(
+        token_ids.iter().filter(|&&token_id| token_id == MEDIA_ID).count(),
+        1
+    );
+
+    let flattened = tokenizer.decode(&token_ids, false).unwrap();
+    let flattened_ids = tokenizer.encode(&flattened, false).unwrap();
+    assert_eq!(
+        flattened_ids.iter().filter(|&&token_id| token_id == OPEN_ID).count(),
+        2
+    );
+    assert_eq!(
+        flattened_ids.iter().filter(|&&token_id| token_id == MEDIA_ID).count(),
+        2
+    );
+}
+
 #[test]
 fn thinking_history_renders_empty_think_channel() {
     let mut request = crate::request::ChatRequest::for_test();
@@ -190,7 +254,9 @@ fn rejects_removed_medium_thinking_effort() {
         .template_kwargs
         .insert("thinking_effort".to_string(), json!("medium"));
 
-    let error = KimiK3ChatRenderer::new().render(&request).unwrap_err();
+    let error = KimiK3ChatRenderer::new(Arc::new(test_tokenizer()))
+        .render(&request)
+        .unwrap_err();
 
     expect![[r#"
         ChatTemplate(
diff --git a/rust/src/chat/src/renderer/test_utils.rs b/rust/src/chat/src/renderer/test_utils.rs
index 138faacc1db..0a4edc81379 100644
--- a/rust/src/chat/src/renderer/test_utils.rs
+++ b/rust/src/chat/src/renderer/test_utils.rs
@@ -48,7 +48,10 @@ pub(crate) struct FixtureRequest {
     messages: Vec<FixtureMessage>,
     add_generation_prompt: Option<bool>,
     reasoning_effort: Option<ReasoningEffort>,
-    /// Extra chat-template kwargs (thinking, preserve_thinking, response_format, …).
+    /// Standard response format passed to model-specific renderers.
+    #[serde(default)]
+    response_format: Option<Value>,
+    /// Extra chat-template kwargs (thinking, preserve_thinking, …).
     #[serde(default)]
     template_kwargs: HashMap<String, Value>,
     /// When omitted, defaults to `auto` if tools are present, otherwise `none`.
@@ -65,6 +68,7 @@ impl FixtureFile {
                 messages,
                 add_generation_prompt: None,
                 reasoning_effort: None,
+                response_format: None,
                 template_kwargs: HashMap::new(),
                 tool_choice: None,
             },
@@ -182,6 +186,7 @@ impl FixtureRequest {
             request.chat_options.generation_prompt_mode = GenerationPromptMode::NoGenerationPrompt;
         }
         request.chat_options.reasoning_effort = self.reasoning_effort;
+        request.chat_options.response_format = self.response_format;
         request.chat_options.template_kwargs.extend(self.template_kwargs);
 
         // Options supply a default thinking toggle only when the fixture did not.
diff --git a/rust/src/chat/src/request.rs b/rust/src/chat/src/request.rs
index eb72557c7d2..56a764f8987 100644
--- a/rust/src/chat/src/request.rs
+++ b/rust/src/chat/src/request.rs
@@ -397,6 +397,10 @@ pub struct ChatOptions {
     /// Effort level exposed to chat templates for reasoning models.
     pub reasoning_effort: Option<ReasoningEffort>,
 
+    /// Standard response format available to model-specific renderers.
+    #[serde(default)]
+    pub response_format: Option<Value>,
+
     /// Additional keyword arguments exposed to the chat template.
     pub template_kwargs: HashMap<String, Value>,
 }
@@ -407,6 +411,7 @@ impl Default for ChatOptions {
             generation_prompt_mode: GenerationPromptMode::StartNewAssistant,
             chat_template: None,
             reasoning_effort: None,
+            response_format: None,
             template_kwargs: HashMap::new(),
         }
     }
diff --git a/rust/src/chat/tests/roundtrip.rs b/rust/src/chat/tests/roundtrip.rs
index 3715ed748ed..9cd670d29d0 100644
--- a/rust/src/chat/tests/roundtrip.rs
+++ b/rust/src/chat/tests/roundtrip.rs
@@ -272,7 +272,7 @@ impl RoundtripCase {
     fn gpt_oss() -> Self {
         Self {
             model_id: "openai/gpt-oss-20b",
-            assistant_stop_suffix: "", // not applicable for token-id cases
+            assistant_stop_suffix: "",
             tool_call_parser: ParserSelection::Auto,
             reasoning_parser: ParserSelection::Auto,
             thinking_behavior: ThinkingBehavior::Always { value: true },
@@ -285,7 +285,7 @@ impl RoundtripCase {
     fn inkling() -> Self {
         Self {
             model_id: "thinkingmachines/Inkling",
-            assistant_stop_suffix: "",
+            assistant_stop_suffix: "<|content_model_end_sampling|>",
             tool_call_parser: ParserSelection::Auto,
             reasoning_parser: ParserSelection::Auto,
             thinking_behavior: ThinkingBehavior::Always { value: true },
@@ -689,14 +689,23 @@ fn decoded_completion_stream(
                 .collect()
         }
         Prompt::TokenIds(token_ids) => {
-            ensure!(
-                assistant_stop_suffix.is_empty(),
-                "token-id roundtrip cases do not support text stop suffixes"
-            );
+            let body = if assistant_stop_suffix.is_empty() {
+                token_ids.as_slice()
+            } else {
+                let stop_token_ids = tokenizer
+                    .encode(assistant_stop_suffix, false)
+                    .context("failed to encode token-id completion stop suffix")?;
+                token_ids.strip_suffix(stop_token_ids.as_slice()).with_context(|| {
+                    format!(
+                        "token-id completion did not end with {:?}: {:?}",
+                        assistant_stop_suffix, token_ids
+                    )
+                })?
+            };
             incremental_decode_chunks(
                 tokenizer,
                 &prompt_token_ids,
-                token_ids,
+                body,
                 TOKEN_COMPLETION_CHUNK_TOKENS,
             )?
         }
diff --git a/rust/src/cmd/Cargo.toml b/rust/src/cmd/Cargo.toml
index a6955059c26..a326a0f9992 100644
--- a/rust/src/cmd/Cargo.toml
+++ b/rust/src/cmd/Cargo.toml
@@ -23,17 +23,16 @@ serde.workspace = true
 serde_json.workspace = true
 serde_with.workspace = true
 thiserror-ext.workspace = true
-time.workspace = true
 tokio = { workspace = true, features = ["signal"] }
 tokio-util.workspace = true
 tracing.workspace = true
-tracing-subscriber.workspace = true
 uuid.workspace = true
 vllm-bench.workspace = true
 vllm-chat.workspace = true
 vllm-engine-core-client.workspace = true
 vllm-managed-engine.workspace = true
 vllm-server.workspace = true
+vllm-tracing.workspace = true
 
 [dev-dependencies]
 expect-test.workspace = true
diff --git a/rust/src/cmd/src/main.rs b/rust/src/cmd/src/main.rs
index 0806f7b75d4..86c9ab6f934 100644
--- a/rust/src/cmd/src/main.rs
+++ b/rust/src/cmd/src/main.rs
@@ -2,7 +2,6 @@
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 mod cli;
-mod logging;
 
 use std::env;
 use std::ffi::OsStr;
@@ -89,7 +88,7 @@ fn main() -> Result<()> {
             "serve" | "frontend" => "RustFrontend",
             _ => "Rust",
         };
-    logging::init_tracing(process_label);
+    vllm_tracing::init_tracing(process_label);
 
     let cli = Cli::parse();
 
diff --git a/rust/src/parser/benches/utils/adapter.rs b/rust/src/parser/benches/utils/adapter.rs
index 243f19e3ce8..9cba325116c 100644
--- a/rust/src/parser/benches/utils/adapter.rs
+++ b/rust/src/parser/benches/utils/adapter.rs
@@ -19,6 +19,10 @@ impl Tokenizer for BenchTokenizer {
         Ok(text.chars().map(|_| u32::MAX).collect())
     }
 
+    fn encode_ordinary(&self, text: &str) -> vllm_tokenizer::Result<Vec<u32>> {
+        self.encode(text, false)
+    }
+
     fn decode(
         &self,
         token_ids: &[u32],
diff --git a/rust/src/parser/src/unified/inkling.rs b/rust/src/parser/src/unified/inkling.rs
index eb78321d126..1ce69d8bfb4 100644
--- a/rust/src/parser/src/unified/inkling.rs
+++ b/rust/src/parser/src/unified/inkling.rs
@@ -414,6 +414,10 @@ mod tests {
             Ok(text.chars().map(u32::from).collect())
         }
 
+        fn encode_ordinary(&self, text: &str) -> vllm_tokenizer::Result<Vec<u32>> {
+            self.encode(text, false)
+        }
+
         fn decode(
             &self,
             token_ids: &[u32],
@@ -733,6 +737,10 @@ mod tests {
                 Ok(vec![])
             }
 
+            fn encode_ordinary(&self, text: &str) -> vllm_tokenizer::Result<Vec<u32>> {
+                self.encode(text, false)
+            }
+
             fn decode(
                 &self,
                 _token_ids: &[u32],
diff --git a/rust/src/server/src/error.rs b/rust/src/server/src/error.rs
index fbd13b77baf..3a8c0ae5c92 100644
--- a/rust/src/server/src/error.rs
+++ b/rust/src/server/src/error.rs
@@ -141,6 +141,22 @@ mod tests {
         assert!(response.error.message.contains("max_tokens=4"));
     }
 
+    #[test]
+    fn sampling_params_validation_maps_to_invalid_request() {
+        let api_error = text_submit_error(
+            "failed to submit completion request",
+            vllm_text::Error::SamplingParams(vllm_text::SamplingParamsError::OutOfRange {
+                parameter: "top_p",
+                value: 0.0,
+                expected: "(0, 1]",
+            }),
+        );
+        assert_eq!(api_error.status_code(), StatusCode::BAD_REQUEST);
+        let response = api_error.to_error_response();
+        assert_eq!(response.error.error_type, "invalid_request_error");
+        assert!(response.error.message.contains("top_p"));
+    }
+
     #[test]
     fn chat_wrapped_prompt_too_long_maps_to_invalid_request() {
         let error = vllm_chat::Error::Text(vllm_text::Error::PromptTooLong {
diff --git a/rust/src/server/src/grpc/tests.rs b/rust/src/server/src/grpc/tests.rs
index 75da670ec3b..808f3422758 100644
--- a/rust/src/server/src/grpc/tests.rs
+++ b/rust/src/server/src/grpc/tests.rs
@@ -648,6 +648,35 @@ async fn unary_generate_min_tokens_above_max_tokens_returns_invalid_argument() {
     server_task.abort();
 }
 
+#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
+#[serial]
+async fn unary_generate_invalid_sampling_params_returns_invalid_argument() {
+    let (mut client, server_task, _engine_task) = grpc_test_server(
+        b"engine-grpc-invalid-sampling",
+        default_stream_output_specs(),
+    )
+    .await;
+
+    let status = client
+        .generate(pb::GenerateRequest {
+            request_id: "test-invalid-sampling".to_string(),
+            model: "test-model".to_string(),
+            prompt: Some(pb::generate_request::Prompt::Text("hi".to_string())),
+            sampling: Some(pb::RandomSampling {
+                top_p: 2.0,
+                ..Default::default()
+            }),
+            ..Default::default()
+        })
+        .await
+        .expect_err("should fail when top_p is out of range");
+
+    assert_eq!(status.code(), tonic::Code::InvalidArgument);
+    assert!(status.message().contains("top_p"));
+
+    server_task.abort();
+}
+
 #[tokio::test(flavor = "multi_thread", worker_threads = 2)]
 #[serial]
 async fn streaming_generate_yields_incremental_responses() {
diff --git a/rust/src/server/src/routes/openai/chat_completions.rs b/rust/src/server/src/routes/openai/chat_completions.rs
index bd86a06136a..8a1e9c4383a 100644
--- a/rust/src/server/src/routes/openai/chat_completions.rs
+++ b/rust/src/server/src/routes/openai/chat_completions.rs
@@ -131,6 +131,7 @@ async fn collect_chat_completion(
         echo,
         return_token_ids,
         return_tokens_as_token_ids,
+        is_named_tool_choice,
     }: ResponseOptions,
 ) -> Result<ChatCompletionResponse, ApiError> {
     let collected = stream.collect_message().await.map_err(|error| {
@@ -157,7 +158,9 @@ async fn collect_chat_completion(
     // When reasoning is hidden, omit them rather than leaking hidden reasoning
     // tokens through per-token metadata.
     let include_output_metadata = include_reasoning || reasoning.is_none();
-    let finish_reason = chat_finish_reason_to_openai(&finish_reason, saw_tool_calls)?.to_string();
+    let finish_reason =
+        chat_finish_reason_to_openai(&finish_reason, saw_tool_calls && !is_named_tool_choice)?
+            .to_string();
     let tool_calls = message
         .tool_calls()
         .map(|call| ToolCall {
@@ -254,6 +257,7 @@ async fn chat_completion_chunk_stream(
         echo,
         return_token_ids,
         return_tokens_as_token_ids,
+        is_named_tool_choice,
     }: ResponseOptions,
     mut y: TryYielder<ChatCompletionStreamResponse, ApiError>,
 ) -> Result<(), ApiError> {
@@ -454,7 +458,7 @@ async fn chat_completion_chunk_stream(
                     &response_model,
                     created,
                     finish_reason,
-                    saw_tool_calls,
+                    saw_tool_calls && !is_named_tool_choice,
                 ) {
                     Ok(chunk) => yield_chunk!(chunk),
                     Err(error) => {
@@ -787,10 +791,10 @@ fn final_chunk(
     response_model: &str,
     created: u64,
     finish_reason: FinishReason,
-    saw_tool_calls: bool,
+    use_tool_calls_finish_reason: bool,
 ) -> Result<ChatCompletionStreamResponse, ApiError> {
     let stop_reason = finish_reason.as_stop_reason().map(stop_reason_to_json);
-    let finish_reason = chat_finish_reason_to_openai(&finish_reason, saw_tool_calls)?;
+    let finish_reason = chat_finish_reason_to_openai(&finish_reason, use_tool_calls_finish_reason)?;
 
     debug!(
         finish_reason = %finish_reason,
@@ -809,10 +813,10 @@ fn final_chunk(
 
 fn chat_finish_reason_to_openai(
     finish_reason: &FinishReason,
-    saw_tool_calls: bool,
+    use_tool_calls_finish_reason: bool,
 ) -> Result<&'static str, ApiError> {
     match finish_reason {
-        FinishReason::Stop(_) if saw_tool_calls => Ok("tool_calls"),
+        FinishReason::Stop(_) if use_tool_calls_finish_reason => Ok("tool_calls"),
         FinishReason::Stop(_) => Ok("stop"),
         FinishReason::Length => Ok("length"),
         FinishReason::Abort => Ok("abort"),
diff --git a/rust/src/server/src/routes/openai/chat_completions/convert.rs b/rust/src/server/src/routes/openai/chat_completions/convert.rs
index 5b2fa3c19ed..2c4050f8f1e 100644
--- a/rust/src/server/src/routes/openai/chat_completions/convert.rs
+++ b/rust/src/server/src/routes/openai/chat_completions/convert.rs
@@ -54,6 +54,8 @@ pub(super) struct ResponseOptions {
     pub return_token_ids: bool,
     /// Whether to format logprob tokens as `token_id:{id}`.
     pub return_tokens_as_token_ids: bool,
+    /// Whether the request forces one named function tool.
+    pub is_named_tool_choice: bool,
 }
 
 /// Validate and lower one OpenAI chat completion request into the internal chat
@@ -87,6 +89,15 @@ pub(super) fn prepare_chat_request(
     )?;
 
     let template_kwargs = request.chat_template_kwargs.unwrap_or_default();
+    let response_format =
+        request.response_format.as_ref().map(serde_json::to_value).transpose().map_err(
+            |error| {
+                ApiError::invalid_request(
+                    format!("failed to serialize response_format: {error}"),
+                    Some("response_format"),
+                )
+            },
+        )?;
 
     let include_usage = (request.stream_options.as_ref())
         .and_then(|options| options.include_usage)
@@ -98,6 +109,7 @@ pub(super) fn prepare_chat_request(
             .and_then(|options| options.continuous_usage_stats)
             .unwrap_or(false);
     let requested_logprobs = request.logprobs;
+    let is_named_tool_choice = matches!(&request.tool_choice, Some(ToolChoice::Function { .. }));
 
     // Auto-enable prompt logprobs for non-streaming echo, matching Python vLLM's
     // behavior.
@@ -147,6 +159,7 @@ pub(super) fn prepare_chat_request(
             generation_prompt_mode,
             chat_template: request.chat_template,
             reasoning_effort: request.reasoning_effort,
+            response_format,
             template_kwargs,
         },
         tools: convert_tools(request.tools)?,
@@ -180,6 +193,7 @@ pub(super) fn prepare_chat_request(
             echo,
             return_token_ids: request.return_token_ids.unwrap_or(false),
             return_tokens_as_token_ids: request.return_tokens_as_token_ids.unwrap_or(false),
+            is_named_tool_choice,
         },
         chat_request,
     })
@@ -396,6 +410,7 @@ fn convert_tool_choice(tool_choice: Option<&ToolChoice>) -> Result<ChatToolChoic
 #[cfg(test)]
 mod tests {
     use std::collections::HashMap;
+    use std::sync::Arc;
 
     use axum::http::HeaderMap;
     use expect_test::expect;
@@ -403,16 +418,18 @@ mod tests {
     use serde_json::json;
     use vllm_chat::{
         AssistantContentBlock, AssistantToolCall, ChatContentPart, ChatMessage as VllmChatMessage,
-        ChatTool as VllmChatTool, ChatToolChoice, GenerationPromptMode,
-        SamplingParams as VllmSamplingParams,
+        ChatRenderer, ChatTool as VllmChatTool, ChatToolChoice, GenerationPromptMode,
+        KimiK3ChatRenderer, SamplingParams as VllmSamplingParams,
     };
-    use vllm_text::output::TextDecodeOptions;
+    use vllm_text::{Prompt, output::TextDecodeOptions};
+    use vllm_tokenizer::{Tokenizer, test_utils::TestTokenizer};
 
     use super::prepare_chat_request;
     use crate::lora::LoraModelResolution;
     use crate::routes::openai::chat_completions::types::{
         AssistantRole, ChatCompletionMessage, ChatCompletionRequest,
     };
+    use crate::routes::openai::utils::structured_outputs::{JsonSchemaFormat, ResponseFormat};
     use crate::routes::openai::utils::types::{
         AudioUrl, ChatMessage, ContentPart, Function, FunctionCallResponse, ImageUrl, InputAudio,
         MessageContent, StreamOptions, Tool, ToolCall, ToolChoice, ToolChoiceValue, VideoUrl,
@@ -469,6 +486,60 @@ mod tests {
         assert!(prepared.chat_request.parallel_tool_calls);
     }
 
+    #[test]
+    fn prepare_chat_request_passes_response_format_to_kimi_k3_renderer() {
+        let mut request = base_request();
+        request.model = "moonshotai/Kimi-K3".to_string();
+        let response_format = ResponseFormat::JsonSchema {
+            json_schema: JsonSchemaFormat {
+                name: "answer".to_string(),
+                description: None,
+                schema: json!({
+                    "type": "object",
+                    "properties": {
+                        "answer": {"type": "string"}
+                    },
+                    "required": ["answer"]
+                }),
+                strict: Some(true),
+            },
+        };
+        request.response_format = Some(response_format.clone());
+        let template_response_format = json!({"type": "json_object"});
+        request.chat_template_kwargs = Some(HashMap::from([(
+            "response_format".to_string(),
+            template_response_format.clone(),
+        )]));
+
+        let prepared = prepare_chat_request(
+            request,
+            &served(&["moonshotai/Kimi-K3"]),
+            ResolvedRequestContext::default(),
+        )
+        .expect("request is valid");
+
+        assert_eq!(
+            prepared.chat_request.chat_options.response_format.as_ref(),
+            Some(&serde_json::to_value(response_format).unwrap())
+        );
+        assert_eq!(
+            prepared.chat_request.chat_options.template_kwargs.get("response_format"),
+            Some(&template_response_format)
+        );
+
+        let tokenizer = Arc::new(TestTokenizer::new());
+        let prompt = KimiK3ChatRenderer::new(tokenizer.clone())
+            .render(&prepared.chat_request)
+            .expect("Kimi K3 rendering succeeds")
+            .prompt;
+        let Prompt::TokenIds(prompt_token_ids) = prompt else {
+            panic!("Kimi K3 renders token IDs");
+        };
+        let prompt = tokenizer.decode(&prompt_token_ids, false).expect("Kimi K3 prompt decodes");
+        assert!(prompt.contains("response_format=json_schema"));
+        assert!(prompt.contains(r#""answer":{"type":"string"}"#));
+    }
+
     #[test]
     fn prepare_chat_request_maps_text_parts() {
         let mut request = base_request();
@@ -1068,6 +1139,7 @@ mod tests {
         .expect("request is valid");
 
         assert_eq!(prepared.chat_request.tool_choice, ChatToolChoice::Required);
+        assert!(!prepared.options.is_named_tool_choice);
     }
 
     #[test]
@@ -1107,6 +1179,7 @@ mod tests {
                 name: "get_weather".to_string(),
             }
         );
+        assert!(prepared.options.is_named_tool_choice);
     }
 
     #[test]
diff --git a/rust/src/server/src/routes/tests.rs b/rust/src/server/src/routes/tests.rs
index 64e904b562b..1c228f99f1f 100644
--- a/rust/src/server/src/routes/tests.rs
+++ b/rust/src/server/src/routes/tests.rs
@@ -153,6 +153,20 @@ fn default_stream_output_specs() -> Vec<(Vec<u32>, Option<EngineCoreFinishReason
     ]
 }
 
+fn weather_tool_call_output_specs() -> Vec<(Vec<u32>, Option<EngineCoreFinishReason>)> {
+    vec![
+        (bytes_to_token_ids(b"<think>Need tool.</think>"), None),
+        (
+            bytes_to_token_ids(b"<tool_call>\n{\"name\":\"get_weather\", "),
+            None,
+        ),
+        (
+            bytes_to_token_ids(b"\"arguments\":{\"city\":\"Paris\"}}\n</tool_call>"),
+            Some(EngineCoreFinishReason::Stop),
+        ),
+    ]
+}
+
 fn assert_adapter_a_lora_request(request: &EngineCoreRequest) {
     let lora = request.lora_request.as_ref().expect("lora request");
     assert_eq!(lora.lora_name, "adapter-a");
@@ -4554,17 +4568,7 @@ async fn include_reasoning_false_suppresses_non_stream_output_metadata() {
 async fn tool_calls_are_mapped_to_tool_call_sse_chunks() {
     let (app, engine_task) = test_app_with_backend_and_stream_output_specs(
         Arc::new(FakeChatBackend::with_model_id("Qwen/Qwen3-0.6B")),
-        vec![
-            (bytes_to_token_ids(b"<think>Need tool.</think>"), None),
-            (
-                bytes_to_token_ids(b"<tool_call>\n{\"name\":\"get_weather\", "),
-                None,
-            ),
-            (
-                bytes_to_token_ids(b"\"arguments\":{\"city\":\"Paris\"}}\n</tool_call>"),
-                Some(EngineCoreFinishReason::Stop),
-            ),
-        ],
+        weather_tool_call_output_specs(),
     )
     .await;
 
@@ -4613,6 +4617,63 @@ async fn tool_calls_are_mapped_to_tool_call_sse_chunks() {
     assert!(text.contains("\"finish_reason\":\"tool_calls\""), "{text}");
 }
 
+#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
+#[serial]
+async fn named_tool_choice_uses_stop_finish_reason() {
+    for stream in [false, true] {
+        let (app, engine_task) = test_app_with_backend_and_stream_output_specs(
+            Arc::new(FakeChatBackend::with_model_id("Qwen/Qwen3-0.6B")),
+            weather_tool_call_output_specs(),
+        )
+        .await;
+
+        let response = app
+            .clone()
+            .call(
+                Request::builder()
+                    .method("POST")
+                    .uri("/v1/chat/completions")
+                    .header("content-type", "application/json")
+                    .body(Body::from(
+                        json!({
+                            "model": "Qwen/Qwen1.5-0.5B-Chat",
+                            "stream": stream,
+                            "messages": [{"role": "user", "content": "hello"}],
+                            "tools": [{
+                                "type": "function",
+                                "function": {
+                                    "name": "get_weather",
+                                    "description": "Get weather",
+                                    "parameters": {
+                                        "type": "object",
+                                        "properties": {"city": {"type": "string"}}
+                                    }
+                                }
+                            }],
+                            "tool_choice": {
+                                "type": "function",
+                                "function": {"name": "get_weather"}
+                            }
+                        })
+                        .to_string(),
+                    ))
+                    .expect("build request"),
+            )
+            .await
+            .expect("call app");
+
+        assert_eq!(response.status(), StatusCode::OK);
+        let body = to_bytes(response.into_body(), usize::MAX).await.expect("read body");
+        engine_task.await.expect("mock engine task");
+        let text = String::from_utf8(body.to_vec()).expect("utf8 body");
+
+        assert!(text.contains("\"tool_calls\":"), "{text}");
+        assert!(text.contains("\"name\":\"get_weather\""), "{text}");
+        assert!(text.contains("\"finish_reason\":\"stop\""), "{text}");
+        assert!(!text.contains("\"finish_reason\":\"tool_calls\""), "{text}");
+    }
+}
+
 #[tokio::test(flavor = "multi_thread", worker_threads = 2)]
 #[serial]
 async fn tool_call_sse_chunks_can_carry_logprobs() {
diff --git a/rust/src/server/src/routes/tokenize/types.rs b/rust/src/server/src/routes/tokenize/types.rs
index 242a7af4733..6684aa55d26 100644
--- a/rust/src/server/src/routes/tokenize/types.rs
+++ b/rust/src/server/src/routes/tokenize/types.rs
@@ -82,6 +82,7 @@ impl TokenizeChatRequest {
                 generation_prompt_mode,
                 chat_template: self.chat_template,
                 reasoning_effort: None,
+                response_format: None,
                 template_kwargs: self.chat_template_kwargs.unwrap_or_default(),
             },
             tools: convert_tools(self.tools)?,
diff --git a/rust/src/text/src/backend/hf/mod.rs b/rust/src/text/src/backend/hf/mod.rs
index 49ae5dbd6b9..4fc7a18753a 100644
--- a/rust/src/text/src/backend/hf/mod.rs
+++ b/rust/src/text/src/backend/hf/mod.rs
@@ -177,6 +177,10 @@ mod tests {
             Ok(vec![])
         }
 
+        fn encode_ordinary(&self, text: &str) -> vllm_tokenizer::Result<Vec<u32>> {
+            self.encode(text, false)
+        }
+
         fn decode(
             &self,
             _token_ids: &[u32],
diff --git a/rust/src/text/src/error.rs b/rust/src/text/src/error.rs
index e72e68196c3..6b385fad68d 100644
--- a/rust/src/text/src/error.rs
+++ b/rust/src/text/src/error.rs
@@ -6,6 +6,7 @@ use vllm_engine_core_client::Error as EngineCoreError;
 use vllm_llm::Error as LlmError;
 
 pub use crate::lower::logprobs::LogprobsError;
+pub use crate::lower::sampling::SamplingParamsError;
 pub use crate::lower::token_ids::TokenIdsError;
 
 #[derive(Debug, Error)]
@@ -23,6 +24,8 @@ pub enum Error {
     Logprobs(#[from] LogprobsError),
     #[error(transparent)]
     TokenIds(#[from] TokenIdsError),
+    #[error(transparent)]
+    SamplingParams(#[from] SamplingParamsError),
     #[error(
         "`min_tokens` must be less than or equal to `max_tokens`, \
          got min_tokens={min_tokens}, max_tokens={max_tokens}"
@@ -50,6 +53,7 @@ impl Error {
             | Self::EmptyPromptTokenIds { .. }
             | Self::Logprobs(_)
             | Self::TokenIds(_)
+            | Self::SamplingParams(_)
             | Self::MinTokensExceedsMaxTokens { .. }
             | Self::InvalidThinkingTokenBudget
             | Self::InvalidRepetitionDetection { .. }
diff --git a/rust/src/text/src/lib.rs b/rust/src/text/src/lib.rs
index b00155999e8..646e4bac9f3 100644
--- a/rust/src/text/src/lib.rs
+++ b/rust/src/text/src/lib.rs
@@ -10,7 +10,7 @@
 use std::mem::take;
 
 pub use backend::{DynTextBackend, SamplingHints, SamplingLimits, TextBackend};
-pub use error::{Error, LogprobsError, Result, TokenIdsError};
+pub use error::{Error, LogprobsError, Result, SamplingParamsError, TokenIdsError};
 use futures::Stream;
 pub use lower::{
     PreparedTextRequest, lower_sampling_params, lower_text_request, resolve_max_tokens,
diff --git a/rust/src/text/src/lower.rs b/rust/src/text/src/lower.rs
index bd43a1d141c..aa54595acda 100644
--- a/rust/src/text/src/lower.rs
+++ b/rust/src/text/src/lower.rs
@@ -4,9 +4,11 @@
 use std::collections::BTreeSet;
 
 pub(crate) mod logprobs;
+pub(crate) mod sampling;
 pub(crate) mod token_ids;
 
 use logprobs::validate_logprobs;
+use sampling::validate_resolved_sampling_params;
 use token_ids::{validate_prompt_token_ids, validate_vocab_range};
 use vllm_engine_core_client::protocol::sampling::{
     EngineCoreSamplingParams, RepetitionDetectionParams,
@@ -186,6 +188,7 @@ pub fn lower_sampling_params(
         skip_reading_prefix_cache,
         extra_args: vllm_xargs,
     };
+    validate_resolved_sampling_params(&params)?;
     validate_vocab_range(&params, &sampling_limits)?;
     Ok(params)
 }
@@ -319,7 +322,7 @@ mod tests {
     use super::*;
     use crate::backend::hf::HfTextBackend;
     use crate::backend::{SamplingHints, TextBackend as _};
-    use crate::error::{LogprobsError, TokenIdsError};
+    use crate::error::{LogprobsError, SamplingParamsError, TokenIdsError};
     use crate::request::{Prompt, TextRequest};
 
     fn stub_tokenizer() -> TestTokenizer {
@@ -482,6 +485,120 @@ mod tests {
         assert!(message.contains("min_count=1"));
     }
 
+    #[test]
+    fn lower_sampling_params_rejects_invalid_sampling_ranges() {
+        let cases = [
+            (
+                "temperature",
+                SamplingParams {
+                    temperature: Some(5.0),
+                    ..SamplingParams::default()
+                },
+            ),
+            (
+                "top_p",
+                SamplingParams {
+                    top_p: Some(0.0),
+                    ..SamplingParams::default()
+                },
+            ),
+            (
+                "min_p",
+                SamplingParams {
+                    min_p: Some(2.0),
+                    ..SamplingParams::default()
+                },
+            ),
+            (
+                "repetition_penalty",
+                SamplingParams {
+                    repetition_penalty: Some(0.0),
+                    ..SamplingParams::default()
+                },
+            ),
+            (
+                "frequency_penalty",
+                SamplingParams {
+                    frequency_penalty: Some(100.0),
+                    ..SamplingParams::default()
+                },
+            ),
+            (
+                "presence_penalty",
+                SamplingParams {
+                    presence_penalty: Some(100.0),
+                    ..SamplingParams::default()
+                },
+            ),
+        ];
+
+        for (expected_parameter, sampling_params) in cases {
+            let error =
+                lower_sampling_params_with_limits(sampling_params, sample_sampling_limits())
+                    .unwrap_err();
+
+            assert!(
+                matches!(
+                    error,
+                    Error::SamplingParams(SamplingParamsError::OutOfRange {
+                        parameter,
+                        ..
+                    }) if parameter == expected_parameter
+                ),
+                "{expected_parameter} should be rejected"
+            );
+        }
+    }
+
+    #[test]
+    fn lower_sampling_params_rejects_non_finite_sampling_values() {
+        for (expected_parameter, sampling_params) in [
+            (
+                "temperature",
+                SamplingParams {
+                    temperature: Some(f32::INFINITY),
+                    ..SamplingParams::default()
+                },
+            ),
+            (
+                "repetition_penalty",
+                SamplingParams {
+                    repetition_penalty: Some(f32::NAN),
+                    ..SamplingParams::default()
+                },
+            ),
+        ] {
+            let error =
+                lower_sampling_params_with_limits(sampling_params, sample_sampling_limits())
+                    .unwrap_err();
+
+            assert!(
+                matches!(
+                    error,
+                    Error::SamplingParams(SamplingParamsError::NotFinite {
+                        parameter,
+                        ..
+                    }) if parameter == expected_parameter
+                ),
+                "{expected_parameter} should reject non-finite values"
+            );
+        }
+    }
+
+    #[test]
+    fn lower_sampling_params_accepts_python_compatible_repetition_penalty_above_two() {
+        let params = lower_sampling_params_with_limits(
+            SamplingParams {
+                repetition_penalty: Some(2.5),
+                ..SamplingParams::default()
+            },
+            sample_sampling_limits(),
+        )
+        .unwrap();
+
+        assert_eq!(params.repetition_penalty, 2.5);
+    }
+
     #[test]
     fn lower_text_request_applies_python_style_eos_hints() {
         let prepared = lower_text_request(
diff --git a/rust/src/text/src/lower/sampling.rs b/rust/src/text/src/lower/sampling.rs
new file mode 100644
index 00000000000..edcdc4b0492
--- /dev/null
+++ b/rust/src/text/src/lower/sampling.rs
@@ -0,0 +1,92 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use thiserror::Error;
+use vllm_engine_core_client::protocol::sampling::EngineCoreSamplingParams;
+
+#[derive(Debug, Error, PartialEq)]
+pub enum SamplingParamsError {
+    #[error("{parameter} must be a finite number, got {value}")]
+    NotFinite { parameter: &'static str, value: f32 },
+    #[error("{parameter} must be in {expected}, got {value}")]
+    OutOfRange {
+        parameter: &'static str,
+        value: f32,
+        expected: &'static str,
+    },
+}
+
+fn validate_frequency_penalty(value: f32) -> Result<(), SamplingParamsError> {
+    validate_closed_range("frequency_penalty", value, -2.0, 2.0, "[-2, 2]")
+}
+
+fn validate_presence_penalty(value: f32) -> Result<(), SamplingParamsError> {
+    validate_closed_range("presence_penalty", value, -2.0, 2.0, "[-2, 2]")
+}
+
+fn validate_temperature(value: f32) -> Result<(), SamplingParamsError> {
+    validate_finite("temperature", value)?;
+    validate_closed_range("temperature", value, 0.0, 2.0, "[0, 2]")
+}
+
+fn validate_top_p(value: f32) -> Result<(), SamplingParamsError> {
+    if value > 0.0 && value <= 1.0 {
+        return Ok(());
+    }
+    Err(SamplingParamsError::OutOfRange {
+        parameter: "top_p",
+        value,
+        expected: "(0, 1]",
+    })
+}
+
+fn validate_min_p(value: f32) -> Result<(), SamplingParamsError> {
+    validate_closed_range("min_p", value, 0.0, 1.0, "[0, 1]")
+}
+
+fn validate_repetition_penalty(value: f32) -> Result<(), SamplingParamsError> {
+    validate_finite("repetition_penalty", value)?;
+    if value > 0.0 {
+        return Ok(());
+    }
+    Err(SamplingParamsError::OutOfRange {
+        parameter: "repetition_penalty",
+        value,
+        expected: "(0, inf)",
+    })
+}
+
+pub(crate) fn validate_resolved_sampling_params(
+    params: &EngineCoreSamplingParams,
+) -> Result<(), SamplingParamsError> {
+    validate_temperature(params.temperature)?;
+    validate_top_p(params.top_p)?;
+    validate_min_p(params.min_p)?;
+    validate_frequency_penalty(params.frequency_penalty)?;
+    validate_presence_penalty(params.presence_penalty)?;
+    validate_repetition_penalty(params.repetition_penalty)
+}
+
+fn validate_finite(parameter: &'static str, value: f32) -> Result<(), SamplingParamsError> {
+    if value.is_finite() {
+        return Ok(());
+    }
+    Err(SamplingParamsError::NotFinite { parameter, value })
+}
+
+fn validate_closed_range(
+    parameter: &'static str,
+    value: f32,
+    min: f32,
+    max: f32,
+    expected: &'static str,
+) -> Result<(), SamplingParamsError> {
+    if value >= min && value <= max {
+        return Ok(());
+    }
+    Err(SamplingParamsError::OutOfRange {
+        parameter,
+        value,
+        expected,
+    })
+}
diff --git a/rust/src/tokenizer/src/hf.rs b/rust/src/tokenizer/src/hf.rs
index 08ec5a22d6b..eb527294a5b 100644
--- a/rust/src/tokenizer/src/hf.rs
+++ b/rust/src/tokenizer/src/hf.rs
@@ -1,13 +1,20 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+use std::borrow::Cow;
 use std::path::Path;
-use std::sync::Arc;
+use std::sync::{Arc, LazyLock};
 
 use fastokens::Tokenizer as FastokensTokenizer;
 use fastokens::decoders::Decoder as FastokensDecoder;
+use fastokens::pre_tokenized::{
+    PreTokenizedString as FastokensPreTokenizedString, Split as FastokensSplit,
+};
+use fastokens::{PreTokenizer as FastokensPreTokenizer, Split as FastokensSplitPreTokenizer};
 use thiserror_ext::AsReport as _;
-use tokenizers::Tokenizer as HfTokenizer;
+use tokenizers::{
+    AddedVocabulary, Model as _, OffsetType, PreTokenizer as _, Tokenizer as HfTokenizer,
+};
 use tracing::{info, warn};
 
 use crate::byte_level_decode::decode_byte_level;
@@ -16,6 +23,8 @@ use crate::{Result, Tokenizer};
 
 mod added_tokens;
 
+static EMPTY_HF_ADDED_VOCABULARY: LazyLock<AddedVocabulary> = LazyLock::new(AddedVocabulary::new);
+
 enum Backend {
     Hf(Box<HfTokenizer>),
     Fastokens(Box<FastokensTokenizer>),
@@ -53,6 +62,85 @@ fn decode_fastokens_byte_level(
     Ok(decode_byte_level(tokens))
 }
 
+fn encode_hf_ordinary(tokenizer: &HfTokenizer, text: &str) -> tokenizers::Result<Vec<u32>> {
+    let mut pretokenized =
+        EMPTY_HF_ADDED_VOCABULARY.extract_and_normalize(tokenizer.get_normalizer(), text);
+
+    if let Some(pre_tokenizer) = tokenizer.get_pre_tokenizer() {
+        pre_tokenizer.pre_tokenize(&mut pretokenized)?;
+    }
+    pretokenized.tokenize(|normalized| tokenizer.get_model().tokenize(normalized.get()))?;
+    let encoding = pretokenized.into_encoding(None, 0, OffsetType::Byte)?;
+    let encoding = tokenizer.post_process(encoding, None, false)?;
+    Ok(encoding.get_ids().to_vec())
+}
+
+fn fastokens_fused_split(tokenizer: &FastokensTokenizer) -> Option<&FastokensSplitPreTokenizer> {
+    // Keep this predicate aligned with fastokens::Tokenizer::detect_fused_byte_level.
+    let FastokensPreTokenizer::Sequence(steps) = tokenizer.pre_tokenizer()? else {
+        return None;
+    };
+    let [
+        FastokensPreTokenizer::Split(split),
+        FastokensPreTokenizer::ByteLevel(byte_level),
+    ] = steps.as_slice()
+    else {
+        return None;
+    };
+    byte_level.is_bulk_only().then_some(split)
+}
+
+fn fastokens_pre_tokenized_ordinary(
+    tokenizer: &FastokensTokenizer,
+    text: &str,
+) -> FastokensPreTokenizedString {
+    // This is fastokens::Tokenizer::build_pre_tokenized with added_tokens = None.
+    let normalized = tokenizer
+        .normalizer()
+        .map_or(Cow::Borrowed(text), |normalizer| normalizer.normalize(text));
+    match normalized {
+        Cow::Borrowed(_) => FastokensPreTokenizedString::from_text(text),
+        Cow::Owned(text) => {
+            let len = text.len();
+            FastokensPreTokenizedString::new(
+                text,
+                vec![FastokensSplit {
+                    range: 0..len,
+                    token_id: None,
+                }],
+            )
+        }
+    }
+}
+
+fn encode_fastokens_ordinary(
+    tokenizer: &FastokensTokenizer,
+    text: &str,
+) -> std::result::Result<Vec<u32>, fastokens::Error> {
+    if text.is_empty() {
+        return Ok(Vec::new());
+    }
+
+    let mut pretokenized = fastokens_pre_tokenized_ordinary(tokenizer, text);
+    let ids = if let Some(split) = fastokens_fused_split(tokenizer) {
+        split.pre_tokenize(&mut pretokenized)?;
+        pretokenized
+            .tokenize_batched(|buffer, splits, output| {
+                tokenizer.model().tokenize_batch_fused(buffer, splits, output)
+            })
+            .map_err(fastokens::Error::Model)?
+    } else {
+        if let Some(pre_tokenizer) = tokenizer.pre_tokenizer() {
+            pre_tokenizer.pre_tokenize(&mut pretokenized)?;
+        }
+        pretokenized
+            .tokenize(|text, output| tokenizer.model().tokenize_into(text, output))
+            .map_err(fastokens::Error::Model)?
+    };
+
+    Ok(tokenizer.post_process(ids, false))
+}
+
 /// Tokenizer from `tokenizer.json` in HuggingFace format.
 ///
 /// This tries to load with `fastokens` first for better performance, then falls
@@ -156,6 +244,17 @@ impl Tokenizer for HuggingFaceTokenizer {
         }
     }
 
+    fn encode_ordinary(&self, text: &str) -> Result<Vec<u32>> {
+        match &self.backend {
+            Backend::Hf(tokenizer) => encode_hf_ordinary(tokenizer, text)
+                .map_err(|error| tokenizer_error!("encoding failed: {}", error.as_report())),
+            Backend::Fastokens(tokenizer) | Backend::FastokensByteLevel(tokenizer) => {
+                encode_fastokens_ordinary(tokenizer, text)
+                    .map_err(|error| tokenizer_error!("encoding failed: {}", error.as_report()))
+            }
+        }
+    }
+
     fn decode(&self, token_ids: &[u32], skip_special_tokens: bool) -> Result<String> {
         match &self.backend {
             Backend::Hf(t) => t
@@ -200,12 +299,19 @@ impl Tokenizer for HuggingFaceTokenizer {
 
 #[cfg(test)]
 mod tests {
+    use std::path::{Path, PathBuf};
+
+    use serde_json::{Value, json};
     use tempfile::tempdir;
     use tokenizers::models::bpe::BPE;
+    use tokenizers::pre_tokenizers::byte_level::ByteLevel;
     use tokenizers::{AddedToken, Tokenizer as HfTokenizer};
 
     use super::{HuggingFaceTokenizer, Tokenizer};
 
+    const REGULAR_TOKEN: &str = "<|regular|>";
+    const SPECIAL_TOKEN: &str = "<|special|>";
+
     fn tiny_bpe_tokenizer() -> HfTokenizer {
         let vocab = [
             ("<unk>".to_string(), 0),
@@ -232,6 +338,186 @@ mod tests {
         HfTokenizer::new(model)
     }
 
+    fn ordinary_test_tokenizer_json(fused: bool, with_added_tokens: bool) -> Value {
+        let mut alphabet: Vec<char> = ByteLevel::alphabet().into_iter().collect();
+        alphabet.sort_unstable();
+        let vocab = alphabet
+            .into_iter()
+            .enumerate()
+            .map(|(id, token)| (token.to_string(), json!(id)))
+            .collect::<serde_json::Map<_, _>>();
+
+        let pre_tokenizer = if fused {
+            json!({
+                "type": "Sequence",
+                "pretokenizers": [
+                    {
+                        "type": "Split",
+                        "pattern": {"Regex": "\\S+|\\s+"},
+                        "behavior": "Isolated",
+                        "invert": false
+                    },
+                    {
+                        "type": "ByteLevel",
+                        "add_prefix_space": false,
+                        "trim_offsets": true,
+                        "use_regex": false
+                    }
+                ]
+            })
+        } else {
+            json!({
+                "type": "ByteLevel",
+                "add_prefix_space": false,
+                "trim_offsets": true,
+                "use_regex": true
+            })
+        };
+        let added_tokens = with_added_tokens.then(|| {
+            json!([
+                {
+                    "id": 256,
+                    "content": REGULAR_TOKEN,
+                    "single_word": false,
+                    "lstrip": false,
+                    "rstrip": false,
+                    "normalized": true,
+                    "special": false
+                },
+                {
+                    "id": 257,
+                    "content": SPECIAL_TOKEN,
+                    "single_word": false,
+                    "lstrip": false,
+                    "rstrip": false,
+                    "normalized": false,
+                    "special": true
+                }
+            ])
+        });
+
+        json!({
+            "version": "1.0",
+            "truncation": {
+                "direction": "Right",
+                "max_length": 24,
+                "strategy": "LongestFirst",
+                "stride": 0
+            },
+            "padding": null,
+            "added_tokens": added_tokens.unwrap_or_else(|| json!([])),
+            "normalizer": {"type": "NFC"},
+            "pre_tokenizer": pre_tokenizer,
+            "post_processor": {
+                "type": "ByteLevel",
+                "add_prefix_space": false,
+                "trim_offsets": true,
+                "use_regex": true
+            },
+            "decoder": {
+                "type": "ByteLevel",
+                "add_prefix_space": false,
+                "trim_offsets": true,
+                "use_regex": true
+            },
+            "model": {
+                "type": "BPE",
+                "dropout": null,
+                "unk_token": null,
+                "continuing_subword_prefix": null,
+                "end_of_word_suffix": null,
+                "fuse_unk": false,
+                "byte_fallback": false,
+                "ignore_merges": false,
+                "vocab": vocab,
+                "merges": []
+            }
+        })
+    }
+
+    fn write_tokenizer_json(dir: &Path, name: &str, value: &Value) -> PathBuf {
+        let path = dir.join(name);
+        std::fs::write(
+            &path,
+            serde_json::to_vec(value).expect("serialize tokenizer"),
+        )
+        .expect("write tokenizer");
+        path
+    }
+
+    fn assert_ordinary_matches_added_empty(
+        constructor: fn(&Path) -> crate::Result<HuggingFaceTokenizer>,
+        fused: bool,
+    ) {
+        let dir = tempdir().expect("create temp dir");
+        let added_path = write_tokenizer_json(
+            dir.path(),
+            "with-added.json",
+            &ordinary_test_tokenizer_json(fused, true),
+        );
+        let empty_path = write_tokenizer_json(
+            dir.path(),
+            "added-empty.json",
+            &ordinary_test_tokenizer_json(fused, false),
+        );
+        let tokenizer = constructor(&added_path).expect("load tokenizer with added tokens");
+        let added_empty = constructor(&empty_path).expect("load tokenizer with empty added tokens");
+
+        if let super::Backend::Fastokens(inner) | super::Backend::FastokensByteLevel(inner) =
+            &tokenizer.backend
+        {
+            assert_eq!(super::fastokens_fused_split(inner).is_some(), fused);
+        }
+
+        assert_eq!(
+            tokenizer.encode(REGULAR_TOKEN, false).unwrap(),
+            vec![tokenizer.token_to_id(REGULAR_TOKEN).unwrap()]
+        );
+        assert_eq!(
+            tokenizer.encode(SPECIAL_TOKEN, false).unwrap(),
+            vec![tokenizer.token_to_id(SPECIAL_TOKEN).unwrap()]
+        );
+
+        for text in [
+            "",
+            "hello",
+            "Cafe\u{301}",
+            REGULAR_TOKEN,
+            SPECIAL_TOKEN,
+            "hello <|regular|> Cafe\u{301} <|special|> tail",
+        ] {
+            assert_eq!(
+                tokenizer.encode_ordinary(text).unwrap(),
+                added_empty.encode(text, false).unwrap(),
+                "fused={fused}, text={text:?}",
+            );
+        }
+        if matches!(&tokenizer.backend, super::Backend::Hf(_)) {
+            assert_eq!(
+                tokenizer
+                    .encode_ordinary("hello <|regular|> Cafe\u{301} <|special|> tail")
+                    .unwrap()
+                    .len(),
+                24,
+                "HF post-processing must retain configured truncation",
+            );
+        }
+    }
+
+    #[test]
+    fn hf_ordinary_matches_original_encode_with_added_empty() {
+        for fused in [false, true] {
+            assert_ordinary_matches_added_empty(HuggingFaceTokenizer::new_hf, fused);
+        }
+    }
+
+    #[test]
+    fn fastokens_ordinary_matches_original_encode_with_added_empty() {
+        for fused in [false, true] {
+            assert_ordinary_matches_added_empty(HuggingFaceTokenizer::new_fastokens, fused);
+        }
+    }
+
     #[test]
     fn hf_constructor_resolves_added_token_ids() {
         let mut tokenizer = tiny_bpe_tokenizer();
diff --git a/rust/src/tokenizer/src/incremental.rs b/rust/src/tokenizer/src/incremental.rs
index 5e470ae4b00..f608fa874b8 100644
--- a/rust/src/tokenizer/src/incremental.rs
+++ b/rust/src/tokenizer/src/incremental.rs
@@ -199,6 +199,10 @@ mod tests {
             unreachable!()
         }
 
+        fn encode_ordinary(&self, _text: &str) -> Result<Vec<u32>> {
+            unreachable!()
+        }
+
         fn decode(&self, token_ids: &[u32], _skip_special_tokens: bool) -> Result<String> {
             let bytes = token_ids.iter().map(|id| *id as u8).collect::<Vec<_>>();
             Ok(String::from_utf8_lossy(&bytes).into_owned())
@@ -273,6 +277,10 @@ mod tests {
             unreachable!()
         }
 
+        fn encode_ordinary(&self, _text: &str) -> Result<Vec<u32>> {
+            unreachable!()
+        }
+
         fn decode(&self, token_ids: &[u32], skip_special_tokens: bool) -> Result<String> {
             let mut text = String::new();
             for &token_id in token_ids {
@@ -410,6 +418,10 @@ mod tests {
             unreachable!()
         }
 
+        fn encode_ordinary(&self, _text: &str) -> Result<Vec<u32>> {
+            unreachable!()
+        }
+
         fn decode(&self, token_ids: &[u32], _skip_special_tokens: bool) -> Result<String> {
             match token_ids {
                 [1] => Ok("abc".into()),
diff --git a/rust/src/tokenizer/src/lib.rs b/rust/src/tokenizer/src/lib.rs
index 6c9fcd3fdea..0f8c7dc16e5 100644
--- a/rust/src/tokenizer/src/lib.rs
+++ b/rust/src/tokenizer/src/lib.rs
@@ -25,6 +25,10 @@ pub trait Tokenizer: Send + Sync {
     /// Encode one prompt string into token IDs.
     fn encode(&self, text: &str, add_special_tokens: bool) -> Result<Vec<u32>>;
 
+    /// Equivalent to `encode(text, false)`, except that every added,
+    /// special, and control-token matcher is bypassed.
+    fn encode_ordinary(&self, text: &str) -> Result<Vec<u32>>;
+
     /// Decode one token sequence into text.
     fn decode(&self, token_ids: &[u32], skip_special_tokens: bool) -> Result<String>;
 
diff --git a/rust/src/tokenizer/src/tekken.rs b/rust/src/tokenizer/src/tekken.rs
index 20b6c26ffc8..5f9342e4a5e 100644
--- a/rust/src/tokenizer/src/tekken.rs
+++ b/rust/src/tokenizer/src/tekken.rs
@@ -35,6 +35,12 @@ impl Tokenizer for TekkenTokenizer {
             .map_err(|error| tokenizer_error!("encoding failed: {error}"))
     }
 
+    fn encode_ordinary(&self, text: &str) -> Result<Vec<u32>> {
+        self.inner
+            .encode(text, false, false)
+            .map_err(|error| tokenizer_error!("encoding failed: {error}"))
+    }
+
     fn decode(&self, token_ids: &[u32], skip_special_tokens: bool) -> Result<String> {
         let policy = if skip_special_tokens {
             tekken::SpecialTokenPolicy::Ignore
@@ -67,3 +73,51 @@ impl Tokenizer for TekkenTokenizer {
         self.inner.is_special_token(token_id)
     }
 }
+
+#[cfg(test)]
+mod tests {
+    use base64::Engine as _;
+    use tekken::config::TokenizerVersion;
+    use tekken::{SpecialTokenInfo, TokenInfo};
+
+    use super::*;
+
+    fn test_tokenizer() -> TekkenTokenizer {
+        let vocab = (0_u8..=255)
+            .map(|byte| TokenInfo {
+                rank: byte as usize,
+                token_bytes: base64::engine::general_purpose::STANDARD.encode([byte]),
+                token_str: None,
+            })
+            .collect();
+        let special_tokens = vec![SpecialTokenInfo {
+            rank: 0,
+            token_str: "<control>".to_string(),
+            is_control: true,
+        }];
+        let inner = Tekkenizer::new(
+            vocab,
+            &special_tokens,
+            r"(?s).",
+            257,
+            1,
+            TokenizerVersion::V3,
+            None,
+        )
+        .expect("build Tekken tokenizer");
+        TekkenTokenizer { inner }
+    }
+
+    #[test]
+    fn ordinary_matches_tekkens_empty_special_encoding() {
+        let tokenizer = test_tokenizer();
+        let text = "user <control> text";
+        let control_id = tokenizer.token_to_id("<control>").unwrap();
+        let ordinary_ids = tokenizer.encode_ordinary(text).unwrap();
+
+        assert_eq!(control_id, 0);
+        assert_eq!(ordinary_ids, tokenizer.encode(text, false).unwrap());
+        assert!(!ordinary_ids.contains(&control_id));
+        assert_eq!(tokenizer.decode(&ordinary_ids, false).unwrap(), text);
+    }
+}
diff --git a/rust/src/tokenizer/src/test_utils.rs b/rust/src/tokenizer/src/test_utils.rs
index 36d1f3d18aa..6fdb6e02721 100644
--- a/rust/src/tokenizer/src/test_utils.rs
+++ b/rust/src/tokenizer/src/test_utils.rs
@@ -208,6 +208,10 @@ impl Tokenizer for TestTokenizer {
         Ok(ids)
     }
 
+    fn encode_ordinary(&self, text: &str) -> Result<Vec<u32>> {
+        Ok(text.as_bytes().iter().copied().map(u32::from).collect())
+    }
+
     fn decode(&self, token_ids: &[u32], skip_special_tokens: bool) -> Result<String> {
         let mut output = String::new();
         let mut pending_bytes = Vec::new();
@@ -374,6 +378,32 @@ mod tests {
         assert!(!tokenizer.is_special_id(0xF002));
     }
 
+    #[test]
+    fn ordinary_encoding_bypasses_all_configured_tokens() {
+        let tokenizer = TestTokenizer::new()
+            .with_bos_token("<bos>", 256)
+            .with_special_token("<control>", 257)
+            .with_regular_token("<visible>", 258);
+        let ordinary_text = "user <control> and <visible>";
+
+        assert_eq!(tokenizer.encode("<control>", false).unwrap(), vec![257]);
+        assert_eq!(tokenizer.encode("<visible>", false).unwrap(), vec![258]);
+        assert_eq!(
+            tokenizer.encode_ordinary(ordinary_text).unwrap(),
+            ordinary_text.as_bytes().iter().copied().map(u32::from).collect::<Vec<_>>()
+        );
+
+        let mut segmented = tokenizer.encode("<control>", false).unwrap();
+        segmented.extend(tokenizer.encode_ordinary(ordinary_text).unwrap());
+        segmented.extend(tokenizer.encode("<visible>", false).unwrap());
+        assert_eq!(segmented.first(), Some(&257));
+        assert_eq!(segmented.last(), Some(&258));
+        assert_eq!(
+            tokenizer.decode(&segmented, false).unwrap(),
+            format!("<control>{ordinary_text}<visible>")
+        );
+    }
+
     #[test]
     #[should_panic(expected = "configured test token id 255 overlaps byte fallback range 0..=255")]
     fn configured_token_id_must_stay_outside_byte_range() {
diff --git a/rust/src/tokenizer/src/tiktoken.rs b/rust/src/tokenizer/src/tiktoken.rs
index b204d2a9b58..fb569d69c35 100644
--- a/rust/src/tokenizer/src/tiktoken.rs
+++ b/rust/src/tokenizer/src/tiktoken.rs
@@ -462,6 +462,13 @@ impl Tokenizer for TiktokenTokenizer {
         })
     }
 
+    fn encode_ordinary(&self, text: &str) -> Result<Vec<u32>> {
+        Ok(match &self.backend {
+            Backend::Riptoken(backend) => backend.inner.encode_ordinary(text),
+            Backend::TiktokenRs(backend) => backend.inner.encode_ordinary(text),
+        })
+    }
+
     fn decode(&self, token_ids: &[u32], skip_special_tokens: bool) -> Result<String> {
         // Filter passes:
         //
@@ -752,6 +759,39 @@ mod tests {
         }
     }
 
+    #[test]
+    fn tiktoken_ordinary_bypasses_every_registered_added_token() {
+        let dir = tempfile::tempdir().expect("create temp dir");
+        let bpe_path = write_synthetic_bpe_file(dir.path());
+        fs::write(
+            dir.path().join("tokenizer_config.json"),
+            r#"{
+                "added_tokens_decoder": {
+                    "257": { "content": "<|im_end|>", "special": true },
+                    "258": { "content": "<|tool_call_begin|>", "special": false }
+                }
+            }"#,
+        )
+        .expect("write tokenizer_config.json");
+        fs::write(dir.path().join("config.json"), r#"{"vocab_size": 260}"#)
+            .expect("write config.json");
+
+        let input = "<|im_end|><|tool_call_begin|><|reserved_token_259|>";
+        let expected: Vec<u32> = input.as_bytes().iter().copied().map(u32::from).collect();
+        for backend in explicit_backends(&bpe_path) {
+            assert_eq!(backend.encode("<|im_end|>", false).unwrap(), vec![257]);
+            assert_eq!(
+                backend.encode("<|tool_call_begin|>", false).unwrap(),
+                vec![258]
+            );
+            assert_eq!(
+                backend.encode("<|reserved_token_259|>", false).unwrap(),
+                vec![259]
+            );
+            assert_eq!(backend.encode_ordinary(input).unwrap(), expected);
+        }
+    }
+
     /// `vocab_size` may live under `text_config` for composite (e.g.
     /// multimodal) configs.
     #[test]
diff --git a/rust/src/tracing/Cargo.toml b/rust/src/tracing/Cargo.toml
new file mode 100644
index 00000000000..e003720814e
--- /dev/null
+++ b/rust/src/tracing/Cargo.toml
@@ -0,0 +1,14 @@
+[package]
+name = "vllm-tracing"
+version.workspace = true
+edition.workspace = true
+description = "Shared tracing subscriber and log formatting for vLLM Rust binaries"
+license.workspace = true
+
+[dependencies]
+time.workspace = true
+tracing.workspace = true
+tracing-subscriber.workspace = true
+
+[lints]
+workspace = true
diff --git a/rust/src/cmd/src/logging.rs b/rust/src/tracing/src/lib.rs
similarity index 98%
rename from rust/src/cmd/src/logging.rs
rename to rust/src/tracing/src/lib.rs
index 936b3692442..eb6da9c0b24 100644
--- a/rust/src/cmd/src/logging.rs
+++ b/rust/src/tracing/src/lib.rs
@@ -1,6 +1,8 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+//! Shared tracing subscriber and log formatting for vLLM Rust binaries.
+
 use std::{env, fmt, process};
 
 use time::UtcOffset;
@@ -26,8 +28,8 @@ const RESET: &str = "\x1b[0m";
 const VLLM_TIME_FORMAT: &[time::format_description::FormatItem<'static>] =
     format_description!("[month]-[day] [hour]:[minute]:[second]");
 
-/// Install the process-wide vLLM-style tracing subscriber for the CLI binary.
-pub(crate) fn init_tracing(process_label: &str) {
+/// Install the process-wide vLLM-style tracing subscriber.
+pub fn init_tracing(process_label: &str) {
     let filter = build_targets_filter(
         env::var("VLLM_LOGGING_LEVEL").ok().as_deref(),
         env::var("RUST_LOG").ok().as_deref(),
diff --git a/tests/compile/passes/distributed/test_fusion_all_reduce.py b/tests/compile/passes/distributed/test_fusion_all_reduce.py
index 1aac4b2bec4..e9c8d0deaa7 100644
--- a/tests/compile/passes/distributed/test_fusion_all_reduce.py
+++ b/tests/compile/passes/distributed/test_fusion_all_reduce.py
@@ -272,12 +272,10 @@ class TestAiterAllReduceRMSNormGroupQuantFP8Model(torch.nn.Module):
         token_num=16,
         eps=1e-6,
         dtype: torch.dtype = torch.bfloat16,
-        use_triton_quant: bool = False,
     ):
         super().__init__()
         self.hidden_size = hidden_size
         self.eps = eps
-        self.use_triton_quant = use_triton_quant
         assert hidden_size % self.quant_group_size == 0, (
             f"hidden_size ({hidden_size}) must be a multiple of "
             f"quant_group_size ({self.quant_group_size}) for per-group FP8 quant"
@@ -289,10 +287,6 @@ class TestAiterAllReduceRMSNormGroupQuantFP8Model(torch.nn.Module):
         ]
 
     def _group_quant(self, rms: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor]:
-        if self.use_triton_quant:
-            return torch.ops.vllm.triton_per_token_group_quant_fp8(
-                rms, self.quant_group_size
-            )
         return torch.ops.vllm.rocm_aiter_group_fp8_quant.default(
             rms, self.quant_group_size
         )
@@ -339,11 +333,7 @@ class TestAiterAllReduceRMSNormGroupQuantFP8Model(torch.nn.Module):
     def ops_in_model_before(self):
         return [
             torch.ops.vllm.all_reduce.default,
-            (
-                torch.ops.vllm.triton_per_token_group_quant_fp8.default
-                if self.use_triton_quant
-                else torch.ops.vllm.rocm_aiter_group_fp8_quant.default
-            ),
+            torch.ops.vllm.rocm_aiter_group_fp8_quant.default,
         ]
 
     def ops_in_model_after(self):
@@ -646,7 +636,6 @@ def all_reduce_fusion_pass_on_test_model(
 
 
 @multi_gpu_test(num_gpus=2)
-@pytest.mark.parametrize("use_triton_quant", [True, False])
 @pytest.mark.parametrize("batch_size", [8])
 @pytest.mark.parametrize("seq_len", [8])
 @pytest.mark.parametrize("hidden_size", [128])
@@ -663,7 +652,6 @@ def test_rocm_aiter_all_reduce_rmsnorm_group_quant_fp8_fusion_pass_replace(
     hidden_size: int,
     dtype: torch.dtype,
     enable_rms_norm_custom_op: bool,
-    use_triton_quant: bool,
     monkeypatch: pytest.MonkeyPatch,
 ):
     """Sibling of ``test_all_reduce_fusion_pass_replace`` for the new
@@ -676,9 +664,9 @@ def test_rocm_aiter_all_reduce_rmsnorm_group_quant_fp8_fusion_pass_replace(
     * ``AiterAllreduceFusedAddRMSNormGroupQuantFP8Pattern`` (with-residual,
       single ``rms`` consumer)
     * ``AiterAllreduceFusedAddRMSNormGroupQuantWithIndexerPattern`` (with-
-      residual, DSv3.2 indexer fan-out; parametrized over both
-      ``triton_per_token_group_quant_fp8`` and ``rocm_aiter_group_fp8_quant``
-      producers).
+      residual, DSv3.2 indexer fan-out; parametrized over
+      ``rocm_aiter_group_fp8_quant``
+      producer).
     """
     with monkeypatch.context() as m:
         m.setenv("VLLM_ROCM_USE_AITER", "1")
@@ -703,7 +691,6 @@ def test_rocm_aiter_all_reduce_rmsnorm_group_quant_fp8_fusion_pass_replace(
                 hidden_size,
                 dtype,
                 enable_rms_norm_custom_op,
-                use_triton_quant,
                 monkeypatch,
             ),
             nprocs=nprocs,
@@ -721,7 +708,6 @@ def rocm_aiter_group_quant_fusion_pass_on_test_model(
     hidden_size: int,
     dtype: torch.dtype,
     enable_rms_norm_custom_op: bool,
-    use_triton_quant: bool,
     monkeypatch: pytest.MonkeyPatch,
 ):
     set_random_seed(0)
@@ -749,10 +735,7 @@ def rocm_aiter_group_quant_fusion_pass_on_test_model(
     custom_ops = []
     if enable_rms_norm_custom_op:
         custom_ops.append("+rms_norm")
-    # ``triton_per_token_group_quant_fp8`` is emitted by ``QuantFP8.forward_hip``
-    # only when QuantFP8 is enabled as a custom op (and ``use_triton=True`` at
-    # the call site). The patterns in this PR are robust to both Triton and
-    # rocm_aiter forms; we always enable +quant_fp8 so the matcher's example
+    # We always enable +quant_fp8 so the matcher's example
     # trace finds the same form the test model uses.
     custom_ops.append("+quant_fp8")
 
@@ -783,9 +766,7 @@ def rocm_aiter_group_quant_fusion_pass_on_test_model(
         )
 
         token_num = batch_size * seq_len
-        model = test_model_cls(
-            hidden_size, token_num, dtype=dtype, use_triton_quant=use_triton_quant
-        )
+        model = test_model_cls(hidden_size, token_num, dtype=dtype)
 
         hidden_states = torch.randn((token_num, hidden_size), requires_grad=False)
 
diff --git a/tests/compile/passes/test_fusion.py b/tests/compile/passes/test_fusion.py
index 92d1902b2c2..591b014d9e2 100644
--- a/tests/compile/passes/test_fusion.py
+++ b/tests/compile/passes/test_fusion.py
@@ -195,8 +195,6 @@ class TestModel(torch.nn.Module):
             # Blockwise path
             if self.use_aiter_fusion and self.use_aiter_quant_op:
                 return [rocm_aiter_ops.get_group_quant_op()]
-            if self.use_aiter_fusion:
-                return [torch.ops.vllm.triton_per_token_group_quant_fp8.default]
         else:
             if self.use_aiter_quant_op:
                 return [rocm_aiter_ops.get_per_token_quant_op()]
diff --git a/tests/compile/passes/test_silu_mul_quant_fusion.py b/tests/compile/passes/test_silu_mul_quant_fusion.py
index bc134ed427a..7d291cc5044 100644
--- a/tests/compile/passes/test_silu_mul_quant_fusion.py
+++ b/tests/compile/passes/test_silu_mul_quant_fusion.py
@@ -158,13 +158,6 @@ class TestSiluMulGroupFp8QuantModel(torch.nn.Module):
             input_dtype=dtype,
         )
 
-        if not current_platform.is_fp8_fnuz():
-            kernel = self.w8a8_block_fp8_linear.kernel
-            orig_quant = kernel.quant_fp8
-            kernel.quant_fp8 = lambda *a, use_triton=False, **kw: orig_quant(
-                *a, use_triton=True, **kw
-            )
-
         self.enable_silu_mul_custom_op = self.silu_and_mul.enabled()
 
     def forward(self, x):
@@ -175,9 +168,7 @@ class TestSiluMulGroupFp8QuantModel(torch.nn.Module):
     def ops_in_model_before(self):
         return [
             SILU_MUL_OP if self.enable_silu_mul_custom_op else torch.ops.aten.mul,
-            rocm_aiter_ops.get_group_quant_op()
-            if current_platform.is_fp8_fnuz()
-            else torch.ops.vllm.triton_per_token_group_quant_fp8.default,
+            rocm_aiter_ops.get_group_quant_op(),
         ]
 
     def ops_in_model_after(self):
diff --git a/tests/compile/test_dynamic_shapes_compilation.py b/tests/compile/test_dynamic_shapes_compilation.py
index 7f725c14f21..3260d5aecc5 100644
--- a/tests/compile/test_dynamic_shapes_compilation.py
+++ b/tests/compile/test_dynamic_shapes_compilation.py
@@ -1,7 +1,6 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import gc
 import tempfile
 from contextlib import contextmanager
 
@@ -9,8 +8,7 @@ import pytest
 import torch
 
 from tests.models.utils import check_logprobs_close
-from tests.utils import wait_for_rocm_memory_to_settle
-from vllm import LLM, SamplingParams
+from vllm import SamplingParams
 from vllm.compilation.decorators import support_torch_compile
 from vllm.config import CompilationConfig, VllmConfig, set_current_vllm_config
 from vllm.config.compilation import (
@@ -49,6 +47,7 @@ def get_test_models():
 @pytest.mark.skipif(not is_torch_equal_or_newer("2.10.0"), reason="requires torch 2.10")
 def test_dynamic_shapes_compilation(
     monkeypatch,
+    vllm_runner,
     model_name,
     shapes_type,
     use_aot_compile,
@@ -79,9 +78,13 @@ def test_dynamic_shapes_compilation(
 
     print(f"Testing {shapes_type.name} dynamic shapes...")
 
-    # Initialize the model with specific dynamic shapes configuration
-    model = LLM(
-        model=model_name,
+    sampling_params = SamplingParams(max_tokens=5, temperature=0, logprobs=10)
+    test_prompts = [prompt, "The capital of France is"]
+
+    # VllmRunner shuts down the engine core on exit, so the eager model
+    # below never races a lingering compiled engine for GPU memory.
+    with vllm_runner(
+        model_name,
         compilation_config={
             "mode": CompilationMode.VLLM_COMPILE,
             "dynamic_shapes_config": {
@@ -90,33 +93,25 @@ def test_dynamic_shapes_compilation(
             },
         },
         max_model_len=1024,
-    )
+        enable_chunked_prefill=None,
+    ) as vllm_model:
+        compiled_outputs = []
+        for p in test_prompts:
+            output = vllm_model.llm.generate(p, sampling_params)[0].outputs[0]
+            assert len(output.text.strip()) > 0, "Compiled model produced empty output"
+            compiled_outputs.append((output.token_ids, output.text, output.logprobs))
 
-    sampling_params = SamplingParams(max_tokens=5, temperature=0, logprobs=10)
-    test_prompts = [prompt, "The capital of France is"]
-
-    compiled_outputs = []
-    for p in test_prompts:
-        output = model.generate(p, sampling_params)[0].outputs[0]
-        assert len(output.text.strip()) > 0, "Compiled model produced empty output"
-        compiled_outputs.append((output.token_ids, output.text, output.logprobs))
-
-    del model
-    gc.collect()
-    torch.accelerator.empty_cache()
-    torch.accelerator.synchronize()
-    wait_for_rocm_memory_to_settle()
-
-    eager_model = LLM(model=model_name, enforce_eager=True, max_model_len=1024)
-    eager_outputs = []
-    for p in test_prompts:
-        output = eager_model.generate(p, sampling_params)[0].outputs[0]
-        assert len(output.text.strip()) > 0, "Eager model produced empty output"
-        eager_outputs.append((output.token_ids, output.text, output.logprobs))
-    del eager_model
-    gc.collect()
-    torch.accelerator.empty_cache()
-    torch.accelerator.synchronize()
+    with vllm_runner(
+        model_name,
+        enforce_eager=True,
+        max_model_len=1024,
+        enable_chunked_prefill=None,
+    ) as vllm_model:
+        eager_outputs = []
+        for p in test_prompts:
+            output = vllm_model.llm.generate(p, sampling_params)[0].outputs[0]
+            assert len(output.text.strip()) > 0, "Eager model produced empty output"
+            eager_outputs.append((output.token_ids, output.text, output.logprobs))
 
     check_logprobs_close(
         outputs_0_lst=eager_outputs,
@@ -241,44 +236,39 @@ def test_model_specialization_with_evaluate_guards(
 
 
 @pytest.mark.skipif(not is_torch_equal_or_newer("2.10.0"), reason="requires torch 2.10")
-def test_piecewise_backend_empty_sym_shape_indices():
+def test_piecewise_backend_empty_sym_shape_indices(vllm_runner):
     """Test that PiecewiseBackend handles empty sym_shape_indices correctly.
 
     When all inputs have static shapes (no torch.SymInt), sym_shape_indices
     will be empty. The fix in PiecewiseBackend.__call__ handles this case
     by using the first compiled range_entry.
     """
-    gc.collect()
-    torch.accelerator.empty_cache()
-    torch.accelerator.synchronize()
-
     # Use small max_model_len and max_num_batched_tokens to encourage
     # static shape compilation with empty sym_shape_indices
-    llm = LLM(
-        model="Qwen/Qwen3-0.6B",
+    with vllm_runner(
+        "Qwen/Qwen3-0.6B",
         max_model_len=512,
         max_num_batched_tokens=1,
+        enable_chunked_prefill=None,
         compilation_config={
             "mode": CompilationMode.VLLM_COMPILE,
             "dynamic_shapes_config": {
                 "type": DynamicShapesType.BACKED.value,
             },
         },
-    )
+    ) as vllm_model:
+        sampling_params = SamplingParams(temperature=0, top_p=0.95, max_tokens=10)
 
-    sampling_params = SamplingParams(temperature=0, top_p=0.95, max_tokens=10)
+        # Generate with static shape inputs
+        output = vllm_model.llm.generate(
+            "Hello, my name is", sampling_params=sampling_params
+        )
+        result = output[0].outputs[0].text
+        assert len(result) > 0, "Should generate non-empty output"
 
-    # Generate with static shape inputs
-    output = llm.generate("Hello, my name is", sampling_params=sampling_params)
-    result = output[0].outputs[0].text
-    assert len(result) > 0, "Should generate non-empty output"
-
-    # Generate again to verify compilation works with empty sym_shape_indices
-    output = llm.generate("The capital of France is", sampling_params=sampling_params)
-    result = output[0].outputs[0].text
-    assert len(result) > 0, "Should generate non-empty output on second run"
-
-    del llm
-    gc.collect()
-    torch.accelerator.empty_cache()
-    torch.accelerator.synchronize()
+        # Generate again to verify compilation works with empty sym_shape_indices
+        output = vllm_model.llm.generate(
+            "The capital of France is", sampling_params=sampling_params
+        )
+        result = output[0].outputs[0].text
+        assert len(result) > 0, "Should generate non-empty output on second run"
diff --git a/tests/distributed/test_elastic_ep.py b/tests/distributed/test_elastic_ep.py
index 4ce7497598a..01c254c2d03 100644
--- a/tests/distributed/test_elastic_ep.py
+++ b/tests/distributed/test_elastic_ep.py
@@ -3,7 +3,9 @@
 
 import os
 import subprocess
+import threading
 import time
+from concurrent.futures import ThreadPoolExecutor
 
 import pytest
 import requests
@@ -40,6 +42,106 @@ def _send_scale_command(server: RemoteOpenAIServer, new_dp_size: int) -> bool:
         return False
 
 
+def _traffic_loop(
+    server: RemoteOpenAIServer,
+    dp_rank: int | None,
+    ready: threading.Barrier,
+    stop: threading.Event,
+    finished: threading.Event,
+    is_probe: bool = False,
+) -> list[tuple[float, float, int | None]]:
+    url = server.url_for("is_scaling_elastic_ep" if is_probe else "v1/completions")
+    payload = {"model": MODEL_NAME, "prompt": "Hello", "max_tokens": 4}
+    headers = None if dp_rank is None else {"X-data-parallel-rank": str(dp_rank)}
+    request_payload = None if is_probe else payload
+    responses = []
+    is_ready = False
+    while not stop.is_set():
+        request_start = time.perf_counter()
+        try:
+            response = requests.post(
+                url, json=request_payload, headers=headers, timeout=120
+            )
+            status_code = response.status_code
+        except requests.exceptions.RequestException:
+            status_code = None
+        responses.append((request_start, time.perf_counter(), status_code))
+        if status_code == 200:
+            if not is_ready:
+                ready.wait(timeout=120)
+                is_ready = True
+            if finished.is_set():
+                return responses
+        time.sleep(0.05)
+    return responses
+
+
+def _downtime(responses: list[tuple[float, float, int | None]]) -> float:
+    rejected = [end for _, end, status in responses if status == 503]
+    if not rejected:
+        return 0
+    recovered = next(
+        end for _, end, status in responses if status == 200 and end > rejected[-1]
+    )
+    return recovered - rejected[0]
+
+
+def _scale_with_traffic(
+    server: RemoteOpenAIServer,
+    source_dp_size: int,
+    new_dp_size: int,
+    traffic_mode: str,
+) -> None:
+    traffic_clients: list[int | None] = []
+    if traffic_mode == "light":
+        traffic_clients = [0]
+    elif traffic_mode == "heavy":
+        traffic_clients = [None] * source_dp_size
+    clients = [(None, True)] + [(rank, False) for rank in traffic_clients]
+    ready = threading.Barrier(len(clients) + 1)
+    stop = threading.Event()
+    finished = threading.Event()
+
+    with ThreadPoolExecutor(max_workers=len(clients)) as executor:
+        futures = [
+            executor.submit(
+                _traffic_loop, server, rank, ready, stop, finished, is_probe
+            )
+            for rank, is_probe in clients
+        ]
+        try:
+            ready.wait(timeout=120)
+            start_time = time.perf_counter()
+            assert _send_scale_command(server, new_dp_size)
+            scale_seconds = time.perf_counter() - start_time
+            finished.set()
+            probe_result, *results = [future.result(timeout=120) for future in futures]
+        finally:
+            stop.set()
+
+    bad_statuses = {
+        status
+        for responses in [probe_result, *results]
+        for _, _, status in responses
+        if status not in (200, 503)
+    }
+    assert not bad_statuses, f"traffic got unexpected statuses {bad_statuses}"
+    probe_503 = [start for start, _, status in probe_result if status == 503]
+    assert probe_503, "Scaling probe did not observe commit"
+    assert not results or any(
+        status == 200 and start_time <= request_start and request_end < probe_503[0]
+        for responses in results
+        for request_start, request_end, status in responses
+    ), "No request completed successfully during preparation"
+
+    print(
+        f"[Elastic EP timing][{source_dp_size}->{new_dp_size}]"
+        f"[traffic={traffic_mode}] "
+        f"scale_seconds={scale_seconds:.3f} "
+        f"downtime_seconds={_downtime(probe_result):.3f}"
+    )
+
+
 def _run_gsm8k_eval(server: RemoteOpenAIServer, stage: str) -> float:
     assert server.port is not None
     result = evaluate_gsm8k(
@@ -59,7 +161,7 @@ def _run_gsm8k_eval(server: RemoteOpenAIServer, stage: str) -> float:
     return accuracy
 
 
-def _base_serve_args(use_async_eplb: bool = False) -> list[str]:
+def _base_serve_args(dp_size: int = 2, enforce_eager: bool = False) -> list[str]:
     args = [
         "--trust-remote-code",
         "--tensor-parallel-size",
@@ -78,57 +180,65 @@ def _base_serve_args(use_async_eplb: bool = False) -> list[str]:
         "--eplb-config.num_redundant_experts",
         "0",
         "--eplb-config.use_async",
-        "true" if use_async_eplb else "false",
+        "true",
         "--eplb-config.step_interval",
-        "10",
+        "300",
         "--eplb-config.window_size",
         "5",
         "--data-parallel-backend",
         "ray",
         "--data-parallel-size",
-        "2",
+        str(dp_size),
         "--api-server-count",
         "1",
+        "--disable-access-log-for-endpoints",
+        "/is_scaling_elastic_ep",
     ]
 
     leader_address = os.environ.get("LEADER_ADDRESS")
     if leader_address:
         args.extend(["--data-parallel-address", leader_address])
+    if enforce_eager:
+        args.append("--enforce-eager")
 
     return args
 
 
 @pytest.mark.parametrize(
-    "use_async_eplb", [False, True], ids=["sync_eplb", "async_eplb"]
+    ("enforce_eager", "traffic_mode"),
+    [
+        pytest.param(True, "none", id="enforce_eager_none"),
+        pytest.param(True, "light", id="enforce_eager_light"),
+        pytest.param(True, "heavy", id="enforce_eager_heavy"),
+        pytest.param(False, "heavy", id="cuda_graphs_heavy"),
+    ],
 )
 @multi_gpu_test(num_gpus=4)
-def test_elastic_ep_scaling(use_async_eplb: bool):
-    if use_async_eplb:
-        from vllm.distributed.eplb.eplb_communicator import has_nixl
+def test_elastic_ep_scaling(enforce_eager: bool, traffic_mode: str):
+    from vllm.distributed.eplb.eplb_communicator import has_nixl
 
-        if not has_nixl():
-            pytest.skip("Async EPLB with elastic EP requires NIXL (not installed)")
+    if not has_nixl():
+        pytest.skip("Async EPLB with elastic EP requires NIXL (not installed)")
 
-    vllm_serve_args = _base_serve_args(use_async_eplb)
+    initial_dp_size = int(os.getenv("VLLM_TEST_ELASTIC_EP_INITIAL_DP", "2"))
+    target_dp_size = int(os.getenv("VLLM_TEST_ELASTIC_EP_TARGET_DP", "4"))
+    assert target_dp_size > initial_dp_size
+    vllm_serve_args = _base_serve_args(initial_dp_size, enforce_eager)
 
     with RemoteOpenAIServer(
         MODEL_NAME, vllm_serve_args, env_dict={}, max_wait_seconds=1200
     ) as server:
-        initial_accuracy = _run_gsm8k_eval(server, "Initial (2 GPUs)")
-
-        assert _send_scale_command(server, 4)
-        time.sleep(10)
-        scale_up_accuracy = _run_gsm8k_eval(server, "After scale up (4 GPUs)")
+        initial_accuracy = _run_gsm8k_eval(server, "Initial")
 
+        _scale_with_traffic(server, initial_dp_size, target_dp_size, traffic_mode)
+        scale_up_accuracy = _run_gsm8k_eval(server, "After scale up")
         assert scale_up_accuracy >= initial_accuracy - ACCURACY_TOL, (
             f"Scale up accuracy {scale_up_accuracy:.3f} dropped more than "
             f"{ACCURACY_TOL} below initial accuracy {initial_accuracy:.3f}"
         )
 
-        assert _send_scale_command(server, 2)
-        time.sleep(5)
-        scale_down_accuracy = _run_gsm8k_eval(server, "After scale down (2 GPUs)")
-
+        _scale_with_traffic(server, target_dp_size, initial_dp_size, traffic_mode)
+        scale_down_accuracy = _run_gsm8k_eval(server, "After scale down")
         assert scale_down_accuracy >= initial_accuracy - ACCURACY_TOL, (
             f"Scale down accuracy {scale_down_accuracy:.3f} dropped more than "
             f"{ACCURACY_TOL} below initial accuracy {initial_accuracy:.3f}"
@@ -147,24 +257,20 @@ def test_elastic_ep_scaling(use_async_eplb: bool):
         print(f"  Tolerance:  {ACCURACY_TOL:.3f}")
 
 
-@pytest.mark.parametrize(
-    "use_async_eplb", [False, True], ids=["sync_eplb", "async_eplb"]
-)
 @multi_gpu_test(num_gpus=4)
-def test_elastic_ep_scaling_uneven(use_async_eplb: bool):
+def test_elastic_ep_scaling_uneven():
     """Test scale up with uneven worker distribution.
 
     This tests the case where num_new_workers % old_dp_size != 0,
     specifically 2 -> 3 where remainder = 1 % 2 = 1.
     This exercises the remainder handling in sender-receiver pairing.
     """
-    if use_async_eplb:
-        from vllm.distributed.eplb.eplb_communicator import has_nixl
+    from vllm.distributed.eplb.eplb_communicator import has_nixl
 
-        if not has_nixl():
-            pytest.skip("Async EPLB with elastic EP requires NIXL (not installed)")
+    if not has_nixl():
+        pytest.skip("Async EPLB with elastic EP requires NIXL (not installed)")
 
-    vllm_serve_args = _base_serve_args(use_async_eplb)
+    vllm_serve_args = _base_serve_args()
 
     with RemoteOpenAIServer(
         MODEL_NAME, vllm_serve_args, env_dict={}, max_wait_seconds=1200
@@ -174,7 +280,6 @@ def test_elastic_ep_scaling_uneven(use_async_eplb: bool):
         # Scale 2 -> 3: This has remainder = 1 % 2 = 1
         # Tests uneven sender-receiver pairing
         assert _send_scale_command(server, 3)
-        time.sleep(10)
         scale_up_accuracy = _run_gsm8k_eval(server, "After scale up (3 GPUs)")
 
         assert scale_up_accuracy >= initial_accuracy - ACCURACY_TOL, (
@@ -184,7 +289,6 @@ def test_elastic_ep_scaling_uneven(use_async_eplb: bool):
 
         # Scale back down to 2
         assert _send_scale_command(server, 2)
-        time.sleep(5)
         scale_down_accuracy = _run_gsm8k_eval(server, "After scale down (2 GPUs)")
 
         assert scale_down_accuracy >= initial_accuracy - ACCURACY_TOL, (
diff --git a/tests/distributed/test_shm_broadcast.py b/tests/distributed/test_shm_broadcast.py
index 17957924051..0b413965032 100644
--- a/tests/distributed/test_shm_broadcast.py
+++ b/tests/distributed/test_shm_broadcast.py
@@ -1,6 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+import io
+import pickle
 import random
 import threading
 import time
@@ -10,12 +12,15 @@ from unittest import mock
 import multiprocess as mp
 import numpy as np
 import pytest
+import torch
 import torch.distributed as dist
 
 from vllm.distributed.device_communicators import shm_broadcast
 from vllm.distributed.device_communicators.shm_broadcast import (
     MessageQueue,
     ShmRingBuffer,
+    _rebuild_tensor,
+    _reduce_tensor,
     check_shm_free_space,
 )
 from vllm.distributed.utils import StatelessProcessGroup
@@ -354,6 +359,150 @@ def test_message_queue_busy_to_idle():
     distributed_run(worker_fn_test_busy_to_idle, 4)
 
 
+@worker_fn_wrapper
+def worker_fn_tensor_broadcast():
+    rank = dist.get_rank()
+    writer_rank = 0
+    message_queue = MessageQueue.create_from_process_group(
+        dist.group.WORLD, 8 * 1024 * 1024, 4, writer_rank
+    )
+
+    # Both ranks construct the identical reference payload.
+    torch.manual_seed(42)
+    payload = {
+        # 2MiB: rides the shm ring as an out-of-band buffer (the receiving
+        # side must copy out of the reusable ring chunk).
+        "mid": torch.randn(1024, 512),
+        # 16MiB > max_chunk_bytes: overflows to the zmq socket (the
+        # receiving side aliases the zmq.Frame zero-copy).
+        "big": torch.randn(4096, 2048, dtype=torch.bfloat16),
+        "nested": ["plain", 123, {"inner": torch.arange(5)}],
+    }
+
+    if rank == writer_rank:
+        with mock.patch(
+            "vllm.distributed.device_communicators.shm_broadcast._reduce_tensor",
+            wraps=_reduce_tensor,
+        ) as wrapped_reduce:
+            message_queue.enqueue(payload)
+        assert wrapped_reduce.call_count == 3
+        # Cycle the ring (max_chunks=4) several times over so that aliased
+        # ring chunks would be overwritten.
+        for i in range(16):
+            message_queue.enqueue({"junk": torch.full((1024, 512), float(i))})
+    else:
+        received = message_queue.dequeue(timeout=30)
+        for key in ("mid", "big"):
+            assert torch.equal(received[key], payload[key]), key
+            assert received[key].dtype == payload[key].dtype, key
+        assert torch.equal(received["nested"][2]["inner"], torch.arange(5))
+
+        snapshot = received["mid"].clone()
+        for i in range(16):
+            junk = message_queue.dequeue(timeout=30)
+            assert torch.equal(junk["junk"], torch.full((1024, 512), float(i)))
+        # Tensors received via the shm ring must not alias chunk memory
+        # that the writer has reused for subsequent messages.
+        assert torch.equal(received["mid"], snapshot)
+        # Rebuilt tensors must be writable, like regular tensors.
+        received["mid"] += 1.0
+        received["big"][0, 0] = 1.0
+
+    dist.barrier()
+    print(f"tensor broadcast passed the test! Rank {rank}")
+
+
+def test_tensor_broadcast():
+    distributed_run(worker_fn_tensor_broadcast, 2)
+
+
+def _dumps_oob(obj) -> tuple[bytes, list]:
+    """Pickle `obj` the same way `MessageQueue.enqueue` does: tensor
+    dispatch table + out-of-band buffers >= 1MiB."""
+    buffers = []
+
+    def callback(buf: pickle.PickleBuffer) -> bool:
+        raw = buf.raw()
+        if raw.nbytes < 1024 * 1024:
+            return True
+        buffers.append(raw)
+        return False
+
+    bio = io.BytesIO()
+    pickler = pickle.Pickler(
+        bio, protocol=pickle.HIGHEST_PROTOCOL, buffer_callback=callback
+    )
+    pickler.dispatch_table = {torch.Tensor: _reduce_tensor}
+    pickler.dump(obj)
+    return bio.getvalue(), buffers
+
+
+@pytest.mark.parametrize(
+    "case",
+    [
+        "small",
+        "mid",
+        "bf16",
+        "fp8",
+        "empty",
+        "scalar",
+        "noncontig",
+        "requires_grad",
+        "conj",
+        "param",
+    ],
+)
+def test_tensor_pickle_roundtrip(case: str):
+    tensor = {
+        # Inlined in-band (< 1MiB) and out-of-band (>= 1MiB) buffers.
+        "small": lambda: torch.randn(100, 10),
+        "mid": lambda: torch.randn(1024, 512),
+        # Dtypes numpy doesn't recognize.
+        "bf16": lambda: torch.randn(512, 512, dtype=torch.bfloat16),
+        "fp8": lambda: torch.randn(32, 32).to(torch.float8_e4m3fn),
+        # Shape edge cases.
+        "empty": lambda: torch.empty(0, 8),
+        "scalar": lambda: torch.tensor(3.14),
+        "noncontig": lambda: torch.randn(64, 64).t(),
+        # These fall back to torch's default reducer.
+        "requires_grad": lambda: torch.randn(8, 8, requires_grad=True),
+        "conj": lambda: torch.randn(4, dtype=torch.complex64).conj(),
+        "param": lambda: torch.nn.Parameter(torch.randn(4), requires_grad=False),
+    }[case]()
+
+    data, buffers = _dumps_oob({"tensor": tensor, "meta": list(range(10))})
+    received = pickle.loads(data, buffers=buffers)["tensor"]
+
+    assert received.shape == tensor.shape
+    assert received.dtype == tensor.dtype
+    if tensor.dtype == torch.float8_e4m3fn:
+        assert torch.equal(received.view(torch.uint8), tensor.view(torch.uint8))
+    else:
+        assert torch.equal(received, tensor)
+    assert received.requires_grad == tensor.requires_grad
+    assert isinstance(received, type(tensor))
+    if tensor.numel() and not tensor.requires_grad:
+        # Rebuilt tensors must be writable, like regular tensors.
+        received.view(-1)[0] = 1.0
+
+
+@pytest.mark.parametrize("case", ["cuda", "requires_grad", "conj"])
+def test_reduce_tensor_fallback(case: str):
+    """Tensors the zero-copy reducer can't safely alias must fall back to
+    torch's default reduction."""
+    if case == "cuda":
+        if not torch.cuda.is_available():
+            pytest.skip("requires CUDA")
+        tensor = torch.randn(4, device="cuda")
+    elif case == "requires_grad":
+        tensor = torch.randn(8, requires_grad=True)
+    else:
+        tensor = torch.randn(4, dtype=torch.complex64).conj()
+
+    reduced = _reduce_tensor(tensor)
+    assert reduced[0] is not _rebuild_tensor
+
+
 @pytest.mark.parametrize("should_warn", [False, True])
 def test_reader_timeout_caps_indefinite_waits(should_warn):
     with (
diff --git a/tests/distributed/test_weight_transfer.py b/tests/distributed/test_weight_transfer.py
index b79aa1974d1..eeeceb95998 100644
--- a/tests/distributed/test_weight_transfer.py
+++ b/tests/distributed/test_weight_transfer.py
@@ -1247,7 +1247,7 @@ class RecordingClient:
         self.order.append("update")
         self.last_update_info = update_info
 
-    def finish_weight_update(self) -> None:
+    def finish_weight_update(self, weight_version: str | None = None) -> None:
         self.order.append("finish")
 
 
@@ -1303,6 +1303,10 @@ class TestTrainerClients:
         assert isinstance(update_req, WeightTransferUpdateRequest)
         assert update_req.update_info == {"names": ["w"]}
 
+        client.finish_weight_update("step-42")
+        handle.finish_weight_update.remote.assert_called_once_with()
+        handle.update_weight_version.remote.assert_called_once_with("step-42")
+
     def test_http_client_pickles_ipc_handles_for_json(self, monkeypatch):
         """HTTP update_weights must encode raw ipc_handles as a base64 pickle."""
         captured = {}
@@ -1334,6 +1338,9 @@ class TestTrainerClients:
         client.update_weights(update_info)
         assert captured["json"]["update_info"] == update_info
 
+        client.finish_weight_update("step-42")
+        assert captured["json"] == {"weight_version": "step-42"}
+
 
 class TestModuleSource:
     """`ModuleSource` metadata vs. materialized iteration (dense, no GPU)."""
diff --git a/tests/engine/test_short_mm_context.py b/tests/engine/test_short_mm_context.py
index 23489c21333..940709c8e53 100644
--- a/tests/engine/test_short_mm_context.py
+++ b/tests/engine/test_short_mm_context.py
@@ -3,6 +3,8 @@
 
 import pytest
 
+from vllm.exceptions import VLLMValidationError
+
 from ..conftest import IMAGE_ASSETS
 
 HF_IMAGE_PROMPTS = IMAGE_ASSETS.prompts(
@@ -19,7 +21,9 @@ models = ["llava-hf/llava-1.5-7b-hf"]
 def test_context_length_too_short(vllm_runner, image_assets, model):
     images = [asset.pil_image for asset in image_assets]
 
-    with pytest.raises(ValueError, match="longer than the maximum model length"):
+    with pytest.raises(
+        VLLMValidationError, match="longer than the maximum model length"
+    ):
         vllm_model = vllm_runner(
             model,
             # LLaVA has a feature size of 576
diff --git a/tests/entrypoints/llm/test_chat.py b/tests/entrypoints/llm/test_chat.py
index 61cdbd3eee2..cbc57b80da4 100644
--- a/tests/entrypoints/llm/test_chat.py
+++ b/tests/entrypoints/llm/test_chat.py
@@ -6,6 +6,7 @@ import pytest
 
 from vllm import LLM
 from vllm.distributed import cleanup_dist_env_and_memory
+from vllm.exceptions import VLLMValidationError
 from vllm.sampling_params import SamplingParams
 
 
@@ -157,7 +158,7 @@ def test_chat_batch_failure_cleanup(llm_for_failure_test):
     batch_2 = [valid_msg, valid_msg]
     sampling_params = SamplingParams(temperature=0, max_tokens=10)
 
-    with pytest.raises(ValueError, match="maximum context length is"):
+    with pytest.raises(VLLMValidationError, match="maximum context length is"):
         llm.chat(batch_1, sampling_params=sampling_params)
     assert llm.llm_engine.get_num_unfinished_requests() == 0
 
diff --git a/tests/entrypoints/llm/test_prompt_validation.py b/tests/entrypoints/llm/test_prompt_validation.py
index c17486d962f..8dd55c6b10e 100644
--- a/tests/entrypoints/llm/test_prompt_validation.py
+++ b/tests/entrypoints/llm/test_prompt_validation.py
@@ -5,17 +5,18 @@ import pytest
 import torch
 
 from vllm import LLM
+from vllm.exceptions import VLLMValidationError
 
 
 def test_empty_prompt():
     llm = LLM(model="openai-community/gpt2", enforce_eager=True)
-    with pytest.raises(ValueError, match="decoder prompt cannot be empty"):
+    with pytest.raises(VLLMValidationError, match="decoder prompt cannot be empty"):
         llm.generate([""])
 
 
 def test_out_of_vocab_token():
     llm = LLM(model="openai-community/gpt2", enforce_eager=True)
-    with pytest.raises(ValueError, match="out of vocabulary"):
+    with pytest.raises(VLLMValidationError, match="out of vocabulary"):
         llm.generate({"prompt_token_ids": [999999]})
 
 
diff --git a/tests/entrypoints/multimodal/llm/test_mm_embeds_only.py b/tests/entrypoints/multimodal/llm/test_mm_embeds_only.py
index 57bec9c1188..0ea18071246 100644
--- a/tests/entrypoints/multimodal/llm/test_mm_embeds_only.py
+++ b/tests/entrypoints/multimodal/llm/test_mm_embeds_only.py
@@ -6,6 +6,7 @@ import pytest
 from tests.entrypoints.multimodal.conftest import managed_llm
 from vllm import LLM, SamplingParams
 from vllm.assets.image import ImageAsset
+from vllm.exceptions import VLLMValidationError
 
 MODEL = "llava-hf/llava-1.5-7b-hf"
 PROMPT = "USER: <image>\nDescribe this image briefly.\nASSISTANT:"
@@ -42,7 +43,7 @@ def test_generate_with_embedding(llm: LLM):
 def test_raw_image_rejected(llm: LLM):
     """Raw image input is still rejected when limit=0."""
     raw_image = ImageAsset("stop_sign").pil_image
-    with pytest.raises(ValueError, match=r"At most 0 image\(s\)"):
+    with pytest.raises(VLLMValidationError, match=r"At most 0 image\(s\)"):
         llm.generate(
             {"prompt": PROMPT, "multi_modal_data": {"image": raw_image}},
             sampling_params=SamplingParams(max_tokens=16),
diff --git a/tests/entrypoints/multimodal/openai/chat_completion/test_audio.py b/tests/entrypoints/multimodal/openai/chat_completion/test_audio.py
index fa0f141afee..a1c13d9e339 100644
--- a/tests/entrypoints/multimodal/openai/chat_completion/test_audio.py
+++ b/tests/entrypoints/multimodal/openai/chat_completion/test_audio.py
@@ -10,6 +10,7 @@ import pytest_asyncio
 from tests.utils import RemoteOpenAIServer
 from vllm.assets.audio import AudioAsset
 from vllm.multimodal.utils import encode_audio_base64, encode_audio_url, fetch_audio
+from vllm.platforms import current_platform
 
 MODEL_NAME = "fixie-ai/ultravox-v0_5-llama-3_2-1b"
 TEST_AUDIO_URLS = [
@@ -18,6 +19,10 @@ TEST_AUDIO_URLS = [
 ]
 MAXIMUM_AUDIOS = 2
 
+# Disable prefix caching on ROCm to reduce non-determinism in
+# streaming-vs-non-streaming comparisons.
+_ROCM_ARGS = ["--no-enable-prefix-caching"] if current_platform.is_rocm() else []
+
 
 @pytest.fixture(scope="module")
 def server():
@@ -32,6 +37,7 @@ def server():
         "--trust-remote-code",
         "--limit-mm-per-prompt",
         json.dumps({"audio": MAXIMUM_AUDIOS}),
+        *_ROCM_ARGS,
     ]
 
     with RemoteOpenAIServer(MODEL_NAME, args) as remote_server:
diff --git a/tests/entrypoints/openai/chat_completion/test_chat.py b/tests/entrypoints/openai/chat_completion/test_chat.py
index 32c72f1ef93..5541b605d99 100644
--- a/tests/entrypoints/openai/chat_completion/test_chat.py
+++ b/tests/entrypoints/openai/chat_completion/test_chat.py
@@ -18,6 +18,7 @@ from tests.utils import RemoteOpenAIServer
 from vllm.entrypoints.openai.chat_completion.protocol import (
     ChatCompletionRequest,
 )
+from vllm.exceptions import VLLMValidationError
 from vllm.sampling_params import SamplingParams
 
 # any model with a chat template should work here
@@ -1074,7 +1075,7 @@ def test_chat_completion_request_n_parameter_exceeds_default_limit(
         max_tokens=10,
     )
 
-    with pytest.raises(ValueError, match="n must be at most"):
+    with pytest.raises(VLLMValidationError, match="n must be at most"):
         request.to_sampling_params(
             max_tokens=10,
             default_sampling_params={},
@@ -1136,7 +1137,7 @@ def test_chat_completion_request_n_parameter_custom_limit(
         max_tokens=10,
     )
 
-    with pytest.raises(ValueError, match="n must be at most 128"):
+    with pytest.raises(VLLMValidationError, match="n must be at most 128"):
         request_over.to_sampling_params(
             max_tokens=10,
             default_sampling_params={},
@@ -1160,7 +1161,7 @@ def test_chat_completion_request_n_parameter_massive_value(
         max_tokens=1,
     )
 
-    with pytest.raises(ValueError, match="n must be at most"):
+    with pytest.raises(VLLMValidationError, match="n must be at most"):
         request.to_sampling_params(
             max_tokens=1,
             default_sampling_params={},
diff --git a/tests/entrypoints/openai/chat_completion/test_chat_error.py b/tests/entrypoints/openai/chat_completion/test_chat_error.py
index 4b42f522e81..9e805254a2f 100644
--- a/tests/entrypoints/openai/chat_completion/test_chat_error.py
+++ b/tests/entrypoints/openai/chat_completion/test_chat_error.py
@@ -6,7 +6,6 @@ from typing import Any
 from unittest.mock import AsyncMock, MagicMock, patch
 
 import pytest
-from pydantic import ValidationError
 
 from vllm.config.multimodal import MultiModalConfig
 from vllm.entrypoints.openai.chat_completion.protocol import (
@@ -18,6 +17,7 @@ from vllm.entrypoints.openai.engine.protocol import GenerationError
 from vllm.entrypoints.openai.models.protocol import BaseModelPath
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
 from vllm.entrypoints.scale_out.render.serving import ServingRender
+from vllm.exceptions import VLLMValidationError
 from vllm.outputs import CompletionOutput, RequestOutput
 from vllm.renderers.hf import HfRenderer
 from vllm.renderers.online_renderer import OnlineRenderer
@@ -479,7 +479,7 @@ def test_json_schema_response_format_missing_schema():
 def test_structural_tag_response_format_invalid(format_value):
     """Malformed structural tags should be rejected during request validation."""
     with pytest.raises(
-        ValidationError,
+        VLLMValidationError,
         match="Invalid response_format structural_tag",
     ):
         ChatCompletionRequest(
@@ -493,7 +493,7 @@ def test_structural_tag_response_format_invalid(format_value):
 def test_batch_structural_tag_response_format_invalid(format_value):
     """Batch chat should reject malformed structural tags at request parsing."""
     with pytest.raises(
-        ValidationError,
+        VLLMValidationError,
         match="Invalid response_format structural_tag",
     ):
         BatchChatCompletionRequest(
@@ -507,7 +507,7 @@ def test_batch_structural_tag_response_format_invalid(format_value):
 def test_structured_outputs_structural_tag_invalid(structural_tag):
     """Malformed direct structured_outputs structural tags should be rejected."""
     with pytest.raises(
-        ValidationError,
+        VLLMValidationError,
         match="Invalid structured_outputs structural_tag",
     ):
         ChatCompletionRequest(
@@ -521,7 +521,7 @@ def test_structured_outputs_structural_tag_invalid(structural_tag):
 def test_non_numeric_logprobs_rejected(field_name):
     """A non-numeric logprobs value must be a clean 400 validation error, not a
     TypeError from the mode='before' comparison (which surfaces as HTTP 500)."""
-    with pytest.raises(ValidationError, match=f"`{field_name}` must be an integer"):
+    with pytest.raises(VLLMValidationError, match=f"`{field_name}` must be an integer"):
         ChatCompletionRequest(
             model=MODEL_NAME,
             messages=[{"role": "user", "content": "hello"}],
diff --git a/tests/entrypoints/openai/chat_completion/test_logprob_token_ids.py b/tests/entrypoints/openai/chat_completion/test_logprob_token_ids.py
index aa04d787ccc..1eea6db48a7 100644
--- a/tests/entrypoints/openai/chat_completion/test_logprob_token_ids.py
+++ b/tests/entrypoints/openai/chat_completion/test_logprob_token_ids.py
@@ -14,11 +14,11 @@ digit-token vocab id).
 import math
 
 import pytest
-from pydantic import ValidationError
 
 from tests.utils import RemoteOpenAIServer
 from vllm.entrypoints.openai.chat_completion.protocol import ChatCompletionRequest
 from vllm.entrypoints.openai.completion.protocol import CompletionRequest
+from vllm.exceptions import VLLMValidationError
 
 MODEL_NAME = "Qwen/Qwen2.5-1.5B-Instruct"
 
@@ -87,7 +87,7 @@ def test_completion_request_decouples_top_k_from_explicit_token_ids():
 
 
 def test_completion_rejects_explicit_token_ids_without_generated_tokens():
-    with pytest.raises(ValidationError, match="no output tokens are generated"):
+    with pytest.raises(VLLMValidationError, match="no output tokens are generated"):
         CompletionRequest(
             model=MODEL_NAME,
             prompt="Hello",
@@ -99,7 +99,7 @@ def test_completion_rejects_explicit_token_ids_without_generated_tokens():
 
 
 def test_requests_reject_explicit_token_ids_with_beam_search():
-    with pytest.raises(ValidationError, match="not supported with beam search"):
+    with pytest.raises(VLLMValidationError, match="not supported with beam search"):
         ChatCompletionRequest(
             model=MODEL_NAME,
             messages=[{"role": "user", "content": "Hello"}],
@@ -108,7 +108,7 @@ def test_requests_reject_explicit_token_ids_with_beam_search():
             use_beam_search=True,
         )
 
-    with pytest.raises(ValidationError, match="not supported with beam search"):
+    with pytest.raises(VLLMValidationError, match="not supported with beam search"):
         CompletionRequest(
             model=MODEL_NAME,
             prompt="Hello",
diff --git a/tests/entrypoints/openai/chat_completion/test_thinking_token_budget_validation.py b/tests/entrypoints/openai/chat_completion/test_thinking_token_budget_validation.py
index e66205b7df2..1b2b76dc093 100644
--- a/tests/entrypoints/openai/chat_completion/test_thinking_token_budget_validation.py
+++ b/tests/entrypoints/openai/chat_completion/test_thinking_token_budget_validation.py
@@ -2,15 +2,15 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import pytest
-from pydantic import ValidationError
 
 from vllm.entrypoints.openai.chat_completion.protocol import ChatCompletionRequest
 from vllm.entrypoints.openai.completion.protocol import CompletionRequest
+from vllm.exceptions import VLLMValidationError
 
 
 @pytest.mark.parametrize("raw_value", [-2, 0.6, 10.5])
 def test_chat_completion_request_rejects_invalid_thinking_token_budget(raw_value):
-    with pytest.raises(ValidationError, match="thinking_token_budget"):
+    with pytest.raises(VLLMValidationError, match="thinking_token_budget"):
         ChatCompletionRequest.model_validate(
             {
                 "model": "qwen",
@@ -44,7 +44,7 @@ def test_chat_completion_request_accepts_minus_one_as_unlimited():
 
 @pytest.mark.parametrize("raw_value", [0.6, 3.14, -2])
 def test_completion_request_rejects_invalid_thinking_token_budget(raw_value):
-    with pytest.raises(ValidationError, match="thinking_token_budget"):
+    with pytest.raises(VLLMValidationError, match="thinking_token_budget"):
         CompletionRequest.model_validate(
             {
                 "model": "qwen",
diff --git a/tests/entrypoints/openai/completion/test_completion_error.py b/tests/entrypoints/openai/completion/test_completion_error.py
index 818cad738d9..2f3db9f697d 100644
--- a/tests/entrypoints/openai/completion/test_completion_error.py
+++ b/tests/entrypoints/openai/completion/test_completion_error.py
@@ -6,7 +6,6 @@ from typing import Any
 from unittest.mock import AsyncMock, MagicMock
 
 import pytest
-from pydantic import ValidationError
 
 from vllm.config.multimodal import MultiModalConfig
 from vllm.entrypoints.openai.completion.protocol import CompletionRequest
@@ -18,6 +17,7 @@ from vllm.entrypoints.openai.engine.protocol import (
 from vllm.entrypoints.openai.models.protocol import BaseModelPath
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
 from vllm.entrypoints.scale_out.render.serving import ServingRender
+from vllm.exceptions import VLLMValidationError
 from vllm.outputs import CompletionOutput, RequestOutput
 from vllm.renderers.hf import HfRenderer
 from vllm.renderers.online_renderer import OnlineRenderer
@@ -430,7 +430,7 @@ def test_json_schema_response_format_missing_schema():
 def test_structural_tag_response_format_invalid(format_value):
     """Malformed structural tags should be rejected during request validation."""
     with pytest.raises(
-        ValidationError,
+        VLLMValidationError,
         match="Invalid response_format structural_tag",
     ):
         CompletionRequest(
@@ -445,7 +445,7 @@ def test_structural_tag_response_format_invalid(format_value):
 def test_structured_outputs_structural_tag_invalid(structural_tag):
     """Malformed direct structured_outputs structural tags should be rejected."""
     with pytest.raises(
-        ValidationError,
+        VLLMValidationError,
         match="Invalid structured_outputs structural_tag",
     ):
         CompletionRequest(
@@ -616,7 +616,7 @@ class TestCompletionPromptListLimit:
 def test_non_numeric_logprobs_rejected(field_name):
     """A non-numeric logprobs value must be a clean 400 validation error, not a
     TypeError from the mode='before' comparison (which surfaces as HTTP 500)."""
-    with pytest.raises(ValidationError, match=f"`{field_name}` must be an integer"):
+    with pytest.raises(VLLMValidationError, match=f"`{field_name}` must be an integer"):
         CompletionRequest(
             model=MODEL_NAME,
             prompt="Test prompt",
diff --git a/tests/entrypoints/openai/completion/test_prompt_validation.py b/tests/entrypoints/openai/completion/test_prompt_validation.py
index 87c6b6e1668..6c40037e07c 100644
--- a/tests/entrypoints/openai/completion/test_prompt_validation.py
+++ b/tests/entrypoints/openai/completion/test_prompt_validation.py
@@ -13,6 +13,7 @@ import torch
 
 from tests.utils import RemoteOpenAIServer
 from vllm.config import ModelConfig
+from vllm.exceptions import VLLMValidationError
 from vllm.renderers.embed_utils import safe_load_prompt_embeds
 
 
@@ -111,5 +112,5 @@ def test_disable_prompt_embeds(dtype: torch.dtype, seq_len: int, hidden_size: in
     buffer.seek(0)
     encoded_tensor = pybase64.b64encode(buffer.getvalue())
 
-    with pytest.raises(ValueError, match="--enable-prompt-embeds"):
+    with pytest.raises(VLLMValidationError, match="--enable-prompt-embeds"):
         safe_load_prompt_embeds(model_config, encoded_tensor)
diff --git a/tests/entrypoints/openai/responses/test_sampling_params.py b/tests/entrypoints/openai/responses/test_sampling_params.py
index 5a68e3a9c0d..6ede3f1f7f2 100644
--- a/tests/entrypoints/openai/responses/test_sampling_params.py
+++ b/tests/entrypoints/openai/responses/test_sampling_params.py
@@ -14,6 +14,7 @@ from vllm.entrypoints.openai.responses.protocol import (
     ResponsesRequest,
     ResponseTextConfig,
 )
+from vllm.exceptions import VLLMValidationError
 from vllm.sampling_params import StructuredOutputsParams
 
 
@@ -163,7 +164,7 @@ class TestResponsesRequestSamplingParams:
             text=text_config,
         )
 
-        with pytest.raises(ValueError) as exc_info:
+        with pytest.raises(VLLMValidationError) as exc_info:
             request.to_sampling_params(default_max_tokens=1000)
 
         assert "Cannot specify both structured_outputs and text.format" in str(
diff --git a/tests/entrypoints/openai/test_dp_supervisor.py b/tests/entrypoints/openai/test_dp_supervisor.py
index 576e7ef16df..df80053deb5 100644
--- a/tests/entrypoints/openai/test_dp_supervisor.py
+++ b/tests/entrypoints/openai/test_dp_supervisor.py
@@ -201,6 +201,7 @@ def test_run_vllm_dp_server_uses_rust_frontend_when_enabled(monkeypatch):
     monkeypatch.setattr(dp_sup.os, "setpgrp", lambda: None)
     monkeypatch.setattr(dp_sup, "set_process_title", lambda *_args: None)
     monkeypatch.setattr(dp_sup, "decorate_logs", lambda *_args: None)
+    monkeypatch.setattr(dp_sup.envs, "VLLM_USE_RUST_FRONTEND", True, raising=False)
     monkeypatch.setattr(
         dp_sup.envs,
         "VLLM_RUST_FRONTEND_PATH",
diff --git a/tests/entrypoints/openai/test_openai_schema.py b/tests/entrypoints/openai/test_openai_schema.py
index 2985c539518..6d3fc2f4474 100644
--- a/tests/entrypoints/openai/test_openai_schema.py
+++ b/tests/entrypoints/openai/test_openai_schema.py
@@ -148,6 +148,7 @@ def test_openapi_stateless(case: schemathesis.Case):
         "/start_draft_weight_update",
         "/update_weights",
         "/finish_weight_update",
+        "/update_weight_version",
     ):
         return
 
diff --git a/tests/entrypoints/serve/instrumentator/test_http_status_metrics.py b/tests/entrypoints/serve/instrumentator/test_http_status_metrics.py
index 0f96bf161d2..060b05bbaca 100644
--- a/tests/entrypoints/serve/instrumentator/test_http_status_metrics.py
+++ b/tests/entrypoints/serve/instrumentator/test_http_status_metrics.py
@@ -8,54 +8,95 @@ PrometheusInstrumentatorMiddleware before being caught by ServerErrorMiddleware.
 """
 
 from argparse import Namespace
-from http import HTTPStatus
 
 import httpx
 import pytest
-from fastapi import FastAPI, HTTPException, Request
-from fastapi.exceptions import RequestValidationError
-from fastapi.responses import JSONResponse
+from fastapi import HTTPException
 from prometheus_client import CollectorRegistry
-from prometheus_fastapi_instrumentator import Instrumentator
 
-from vllm.entrypoints.serve.utils.server_utils import exception_handler
-from vllm.exceptions import VLLMNotFoundError, VLLMValidationError
+from vllm.entrypoints.openai.api_server import build_app
+from vllm.exceptions import (
+    VLLMNotFoundError,
+    VLLMServerError,
+    VLLMValidationError,
+)
 
 
-@pytest.fixture
+@pytest.fixture(scope="module")
+def should_do_global_cleanup_after_test() -> bool:
+    # This suite never initializes distributed/accelerator state.
+    return False
+
+
+def _build_args() -> Namespace:
+    """Minimal args for ``build_app``; avoids ``make_arg_parser`` device probing."""
+    return Namespace(
+        disable_fastapi_docs=True,
+        enable_offline_docs=False,
+        root_path=None,
+        allowed_origins=["*"],
+        allow_credentials=False,
+        allowed_methods=["*"],
+        allowed_headers=["*"],
+        api_key=None,
+        enable_request_id_headers=False,
+        enable_fault_tolerance=False,
+        middleware=[],
+        log_error_stack=False,
+    )
+
+
+@pytest.fixture(scope="module")
 def registry():
-    """Create a fresh Prometheus registry for each test."""
+    """Shared Prometheus registry for the module-scoped app."""
     return CollectorRegistry()
 
 
-@pytest.fixture
+@pytest.fixture(scope="module")
 def app(registry):
-    """Create a minimal FastAPI app that mirrors vLLM's exception handler
-    and Prometheus middleware setup."""
+    """Build the real vLLM FastAPI app once and attach probe routes that raise.
 
-    app = FastAPI()
+    Patch the name used by ``attach_router`` (imported into the instrumentator
+    metrics module), not ``vllm.v1.metrics.prometheus`` alone — that binding is
+    captured at import time.
+    """
+    import vllm.entrypoints.serve.instrumentator.metrics as metrics_mod
 
-    # Mock app state that exception_handler needs
-    app.state.args = Namespace(log_error_stack=False)
+    original = metrics_mod.get_prometheus_registry
+    metrics_mod.get_prometheus_registry = lambda: registry
+    try:
+        app = build_app(_build_args(), supported_tasks=())
+    finally:
+        metrics_mod.get_prometheus_registry = original
 
-    # Register exception handlers exactly as vLLM does in build_app()
-    app.exception_handler(HTTPException)(_http_exception_handler)
-    app.exception_handler(RequestValidationError)(_validation_exception_handler)
-    app.exception_handler(ValueError)(exception_handler)
-    app.exception_handler(TypeError)(exception_handler)
-    app.exception_handler(OverflowError)(exception_handler)
-    app.exception_handler(NotImplementedError)(exception_handler)
-    app.exception_handler(VLLMValidationError)(exception_handler)
-    app.exception_handler(VLLMNotFoundError)(exception_handler)
-    app.exception_handler(Exception)(exception_handler)
+    @app.get("/raise_http_exception_400")
+    async def raise_http_exception_400():
+        raise HTTPException(status_code=400, detail="bad request")
 
-    # Instrument with Prometheus (same as vLLM's attach_router)
-    Instrumentator(
-        excluded_handlers=["/metrics"],
-        registry=registry,
-    ).add().instrument(app)
+    @app.get("/raise_http_exception_404")
+    async def raise_http_exception_404():
+        raise HTTPException(status_code=404, detail="not found")
+
+    @app.get("/raise_request_validation_error")
+    async def raise_request_validation_error(n: int):
+        # Invalid ``n`` triggers FastAPI's RequestValidationError.
+        return {"n": n}
+
+    @app.get("/raise_vllm_validation_error")
+    async def raise_vllm_validation_error():
+        raise VLLMValidationError("bad parameter", parameter="temperature")
+
+    @app.get("/raise_vllm_not_found_error")
+    async def raise_vllm_not_found_error():
+        raise VLLMNotFoundError("model not found")
+
+    @app.get("/raise_vllm_server_error")
+    async def raise_vllm_server_error():
+        # Bare VLLMServerError goes through vllm_error_handler → 500.
+        # EngineGenerateError / EngineDeadError are not used here: they call
+        # terminate_if_errored and need engine/server state.
+        raise VLLMServerError("internal server failure")
 
-    # Test routes that raise different exception types
     @app.get("/raise_value_error")
     async def raise_value_error():
         raise ValueError("invalid input value")
@@ -72,22 +113,6 @@ def app(registry):
     async def raise_not_implemented_error():
         raise NotImplementedError("feature not supported")
 
-    @app.get("/raise_vllm_validation_error")
-    async def raise_vllm_validation_error():
-        raise VLLMValidationError("bad parameter", parameter="temperature")
-
-    @app.get("/raise_vllm_not_found_error")
-    async def raise_vllm_not_found_error():
-        raise VLLMNotFoundError("model not found")
-
-    @app.get("/raise_http_exception_400")
-    async def raise_http_exception_400():
-        raise HTTPException(status_code=400, detail="bad request")
-
-    @app.get("/raise_http_exception_404")
-    async def raise_http_exception_404():
-        raise HTTPException(status_code=404, detail="not found")
-
     @app.get("/raise_runtime_error")
     async def raise_runtime_error():
         raise RuntimeError("unexpected server error")
@@ -99,14 +124,6 @@ def app(registry):
     return app
 
 
-async def _http_exception_handler(req: Request, exc: HTTPException):
-    return JSONResponse({"error": exc.detail}, status_code=exc.status_code)
-
-
-async def _validation_exception_handler(req: Request, exc: RequestValidationError):
-    return JSONResponse({"error": str(exc)}, status_code=HTTPStatus.BAD_REQUEST)
-
-
 def _get_http_requests_total(registry, method: str, handler: str):
     """Extract the http_requests_total metric values grouped by status.
 
@@ -128,31 +145,31 @@ def _get_http_requests_total(registry, method: str, handler: str):
 
 @pytest.mark.asyncio
 @pytest.mark.parametrize(
-    "endpoint,expected_status_group,expected_http_code",
+    "endpoint,expected_status_group,expected_http_code,request_kwargs",
     [
-        # These should record as 4xx in Prometheus
-        ("/raise_value_error", "4xx", 400),
-        ("/raise_type_error", "4xx", 400),
-        ("/raise_overflow_error", "4xx", 400),
-        ("/raise_vllm_validation_error", "4xx", 400),
-        ("/raise_vllm_not_found_error", "4xx", 404),
-        ("/raise_http_exception_400", "4xx", 400),
-        ("/raise_http_exception_404", "4xx", 404),
-        # NotImplementedError returns 501 which is still 5xx group
-        ("/raise_not_implemented_error", "5xx", 501),
-        # These should record as 5xx in Prometheus (genuine server errors)
-        ("/raise_runtime_error", "5xx", 500),
-        # Successful requests should record as 2xx
-        ("/success", "2xx", 200),
+        ("/raise_http_exception_400", "4xx", 400, {}),
+        ("/raise_http_exception_404", "4xx", 404, {}),
+        ("/raise_request_validation_error", "4xx", 400, {"params": {"n": "x"}}),
+        ("/raise_vllm_validation_error", "4xx", 400, {}),
+        ("/raise_vllm_not_found_error", "4xx", 404, {}),
+        ("/raise_vllm_server_error", "5xx", 500, {}),
+        ("/raise_value_error", "4xx", 400, {}),
+        ("/raise_type_error", "4xx", 400, {}),
+        ("/raise_overflow_error", "4xx", 400, {}),
+        ("/raise_not_implemented_error", "5xx", 501, {}),
+        ("/raise_runtime_error", "5xx", 500, {}),
+        ("/success", "2xx", 200, {}),
     ],
     ids=[
+        "HTTPException(400)->4xx",
+        "HTTPException(404)->4xx",
+        "RequestValidationError->4xx",
+        "VLLMValidationError->4xx",
+        "VLLMNotFoundError->4xx",
+        "VLLMServerError->5xx",
         "ValueError->4xx",
         "TypeError->4xx",
         "OverflowError->4xx",
-        "VLLMValidationError->4xx",
-        "VLLMNotFoundError->4xx",
-        "HTTPException(400)->4xx",
-        "HTTPException(404)->4xx",
         "NotImplementedError->5xx",
         "RuntimeError->5xx",
         "success->2xx",
@@ -164,6 +181,7 @@ async def test_http_requests_total_records_correct_status(
     endpoint,
     expected_status_group,
     expected_http_code,
+    request_kwargs,
 ):
     """Verify that http_requests_total records the correct status group.
 
@@ -177,7 +195,7 @@ async def test_http_requests_total_records_correct_status(
     async with httpx.AsyncClient(
         transport=transport, base_url="http://testserver"
     ) as client:
-        response = await client.get(endpoint)
+        response = await client.get(endpoint, **request_kwargs)
 
     # Verify the HTTP response code returned to the client is correct
     assert response.status_code == expected_http_code, (
diff --git a/tests/entrypoints/unit_tests/test_chat_utils.py b/tests/entrypoints/unit_tests/test_chat_utils.py
index 82b262321d6..b5a35c2cb25 100644
--- a/tests/entrypoints/unit_tests/test_chat_utils.py
+++ b/tests/entrypoints/unit_tests/test_chat_utils.py
@@ -21,6 +21,7 @@ from vllm.entrypoints.chat_utils import (
     parse_chat_messages,
     parse_chat_messages_async,
 )
+from vllm.exceptions import VLLMValidationError
 from vllm.inputs import MultiModalDataDict, MultiModalUUIDDict
 from vllm.multimodal.utils import (
     encode_audio_url,
@@ -1504,7 +1505,7 @@ def test_parse_chat_messages_rejects_too_many_images_in_one_message(
             "ignore",
             message="coroutine 'async_get_and_parse_image' was never awaited",
         )
-        with pytest.raises(ValueError, match="At most"):
+        with pytest.raises(VLLMValidationError, match="At most"):
             parse_chat_messages(
                 [
                     {
@@ -1540,7 +1541,7 @@ def test_parse_chat_messages_rejects_too_many_images_across_messages(
             "ignore",
             message="coroutine 'async_get_and_parse_image' was never awaited",
         )
-        with pytest.raises(ValueError, match="At most"):
+        with pytest.raises(VLLMValidationError, match="At most"):
             parse_chat_messages(
                 [
                     {
diff --git a/tests/entrypoints/weight_transfer/test_weight_transfer_llm.py b/tests/entrypoints/weight_transfer/test_weight_transfer_llm.py
index 9088b3c5e8d..31d562e5ccf 100644
--- a/tests/entrypoints/weight_transfer/test_weight_transfer_llm.py
+++ b/tests/entrypoints/weight_transfer/test_weight_transfer_llm.py
@@ -234,6 +234,7 @@ def test_update_weights_calls_engine():
             assert shapes == test_shapes
 
         llm.finish_weight_update()
+        assert llm.get_weight_version() == "default"
 
 
 @create_new_process_for_each_test()
@@ -259,6 +260,8 @@ def test_full_weight_transfer_flow():
             weight_transfer_config=WeightTransferConfig(backend="nccl"),
         )
 
+        assert llm.get_weight_version() == "default"
+
         # Step 1: Initialize weight transfer engine
         llm.init_weight_transfer_engine(
             WeightTransferInitRequest(init_info={"test_param": "flow_test"})
@@ -278,8 +281,15 @@ def test_full_weight_transfer_flow():
             )
         )
 
+        assert llm.get_weight_version() == "default"
+
         # Step 4: Finish weight update
-        llm.finish_weight_update()
+        llm.finish_weight_update("step-42")
+
+        assert llm.get_weight_version() == "step-42"
+
+        llm.update_weight_version("manual-version")
+        assert llm.get_weight_version() == "manual-version"
 
         # Verify the full flow completed
         def check_flow(self):
diff --git a/tests/evals/gpt_oss/configs/gpt-oss-20b-xpu-baseline.yaml b/tests/evals/gpt_oss/configs/gpt-oss-20b-xpu-baseline.yaml
new file mode 100644
index 00000000000..78a583888ba
--- /dev/null
+++ b/tests/evals/gpt_oss/configs/gpt-oss-20b-xpu-baseline.yaml
@@ -0,0 +1,5 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+model_name: openai/gpt-oss-20b
+metric_threshold: 0.568
+reasoning_effort: low
diff --git a/tests/evals/gpt_oss/configs/gpt-oss-20b-xpu-triton-attn.yaml b/tests/evals/gpt_oss/configs/gpt-oss-20b-xpu-triton-attn.yaml
new file mode 100644
index 00000000000..e711ffb331e
--- /dev/null
+++ b/tests/evals/gpt_oss/configs/gpt-oss-20b-xpu-triton-attn.yaml
@@ -0,0 +1,6 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+model_name: openai/gpt-oss-20b
+metric_threshold: 0.568
+reasoning_effort: low
+server_args: "--attention-backend TRITON_ATTN"
diff --git a/tests/evals/gpt_oss/configs/models-xpu.txt b/tests/evals/gpt_oss/configs/models-xpu.txt
new file mode 100644
index 00000000000..a9de9266b14
--- /dev/null
+++ b/tests/evals/gpt_oss/configs/models-xpu.txt
@@ -0,0 +1,3 @@
+# Intel XPU model configurations for GPQA evaluation
+gpt-oss-20b-xpu-baseline.yaml
+gpt-oss-20b-xpu-triton-attn.yaml
diff --git a/tests/evals/gsm8k/configs/Qwen3.5-397B-A17B-NVFP4-DEP2-MTP.yaml b/tests/evals/gsm8k/configs/Qwen3.5-397B-A17B-NVFP4-DEP2-MTP.yaml
index d247515a0f0..921365ae686 100644
--- a/tests/evals/gsm8k/configs/Qwen3.5-397B-A17B-NVFP4-DEP2-MTP.yaml
+++ b/tests/evals/gsm8k/configs/Qwen3.5-397B-A17B-NVFP4-DEP2-MTP.yaml
@@ -3,8 +3,9 @@ accuracy_threshold: 0.88
 tolerance: 0.03
 num_questions: 1319
 num_fewshot: 5
+max_tokens: 12000
 server_args: >-
-  --max-model-len 4096
+  --max-model-len 16384
   --data-parallel-size 2
   --enable-expert-parallel
   --max-num-seqs 384
diff --git a/tests/kernels/attention/test_rocm_triton_attn_dsv4.py b/tests/kernels/attention/test_rocm_triton_attn_dsv4.py
index 77e068ab171..6fe2a3e7758 100644
--- a/tests/kernels/attention/test_rocm_triton_attn_dsv4.py
+++ b/tests/kernels/attention/test_rocm_triton_attn_dsv4.py
@@ -1,6 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+from types import SimpleNamespace
+
 import pytest
 import torch
 
@@ -196,6 +198,68 @@ def _ragged_from_rows(
     )
 
 
+@torch.inference_mode()
+def test_paged_mqa_logits_do_not_contain_nan(monkeypatch) -> None:
+    from vllm._aiter_ops import rocm_aiter_ops
+    from vllm.v1.attention.ops import rocm_aiter_mla_sparse as mod
+
+    device = torch.device("cuda")
+
+    class FakeWorkspaceManager:
+        def get_simultaneous(self, *shapes_and_dtypes):
+            return [
+                torch.empty(shape, dtype=dtype, device=device)
+                for shape, dtype in shapes_and_dtypes
+            ]
+
+    def fake_paged_mqa_logits(
+        q_fp8,
+        kv_cache_fp8,
+        weights,
+        out_logits,
+        context_lens,
+        block_tables,
+        max_seq_len,
+        **kwargs,
+    ):
+        del (
+            q_fp8,
+            kv_cache_fp8,
+            weights,
+            context_lens,
+            block_tables,
+            max_seq_len,
+            kwargs,
+        )
+        out_logits.fill_(float("nan"))
+
+    monkeypatch.setattr(mod, "_ON_GFX942", False)
+    monkeypatch.setattr(mod, "_ON_GFX950", True)
+    monkeypatch.setattr(rocm_aiter_ops, "is_enabled", lambda: True)
+    monkeypatch.setattr(
+        mod,
+        "paged_mqa_logits_module",
+        lambda: SimpleNamespace(deepgemm_fp8_paged_mqa_logits=fake_paged_mqa_logits),
+    )
+    monkeypatch.setattr(
+        mod, "current_workspace_manager", lambda: FakeWorkspaceManager()
+    )
+
+    q_fp8 = torch.empty((1, 1, 1, 1), dtype=torch.uint8, device=device)
+    kv_cache_fp8 = torch.empty((1, 1, 1, 5), dtype=torch.uint8, device=device)
+    logits = mod.rocm_fp8_paged_mqa_logits(
+        q_fp8,
+        kv_cache_fp8,
+        torch.empty((1, 1), dtype=torch.float32, device=device),
+        torch.ones(1, dtype=torch.int32, device=device),
+        torch.zeros((1, 1), dtype=torch.int32, device=device),
+        torch.empty(0, dtype=torch.int32, device=device),
+        1,
+    )
+
+    assert not torch.isnan(logits).any()
+
+
 @torch.inference_mode()
 def test_compute_global_topk_ragged_indices_and_indptr() -> None:
     from vllm.models.deepseek_v4.amd.rocm import (
diff --git a/tests/kernels/core/test_mrope.py b/tests/kernels/core/test_mrope.py
index 29051b4a00c..6f64fabfe9b 100644
--- a/tests/kernels/core/test_mrope.py
+++ b/tests/kernels/core/test_mrope.py
@@ -38,6 +38,7 @@ def generate_test_data(
 
 class MRoPETestInfo(NamedTuple):
     model_name: str
+    is_neox_style: bool = True
     # https://github.com/pytorch/pytorch/blob/main/torch/testing/_comparison.py#L1317
     atol: float = 1e-2
     rtol: float = 1.6e-2
@@ -45,7 +46,10 @@ class MRoPETestInfo(NamedTuple):
 
 
 MODELS_TO_TEST = [
-    MRoPETestInfo(model_name="zai-org/GLM-4.1V-9B-Thinking"),
+    MRoPETestInfo(
+        model_name="zai-org/GLM-4.1V-9B-Thinking",
+        is_neox_style=False,
+    ),
     MRoPETestInfo(model_name="Qwen/Qwen2-VL-7B-Instruct"),
     MRoPETestInfo(model_name="Qwen/Qwen2-VL-72B-Instruct"),
     MRoPETestInfo(model_name="Qwen/Qwen2.5-VL-72B-Instruct"),
@@ -92,7 +96,7 @@ def test_mrope(
         if hasattr(config, "head_dim")
         else config.hidden_size // total_num_heads
     )
-    is_neox_style = True
+    is_neox_style = model_info.is_neox_style
 
     max_position = config.max_position_embeddings
 
@@ -162,7 +166,7 @@ def test_mrope_torch_compile_tracing(
         if hasattr(config, "head_dim")
         else config.hidden_size // total_num_heads
     )
-    is_neox_style = True
+    is_neox_style = model_info.is_neox_style
     max_position = config.max_position_embeddings
 
     mrope_helper_class = get_rope(
diff --git a/tests/kernels/moe/test_flashinfer_cutedsl_nvfp4_moe.py b/tests/kernels/moe/test_flashinfer_cutedsl_nvfp4_moe.py
new file mode 100644
index 00000000000..a7a7c5251bc
--- /dev/null
+++ b/tests/kernels/moe/test_flashinfer_cutedsl_nvfp4_moe.py
@@ -0,0 +1,230 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for FlashInfer CuTeDSL NVFP4 MoE."""
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+import vllm.model_executor.layers.fused_moe.modular_kernel as mk
+from tests.kernels.quantization.nvfp4_utils import (
+    FLOAT4_E2M1_MAX,
+    FLOAT8_E4M3_MAX,
+    break_fp4_bytes,
+)
+from tests.kernels.utils import torch_moe
+from vllm import _custom_ops as ops
+from vllm.config import ParallelConfig, VllmConfig, set_current_vllm_config
+from vllm.model_executor.layers.fused_moe import fused_topk
+from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+from vllm.model_executor.layers.fused_moe.all2all_utils import (
+    maybe_make_prepare_finalize,
+)
+from vllm.model_executor.layers.fused_moe.config import (
+    FusedMoEConfig,
+    FusedMoEParallelConfig,
+    RoutingMethodType,
+    nvfp4_moe_quant_config,
+)
+from vllm.model_executor.layers.fused_moe.experts.flashinfer_cutedsl_moe import (
+    FlashInferCuteDSLExperts,
+)
+from vllm.model_executor.layers.quantization.utils.flashinfer_fp4_moe import (
+    prepare_nvfp4_moe_layer_for_flashinfer_cutedsl,
+)
+from vllm.platforms import current_platform
+from vllm.utils.flashinfer import has_flashinfer_cutedsl_moe_nvfp4
+from vllm.utils.math_utils import next_power_of_2
+from vllm.utils.torch_utils import set_random_seed
+
+if not has_flashinfer_cutedsl_moe_nvfp4() or not (
+    current_platform.is_device_capability_family(100)
+):
+    pytest.skip(
+        "Requires FlashInfer CuTeDSL NVFP4 MoE on SM100",
+        allow_module_level=True,
+    )
+
+
+def _quantize_nvfp4_linear(
+    weight: torch.Tensor,
+) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    weights_q = []
+    scales = []
+    global_scales = []
+    for expert_weight in weight:
+        global_scale = (
+            FLOAT8_E4M3_MAX * FLOAT4_E2M1_MAX / expert_weight.abs().max()
+        ).to(torch.float32)
+        weight_q, scale = ops.scaled_fp4_quant(
+            expert_weight,
+            global_scale,
+            is_sf_swizzled_layout=False,
+        )
+        weights_q.append(weight_q)
+        scales.append(scale)
+        global_scales.append(global_scale)
+    return torch.stack(weights_q), torch.stack(scales), torch.stack(global_scales)
+
+
+def _dequantize_nvfp4_linear(
+    tensor_fp4: torch.Tensor,
+    tensor_sf: torch.Tensor,
+    global_scale: torch.Tensor,
+    dtype: torch.dtype,
+) -> torch.Tensor:
+    assert tensor_fp4.dtype == torch.uint8
+    m, packed_k = tensor_fp4.shape
+    k = packed_k * 2
+    tensor_f32 = break_fp4_bytes(tensor_fp4, torch.float32)
+    tensor_f32 = tensor_f32.reshape(m, k // 16, 16)
+    tensor_sf = tensor_sf.view(torch.float8_e4m3fn).to(torch.float32)
+    tensor_sf = tensor_sf[:, : k // 16] / global_scale
+    return (tensor_f32 * tensor_sf.unsqueeze(-1)).reshape(m, k).to(dtype)
+
+
+@pytest.mark.parametrize("m,n,k,e,topk", [(16, 128, 512, 4, 2)])
+@pytest.mark.parametrize("dtype", [torch.bfloat16])
+@torch.inference_mode()
+def test_flashinfer_cutedsl_fp4_moe_relu2_no_mul(
+    m: int,
+    n: int,
+    k: int,
+    e: int,
+    topk: int,
+    dtype: torch.dtype,
+    workspace_init,
+):
+    set_random_seed(7)
+    with set_current_vllm_config(
+        VllmConfig(parallel_config=ParallelConfig(pipeline_parallel_size=1))
+    ):
+        hidden_states = torch.randn((m, k), device="cuda", dtype=dtype) / 10
+
+        w1 = torch.randn((e, n, k), device="cuda", dtype=dtype) / 15
+        w2 = torch.randn((e, k, n), device="cuda", dtype=dtype) / 15
+        w1_q, w1_scale, w1_global_scale = _quantize_nvfp4_linear(w1)
+        w2_q, w2_scale, w2_global_scale = _quantize_nvfp4_linear(w2)
+
+        score = torch.randn((m, e), device="cuda", dtype=dtype)
+        topk_weights, topk_ids, _ = fused_topk(
+            hidden_states, score, topk, renormalize=False
+        )
+
+        activation = MoEActivation.RELU2_NO_MUL
+        fake_layer = SimpleNamespace(activation=activation)
+        a1_scale = torch.ones(1, device="cuda", dtype=torch.float32)
+        a2_scale = torch.ones(1, device="cuda", dtype=torch.float32)
+        (
+            w1_cutedsl,
+            w1_scale_cutedsl,
+            w1_alpha,
+            a1_scale,
+            w2_cutedsl,
+            w2_scale_cutedsl,
+            w2_alpha,
+            a2_scale,
+        ) = prepare_nvfp4_moe_layer_for_flashinfer_cutedsl(
+            layer=fake_layer,
+            w13=w1_q,
+            w13_scale=w1_scale,
+            w13_scale_2=(1.0 / w1_global_scale),
+            a13_scale=a1_scale,
+            w2=w2_q,
+            w2_scale=w2_scale,
+            w2_scale_2=(1.0 / w2_global_scale),
+            a2_scale=a2_scale,
+        )
+        quant_config = nvfp4_moe_quant_config(
+            g1_alphas=w1_alpha,
+            g2_alphas=w2_alpha,
+            a1_gscale=(1.0 / a1_scale),
+            a2_gscale=(1.0 / a2_scale),
+            w1_scale=w1_scale_cutedsl,
+            w2_scale=w2_scale_cutedsl,
+            is_scale_swizzled=False,
+        )
+        moe_config = FusedMoEConfig(
+            num_experts=e,
+            experts_per_token=topk,
+            hidden_dim=k,
+            intermediate_size=n,
+            num_local_experts=e,
+            num_logical_experts=e,
+            activation=activation,
+            device="cuda",
+            moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
+            in_dtype=dtype,
+            routing_method=RoutingMethodType.TopK,
+            max_num_tokens=next_power_of_2(m),
+        )
+
+        cutedsl_experts = mk.FusedMoEKernel(
+            maybe_make_prepare_finalize(
+                moe=moe_config,
+                quant_config=quant_config,
+                allow_new_interface=True,
+                use_monolithic=False,
+            ),
+            FlashInferCuteDSLExperts(
+                moe_config=moe_config,
+                quant_config=quant_config,
+            ),
+        )
+
+        cutedsl_output = cutedsl_experts.apply(
+            hidden_states=hidden_states,
+            w1=w1_cutedsl,
+            w2=w2_cutedsl,
+            topk_weights=topk_weights,
+            topk_ids=topk_ids,
+            activation=activation,
+            global_num_experts=e,
+            expert_map=None,
+            apply_router_weight_on_input=False,
+        )
+
+        a_global_scale = torch.ones(1, device="cuda", dtype=torch.float32)
+        a_q, a_scale = ops.scaled_fp4_quant(
+            hidden_states,
+            a_global_scale,
+            is_sf_swizzled_layout=False,
+        )
+        a_in_dtype = _dequantize_nvfp4_linear(
+            a_q,
+            a_scale,
+            a_global_scale,
+            dtype=dtype,
+        )
+
+        w1_d = torch.empty((e, n, k), device="cuda", dtype=dtype)
+        w2_d = torch.empty((e, k, n), device="cuda", dtype=dtype)
+        for idx in range(e):
+            w1_d[idx] = _dequantize_nvfp4_linear(
+                w1_q[idx],
+                w1_scale[idx],
+                w1_global_scale[idx],
+                dtype=dtype,
+            )
+            w2_d[idx] = _dequantize_nvfp4_linear(
+                w2_q[idx],
+                w2_scale[idx],
+                w2_global_scale[idx],
+                dtype=dtype,
+            )
+
+        torch_output = torch_moe(
+            a_in_dtype,
+            w1_d,
+            w2_d,
+            score,
+            topk,
+            activation=activation,
+        )
+        torch.testing.assert_close(
+            torch_output,
+            cutedsl_output,
+            atol=2e-1,
+            rtol=2e-1,
+        )
diff --git a/tests/kernels/moe/test_moe.py b/tests/kernels/moe/test_moe.py
index 9c43aa97409..a519980dead 100644
--- a/tests/kernels/moe/test_moe.py
+++ b/tests/kernels/moe/test_moe.py
@@ -36,6 +36,9 @@ from vllm.model_executor.layers.fused_moe.experts.marlin_moe import (
     batched_fused_marlin_moe,
     fused_marlin_moe,
 )
+from vllm.model_executor.layers.fused_moe.utils import (
+    moe_use_td_hw_supported,
+)
 from vllm.model_executor.layers.quantization.utils.marlin_utils import (
     marlin_permute_bias,
 )
@@ -53,9 +56,12 @@ from vllm.model_executor.layers.quantization.utils.marlin_utils_test import (
 from vllm.model_executor.layers.quantization.utils.quant_utils import quantize_weights
 from vllm.platforms import current_platform
 from vllm.scalar_type import ScalarType, scalar_types
+from vllm.triton_utils import tl
 from vllm.utils.math_utils import next_power_of_2
 from vllm.utils.torch_utils import set_random_seed
 
+DEVICE_TYPE = current_platform.device_type
+
 
 def iterative_moe(
     hidden_states: torch.Tensor,
@@ -289,6 +295,7 @@ def run_moe_test(
 @pytest.mark.parametrize("ep_size", EP_SIZE)
 @pytest.mark.parametrize("dtype", [torch.bfloat16])
 @pytest.mark.parametrize("padding", [True, False])
+@pytest.mark.parametrize("use_td", [False, True])
 def test_fused_moe(
     m: int,
     n: int,
@@ -298,9 +305,19 @@ def test_fused_moe(
     ep_size: int,
     dtype: torch.dtype,
     padding: bool,
+    use_td: bool,
     monkeypatch,
     workspace_init,
 ):
+    if use_td and not hasattr(tl, "make_tensor_descriptor"):
+        pytest.skip("Triton < 3.6 lacks tl.make_tensor_descriptor")
+    if use_td and not moe_use_td_hw_supported():
+        pytest.skip(
+            "tensor_descriptor.gather requires XPU or NVIDIA Blackwell "
+            "(sm100+); lowers to tile::gather4 (tcgen05/TMEM), which ptxas "
+            "rejects on Hopper (sm90) and earlier"
+        )
+    monkeypatch.setenv("VLLM_TRITON_USE_TD", "1" if use_td else "0")
     set_random_seed(7)
 
     #
@@ -311,17 +328,17 @@ def test_fused_moe(
     # Setup test data
     #
 
-    a = torch.randn((m, k), device="cuda", dtype=dtype) / 10
-    w1 = torch.randn((e, 2 * n, k), device="cuda", dtype=dtype) / 10
-    w2 = torch.randn((e, k, n), device="cuda", dtype=dtype) / 10
+    a = torch.randn((m, k), device=DEVICE_TYPE, dtype=dtype) / 10
+    w1 = torch.randn((e, 2 * n, k), device=DEVICE_TYPE, dtype=dtype) / 10
+    w2 = torch.randn((e, k, n), device=DEVICE_TYPE, dtype=dtype) / 10
 
-    score = torch.randn((m, e), device="cuda", dtype=dtype)
+    score = torch.randn((m, e), device=DEVICE_TYPE, dtype=dtype)
 
     if ep_size > 1:
         local_e = e // ep_size
-        e_ids = torch.randint(0, e, (local_e,), device="cuda", dtype=torch.int32)
-        e_map = torch.full((e,), -1, device="cuda", dtype=torch.int32)
-        e_map[e_ids] = torch.arange(local_e, device="cuda", dtype=torch.int32)
+        e_ids = torch.randint(0, e, (local_e,), device=DEVICE_TYPE, dtype=torch.int32)
+        e_map = torch.full((e,), -1, device=DEVICE_TYPE, dtype=torch.int32)
+        e_map[e_ids] = torch.arange(local_e, device=DEVICE_TYPE, dtype=torch.int32)
         w1 = w1[e_ids]
         w2 = w2[e_ids]
     else:
diff --git a/tests/kernels/moe/test_mxfp8_aiter_backend_selection.py b/tests/kernels/moe/test_mxfp8_aiter_backend_selection.py
index 7c2fdbabe29..aa1747f4153 100644
--- a/tests/kernels/moe/test_mxfp8_aiter_backend_selection.py
+++ b/tests/kernels/moe/test_mxfp8_aiter_backend_selection.py
@@ -19,9 +19,17 @@ if not current_platform.is_rocm():
     pytest.skip("This test can only run on ROCm.", allow_module_level=True)
 
 from tests.kernels.moe.utils import make_dummy_moe_config  # noqa: E402
+from vllm.model_executor.layers.fused_moe.activation import (  # noqa: E402
+    MoEActivation,
+)
 from vllm.model_executor.layers.fused_moe.experts.aiter_mxfp8_moe import (  # noqa: E402
+    _AITER_SWIGLU_ALPHA,
+    _AITER_SWIGLU_BETA,
     AiterMxfp8Experts,
 )
+from vllm.model_executor.layers.fused_moe.experts.mxfp8_native_moe import (  # noqa: E402
+    Mxfp8NativeTritonExperts,
+)
 from vllm.model_executor.layers.fused_moe.modular_kernel import (  # noqa: E402
     FusedMoEActivationFormat,
 )
@@ -33,6 +41,7 @@ from vllm.model_executor.layers.fused_moe.oracle.mxfp8 import (  # noqa: E402
     _SUPPORTED_BACKENDS,
     _mxfp8_backend_to_kernel_cls,
     _select_kernel_cls,
+    select_mxfp8_moe_backend,
 )
 from vllm.model_executor.layers.quantization.utils.quant_utils import (  # noqa: E402
     kMxfp8Dynamic,
@@ -43,7 +52,17 @@ _AITER_MOD = "vllm.model_executor.layers.fused_moe.experts.aiter_mxfp8_moe"
 
 
 def _config(ep_size: int = 1):
-    cfg = make_dummy_moe_config(num_experts=128, experts_per_token=4, hidden_dim=6144)
+    # AiterMxfp8Experts hardcodes SwiGLU-OAI: match its required activation and
+    # alpha/beta so is_supported_config doesn't reject the config on those grounds.
+    cfg = make_dummy_moe_config(
+        num_experts=128,
+        experts_per_token=4,
+        hidden_dim=6144,
+        activation=MoEActivation.SWIGLUOAI_UNINTERLEAVE,
+    )
+    cfg = dataclasses.replace(
+        cfg, swiglu_alpha=_AITER_SWIGLU_ALPHA, swiglu_beta=_AITER_SWIGLU_BETA
+    )
     if ep_size != 1:
         cfg = dataclasses.replace(
             cfg,
@@ -76,12 +95,6 @@ def test_aiter_mxfp8_registered():
     ]
 
 
-def test_triton_selectable():
-    assert _BACKEND_NAME_MAP["triton"] is Fp8MoeBackend.TRITON_MXFP8
-    # Not auto-selected (only reachable explicitly), so FlyDSL still wins auto.
-    assert Fp8MoeBackend.TRITON_MXFP8 not in _SUPPORTED_BACKENDS
-
-
 @pytest.mark.parametrize("ep_size", [1, 2])
 def test_ep_supported(ep_size):
     """FlyDSL accepts both TP and EP: apply() forwards expert_map as expert_mask."""
@@ -133,3 +146,20 @@ def test_explicit_moe_backend_aiter():
         pytest.raises(ValueError, match="flydsl package"),
     ):
         _select_kernel_cls(Fp8MoeBackend.AITER_MXFP8, _config(1))
+
+
+def test_gfx950_picks_aiter():
+    """Auto-select on real ROCm hardware with flydsl usable -> FlyDSL wins."""
+    with _flydsl_installed(True):
+        backend, experts_cls = select_mxfp8_moe_backend(_config())
+    assert backend is Fp8MoeBackend.AITER_MXFP8
+    assert experts_cls is AiterMxfp8Experts
+
+
+def test_gfx942_picks_triton():
+    """flydsl unusable (e.g. gfx942, no FlyDSL support) -> native Triton
+    dot_scaled backend wins instead."""
+    with _flydsl_installed(False):
+        backend, experts_cls = select_mxfp8_moe_backend(_config())
+    assert backend is Fp8MoeBackend.TRITON_MXFP8
+    assert experts_cls is Mxfp8NativeTritonExperts
diff --git a/tests/kernels/moe/test_ocp_mx_moe.py b/tests/kernels/moe/test_ocp_mx_moe.py
index 2f819c09aaa..7e8d4e3028a 100644
--- a/tests/kernels/moe/test_ocp_mx_moe.py
+++ b/tests/kernels/moe/test_ocp_mx_moe.py
@@ -1558,3 +1558,51 @@ def test_mxfp4_emulation_rounds_up_to_block_size(
     # The block-scale buffer (dim // OCP_MX_BLOCK_SIZE) must not floor-truncate.
     assert rounded_hidden % OCP_MX_BLOCK_SIZE == 0
     assert rounded_intermediate % OCP_MX_BLOCK_SIZE == 0
+
+
+def test_select_mxfp4_moe_backend_raises_with_unsupported_reasons(
+    monkeypatch: pytest.MonkeyPatch,
+):
+    """
+    select_mxfp4_moe_backend() must raise NotImplementedError, with the
+    collected per-backend unsupported reasons in the message, when no
+    backend supports the requested deployment configuration.
+    """
+    import vllm.model_executor.layers.fused_moe.oracle.mxfp4 as mxfp4_oracle
+    from vllm.model_executor.layers.fused_moe import FusedMoEConfig
+    from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+    from vllm.model_executor.layers.fused_moe.config import (
+        FusedMoEParallelConfig,
+        RoutingMethodType,
+    )
+
+    class UnsupportedExperts:
+        @staticmethod
+        def is_supported_config(
+            cls, moe_config, weight_key, activation_key, activation_format
+        ):
+            return False, f"unsupported reason for {cls.__name__}"
+
+    monkeypatch.setattr(
+        mxfp4_oracle, "backend_to_kernel_cls", lambda backend: [UnsupportedExperts]
+    )
+    monkeypatch.setattr(mxfp4_oracle, "_user_moe_activation_override", lambda: None)
+    monkeypatch.setattr(current_platform, "is_xpu", lambda: False)
+    monkeypatch.setattr(current_platform, "is_cpu", lambda: False)
+
+    moe_config = FusedMoEConfig(
+        num_experts=8,
+        experts_per_token=2,
+        hidden_dim=256,
+        intermediate_size=256,
+        num_local_experts=8,
+        num_logical_experts=8,
+        moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
+        activation=MoEActivation.SILU,
+        in_dtype=torch.bfloat16,
+        device="cpu",
+        routing_method=RoutingMethodType.Renormalize,
+    )
+
+    with pytest.raises(NotImplementedError, match="Unsupported reasons"):
+        mxfp4_oracle.select_mxfp4_moe_backend(moe_config)
diff --git a/tests/kernels/moe/test_routing.py b/tests/kernels/moe/test_routing.py
index 62a4968a0d1..9b12fef0454 100644
--- a/tests/kernels/moe/test_routing.py
+++ b/tests/kernels/moe/test_routing.py
@@ -8,9 +8,19 @@ import torch
 
 from vllm._aiter_ops import rocm_aiter_ops
 from vllm.distributed.eplb.eplb_state import EplbLayerState
+from vllm.model_executor.layers.fused_moe.config import RoutingMethodType
 from vllm.model_executor.layers.fused_moe.router.base_router import (
     eplb_map_to_physical_and_record,
 )
+from vllm.model_executor.layers.fused_moe.router.fused_topk_bias_router import (
+    FusedTopKBiasRouter,
+)
+from vllm.model_executor.layers.fused_moe.router.fused_topk_router import (
+    FusedTopKRouter,
+)
+from vllm.model_executor.layers.fused_moe.router.grouped_topk_router import (
+    GroupedTopKRouter,
+)
 from vllm.model_executor.layers.fused_moe.router.router_factory import (
     create_fused_moe_router,
 )
@@ -36,6 +46,112 @@ TOP_KS = [2, 4, 6]
 NUM_EXPERTS = [8, 16, 64]
 
 
+def test_degenerate_grouped_config_uses_standard_topk() -> None:
+    router = create_fused_moe_router(
+        top_k=4,
+        global_num_experts=128,
+        use_grouped_topk=True,
+        num_expert_group=1,
+        topk_group=1,
+        scoring_func="softmax",
+        renormalize=True,
+    )
+
+    assert isinstance(router, FusedTopKRouter)
+    hidden_states, router_logits = make_test_data(32, 256, 128)
+
+    topk_weights, topk_ids = router.select_experts(hidden_states, router_logits)
+    baseline_weights, baseline_ids = baseline_fused_topk(
+        router_logits,
+        top_k=4,
+        renormalize=True,
+    )
+
+    assert_routing_results_close(
+        topk_weights,
+        topk_ids,
+        baseline_weights,
+        baseline_ids,
+    )
+
+
+def test_multiple_expert_groups_use_grouped_topk() -> None:
+    router = create_fused_moe_router(
+        top_k=4,
+        global_num_experts=128,
+        use_grouped_topk=True,
+        num_expert_group=8,
+        topk_group=4,
+        scoring_func="softmax",
+        renormalize=True,
+    )
+
+    assert isinstance(router, GroupedTopKRouter)
+
+
+def test_degenerate_grouped_config_with_bias_uses_topk_bias() -> None:
+    router = create_fused_moe_router(
+        top_k=4,
+        global_num_experts=128,
+        use_grouped_topk=True,
+        num_expert_group=1,
+        topk_group=1,
+        scoring_func="softmax",
+        renormalize=True,
+        e_score_correction_bias=torch.empty(128),
+    )
+
+    assert isinstance(router, FusedTopKBiasRouter)
+
+
+def test_degenerate_grouped_config_with_bias_keeps_routed_scale() -> None:
+    router = create_fused_moe_router(
+        top_k=4,
+        global_num_experts=128,
+        use_grouped_topk=True,
+        num_expert_group=1,
+        topk_group=1,
+        scoring_func="softmax",
+        renormalize=True,
+        routed_scaling_factor=1.1,
+        e_score_correction_bias=torch.empty(128),
+    )
+
+    assert isinstance(router, FusedTopKBiasRouter)
+    assert router.routed_scaling_factor == 1.1
+
+
+def test_degenerate_deepseek_v3_routing_stays_grouped() -> None:
+    router = create_fused_moe_router(
+        top_k=4,
+        global_num_experts=128,
+        use_grouped_topk=True,
+        num_expert_group=1,
+        topk_group=1,
+        scoring_func="sigmoid",
+        renormalize=True,
+        e_score_correction_bias=torch.empty(128),
+    )
+
+    assert isinstance(router, GroupedTopKRouter)
+    assert router.routing_method_type == RoutingMethodType.DeepSeekV3
+
+
+def test_single_expert_group_with_non_unit_scale_uses_grouped_topk() -> None:
+    router = create_fused_moe_router(
+        top_k=4,
+        global_num_experts=128,
+        use_grouped_topk=True,
+        num_expert_group=1,
+        topk_group=1,
+        scoring_func="softmax",
+        renormalize=True,
+        routed_scaling_factor=1.1,
+    )
+
+    assert isinstance(router, GroupedTopKRouter)
+
+
 def setup_eplb_state(
     enable_eplb: bool, global_num_experts: int
 ) -> EplbLayerState | None:
diff --git a/tests/kernels/test_fused_minimax_m3_qknorm_rope_kv_insert.py b/tests/kernels/test_fused_minimax_m3_qknorm_rope_kv_insert.py
index 626b06290e0..9c4a996438e 100644
--- a/tests/kernels/test_fused_minimax_m3_qknorm_rope_kv_insert.py
+++ b/tests/kernels/test_fused_minimax_m3_qknorm_rope_kv_insert.py
@@ -147,8 +147,11 @@ def test_dense_norm_rope(num_tokens, num_heads, num_kv_heads):
         eps,
     ).view(num_tokens, kvsz)
 
-    torch.testing.assert_close(q_out, q_ref, rtol=1e-2, atol=1e-2)
-    torch.testing.assert_close(k_out, k_ref, rtol=1e-2, atol=1e-2)
+    # The fused kernel keeps an fp32 intermediate across norm->rope, while the
+    # reference materializes bf16 after the norm (the unfused boundary), so
+    # rounding-boundary elements can differ by ~1 bf16 ulp.
+    torch.testing.assert_close(q_out, q_ref, rtol=2e-2, atol=2e-2)
+    torch.testing.assert_close(k_out, k_ref, rtol=2e-2, atol=2e-2)
     # V is untouched.
     torch.testing.assert_close(v_out, v_in, rtol=0, atol=0)
 
@@ -255,8 +258,11 @@ def test_sparse_full(num_tokens, block_size, kv_cache_dtype):
         ik_orig.view(num_tokens, 1, HEAD_DIM), ik_w, positions, cos_sin, eps
     ).view(num_tokens, HEAD_DIM)
 
-    torch.testing.assert_close(q_out, q_ref, rtol=1e-2, atol=1e-2)
-    torch.testing.assert_close(k_out, k_ref, rtol=1e-2, atol=1e-2)
+    # The fused kernel keeps an fp32 intermediate across norm->rope, while the
+    # reference materializes bf16 after the norm (the unfused boundary), so
+    # rounding-boundary elements can differ by ~1 bf16 ulp.
+    torch.testing.assert_close(q_out, q_ref, rtol=2e-2, atol=2e-2)
+    torch.testing.assert_close(k_out, k_ref, rtol=2e-2, atol=2e-2)
     torch.testing.assert_close(index_q, iq_ref, rtol=1e-2, atol=1e-2)
     torch.testing.assert_close(index_k, ik_ref, rtol=1e-2, atol=1e-2)
 
@@ -376,8 +382,11 @@ def test_sparse_skip_index_branch(num_tokens, block_size, kv_cache_dtype):
         eps,
     ).view(num_tokens, kvsz)
 
-    torch.testing.assert_close(q_out, q_ref, rtol=1e-2, atol=1e-2)
-    torch.testing.assert_close(k_out, k_ref, rtol=1e-2, atol=1e-2)
+    # The fused kernel keeps an fp32 intermediate across norm->rope, while the
+    # reference materializes bf16 after the norm (the unfused boundary), so
+    # rounding-boundary elements can differ by ~1 bf16 ulp.
+    torch.testing.assert_close(q_out, q_ref, rtol=2e-2, atol=2e-2)
+    torch.testing.assert_close(k_out, k_ref, rtol=2e-2, atol=2e-2)
     torch.testing.assert_close(v_out, v_in, rtol=0, atol=0)
     torch.testing.assert_close(index_q_out, index_q_in, rtol=0, atol=0)
     torch.testing.assert_close(index_k_out, index_k_in, rtol=0, atol=0)
diff --git a/tests/lora/test_qwenvl.py b/tests/lora/test_qwenvl.py
index 3cbb534bdad..3362aa46ed2 100644
--- a/tests/lora/test_qwenvl.py
+++ b/tests/lora/test_qwenvl.py
@@ -186,6 +186,18 @@ QWEN25VL_MODEL_PATH = "Qwen/Qwen2.5-VL-3B-Instruct"
 QWEN3VL_MODEL_PATH = "Qwen/Qwen3-VL-4B-Instruct"
 
 
+def _enable_deterministic_lora_shrink(monkeypatch: pytest.MonkeyPatch) -> None:
+    # These tests assert exact greedy outputs. Force the Triton LoRA shrink
+    # kernel to use SPLIT_K=1 so it stores the complete reduction directly
+    # instead of accumulating split-K partial results with atomic_add. This
+    # targets reduction determinism, not full batch invariance.
+    monkeypatch.setenv("VLLM_BATCH_INVARIANT", "1")
+    # The kernel configuration reads VLLM_BATCH_INVARIANT at import time.
+    # Spawn the engine process so it observes this setting even if the LoRA
+    # Triton utilities were already imported during test collection.
+    monkeypatch.setenv("VLLM_WORKER_MULTIPROC_METHOD", "spawn")
+
+
 def test_qwen2vl_lora(qwen2vl_lora_files):
     """Test Qwen 2.0 VL model with LoRA"""
     config = TestConfig(model_path=QWEN2VL_MODEL_PATH, lora_path=qwen2vl_lora_files)
@@ -250,7 +262,12 @@ def test_qwen25vl_vision_lora(qwen25vl_vision_lora_files):
             )
 
 
-def test_qwen3vl_vision_lora(qwen3vl_vision_lora_files):
+def test_qwen3vl_vision_lora(
+    qwen3vl_vision_lora_files,
+    monkeypatch: pytest.MonkeyPatch,
+):
+    _enable_deterministic_lora_shrink(monkeypatch)
+
     config = TestConfig(
         model_path=QWEN3VL_MODEL_PATH,
         lora_path=qwen3vl_vision_lora_files,
@@ -273,6 +290,7 @@ def test_qwen2vl_multiple_lora_types(
     qwen2vl_language_lora_files,
     qwen2vl_vision_tower_connector_lora_files,
     qwen2vl_vision_tower_lora_files,
+    monkeypatch: pytest.MonkeyPatch,
 ):
     """
     Test multiple LoRA adapter types (language, vision tower + connector,
@@ -283,6 +301,8 @@ def test_qwen2vl_multiple_lora_types(
     the multimodal encoder cache correctly manages state transitions between
     language-only and vision-enabled LoRA adapters.
     """
+    _enable_deterministic_lora_shrink(monkeypatch)
+
     config = TestConfig(
         model_path=QWEN2VL_MODEL_PATH,
         # We'll override the lora_path for each specific test, but need to provide
diff --git a/tests/model_executor/layers/test_mla_short_prefill_indexer.py b/tests/model_executor/layers/test_mla_short_prefill_indexer.py
new file mode 100644
index 00000000000..6e1e10e8b45
--- /dev/null
+++ b/tests/model_executor/layers/test_mla_short_prefill_indexer.py
@@ -0,0 +1,165 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+import vllm.model_executor.layers.sparse_attn_indexer as sparse_indexer
+from vllm.config import CUDAGraphMode
+from vllm.v1.attention.backends.mla.indexer import DeepseekV32IndexerMetadata
+
+INDEXER_LAYER = "model.layers.0.self_attn.indexer.k_cache"
+MLA_LAYER = "model.layers.0.self_attn.attn"
+
+
+def make_indexer_metadata(
+    *,
+    num_decodes: int = 0,
+    num_decode_tokens: int = 0,
+    num_prefills: int = 1,
+    num_prefill_tokens: int = 1,
+    slot_mapping: torch.Tensor | None = None,
+) -> DeepseekV32IndexerMetadata:
+    if slot_mapping is None:
+        slot_mapping = torch.zeros(num_prefill_tokens, dtype=torch.long)
+    return DeepseekV32IndexerMetadata(
+        seq_lens=torch.empty(0, dtype=torch.int32),
+        max_seq_len=2048,
+        slot_mapping=slot_mapping,
+        num_decodes=num_decodes,
+        num_decode_tokens=num_decode_tokens,
+        num_prefills=num_prefills,
+        num_prefill_tokens=num_prefill_tokens,
+        prefill=SimpleNamespace(chunks=[]) if num_prefills else None,
+    )
+
+
+def make_mla_metadata(*, use_dense_mha: bool = True, num_decode_tokens: int = 0):
+    return SimpleNamespace(
+        num_decode_tokens=num_decode_tokens,
+        prefill=SimpleNamespace(use_dense_mha=use_dense_mha),
+    )
+
+
+@pytest.mark.parametrize(
+    "batch_kind",
+    ["short", "threshold_mismatch", "force_mqa", "mla_decode", "capture", "full"],
+)
+def test_short_prefill_updates_k_cache_before_scoring_decision(
+    monkeypatch: pytest.MonkeyPatch,
+    batch_kind: str,
+):
+    slot_mapping = torch.tensor([63, 64, 127, 128, -1])
+    mla_num_decode_tokens = 1 if batch_kind == "mla_decode" else 0
+    runtime_mode = (
+        CUDAGraphMode.FULL if batch_kind == "full" else CUDAGraphMode.PIECEWISE
+    )
+    should_skip = batch_kind in ("short", "threshold_mismatch")
+    num_decodes = int(batch_kind == "threshold_mismatch")
+    num_decode_tokens = 3 if batch_kind == "threshold_mismatch" else 0
+    num_prefills = 0 if batch_kind == "threshold_mismatch" else 2
+    num_prefill_tokens = 0 if batch_kind == "threshold_mismatch" else 5
+    if batch_kind == "threshold_mismatch":
+        # With MTP=3 the indexer threshold is four. A main MLA backend whose
+        # threshold is one (for example FlashMLA under DCP) still routes this
+        # three-token extend through dense prefill attention.
+        slot_mapping = slot_mapping[:3]
+    indexer_metadata = make_indexer_metadata(
+        num_decodes=num_decodes,
+        num_decode_tokens=num_decode_tokens,
+        num_prefills=num_prefills,
+        num_prefill_tokens=num_prefill_tokens,
+        slot_mapping=slot_mapping,
+    )
+    if indexer_metadata.num_decodes:
+        indexer_metadata.decode = object()
+    mla_metadata = make_mla_metadata(
+        use_dense_mha=batch_kind != "force_mqa",
+        num_decode_tokens=mla_num_decode_tokens,
+    )
+
+    observed: dict[str, object] = {}
+
+    monkeypatch.setattr(
+        sparse_indexer,
+        "get_forward_context",
+        lambda: SimpleNamespace(
+            attn_metadata={
+                INDEXER_LAYER: indexer_metadata,
+                MLA_LAYER: mla_metadata,
+            },
+            cudagraph_runtime_mode=runtime_mode,
+        ),
+    )
+    monkeypatch.setattr(
+        sparse_indexer.current_platform, "fp8_dtype", lambda: torch.float16
+    )
+    monkeypatch.setattr(
+        torch.cuda,
+        "is_current_stream_capturing",
+        lambda: batch_kind == "capture",
+    )
+
+    def record_cache_update(k, kv_cache, slots, block_size, scale_fmt):
+        observed.update(k=k.clone(), slots=slots)
+
+    monkeypatch.setattr(
+        sparse_indexer.ops, "indexer_k_quant_and_cache", record_cache_update
+    )
+
+    class ScoringReached(Exception):
+        pass
+
+    def scoring_trigger():
+        if should_skip:
+            pytest.fail("short dense-MHA prefill must not enter indexer scoring")
+        raise ScoringReached
+
+    def scoring_decode(*args):
+        raise ScoringReached
+
+    monkeypatch.setattr(sparse_indexer, "current_workspace_manager", scoring_trigger)
+    monkeypatch.setattr(
+        sparse_indexer,
+        "kv_cache_as_quant_view",
+        scoring_decode,
+    )
+
+    hidden_states = torch.full((7, 1), float("inf"))
+    k = torch.arange(28, dtype=torch.float32).reshape(7, 4)
+    topk_indices = torch.full((7, 2048), 17, dtype=torch.int32)
+
+    def run_indexer():
+        return sparse_indexer.sparse_attn_indexer(
+            hidden_states,
+            INDEXER_LAYER,
+            torch.empty(1),
+            torch.full((7, 1), float("inf")),
+            None,
+            k,
+            torch.full((7, 1), float("inf")),
+            128,
+            "ue8m0",
+            2048,
+            4,
+            4096,
+            4096,
+            topk_indices,
+            False,
+            False,
+            MLA_LAYER,
+        )
+
+    if should_skip:
+        assert run_indexer() is topk_indices
+        assert torch.all(topk_indices == 17)
+    else:
+        with pytest.raises(ScoringReached):
+            run_indexer()
+        assert torch.all(topk_indices == -1)
+
+    # K cache is always updated before the scoring decision.
+    torch.testing.assert_close(observed["k"], k[: slot_mapping.numel()])
+    assert observed["slots"] is slot_mapping
diff --git a/tests/models/inkling/test_moe_weight_layout.py b/tests/models/inkling/test_moe_weight_layout.py
index a7307e167a1..e15a573b845 100644
--- a/tests/models/inkling/test_moe_weight_layout.py
+++ b/tests/models/inkling/test_moe_weight_layout.py
@@ -187,6 +187,32 @@ def test_moe_loads_compressed_tensors_global_scale(
     assert loaded == [f"experts.routed_experts.{projection}_{scale_kind}_global_scale"]
 
 
+@pytest.mark.parametrize(("projection", "checkpoint_rows"), [("w13", 8), ("w2", 4)])
+def test_moe_loads_channelwise_scale_for_tp(
+    projection: str, checkpoint_rows: int
+) -> None:
+    param = torch.nn.Parameter(torch.empty(2, 4, 1))
+    experts = SimpleNamespace(
+        **{f"{projection}_weight_scale": param},
+        moe_config=SimpleNamespace(moe_parallel_config=SimpleNamespace(tp_rank=1)),
+    )
+    layer = SimpleNamespace(
+        experts=SimpleNamespace(routed_experts=experts),
+        _local_expert_slots=lambda: {0: 0, 2: 1},
+    )
+    checkpoint_scale = torch.arange(3 * checkpoint_rows).reshape(3, checkpoint_rows, 1)
+
+    loaded = moe.InklingMoE.load_expert_weight(
+        layer, f"experts.{projection}_weight_scale", checkpoint_scale
+    )
+
+    expected = checkpoint_scale[[0, 2]]
+    if projection == "w13":
+        expected = expected[:, 4:].reshape(2, 2, 2, 1).transpose(1, 2).flatten(1, 2)
+    torch.testing.assert_close(param, expected.float())
+    assert loaded == [f"experts.routed_experts.{projection}_weight_scale"]
+
+
 def test_sink_down_projection_is_packed_during_load(monkeypatch) -> None:
     monkeypatch.setattr(moe, "get_tensor_model_parallel_world_size", lambda: 2)
     monkeypatch.setattr(moe, "get_tensor_model_parallel_rank", lambda: 1)
diff --git a/tests/models/kimi_k3/test_sequence_parallel.py b/tests/models/kimi_k3/test_sequence_parallel.py
index 24da286b2ab..c65f9bd9532 100644
--- a/tests/models/kimi_k3/test_sequence_parallel.py
+++ b/tests/models/kimi_k3/test_sequence_parallel.py
@@ -12,6 +12,7 @@ from vllm.config import ParallelConfig
 from vllm.models.kimi_k3.nvidia import model as kimi_model
 from vllm.models.kimi_k3.nvidia import mtp as kimi_mtp
 from vllm.models.kimi_k3.nvidia.ops import sequence_parallel as sp_ops
+from vllm.platforms import current_platform
 
 
 class _IdentityNorm(nn.Module):
@@ -108,15 +109,27 @@ def test_sp_padding_mask_marks_added_rows(
     torch.testing.assert_close(actual, torch.tensor(expected))
 
 
-def test_moe_sequence_parallel_is_available_without_data_parallel():
+@pytest.mark.parametrize(
+    ("data_parallel_size", "expected"),
+    [
+        (1, False),
+        (2, True),
+    ],
+)
+def test_moe_sequence_parallel_requires_data_parallel(
+    monkeypatch,
+    data_parallel_size: int,
+    expected: bool,
+):
+    monkeypatch.setattr(current_platform, "device_count", lambda: 2)
     parallel_config = ParallelConfig(
         tensor_parallel_size=2,
-        data_parallel_size=1,
+        data_parallel_size=data_parallel_size,
         enable_expert_parallel=True,
         all2all_backend="allgather_reducescatter",
     )
 
-    assert parallel_config.use_sequence_parallel_moe
+    assert parallel_config.use_sequence_parallel_moe is expected
 
 
 def test_kimi_decoder_layer_keeps_moe_states_sequence_sharded(monkeypatch):
diff --git a/tests/models/multimodal/generation/test_common.py b/tests/models/multimodal/generation/test_common.py
index ecf1ad26d0c..5c431140b04 100644
--- a/tests/models/multimodal/generation/test_common.py
+++ b/tests/models/multimodal/generation/test_common.py
@@ -1242,7 +1242,7 @@ def test_custom_inputs_models(
         create_new_process_for_each_test=True,
     ),
 )
-@create_new_process_for_each_test()
+@create_new_process_for_each_test("spawn")
 def test_single_image_models_heavy(
     tmp_path: PosixPath,
     model_type: str,
diff --git a/tests/models/multimodal/processing/test_gemma4.py b/tests/models/multimodal/processing/test_gemma4.py
index a355501fdd8..f30afe47dde 100644
--- a/tests/models/multimodal/processing/test_gemma4.py
+++ b/tests/models/multimodal/processing/test_gemma4.py
@@ -7,6 +7,7 @@ import pytest
 import torch
 from PIL import Image as PILImage
 
+from vllm.exceptions import VLLMValidationError
 from vllm.model_executor.models.gemma4_mm import (
     Gemma4ForConditionalGeneration,
     Gemma4ImagePixelInputs,
@@ -222,7 +223,7 @@ def test_limit_mm_per_prompt(
     mm_data = {"image": images}
 
     # Expect ValueError when exceeding limit
-    with pytest.raises(ValueError, match="At most 1 image"):
+    with pytest.raises(VLLMValidationError, match="At most 1 image"):
         processor(
             prompt,
             mm_items=processor.info.parse_mm_data(mm_data),
diff --git a/tests/models/multimodal/processing/test_gemma4_unified.py b/tests/models/multimodal/processing/test_gemma4_unified.py
index 473ba729b85..67a81ddb7b4 100644
--- a/tests/models/multimodal/processing/test_gemma4_unified.py
+++ b/tests/models/multimodal/processing/test_gemma4_unified.py
@@ -7,6 +7,7 @@ import pytest
 import torch
 from PIL import Image as PILImage
 
+from vllm.exceptions import VLLMValidationError
 from vllm.model_executor.models.gemma4_mm import Gemma4ImagePixelInputs
 from vllm.multimodal import MULTIMODAL_REGISTRY
 from vllm.multimodal.inputs import MultiModalFieldConfig
@@ -197,7 +198,7 @@ def test_limit_mm_per_prompt(
 
     mm_data = {"image": images}
 
-    with pytest.raises(ValueError, match="At most 1 image"):
+    with pytest.raises(VLLMValidationError, match="At most 1 image"):
         processor(
             prompt,
             mm_items=processor.info.parse_mm_data(mm_data),
diff --git a/tests/models/multimodal/processing/test_glm4_1v.py b/tests/models/multimodal/processing/test_glm4_1v.py
index 8a777832826..e45e741c5b4 100644
--- a/tests/models/multimodal/processing/test_glm4_1v.py
+++ b/tests/models/multimodal/processing/test_glm4_1v.py
@@ -63,6 +63,31 @@ def test_encoder_cudagraph_uses_model_video_frame_limit():
     assert Glm4vForConditionalGeneration.get_max_frames_per_video(model) == 600
 
 
+@pytest.mark.parametrize(
+    ("temporal_patch_size", "expected_grid_t"),
+    [(2, 9), (4, 5), (8, 3)],
+)
+def test_vision_info_rounds_up_temporal_frames(
+    temporal_patch_size: int,
+    expected_grid_t: int,
+):
+    info = Mock(spec=Glm4vProcessingInfo)
+    vision_config = info.get_hf_config.return_value.vision_config
+    vision_config.patch_size = 14
+    vision_config.spatial_merge_size = 2
+    vision_config.temporal_patch_size = temporal_patch_size
+
+    _, num_vision_tokens = Glm4vProcessingInfo._get_vision_info(
+        info,
+        image_width=28,
+        image_height=28,
+        num_frames=17,
+        do_resize=False,
+    )
+
+    assert num_vision_tokens == expected_grid_t
+
+
 @pytest.mark.parametrize("model_id", ["zai-org/GLM-4.1V-9B-Thinking"])
 @pytest.mark.parametrize("expected_toks_per_frame", [299])
 @pytest.mark.parametrize(
diff --git a/tests/models/quantization/test_mxfp8.py b/tests/models/quantization/test_mxfp8.py
index 7c250d11576..c12a72a09c0 100644
--- a/tests/models/quantization/test_mxfp8.py
+++ b/tests/models/quantization/test_mxfp8.py
@@ -17,8 +17,10 @@ diverse prompts from ``tests/prompts/example.txt``.
 """
 
 import pytest
+import torch
 
 from tests.quantization.utils import is_quant_method_supported
+from vllm.platforms import current_platform
 
 from ..utils import check_logprobs_close
 
@@ -81,6 +83,170 @@ def test_mxfp8_logprobs(
         )
 
 
+@pytest.mark.skipif(
+    not is_quant_method_supported("mxfp8"),
+    reason="mxfp8 is not supported on this GPU type (requires sm_100+).",
+)
+@pytest.mark.skipif(
+    not current_platform.is_rocm(),
+    reason="AITER MXFP8 MoE backend is ROCm-only.",
+)
+@pytest.mark.quant_model
+def test_mxfp8_aiter_requires_swigluoai_activation(
+    monkeypatch: pytest.MonkeyPatch,
+) -> None:
+    from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+    from vllm.model_executor.layers.fused_moe.config import (
+        FusedMoEConfig,
+        FusedMoEParallelConfig,
+        RoutingMethodType,
+    )
+    from vllm.model_executor.layers.fused_moe.experts import aiter_mxfp8_moe
+    from vllm.model_executor.layers.fused_moe.oracle.mxfp8 import (
+        select_mxfp8_moe_backend,
+    )
+
+    monkeypatch.setattr(
+        aiter_mxfp8_moe.AiterMxfp8Experts,
+        "_supports_current_device",
+        staticmethod(lambda: True),
+    )
+    monkeypatch.setattr(
+        aiter_mxfp8_moe,
+        "is_aiter_mxfp8_moe_available",
+        lambda: True,
+    )
+
+    config = FusedMoEConfig(
+        num_experts=8,
+        experts_per_token=2,
+        hidden_dim=256,
+        intermediate_size=256,
+        num_local_experts=8,
+        num_logical_experts=8,
+        moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
+        activation=MoEActivation.SILU,
+        in_dtype=torch.bfloat16,
+        device="cuda",
+        routing_method=RoutingMethodType.Renormalize,
+        moe_backend="aiter",
+    )
+
+    with pytest.raises(ValueError, match="requires activation=swigluoai_uninterleave"):
+        select_mxfp8_moe_backend(config)
+
+
+@pytest.mark.skipif(
+    not is_quant_method_supported("mxfp8"),
+    reason="mxfp8 is not supported on this GPU type (requires sm_100+).",
+)
+@pytest.mark.skipif(
+    not current_platform.is_rocm(),
+    reason="AITER MXFP8 MoE backend is ROCm-only.",
+)
+@pytest.mark.quant_model
+def test_mxfp8_aiter_requires_swigluoai_params(
+    monkeypatch: pytest.MonkeyPatch,
+) -> None:
+    from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+    from vllm.model_executor.layers.fused_moe.config import (
+        FusedMoEConfig,
+        FusedMoEParallelConfig,
+        RoutingMethodType,
+    )
+    from vllm.model_executor.layers.fused_moe.experts import aiter_mxfp8_moe
+    from vllm.model_executor.layers.fused_moe.oracle.mxfp8 import (
+        select_mxfp8_moe_backend,
+    )
+
+    monkeypatch.setattr(
+        aiter_mxfp8_moe.AiterMxfp8Experts,
+        "_supports_current_device",
+        staticmethod(lambda: True),
+    )
+    monkeypatch.setattr(
+        aiter_mxfp8_moe,
+        "is_aiter_mxfp8_moe_available",
+        lambda: True,
+    )
+
+    config = FusedMoEConfig(
+        num_experts=8,
+        experts_per_token=2,
+        hidden_dim=256,
+        intermediate_size=256,
+        num_local_experts=8,
+        num_logical_experts=8,
+        moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
+        activation=MoEActivation.SWIGLUOAI_UNINTERLEAVE,
+        in_dtype=torch.bfloat16,
+        device="cuda",
+        routing_method=RoutingMethodType.Renormalize,
+        moe_backend="aiter",
+    )
+
+    with pytest.raises(ValueError, match="hardcodes SwiGLU-OAI"):
+        select_mxfp8_moe_backend(config)
+
+
+@pytest.mark.skipif(
+    not is_quant_method_supported("mxfp8"),
+    reason="mxfp8 is not supported on this GPU type (requires sm_100+).",
+)
+@pytest.mark.skipif(
+    not current_platform.is_rocm(),
+    reason="AITER MXFP8 MoE backend is ROCm-only.",
+)
+@pytest.mark.quant_model
+def test_mxfp8_aiter_accepts_swigluoai_params(
+    monkeypatch: pytest.MonkeyPatch,
+) -> None:
+    from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+    from vllm.model_executor.layers.fused_moe.config import (
+        FusedMoEConfig,
+        FusedMoEParallelConfig,
+        RoutingMethodType,
+    )
+    from vllm.model_executor.layers.fused_moe.experts import aiter_mxfp8_moe
+    from vllm.model_executor.layers.fused_moe.oracle.fp8 import Fp8MoeBackend
+    from vllm.model_executor.layers.fused_moe.oracle.mxfp8 import (
+        select_mxfp8_moe_backend,
+    )
+
+    monkeypatch.setattr(
+        aiter_mxfp8_moe.AiterMxfp8Experts,
+        "_supports_current_device",
+        staticmethod(lambda: True),
+    )
+    monkeypatch.setattr(
+        aiter_mxfp8_moe,
+        "is_aiter_mxfp8_moe_available",
+        lambda: True,
+    )
+
+    config = FusedMoEConfig(
+        num_experts=8,
+        experts_per_token=2,
+        hidden_dim=256,
+        intermediate_size=256,
+        num_local_experts=8,
+        num_logical_experts=8,
+        moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
+        activation=MoEActivation.SWIGLUOAI_UNINTERLEAVE,
+        in_dtype=torch.bfloat16,
+        device="cuda",
+        routing_method=RoutingMethodType.Renormalize,
+        moe_backend="aiter",
+        swiglu_alpha=aiter_mxfp8_moe._AITER_SWIGLU_ALPHA,
+        swiglu_beta=aiter_mxfp8_moe._AITER_SWIGLU_BETA,
+    )
+
+    backend, experts_cls = select_mxfp8_moe_backend(config)
+
+    assert backend == Fp8MoeBackend.AITER_MXFP8
+    assert experts_cls is aiter_mxfp8_moe.AiterMxfp8Experts
+
+
 @pytest.mark.skipif(
     not is_quant_method_supported("mxfp8"),
     reason="mxfp8 is not supported on this GPU type (requires sm_100+).",
diff --git a/tests/multimodal/media/test_unprocessable_entity_error.py b/tests/multimodal/media/test_unprocessable_entity_error.py
index 8be70383b8a..7cad4295575 100644
--- a/tests/multimodal/media/test_unprocessable_entity_error.py
+++ b/tests/multimodal/media/test_unprocessable_entity_error.py
@@ -14,7 +14,7 @@ import aiohttp
 import pytest
 
 from vllm.entrypoints.serve.utils.error_response import create_error_response
-from vllm.exceptions import VLLMUnprocessableEntityError
+from vllm.exceptions import VLLMClientError, VLLMUnprocessableEntityError
 from vllm.multimodal.media import MediaConnector
 
 
@@ -35,9 +35,9 @@ class TestVLLMUnprocessableEntityError:
         assert "parameter=image_url" in str(exc)
         assert "value=https://example.com/image.jpg" in str(exc)
 
-    def test_is_value_error_subclass(self):
+    def test_is_client_error_subclass(self):
         exc = VLLMUnprocessableEntityError("Test")
-        assert isinstance(exc, ValueError)
+        assert isinstance(exc, VLLMClientError)
 
 
 class TestMediaConnectorErrorHandling:
diff --git a/tests/multimodal/test_processing.py b/tests/multimodal/test_processing.py
index 66acdbe62ff..2153cd2c259 100644
--- a/tests/multimodal/test_processing.py
+++ b/tests/multimodal/test_processing.py
@@ -8,6 +8,7 @@ import numpy as np
 import pytest
 
 from vllm.config import ModelConfig
+from vllm.exceptions import VLLMValidationError
 from vllm.multimodal import MULTIMODAL_REGISTRY
 from vllm.multimodal.processing.context import InputProcessingContext
 from vllm.multimodal.processing.processor import (
@@ -931,7 +932,11 @@ def test_limit_mm_per_prompt_apply(model_id, num_images, limit, is_valid):
     else:
         mm_data = {"image": [image] * num_images}
 
-    exc_ctx = nullcontext() if is_valid else pytest.raises(ValueError, match="At most")
+    exc_ctx = (
+        nullcontext()
+        if is_valid
+        else pytest.raises(VLLMValidationError, match="At most")
+    )
 
     with exc_ctx:
         processor(
diff --git a/tests/quantization/test_auto_round.py b/tests/quantization/test_auto_round.py
index 732080fc967..59c6a1e326b 100644
--- a/tests/quantization/test_auto_round.py
+++ b/tests/quantization/test_auto_round.py
@@ -243,6 +243,27 @@ def test_inc_config_parser_parallel_lm_head_defaults_to_unquantized() -> None:
     assert layer_config.bits == 16
 
 
+def test_inc_config_parser_suffix_match_for_lm_head() -> None:
+    """Short extra_config key should match fully-qualified lm_head layer name."""
+    layer = object.__new__(ParallelLMHead)
+    config = make_config(
+        extra_config={
+            "lm_head": {
+                "bits": 4,
+                "group_size": 128,
+                "sym": True,
+            }
+        }
+    )
+
+    layer_config = config.config_parser.resolve(layer, "model.language_model.lm_head")
+
+    assert layer_config.quantized is True
+    assert layer_config.bits == 4
+    assert layer_config.group_size == 128
+    assert layer_config.sym is True
+
+
 def test_inc_config_parser_fused_moe_requires_consistent_configs() -> None:
     config = make_config(
         extra_config={
@@ -790,6 +811,34 @@ def test_inc_get_quant_method_linear_uses_resolved_scheme(monkeypatch) -> None:
     assert method is sentinel
 
 
+def test_inc_get_quant_method_lm_head_uses_suffix_match(monkeypatch) -> None:
+    """lm_head extra_config should apply to fully-qualified prefix."""
+    config = make_config(
+        extra_config={
+            "lm_head": {
+                "bits": 4,
+                "group_size": 128,
+                "sym": True,
+            }
+        }
+    )
+    layer = object.__new__(ParallelLMHead)
+    sentinel = object()
+
+    class DummyScheme:
+        def get_linear_method(self, _config, _layer, _prefix, _layer_config):
+            return sentinel
+
+    monkeypatch.setattr(
+        "vllm.model_executor.layers.quantization.inc.schemes.factory.resolve_scheme",
+        lambda _layer_config: DummyScheme(),
+    )
+
+    method = config.get_quant_method(layer, "model.language_model.lm_head")
+
+    assert method is sentinel
+
+
 def test_inc_get_quant_method_moe_uses_resolved_scheme(monkeypatch) -> None:
     config = make_config()
     layer = object.__new__(RoutedExperts)
diff --git a/tests/quantization/test_compressed_tensors.py b/tests/quantization/test_compressed_tensors.py
index 626717cd4a3..8c509888700 100644
--- a/tests/quantization/test_compressed_tensors.py
+++ b/tests/quantization/test_compressed_tensors.py
@@ -473,6 +473,11 @@ def test_compressed_tensors_w4a8_fp8(vllm_runner, args):
             "Flat is better than nested.\nSparse is better than dense.",
             150.0,
         ),
+        (
+            "nm-testing/Llama-3.2-1B-Instruct-quipv16-nvfp4",
+            "Flat is better than nested.\nSparse is better than dense.",
+            150.0,
+        ),
     ],
 )
 def test_compressed_tensors_transforms_perplexity(
diff --git a/tests/quantization/test_gfx950_moe.py b/tests/quantization/test_gfx950_moe.py
index 0efcc8a3c62..c8d34bb0ab5 100644
--- a/tests/quantization/test_gfx950_moe.py
+++ b/tests/quantization/test_gfx950_moe.py
@@ -79,21 +79,6 @@ def test_w4a4_dispatches_to_aiter(mxfp4_oracle_config):
     assert experts_cls is not None
 
 
-@pytest.mark.skipif(not ROCM_GFX950, reason="Requires GFX950 (mi355x)")
-@pytest.mark.skipif(
-    ROCM_AITER_AVAILABLE,
-    reason="Test requires AITER disabled (unset VLLM_ROCM_USE_AITER)",
-)
-def test_w4a4_falls_back_to_triton_unfused_without_aiter(mxfp4_oracle_config):
-    """Without AITER and no --moe-backend, ROCm falls back to TRITON_UNFUSED."""
-    config = _make_w4a4_moe_config()
-    backend, experts_cls = select_mxfp4_moe_backend(
-        config, activation_key=kMxfp4Dynamic
-    )
-    assert backend == Mxfp4MoeBackend.TRITON_UNFUSED
-    assert experts_cls is not None
-
-
 @pytest.mark.skipif(not ROCM_GFX950, reason="Requires GFX950 (mi355x)")
 def test_w4a4_dispatches_to_emulation_with_moe_backend(mxfp4_oracle_config):
     """With --moe-backend emulation, W4A4 selects EMULATION."""
diff --git a/tests/renderers/test_chat_utils_prompt_embeds.py b/tests/renderers/test_chat_utils_prompt_embeds.py
index 2238c41f498..3f08194b9be 100644
--- a/tests/renderers/test_chat_utils_prompt_embeds.py
+++ b/tests/renderers/test_chat_utils_prompt_embeds.py
@@ -26,6 +26,7 @@ from vllm.entrypoints.chat_utils import (
     parse_chat_messages,
     parse_chat_messages_async,
 )
+from vllm.exceptions import VLLMValidationError
 from vllm.renderers.hf import (
     _PROMPT_EMBEDS_PLACEHOLDER_SPAN_MISMATCH_ERROR,
     _build_mixed_prompt_embeds,
@@ -264,7 +265,7 @@ def test_parse_chat_messages_requires_flag():
             "content": [{"type": "prompt_embeds", "data": b64}],
         }
     ]
-    with pytest.raises(ValueError, match=_ENABLE_PROMPT_EMBEDS_ERROR):
+    with pytest.raises(VLLMValidationError, match=_ENABLE_PROMPT_EMBEDS_ERROR):
         parse_chat_messages(
             messages,
             mc,
@@ -283,7 +284,7 @@ def test_parse_chat_messages_rejects_missing_data():
             "content": [{"type": "prompt_embeds"}],  # no `data`
         }
     ]
-    with pytest.raises(ValueError, match=_PROMPT_EMBEDS_MISSING_DATA_ERROR):
+    with pytest.raises(VLLMValidationError, match=_PROMPT_EMBEDS_MISSING_DATA_ERROR):
         parse_chat_messages(
             messages,
             mc,
diff --git a/tests/renderers/test_completions.py b/tests/renderers/test_completions.py
index d184eb8621c..1849eeac7dc 100644
--- a/tests/renderers/test_completions.py
+++ b/tests/renderers/test_completions.py
@@ -11,6 +11,7 @@ import pytest
 import torch
 
 from vllm.config import ModelConfig
+from vllm.exceptions import VLLMValidationError
 from vllm.inputs import SingletonPrompt
 from vllm.renderers import TokenizeParams
 from vllm.renderers.hf import HfRenderer
@@ -286,7 +287,7 @@ class TestRenderPrompt:
         )
 
         with pytest.raises(
-            ValueError,
+            VLLMValidationError,
             match="maximum context length is",
         ):
             renderer.tokenize_prompts(
@@ -307,7 +308,7 @@ class TestRenderPrompt:
         )
 
         with pytest.raises(
-            ValueError,
+            VLLMValidationError,
             match="maximum context length is",
         ):
             renderer.tokenize_prompts(
@@ -328,7 +329,7 @@ class TestRenderPrompt:
         )
 
         with pytest.raises(
-            ValueError,
+            VLLMValidationError,
             match="maximum context length is",
         ):
             renderer.tokenize_prompts(
diff --git a/tests/samplers/test_non_finite_params.py b/tests/samplers/test_non_finite_params.py
index 57fe90f314c..f982953d608 100644
--- a/tests/samplers/test_non_finite_params.py
+++ b/tests/samplers/test_non_finite_params.py
@@ -42,7 +42,7 @@ class TestNonFiniteRepetitionPenalty:
         ids=["nan", "inf", "-inf", "math.nan", "math.inf"],
     )
     def test_non_finite_repetition_penalty_rejected(self, value: float):
-        with pytest.raises(ValueError, match="repetition_penalty"):
+        with pytest.raises(VLLMValidationError, match="repetition_penalty"):
             SamplingParams(repetition_penalty=value)
 
     def test_finite_repetition_penalty_accepted(self):
diff --git a/tests/test_cmake_utils.py b/tests/test_cmake_utils.py
index 227ec231eb2..d0673bc462e 100644
--- a/tests/test_cmake_utils.py
+++ b/tests/test_cmake_utils.py
@@ -21,3 +21,28 @@ endif()
     )
 
     subprocess.run(["cmake", "-P", script], check=True)
+
+
+def test_extract_archs_prefers_sass_target_over_corrupted_virtual_arch(
+    tmp_path: Path,
+):
+    """torch's autodetection can emit a bogus arch=compute_* half (e.g.
+    capability 12.1 corrupted to arch=compute_20,code=sm_121); the SASS
+    target must win, while PTX-only entries keep the virtual arch."""
+    repo_root = Path(__file__).parents[1]
+    script = tmp_path / "test_extract_archs.cmake"
+    script.write_text(
+        f"""
+cmake_minimum_required(VERSION 3.26)
+include("{repo_root / "cmake" / "utils.cmake"}")
+extract_unique_cuda_archs_ascending(actual
+  "-gencode arch=compute_20,code=sm_121;\
+-gencode arch=compute_80,code=sm_80;\
+-gencode arch=compute_80,code=compute_80")
+if(NOT "${{actual}}" STREQUAL "8.0;12.1")
+  message(FATAL_ERROR "Expected '8.0;12.1', got '${{actual}}'")
+endif()
+"""
+    )
+
+    subprocess.run(["cmake", "-P", script], check=True)
diff --git a/tests/test_envs.py b/tests/test_envs.py
index 56c04dd6f2e..5e0363e33a1 100644
--- a/tests/test_envs.py
+++ b/tests/test_envs.py
@@ -15,6 +15,7 @@ from vllm.envs import (
     env_with_choices,
     environment_variables,
 )
+from vllm.exceptions import VLLMValidationError
 
 
 def test_getattr_without_cache(monkeypatch: pytest.MonkeyPatch):
@@ -145,6 +146,15 @@ def test_precompiled_install_flags_are_orthogonal() -> None:
         assert environment_variables["VLLM_USE_PRECOMPILED_RUST"]() is True
 
 
+def test_rust_bench_auto_path_missing_fails_fast() -> None:
+    with (
+        patch.dict(os.environ, {"VLLM_USE_RUST_BENCH": "1"}, clear=True),
+        patch("vllm.envs.os.path.isfile", return_value=False),
+        pytest.raises(FileNotFoundError, match="vllm-rs binary was not found"),
+    ):
+        environment_variables["VLLM_RUST_FRONTEND_PATH"]()
+
+
 class TestEnvWithChoices:
     """Test cases for env_with_choices function."""
 
@@ -538,7 +548,7 @@ class TestVllmMaxNSequences:
         max_n = envs.VLLM_MAX_N_SEQUENCES
         SamplingParams(n=max_n)
 
-        with pytest.raises(ValueError, match="n must be at most"):
+        with pytest.raises(VLLMValidationError, match="n must be at most"):
             SamplingParams(n=max_n + 1)
 
     def test_sampling_params_respects_custom_limit(
@@ -554,5 +564,5 @@ class TestVllmMaxNSequences:
 
         SamplingParams(n=128)
 
-        with pytest.raises(ValueError, match="n must be at most 128"):
+        with pytest.raises(VLLMValidationError, match="n must be at most 128"):
             SamplingParams(n=129)
diff --git a/tests/test_pooling_params.py b/tests/test_pooling_params.py
index 17d04078b4e..f34270d0da5 100644
--- a/tests/test_pooling_params.py
+++ b/tests/test_pooling_params.py
@@ -52,13 +52,19 @@ class MockModelConfig:
 def test_removed_pooling_parameters(parameter: str, value: Any, message: str):
     data = {"input": "hello", parameter: value}
     for request_type in (EmbeddingRequest, ClassificationRequest, PoolingRequest):
-        with pytest.raises(ValidationError, match=message) as exc_info:
+        with pytest.raises(VLLMValidationError, match=message):
             TypeAdapter(request_type).validate_python(data)
-        assert len(exc_info.value.errors()) == 1
 
-    with pytest.raises(ValidationError, match=message) as exc_info:
-        TypeAdapter(PoolerConfig).validate_python({parameter: value})
-    assert len(exc_info.value.errors()) == 1
+    # PoolerConfig still raises bare ValueError for `normalize`
+    # (wrapped to ValidationError by Pydantic), but `check_removed_pooling_task`
+    # raises VLLMValidationError for removed tasks.
+    if parameter == "normalize":
+        with pytest.raises(ValidationError, match=message) as exc_info:
+            TypeAdapter(PoolerConfig).validate_python({parameter: value})
+        assert len(exc_info.value.errors()) == 1
+    else:
+        with pytest.raises(VLLMValidationError, match=message):
+            TypeAdapter(PoolerConfig).validate_python({parameter: value})
 
     if parameter == "task":
         with pytest.raises(VLLMValidationError, match=message):
@@ -80,7 +86,7 @@ def test_embed():
 
     invalid_parameters = classify_parameters + step_pooling_parameters
     for p in set(invalid_parameters) - set(embed_parameters):
-        with pytest.raises(ValueError):
+        with pytest.raises(VLLMValidationError):
             pooling_params = PoolingParams(task=task, **{p: True})
             pooling_params.verify(model_config)
 
@@ -100,7 +106,7 @@ def test_embed_dimensions(model_info: EmbedModelInfo):
     pooling_params = PoolingParams(task=task, dimensions=None)
     pooling_params.verify(model_config)
 
-    with pytest.raises(ValueError):
+    with pytest.raises(VLLMValidationError):
         pooling_params = PoolingParams(task=task, dimensions=1)
         pooling_params.verify(model_config)
 
@@ -131,7 +137,7 @@ def test_embed_dimensions_matryoshka_without_list_upper_bound():
 
     PoolingParams(task=task, dimensions=16).verify(model_config)
 
-    with pytest.raises(ValueError):
+    with pytest.raises(VLLMValidationError):
         PoolingParams(task=task, dimensions=64).verify(model_config)
 
 
@@ -150,7 +156,7 @@ def test_classify(task):
 
     invalid_parameters = embed_parameters + step_pooling_parameters
     for p in set(invalid_parameters) - set(classify_parameters):
-        with pytest.raises(ValueError):
+        with pytest.raises(VLLMValidationError):
             pooling_params = PoolingParams(task=task, **{p: True})
             pooling_params.verify(model_config)
 
@@ -176,7 +182,7 @@ def test_token_embed(pooling_type: str):
         invalid_parameters = classify_parameters + step_pooling_parameters
 
     for p in set(invalid_parameters) - set(embed_parameters):
-        with pytest.raises(ValueError):
+        with pytest.raises(VLLMValidationError):
             pooling_params = PoolingParams(task=task, **{p: True})
             pooling_params.verify(model_config)
 
@@ -202,6 +208,6 @@ def test_token_classify(pooling_type: str):
         invalid_parameters = embed_parameters + step_pooling_parameters
 
     for p in set(invalid_parameters) - set(classify_parameters):
-        with pytest.raises(ValueError):
+        with pytest.raises(VLLMValidationError):
             pooling_params = PoolingParams(task=task, **{p: True})
             pooling_params.verify(model_config)
diff --git a/tests/test_sampling_params.py b/tests/test_sampling_params.py
index e5d811fbb13..65ab0738c96 100644
--- a/tests/test_sampling_params.py
+++ b/tests/test_sampling_params.py
@@ -5,6 +5,7 @@ from dataclasses import dataclass
 import pytest
 
 from vllm import SamplingParams
+from vllm.exceptions import VLLMValidationError
 
 
 @dataclass
@@ -32,7 +33,7 @@ class MockModelConfig:
 )
 def test_diffusion_rejects_unsupported_params(kwargs: dict):
     params = SamplingParams(**kwargs)
-    with pytest.raises(ValueError, match="not yet supported with diffusion"):
+    with pytest.raises(VLLMValidationError, match="not yet supported with diffusion"):
         params.verify(MockModelConfig(is_diffusion=True), None, None, None)
 
 
diff --git a/tests/tool_use/test_chat_completion_request_validations.py b/tests/tool_use/test_chat_completion_request_validations.py
index 7adf4beb9d8..1def2acc56c 100644
--- a/tests/tool_use/test_chat_completion_request_validations.py
+++ b/tests/tool_use/test_chat_completion_request_validations.py
@@ -4,6 +4,7 @@
 import pytest
 
 from vllm.entrypoints.openai.chat_completion.protocol import ChatCompletionRequest
+from vllm.exceptions import VLLMValidationError
 
 
 def test_chat_completion_request_with_no_tools():
@@ -27,7 +28,7 @@ def test_chat_completion_request_with_no_tools():
     assert request.tool_choice == "none"
 
     # tools key present but empty -- should be rejected
-    with pytest.raises(ValueError, match="must not be an empty array"):
+    with pytest.raises(VLLMValidationError, match="must not be an empty array"):
         ChatCompletionRequest.model_validate(
             {
                 "messages": [{"role": "user", "content": "Hello"}],
@@ -40,7 +41,7 @@ def test_chat_completion_request_with_no_tools():
 @pytest.mark.parametrize("tool_choice", ["auto", "required"])
 def test_chat_completion_request_with_tool_choice_but_no_tools(tool_choice):
     with pytest.raises(
-        ValueError, match="When using `tool_choice`, `tools` must be set."
+        VLLMValidationError, match="When using `tool_choice`, `tools` must be set."
     ):
         ChatCompletionRequest.model_validate(
             {
@@ -51,7 +52,7 @@ def test_chat_completion_request_with_tool_choice_but_no_tools(tool_choice):
         )
 
     with pytest.raises(
-        ValueError, match="When using `tool_choice`, `tools` must be set."
+        VLLMValidationError, match="When using `tool_choice`, `tools` must be set."
     ):
         ChatCompletionRequest.model_validate(
             {
@@ -134,7 +135,7 @@ SAMPLE_TOOL = {
 def test_structured_outputs_with_named_tool_choice_rejected():
     """structured_outputs cannot be combined with a named tool_choice."""
     with pytest.raises(
-        ValueError,
+        VLLMValidationError,
         match="structured outputs or tools, not both",
     ):
         ChatCompletionRequest.model_validate(
@@ -168,7 +169,7 @@ def test_structured_outputs_with_auto_tool_choice_allowed():
 def test_multiple_structured_outputs_rejected():
     """Only one kind of structured output constraint is allowed."""
     with pytest.raises(
-        ValueError,
+        VLLMValidationError,
         match="You can only use one kind of constraints",
     ):
         ChatCompletionRequest.model_validate(
diff --git a/tests/tool_use/test_responses_request_validations.py b/tests/tool_use/test_responses_request_validations.py
index 59b156b76a1..b50482960c8 100644
--- a/tests/tool_use/test_responses_request_validations.py
+++ b/tests/tool_use/test_responses_request_validations.py
@@ -2,12 +2,12 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import pytest
-from pydantic import ValidationError
 
 from vllm.entrypoints.openai.responses.protocol import (
     ResponsesRequest,
     ResponsesResponse,
 )
+from vllm.exceptions import VLLMValidationError
 
 SAMPLE_TOOL = {
     "type": "function",
@@ -58,13 +58,13 @@ def test_responses_request_required_without_tools(tools):
     if tools is not None:
         kwargs["tools"] = tools
     with pytest.raises(
-        ValidationError, match="Tool choice 'required' must be specified"
+        VLLMValidationError, match="Tool choice 'required' must be specified"
     ):
         ResponsesRequest.model_validate(kwargs)
 
 
 def test_responses_request_named_tool_choice_without_tools():
-    with pytest.raises(ValidationError, match="not found in 'tools' parameter"):
+    with pytest.raises(VLLMValidationError, match="not found in 'tools' parameter"):
         ResponsesRequest.model_validate(
             {
                 "input": "Hello",
@@ -107,7 +107,7 @@ def test_responses_request_named_tool_choice_matching():
 
 
 def test_responses_request_named_tool_choice_not_matching():
-    with pytest.raises(ValidationError, match="not found in 'tools' parameter"):
+    with pytest.raises(VLLMValidationError, match="not found in 'tools' parameter"):
         ResponsesRequest.model_validate(
             {
                 "input": "Hello",
@@ -164,7 +164,7 @@ def test_responses_request_empty_tools_tool_choice_auto():
     ],
 )
 def test_responses_request_named_tool_choice_missing_name(tool_choice):
-    with pytest.raises(ValidationError, match="not found in 'tools' parameter"):
+    with pytest.raises(VLLMValidationError, match="not found in 'tools' parameter"):
         ResponsesRequest.model_validate(
             {
                 "input": "Hello",
@@ -176,7 +176,7 @@ def test_responses_request_named_tool_choice_missing_name(tool_choice):
 
 
 def test_responses_request_empty_tools_named_tool_choice():
-    with pytest.raises(ValidationError, match="not found in 'tools' parameter"):
+    with pytest.raises(VLLMValidationError, match="not found in 'tools' parameter"):
         ResponsesRequest.model_validate(
             {
                 "input": "Hello",
diff --git a/tests/v1/attention/test_mla_prefill_selector.py b/tests/v1/attention/test_mla_prefill_selector.py
index f5985e7bc8e..e6d9f939ea5 100644
--- a/tests/v1/attention/test_mla_prefill_selector.py
+++ b/tests/v1/attention/test_mla_prefill_selector.py
@@ -8,6 +8,7 @@ import pytest
 import torch
 
 from vllm.config import AttentionConfig, ModelConfig, VllmConfig
+from vllm.platforms import current_platform
 from vllm.platforms.interface import DeviceCapability
 from vllm.v1.attention.backends.mla.prefill.base import MLADimensions
 from vllm.v1.attention.backends.mla.prefill.registry import MLAPrefillBackendEnum
@@ -287,6 +288,11 @@ class TestBackendValidation:
             assert invalid_reasons == []
 
 
+@pytest.mark.skipif(
+    not current_platform.is_cuda_alike(),
+    reason="Imports vllm.platforms.rocm, whose module init requires a CUDA or "
+    "ROCm torch build; not importable on XPU/CPU/TPU.",
+)
 class TestROCmAiterFAPrefillSelection:
     """Tests for the ROCm AITER FlashAttention MLA prefill backend."""
 
diff --git a/tests/v1/determinism/test_batch_invariance.py b/tests/v1/determinism/test_batch_invariance.py
index b2706ed89b7..37fd5cba6a5 100644
--- a/tests/v1/determinism/test_batch_invariance.py
+++ b/tests/v1/determinism/test_batch_invariance.py
@@ -27,8 +27,10 @@ from vllm.platforms import current_platform
     "backend",
     BACKENDS,
 )
+@pytest.mark.parametrize("rms_norm_impl", ["default", "vllm_c"])
 def test_v1_generation_is_deterministic_across_batch_sizes_with_needle(
     backend,
+    rms_norm_impl,
 ):
     """
     Ensures that the same request (the 'needle' prompt) yields identical output
@@ -60,6 +62,16 @@ def test_v1_generation_is_deterministic_across_batch_sizes_with_needle(
     random.seed(seed)
 
     attention_config = {"backend": backend}
+    # Force the C++ RMSNorm implementation so we actually exercise the
+    # num_tokens-dependent block-size branches.
+    kernel_config = None
+    if rms_norm_impl == "vllm_c":
+        kernel_config = {
+            "ir_op_priority": {
+                "rms_norm": ["vllm_c"],
+                "fused_add_rms_norm": ["vllm_c"],
+            }
+        }
     # Allow overrides from environment (useful for CI tuning)
     # "facebook/opt-125m" is too small, doesn't reliably test determinism
     model = TEST_MODEL
@@ -96,6 +108,7 @@ def test_v1_generation_is_deterministic_across_batch_sizes_with_needle(
             gpu_memory_utilization=gpu_mem_util,
             max_model_len=max_model_len,
             attention_config=attention_config,
+            kernel_config=kernel_config,
         )
 
         # Baseline generation for the needle prompt alone.
@@ -923,11 +936,15 @@ def LLM_with_max_seqs(
     gpu_memory_utilization: float,
     max_model_len: int,
     attention_config: dict | None = None,
+    kernel_config: dict | None = None,
 ) -> LLM:
     """
     Helper to construct an LLM with a specific max_num_seqs (batch-size limit)
     using the high-level v1 LLM API, while constraining memory usage.
     """
+    extra_kwargs: dict = {}
+    if kernel_config is not None:
+        extra_kwargs["kernel_config"] = kernel_config
     return LLM(
         model=model,
         max_num_seqs=max_num_seqs,
@@ -939,4 +956,5 @@ def LLM_with_max_seqs(
         attention_config=attention_config,
         # Enable for MOE models
         # enable_expert_parallel=True,
+        **extra_kwargs,
     )
diff --git a/tests/v1/determinism/test_rms_norm_batch_invariant.py b/tests/v1/determinism/test_rms_norm_batch_invariant.py
index dfd08351277..232a43b1f98 100644
--- a/tests/v1/determinism/test_rms_norm_batch_invariant.py
+++ b/tests/v1/determinism/test_rms_norm_batch_invariant.py
@@ -28,16 +28,18 @@ def _rms_norm_reference(
 
 
 @skip_if_not_cuda
-@pytest.mark.parametrize("batch_size", [1, 4, 16, 64])
+@pytest.mark.parametrize("batch_size", [1, 4, 64, 300])
 @pytest.mark.parametrize("hidden_size", [512, 2048, 4096, 8192])
 @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
 @pytest.mark.parametrize("eps", [1e-6, 1e-5])
+@pytest.mark.parametrize("seed", list(range(4)))
 def test_rms_norm_batch_invariant_vs_reference(
     default_vllm_config,
     batch_size: int,
     hidden_size: int,
     dtype: torch.dtype,
     eps: float,
+    seed: int,
 ):
     """
     Compare batch-invariant Triton RMS norm against a PyTorch reference.
@@ -48,7 +50,7 @@ def test_rms_norm_batch_invariant_vs_reference(
     device = torch.device(DEVICE_TYPE)
 
     # Create test input and weight
-    torch.manual_seed(42)
+    torch.manual_seed(seed)
     input_tensor = torch.randn(batch_size, hidden_size, dtype=dtype, device=device)
     weight = torch.randn(hidden_size, dtype=dtype, device=device)
 
@@ -71,7 +73,7 @@ def test_rms_norm_batch_invariant_vs_reference(
         atol=atol,
         msg=f"RMS norm mismatch for batch_size={batch_size}, "
         f"hidden_size={hidden_size}, "
-        f"dtype={dtype}, eps={eps}",
+        f"dtype={dtype}, eps={eps}, seed={seed}",
     )
 
 
@@ -79,17 +81,21 @@ def test_rms_norm_batch_invariant_vs_reference(
 @pytest.mark.parametrize("hidden_size", [512, 4096])
 @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
 @pytest.mark.parametrize("eps", [1e-6])
+@pytest.mark.parametrize("n_extra", [3, 299])
+@pytest.mark.parametrize("seed", list(range(16)))
 def test_fused_add_rms_norm_batch_invariant_residual_path(
     hidden_size: int,
     dtype: torch.dtype,
     eps: float,
+    n_extra: int,
+    seed: int,
 ):
     """
     Test the batch-invariant fused residual-add + RMSNorm helper directly.
     """
     device = torch.device(DEVICE_TYPE)
 
-    torch.manual_seed(42)
+    torch.manual_seed(seed)
     x_single = torch.randn(1, hidden_size, dtype=dtype, device=device)
     residual_single = torch.randn(1, hidden_size, dtype=dtype, device=device)
     weight = torch.randn(hidden_size, dtype=dtype, device=device)
@@ -97,14 +103,14 @@ def test_fused_add_rms_norm_batch_invariant_residual_path(
     x_batch = torch.cat(
         [
             x_single,
-            torch.randn(3, hidden_size, dtype=dtype, device=device),
+            torch.randn(n_extra, hidden_size, dtype=dtype, device=device),
         ],
         dim=0,
     )
     residual_batch = torch.cat(
         [
             residual_single,
-            torch.randn(3, hidden_size, dtype=dtype, device=device),
+            torch.randn(n_extra, hidden_size, dtype=dtype, device=device),
         ],
         dim=0,
     )
@@ -168,6 +174,138 @@ def test_fused_add_rms_norm_batch_invariant_residual_path(
     )
 
 
+FP8_DTYPE = current_platform.fp8_dtype()
+
+# The large launch (num_tokens=300 >= 256) drops an un-pinned kernel to block
+# 256, while the small launch (255 rows) stays under the threshold and keeps the
+# larger block (1024, or 512 for per-block quant). Under the pin the two launches
+# use the same block, so the shared first 255 rows must match bit-for-bit; 255 is
+# the most rows a single small launch can hold (< 256, and <= 256 for per-block).
+_LARGE_TOKENS = 300
+_SMALL_TOKENS = 255
+
+
+def _assert_rows_bit_identical(small, large, msg):
+    if small.dtype == FP8_DTYPE:
+        assert torch.equal(small.view(torch.uint8), large.view(torch.uint8)), msg
+    else:
+        torch.testing.assert_close(small, large, rtol=0.0, atol=0.0, msg=msg)
+
+
+@skip_if_not_cuda
+@pytest.mark.parametrize("hidden_size", [512, 4096])
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
+@pytest.mark.parametrize("seed", list(range(4)))
+def test_rms_norm_batch_invariant_nonresidual_kernel(
+    hidden_size: int, dtype: torch.dtype, seed: int
+):
+    """C++ ``rms_norm`` (no residual) must be batch invariant across the block
+    threshold. Reached in compiled mode with ``ir_op_priority.rms_norm=["vllm_c"]``
+    (default priority is ``native``/inductor codegen when compiling).
+    """
+    import vllm._custom_ops as ops
+
+    device = torch.device(DEVICE_TYPE)
+    torch.manual_seed(seed)
+    rows = torch.randn(_LARGE_TOKENS, hidden_size, dtype=dtype, device=device)
+    weight = torch.randn(hidden_size, dtype=dtype, device=device)
+
+    def rms_norm(x):
+        out = torch.empty_like(x)
+        ops.rms_norm(out, x, weight, 1e-6)
+        return out
+
+    large = rms_norm(rows.clone())
+    small = rms_norm(rows[:_SMALL_TOKENS].clone())
+    _assert_rows_bit_identical(
+        small,
+        large[:_SMALL_TOKENS],
+        "rms_norm output depends on num_tokens (block size)",
+    )
+
+
+@skip_if_not_cuda
+@pytest.mark.parametrize("hidden_size", [512, 4096])
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
+@pytest.mark.parametrize("seed", list(range(4)))
+@pytest.mark.parametrize("add_residual", [False, True])
+def test_rms_norm_static_fp8_quant_batch_invariant(
+    hidden_size: int, dtype: torch.dtype, seed: int, add_residual: bool
+):
+    """C++ static per-tensor fp8-quant RMSNorm must be batch invariant across
+    the block threshold. Covers ``rms_norm_static_fp8_quant`` and, with
+    ``add_residual``, ``fused_add_rms_norm_static_fp8_quant`` (the compiled fp8
+    path where ``RMSNormQuantFusionPass`` rewrites norm + quant into them).
+    """
+    device = torch.device(DEVICE_TYPE)
+    torch.manual_seed(seed)
+    rows = torch.randn(_LARGE_TOKENS, hidden_size, dtype=dtype, device=device)
+    residual = (
+        torch.randn(_LARGE_TOKENS, hidden_size, dtype=dtype, device=device)
+        if add_residual
+        else None
+    )
+    weight = torch.randn(hidden_size, dtype=dtype, device=device)
+    quant_scale = torch.tensor(1.0, dtype=torch.float32, device=device)
+
+    def quant(x, res):
+        out = torch.empty_like(x, dtype=FP8_DTYPE)
+        if add_residual:
+            torch.ops._C.fused_add_rms_norm_static_fp8_quant(
+                out, x, res, weight, quant_scale, 1e-6
+            )
+        else:
+            torch.ops._C.rms_norm_static_fp8_quant(out, x, weight, quant_scale, 1e-6)
+        return out
+
+    large = quant(rows.clone(), residual.clone() if residual is not None else None)
+    small = quant(
+        rows[:_SMALL_TOKENS].clone(),
+        residual[:_SMALL_TOKENS].clone() if residual is not None else None,
+    )
+    _assert_rows_bit_identical(
+        small,
+        large[:_SMALL_TOKENS],
+        "static-fp8-quant RMSNorm output depends on num_tokens (block size)",
+    )
+
+
+@skip_if_not_cuda
+@pytest.mark.parametrize("hidden_size", [512, 4096])
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
+@pytest.mark.parametrize("seed", list(range(4)))
+def test_rms_norm_per_block_quant_batch_invariant(
+    hidden_size: int, dtype: torch.dtype, seed: int
+):
+    """C++ ``rms_norm_per_block_quant`` must be batch invariant across the
+    block threshold (compiled fp8 block-quant path; block pinned to 512)."""
+    import vllm._custom_ops as ops
+
+    device = torch.device(DEVICE_TYPE)
+    torch.manual_seed(seed)
+    rows = torch.randn(_LARGE_TOKENS, hidden_size, dtype=dtype, device=device)
+    weight = torch.randn(hidden_size, dtype=dtype, device=device)
+    group_size = [1, 128]
+
+    def per_block_quant(x):
+        return ops.rms_norm_per_block_quant(x, weight, 1e-6, FP8_DTYPE, group_size)
+
+    out_large, scale_large = per_block_quant(rows.clone())
+    out_small, scale_small = per_block_quant(rows[:_SMALL_TOKENS].clone())
+    _assert_rows_bit_identical(
+        out_small,
+        out_large[:_SMALL_TOKENS],
+        "rms_norm_per_block_quant output depends on num_tokens (block size)",
+    )
+    torch.testing.assert_close(
+        scale_small,
+        scale_large[:_SMALL_TOKENS],
+        rtol=0.0,
+        atol=0.0,
+        msg="rms_norm_per_block_quant scales depend on num_tokens (block size)",
+    )
+
+
 @skip_if_not_cuda
 @pytest.mark.parametrize("batch_size", [1, 16, 128])
 @pytest.mark.parametrize("seq_len", [1, 32, 512])
diff --git a/tests/v1/e2e/general/test_context_length.py b/tests/v1/e2e/general/test_context_length.py
index cd0aff79de8..955e5c5bf02 100644
--- a/tests/v1/e2e/general/test_context_length.py
+++ b/tests/v1/e2e/general/test_context_length.py
@@ -59,7 +59,7 @@ def test_decoder_max_context_length_validation(
                 "Make sure that `max_model_len` is no smaller than the number of "
                 "text tokens (prompt + requested output tokens)."
             )
-            with pytest.raises(ValueError) as excinfo:
+            with pytest.raises(VLLMValidationError) as excinfo:
                 vllm_model.generate_greedy(prompt_ids, max_tokens)
             assert expected_msg in str(excinfo.value)
 
diff --git a/tests/v1/e2e/general/test_min_tokens.py b/tests/v1/e2e/general/test_min_tokens.py
index bb041cd3862..c5b6341fb98 100644
--- a/tests/v1/e2e/general/test_min_tokens.py
+++ b/tests/v1/e2e/general/test_min_tokens.py
@@ -16,6 +16,7 @@ Covers:
 import pytest
 
 from vllm import LLM, SamplingParams
+from vllm.exceptions import VLLMValidationError
 from vllm.outputs import RequestOutput
 
 # Test configuration
@@ -479,13 +480,13 @@ def test_min_tokens_validation():
 
     # Invalid cases
     with pytest.raises(
-        ValueError,
+        VLLMValidationError,
         match="min_tokens must be greater than or equal to 0",
     ):
         SamplingParams(min_tokens=-1, max_tokens=10)
 
     with pytest.raises(
-        ValueError,
+        VLLMValidationError,
         match="min_tokens must be less than or equal to max_tokens",
     ):
         SamplingParams(min_tokens=15, max_tokens=10)
diff --git a/tests/v1/e2e/general/test_streaming_input.py b/tests/v1/e2e/general/test_streaming_input.py
index 01c5fe6f8eb..1954ce6a7dc 100644
--- a/tests/v1/e2e/general/test_streaming_input.py
+++ b/tests/v1/e2e/general/test_streaming_input.py
@@ -20,6 +20,7 @@ import pytest_asyncio
 
 from vllm import SamplingParams
 from vllm.engine.protocol import StreamingInput
+from vllm.exceptions import VLLMValidationError
 from vllm.outputs import RequestOutput
 from vllm.platforms import current_platform
 from vllm.sampling_params import RequestOutputKind
@@ -571,13 +572,17 @@ async def test_streaming_input_validation_errors(engine: AsyncLLM):
         yield StreamingInput(prompt="test")
 
     # Test n > 1 is rejected
-    with pytest.raises(ValueError, match="Input streaming not currently supported"):
+    with pytest.raises(
+        VLLMValidationError, match="Input streaming not currently supported"
+    ):
         params_n2 = SamplingParams(max_tokens=10, n=2)
         async for _ in engine.generate(dummy_generator(), params_n2, "test_n2"):
             pass
 
     # Test FINAL_ONLY is rejected
-    with pytest.raises(ValueError, match="Input streaming not currently supported"):
+    with pytest.raises(
+        VLLMValidationError, match="Input streaming not currently supported"
+    ):
         params_final = SamplingParams(
             max_tokens=10, output_kind=RequestOutputKind.FINAL_ONLY
         )
@@ -585,7 +590,9 @@ async def test_streaming_input_validation_errors(engine: AsyncLLM):
             pass
 
     # Test stop strings are rejected
-    with pytest.raises(ValueError, match="Input streaming not currently supported"):
+    with pytest.raises(
+        VLLMValidationError, match="Input streaming not currently supported"
+    ):
         params_stop = SamplingParams(max_tokens=10, stop=["stop"])
         async for _ in engine.generate(dummy_generator(), params_stop, "test_stop"):
             pass
diff --git a/tests/v1/e2e/spec_decode/test_spec_decode.py b/tests/v1/e2e/spec_decode/test_spec_decode.py
index 17c48721f80..482cbfb82ed 100644
--- a/tests/v1/e2e/spec_decode/test_spec_decode.py
+++ b/tests/v1/e2e/spec_decode/test_spec_decode.py
@@ -23,7 +23,7 @@ from vllm import LLM, SamplingParams
 from vllm.assets.base import VLLM_S3_BUCKET_URL
 from vllm.assets.image import VLM_IMAGES_DIR
 from vllm.benchmarks.datasets import InstructCoderDataset
-from vllm.config import VllmConfig, replace
+from vllm.config import CompilationConfig, VllmConfig, replace
 from vllm.distributed import cleanup_dist_env_and_memory
 from vllm.engine.arg_utils import EngineArgs
 from vllm.platforms import current_platform
@@ -160,6 +160,12 @@ def reset_torch_dynamo():
     torch._dynamo.reset()
 
 
+@pytest.fixture
+def disable_vllm_compile_cache_on_rocm(request: pytest.FixtureRequest) -> None:
+    if current_platform.is_rocm():
+        request.getfixturevalue("disable_vllm_compile_cache")
+
+
 @pytest.mark.parametrize(
     "speculative_config",
     [
@@ -175,21 +181,26 @@ def reset_torch_dynamo():
         },
     ],
 )
+@pytest.mark.usefixtures("disable_vllm_compile_cache_on_rocm")
 @single_gpu_only
 @large_gpu_mark(min_gb=20)
 def test_ngram_and_suffix_correctness(
     speculative_config: dict,
     model_name: str,
+    vllm_runner,
 ):
-    spec_llm = LLM(
-        model=model_name,
+    with vllm_runner(
+        model_name,
+        # Keep LLM defaults; VllmRunner only provides lifecycle cleanup here.
+        trust_remote_code=False,
+        enable_chunked_prefill=None,
         speculative_config=speculative_config,
         max_model_len=4096,
-    )
-    evaluate_llm_for_gsm8k(spec_llm)
-    del spec_llm
-    torch.accelerator.empty_cache()
-    cleanup_dist_env_and_memory()
+        # Preserve LLM's default compilation/cudagraph configuration. Without
+        # this, VllmRunner injects its reduced test-only capture sizes.
+        compilation_config=CompilationConfig(),
+    ) as runner:
+        evaluate_llm_for_gsm8k(runner.llm)
 
 
 @pytest.mark.parametrize("async_scheduling", [True], ids=["async"])
diff --git a/tests/v1/ec_connector/integration/run_epd_correctness_test.sh b/tests/v1/ec_connector/integration/run_epd_correctness_test.sh
index 65716444a57..c58df0c076a 100644
--- a/tests/v1/ec_connector/integration/run_epd_correctness_test.sh
+++ b/tests/v1/ec_connector/integration/run_epd_correctness_test.sh
@@ -21,12 +21,19 @@ GIT_ROOT="${GIT_ROOT:-$(cd -- "${SCRIPT_DIR}/../../../.." && pwd -P)}"
 
 # Model to test
 MODEL="${MODEL:-Qwen/Qwen2.5-VL-3B-Instruct}"
+MAX_MODEL_LEN="${MAX_MODEL_LEN:-10240}"
+GPU_MEMORY_UTILIZATION="${GPU_MEMORY_UTILIZATION:-0.7}"
+MAX_NUM_SEQS="${MAX_NUM_SEQS:-128}"
 
 # Set 1 to use multimodal prompts; else to use text-only
 USE_MM_PROMPTS="${USE_MM_PROMPTS:-1}"
-MM_FLAG=""
-if [ "$USE_MM_PROMPTS" = "1" ]; then
-    MM_FLAG="--use_mm_prompts"
+USE_TWO_IMAGE_PROMPT="${USE_TWO_IMAGE_PROMPT:-1}"
+TEST_FLAGS=()
+if [[ "$USE_MM_PROMPTS" == "1" ]]; then
+    TEST_FLAGS+=(--use_mm_prompts)
+fi
+if [[ "$USE_TWO_IMAGE_PROMPT" != "1" ]]; then
+    TEST_FLAGS+=(--skip_two_image_prompt)
 fi
 
 # GPU configuration
@@ -36,6 +43,16 @@ GPU_D="${GPU_D:-2}"
 GPU_SINGLE="${GPU_SINGLE:-$GPU_P}"
 GPU_PD="${GPU_PD:-$GPU_P}"
 
+# Device platform and affinity environment variable
+DEVICE_PLATFORM="${DEVICE_PLATFORM:-cuda}"
+if [[ -z "${DEVICE_AFFINITY_ENV:-}" ]]; then
+    if [[ "${DEVICE_PLATFORM,,}" == "xpu" ]]; then
+        DEVICE_AFFINITY_ENV="ZE_AFFINITY_MASK"
+    else
+        DEVICE_AFFINITY_ENV="CUDA_VISIBLE_DEVICES"
+    fi
+fi
+
 # Port
 ENCODE_PORT="${ENCODE_PORT:-19534}"
 PREFILL_PORT="${PREFILL_PORT:-19535}"
@@ -87,11 +104,12 @@ run_baseline() {
     
     # Start baseline instance
     echo "Starting baseline instance on GPU $GPU_SINGLE, port $PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_SINGLE" vllm serve "$MODEL" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_SINGLE" vllm serve "$MODEL" \
         --port "$PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
-        --gpu-memory-utilization 0.7 \
-        --max-num-seqs 128 \
+        --gpu-memory-utilization 0.9 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         > "$LOG_PATH"/baseline.log 2>&1 &
     
@@ -112,7 +130,7 @@ run_baseline() {
         --model_name "$MODEL" \
         --mode baseline \
         --baseline_file "$BASELINE_FILE" \
-        $MM_FLAG
+        "${TEST_FLAGS[@]}"
     
     # Cleanup baseline
     echo "Stopping baseline instance..."
@@ -139,14 +157,15 @@ run_epd_1e_1pd() {
     
     # Start encoder instance
     echo "Starting encoder instance on GPU $GPU_E, port $ENCODE_PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_E" vllm serve "$MODEL" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_E" vllm serve "$MODEL" \
         --port "$ENCODE_PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
         --gpu-memory-utilization 0.01 \
         --enable-request-id-headers \
         --no-enable-prefix-caching \
         --max-num-batched-tokens 114688 \
-        --max-num-seqs 128 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         --ec-transfer-config '{
             "ec_connector": "ECExampleConnector",
@@ -160,12 +179,13 @@ run_epd_1e_1pd() {
     
     # Start prefill+decode instance
     echo "Starting PD instance on GPU $GPU_PD, port $PREFILL_DECODE_PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_PD" vllm serve "$MODEL" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_PD" vllm serve "$MODEL" \
         --port "$PREFILL_DECODE_PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
-        --gpu-memory-utilization 0.7 \
+        --gpu-memory-utilization "$GPU_MEMORY_UTILIZATION" \
         --enable-request-id-headers \
-        --max-num-seqs 128 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         --ec-transfer-config '{
             "ec_connector": "ECExampleConnector",
@@ -212,7 +232,7 @@ run_epd_1e_1pd() {
         --model_name "$MODEL" \
         --mode disagg \
         --baseline_file "$BASELINE_FILE" \
-        $MM_FLAG
+        "${TEST_FLAGS[@]}"
     
     # Cleanup
     echo "✓✓ 1E+1PD Correctness Test finished"
@@ -242,14 +262,15 @@ run_baseline_1p_1d() {
     
     # Start prefill instance
     echo "Starting prefill instance on GPU $GPU_P, port $PREFILL_PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_P" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_P" \
     VLLM_NIXL_SIDE_CHANNEL_PORT=5559 \
     vllm serve "$MODEL" \
         --port "$PREFILL_PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
-        --gpu-memory-utilization 0.7 \
+        --gpu-memory-utilization "$GPU_MEMORY_UTILIZATION" \
         --enable-request-id-headers \
-        --max-num-seqs 128 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         --kv-transfer-config '{
             "kv_connector": "NixlConnector",
@@ -260,14 +281,15 @@ run_baseline_1p_1d() {
     
     # Start decode instance
     echo "Starting decode instance on GPU $GPU_D, port $DECODE_PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_D" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_D" \
     VLLM_NIXL_SIDE_CHANNEL_PORT=6000 \
     vllm serve "$MODEL" \
         --port "$DECODE_PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
-        --gpu-memory-utilization 0.7 \
+        --gpu-memory-utilization "$GPU_MEMORY_UTILIZATION" \
         --enable-request-id-headers \
-        --max-num-seqs 128 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         --kv-transfer-config '{
             "kv_connector": "NixlConnector",
@@ -309,7 +331,7 @@ run_baseline_1p_1d() {
         --model_name "$MODEL" \
         --mode baseline_pd \
         --baseline_file "$BASELINE_PD_FILE" \
-        $MM_FLAG
+        "${TEST_FLAGS[@]}"
     
     # Cleanup
     echo "Stopping PD (1P+1D) instances..."
@@ -339,14 +361,15 @@ run_epd_1e_1p_1d() {
     
     # Start encoder instance
     echo "Starting encoder instance on GPU $GPU_E, port $ENCODE_PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_E" vllm serve "$MODEL" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_E" vllm serve "$MODEL" \
         --port "$ENCODE_PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
         --gpu-memory-utilization 0.01 \
         --enable-request-id-headers \
         --no-enable-prefix-caching \
         --max-num-batched-tokens 114688 \
-        --max-num-seqs 128 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         --ec-transfer-config '{
             "ec_connector": "ECExampleConnector",
@@ -360,14 +383,15 @@ run_epd_1e_1p_1d() {
     
     # Start prefill instance
     echo "Starting prefill instance on GPU $GPU_P, port $PREFILL_PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_P" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_P" \
     VLLM_NIXL_SIDE_CHANNEL_PORT=5559 \
     vllm serve "$MODEL" \
         --port "$PREFILL_PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
-        --gpu-memory-utilization 0.7 \
+        --gpu-memory-utilization "$GPU_MEMORY_UTILIZATION" \
         --enable-request-id-headers \
-        --max-num-seqs 128 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         --ec-transfer-config '{
             "ec_connector": "ECExampleConnector",
@@ -385,14 +409,15 @@ run_epd_1e_1p_1d() {
     
     # Start decode instance
     echo "Starting decode instance on GPU $GPU_D, port $DECODE_PORT"
-    CUDA_VISIBLE_DEVICES="$GPU_D" \
+    env "$DEVICE_AFFINITY_ENV=$GPU_D" \
     VLLM_NIXL_SIDE_CHANNEL_PORT=6000 \
     vllm serve "$MODEL" \
         --port "$DECODE_PORT" \
+        --max-model-len "$MAX_MODEL_LEN" \
         --enforce-eager \
-        --gpu-memory-utilization 0.7 \
+        --gpu-memory-utilization "$GPU_MEMORY_UTILIZATION" \
         --enable-request-id-headers \
-        --max-num-seqs 128 \
+        --max-num-seqs "$MAX_NUM_SEQS" \
         --allowed-local-media-path "${GIT_ROOT}"/tests/v1/ec_connector/integration \
         --kv-transfer-config '{
             "kv_connector": "NixlConnector",
@@ -438,7 +463,7 @@ run_epd_1e_1p_1d() {
         --model_name "$MODEL" \
         --mode disagg \
         --baseline_file "$BASELINE_PD_FILE" \
-        $MM_FLAG
+        "${TEST_FLAGS[@]}"
     
     # Cleanup
     echo "✓✓ 1E+1P+1D Correctness Test finished"
@@ -465,7 +490,7 @@ run_epd_1e_1pd
 # Step 3: Test baseline 1P + 1D
 run_baseline_1p_1d
 
-# Step 4: Test 1E + 1P + 1D
+# # Step 4: Test 1E + 1P + 1D
 run_epd_1e_1p_1d
 
 # Cleanup output file
diff --git a/tests/v1/ec_connector/integration/test_epd_correctness.py b/tests/v1/ec_connector/integration/test_epd_correctness.py
index eae4b742724..ece73efa42d 100644
--- a/tests/v1/ec_connector/integration/test_epd_correctness.py
+++ b/tests/v1/ec_connector/integration/test_epd_correctness.py
@@ -192,6 +192,12 @@ def main():
         help="Use multimodal prompts (default: use text-only for quick testing)",
     )
 
+    parser.add_argument(
+        "--skip_two_image_prompt",
+        action="store_true",
+        help="Skip the two-image multimodal prompt",
+    )
+
     args = parser.parse_args()
 
     print(f"Service URL: {args.service_url}")
@@ -221,7 +227,9 @@ def main():
 
     # Select prompts to use
     if args.use_mm_prompts:
-        test_prompts = SAMPLE_PROMPTS_MM
+        test_prompts = (
+            SAMPLE_PROMPTS_MM[:1] if args.skip_two_image_prompt else SAMPLE_PROMPTS_MM
+        )
         print("Using multimodal prompts")
     else:
         test_prompts = SAMPLE_PROMPTS_TEXT
diff --git a/tests/v1/engine/test_async_llm.py b/tests/v1/engine/test_async_llm.py
index afb6e4c98b7..3845d6297a7 100644
--- a/tests/v1/engine/test_async_llm.py
+++ b/tests/v1/engine/test_async_llm.py
@@ -19,6 +19,7 @@ from vllm.entrypoints.openai.chat_completion.protocol import (
 from vllm.entrypoints.openai.chat_completion.serving import OpenAIServingChat
 from vllm.entrypoints.openai.models.protocol import BaseModelPath
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
+from vllm.exceptions import VLLMValidationError
 from vllm.inputs import PromptType
 from vllm.outputs import RequestOutput
 from vllm.platforms import current_platform
@@ -485,7 +486,7 @@ async def test_dp_rank_argument():
             pass
 
         # Test with out-of-range DP rank.
-        with pytest.raises(ValueError):
+        with pytest.raises(VLLMValidationError):
             async for _ in engine.generate(
                 request_id="request-35",
                 prompt=TEXT_PROMPT,
@@ -554,8 +555,8 @@ async def test_header_dp_rank_argument():
         # Test 2: Out-of-range DP rank (1)
         mock_raw_request.headers = {"X-data-parallel-rank": "1"}
 
-        # should raise ValueError for out-of-range rank
-        with pytest.raises(ValueError):
+        # should raise VLLMValidationError for out-of-range rank
+        with pytest.raises(VLLMValidationError):
             await serving_chat.create_chat_completion(req, mock_raw_request)
 
 
diff --git a/tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh b/tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh
index 22682e02bea..fc5c04a1ad0 100755
--- a/tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh
+++ b/tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh
@@ -227,9 +227,12 @@ run_tests_for_model() {
     # Calculate side channel port
     SIDE_CHANNEL_PORT=$((5659 + i * $DECODER_TP_SIZE))
     INTERNAL_PORT=$((DECODER_INTERNAL_PORT_BASE + i * INTERNAL_PORT_STRIDE))
-    DECODER_INTERNAL_PORT_ENV=
+    # For non-DP mode, set VLLM_PORT to pin the internal port;
+    # For DP mode, set VLLM_DP_MASTER_PORT instead to avoid race condition.
     if [[ -z "${DP_EP:-}" ]]; then
       DECODER_INTERNAL_PORT_ENV="VLLM_PORT=$INTERNAL_PORT"
+    else
+      DECODER_INTERNAL_PORT_ENV="VLLM_DP_MASTER_PORT=$INTERNAL_PORT"
     fi
 
     echo "Starting decode instance $i on GPU $GPU_ID, port $PORT"
diff --git a/tests/v1/kv_connector/nixl_integration/run_edge_case_test.sh b/tests/v1/kv_connector/nixl_integration/run_edge_case_test.sh
index 9d8e4df8c53..c3240ab5c17 100755
--- a/tests/v1/kv_connector/nixl_integration/run_edge_case_test.sh
+++ b/tests/v1/kv_connector/nixl_integration/run_edge_case_test.sh
@@ -3,8 +3,8 @@ set -xe
 
 # Parse command line arguments
 KV_BUFFER_DEVICE="cuda"  # Default to cuda
-PREFILL_GPU_ID=4         # Default GPU IDs
-DECODE_GPU_ID=5
+PREFILL_GPU_ID="${PREFILL_GPU_ID:-4}"  # Default GPU IDs
+DECODE_GPU_ID="${DECODE_GPU_ID:-5}"
 while [[ $# -gt 0 ]]; do
   case $1 in
     --kv_buffer_device)
@@ -70,6 +70,7 @@ run_tests_for_model() {
   --port $PREFILL_PORT \
   --enforce-eager \
   --gpu-memory-utilization 0.2 \
+  --max-model-len 8192 \
   --kv-transfer-config '$KV_CONFIG'"
 
   FULL_CMD="$BASE_CMD"
@@ -84,6 +85,7 @@ run_tests_for_model() {
   --port $DECODE_PORT \
   --enforce-eager \
   --gpu-memory-utilization 0.2 \
+  --max-model-len 8192 \
   --kv-transfer-config '$KV_CONFIG'"
 
   FULL_CMD="$BASE_CMD"
@@ -98,7 +100,7 @@ run_tests_for_model() {
 
   # Build the command for the proxy server with all the hosts and ports
   PROXY_PORT=8192
-  PROXY_CMD="python ${GIT_ROOT}/tests/v1/kv_connector/nixl_integration/toy_proxy_server.py --port $PROXY_PORT"
+  PROXY_CMD="python3 ${GIT_ROOT}/tests/v1/kv_connector/nixl_integration/toy_proxy_server.py --port $PROXY_PORT"
   PROXY_CMD+=" --prefiller-ports ${PREFILL_PORT}"
   PROXY_CMD+=" --decoder-ports ${DECODE_PORT}"
   # Start the proxy server
@@ -110,7 +112,7 @@ run_tests_for_model() {
 
   # Run lm eval for this model
   echo "Running tests for $model_name"
-  PREFILL_PORT=$PREFILL_PORT DECODE_PORT=$DECODE_PORT PROXY_PORT=$PROXY_PORT python -m pytest -s -v "${GIT_ROOT}"/tests/v1/kv_connector/nixl_integration/test_edge_cases.py
+  PREFILL_PORT=$PREFILL_PORT DECODE_PORT=$DECODE_PORT PROXY_PORT=$PROXY_PORT python3 -m pytest -s -v "${GIT_ROOT}"/tests/v1/kv_connector/nixl_integration/test_edge_cases.py
 
   # Clean up before running next model
   cleanup_instances
diff --git a/tests/v1/kv_connector/unit/offloading_connector/test_config.py b/tests/v1/kv_connector/unit/offloading_connector/test_config.py
index fc426ff318d..5a66b463e30 100644
--- a/tests/v1/kv_connector/unit/offloading_connector/test_config.py
+++ b/tests/v1/kv_connector/unit/offloading_connector/test_config.py
@@ -8,10 +8,15 @@ from unittest.mock import MagicMock, patch
 import pytest
 import torch
 
+from tests.v1.kv_connector.unit.offloading_connector.utils import MockOffloadingSpec
 from vllm.config import KVTransferConfig, ParallelConfig, VllmConfig
 from vllm.distributed.kv_transfer.kv_connector.v1.offloading.config import (
     build_offloading_config,
 )
+from vllm.distributed.kv_transfer.kv_connector.v1.offloading.scheduler import (
+    SchedulerOffloadConfig,
+    is_store_reachable_swa_chunk,
+)
 from vllm.platforms import current_platform
 from vllm.v1.kv_cache_interface import (
     FullAttentionSpec,
@@ -179,6 +184,25 @@ def _make_hybrid_kv_cache_config() -> KVCacheConfig:
     )
 
 
+def _make_mamba_hybrid_kv_cache_config() -> KVCacheConfig:
+    return KVCacheConfig(
+        num_blocks=4,
+        kv_cache_tensors=[],
+        kv_cache_groups=[
+            KVCacheGroupSpec(["full_layer"], _full_attention_spec()),
+            KVCacheGroupSpec(
+                ["mamba_layer"],
+                MambaSpec(
+                    block_size=16,
+                    shapes=((1, 1),),
+                    dtypes=(torch.float32,),
+                    mamba_cache_mode="align",
+                ),
+            ),
+        ],
+    )
+
+
 def _parallelism_agnostic(kv_cache_groups: list[KVCacheGroupSpec]) -> bool:
     config = _make_vllm_config()
     kv_cache_config = KVCacheConfig(
@@ -267,6 +291,38 @@ def test_prefill_context_parallelism_does_not_scale_group_blocks():
     assert offloading_config.cache.blocks_per_chunk == 4
 
 
+def test_dcp_scales_attention_but_not_mamba_group_blocks():
+    config = _make_vllm_config(tensor_parallel_size=2, decode_context_parallel_size=2)
+    config.speculative_config = None
+
+    offloading_config = build_offloading_config(
+        config, _make_mamba_hybrid_kv_cache_config()
+    )
+
+    assert tuple(group.tokens_per_block for group in offloading_config.groups) == (
+        32,
+        16,
+    )
+    scheduler_config = SchedulerOffloadConfig.from_spec(
+        MockOffloadingSpec(offloading_config),
+        config,
+        _make_mamba_hybrid_kv_cache_config(),
+    )
+    mamba_group = scheduler_config.kv_group_configs[1]
+    assert mamba_group.alignment_chunk_count == 2
+    assert [
+        chunk_idx
+        for chunk_idx in range(4)
+        if is_store_reachable_swa_chunk(
+            chunk_idx,
+            4,
+            mamba_group.alignment_chunk_count,
+            mamba_group.sliding_window_size_in_chunks,
+            mamba_group.is_eagle_group,
+        )
+    ] == [1, 3]
+
+
 def test_preserves_data_parallel_index():
     config = _make_vllm_config()
     config.parallel_config.data_parallel_index = 2
diff --git a/tests/v1/kv_connector/unit/test_nixl_connector_hma.py b/tests/v1/kv_connector/unit/test_nixl_connector_hma.py
index 7557b183f6f..0124f6516ef 100644
--- a/tests/v1/kv_connector/unit/test_nixl_connector_hma.py
+++ b/tests/v1/kv_connector/unit/test_nixl_connector_hma.py
@@ -1447,6 +1447,11 @@ def test_register_kv_caches_hybrid_mla_dual_purpose_regions():
     kv_cache_config = _make_hybrid_mla_kv_cache_config()
     unified_page = kv_cache_config.kv_cache_groups[0].kv_cache_spec.page_size_bytes
     vllm_config = create_vllm_config(block_size=12)
+    # kv_buffer_device defaults to the *real* platform's device type, which on
+    # a CPU-only test host would make this a host-buffer worker: host xfer
+    # buffers are per-layer, so the HMA shared tensors would not be
+    # deduplicated. Pin it to the faked device type.
+    vllm_config.kv_transfer_config.kv_buffer_device = "cuda"
 
     fake_backend = MagicMock()
     fake_backend.get_supported_kernel_block_sizes.return_value = [4]
diff --git a/tests/v1/kv_connector/unit/test_nixl_desc_geometry.py b/tests/v1/kv_connector/unit/test_nixl_desc_geometry.py
index 57d3e028432..6ac9ce030d0 100644
--- a/tests/v1/kv_connector/unit/test_nixl_desc_geometry.py
+++ b/tests/v1/kv_connector/unit/test_nixl_desc_geometry.py
@@ -147,6 +147,11 @@ def _make_mla_hybrid_worker(local_block_size, kernel_block_size, num_logical_blo
 
     vllm_config = create_vllm_config(block_size=local_block_size)
     vllm_config.cache_config.enable_prefix_caching = False
+    # kv_buffer_device defaults to the *real* platform's device type, which on
+    # a CPU-only test host would make this a host-buffer worker: host xfer
+    # buffers are per-layer, so the HMA shared-tensor regions this test builds
+    # would not be deduplicated. Pin it to the faked device type.
+    vllm_config.kv_transfer_config.kv_buffer_device = "cuda"
 
     from unittest.mock import MagicMock
 
@@ -214,8 +219,10 @@ def _make_remote_meta(
 
     remote_ppl = remote_block_size // remote_kernel_block_size
     # Kernel-granularity pages are TP-independent for MLA hybrids and must
-    # match the local ones for the handshake to pass.
-    kernel_page = worker.block_len_per_layer[0]
+    # match the local ones for the handshake to pass, scaled down by the
+    # block-size ratio when the remote's kernel block is smaller.
+    block_size_ratio = worker.block_size // remote_kernel_block_size
+    kernel_page = worker.block_len_per_layer[0] // block_size_ratio
     return NixlAgentMetadata(
         engine_id="remote-engine",
         agent_metadata=b"remote-agent-meta",
@@ -327,37 +334,49 @@ def test_hetero_ppl_multi_read_writes_stay_within_request_blocks():
 
 
 def _resolve(
-    desc_arr, idx, bases, region_size, unified_page, kernel_page, logical_ids_attn
+    desc_arr,
+    idx,
+    bases,
+    region_size,
+    unified_page,
+    desc_page,
+    logical_ids_attn,
+    block_tokens,
 ):
     """Resolve a desc id to (region, kind, token_start) where kind is 'attn'
-    (kernel-page sized, block-aligned, in the request's attention blocks) or
-    'mamba'. token_start is the request-relative kernel-block index."""
+    (desc-page sized, sub-block-aligned, in the request's attention blocks)
+    or 'mamba'. token_start is the request-relative token offset, so local
+    and remote are comparable even when their kernel blocks differ in size."""
     addr, length, _ = (int(x) for x in desc_arr[int(idx)])
     for region, base in enumerate(bases):
         off = addr - base
         if 0 <= off < region_size:
             b = off // unified_page
             rem = off % unified_page
-            if (
-                length == kernel_page
-                and rem % kernel_page == 0
-                and (b in logical_ids_attn)
-            ):
+            if length == desc_page and rem % desc_page == 0 and b in logical_ids_attn:
                 pos = logical_ids_attn.index(b)
-                tokens_per_block = unified_page // kernel_page
-                sub = rem // kernel_page
-                return (region, "attn", (pos * tokens_per_block + sub))
+                tokens_per_desc = block_tokens * desc_page // unified_page
+                sub = rem // desc_page
+                return (region, "attn", pos * block_tokens + sub * tokens_per_desc)
             return (region, "mamba", None)
     raise AssertionError(f"desc {idx} addr {addr:#x} not in any region")
 
 
-def _run_hetero_case(local_block, kernel, remote_block, num_tokens, tp_size=2):
-    """Full pull-path run for one geometry; returns pairing records."""
+def _run_hetero_case(
+    local_block, kernel, remote_block, num_tokens, tp_size=2, remote_kernel=None
+):
+    """Full pull-path run for one geometry; returns pairing records.
+
+    ``remote_kernel`` defaults to the local kernel block size; a smaller
+    value additionally exercises block_size_ratio > 1.
+    """
     from vllm.distributed.kv_transfer.kv_connector.v1.nixl.metadata import (
         NixlConnectorMetadata,
     )
 
-    remote_ppl = remote_block // kernel
+    remote_kernel = remote_kernel or kernel
+    block_size_ratio = kernel // remote_kernel
+    remote_ppl = remote_block // remote_kernel
     matched = num_tokens - 1  # mamba N-1 rule
     n_local = -(-num_tokens // local_block)
     n_remote = -(-matched // remote_block)
@@ -372,7 +391,7 @@ def _run_hetero_case(local_block, kernel, remote_block, num_tokens, tp_size=2):
     meta_r = _make_remote_meta(
         worker,
         remote_block_size=remote_block,
-        remote_kernel_block_size=kernel,
+        remote_kernel_block_size=remote_kernel,
         remote_num_logical=max(2 * n_remote + 4, 8),
         remote_ssm_sizes=(48 // tp_size, 64 // tp_size),
     )
@@ -420,7 +439,9 @@ def _run_hetero_case(local_block, kernel, remote_block, num_tokens, tp_size=2):
     remote_bases = [0x10_000_000, 0x20_000_000]
     local_unified = worker._test_unified_page
     remote_unified = (local_unified // local_block) * remote_block
-    kernel_page = worker.block_len_per_layer[0]
+    # With block_size_ratio > 1 the local page is split into ratio sub-descs,
+    # each the size of a whole remote kernel page.
+    desc_page = worker.block_len_per_layer[0] // block_size_ratio
     meta_r_num_blocks_bytes = (meta_r.num_blocks // remote_ppl) * remote_unified
     covered_tokens = set()
     for op, lh, lids, rh, rids in nixl.xfers:
@@ -432,8 +453,9 @@ def _run_hetero_case(local_block, kernel, remote_block, num_tokens, tp_size=2):
                 local_bases,
                 len(worker._test_tensors[0]),
                 local_unified,
-                kernel_page,
+                desc_page,
                 local_attn,
+                local_block,
             )
             rreg, rkind, rtok = _resolve(
                 rarr,
@@ -441,8 +463,9 @@ def _run_hetero_case(local_block, kernel, remote_block, num_tokens, tp_size=2):
                 remote_bases,
                 meta_r_num_blocks_bytes,
                 remote_unified,
-                kernel_page,
+                desc_page,
                 remote_attn,
+                remote_block,
             )
             assert lkind == rkind, (
                 f"pair kind mismatch: local {lkind} vs remote {rkind} "
@@ -454,16 +477,16 @@ def _run_hetero_case(local_block, kernel, remote_block, num_tokens, tp_size=2):
             )
             if lkind == "attn":
                 assert ltok == rtok, (
-                    f"TOKEN MISALIGNMENT: local kernel block holds tokens "
-                    f"[{ltok * kernel}..) but receives remote tokens "
-                    f"[{rtok * kernel}..) "
+                    f"TOKEN MISALIGNMENT: local sub-block holds tokens "
+                    f"[{ltok}..) but receives remote tokens [{rtok}..) "
                     f"(geometry local_block={local_block}, "
                     f"remote_block={remote_block}, N={num_tokens})"
                 )
-                covered_tokens.add(ltok * kernel)
+                covered_tokens.add(ltok)
 
-    # Invariant 3: full coverage of the matched tokens.
-    needed = {t for t in range(0, matched - matched % kernel, kernel)}
+    # Invariant 3: full coverage of the matched tokens, at the finest
+    # transfer granularity (the remote kernel block).
+    needed = {t for t in range(0, matched - matched % remote_kernel, remote_kernel)}
     missing = needed - covered_tokens
     assert not missing, (
         f"tokens never transferred: {sorted(missing)[:8]} "
@@ -524,6 +547,29 @@ def test_hetero_ppl_token_alignment_sweep(local_block, remote_block, num_tokens)
     )
 
 
+@pytest.mark.cpu_test
+@pytest.mark.parametrize(
+    "num_tokens",
+    # Residues around the remote kernel block (4), the local kernel block
+    # (8), the remote logical block (8) and the local logical block (24).
+    [2, 5, 8, 9, 13, 16, 17, 21, 24, 25, 29, 32, 33, 41, 48, 49],
+)
+def test_hetero_ppl_with_block_size_ratio(num_tokens):
+    """Both hetero regimes at once: kernel blocks differ (local 8 / remote
+    4, block_size_ratio=2) *and* physical_blocks_per_logical differs (3 vs
+    2). The transfer is clipped at remote sub-block granularity by the
+    pairing and front-trimmed by _apply_prefix_caching, so the
+    untransferred tail can span both a partial block and whole blocks —
+    the case each of the two former zeroing paths handled only half of."""
+    _run_hetero_case(
+        local_block=24,
+        kernel=8,
+        remote_block=8,
+        remote_kernel=4,
+        num_tokens=num_tokens,
+    )
+
+
 @pytest.mark.cpu_test
 @pytest.mark.parametrize(
     "num_tokens",
@@ -571,3 +617,24 @@ def test_mla_hybrid_large_ppl_geometry(num_tokens):
         num_tokens=num_tokens,
         tp_size=8,
     )
+
+
+@pytest.mark.cpu_test
+def test_mismatched_mla_kernel_page_rejected_for_mla_hybrid():
+    """The MLA per-token page is TP-independent, so kernel block lengths
+    differing by anything other than the block-size ratio must fail the
+    handshake loudly rather than transfer at mismatched geometry."""
+    worker = _make_mla_hybrid_worker(
+        local_block_size=12, kernel_block_size=4, num_logical_blocks=8
+    )
+    meta_r = _make_remote_meta(
+        worker,
+        remote_block_size=8,
+        remote_kernel_block_size=4,
+        remote_num_logical=12,
+        remote_ssm_sizes=(24, 32),
+    )
+    # Equal kernel block sizes (ratio 1), but a half-sized per-token page.
+    meta_r.block_lens = [x // 2 for x in worker.block_len_per_layer]
+    with pytest.raises((AssertionError, RuntimeError)):
+        worker.add_remote_agent(meta_r, remote_tp_rank=0, remote_tp_size=2)
diff --git a/tests/v1/kv_connector/unit/test_nixl_push_connector.py b/tests/v1/kv_connector/unit/test_nixl_push_connector.py
index d670e5563b6..adf0063701c 100644
--- a/tests/v1/kv_connector/unit/test_nixl_push_connector.py
+++ b/tests/v1/kv_connector/unit/test_nixl_push_connector.py
@@ -43,6 +43,7 @@ from vllm.distributed.kv_transfer.kv_connector.v1.nixl.push_worker import (
 from vllm.distributed.kv_transfer.kv_connector.v1.nixl.utils import (
     get_base_request_id,
 )
+from vllm.v1.kv_cache_interface import FullAttentionSpec
 from vllm.v1.outputs import KVConnectorOutput
 
 from .utils import make_nixl_push_scheduler
@@ -337,6 +338,9 @@ class _StubWriterWorker(NixlPushConnectorWorker):
         w.engine_id = "test-decode-engine"
         w._remote_agents = {}
         w._physical_blocks_per_logical_kv_block = 1
+        # Single non-hybrid attention group, matching the stub block id lists.
+        w._has_mamba = False
+        w._group_spec_types = (FullAttentionSpec,)
 
         # Track _do_start_push_kv invocations.
         calls: list[tuple[str, Any, dict[str, Any]]] = []
diff --git a/tests/v1/kv_offload/cpu/__init__.py b/tests/v1/kv_offload/cpu/__init__.py
new file mode 100644
index 00000000000..e69de29bb2d
diff --git a/tests/v1/kv_offload/cpu/policies/__init__.py b/tests/v1/kv_offload/cpu/policies/__init__.py
new file mode 100644
index 00000000000..e69de29bb2d
diff --git a/tests/v1/kv_offload/cpu/policies/test_factory.py b/tests/v1/kv_offload/cpu/policies/test_factory.py
new file mode 100644
index 00000000000..14ccf8b67e7
--- /dev/null
+++ b/tests/v1/kv_offload/cpu/policies/test_factory.py
@@ -0,0 +1,111 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+from collections.abc import Iterable
+
+import pytest
+
+from vllm.v1.kv_offload.base import OffloadKey, ReqContext
+from vllm.v1.kv_offload.cpu.manager import CPUOffloadingManager
+from vllm.v1.kv_offload.cpu.policies.arc import ARCCachePolicy
+from vllm.v1.kv_offload.cpu.policies.base import BlockStatus, CachePolicy
+from vllm.v1.kv_offload.cpu.policies.factory import CachePolicyFactory
+from vllm.v1.kv_offload.cpu.policies.lru import LRUCachePolicy
+
+
+class _DummyCachePolicy(CachePolicy):
+    """Minimal CachePolicy for CachePolicyFactory registration tests. Loaded
+    by module path, so it must be importable at module scope (mirrors
+    tests/v1/kv_offload/test_factory.py's SingleArgExternalOffloadingSpec)."""
+
+    def __init__(self, cache_capacity: int) -> None:
+        self.cache_capacity = cache_capacity
+
+    def get(self, key: OffloadKey) -> BlockStatus | None:
+        return None
+
+    def insert(self, key: OffloadKey, block: BlockStatus) -> None:
+        pass
+
+    def remove(self, key: OffloadKey) -> None:
+        pass
+
+    def touch(self, keys: Iterable[OffloadKey], req_context: ReqContext) -> None:
+        pass
+
+    def evict(
+        self, n: int, protected: set[OffloadKey]
+    ) -> list[tuple[OffloadKey, BlockStatus]] | None:
+        return None
+
+    def clear(self) -> None:
+        pass
+
+
+@pytest.fixture(autouse=True)
+def restore_cache_policy_registry():
+    """Save and restore CachePolicyFactory._registry between tests."""
+    original = dict(CachePolicyFactory._registry)
+    yield
+    CachePolicyFactory._registry = original
+
+
+class TestCachePolicyFactory:
+    """Unit tests for CachePolicyFactory (registration/resolution by name)."""
+
+    def test_pre_registered_policies_can_be_imported(self):
+        """If someone moves a policy module but forgets to update
+        factory.py, CI fails."""
+        for name in CachePolicyFactory._registry:
+            cls = CachePolicyFactory._registry[name]()
+            assert issubclass(cls, CachePolicy)
+
+    def test_lru_and_arc_registered(self):
+        assert CachePolicyFactory.get_cache_policy_cls("lru") is LRUCachePolicy
+        assert CachePolicyFactory.get_cache_policy_cls("arc") is ARCCachePolicy
+
+    def test_register_and_resolve_custom_policy(self):
+        CachePolicyFactory.register_cache_policy(
+            "dummy",
+            "tests.v1.kv_offload.cpu.policies.test_factory",
+            "_DummyCachePolicy",
+        )
+        policy_cls = CachePolicyFactory.get_cache_policy_cls("dummy")
+        assert policy_cls is _DummyCachePolicy
+
+        manager = CPUOffloadingManager(num_blocks=4, cache_policy="dummy")
+        assert isinstance(manager._policy, _DummyCachePolicy)
+
+    def test_unregistered_policy_raises(self):
+        with pytest.raises(ValueError, match="Unknown cache policy"):
+            CachePolicyFactory.get_cache_policy_cls("nonexistent")
+
+    def test_duplicate_registration_raises(self):
+        with pytest.raises(ValueError, match="is already registered"):
+            CachePolicyFactory.register_cache_policy("lru", "some.module", "SomeClass")
+
+    def test_dynamic_load_via_cache_policy_module_path(self):
+        """Out-of-tree policy loaded via cache_policy_module_path, no
+        register_cache_policy() call -- this is how external projects
+        integrate a custom CachePolicy without forking/patching vLLM.
+        Mirrors tests/v1/kv_offload/test_factory.py's
+        test_dynamic_load_via_spec_module_path."""
+        policy_cls = CachePolicyFactory.get_cache_policy_cls(
+            "_DummyCachePolicy", "tests.v1.kv_offload.cpu.policies.test_factory"
+        )
+        assert policy_cls is _DummyCachePolicy
+
+    def test_manager_resolves_policy_via_module_path(self):
+        """End-to-end: CPUOffloadingManager resolves an unregistered policy
+        purely from cache_policy_module_path."""
+        manager = CPUOffloadingManager(
+            num_blocks=4,
+            cache_policy="_DummyCachePolicy",
+            cache_policy_module_path="tests.v1.kv_offload.cpu.policies.test_factory",
+        )
+        assert isinstance(manager._policy, _DummyCachePolicy)
+
+    def test_unregistered_policy_without_module_path_raises(self):
+        """eviction_policy not in registry + no cache_policy_module_path ->
+        ValueError, same shape as the OffloadingSpecFactory error path."""
+        with pytest.raises(ValueError, match="Unknown cache policy"):
+            CachePolicyFactory.get_cache_policy_cls("nonexistent", None)
diff --git a/tests/v1/kv_offload/cpu/test_manager.py b/tests/v1/kv_offload/cpu/test_manager.py
index 6520a93fd1a..5a36be16bf2 100644
--- a/tests/v1/kv_offload/cpu/test_manager.py
+++ b/tests/v1/kv_offload/cpu/test_manager.py
@@ -37,6 +37,7 @@ _EMPTY_REQ_CTX = make_req_context()
 def make_cpu_manager(
     num_blocks: int = 4,
     cache_policy: str = "lru",
+    cache_policy_module_path: str | None = None,
     enable_events: bool = False,
     store_threshold: int = 0,
     max_tracker_size: int = 64_000,
@@ -44,6 +45,7 @@ def make_cpu_manager(
     return CPUOffloadingManager(
         num_blocks=num_blocks,
         cache_policy=cache_policy,
+        cache_policy_module_path=cache_policy_module_path,
         enable_events=enable_events,
         store_threshold=store_threshold,
         max_tracker_size=max_tracker_size,
diff --git a/tests/v1/kv_offload/test_file_mapper.py b/tests/v1/kv_offload/test_file_mapper.py
index c2c4427e184..027f01133b2 100644
--- a/tests/v1/kv_offload/test_file_mapper.py
+++ b/tests/v1/kv_offload/test_file_mapper.py
@@ -51,6 +51,7 @@ def make_mapper_from_offloading_spec(**kwargs) -> FileMapper:
             data_parallel_index=0,
             is_parallelism_agnostic=kwargs.get("is_parallelism_agnostic", False),
         ),
+        replicated_layout=kwargs.get("replicated_layout", False),
     )
     spec = MagicMock(spec=OffloadingSpec)
     spec.config = config
@@ -205,3 +206,107 @@ def test_parallel_agnostic_separates_persistent_layouts():
     assert agnostic.base_path != specific.base_path
     assert "parallel_agnostic" not in agnostic.fields
     assert specific.fields["parallel_agnostic"] is False
+
+
+# ---------------------------------------------------------------------------
+# replicated_layout: OR'd into parallel-agnostic identity for compact rows
+# ---------------------------------------------------------------------------
+
+
+def test_replicated_layout_collapses_parallel_identity():
+    shared = dict(
+        model_name="mla-model",
+        groups=((16, "mla_layer"),),
+        replicated_layout=True,
+        parallel_agnostic=True,
+    )
+    tp2 = make_mapper_from_offloading_spec(tp_size=2, world_size=2, rank=1, **shared)
+    tp4 = make_mapper_from_offloading_spec(tp_size=4, world_size=4, rank=3, **shared)
+
+    assert tp2.base_path == tp4.base_path
+    for fm in (tp2, tp4):
+        assert fm.fields["tp_size"] == 1
+        assert fm.fields["pp_size"] == 1
+        assert fm.fields["pcp_size"] == 1
+        assert fm.fields["dcp_size"] == 1
+        assert fm.rank == 0
+        assert "parallel_agnostic" not in fm.fields
+        assert fm.fields["replicated_layout"] is True
+        assert fm.get_file_name(make_offload_key(b"\x01" * 8, 0)).startswith(
+            f"{fm.base_path}_r0/"
+        )
+
+
+def test_replicated_layout_requires_caller_opt_in():
+    fm = make_mapper_from_offloading_spec(
+        tp_size=2,
+        world_size=2,
+        rank=1,
+        replicated_layout=True,
+        parallel_agnostic=False,
+    )
+    assert fm.fields["tp_size"] == 2
+    assert fm.rank == 1
+    assert fm.fields["parallel_agnostic"] is False
+    assert "replicated_layout" not in fm.fields
+    baseline = make_mapper_from_offloading_spec(
+        tp_size=2,
+        world_size=2,
+        rank=1,
+        replicated_layout=False,
+        parallel_agnostic=False,
+    )
+    assert fm.base_path == baseline.base_path
+
+
+def test_non_replicated_keeps_parallel_identity():
+    fm = make_mapper_from_offloading_spec(
+        tp_size=4,
+        world_size=4,
+        rank=2,
+        replicated_layout=False,
+        is_parallelism_agnostic=False,
+        parallel_agnostic=True,
+    )
+    assert fm.fields["tp_size"] == 4
+    assert fm.rank == 2
+    assert fm.fields["parallel_agnostic"] is False
+    assert fm.get_file_name(make_offload_key(b"\x02" * 8, 0)).startswith(
+        f"{fm.base_path}_r2/"
+    )
+
+
+def test_replicated_and_parallelism_agnostic_separate_layouts():
+    shared = dict(
+        model_name="shared-model",
+        groups=((16, "layer0"),),
+        tp_size=2,
+        world_size=2,
+        rank=1,
+        parallel_agnostic=True,
+    )
+    via_agnostic = make_mapper_from_offloading_spec(
+        is_parallelism_agnostic=True,
+        replicated_layout=False,
+        **shared,
+    )
+    via_replicated = make_mapper_from_offloading_spec(
+        is_parallelism_agnostic=False,
+        replicated_layout=True,
+        **shared,
+    )
+    assert via_agnostic.base_path != via_replicated.base_path
+    assert "replicated_layout" not in via_agnostic.fields
+    assert via_replicated.fields["replicated_layout"] is True
+
+
+def test_replicated_layout_run_config_tp_invariant():
+    shared = dict(
+        model_name="mla-model",
+        groups=((16, "mla_layer"),),
+        replicated_layout=True,
+        parallel_agnostic=True,
+    )
+    tp2 = make_mapper_from_offloading_spec(tp_size=2, world_size=2, rank=0, **shared)
+    tp4 = make_mapper_from_offloading_spec(tp_size=4, world_size=4, rank=2, **shared)
+    assert tp2.get_run_config() == tp4.get_run_config()
diff --git a/tests/v1/kv_offload/tiering/p2p/test_manager.py b/tests/v1/kv_offload/tiering/p2p/test_manager.py
index cd6e8f382a5..c8a9d7f3ea3 100644
--- a/tests/v1/kv_offload/tiering/p2p/test_manager.py
+++ b/tests/v1/kv_offload/tiering/p2p/test_manager.py
@@ -1612,16 +1612,17 @@ class TestBindHostPortDefaults:
         monkeypatch.setattr(
             manager_module,
             "NixlTransport",
-            lambda agent_name, *a, **k: calls.update(nixl_name=agent_name)
-            or SimpleNamespace(),
+            lambda agent_name, *a, **k: (
+                calls.update(nixl_name=agent_name) or SimpleNamespace()
+            ),
         )
         monkeypatch.setattr(
             manager_module,
             "ZmqTransport",
-            lambda local_id, host, port, *a, **k: calls.update(
-                zmq_id=local_id, zmq_host=host, zmq_port=port
-            )
-            or SimpleNamespace(),
+            lambda local_id, host, port, *a, **k: (
+                calls.update(zmq_id=local_id, zmq_host=host, zmq_port=port)
+                or SimpleNamespace()
+            ),
         )
         spec = SimpleNamespace(
             blocks_per_chunk=1,
diff --git a/tests/v1/kv_offload/tiering/p2p/test_sessions.py b/tests/v1/kv_offload/tiering/p2p/test_sessions.py
index cab4ce6705a..20e4bad2186 100644
--- a/tests/v1/kv_offload/tiering/p2p/test_sessions.py
+++ b/tests/v1/kv_offload/tiering/p2p/test_sessions.py
@@ -49,6 +49,7 @@ from vllm.v1.kv_offload.tiering.p2p.session.protocol import (
 from vllm.v1.kv_offload.tiering.p2p.session.server import (
     _CANCEL_DRAIN_TIMEOUT_S,
     _InflightXfer,
+    _OutboundRequestState,
 )
 from vllm.v1.kv_offload.tiering.p2p.session.session import (
     _MAX_CONSECUTIVE_DISPATCH_ERRORS,
@@ -315,10 +316,26 @@ def _activate(
 # either a missing entry or a None field. These helpers paper over that.
 
 
+def _client_load(session: P2PSession, kv_request_id: str):
+    """The single in-flight load of a kv_request_id (loads are per-round)."""
+    loads = session._client._requests[kv_request_id].loads
+    assert len(loads) == 1
+    return next(iter(loads.values()))
+
+
 def _srv_outbound(session: P2PSession, kv_request_id: str):
-    """Outbound serve state for a kv_request_id, or None (idle / GC'd)."""
+    """Serve-side round for a kv_request_id, or None (idle / GC'd).
+
+    Rounds are keyed by wire round_seq; surfaces the demanded round when
+    a fetch has bound one, else any parked supply round.
+    """
     st = session._server._requests.get(kv_request_id)
-    return st.outbound if st is not None else None
+    if st is None or not st.outbound:
+        return None
+    for rnd in st.outbound.values():
+        if rnd.demand_received:
+            return rnd
+    return next(iter(st.outbound.values()))
 
 
 def _srv_lookups(session: P2PSession) -> list:
@@ -330,8 +347,10 @@ def _srv_lookups(session: P2PSession) -> list:
 
 def _srv_abort_started(session: P2PSession, kv_request_id: str) -> float | None:
     """Pending-abort start time for a kv_request_id, or None."""
-    st = session._server._requests.get(kv_request_id)
-    return st.abort_started_at if st is not None else None
+    for (kv, _), started in session._server._pending_aborts.items():
+        if kv == kv_request_id:
+            return started
+    return None
 
 
 def _srv_inflight_count(session: P2PSession, kv_request_id: str) -> int:
@@ -479,6 +498,7 @@ class TestClientFlows:
         conn.enqueue(
             {
                 TYPE_KEY: TransferDoneMsg.TYPE,
+                TransferDoneMsg.ROUND_SEQ: 0,
                 TransferDoneMsg.KV_REQUEST_ID: "req-1",
                 TransferDoneMsg.SUCCESS: True,
             }
@@ -495,6 +515,7 @@ class TestClientFlows:
         conn.enqueue(
             {
                 TYPE_KEY: TransferDoneMsg.TYPE,
+                TransferDoneMsg.ROUND_SEQ: 0,
                 TransferDoneMsg.KV_REQUEST_ID: "req-1",
                 TransferDoneMsg.SUCCESS: False,
             }
@@ -538,6 +559,7 @@ class TestClientFlows:
         conn.enqueue(
             {
                 TYPE_KEY: TransferDoneMsg.TYPE,
+                TransferDoneMsg.ROUND_SEQ: 0,
                 TransferDoneMsg.KV_REQUEST_ID: "req-1",
                 TransferDoneMsg.SUCCESS: True,
             }
@@ -552,7 +574,7 @@ class TestClientFlows:
         session.request_blocks(
             job_id=1, kv_request_id="req-1", keys=[b"k"], block_ids=[0]
         )
-        session._client._requests["req-1"].load.submitted_at = time.monotonic() - 60.0
+        _client_load(session, "req-1").submitted_at = time.monotonic() - 60.0
         session.poll()
         abort = conn._sent[-1]
         assert abort[TYPE_KEY] == AbortFetchMsg.TYPE
@@ -569,7 +591,7 @@ class TestClientFlows:
             job_id=7, kv_request_id="req-7", keys=[b"k"], block_ids=[0]
         )
         # 1) Trip the load timeout to send AbortFetch and stamp aborted_at.
-        session._client._requests["req-7"].load.submitted_at = (
+        _client_load(session, "req-7").submitted_at = (
             time.monotonic() - _LOAD_TIMEOUT_S - 1.0
         )
         loads = session.poll().loads
@@ -579,11 +601,11 @@ class TestClientFlows:
             and m[AbortFetchMsg.KV_REQUEST_ID] == "req-7"
             for m in conn._sent
         )
-        assert session._client._requests["req-7"].load.aborted_at is not None
+        assert _client_load(session, "req-7").aborted_at is not None
 
         # 2) Now backdate aborted_at past the abort-ack timeout. No ack ever
         # arrived from the peer.
-        session._client._requests["req-7"].load.aborted_at = (
+        _client_load(session, "req-7").aborted_at = (
             time.monotonic() - _ABORT_ACK_TIMEOUT_S - 1.0
         )
         loads = session.poll().loads
@@ -599,17 +621,18 @@ class TestClientFlows:
         session.request_blocks(
             job_id=8, kv_request_id="req-8", keys=[b"k"], block_ids=[0]
         )
-        session._client._requests["req-8"].load.submitted_at = (
+        _client_load(session, "req-8").submitted_at = (
             time.monotonic() - _LOAD_TIMEOUT_S - 1.0
         )
         # First poll: AbortFetch goes out.
         session.poll()
-        assert session._client._requests["req-8"].load.aborted_at is not None
+        assert _client_load(session, "req-8").aborted_at is not None
 
         # Peer acks the abort.
         conn.enqueue(
             {
                 TYPE_KEY: AbortAckMsg.TYPE,
+                AbortAckMsg.ROUND_SEQ: 0,
                 AbortAckMsg.KV_REQUEST_ID: "req-8",
             }
         )
@@ -915,6 +938,7 @@ class TestLookupFlow:
         conn.enqueue(
             {
                 TYPE_KEY: LookupMsg.TYPE,
+                LookupMsg.ROUND_SEQ: 0,
                 LookupMsg.KV_REQUEST_ID: "req-1",
                 LookupMsg.KEYS: [b"hX", b"hY", b"hZ"],
             }
@@ -951,6 +975,7 @@ def _send_lookup(conn: FakeConnection, kv_request_id: str, keys: list[bytes]):
     conn.enqueue(
         {
             TYPE_KEY: LookupMsg.TYPE,
+            LookupMsg.ROUND_SEQ: 0,
             LookupMsg.KV_REQUEST_ID: kv_request_id,
             LookupMsg.KEYS: list(keys),
         }
@@ -1201,6 +1226,7 @@ class TestServerLookupHandling:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [],
                 FetchMsg.BLOCK_INDEXES: [],
@@ -1240,6 +1266,7 @@ class TestServerLookupHandling:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"hA", b"hB"],
                 FetchMsg.BLOCK_INDEXES: [20, 21],
@@ -1277,6 +1304,7 @@ class TestServerFlows:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1", b"k2"],
                 FetchMsg.BLOCK_INDEXES: [10, 11],
@@ -1295,6 +1323,7 @@ class TestServerFlows:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [5],
@@ -1313,6 +1342,7 @@ class TestServerFlows:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [5],
@@ -1331,6 +1361,7 @@ class TestServerFlows:
         conn.enqueue(
             {
                 TYPE_KEY: AbortFetchMsg.TYPE,
+                AbortFetchMsg.ROUND_SEQ: 0,
                 AbortFetchMsg.KV_REQUEST_ID: "req-1",
             }
         )
@@ -1349,13 +1380,19 @@ class TestServerFlows:
         tid = 42
         session._server._inflight_add(
             tid,
-            _InflightXfer(kv_request_id="req-1", block_count=1, job_ids={1}),
+            _InflightXfer(
+                kv_request_id="req-1",
+                block_count=1,
+                job_ids={1},
+                round=_OutboundRequestState(inflight=1),
+            ),
         )
         transport._cancel_still_inflight.add(tid)
 
         conn.enqueue(
             {
                 TYPE_KEY: AbortFetchMsg.TYPE,
+                AbortFetchMsg.ROUND_SEQ: 0,
                 AbortFetchMsg.KV_REQUEST_ID: "req-1",
             }
         )
@@ -1376,13 +1413,19 @@ class TestServerFlows:
         tid = 42
         session._server._inflight_add(
             tid,
-            _InflightXfer(kv_request_id="req-1", block_count=1, job_ids={1}),
+            _InflightXfer(
+                kv_request_id="req-1",
+                block_count=1,
+                job_ids={1},
+                round=_OutboundRequestState(inflight=1),
+            ),
         )
         transport._cancel_still_inflight.add(tid)
 
         conn.enqueue(
             {
                 TYPE_KEY: AbortFetchMsg.TYPE,
+                AbortFetchMsg.ROUND_SEQ: 0,
                 AbortFetchMsg.KV_REQUEST_ID: "req-1",
             }
         )
@@ -1409,20 +1452,26 @@ class TestServerFlows:
         tid = 42
         session._server._inflight_add(
             tid,
-            _InflightXfer(kv_request_id="req-1", block_count=1, job_ids={1}),
+            _InflightXfer(
+                kv_request_id="req-1",
+                block_count=1,
+                job_ids={1},
+                round=_OutboundRequestState(inflight=1),
+            ),
         )
         transport._cancel_still_inflight.add(tid)
 
         conn.enqueue(
             {
                 TYPE_KEY: AbortFetchMsg.TYPE,
+                AbortFetchMsg.ROUND_SEQ: 0,
                 AbortFetchMsg.KV_REQUEST_ID: "req-1",
             }
         )
         session.poll()
         assert _srv_abort_started(session, "req-1") is not None
         # Backdate past the drain deadline.
-        session._server._requests["req-1"].abort_started_at = (
+        session._server._pending_aborts[("req-1", 0)] = (
             time.monotonic() - _CANCEL_DRAIN_TIMEOUT_S - 1.0
         )
         # Even if the transport still claims it can't cancel, the
@@ -1445,13 +1494,19 @@ class TestServerFlows:
         tid = 42
         session._server._inflight_add(
             tid,
-            _InflightXfer(kv_request_id="req-1", block_count=1, job_ids={1}),
+            _InflightXfer(
+                kv_request_id="req-1",
+                block_count=1,
+                job_ids={1},
+                round=_OutboundRequestState(inflight=1),
+            ),
         )
         transport._cancel_still_inflight.add(tid)
 
         conn.enqueue(
             {
                 TYPE_KEY: AbortFetchMsg.TYPE,
+                AbortFetchMsg.ROUND_SEQ: 0,
                 AbortFetchMsg.KV_REQUEST_ID: "req-1",
             }
         )
@@ -1463,6 +1518,7 @@ class TestServerFlows:
         conn.enqueue(
             {
                 TYPE_KEY: AbortFetchMsg.TYPE,
+                AbortFetchMsg.ROUND_SEQ: 0,
                 AbortFetchMsg.KV_REQUEST_ID: "req-1",
             }
         )
@@ -1497,6 +1553,7 @@ class TestServerFlows:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [5],
@@ -1529,6 +1586,7 @@ class TestServerFlows:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [5],
@@ -1571,6 +1629,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [5],
@@ -1596,6 +1655,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1", b"k2"],
                 FetchMsg.BLOCK_INDEXES: [10, 11],
@@ -1631,6 +1691,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [10],
@@ -1665,6 +1726,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [10],
@@ -1686,6 +1748,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-2",
                 FetchMsg.KEYS: [b"k1", b"k2"],
                 FetchMsg.BLOCK_INDEXES: [10, 11],
@@ -1719,6 +1782,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"demand"],
                 FetchMsg.BLOCK_INDEXES: [5],
@@ -1753,6 +1817,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [10],
@@ -1790,6 +1855,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [10],
@@ -1824,6 +1890,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1", b"k2", b"k3"],
                 FetchMsg.BLOCK_INDEXES: [10, 11, 12],
@@ -1885,6 +1952,7 @@ class TestFinishRequestServerSide:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1", b"k2"],
                 FetchMsg.BLOCK_INDEXES: [10, 11],
@@ -1951,6 +2019,7 @@ class TestBidirectional:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-srv",
                 FetchMsg.KEYS: [b"served"],
                 FetchMsg.BLOCK_INDEXES: [7],
@@ -1981,6 +2050,7 @@ class TestBidirectional:
         conn.enqueue(
             {
                 TYPE_KEY: TransferDoneMsg.TYPE,
+                TransferDoneMsg.ROUND_SEQ: 0,
                 TransferDoneMsg.KV_REQUEST_ID: "req-cli",
                 TransferDoneMsg.SUCCESS: True,
             }
@@ -2168,6 +2238,7 @@ class TestAdversarial:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-bad",
                 FetchMsg.KEYS: [b"k1", b"k2"],
                 FetchMsg.BLOCK_INDEXES: [1],
@@ -2216,6 +2287,7 @@ class TestDispatchErrorHandling:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-bad",
                 FetchMsg.KEYS: [b"k1", b"k2"],
                 FetchMsg.BLOCK_INDEXES: [1],
@@ -2246,6 +2318,7 @@ class TestDispatchErrorHandling:
         conn.enqueue(
             {
                 TYPE_KEY: FetchMsg.TYPE,
+                FetchMsg.ROUND_SEQ: 0,
                 FetchMsg.KV_REQUEST_ID: "req-1",
                 FetchMsg.KEYS: [b"k1"],
                 FetchMsg.BLOCK_INDEXES: [0],
@@ -2269,6 +2342,7 @@ class TestDispatchErrorHandling:
             conn.enqueue(
                 {
                     TYPE_KEY: FetchMsg.TYPE,
+                    FetchMsg.ROUND_SEQ: 0,
                     FetchMsg.KV_REQUEST_ID: "req-1",
                     FetchMsg.KEYS: [b"k1"],
                     FetchMsg.BLOCK_INDEXES: [0],
@@ -2296,6 +2370,7 @@ class TestDispatchErrorHandling:
             conn.enqueue(
                 {
                     TYPE_KEY: FetchMsg.TYPE,
+                    FetchMsg.ROUND_SEQ: 0,
                     FetchMsg.KV_REQUEST_ID: "req-1",
                     FetchMsg.KEYS: [b"k1"],
                     FetchMsg.BLOCK_INDEXES: [0],
@@ -2341,6 +2416,7 @@ class TestInflightPerReqInvariant:
             conn.enqueue(
                 {
                     TYPE_KEY: FetchMsg.TYPE,
+                    FetchMsg.ROUND_SEQ: 0,
                     FetchMsg.KV_REQUEST_ID: kv_id,
                     FetchMsg.KEYS: keys,
                     FetchMsg.BLOCK_INDEXES: indexes,
@@ -2394,7 +2470,12 @@ class TestInflightPerReqInvariant:
                 tid = kv_id_idx * 10 + j
                 session._server._inflight_add(
                     tid,
-                    _InflightXfer(kv_request_id=kv_id, block_count=1, job_ids={tid}),
+                    _InflightXfer(
+                        kv_request_id=kv_id,
+                        block_count=1,
+                        job_ids={tid},
+                        round=_OutboundRequestState(inflight=1),
+                    ),
                 )
         assert _srv_total_inflight(session) == len(session._server._inflight)
         assert session._server._has_inflight_for("req-0")
@@ -2488,6 +2569,7 @@ class TestFetchMsgValidation:
     def _valid_msg(self) -> dict:
         return {
             TYPE_KEY: FetchMsg.TYPE,
+            FetchMsg.ROUND_SEQ: 0,
             FetchMsg.KV_REQUEST_ID: "req-1",
             FetchMsg.KEYS: [b"k1", b"k2"],
             FetchMsg.BLOCK_INDEXES: [0, 1],
@@ -2513,6 +2595,7 @@ class TestTransferDoneMsgValidation:
     def test_valid_message_passes(self):
         msg = {
             TYPE_KEY: TransferDoneMsg.TYPE,
+            TransferDoneMsg.ROUND_SEQ: 0,
             TransferDoneMsg.KV_REQUEST_ID: "req-1",
             TransferDoneMsg.SUCCESS: True,
         }
@@ -2521,6 +2604,7 @@ class TestTransferDoneMsgValidation:
     def test_success_wrong_type(self):
         msg = {
             TYPE_KEY: TransferDoneMsg.TYPE,
+            TransferDoneMsg.ROUND_SEQ: 0,
             TransferDoneMsg.KV_REQUEST_ID: "req-1",
             TransferDoneMsg.SUCCESS: 1,
         }
diff --git a/tests/v1/kv_offload/tiering/test_fs_tier.py b/tests/v1/kv_offload/tiering/test_fs_tier.py
index 66533e3676b..2959ac1aa03 100644
--- a/tests/v1/kv_offload/tiering/test_fs_tier.py
+++ b/tests/v1/kv_offload/tiering/test_fs_tier.py
@@ -52,8 +52,18 @@ _DTYPE: torch.dtype = torch.float32
 _CTX = ReqContext(req_id="test")
 
 
-def _make_offloading_spec(enable_kv_cache_events: bool) -> MagicMock:
+def _make_offloading_spec(
+    enable_kv_cache_events: bool = False,
+    *,
+    tp_size: int = 1,
+    rank: int = 0,
+    world_size: int | None = None,
+    replicated_layout: bool = False,
+    is_parallelism_agnostic: bool = False,
+) -> MagicMock:
     """Mock spec with an explicit global KV events flag."""
+    if world_size is None:
+        world_size = tp_size
     spec = MagicMock()
     spec.config = OffloadingConfig(
         groups=(),
@@ -64,15 +74,16 @@ def _make_offloading_spec(enable_kv_cache_events: bool) -> MagicMock:
         model=OffloadingModelConfig(name="test-model", dtype="float32"),
         cache=OffloadingCacheConfig(tokens_per_hash=16, blocks_per_chunk=1),
         parallel=OffloadingParallelConfig(
-            rank=0,
-            world_size=1,
-            tp_size=1,
+            rank=rank,
+            world_size=world_size,
+            tp_size=tp_size,
             pp_size=1,
             pcp_size=1,
             dcp_size=1,
             data_parallel_index=0,
-            is_parallelism_agnostic=False,
+            is_parallelism_agnostic=is_parallelism_agnostic,
         ),
+        replicated_layout=replicated_layout,
     )
     spec.blocks_per_chunk = 1
     spec.kv_events_config = OffloadingKVEventsConfig(
@@ -725,3 +736,48 @@ def test_cascade_store_emits_fs_event_through_tiering_manager(tmp_path):
         assert not fs_events[0].removed
     finally:
         tier.shutdown()
+
+
+def test_fs_tier_cross_tp_round_trip(tmp_path):
+    """TP=2 replicated writer and TP=4 reader share namespace and bytes."""
+    root = str(tmp_path)
+    writer_tensor = _page_aligned_rand_tensor(4, _BLOCK_ELEMENTS)
+    expected = writer_tensor[0].clone()
+    writer = FileSystemTierManager(
+        offloading_spec=_make_offloading_spec(
+            tp_size=2, world_size=2, rank=0, replicated_layout=True
+        ),
+        primary_kv_view=memoryview(writer_tensor.numpy()),
+        tier_type="fs",
+        root_dir=root,
+        n_read_threads=2,
+        n_write_threads=2,
+    )
+    try:
+        writer.submit_store(make_job(1, [key(7)], [0]))
+        assert all(r.success for r in drain(writer))
+        writer_base = writer.file_mapper.base_path
+        writer_path = writer.file_mapper.get_file_name(key(7))
+    finally:
+        writer.shutdown()
+
+    reader_tensor = _page_aligned_zero_tensor(4, _BLOCK_ELEMENTS)
+    reader = FileSystemTierManager(
+        offloading_spec=_make_offloading_spec(
+            tp_size=4, world_size=4, rank=3, replicated_layout=True
+        ),
+        primary_kv_view=memoryview(reader_tensor.numpy()),
+        tier_type="fs",
+        root_dir=root,
+        n_read_threads=2,
+        n_write_threads=2,
+    )
+    try:
+        assert reader.file_mapper.base_path == writer_base
+        assert reader.file_mapper.get_file_name(key(7)) == writer_path
+        assert lookup_and_wait(reader, [key(7)]) == [LookupResult.HIT]
+        reader.submit_load(make_job(2, [key(7)], [1], is_promotion=True))
+        assert all(r.success for r in drain(reader))
+        assert torch.allclose(reader_tensor[1], expected)
+    finally:
+        reader.shutdown()
diff --git a/tests/v1/kv_offload/tiering/test_obj_tier.py b/tests/v1/kv_offload/tiering/test_obj_tier.py
index fc30e1437a7..661438dce63 100644
--- a/tests/v1/kv_offload/tiering/test_obj_tier.py
+++ b/tests/v1/kv_offload/tiering/test_obj_tier.py
@@ -35,6 +35,10 @@ from vllm.v1.kv_offload.config import (
     OffloadingParallelConfig,
 )
 from vllm.v1.kv_offload.tiering.base import JobMetadata, JobResult
+from vllm.v1.kv_offload.tiering.manager import (
+    CPUPrimaryTierOffloadingManager,
+    TieringOffloadingManager,
+)
 from vllm.v1.kv_offload.tiering.obj.config import ObjStoreConfig
 from vllm.v1.kv_offload.tiering.obj.manager import ObjectStoreSecondaryTierManager
 
@@ -43,7 +47,17 @@ from vllm.v1.kv_offload.tiering.obj.manager import ObjectStoreSecondaryTierManag
 # ---------------------------------------------------------------------------
 
 
-def _make_offloading_config(enable_kv_cache_events: bool) -> OffloadingConfig:
+def _make_offloading_config(
+    enable_kv_cache_events: bool,
+    *,
+    tp_size: int = 1,
+    rank: int = 0,
+    world_size: int | None = None,
+    replicated_layout: bool = False,
+    is_parallelism_agnostic: bool = False,
+) -> OffloadingConfig:
+    if world_size is None:
+        world_size = tp_size
     return OffloadingConfig(
         groups=(),
         worker_kv_bytes_per_block=0,
@@ -53,15 +67,16 @@ def _make_offloading_config(enable_kv_cache_events: bool) -> OffloadingConfig:
         model=OffloadingModelConfig(name="test/model", dtype="float16"),
         cache=OffloadingCacheConfig(tokens_per_hash=16, blocks_per_chunk=1),
         parallel=OffloadingParallelConfig(
-            rank=0,
-            world_size=1,
-            tp_size=1,
+            rank=rank,
+            world_size=world_size,
+            tp_size=tp_size,
             pp_size=1,
             pcp_size=1,
             dcp_size=1,
             data_parallel_index=0,
-            is_parallelism_agnostic=False,
+            is_parallelism_agnostic=is_parallelism_agnostic,
         ),
+        replicated_layout=replicated_layout,
     )
 
 
@@ -209,12 +224,14 @@ def _make_events_spec(enable_kv_cache_events: bool) -> SimpleNamespace:
 def _make_tier(
     num_blocks: int = 4,
     offloading_spec: SimpleNamespace = _OFFLOADING_SPEC,
+    primary_kv_view: memoryview | None = None,
     **tier_kwargs,
 ) -> tuple[ObjectStoreSecondaryTierManager, MockNixlAgent]:
     """Create a tier backed by a fresh MockNixlAgent."""
     mock_agent = MockNixlAgent()
-    tensor = torch.zeros((num_blocks, _BLOCK_ELEMENTS), dtype=_DTYPE)
-    view = memoryview(tensor.numpy())
+    if primary_kv_view is None:
+        tensor = torch.zeros((num_blocks, _BLOCK_ELEMENTS), dtype=_DTYPE)
+        primary_kv_view = memoryview(tensor.numpy())
     with (
         patch("vllm.v1.kv_offload.tiering.obj.manager.nixl_agent_config"),
         patch(
@@ -224,7 +241,7 @@ def _make_tier(
     ):
         tier = ObjectStoreSecondaryTierManager(
             offloading_spec=offloading_spec,
-            primary_kv_view=view,
+            primary_kv_view=primary_kv_view,
             tier_type="obj",
             store_config=_STORE_CONFIG,
             prefix=_RUN_PREFIX,
@@ -438,6 +455,105 @@ class TestMockObjTierFailures:
         assert not by_id[1].success
         assert by_id[2].success
 
+    def test_release_xfer_failure_retries_without_losing_result(self, monkeypatch):
+        tier, agent = _make_tier(num_blocks=4)
+        agent.check_xfer_state = MagicMock(side_effect=RuntimeError("poll failed"))
+        release_xfer = MagicMock(
+            side_effect=[RuntimeError("transfer is still active"), None]
+        )
+        monkeypatch.setattr(agent, "release_xfer_handle", release_xfer)
+
+        tier.submit_store(make_job(1, [key(1)], [0]))
+
+        # The transfer handle could not be released safely, so the job must
+        # remain tracked and must not be finalized yet.
+        assert list(tier.get_finished_jobs()) == []
+        assert 1 in tier._transfers
+
+        # Cleanup is retried without polling again or changing the failure
+        # verdict. The completion is then returned exactly once.
+        results = list(tier.get_finished_jobs())
+        assert len(results) == 1
+        assert results[0].job_id == 1
+        assert not results[0].success
+        assert agent.check_xfer_state.call_count == 2
+        assert release_xfer.call_count == 2
+        assert not tier._transfers
+        assert list(tier.get_finished_jobs()) == []
+
+    @pytest.mark.parametrize(
+        "cleanup_method", ["release_dlist_handle", "deregister_memory"]
+    )
+    def test_post_transfer_cleanup_failure_does_not_lose_result(
+        self, monkeypatch, cleanup_method
+    ):
+        tier, agent = _make_tier(num_blocks=4)
+        monkeypatch.setattr(
+            agent,
+            cleanup_method,
+            MagicMock(side_effect=RuntimeError("cleanup failed")),
+        )
+
+        tier.submit_store(make_job(1, [key(1)], [0]))
+        results = list(tier.get_finished_jobs())
+
+        assert len(results) == 1
+        assert results[0].job_id == 1
+        assert results[0].success
+        assert not tier._transfers
+        assert list(tier.get_finished_jobs()) == []
+
+    def test_xfer_cleanup_retry_finalizes_parent_job_and_primary_pin(self, monkeypatch):
+        num_blocks = 4
+        tensor = torch.zeros((num_blocks, _BLOCK_ELEMENTS), dtype=_DTYPE)
+        primary_kv_view = memoryview(tensor.numpy())
+        mmap_region = MagicMock()
+        mmap_region.create_kv_memoryview.return_value = primary_kv_view
+        primary_tier = CPUPrimaryTierOffloadingManager(
+            num_blocks=num_blocks, mmap_region=mmap_region
+        )
+        obj_tier, agent = _make_tier(
+            num_blocks=num_blocks, primary_kv_view=primary_kv_view
+        )
+        manager = TieringOffloadingManager(
+            primary_tier=primary_tier, secondary_tiers=[obj_tier]
+        )
+
+        keys = [key(1)]
+        primary_result = primary_tier.prepare_store(keys, _CTX)
+        assert primary_result is not None
+        primary_tier.complete_store(keys, _CTX, success=True)
+        job = manager.create_store_job(keys, _CTX)
+        obj_tier.submit_store(job)
+
+        block = primary_tier._policy.get(keys[0])
+        assert block is not None
+        assert block.ref_cnt == 1
+        assert len(manager._transfer_jobs) == 1
+
+        agent.check_xfer_state = MagicMock(side_effect=RuntimeError("poll failed"))
+        release_xfer = MagicMock(
+            side_effect=[RuntimeError("transfer is still active"), None]
+        )
+        monkeypatch.setattr(agent, "release_xfer_handle", release_xfer)
+        schedule_context = ScheduleEndContext(new_req_ids=[], preempted_req_ids=())
+
+        manager.on_schedule_end(schedule_context)
+
+        assert len(manager._transfer_jobs) == 1
+        assert block.ref_cnt == 1
+        assert len(obj_tier._transfers) == 1
+        assert manager.has_pending_work()
+
+        manager.on_schedule_end(schedule_context)
+
+        assert manager._transfer_jobs == {}
+        assert block.ref_cnt == 0
+        assert obj_tier._transfers == {}
+        assert not manager.has_pending_work()
+        assert agent.check_xfer_state.call_count == 2
+        assert release_xfer.call_count == 2
+
 
 class TestMockObjTierShutdown:
     def test_shutdown_clears_in_flight_transfers(self):
@@ -617,3 +733,37 @@ class TestObjStoreConfig:
         params = cfg.to_nixl_params()
         assert params["ca_bundle"] == "/path/to/ca.pem"
         assert "access_key" not in params
+
+
+def test_obj_tier_replicated_layout_collapses_mapper_identity():
+    """TP=2 and TP=4 replicated configs share the obj FileMapper namespace."""
+    tp2_spec = SimpleNamespace(
+        config=_make_offloading_config(
+            False, tp_size=2, world_size=2, rank=1, replicated_layout=True
+        ),
+        kv_events_config=OffloadingKVEventsConfig(
+            enable_kv_cache_events=False,
+            self_describing_kv_events=False,
+        ),
+    )
+    tp4_spec = SimpleNamespace(
+        config=_make_offloading_config(
+            False, tp_size=4, world_size=4, rank=3, replicated_layout=True
+        ),
+        kv_events_config=OffloadingKVEventsConfig(
+            enable_kv_cache_events=False,
+            self_describing_kv_events=False,
+        ),
+    )
+    tp2_tier, _ = _make_tier(offloading_spec=tp2_spec)
+    tp4_tier, _ = _make_tier(offloading_spec=tp4_spec)
+    try:
+        assert tp2_tier._file_mapper.base_path == tp4_tier._file_mapper.base_path
+        assert tp2_tier._file_mapper.rank == 0
+        assert tp4_tier._file_mapper.rank == 0
+        assert tp2_tier._file_mapper.get_run_config() == (
+            tp4_tier._file_mapper.get_run_config()
+        )
+    finally:
+        tp2_tier.shutdown()
+        tp4_tier.shutdown()
diff --git a/tests/v1/logits_processors/utils.py b/tests/v1/logits_processors/utils.py
index fc8ce50c05f..f57ea285eb7 100644
--- a/tests/v1/logits_processors/utils.py
+++ b/tests/v1/logits_processors/utils.py
@@ -11,6 +11,7 @@ import torch
 
 from tests.utils import requires_spawn_multiprocessing
 from vllm.config import VllmConfig
+from vllm.exceptions import VLLMValidationError
 from vllm.logger import init_logger
 from vllm.sampling_params import SamplingParams
 from vllm.v1.sample.logits_processor import (
@@ -61,7 +62,7 @@ class DummyLogitsProcessor(LogitsProcessor):
             "target_token"
         )
         if target_token is not None and not isinstance(target_token, int):
-            raise ValueError(
+            raise VLLMValidationError(
                 f"target_token value {target_token} {type(target_token)} is not int"
             )
 
diff --git a/tests/v1/sample/test_logprobs.py b/tests/v1/sample/test_logprobs.py
index aa17d2a1004..d643b0b6fde 100644
--- a/tests/v1/sample/test_logprobs.py
+++ b/tests/v1/sample/test_logprobs.py
@@ -405,7 +405,7 @@ def test_max_logprobs():
         runner.generate(["Hello world"], sampling_params=vllm_sampling_params)
 
         bad_sampling_params = SamplingParams(logprobs=2)
-        with pytest.raises(ValueError):
+        with pytest.raises(VLLMValidationError):
             runner.generate(["Hello world"], sampling_params=bad_sampling_params)
 
 
diff --git a/tests/v1/sample/test_sampling_params_e2e.py b/tests/v1/sample/test_sampling_params_e2e.py
index 56b93ea1e01..d385e96b7a2 100644
--- a/tests/v1/sample/test_sampling_params_e2e.py
+++ b/tests/v1/sample/test_sampling_params_e2e.py
@@ -4,6 +4,7 @@
 import pytest
 
 from vllm import LLM, SamplingParams
+from vllm.exceptions import VLLMValidationError
 
 MODEL = "hmellor/tiny-random-LlamaForCausalLM"
 PROMPT = "Hello my name is Robert and I"
@@ -161,15 +162,15 @@ def test_allowed_token_ids(llm):
         assert output[0].outputs[0].token_ids[-1] == token_id
 
     # Reject empty allowed_token_ids.
-    with pytest.raises(ValueError):
+    with pytest.raises(VLLMValidationError):
         _ = llm.generate(PROMPT, SamplingParams(allowed_token_ids=[]))
 
     # Reject negative token id.
-    with pytest.raises(ValueError):
+    with pytest.raises(VLLMValidationError):
         _ = llm.generate(PROMPT, SamplingParams(allowed_token_ids=[-1]))
 
     # Reject out of vocabulary.
-    with pytest.raises(ValueError):
+    with pytest.raises(VLLMValidationError):
         _ = llm.generate(PROMPT, SamplingParams(allowed_token_ids=[10000000]))
 
 
diff --git a/tests/v1/structured_output/test_validation.py b/tests/v1/structured_output/test_validation.py
index 31ce961ff61..7ea60cc6609 100644
--- a/tests/v1/structured_output/test_validation.py
+++ b/tests/v1/structured_output/test_validation.py
@@ -5,6 +5,7 @@
 import pytest
 
 from vllm.config import StructuredOutputsConfig
+from vllm.exceptions import VLLMValidationError
 from vllm.sampling_params import SamplingParams, StructuredOutputsParams
 
 pytestmark = pytest.mark.cpu_test
@@ -32,7 +33,7 @@ def test_structured_outputs_rejected_for_diffusion_models():
     params = SamplingParams(
         structured_outputs=StructuredOutputsParams(json=JSON_SCHEMA)
     )
-    with pytest.raises(ValueError, match="not yet supported for diffusion"):
+    with pytest.raises(VLLMValidationError, match="not yet supported for diffusion"):
         params._validate_structured_outputs(
             _StubModelConfig(is_diffusion=True),
             StructuredOutputsConfig(),
@@ -63,7 +64,7 @@ def test_degenerate_structured_outputs_rejected(structured_outputs, match):
     rejected at request validation (-> 400) instead of reaching and crashing
     the engine."""
     params = SamplingParams(structured_outputs=structured_outputs)
-    with pytest.raises(ValueError, match=match):
+    with pytest.raises(VLLMValidationError, match=match):
         params._validate_structured_outputs(
             _StubModelConfig(is_diffusion=False),
             StructuredOutputsConfig(),
diff --git a/tests/v1/test_serial_utils.py b/tests/v1/test_serial_utils.py
index 4ed8724e60f..9f7761c22b5 100644
--- a/tests/v1/test_serial_utils.py
+++ b/tests/v1/test_serial_utils.py
@@ -423,3 +423,139 @@ def test_multiple_senders_single_receiver_ipc():
         assert torch.allclose(decoded.prompt_embeds, original_tensor), (
             f"Value mismatch for sender {sender_idx} msg {msg_idx}"
         )
+
+
+def _logprobs_outputs(num_reqs: int, num_prompt_tokens: int):
+    """An EngineCoreOutputs carrying prompt logprobs, as the engine core sends
+    it: many requests, each with per-token tensors small enough that pyzmq
+    copies their frames, while the accumulated payload frame is large enough
+    that pyzmq sends it zero-copy."""
+    from vllm.v1.engine import EngineCoreOutput, EngineCoreOutputs
+    from vllm.v1.outputs import LogprobsTensors
+
+    outputs = []
+    for req in range(num_reqs):
+        num_tokens = num_prompt_tokens + req % 4
+        outputs.append(
+            EngineCoreOutput(
+                request_id=f"req-{req:08d}",
+                new_token_ids=[req],
+                new_prompt_logprobs_tensors=LogprobsTensors(
+                    logprob_token_ids=torch.arange(
+                        num_tokens * 2, dtype=torch.int64
+                    ).view(num_tokens, 2),
+                    logprobs=torch.zeros(num_tokens, 2, dtype=torch.float32),
+                    selected_token_ranks=torch.zeros(num_tokens, dtype=torch.int32),
+                ),
+            )
+        )
+    return EngineCoreOutputs(outputs=outputs)
+
+
+def test_payload_buffer_reuse_does_not_corrupt_in_flight_messages():
+    """The engine core recycles the msgpack payload buffer across messages
+    (`MsgpackEncoder.encode_into`). It may only do so once zmq has finished
+    sending that buffer, otherwise a newer payload is delivered alongside the
+    older message's zero-copy tensor frames.
+
+    `Socket.send_multipart(track=True)` cannot be used to detect this: it
+    returns a tracker for the last frame only, and pyzmq copies frames below
+    `zmq.COPY_THRESHOLD` and reports them as already-sent.
+    """
+    import zmq
+
+    from vllm.v1.engine import EngineCoreOutputs
+    from vllm.v1.engine.core import EngineCoreProc
+
+    num_msgs = 100
+    encoder = MsgpackEncoder()
+    decoder = MsgpackDecoder(EngineCoreOutputs)
+    # Enough requests that the payload frame is zero-copied rather than copied
+    # by pyzmq, which is what makes early reuse observable.
+    messages = [_logprobs_outputs(300, 24 + i % 8) for i in range(num_msgs)]
+    assert len(encoder.encode(messages[0])[0]) >= zmq.COPY_THRESHOLD
+
+    reuse_buffers: list[bytearray] = []
+    pending: list[tuple[zmq.MessageTracker, bytearray]] = []
+    with zmq.Context() as ctx:
+        push = ctx.socket(zmq.PUSH)
+        push.bind("inproc://test-payload-reuse")
+        pull = ctx.socket(zmq.PULL)
+        pull.connect("inproc://test-payload-reuse")
+
+        for outputs in messages:
+            while pending and pending[0][0].done:
+                reuse_buffers.append(pending.pop(0)[1])
+            buffer = reuse_buffers.pop() if reuse_buffers else bytearray()
+            buffers = encoder.encode_into(outputs, buffer)
+            tracker = EngineCoreProc._send_msg_tracking_payload(push, buffers)
+            if tracker.done:
+                reuse_buffers.append(buffer)
+            else:
+                pending.append((tracker, buffer))
+
+        for i, sent in enumerate(messages):
+            received = decoder.decode(pull.recv_multipart(copy=False))
+            assert len(received.outputs) == len(sent.outputs), f"message {i}"
+            for expected, actual in zip(sent.outputs, received.outputs):
+                sent_ids = expected.new_prompt_logprobs_tensors.logprob_token_ids
+                got_ids = actual.new_prompt_logprobs_tensors.logprob_token_ids
+                assert actual.request_id == expected.request_id, f"message {i}"
+                assert torch.equal(got_ids, sent_ids), (
+                    f"message {i} request {actual.request_id}: corrupted "
+                    f"prompt logprobs, {got_ids.shape} vs {sent_ids.shape}"
+                )
+        push.close(linger=0)
+        pull.close(linger=0)
+
+
+def test_zero_copy_frames_survive_without_caller_side_references():
+    """Callers don't need to retain the encoded object until zmq has sent it:
+    for a zero-copy frame, zmq holds its own reference to the backing buffer.
+
+    The engine core clients rely on this when sending requests that carry
+    tensors (e.g. prompt embeds) without tracking the messages.
+
+    What makes that safe is that `tensor_data()` hands zmq a memoryview which
+    transitively references the source tensor, so refcounting - not timing -
+    keeps the memory from being freed and reused underneath zmq.
+    """
+    import gc
+
+    import zmq
+
+    from vllm.v1.utils import tensor_data
+
+    num_elems = 100_000  # comfortably over zmq.COPY_THRESHOLD
+    expected = torch.arange(num_elems, dtype=torch.int64)
+    encoder = MsgpackEncoder()
+    decoder = MsgpackDecoder(RequestWithTensor)
+
+    # The buffer handed to zmq must keep the tensor's storage alive by itself.
+    holder = tensor_data(expected).obj
+    while getattr(holder, "base", None) is not None:
+        holder = holder.base
+    assert isinstance(holder, torch.Tensor)
+    assert holder.data_ptr() == expected.data_ptr()
+
+    with zmq.Context() as ctx:
+        push = ctx.socket(zmq.PUSH)
+        push.bind("inproc://test-zero-copy-lifetime")
+        pull = ctx.socket(zmq.PULL)
+        pull.connect("inproc://test-zero-copy-lifetime")
+
+        request = RequestWithTensor(prompt_embeds=expected.clone(), data="req")
+        buffers = encoder.encode(request)
+        assert max(len(buf) for buf in buffers) >= zmq.COPY_THRESHOLD
+        push.send_multipart(buffers, copy=False)
+
+        # Drop every reference the sender holds, then churn the allocator.
+        del request, buffers
+        gc.collect()
+        torch.arange(num_elems * 4, dtype=torch.int64)
+
+        decoded = decoder.decode(pull.recv_multipart(copy=False))
+        assert decoded.prompt_embeds is not None
+        assert torch.equal(decoded.prompt_embeds, expected)
+        push.close(linger=0)
+        pull.close(linger=0)
diff --git a/tests/v1/worker/test_gpu_model_runner.py b/tests/v1/worker/test_gpu_model_runner.py
index 79e3a60e981..d2d89a74fce 100644
--- a/tests/v1/worker/test_gpu_model_runner.py
+++ b/tests/v1/worker/test_gpu_model_runner.py
@@ -1669,3 +1669,59 @@ def test_mamba_cache_raises_when_max_num_seqs_exceeds_blocks():
 
         with pytest.raises(ValueError, match="max_num_seqs"):
             runner.initialize_kv_cache(kv_cache_config)
+
+
+class TestInitFp8KvScalesHybridModels:
+    """Verify init_fp8_kv_scales handles heterogeneous kv_caches entries.
+
+    Hybrid models (Mamba, DeltaNet) store per-layer state as a list of tensors
+    rather than a single tensor. init_fp8_kv_scales must iterate both forms.
+    """
+
+    @staticmethod
+    def _make_runner_stub(kv_caches):
+        runner = Mock(spec=GPUModelRunner)
+        runner.cache_config = SimpleNamespace(cache_dtype="fp8_e4m3")
+        runner.kv_caches = kv_caches
+        runner.compilation_config = SimpleNamespace(static_forward_context={})
+        runner.init_fp8_kv_scales = GPUModelRunner.init_fp8_kv_scales.__get__(
+            runner, GPUModelRunner
+        )
+        return runner
+
+    def test_zeroes_both_tensor_and_list_entries(self):
+        single_tensor = torch.ones(4, 8)
+        list_tensors = [torch.ones(2, 4), torch.ones(3, 6)]
+
+        runner = self._make_runner_stub([single_tensor, list_tensors])
+        runner.init_fp8_kv_scales()
+
+        assert (single_tensor == 0).all()
+        assert all((t == 0).all() for t in list_tensors)
+
+    def test_skips_none_entries(self):
+        tensor = torch.ones(4, 8)
+        runner = self._make_runner_stub([None, tensor, None])
+        runner.init_fp8_kv_scales()
+
+        assert (tensor == 0).all()
+
+    def test_noop_when_kv_cache_not_quantized(self):
+        tensor = torch.ones(4, 8)
+        runner = self._make_runner_stub([tensor])
+        runner.cache_config.cache_dtype = "auto"
+        runner.init_fp8_kv_scales()
+
+        assert (tensor == 1).all()
+
+    def test_mixed_none_tensor_and_list(self):
+        t1 = torch.ones(2, 2)
+        t2 = torch.ones(3, 3)
+        list_entry = [torch.ones(1, 1), torch.ones(1, 1)]
+
+        runner = self._make_runner_stub([None, t1, list_entry, None, t2])
+        runner.init_fp8_kv_scales()
+
+        assert (t1 == 0).all()
+        assert (t2 == 0).all()
+        assert all((t == 0).all() for t in list_entry)
diff --git a/tests/v1/worker/test_kv_block_zeroer.py b/tests/v1/worker/test_kv_block_zeroer.py
index b212e3ae17b..17aa1bf38d4 100644
--- a/tests/v1/worker/test_kv_block_zeroer.py
+++ b/tests/v1/worker/test_kv_block_zeroer.py
@@ -4,7 +4,7 @@
 import pytest
 import torch
 
-from vllm.v1.worker.utils import KVBlockZeroer
+from vllm.v1.worker.utils import KVBlockZeroer, _zero_kv_blocks_kernel
 
 
 @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
@@ -86,3 +86,67 @@ def test_non_uniform_page_sizes():
         assert torch.all(storage[1] == 0)
         assert torch.all(storage[2] == 0)
         assert torch.all(storage[3] == 1)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+def test_warmup_compiles_every_n_blocks_specialization():
+    """After warmup, no launch should trigger a first-request JIT compile.
+
+    ``n_blocks`` is ``do_not_specialize``, so a single warmup launch must
+    cover every block count.
+    """
+    device = torch.device("cuda")
+    num_blocks = 64
+    page_size_el = 4
+    storage = torch.ones((num_blocks, page_size_el), dtype=torch.int32, device=device)
+
+    zeroer = KVBlockZeroer.__new__(KVBlockZeroer)
+    zeroer.device = device
+    zeroer._meta = (
+        torch.tensor([storage.data_ptr()], dtype=torch.uint64, device=device),
+        torch.tensor([page_size_el], dtype=torch.int64, device=device),
+        1,  # max_chunks
+        page_size_el,  # blk_size
+        1,  # n_segs
+    )
+
+    def compiled_variants() -> set:
+        return {
+            key
+            for caches in _zero_kv_blocks_kernel.device_caches.values()
+            for key in caches[0]
+        }
+
+    zeroer.warmup(num_blocks)
+    torch.accelerator.synchronize()
+    warmed = compiled_variants()
+    assert warmed
+
+    for n_blocks in (1, 2, 3, 16, 32):
+        zeroer.zero_block_ids(list(range(n_blocks)))
+    torch.accelerator.synchronize()
+
+    assert compiled_variants() == warmed
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+def test_warmup_respects_available_block_count():
+    """An empty KV cache must not be warmed with out-of-range block IDs."""
+    device = torch.device("cuda")
+    page_size_el = 4
+    storage = torch.ones((1, page_size_el), dtype=torch.int32, device=device)
+
+    zeroer = KVBlockZeroer.__new__(KVBlockZeroer)
+    zeroer.device = device
+    zeroer._meta = (
+        torch.tensor([storage.data_ptr()], dtype=torch.uint64, device=device),
+        torch.tensor([page_size_el], dtype=torch.int64, device=device),
+        1,
+        page_size_el,
+        1,
+    )
+
+    zeroer.warmup(0)
+    torch.accelerator.synchronize()
+
+    assert torch.all(storage == 1)
diff --git a/tools/ep_kernels/install_python_libraries.sh b/tools/ep_kernels/install_python_libraries.sh
index 739f031c9ef..5f5a597baca 100755
--- a/tools/ep_kernels/install_python_libraries.sh
+++ b/tools/ep_kernels/install_python_libraries.sh
@@ -197,6 +197,21 @@ do_build() {
 #endif' csrc/kernels/backend/symmetric.hpp
     fi
 
+    if [[ "$name" == "DeepEP" ]]; then
+        # DeepEP links against the CUDA driver API in driverless build images.
+        local cuda_driver_stub
+        local cuda_driver_stub_dir
+        cuda_driver_stub=$(
+            find -H "$CUDA_HOME" -path "*/stubs/libcuda.so" -print -quit
+        )
+        if [[ -z "$cuda_driver_stub" ]]; then
+            echo "CUDA driver stub not found under $CUDA_HOME" >&2
+            exit 1
+        fi
+        cuda_driver_stub_dir=$(dirname "$cuda_driver_stub")
+        export LIBRARY_PATH="${cuda_driver_stub_dir}${LIBRARY_PATH:+:$LIBRARY_PATH}"
+    fi
+
     if [ "$MODE" = "install" ]; then
         echo "Installing $name into environment"
         eval "$extra_env" uv pip install --no-build-isolation -vvv .
diff --git a/tools/pre_commit/check_forbidden_imports.py b/tools/pre_commit/check_forbidden_imports.py
index a788cecc6ce..52a95ce1d8d 100644
--- a/tools/pre_commit/check_forbidden_imports.py
+++ b/tools/pre_commit/check_forbidden_imports.py
@@ -48,6 +48,7 @@ CHECK_IMPORTS = {
             "vllm/distributed/device_communicators/shm_object_storage.py",
             "vllm/distributed/weight_transfer/ipc_engine.py",
             "vllm/distributed/weight_transfer/clients.py",
+            "tests/distributed/test_shm_broadcast.py",
             "tests/distributed/test_weight_transfer.py",
             "vllm/utils/hashing.py",
             "tests/multimodal/media/test_base.py",
diff --git a/vllm/_custom_ops.py b/vllm/_custom_ops.py
index bfbf03ac814..bc6d6d4afa4 100644
--- a/vllm/_custom_ops.py
+++ b/vllm/_custom_ops.py
@@ -2779,33 +2779,6 @@ def fused_kda_decode(
     return out
 
 
-def kimi_k3_attn_res(
-    prefix: torch.Tensor,
-    delta: torch.Tensor,
-    blocks: torch.Tensor,
-    norm_weight: torch.Tensor,
-    qk_weight: torch.Tensor,
-    output_norm_weight: torch.Tensor,
-    num_blocks: int,
-    eps: float,
-    output_norm_eps: float,
-) -> torch.Tensor:
-    output = torch.empty_like(prefix)
-    torch.ops._C.kimi_k3_attn_res(
-        prefix,
-        delta,
-        blocks,
-        norm_weight,
-        qk_weight,
-        output_norm_weight,
-        output,
-        num_blocks,
-        eps,
-        output_norm_eps,
-    )
-    return output
-
-
 def concat_and_cache_mla(
     kv_c: torch.Tensor,
     k_pe: torch.Tensor,
@@ -2839,6 +2812,33 @@ def concat_and_cache_mla_grouped(
     )
 
 
+def kimi_k3_attn_res(
+    prefix: torch.Tensor,
+    delta: torch.Tensor,
+    blocks: torch.Tensor,
+    norm_weight: torch.Tensor,
+    qk_weight: torch.Tensor,
+    output_norm_weight: torch.Tensor,
+    num_blocks: int,
+    eps: float,
+    output_norm_eps: float,
+) -> torch.Tensor:
+    output = torch.empty_like(prefix)
+    torch.ops._C.kimi_k3_attn_res(
+        prefix,
+        delta,
+        blocks,
+        norm_weight,
+        qk_weight,
+        output_norm_weight,
+        output,
+        num_blocks,
+        eps,
+        output_norm_eps,
+    )
+    return output
+
+
 def concat_and_cache_mla_rope_fused(
     positions: torch.Tensor,
     q_pe: torch.Tensor,
@@ -4117,9 +4117,40 @@ def fusedQuantizeNv(
         padded_rows, padded_cols, dtype=torch.float8_e4m3fn, device=a.device
     )
 
-    return torch.ops._qutlass_C.fusedQuantizeNvAbsMax(
-        a, b, xh_e2m1, xh_e4m3, global_scale
-    )
+    safeFusedQuantizeNv(a, b, xh_e2m1, xh_e4m3, global_scale)
+    return xh_e2m1, xh_e4m3
+
+
+@torch.library.custom_op(
+    "vllm::safeFusedQuantizeNv", mutates_args=("xh_e2m1", "xh_e4m3")
+)
+def safeFusedQuantizeNv(
+    a: torch.Tensor,
+    b: torch.Tensor,
+    xh_e2m1: torch.Tensor,
+    xh_e4m3: torch.Tensor,
+    global_scale: torch.Tensor,
+) -> None:
+    """
+    Wrapper for QUTLASS fusedQuantizeNv method that operates on tensors in-place
+    rather than returning them, to prevent torch 2.12+ errors that outputs of custom
+    operators may not alias any inputs to the custom operator.
+    """
+    torch.ops._qutlass_C.fusedQuantizeNvAbsMax(a, b, xh_e2m1, xh_e4m3, global_scale)
+    return
+
+
+if hasattr(torch.ops._qutlass_C, "fusedQuantizeNv"):
+
+    @register_fake("vllm::safeFusedQuantizeNv")
+    def _fake_fused_quantize_nv(
+        a: torch.Tensor,
+        b: torch.Tensor,
+        xh_e2m1: torch.Tensor,
+        xh_e4m3: torch.Tensor,
+        global_scale: torch.Tensor,
+    ) -> None:
+        return
 
 
 def hadacore_transform(x: torch.Tensor, inplace: bool = True) -> torch.Tensor:
diff --git a/vllm/compilation/passes/fusion/allreduce_rms_fusion.py b/vllm/compilation/passes/fusion/allreduce_rms_fusion.py
index f7ff7df66cd..1722b524eeb 100644
--- a/vllm/compilation/passes/fusion/allreduce_rms_fusion.py
+++ b/vllm/compilation/passes/fusion/allreduce_rms_fusion.py
@@ -1416,8 +1416,7 @@ class AiterAllreduceFusedAddRMSNormGroupQuantWithIndexerPattern(
     The trailing FP8 group-quant is matched via ``MatcherQuantFP8`` (consistent
     with the sibling patterns above), which traces both ``QuantFP8.forward_hip``
     and ``forward_native`` paths and so matches whichever op the call site
-    lowers to (``vllm.triton_per_token_group_quant_fp8`` or
-    ``vllm.rocm_aiter_group_fp8_quant``).
+    lowers to (``vllm.rocm_aiter_group_fp8_quant``).
     """
 
     def __init__(
diff --git a/vllm/config/parallel.py b/vllm/config/parallel.py
index 814e73eac22..2e908ead176 100644
--- a/vllm/config/parallel.py
+++ b/vllm/config/parallel.py
@@ -686,6 +686,14 @@ class ParallelConfig:
             and self.data_parallel_size > 1
         )
 
+    @property
+    def use_all2all(self) -> bool:
+        return (
+            self.data_parallel_size > 1
+            or self.use_sequence_parallel_moe
+            or (self.enable_expert_parallel and self.prefill_context_parallel_size > 1)
+        )
+
     @property
     def use_batched_dp_moe(self) -> bool:
         return (
@@ -786,6 +794,7 @@ class ParallelConfig:
             "data_parallel_master_ip",
             "data_parallel_master_port",
             "_data_parallel_master_port_list",
+            "_coord_store_port",
             "data_parallel_rpc_port",
             "rank",
             "master_addr",
diff --git a/vllm/distributed/device_communicators/base_device_communicator.py b/vllm/distributed/device_communicators/base_device_communicator.py
index 45438a54691..73fd1331f5c 100644
--- a/vllm/distributed/device_communicators/base_device_communicator.py
+++ b/vllm/distributed/device_communicators/base_device_communicator.py
@@ -175,6 +175,7 @@ class DeviceCommunicatorBase:
         unique_name: str = "",
         global_ranks: list[int] | None = None,
         global_world_size: int | None = None,
+        use_all2all: bool = False,
     ):
         self.device = device or torch.device("cpu")
         self.cpu_group = cpu_group
@@ -204,26 +205,15 @@ class DeviceCommunicatorBase:
             self.global_world_size = dist.get_world_size()
             self.rank_in_group = dist.get_group_rank(self.cpu_group, self.global_rank)
 
-        use_ep = False
         all2all_backend = None
         from vllm.config import get_current_vllm_config_or_none
 
         config = get_current_vllm_config_or_none()
         if config is not None:
-            # initialize the all2all manager for DP or sequence-parallel EP.
-            parallel_config = config.parallel_config
-            use_ep = (
-                parallel_config.data_parallel_size > 1
-                or parallel_config.use_sequence_parallel_moe
-                or (
-                    parallel_config.enable_expert_parallel
-                    and parallel_config.prefill_context_parallel_size > 1
-                )
-            )
-            all2all_backend = parallel_config.all2all_backend
+            all2all_backend = config.parallel_config.all2all_backend
 
         self.is_ep_communicator = unique_name.split(":")[0] == "ep"
-        self.use_all2all = self.is_ep_communicator and use_ep
+        self.use_all2all = self.is_ep_communicator and use_all2all
         self.all2all_backend = all2all_backend
         self.all2all_manager: All2AllManagerBase | None = None
 
diff --git a/vllm/distributed/device_communicators/cpu_communicator.py b/vllm/distributed/device_communicators/cpu_communicator.py
index 9ec4b72f80d..8ea12d9255a 100644
--- a/vllm/distributed/device_communicators/cpu_communicator.py
+++ b/vllm/distributed/device_communicators/cpu_communicator.py
@@ -24,8 +24,11 @@ class CpuCommunicator(DeviceCommunicatorBase):
         device: torch.device | None = None,
         device_group: ProcessGroup | None = None,
         unique_name: str = "",
+        use_all2all: bool = False,
     ):
-        super().__init__(cpu_group, device, device_group, unique_name)
+        super().__init__(
+            cpu_group, device, device_group, unique_name, use_all2all=use_all2all
+        )
         self.dist_module = torch.distributed
 
         if (
diff --git a/vllm/distributed/device_communicators/cuda_communicator.py b/vllm/distributed/device_communicators/cuda_communicator.py
index 804bf84d742..06b441c5a41 100644
--- a/vllm/distributed/device_communicators/cuda_communicator.py
+++ b/vllm/distributed/device_communicators/cuda_communicator.py
@@ -36,6 +36,7 @@ class CudaCommunicator(DeviceCommunicatorBase):
         global_ranks: list[int] | None = None,
         global_world_size: int | None = None,
         tcp_store_group: StatelessProcessGroup | None = None,
+        use_all2all: bool = False,
     ):
         super().__init__(
             cpu_group,
@@ -44,6 +45,7 @@ class CudaCommunicator(DeviceCommunicatorBase):
             unique_name,
             global_ranks,
             global_world_size,
+            use_all2all=use_all2all,
         )
         if "tp" not in unique_name:
             # custom allreduce or torch symm mem can be used only by tp
diff --git a/vllm/distributed/device_communicators/shm_broadcast.py b/vllm/distributed/device_communicators/shm_broadcast.py
index afabdf18c80..e59b14f7a6b 100644
--- a/vllm/distributed/device_communicators/shm_broadcast.py
+++ b/vllm/distributed/device_communicators/shm_broadcast.py
@@ -1,6 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import copyreg
 import functools
+import io
 import os
 import pickle
 import shutil
@@ -385,6 +387,70 @@ class ShmRingBuffer:
             yield buf
 
 
+def _rebuild_tensor(buf: Any, shape: tuple[int, ...], dtype_str: str) -> torch.Tensor:
+    """Rebuild a tensor from an out-of-band pickle buffer.
+
+    Counterpart of `_reduce_tensor`. Note that pickle passes the original
+    buffer-providing object from `loads(buffers=...)` straight to this
+    function (no `PickleBuffer` wrapper on the receiving side), so `buf` is
+    a `zmq.Frame`, a `memoryview` of a shared-memory ring chunk, or `bytes`
+    if the buffer was serialized in-band.
+    """
+    dtype = getattr(torch, dtype_str)
+    assert isinstance(dtype, torch.dtype)
+    if isinstance(buf, zmq.Frame):
+        # ZMQ frames own their message memory independently of any context,
+        # so the tensor can safely alias it with zero copies. The tensor's
+        # storage keeps the frame (and thus its bytes) alive via a strong
+        # reference for as long as the tensor is.
+        try:
+            return torch.frombuffer(buf, dtype=torch.uint8).view(dtype).view(shape)
+        except ValueError:
+            # Empty or read-only frame buffer; fall through to the copy path.
+            pass
+    # Shared-memory ring buffer chunks are reused by the writer once all
+    # readers have marked them read, so we must copy out of them. bytearray
+    # (vs bytes) keeps the resulting tensor writable, matching normal tensor
+    # semantics.
+    raw = bytearray(buf)
+    if not raw:
+        assert 0 in shape
+        return torch.empty(shape, dtype=dtype)
+    return torch.frombuffer(raw, dtype=torch.uint8).view(dtype).view(shape)
+
+
+def _reduce_tensor(tensor: torch.Tensor):
+    """Reduce a CPU tensor to a `PickleBuffer` for out-of-band pickling.
+
+    `torch.Tensor.__reduce_ex__` copies the tensor bytes into the pickle
+    byte stream via `torch.serialization` and never emits a `PickleBuffer`,
+    which defeats the out-of-band buffer handling in `MessageQueue.enqueue`.
+    This reducer instead exposes the tensor's memory directly, so large
+    tensors (e.g. `prompt_embeds` in `SchedulerOutput`) traverse the queue
+    without being copied into and back out of the pickled message.
+    """
+    if (
+        tensor.device.type == "cpu"
+        and tensor.layout == torch.strided
+        and not tensor.requires_grad
+    ):
+        try:
+            # The uint8 view exposes the raw bytes via the buffer protocol,
+            # including for dtypes numpy doesn't recognize (bfloat16, fp8, ...).
+            # reshape(-1) first so that 0-dim tensors can be viewed as well.
+            raw = tensor.contiguous().reshape(-1).view(torch.uint8).numpy()
+        except RuntimeError:
+            # Exotic tensors (e.g. with the conjugate bit set) that don't
+            # support aliasing views; let torch handle them.
+            pass
+        else:
+            dtype_str = str(tensor.dtype).removeprefix("torch.")
+            return _rebuild_tensor, (PickleBuffer(raw), tuple(tensor.shape), dtype_str)
+
+    # Fall back to torch's default (copying) reduction.
+    return tensor.__reduce_ex__(pickle.HIGHEST_PROTOCOL)
+
+
 @dataclass
 class Handle:
     local_reader_ranks: list[int] = field(default_factory=list)
@@ -771,9 +837,23 @@ class MessageQueue:
             total_bytes += len(raw_buf) + 4
             return False
 
-        all_buffers[0] = pickle.dumps(
-            obj, protocol=pickle.HIGHEST_PROTOCOL, buffer_callback=oob_callback
-        )
+        # CPU tensors are routed through `_reduce_tensor` so that their
+        # bytes are emitted as out-of-band buffers instead of being
+        # copied into the pickle stream by torch's default reducer.
+        # Start from `copyreg.dispatch_table` to preserve globally
+        # registered reducers (e.g. `re.Pattern`); the per-pickler
+        # dispatch table would otherwise shadow them.
+        dispatch_table = dict(copyreg.dispatch_table)
+        dispatch_table[torch.Tensor] = _reduce_tensor
+        with io.BytesIO() as bio:
+            pickler = pickle.Pickler(
+                bio,
+                protocol=pickle.HIGHEST_PROTOCOL,
+                buffer_callback=oob_callback,
+            )
+            pickler.dispatch_table = dispatch_table
+            pickler.dump(obj)
+            all_buffers[0] = bio.getvalue()
         if self.n_local_reader > 0:
             if total_bytes + len(all_buffers[0]) >= self.buffer.max_chunk_bytes:
                 with self.acquire_write(timeout) as buf:
diff --git a/vllm/distributed/device_communicators/xpu_communicator.py b/vllm/distributed/device_communicators/xpu_communicator.py
index 1b6ce9e8aae..7ca132824ec 100644
--- a/vllm/distributed/device_communicators/xpu_communicator.py
+++ b/vllm/distributed/device_communicators/xpu_communicator.py
@@ -20,8 +20,11 @@ class XpuCommunicator(DeviceCommunicatorBase):
         device: torch.device | None = None,
         device_group: ProcessGroup | None = None,
         unique_name: str = "",
+        use_all2all: bool = False,
     ):
-        super().__init__(cpu_group, device, device_group, unique_name)
+        super().__init__(
+            cpu_group, device, device_group, unique_name, use_all2all=use_all2all
+        )
         self.ca_comm: None = None
         if self.use_all2all:
             if self.all2all_backend in ("naive", "allgather_reducescatter"):
diff --git a/vllm/distributed/elastic_ep/elastic_execute.py b/vllm/distributed/elastic_ep/elastic_execute.py
index b0c3740f57e..cea7fcb2f01 100644
--- a/vllm/distributed/elastic_ep/elastic_execute.py
+++ b/vllm/distributed/elastic_ep/elastic_execute.py
@@ -1,9 +1,9 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import copy
 import gc
 import weakref
 from collections.abc import Iterable, Sequence
+from concurrent.futures import Future, ThreadPoolExecutor
 from dataclasses import replace
 from typing import TYPE_CHECKING
 
@@ -43,6 +43,7 @@ from vllm.model_executor.layers.fused_moe.config import FusedMoEParallelConfig
 from vllm.model_executor.layers.fused_moe.eep_reconfigure import (
     make_eep_staged_quant_method,
 )
+from vllm.model_executor.warmup.kernel_warmup import kernel_warmup
 from vllm.utils import is_moe_layer
 from vllm.v1.engine import ReconfigureDistributedRequest, ReconfigureRankType
 from vllm.v1.worker.gpu_ubatch_wrapper import UBatchWrapper
@@ -145,6 +146,10 @@ class ElasticEPScalingExecutor:
         self.worker_ref = weakref.ref(worker)
         self.reconfig_request = None
         self._staged_moe_quant_methods: dict[nn.Module, FusedMoEMethodBase] = {}
+        self._async_executor = ThreadPoolExecutor(
+            max_workers=1, thread_name_prefix="ElasticEPAsync"
+        )
+        self._async_future: Future[None] | None = None
 
     @property
     def worker(self):
@@ -159,59 +164,68 @@ class ElasticEPScalingExecutor:
             raise ValueError(f"Unknown execute method: {execute_method}")
         return method(*args, **kwargs)
 
-    def _set_eplb_suppressed(self, suppressed: bool) -> None:
-        self.worker.model_runner.eep_eplb_suppressed = suppressed
-        ep_group = get_standby_ep_group() or get_ep_group()
-        if ep_group.rank == 0:
-            logger.info(
-                "[Elastic EP] EPLB %s elastic scaling transition",
-                "disabled during" if suppressed else "re-enabled after",
-            )
+    def start_async(self, execute_method: str, *args, **kwargs) -> str:
+        if self._async_future is not None:
+            raise RuntimeError("Another Elastic EP async method is active")
+        if args and isinstance(args[0], ReconfigureDistributedRequest):
+            self.reconfig_request = args[0]
+        dp_rank = self.worker.vllm_config.parallel_config.data_parallel_rank
+        done_key = f"eep_async/{execute_method}/{dp_rank}/{self.worker.rank}"
+        self._async_future = self._async_executor.submit(
+            self._run_async, execute_method, *args, **kwargs
+        )
+        self._async_future.add_done_callback(lambda _: self._mark_async_done(done_key))
+        return done_key
+
+    def _run_async(self, execute_method: str, *args, **kwargs) -> None:
+        from vllm.platforms import current_platform
+
+        self.worker.vllm_config.enable_trace_function_call_for_thread()
+        assert hasattr(self.worker, "device")
+        current_platform.set_device(self.worker.device)
+        with set_current_vllm_config(self.worker.vllm_config):
+            self.execute(execute_method, *args, **kwargs)
+
+    def _mark_async_done(self, done_key: str) -> None:
+        from vllm.distributed.utils import get_cached_tcp_store_client
+
+        assert self.reconfig_request is not None
+        get_cached_tcp_store_client(
+            self.reconfig_request.new_data_parallel_master_ip,
+            self.reconfig_request.coord_store_port,
+        ).set(done_key, b"1")
+
+    def clear_async(self) -> None:
+        future = self._async_future
+        if future is None:
+            raise RuntimeError("No Elastic EP async method is active")
+        if not future.done():
+            raise RuntimeError("Elastic EP async method is not done")
+        self._async_future = None
+        future.result()
 
     def load_model(self) -> None:
-        (
-            expanded_physical_to_logical,
-            num_logical_experts,
-            old_num_physical_experts,
-        ) = self.receive_expert_mapping()
-        num_physical_experts = expanded_physical_to_logical.shape[1]
-        self.worker.parallel_config.eplb_config.num_redundant_experts = (
-            num_physical_experts - num_logical_experts
-        )
         self.worker.load_model(load_dummy_weights=True)
-        self.worker.model_runner.setup_eplb_from_mapping(
-            expanded_physical_to_logical, old_num_physical_experts
-        )
-        self._set_eplb_suppressed(True)
 
     def create_standby_groups(
-        self, reconfig_request: ReconfigureDistributedRequest
+        self, reconfig_request: ReconfigureDistributedRequest, use_all2all: bool
     ) -> None:
         self.reconfig_request = reconfig_request
         new_dp_size = reconfig_request.new_data_parallel_size
         old_dp_size = get_dp_group().world_size
-        world_size = self.worker.vllm_config.parallel_config.world_size
+        parallel_config = self.worker.vllm_config.parallel_config
+        world_size = parallel_config.world_size
         new_world_size_across_dp = world_size * new_dp_size
-        updated_config = copy.copy(self.worker.vllm_config)
-        updated_config.parallel_config = copy.deepcopy(
-            self.worker.vllm_config.parallel_config
+        create_standby_groups(
+            new_dp_size=new_dp_size,
+            new_world_size_across_dp=new_world_size_across_dp,
+            master_ip=reconfig_request.new_data_parallel_master_ip,
+            coord_store_port=reconfig_request.coord_store_port,
+            use_all2all=use_all2all,
+            enable_eplb=parallel_config.enable_eplb,
         )
-        updated_config.parallel_config.data_parallel_size = new_dp_size
-        with set_current_vllm_config(updated_config):
-            create_standby_groups(
-                new_dp_size=new_dp_size,
-                new_world_size_across_dp=new_world_size_across_dp,
-                master_ip=reconfig_request.new_data_parallel_master_ip,
-                coord_store_port=reconfig_request.coord_store_port,
-                enable_eplb=updated_config.parallel_config.enable_eplb,
-            )
-        if new_dp_size > old_dp_size:
-            self._set_eplb_suppressed(True)
-            eplb_state = self.worker.model_runner.eplb_state
-            if eplb_state is not None:
-                eplb_state.drain_async()
-        elif new_dp_size < old_dp_size:
-            self._stage_standby_moe_quant_methods()
+        if new_dp_size < old_dp_size:
+            self.stage_standby_moe_quant_methods()
 
     def transfer_weights(self, old_dp_size: int, new_dp_size: int) -> None:
         standby_dp_group = get_standby_dp_group()
@@ -265,6 +279,7 @@ class ElasticEPScalingExecutor:
         model_config = self.worker.model_runner.model_config
         eplb_state = self.worker.model_runner.eplb_state
         assert eplb_state is not None
+        eplb_state.drain_async()
         eplb_model_state = eplb_state.model_states[model_config.compute_hash()]
         physical_to_logical = eplb_model_state.physical_to_logical_map
         num_physical_experts = physical_to_logical.shape[1]
@@ -278,10 +293,6 @@ class ElasticEPScalingExecutor:
             src_rank=0,
             device=self.worker.device,
         )
-        # New workers enter load_model after receiving the expert mapping.
-        # Stage replacement MoE kernels before returning to the state machine
-        # so existing ranks can participate in collective EP comm creation.
-        self._stage_standby_moe_quant_methods()
 
     def _make_eep_moe_config(self, module, dp_group, ep_group):
         parallel_config = self.worker.vllm_config.parallel_config
@@ -300,7 +311,7 @@ class ElasticEPScalingExecutor:
             moe_parallel_config=moe_parallel_config,
         )
 
-    def _stage_standby_moe_quant_methods(self) -> None:
+    def stage_standby_moe_quant_methods(self) -> None:
         standby_dp_group = get_standby_dp_group()
         standby_ep_group = get_standby_ep_group()
         model = self.worker.model_runner.get_model()
@@ -500,26 +511,6 @@ class ElasticEPScalingExecutor:
             compilation_counter.stock_torch_compile_count += 1
             self.worker.model_runner.model.compile(fullgraph=True, backend=backend)
 
-        multi_block_table = self.worker.model_runner.input_batch.block_table
-        saved_block_tables: list[tuple[torch.Tensor, torch.Tensor]] = []
-        for bt in multi_block_table.block_tables:
-            saved_block_tables.append(
-                (bt.block_table.gpu.clone(), bt.block_table.cpu.clone())
-            )
-        multi_block_table.clear()
-
-        unlock_workspace()
-        self.worker.compile_or_warm_up_model()
-        lock_workspace()
-
-        for bt, (saved_gpu, saved_cpu) in zip(
-            multi_block_table.block_tables, saved_block_tables
-        ):
-            bt.block_table.gpu.copy_(saved_gpu)
-            bt.block_table.cpu.copy_(saved_cpu)
-        if new_dp_size < old_dp_size:
-            self._set_eplb_suppressed(False)
-
     def _perform_eplb_reshuffle(
         self, rank_mapping: dict[int, int] | None = None
     ) -> None:
@@ -553,12 +544,25 @@ class ElasticEPScalingExecutor:
         if get_ep_group().rank == 0:
             logger.info("[Elastic EP] Expert resharding completed")
 
-    def perform_eplb_reshuffle(self) -> None:
+    def commit_scale_up(self, is_existing_worker: bool) -> None:
+        if is_existing_worker:
+            self.broadcast_expert_mapping()
+            self.switch_and_prepare()
+        else:
+            mapping, _, num_valid_experts = self.receive_expert_mapping()
+            self.worker.model_runner.setup_eplb_from_mapping(mapping, num_valid_experts)
         self._perform_eplb_reshuffle()
-        self._set_eplb_suppressed(False)
+        self.warm_and_capture()
+
+    def commit_scale_down(self, new_dp_size: int, removing: bool) -> None:
+        self.perform_scale_down_eplb_reshuffle(new_dp_size)
+        if removing:
+            self.switch_and_remove()
+        else:
+            self.switch_and_prepare()
+            self.warm_and_capture()
 
     def perform_scale_down_eplb_reshuffle(self, new_dp_size: int) -> None:
-        self._set_eplb_suppressed(True)
         eplb_state = self.worker.model_runner.eplb_state
         if eplb_state is not None:
             eplb_state.drain_async()
@@ -599,12 +603,17 @@ class ElasticEPScalingExecutor:
             )
 
         model = self.worker.model_runner.get_model()
+        expert_weights = [
+            module.get_expert_weights()
+            for module in model.modules()
+            if is_moe_layer(module)
+        ]
         batch_transfer_weights(
             model=model,
             is_sender=False,
             peer_rank=sender_rank,
             dp_group=dp_group,
-            expert_weights=model.expert_weights,
+            expert_weights=expert_weights,
         )
         torch.accelerator.synchronize()
 
@@ -643,14 +652,17 @@ class ElasticEPScalingExecutor:
         with set_current_vllm_config(self.worker.vllm_config):
             prepare_communication_buffer_for_model(self.worker.model_runner.get_model())
 
-    def rewarm_workspace(self) -> None:
+    def warmup_local_kernels(self) -> None:
+        with set_current_vllm_config(self.worker.vllm_config):
+            kernel_warmup(self.worker, process_local_only=True)
+
+    def warm_and_capture(self) -> None:
         # Must run on every DP sibling in lockstep: _dummy_run calls
         # coordinate_batch_across_dp whenever data_parallel_size > 1
         # (gpu_model_runner.py:3663), which deadlocks if any rank skips it.
 
-        # Save and clear block tables so profile_run/compile_or_warm_up_model
-        # don't write dummy slot mappings into real KV-cache blocks (mirrors
-        # switch_and_prepare's pattern).
+        # Save and clear block tables so the dummy MoE forward doesn't
+        # write dummy slot mappings into real KV-cache blocks.
         multi_block_table = self.worker.model_runner.input_batch.block_table
         saved_block_tables: list[tuple[torch.Tensor, torch.Tensor]] = []
         for bt in multi_block_table.block_tables:
@@ -660,19 +672,16 @@ class ElasticEPScalingExecutor:
         multi_block_table.clear()
 
         # _ensure_workspace_size allocates a fresh tensor on grow, leaving
-        # captured CUDA graphs with stale data pointers; drop graphs before
-        # re-warm so captures realign with the resized buffer.
+        # any captured CUDA graph with a stale data pointer; drop graphs
+        # before re-warm so captures realign with the resized buffer.
         self._release_cuda_graphs()
         unlock_workspace()
 
-        # Grow the MoE workspace at max_num_tokens.
-        # compile_or_warm_up_model alone only exercises cudagraph-capture
-        # sizes (≤64 tokens for this test) and leaves the workspace at
-        # ~10-14 MB; the post-all-to-all per-rank token count under real
-        # post-reshuffle routing needs hundreds of MB. Use _dummy_run
-        # directly (rather than profile_run) with skip_eplb=True so dummy
-        # routing doesn't pollute the just-rebalanced EPLB stats — same
-        # convention compile_or_warm_up_model itself uses.
+        # Grow the MoE workspace at max_num_tokens. compile_or_warm_up_model
+        # alone only exercises cudagraph-capture sizes and can leave the
+        # workspace too small for post-reshuffle routing. Use _dummy_run
+        # directly with skip_eplb=True so dummy routing doesn't pollute the
+        # just-rebalanced EPLB stats.
         runner = self.worker.model_runner
         runner._dummy_run(runner.max_num_tokens, is_profile=True, skip_eplb=True)
         self.worker.compile_or_warm_up_model()
diff --git a/vllm/distributed/elastic_ep/elastic_state.py b/vllm/distributed/elastic_ep/elastic_state.py
index 256efe46a4a..f33fd90e863 100644
--- a/vllm/distributed/elastic_ep/elastic_state.py
+++ b/vllm/distributed/elastic_ep/elastic_state.py
@@ -1,18 +1,17 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 import enum
-import time
 import weakref
-from datetime import timedelta
-from typing import TYPE_CHECKING, Literal, TypeAlias
+from concurrent.futures import Future, ThreadPoolExecutor
+from typing import TYPE_CHECKING, Any, Literal, TypeAlias
 
 import torch.distributed
 
 from vllm.config import ParallelConfig
 from vllm.distributed import (
-    sched_yield,
     stateless_destroy_torch_distributed_process_group,
 )
+from vllm.distributed.utils import get_cached_tcp_store_client
 from vllm.logger import init_logger
 from vllm.v1.engine import (
     EEPNotificationType,
@@ -31,35 +30,29 @@ WorkerType = Literal["existing", "new", "removing"]
 
 
 class ScaleUpExistingEngineState(enum.IntEnum):
-    WAIT_NEW_CORE_ENGINES_INIT = 0
-    CREATE_STANDBY_GROUPS = 1
-    TRANSFER_EXPERT_MAPPING = 2
-    WAIT_NEW_CORE_ENGINES_WEIGHTS_INIT = 3
-    TRANSFER_WEIGHTS = 4
-    SYNC_KV_CACHE_MEMORY_SIZE = 5
-    SWITCH_AND_PREPARE = 6
-    EPLB_RESHUFFLE = 7
-    COMPLETE = 8
+    CREATE_STANDBY_GROUPS = 0
+    STAGE_QUANT_METHODS = 1
+    TRANSFER_WEIGHTS = 2
+    SYNC_KV_CACHE_MEMORY_SIZE = 3
+    COMMIT_SCALE_UP = 4  # Blocks forward passes.
+    COMPLETE = 5
 
 
 class ScaleUpNewEngineState(enum.IntEnum):
     PRE_KV_INIT = 0
     PREPARE = 1
-    EPLB_RESHUFFLE = 2
-    COMPLETE = 3
+    COMPLETE = 2
 
 
 class ScaleDownRemainingEngineState(enum.IntEnum):
     PREPARE = 0
-    EPLB_RESHUFFLE = 1
-    SWITCH_AND_PREPARE = 2
-    COMPLETE = 3
+    COMMIT_SCALE_DOWN = 1  # Blocks forward passes.
+    COMPLETE = 2
 
 
 class ScaleDownRemovingEngineState(enum.IntEnum):
     PREPARE = 0
-    EPLB_RESHUFFLE = 1
-    COMPLETE = 2
+    COMPLETE = 1
 
 
 EngineState: TypeAlias = (
@@ -70,15 +63,6 @@ EngineState: TypeAlias = (
 )
 
 
-class _BarrierTimeoutError(RuntimeError):
-    """
-    Exception raised for timeout
-    in the first stage of our two-staged
-    TCPStore based barrier to synchronize the
-    execution of all engines in the DP group.
-    """
-
-
 class ElasticEPScalingState:
     def __init__(
         self,
@@ -94,20 +78,24 @@ class ElasticEPScalingState:
         self.engine_core_ref = weakref.ref(engine_core)
         self.vllm_config = vllm_config
         self.old_dp_group = self.engine_core.dp_group if worker_type != "new" else None
-        self.old_dp_store = self.engine_core.dp_store if worker_type != "new" else None
         self.new_parallel_config: ParallelConfig = new_parallel_config
         self.new_dp_group = self.engine_core.dp_group if worker_type == "new" else None
         self.new_dp_store = self.engine_core.dp_store if worker_type == "new" else None
         self.worker_type = worker_type
         self.scale_type = scale_type
         self.reconfig_request = reconfig_request
-
+        self.commit_requested = False
+        self._prepare_executor = ThreadPoolExecutor(
+            max_workers=1, thread_name_prefix="ElasticEPPrepare"
+        )
+        self._prepare_future: Future[Any] | None = None
+        self._new_dp_sync: tuple[object, Any] | None = None
         self.state: EngineState
         if scale_type == "scale_up":
             self.state = (
                 ScaleUpNewEngineState.PRE_KV_INIT
                 if worker_type == "new"
-                else ScaleUpExistingEngineState.WAIT_NEW_CORE_ENGINES_INIT
+                else ScaleUpExistingEngineState.CREATE_STANDBY_GROUPS
             )
         else:
             self.state = (
@@ -130,6 +118,31 @@ class ElasticEPScalingState:
             raise RuntimeError("Engine core has been garbage collected")
         return engine_core
 
+    def _collective_rpc(self, *args, **kwargs):
+        return self.model_executor.collective_rpc(*args, **kwargs)
+
+    def _execute_async(self, execute_method: str, *args) -> bool:
+        if self._prepare_future is None:
+            done_keys = self._collective_rpc(
+                "elastic_ep_execute",
+                args=("start_async", execute_method, *args),
+            )
+            assert self.reconfig_request is not None
+            coord_store = get_cached_tcp_store_client(
+                self.reconfig_request.new_data_parallel_master_ip,
+                self.reconfig_request.coord_store_port,
+            )
+            self._prepare_future = self._prepare_executor.submit(
+                coord_store.wait, done_keys
+            )
+        if not self._prepare_future.done():
+            return False
+
+        self._prepare_future.result()
+        self._collective_rpc("elastic_ep_execute", args=("clear_async",))
+        self._prepare_future = None
+        return True
+
     def progress(self) -> bool:
         if self.scale_type == "scale_up":
             return (
@@ -149,157 +162,43 @@ class ElasticEPScalingState:
         assert self.progress()
         assert self.state == ScaleUpNewEngineState.PREPARE
 
-    def _execute_tcp_store_barrier(
-        self, dp_store, group_rank, group_size, barrier_id, timeout=None
-    ):
-        arrival_key = f"arrival_{barrier_id}_{group_rank}"
-        dp_store.set(arrival_key, b"1")
-
-        start_time = time.time()
-        processes_arrived: set[int] = set()
-
-        while len(processes_arrived) < group_size:
-            if (
-                timeout is not None
-                and time.time() - start_time > timeout.total_seconds()
-            ):
-                raise _BarrierTimeoutError(
-                    f"Barrier timed out after {timeout.total_seconds()} seconds"
-                )
-
-            for i in range(group_size):
-                if i in processes_arrived:
-                    continue
-
-                key = f"arrival_{barrier_id}_{i}"
-                present = dp_store.check([key])
-                if present:
-                    processes_arrived.add(i)
-
-            if len(processes_arrived) < group_size:
-                sched_yield()
-
-    def _staged_barrier(self, use_new_group: bool, barrier_name: str) -> bool:
-        """
-        Execute a two-staged barrier to synchronize all engines in the DP group.
-
-        Some DP EngineCores may receive the reconfiguration notifications
-        later than others, and already proceed to engine step (model forward)
-        in the busy loop.
-        In this case, EngineCores that already proceed to reconfiguration
-        should skip reconfiguration and execute model forward for one more
-        step, so in the next step, all EngineCores will be synchronized.
-        We use a two-staged barrier to achieve this. The first time each
-        EngineCore executes the barrier, if a timeout is reached before the
-        barrier completes, that means some EngineCores have already entered
-        engine step. The EngineCores that timed out will then proceed to
-        engine step, and will synchronize with the other EngineCores in the
-        next step with a barrier without timeout.
-        """
-        dp_group = self.new_dp_group if use_new_group else self.old_dp_group
-        dp_store = self.new_dp_store if use_new_group else self.old_dp_store
-        assert dp_group is not None and dp_store is not None
-
-        group_rank = dp_group.rank()
-        group_size = dp_group.size()
-        barrier_id = f"eep_barrier_{barrier_name}"
-        sync_key = f"{barrier_id}_sync"
-
-        # TODO(yongji): figure out appropriate timeout for the barrier
-        timeout = None if dp_store.check([sync_key]) else timedelta(seconds=5)
-
-        try:
-            self._execute_tcp_store_barrier(
-                dp_store, group_rank, group_size, barrier_id, timeout=timeout
-            )
-            torch.distributed.barrier(dp_group)
-            if group_rank == 0:
-                dp_store.delete_key(sync_key)
-                for i in range(group_size):
-                    dp_store.delete_key(f"arrival_{barrier_id}_{i}")
-            return True
-        except _BarrierTimeoutError as e:
-            if timeout is None:
-                raise RuntimeError("Unexpected timeout encountered") from e
-            dp_store.compare_set(sync_key, "", b"1")
-            return False
-
     def _progress_existing_engine(self) -> bool:
         state = self.state
-        assert self.old_dp_group is not None and self.old_dp_store is not None
+        assert self.old_dp_group is not None
 
-        if state == ScaleUpExistingEngineState.WAIT_NEW_CORE_ENGINES_INIT:
-            return False
-
-        elif state == ScaleUpExistingEngineState.CREATE_STANDBY_GROUPS:
-            # NOTE(yongji): wait for all existing workers to receive the request
-            if (
-                int(self.old_dp_store.get("eep_barrier_engine_count"))
-                < self.old_dp_group.size()
-            ):
+        if state == ScaleUpExistingEngineState.CREATE_STANDBY_GROUPS:
+            if not self._create_standby_groups():
                 return False
-            if not self._staged_barrier(
-                use_new_group=False, barrier_name="create_standby_groups"
-            ):
-                return False
-            if self.old_dp_group.rank() == 0:
-                self.old_dp_store.delete_key("eep_barrier_engine_count")
-            self._create_standby_groups()
-            self.state = ScaleUpExistingEngineState.TRANSFER_EXPERT_MAPPING
+            self.state = ScaleUpExistingEngineState.STAGE_QUANT_METHODS
             return True
 
-        elif state == ScaleUpExistingEngineState.TRANSFER_EXPERT_MAPPING:
-            self._transfer_expert_mapping()
-            self.state = ScaleUpExistingEngineState.WAIT_NEW_CORE_ENGINES_WEIGHTS_INIT
+        elif state == ScaleUpExistingEngineState.STAGE_QUANT_METHODS:
+            if not self._execute_async("stage_standby_moe_quant_methods"):
+                return False
+            self.state = ScaleUpExistingEngineState.TRANSFER_WEIGHTS
             return True
 
-        elif state == ScaleUpExistingEngineState.WAIT_NEW_CORE_ENGINES_WEIGHTS_INIT:
-            return False
-
         elif state == ScaleUpExistingEngineState.TRANSFER_WEIGHTS:
-            if (
-                int(self.old_dp_store.get("eep_barrier_engine_count"))
-                < self.old_dp_group.size()
-            ):
+            if not self._transfer_weights():
                 return False
-            if not self._staged_barrier(
-                use_new_group=False, barrier_name="transfer_weights"
-            ):
-                return False
-            if self.old_dp_group.rank() == 0:
-                self.old_dp_store.delete_key("eep_barrier_engine_count")
-            self._transfer_weights()
             self.state = ScaleUpExistingEngineState.SYNC_KV_CACHE_MEMORY_SIZE
             return True
 
         elif state == ScaleUpExistingEngineState.SYNC_KV_CACHE_MEMORY_SIZE:
-            self._sync_kv_cache_memory_size()
-            self.state = ScaleUpExistingEngineState.SWITCH_AND_PREPARE
+            if not self._sync_kv_cache_memory_size():
+                return False
+            self.state = ScaleUpExistingEngineState.COMMIT_SCALE_UP
+            self._mark_ready_for_switch()
             return True
 
-        elif state == ScaleUpExistingEngineState.SWITCH_AND_PREPARE:
-            self._switch_and_prepare()
-            self.state = ScaleUpExistingEngineState.EPLB_RESHUFFLE
-            assert self.new_dp_store is not None
-            self.new_dp_store.add("eep_barrier_engine_count", 1)
-            return True
-
-        elif state == ScaleUpExistingEngineState.EPLB_RESHUFFLE:
-            assert self.new_dp_group is not None and self.new_dp_store is not None
-            if (
-                int(self.new_dp_store.get("eep_barrier_engine_count"))
-                < self.new_dp_group.size()
-            ):
+        elif state == ScaleUpExistingEngineState.COMMIT_SCALE_UP:
+            if not self.commit_requested:
                 return False
-            if not self._staged_barrier(
-                use_new_group=True, barrier_name="eplb_reshuffle"
-            ):
-                return False
-            if self.new_dp_group.rank() == 0:
-                self.new_dp_store.delete_key("eep_barrier_engine_count")
-            self._eplb_reshuffle()
+            self._commit_new_dp_group()
+            self._collective_rpc("elastic_ep_execute", args=("commit_scale_up", True))
             self.state = ScaleUpExistingEngineState.COMPLETE
             self._update_parallel_config()
+            self._send_reconfigure_finished()
             return True
 
         else:
@@ -311,22 +210,17 @@ class ElasticEPScalingState:
         assert self.new_dp_group is not None and self.new_dp_store is not None
 
         if state == ScaleUpNewEngineState.PRE_KV_INIT:
-            self.engine_core._eep_send_engine_core_notification(
-                EEPNotificationType.NEW_CORE_ENGINES_WEIGHTS_INIT_READY
-            )
-            self.model_executor.collective_rpc(
-                "elastic_ep_execute", args=("receive_weights",)
-            )
+            self._collective_rpc("elastic_ep_execute", args=("receive_weights",))
             self.engine_core.available_gpu_memory_for_kv_cache = (
                 ParallelConfig.sync_kv_cache_memory_size(self.new_dp_group, -1)
             )
-            self.model_executor.collective_rpc(
-                "elastic_ep_execute", args=("prepare_new_worker",)
-            )
+            self._collective_rpc("elastic_ep_execute", args=("prepare_new_worker",))
             self.state = ScaleUpNewEngineState.PREPARE
             return True
 
         elif state == ScaleUpNewEngineState.PREPARE:
+            self._collective_rpc("elastic_ep_execute", args=("warmup_local_kernels",))
+            self._mark_ready_for_switch()
             tensor = torch.tensor([0, 0, 0], dtype=torch.int32, device="cpu")
             torch.distributed.all_reduce(
                 tensor,
@@ -337,22 +231,7 @@ class ElasticEPScalingState:
             self.engine_core.engines_running = bool(data[0])
             self.engine_core.current_wave = int(data[1])
             self.engine_core.step_counter = int(data[2])
-            self.state = ScaleUpNewEngineState.EPLB_RESHUFFLE
-            self.new_dp_store.add("eep_barrier_engine_count", 1)
-            return True
-
-        elif state == ScaleUpNewEngineState.EPLB_RESHUFFLE:
-            if (
-                int(self.new_dp_store.get("eep_barrier_engine_count"))
-                < self.new_dp_group.size()
-            ):
-                return False
-            if not self._staged_barrier(
-                use_new_group=True, barrier_name="eplb_reshuffle"
-            ):
-                return False
-            assert self.new_dp_group.rank() > 0
-            self._eplb_reshuffle()
+            self._collective_rpc("elastic_ep_execute", args=("commit_scale_up", False))
             self.state = ScaleUpNewEngineState.COMPLETE
             return True
 
@@ -362,38 +241,23 @@ class ElasticEPScalingState:
 
     def _progress_remaining_engine(self) -> bool:
         state = self.state
-        assert self.old_dp_group is not None and self.old_dp_store is not None
+        assert self.old_dp_group is not None
 
         if state == ScaleDownRemainingEngineState.PREPARE:
-            self.state = ScaleDownRemainingEngineState.EPLB_RESHUFFLE
-            self.old_dp_store.add("eep_barrier_engine_count", 1)
-            return True
+            if self._create_standby_groups():
+                self.state = ScaleDownRemainingEngineState.COMMIT_SCALE_DOWN
+                self._mark_ready_for_switch()
+                return True
+            return False
 
-        elif state == ScaleDownRemainingEngineState.EPLB_RESHUFFLE:
-            if (
-                int(self.old_dp_store.get("eep_barrier_engine_count"))
-                < self.old_dp_group.size()
-            ):
+        elif state == ScaleDownRemainingEngineState.COMMIT_SCALE_DOWN:
+            if not self.commit_requested:
                 return False
-            if not self._staged_barrier(
-                use_new_group=False, barrier_name="eplb_reshuffle"
-            ):
-                return False
-            if self.old_dp_group.rank() == 0:
-                self.old_dp_store.delete_key("eep_barrier_engine_count")
-            self._eplb_reshuffle_before_scale_down()
-            self.state = ScaleDownRemainingEngineState.SWITCH_AND_PREPARE
-            # NOTE(yongji): currently, after EPLB reshuffle
-            # that redistributes experts to remaining workers, workers
-            # to be removed will immediately initiate shutdown;
-            # existing workers can no longer execute forward steps using
-            # the old setup. In the future, we may keep
-            # the removing workers alive a bit longer,
-            # e.g., to drain in-batch requests.
-            self._create_standby_groups()
-            self._switch_and_prepare()
+            self._commit_scale_down(removing=False)
+            self._commit_new_dp_group()
             self._update_parallel_config()
             self.state = ScaleDownRemainingEngineState.COMPLETE
+            self._send_reconfigure_finished()
             return True
 
         else:
@@ -402,26 +266,11 @@ class ElasticEPScalingState:
 
     def _progress_removing_engine(self) -> bool:
         state = self.state
-        assert self.old_dp_group is not None and self.old_dp_store is not None
+        assert self.old_dp_group is not None
 
         if state == ScaleDownRemovingEngineState.PREPARE:
-            self.state = ScaleDownRemovingEngineState.EPLB_RESHUFFLE
-            self.old_dp_store.add("eep_barrier_engine_count", 1)
-            return True
-
-        if state == ScaleDownRemovingEngineState.EPLB_RESHUFFLE:
-            if (
-                int(self.old_dp_store.get("eep_barrier_engine_count"))
-                < self.old_dp_group.size()
-            ):
-                return False
-            if not self._staged_barrier(
-                use_new_group=False, barrier_name="eplb_reshuffle"
-            ):
-                return False
             assert self.old_dp_group.rank() > 0
-            self._eplb_reshuffle_before_scale_down()
-            self._switch_and_remove()
+            self._commit_scale_down(removing=True)
             self.state = ScaleDownRemovingEngineState.COMPLETE
             self.engine_core._eep_send_engine_core_notification(
                 EEPNotificationType.SHUTDOWN_COMPLETE
@@ -432,22 +281,22 @@ class ElasticEPScalingState:
             assert self.state == ScaleDownRemovingEngineState.COMPLETE
             return True
 
-    def handle_notification(self, notification_type: EEPNotificationType):
-        assert self.worker_type != "new"
-        assert self.old_dp_store is not None
-        if (
-            notification_type == EEPNotificationType.NEW_CORE_ENGINES_INIT_READY
-            and self.state == ScaleUpExistingEngineState.WAIT_NEW_CORE_ENGINES_INIT
-        ):
-            self.old_dp_store.add("eep_barrier_engine_count", 1)
-            self.state = ScaleUpExistingEngineState.CREATE_STANDBY_GROUPS
-        elif (
-            notification_type == EEPNotificationType.NEW_CORE_ENGINES_WEIGHTS_INIT_READY
-            and self.state
-            == ScaleUpExistingEngineState.WAIT_NEW_CORE_ENGINES_WEIGHTS_INIT
-        ):
-            self.old_dp_store.add("eep_barrier_engine_count", 1)
-            self.state = ScaleUpExistingEngineState.TRANSFER_WEIGHTS
+    def is_ready_for_switch(self) -> bool:
+        return self.worker_type == "existing" and (
+            self.state is ScaleUpExistingEngineState.COMMIT_SCALE_UP
+            or self.state is ScaleDownRemainingEngineState.COMMIT_SCALE_DOWN
+        )
+
+    @property
+    def ready_key(self) -> str:
+        return f"eep_ready/{self.engine_core.dp_rank}"
+
+    def _mark_ready_for_switch(self) -> None:
+        parallel_config = self.new_parallel_config
+        get_cached_tcp_store_client(
+            parallel_config.data_parallel_master_ip,
+            parallel_config._coord_store_port,
+        ).set(self.ready_key, b"1")
 
     def is_complete(self) -> bool:
         if self.scale_type == "scale_up":
@@ -462,50 +311,78 @@ class ElasticEPScalingState:
             else self.state == ScaleDownRemainingEngineState.COMPLETE
         )
 
-    def _create_standby_groups(self):
+    def _init_new_dp_group(self) -> tuple[Any, Any]:
+        return self.new_parallel_config.stateless_init_dp_group(return_store=True)
+
+    def _ensure_new_dp_group(self) -> bool:
+        if self.new_dp_group is not None:
+            return True
+
+        if self._prepare_future is None:
+            self._prepare_future = self._prepare_executor.submit(
+                self._init_new_dp_group
+            )
+        if not self._prepare_future.done():
+            return False
+
+        self.new_dp_group, self.new_dp_store = self._prepare_future.result()
+        self._prepare_future = None
+        return True
+
+    def _create_standby_groups(self) -> bool:
         assert self.old_dp_group is not None
-        self.new_dp_group, self.new_dp_store = (
-            self.new_parallel_config.stateless_init_dp_group(return_store=True)
-        )
-        self.model_executor.collective_rpc(
-            "elastic_ep_execute", args=("create_standby_groups", self.reconfig_request)
-        )
+        if not self._ensure_new_dp_group():
+            return False
+        if not self._execute_async(
+            "create_standby_groups",
+            self.reconfig_request,
+            self.new_parallel_config.use_all2all,
+        ):
+            return False
         if self.old_dp_group.rank() == 0:
             logger.info("[Elastic EP] Created standby communication groups")
+        return True
 
-    def _transfer_weights(self):
+    def _transfer_weights(self) -> bool:
         assert self.reconfig_request is not None and self.old_dp_group is not None
         old_dp_size = self.old_dp_group.size()
         new_dp_size = self.reconfig_request.new_data_parallel_size
 
-        self.model_executor.collective_rpc(
-            "elastic_ep_execute", args=("transfer_weights", old_dp_size, new_dp_size)
-        )
+        if not self._execute_async("transfer_weights", old_dp_size, new_dp_size):
+            return False
         if self.old_dp_group.rank() == 0:
             logger.info("[Elastic EP] Transferred weights to new workers")
+        return True
 
-    def _transfer_expert_mapping(self):
-        assert self.old_dp_group is not None
-        self.model_executor.collective_rpc(
-            "elastic_ep_execute", args=("broadcast_expert_mapping",)
-        )
-        if self.old_dp_group.rank() == 0:
-            logger.info("[Elastic EP] Broadcasted expert mapping to new workers")
-
-    def _sync_kv_cache_memory_size(self):
+    def _sync_kv_cache_memory_size(self) -> bool:
         assert self.engine_core.available_gpu_memory_for_kv_cache > 0
         assert self.new_dp_group is not None and self.old_dp_group is not None
-        ParallelConfig.sync_kv_cache_memory_size(
-            self.new_dp_group,
-            self.engine_core.available_gpu_memory_for_kv_cache,
-        )
+
+        if self._new_dp_sync is None:
+            tensor = torch.tensor(
+                [self.engine_core.available_gpu_memory_for_kv_cache],
+                dtype=torch.int64,
+                device="cpu",
+            )
+            work = torch.distributed.all_reduce(
+                tensor,
+                op=torch.distributed.ReduceOp.MIN,
+                group=self.new_dp_group,
+                async_op=True,
+            )
+            self._new_dp_sync = (tensor, work)
+            return False
+
+        _, work = self._new_dp_sync
+        if not work.is_completed():
+            return False
+        work.wait()
+        self._new_dp_sync = None
         if self.old_dp_group.rank() == 0:
             logger.info("[Elastic EP] Synced KV cache memory size to new workers")
+        return True
 
-    def _switch_and_prepare(self):
-        self.model_executor.collective_rpc(
-            "elastic_ep_execute", args=("switch_and_prepare",)
-        )
+    def _commit_new_dp_group(self):
         old_dp_group = self.old_dp_group
         stateless_destroy_torch_distributed_process_group(old_dp_group)
         assert self.new_dp_group is not None
@@ -529,41 +406,28 @@ class ElasticEPScalingState:
         self.engine_core.current_wave = int(data[1])
         self.engine_core.step_counter = int(data[2])
         if new_dp_group.rank() == 0:
+            logger.info("[Elastic EP] Switched to new setup")
+
+    def _send_reconfigure_finished(self):
+        assert self.new_dp_group is not None
+        if self.new_dp_group.rank() == 0:
             self.engine_core._eep_send_engine_core_notification(
                 EEPNotificationType.RECONFIGURE_FINISHED
             )
-            logger.info("[Elastic EP] Switched to new setup")
 
-    def _eplb_reshuffle(self):
-        self.model_executor.collective_rpc(
-            "elastic_ep_execute", args=("perform_eplb_reshuffle",)
-        )
-        # Reshuffle changes per-rank token routing; the locked MoE workspace
-        # may now be too small. Rewarm covers both new and existing engines.
-        self.model_executor.collective_rpc(
-            "elastic_ep_execute", args=("rewarm_workspace",)
-        )
-        assert self.new_dp_group is not None
-        if self.new_dp_group.rank() == 0:
-            logger.info("[Elastic EP] EPLB reshuffle completed")
-
-    def _eplb_reshuffle_before_scale_down(self):
+    def _commit_scale_down(self, removing: bool):
         assert self.reconfig_request is not None and self.old_dp_group is not None
-        self.model_executor.collective_rpc(
+        self._collective_rpc(
             "elastic_ep_execute",
             args=(
-                "perform_scale_down_eplb_reshuffle",
+                "commit_scale_down",
                 self.reconfig_request.new_data_parallel_size,
+                removing,
             ),
         )
         if self.old_dp_group.rank() == 0:
             logger.info("[Elastic EP] EPLB reshuffle completed")
 
-    def _switch_and_remove(self):
-        self.model_executor.collective_rpc(
-            "elastic_ep_execute", args=("switch_and_remove",)
-        )
-
     def _update_parallel_config(self):
         assert self.reconfig_request is not None
         reconfig_request = self.reconfig_request
diff --git a/vllm/distributed/elastic_ep/standby_state.py b/vllm/distributed/elastic_ep/standby_state.py
index 846793a955f..1892f3e7942 100644
--- a/vllm/distributed/elastic_ep/standby_state.py
+++ b/vllm/distributed/elastic_ep/standby_state.py
@@ -39,6 +39,7 @@ def create_standby_groups(
     new_world_size_across_dp: int,
     master_ip: str,
     coord_store_port: int,
+    use_all2all: bool,
     enable_eplb: bool = True,
     backend: str | None = None,
 ) -> None:
@@ -86,7 +87,7 @@ def create_standby_groups(
     )
     standby_ep_ranks = [x.tolist() for x in standby_ep_ranks]
     _STANDBY_EP = _init_stateless_group(
-        standby_ep_ranks, "ep", master_ip, backend, coord_store=coord_store
+        standby_ep_ranks, "ep", master_ip, backend, coord_store, use_all2all=use_all2all
     )
 
     if enable_eplb:
diff --git a/vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py b/vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py
index e5edad8861c..12b9f0d937f 100644
--- a/vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py
+++ b/vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py
@@ -68,6 +68,7 @@ from vllm.distributed.parallel_state import (
 from vllm.logger import init_logger
 from vllm.platforms import current_platform
 from vllm.utils.network_utils import make_zmq_path
+from vllm.utils.torch_utils import async_tensor_h2d
 from vllm.v1.attention.backends.utils import get_kv_cache_layout
 from vllm.v1.kv_cache_interface import (
     FullAttentionSpec,
@@ -172,11 +173,6 @@ class NixlBaseConnectorWorker:
         Split counts are derived from source_ranks_per_group lengths.
         FA uses rank_to_attention_slot for the slot offset;
         SSM uses the rank's positional index.
-
-        With ``block_size_ratio`` > 1 the FA descriptors are remote-granularity
-        sub-blocks; replicated regions and single-source FA pass through whole,
-        and SSM descriptors are never ratio-expanded, so only genuinely
-        head-sharded FA reads are incompatible with a block-size mismatch.
         """
         fa_idx = next(
             i for i, t in enumerate(self._group_spec_types) if _is_attention_spec(t)
@@ -458,7 +454,7 @@ class NixlBaseConnectorWorker:
         self.src_xfer_handles_by_block_size: dict[int, int] = {}
         # Local descriptor arrays per remote block size (block_size_ratio>1),
         # kept for building per-tp-ratio splits at the same granularity.
-        self._src_blocks_data_by_block_size: dict[int, np.ndarray] = {}
+        self.src_blocks_data_by_block_size: dict[int, np.ndarray] = {}
         # Populated dynamically during handshake based on remote configuration.
         # Per-source split handles, keyed by (tp_ratio, remote_block_size).
         self.src_xfer_handles_by_tp_ratio: dict[tuple[int, int], list[int]] = {}
@@ -717,8 +713,9 @@ class NixlBaseConnectorWorker:
                     )
                 setup_agent_time = time.perf_counter()
                 logger.debug(
-                    "NIXL handshake: add agent took: %s",
+                    "NIXL handshake: add agent took: %s (notif_agents_only=%s)",
                     setup_agent_time - got_metadata_time,
+                    notif_agents_only,
                 )
                 remote_ranks = (remote_pp_rank, remote_rank)
                 remote_rank_to_agent_name[remote_ranks] = remote_agent_name
@@ -1141,19 +1138,20 @@ class NixlBaseConnectorWorker:
             # [`num_blocks` * `page_size`]
             curr_tensor_size_bytes = num_blocks * physical_page_size
 
+            base_addr = cache.data_ptr()
             is_mla_region = isinstance(
                 layer_spec, (MLAAttentionSpec, SlidingWindowMLASpec)
             )
-            base_addr = cache.data_ptr()
             if base_addr in seen_base_addresses:
-                region_idx = seen_base_addresses.index(base_addr)
-                self._region_is_mla[region_idx] = (
-                    self._region_is_mla[region_idx] or is_mla_region
-                )
                 # NOTE (NickLucche) HMA employs memory pooling to share tensors
                 # across groups. This results in skipping all tensors but the ones
                 # pointed to by group0. Also, generally we will have more blocks
                 # per tensor but fewer regions.
+                # A shared tensor may back both SSM and attention layers (e.g.
+                # KDA+MLA in KimiLinear); the region's FA view is MLA whichever
+                # layer registered it first.
+                idx = seen_base_addresses.index(base_addr)
+                self._region_is_mla[idx] |= is_mla_region
                 logger.debug("Skipping %s because it's already seen", layer_name)
                 continue
             logger.debug(
@@ -1632,8 +1630,8 @@ class NixlBaseConnectorWorker:
                     remote_block_size
                 )
                 self.src_xfer_handles_by_block_size[remote_block_size] = handle
-                self._src_blocks_data_by_block_size[remote_block_size] = blocks_data
-            src_blocks_data = self._src_blocks_data_by_block_size[remote_block_size]
+                self.src_blocks_data_by_block_size[remote_block_size] = blocks_data
+            src_blocks_data = self.src_blocks_data_by_block_size[remote_block_size]
 
         ### (Optional) Register local agent memory regions. MLA is not split.
         split_key = (tp_ratio, remote_block_size)
@@ -1811,12 +1809,16 @@ class NixlBaseConnectorWorker:
         if self._has_mamba and self.use_mla:
             # Hybrid MLA+SSM (e.g. KimiLinear's KDA+MLA): regions are
             # kernel-granularity views of the mamba-unified page. The MLA
-            # per-token page and the kernel block size are TP-independent,
-            # so block_lens must match exactly even under heterogeneous TP.
+            # per-token page is TP-independent, so the block lengths must
+            # match up to the kernel block size ratio even under
+            # heterogeneous TP (remote kernel blocks may be smaller).
             # SSM geometry is validated via ssm_sizes/conv offsets instead.
-            assert self.block_len_per_layer == nixl_agent_meta.block_lens, (
+            assert self.block_len_per_layer == [
+                block_len * block_size_ratio for block_len in nixl_agent_meta.block_lens
+            ], (
                 "Hybrid MLA kernel-granularity block lengths must match "
-                f"between P and D: local={self.block_len_per_layer}, "
+                f"between P and D (block_size_ratio={block_size_ratio}): "
+                f"local={self.block_len_per_layer}, "
                 f"remote={nixl_agent_meta.block_lens}."
             )
         elif not self._has_mamba:
@@ -1932,29 +1934,38 @@ class NixlBaseConnectorWorker:
     def post_process_device_kv_on_receive(
         self,
         block_size_ratio: int,
-        block_ids_list: list[tuple[list[int], int | None]],
+        block_ids_list: list[tuple[list[int], int]],
+        convert: bool = True,
     ):
         """
         Post process device kv cache after receiving from remote.
 
-        3 types of post processing supported:
+        3 types of conversion supported (``convert``):
             * kv_cache_postprocess_layout => convert from HND to NHD
             * kv_cache_postprocess_blksize => convert from small block size
               to large block size
             * kv_cache_postprocess_blksize_and_layout => convert from small
               block size to large block size and convert from HND to NHD
 
-        With a block-size ratio, the last local block of a request may have
-        received fewer than ``block_size_ratio`` remote sub-blocks; its
-        untransferred token tail is zeroed here, since these freshly
-        allocated blocks were excluded from the scheduler's KV zeroing
-        (stale bytes could otherwise surface as NaNs on hybrid models).
+        The transfer only covers ``covered_sub_blocks`` remote-sized
+        sub-blocks of each request's local attention blocks; the rest was
+        clipped, either by remote-block pairing (block-size ratio) or by the
+        hetero-ppl front trim in ``_apply_prefix_caching``. Those blocks were
+        excluded from the scheduler's alloc-time KV zeroing (which would race
+        the RDMA write), so everything past the covered range is zeroed here.
+        Stale bytes would otherwise surface as garbage or NaNs once decode
+        grows into the untransferred tail.
         """
         if len(self.device_kv_caches) == 0:
             return
         assert block_size_ratio >= 1, "Only nP < nD supported currently."
         assert self.transfer_topo is not None
-        if self.enable_permute_local_kv and block_size_ratio > 1:
+        if not convert:
+            logger.debug(
+                "Post-processing device kv cache on receive by zeroing "
+                "untransferred blocks."
+            )
+        elif self.enable_permute_local_kv and block_size_ratio > 1:
             logger.debug(
                 "Post-processing device kv cache on receive by converting "
                 "block_size with %sx bigger and permuting layout from HND"
@@ -1974,63 +1985,44 @@ class NixlBaseConnectorWorker:
             )
 
         attn_caches = self._attention_kv_caches
+        device = attn_caches[0].device
         for block_ids, covered_sub_blocks in block_ids_list:
-            indices = torch.tensor(block_ids, device=self.device_type, dtype=torch.long)
+            # Blocks the transfer didn't write: the token tail of the last
+            # partially covered block, then everything beyond it.
+            covered_blocks, sub_blocks_in_last = divmod(
+                covered_sub_blocks, block_size_ratio
+            )
+            first_stale = covered_blocks + (1 if sub_blocks_in_last else 0)
+            has_stale = first_stale < len(block_ids)
+            indices = None
+            if convert or has_stale:
+                indices = async_tensor_h2d(block_ids, device, torch.long)
 
-            for cache in attn_caches:
-                if self.enable_permute_local_kv and block_size_ratio > 1:
-                    kv_postprocess_blksize_and_layout_on_receive(
-                        cache, indices, block_size_ratio
-                    )
-                elif self.enable_permute_local_kv:
-                    kv_postprocess_layout_on_receive(cache, indices)
-                else:
-                    kv_postprocess_blksize_on_receive(cache, indices, block_size_ratio)
+            if convert:
+                for cache in attn_caches:
+                    if self.enable_permute_local_kv and block_size_ratio > 1:
+                        kv_postprocess_blksize_and_layout_on_receive(
+                            cache, indices, block_size_ratio
+                        )
+                    elif self.enable_permute_local_kv:
+                        kv_postprocess_layout_on_receive(cache, indices)
+                    else:
+                        kv_postprocess_blksize_on_receive(
+                            cache, indices, block_size_ratio
+                        )
 
-            if covered_sub_blocks is None:
-                continue
-            # Zero the untransferred token tail of the last covered block
-            # (blocks wholly beyond the data are never read and stay as-is).
-            last_idx = (covered_sub_blocks - 1) // block_size_ratio
-            covered_in_last = covered_sub_blocks - last_idx * block_size_ratio
-            if covered_in_last == block_size_ratio:
-                continue
-            last_block_id = block_ids[last_idx]
-            for cache in attn_caches:
-                # Both post-processed layouts leave tokens on dim 1.
-                sub_block_tokens = cache.shape[1] // block_size_ratio
-                cache[last_block_id, covered_in_last * sub_block_tokens :].zero_()
-
-    def _zero_untransferred_hetero_ppl_tail(self, meta: ReqMeta, remote_info) -> None:
-        """Zero attention kernel blocks the hetero-ppl transfer clipped.
-
-        With equal kernel pages but differing logical block sizes (hybrid
-        heterogeneous TP), the transfer is front-trimmed to
-        min(local, remote) kernel blocks, leaving the tail of the last
-        local logical block unwritten. Those blocks were excluded from the
-        scheduler's alloc-time KV zeroing (it would race the RDMA write),
-        so stale bytes would otherwise surface as garbage once decode
-        grows into them.
-        """
-        assert meta.remote is not None
-        if (
-            not self._has_mamba
-            or remote_info.remote_physical_blocks_per_logical
-            == self._physical_blocks_per_logical_kv_block
-        ):
-            return
-        stale_ids: list[int] = []
-        for g, local_group in enumerate(meta.local_physical_block_ids):
-            if not local_group or _is_ssm_spec(self._group_spec_types[g]):
-                continue
-            covered = min(len(local_group), len(meta.remote.block_ids[g]))
-            stale_ids.extend(local_group[covered:])
-        if not stale_ids:
-            return
-        caches = self._attention_kv_caches
-        indices = torch.tensor(stale_ids, device=caches[0].device, dtype=torch.long)
-        for cache in caches:
-            cache.index_fill_(0, indices, 0)
+            if sub_blocks_in_last:
+                last_block_id = block_ids[covered_blocks]
+                for cache in attn_caches:
+                    # Both post-processed layouts leave tokens on dim 1.
+                    sub_block_tokens = cache.shape[1] // block_size_ratio
+                    zero_from = sub_blocks_in_last * sub_block_tokens
+                    cache[last_block_id, zero_from:].zero_()
+            if has_stale:
+                assert indices is not None
+                stale_ids = indices[first_stale:]
+                for cache in attn_caches:
+                    cache.index_fill_(0, stale_ids, 0)
 
     def post_process_device_kv_on_receive_heterogeneous_attn(
         self, block_ids: list[int]
@@ -2100,30 +2092,33 @@ class NixlBaseConnectorWorker:
             if self.use_host_buffer:
                 self.sync_recved_kv_to_device(req_id, meta)
 
-            # post processing for heteroblocksize
+            # Post processing for heteroblocksize/layout, and for blocks the
+            # transfer clipped. The latter happens either at remote-block
+            # granularity (block_size_ratio > 1) or at kernel-block
+            # granularity, when equal kernel pages meet differing logical
+            # block sizes and _apply_prefix_caching front-trims to the
+            # minimum count (hybrid heterogeneous TP).
             remote_info = self.transfer_topo.get_engine_info(meta.remote.engine_id)
             block_size_ratio = self.transfer_topo.block_size_ratio(
                 remote_info.remote_block_size
             )
-            if not self.use_mla and (
-                block_size_ratio > 1 or self.enable_permute_local_kv
-            ):
+            hetero_ppl = (
+                remote_info.remote_physical_blocks_per_logical
+                != self._physical_blocks_per_logical_kv_block
+            )
+            if block_size_ratio > 1 or self.enable_permute_local_kv or hetero_ppl:
                 for g, local_group in enumerate(meta.local_physical_block_ids):
                     if not local_group or _is_ssm_spec(self._group_spec_types[g]):
                         continue
                     # Number of remote-sized sub-blocks the transfer covered;
-                    # the remainder of the last local block was clipped from
-                    # the transfer and must be zeroed.
-                    covered_sub_blocks = None
-                    if block_size_ratio > 1:
-                        covered_sub_blocks = min(
-                            len(local_group) * block_size_ratio,
-                            len(meta.remote.block_ids[g]),
-                        )
+                    # everything past this was clipped and must be zeroed.
+                    covered_sub_blocks = min(
+                        len(local_group) * block_size_ratio,
+                        len(meta.remote.block_ids[g]),
+                    )
                     block_ids_for_blocksize_post_process[block_size_ratio].append(
                         (local_group, covered_sub_blocks)
                     )
-            self._zero_untransferred_hetero_ppl_tail(meta, remote_info)
             # post processing for heterogeneous attention
             if self.enable_heterogeneous_attn_post_process:
                 block_ids_for_heterogeneous_attn_post_process.append(
@@ -2133,7 +2128,14 @@ class NixlBaseConnectorWorker:
             block_size_ratio,
             block_ids_list,
         ) in block_ids_for_blocksize_post_process.items():
-            self.post_process_device_kv_on_receive(block_size_ratio, block_ids_list)
+            # MLA never needs the block-size/layout conversion, but its
+            # clipped blocks still need zeroing.
+            convert = not self.use_mla and (
+                block_size_ratio > 1 or self.enable_permute_local_kv
+            )
+            self.post_process_device_kv_on_receive(
+                block_size_ratio, block_ids_list, convert
+            )
 
         for block_ids in block_ids_for_heterogeneous_attn_post_process:
             self.post_process_device_kv_on_receive_heterogeneous_attn(block_ids)
diff --git a/vllm/distributed/kv_transfer/kv_connector/v1/nixl/push_worker.py b/vllm/distributed/kv_transfer/kv_connector/v1/nixl/push_worker.py
index 0d2f55234df..13af0314f0e 100644
--- a/vllm/distributed/kv_transfer/kv_connector/v1/nixl/push_worker.py
+++ b/vllm/distributed/kv_transfer/kv_connector/v1/nixl/push_worker.py
@@ -308,8 +308,10 @@ class NixlPushConnectorWorker(NixlBaseConnectorWorker):
             reg_data["remote_port"],
             reg_data["remote_tp_size"],
             pp_size=remote_pp_size,
-            # D never addresses P memory in push mode; just load P's agents.
-            notif_agents_only=remote_pp_size > 1,
+            # D only ever sends PUSH_REG notifs to P and never reads or writes
+            # P's memory in push mode, so it never needs the transfer
+            # descriptors set up by the full add_remote_agent path.
+            notif_agents_only=True,
         )
         if fut is None:
             self._do_send_reg_notif(req_id, reg_data)
diff --git a/vllm/distributed/kv_transfer/kv_connector/v1/offloading/config.py b/vllm/distributed/kv_transfer/kv_connector/v1/offloading/config.py
index b86ebf96bb6..b9837bcb4b0 100644
--- a/vllm/distributed/kv_transfer/kv_connector/v1/offloading/config.py
+++ b/vllm/distributed/kv_transfer/kv_connector/v1/offloading/config.py
@@ -5,7 +5,11 @@
 from typing import TYPE_CHECKING
 
 from vllm.v1.core.kv_cache_utils import resolve_kv_cache_block_sizes
-from vllm.v1.kv_cache_interface import FullAttentionSpec, MLAAttentionSpec
+from vllm.v1.kv_cache_interface import (
+    AttentionSpec,
+    FullAttentionSpec,
+    MLAAttentionSpec,
+)
 from vllm.v1.kv_offload.config import (
     OffloadingCacheConfig,
     OffloadingConfig,
@@ -40,7 +44,11 @@ def build_offloading_config(
         OffloadingGroupConfig(
             tokens_per_block=(
                 group.kv_cache_spec.block_size
-                * parallel_config.decode_context_parallel_size
+                * (
+                    parallel_config.decode_context_parallel_size
+                    if isinstance(group.kv_cache_spec, AttentionSpec)
+                    else 1
+                )
             ),
             layer_names=tuple(group.layer_names),
         )
diff --git a/vllm/distributed/parallel_state.py b/vllm/distributed/parallel_state.py
index 4284a609d67..a90e8acbcad 100644
--- a/vllm/distributed/parallel_state.py
+++ b/vllm/distributed/parallel_state.py
@@ -414,6 +414,7 @@ class GroupCoordinator:
         use_device_communicator: bool,  # whether to use device communicator
         use_message_queue_broadcaster: bool = False,
         group_name: str | None = None,
+        use_all2all: bool = False,
     ):
         group_name = group_name or "anonymous"
         self.unique_name = _get_unique_name(group_name)
@@ -508,6 +509,7 @@ class GroupCoordinator:
                 device=self.device,
                 device_group=self.device_group,
                 unique_name=self.unique_name,
+                use_all2all=use_all2all,
             )
 
         from vllm.distributed.device_communicators.shm_broadcast import MessageQueue
@@ -1321,6 +1323,7 @@ def init_model_parallel_group(
     use_message_queue_broadcaster: bool = False,
     group_name: str | None = None,
     use_device_communicator: bool = True,
+    use_all2all: bool = False,
 ) -> GroupCoordinator:
     return GroupCoordinator(
         group_ranks=group_ranks,
@@ -1329,6 +1332,7 @@ def init_model_parallel_group(
         use_device_communicator=use_device_communicator,
         use_message_queue_broadcaster=use_message_queue_broadcaster,
         group_name=group_name,
+        use_all2all=use_all2all,
     )
 
 
@@ -1339,6 +1343,7 @@ def _init_stateless_group(
     backend: str,
     coord_store: Store,
     use_device_communicator: bool = True,
+    use_all2all: bool = False,
 ) -> "StatelessGroupCoordinator":
     """Create a StatelessGroupCoordinator with the given parameters."""
     from vllm.distributed.stateless_coordinator import StatelessGroupCoordinator
@@ -1354,6 +1359,7 @@ def _init_stateless_group(
         coord_store=coord_store,
         global_rank=world.rank,
         global_world_size=world.world_size,
+        use_all2all=use_all2all,
     )
 
 
@@ -1924,6 +1930,7 @@ def initialize_model_parallel(
             .unbind(0)
         )
         group_ranks = [x.tolist() for x in group_ranks]
+        use_all2all = parallel_config.use_all2all
         if enable_elastic_ep:
             _EP = _init_stateless_group(
                 group_ranks,
@@ -1931,10 +1938,15 @@ def initialize_model_parallel(
                 parallel_config.data_parallel_master_ip,
                 backend,
                 coord_store=coord_store,
+                use_all2all=use_all2all,
             )
         else:
             _EP = init_model_parallel_group(
-                group_ranks, get_world_group().local_rank, backend, group_name="ep"
+                group_ranks,
+                get_world_group().local_rank,
+                backend,
+                group_name="ep",
+                use_all2all=use_all2all,
             )
 
         # Create EPLB group with the same ranks as EP if EPLB is enabled.
diff --git a/vllm/distributed/stateless_coordinator.py b/vllm/distributed/stateless_coordinator.py
index 5f4597d07cb..38c74a97c55 100644
--- a/vllm/distributed/stateless_coordinator.py
+++ b/vllm/distributed/stateless_coordinator.py
@@ -79,6 +79,7 @@ class StatelessGroupCoordinator(GroupCoordinator):
         host: str = "127.0.0.1",
         global_rank: int = 0,
         global_world_size: int = 1,
+        use_all2all: bool = False,
     ):
         group_name = group_name or "anonymous"
         self.unique_name = _get_unique_name(group_name)
@@ -191,6 +192,7 @@ class StatelessGroupCoordinator(GroupCoordinator):
                 global_ranks=self.ranks,
                 global_world_size=global_world_size,
                 tcp_store_group=self.tcp_store_group,
+                use_all2all=use_all2all,
             )
 
         self.mq_broadcaster = None
diff --git a/vllm/distributed/weight_transfer/base.py b/vllm/distributed/weight_transfer/base.py
index 2e377e29253..adddf41ff4e 100644
--- a/vllm/distributed/weight_transfer/base.py
+++ b/vllm/distributed/weight_transfer/base.py
@@ -370,7 +370,7 @@ class VLLMWeightSyncClient(Protocol):
 
     def update_weights(self, update_info: dict[str, Any]) -> None: ...
 
-    def finish_weight_update(self) -> None: ...
+    def finish_weight_update(self, weight_version: str | None = None) -> None: ...
 
 
 class TrainerWeightTransferEngine(ABC, Generic[TConfig, TInitInfo]):
diff --git a/vllm/distributed/weight_transfer/clients.py b/vllm/distributed/weight_transfer/clients.py
index 4f54a6e291e..12dd0c9eacc 100644
--- a/vllm/distributed/weight_transfer/clients.py
+++ b/vllm/distributed/weight_transfer/clients.py
@@ -77,8 +77,11 @@ class HTTPVLLMWeightSyncClient:
             "update_weights", {"update_info": _json_safe_update_info(update_info)}
         )
 
-    def finish_weight_update(self) -> None:
-        self._post("finish_weight_update")
+    def finish_weight_update(self, weight_version: str | None = None) -> None:
+        json = (
+            {"weight_version": weight_version} if weight_version is not None else None
+        )
+        self._post("finish_weight_update", json)
 
 
 class RayVLLMWeightSyncClient:
@@ -108,7 +111,11 @@ class RayVLLMWeightSyncClient:
         request = WeightTransferUpdateRequest(update_info=update_info)
         ray.get([h.update_weights.remote(request) for h in self.handles])
 
-    def finish_weight_update(self) -> None:
+    def finish_weight_update(self, weight_version: str | None = None) -> None:
         import ray
 
         ray.get([h.finish_weight_update.remote() for h in self.handles])
+        if weight_version is not None:
+            ray.get(
+                [h.update_weight_version.remote(weight_version) for h in self.handles]
+            )
diff --git a/vllm/engine/protocol.py b/vllm/engine/protocol.py
index ef3be178ac8..5a9b9f96d2c 100644
--- a/vllm/engine/protocol.py
+++ b/vllm/engine/protocol.py
@@ -267,6 +267,14 @@ class EngineClient(ABC):
         """Batched weight update for RL training."""
         raise NotImplementedError
 
-    async def finish_weight_update(self) -> None:
-        """Finish the current weight update."""
+    async def finish_weight_update(self, weight_version: str | None = None) -> None:
+        """Finish the weight update and set its version if provided."""
+        raise NotImplementedError
+
+    async def update_weight_version(self, new_version: str) -> None:
+        """Set the weight version without updating weights."""
+        raise NotImplementedError
+
+    async def get_weight_version(self) -> str:
+        """Return the latest committed weight version."""
         raise NotImplementedError
diff --git a/vllm/entrypoints/cli/benchmark/main.py b/vllm/entrypoints/cli/benchmark/main.py
index 1afac64b148..9ea49987091 100644
--- a/vllm/entrypoints/cli/benchmark/main.py
+++ b/vllm/entrypoints/cli/benchmark/main.py
@@ -2,18 +2,38 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import argparse
+import os
 import sys
 import typing
 
+from vllm import envs
 from vllm.entrypoints.cli.benchmark.base import BenchmarkSubcommandBase
 from vllm.entrypoints.cli.types import CLISubcommand
 from vllm.entrypoints.serve.utils.api_utils import VLLM_SUBCMD_PARSER_EPILOG
+from vllm.logger import init_logger
 
 if typing.TYPE_CHECKING:
     from vllm.utils.argparse_utils import FlexibleArgumentParser
 else:
     FlexibleArgumentParser = argparse.ArgumentParser
 
+logger = init_logger(__name__)
+
+
+def maybe_exec_rust_bench() -> None:
+    if sys.argv[1:3] != ["bench", "serve"] or not envs.VLLM_USE_RUST_BENCH:
+        return
+
+    rust_cli = envs.VLLM_RUST_FRONTEND_PATH
+    if rust_cli is None:
+        raise RuntimeError(
+            "VLLM_USE_RUST_BENCH=1 requires VLLM_RUST_FRONTEND_PATH "
+            "to resolve to the vllm-rs binary."
+        )
+
+    logger.info("Delegating `vllm bench serve` to Rust binary at %s.", rust_cli)
+    os.execv(rust_cli, [rust_cli, "bench", "serve", *sys.argv[3:]])
+
 
 def _import_bench_subcommand_modules() -> None:
     # Imported lazily so `BenchmarkSubcommandBase` subclasses register only
diff --git a/vllm/entrypoints/cli/benchmark/serve.py b/vllm/entrypoints/cli/benchmark/serve.py
index 41a65273ba8..188afd6c703 100644
--- a/vllm/entrypoints/cli/benchmark/serve.py
+++ b/vllm/entrypoints/cli/benchmark/serve.py
@@ -1,68 +1,11 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 import argparse
-import os
-import sys
-from pathlib import Path
 
-from vllm.benchmarks.serve import add_cli_args
-from vllm.benchmarks.serve import main as python_main
+from vllm.benchmarks.serve import add_cli_args, main
 from vllm.entrypoints.cli.benchmark.base import BenchmarkSubcommandBase
-from vllm.logger import init_logger
 from vllm.utils.argparse_utils import FlexibleArgumentParser
 
-logger = init_logger(__name__)
-_RUST_CLI_PATH = Path(__file__).resolve().parents[3] / "vllm-rs"
-_RUST_SUPPORTED_DATASETS = frozenset(
-    {
-        "custom",
-        "hf",
-        "prefix_repetition",
-        "random",
-        "random-mm",
-        "random-rerank",
-        "sharegpt",
-        "sonnet",
-        "speed_bench",
-    }
-)
-_RUST_SUPPORTED_BACKENDS = frozenset(
-    {
-        "openai",
-        "openai-chat",
-        "openai-embeddings",
-        "openai-embeddings-chat",
-        "vllm",
-        "vllm-pooling",
-        "vllm-rerank",
-    }
-)
-
-
-def _rust_unsupported_reason(args: argparse.Namespace) -> str | None:
-    if args.dataset_name not in _RUST_SUPPORTED_DATASETS:
-        return f"dataset {args.dataset_name!r} is not supported by the Rust benchmark"
-    if args.backend not in _RUST_SUPPORTED_BACKENDS:
-        return f"backend {args.backend!r} is not supported by the Rust benchmark"
-    return None
-
-
-def _maybe_exec_rust_bench(args: argparse.Namespace) -> None:
-    if reason := _rust_unsupported_reason(args):
-        logger.info("Using Python benchmark: %s.", reason)
-        return
-
-    if not _RUST_CLI_PATH.is_file():
-        logger.warning(
-            "Rust benchmark binary not found at %s; falling back to Python.",
-            _RUST_CLI_PATH,
-        )
-        return
-
-    rust_cli = str(_RUST_CLI_PATH)
-    logger.info("Delegating `vllm bench serve` to Rust binary at %s.", rust_cli)
-    os.execv(rust_cli, [rust_cli, "bench", "serve", *sys.argv[3:]])
-
 
 class BenchmarkServingSubcommand(BenchmarkSubcommandBase):
     """The `serve` subcommand for `vllm bench`."""
@@ -76,5 +19,4 @@ class BenchmarkServingSubcommand(BenchmarkSubcommandBase):
 
     @staticmethod
     def cmd(args: argparse.Namespace) -> None:
-        _maybe_exec_rust_bench(args)
-        python_main(args)
+        main(args)
diff --git a/vllm/entrypoints/cli/main.py b/vllm/entrypoints/cli/main.py
index fe0b339b3ed..3dc69dd3ad2 100644
--- a/vllm/entrypoints/cli/main.py
+++ b/vllm/entrypoints/cli/main.py
@@ -54,6 +54,8 @@ def main():
         logger.info("Delegating entrypoint handling to vllm-omni")
         omni_main()
     else:
+        vllm.entrypoints.cli.benchmark.main.maybe_exec_rust_bench()
+
         # For 'vllm bench *': use CPU instead of UnspecifiedPlatform by default
         if len(sys.argv) > 1 and sys.argv[1] == "bench":
             logger.debug(
diff --git a/vllm/entrypoints/cli/serve.py b/vllm/entrypoints/cli/serve.py
index d5e9b2bc874..08cb79f2081 100644
--- a/vllm/entrypoints/cli/serve.py
+++ b/vllm/entrypoints/cli/serve.py
@@ -58,6 +58,10 @@ class ServeSubcommand(CLISubcommand):
             uvloop.run(serve_grpc(args))
             return
 
+        rust_frontend_path = (
+            envs.VLLM_RUST_FRONTEND_PATH if envs.VLLM_USE_RUST_FRONTEND else None
+        )
+
         if args.headless:
             if args.api_server_count is not None and args.api_server_count > 0:
                 raise ValueError(
@@ -103,7 +107,7 @@ class ServeSubcommand(CLISubcommand):
         # - Hybrid LB: Use local DP size (internal LB for local ranks only)
         # - Internal LB: Use full DP size
         if args.api_server_count is None:
-            if is_multi_port or is_external_lb or envs.VLLM_RUST_FRONTEND_PATH:
+            if is_multi_port or is_external_lb or rust_frontend_path:
                 args.api_server_count = 1
             elif is_hybrid_lb:
                 args.api_server_count = args.data_parallel_size_local or 1
@@ -120,7 +124,7 @@ class ServeSubcommand(CLISubcommand):
                         "Defaulting api_server_count to data_parallel_size (%d).",
                         args.api_server_count,
                     )
-        elif envs.VLLM_RUST_FRONTEND_PATH and args.api_server_count > 1:
+        elif rust_frontend_path and args.api_server_count > 1:
             logger.warning(
                 "Ignoring --api-server-count=%d when using rust front-end process",
                 args.api_server_count,
@@ -140,7 +144,7 @@ class ServeSubcommand(CLISubcommand):
             run_dp_supervisor(args)
         elif args.api_server_count < 1:
             run_headless(args)
-        elif args.api_server_count > 1 or envs.VLLM_RUST_FRONTEND_PATH:
+        elif args.api_server_count > 1 or rust_frontend_path:
             run_multi_api_server(args)
         else:
             # Single API server (this process).
@@ -256,7 +260,9 @@ def run_headless(args: argparse.Namespace):
 
 def run_multi_api_server(args: argparse.Namespace):
     assert not args.headless
-    rust_frontend_path = envs.VLLM_RUST_FRONTEND_PATH
+    rust_frontend_path = (
+        envs.VLLM_RUST_FRONTEND_PATH if envs.VLLM_USE_RUST_FRONTEND else None
+    )
     num_api_servers: int = args.api_server_count
     assert num_api_servers > 0
 
diff --git a/vllm/entrypoints/llm.py b/vllm/entrypoints/llm.py
index b3205728e49..4274819d988 100644
--- a/vllm/entrypoints/llm.py
+++ b/vllm/entrypoints/llm.py
@@ -885,9 +885,19 @@ class LLM(BeamSearchOfflineMixin, PoolingOfflineMixin, OfflineInferenceMixin):
             "update_weights", kwargs={"update_info": update_info_dict}
         )
 
-    def finish_weight_update(self) -> None:
-        """Finish the current weight update."""
+    def finish_weight_update(self, weight_version: str | None = None) -> None:
+        """Finish the weight update and set its version if provided."""
         self.llm_engine.collective_rpc("finish_weight_update")
+        if weight_version is not None:
+            self.llm_engine.set_weight_version(weight_version)
+
+    def update_weight_version(self, new_version: str) -> None:
+        """Set the weight version without updating weights."""
+        self.llm_engine.set_weight_version(new_version)
+
+    def get_weight_version(self) -> str:
+        """Return the latest committed weight version."""
+        return self.llm_engine.get_weight_version()
 
     def __repr__(self) -> str:
         """Return a transformers-style hierarchical view of the model."""
diff --git a/vllm/entrypoints/openai/api_server.py b/vllm/entrypoints/openai/api_server.py
index 9103dd7fae9..f57e320a906 100644
--- a/vllm/entrypoints/openai/api_server.py
+++ b/vllm/entrypoints/openai/api_server.py
@@ -28,7 +28,6 @@ from vllm.engine.protocol import EngineClient
 from vllm.entrypoints.chat_utils import load_chat_template
 from vllm.entrypoints.launcher import serve_http
 from vllm.entrypoints.openai.cli_args import make_arg_parser, validate_parsed_serve_args
-from vllm.entrypoints.openai.engine.protocol import GenerationError
 from vllm.entrypoints.openai.models.protocol import BaseModelPath
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
 from vllm.entrypoints.serve.elastic_ep.middleware import ScalingMiddleware
@@ -42,20 +41,15 @@ from vllm.entrypoints.serve.utils.api_utils import (
 )
 from vllm.entrypoints.serve.utils.request_logger import RequestLogger
 from vllm.entrypoints.serve.utils.server_utils import (
-    engine_error_handler,
     exception_handler,
-    generation_error_handler,
     get_uvicorn_log_config,
     http_exception_handler,
     lifespan,
     log_response,
     validation_exception_handler,
+    vllm_error_handler,
 )
-from vllm.exceptions import (
-    VLLMNotFoundError,
-    VLLMUnprocessableEntityError,
-    VLLMValidationError,
-)
+from vllm.exceptions import VLLMError
 from vllm.logger import init_logger
 from vllm.reasoning import ReasoningParserManager
 from vllm.renderers.online_derenderer import OnlineDerenderer
@@ -67,7 +61,6 @@ from vllm.usage.usage_lib import UsageContext
 from vllm.utils.argparse_utils import FlexibleArgumentParser
 from vllm.utils.network_utils import is_valid_ipv6_address
 from vllm.utils.system_utils import decorate_logs, set_ulimit
-from vllm.v1.engine.exceptions import EngineDeadError, EngineGenerateError
 from vllm.version import __version__ as VLLM_VERSION
 
 prometheus_multiproc_dir: tempfile.TemporaryDirectory
@@ -291,23 +284,26 @@ def build_app(
         allow_headers=args.allowed_headers,
     )
 
+    # Exception handlers are registered in four layers:
+    #   1. framework errors raised by FastAPI/Starlette
+    #   2. vLLM-specific errors dispatched via a single ``VLLMError`` handler
+    #   3. fallback handlers for raw exceptions not yet migrated to ``VLLMError``
+    #   4. the raw ``Exception`` handler as a safety net
+    # Registering specific exception types (rather than only ``Exception``)
+    # ensures they are handled by ``ExceptionMiddleware`` (inside the Prometheus
+    # middleware) rather than ``ServerErrorMiddleware`` (outside it), so their
+    # status codes are recorded correctly.
     app.exception_handler(HTTPException)(http_exception_handler)
     app.exception_handler(RequestValidationError)(validation_exception_handler)
-    app.exception_handler(EngineGenerateError)(engine_error_handler)
-    app.exception_handler(EngineDeadError)(engine_error_handler)
-    app.exception_handler(GenerationError)(generation_error_handler)
-    # Register specific exception types so they are handled by
-    # ExceptionMiddleware (inside the Prometheus middleware) rather than
-    # ServerErrorMiddleware (outside it). Without this, these exceptions
-    # propagate through Prometheus as unhandled and get recorded as 5xx
-    # even though they result in 4xx responses to the client.
-    app.exception_handler(VLLMValidationError)(exception_handler)
-    app.exception_handler(VLLMUnprocessableEntityError)(exception_handler)
-    app.exception_handler(VLLMNotFoundError)(exception_handler)
+
+    app.exception_handler(VLLMError)(vllm_error_handler)
+
+    # TODO(zqzten): remove these fallback handlers after migration to VLLMError
     app.exception_handler(ValueError)(exception_handler)
     app.exception_handler(TypeError)(exception_handler)
     app.exception_handler(OverflowError)(exception_handler)
     app.exception_handler(NotImplementedError)(exception_handler)
+
     app.exception_handler(Exception)(exception_handler)
 
     # Ensure --api-key option from CLI takes precedence over VLLM_API_KEY
diff --git a/vllm/entrypoints/openai/dp_supervisor.py b/vllm/entrypoints/openai/dp_supervisor.py
index d669ec4d1d5..8ce6233c1ba 100644
--- a/vllm/entrypoints/openai/dp_supervisor.py
+++ b/vllm/entrypoints/openai/dp_supervisor.py
@@ -257,7 +257,7 @@ def _run_vllm_dp_server(child_args: argparse.Namespace) -> None:
     name = f"APIServer_DP{child_args.data_parallel_rank}"
     set_process_title(name)
     decorate_logs(name)
-    if envs.VLLM_RUST_FRONTEND_PATH:
+    if envs.VLLM_USE_RUST_FRONTEND and envs.VLLM_RUST_FRONTEND_PATH:
         _run_rust_vllm_dp_server(child_args)
     else:
         _run_python_vllm_dp_server(child_args)
diff --git a/vllm/entrypoints/openai/engine/protocol.py b/vllm/entrypoints/openai/engine/protocol.py
index 95190659e05..203cc4d81cc 100644
--- a/vllm/entrypoints/openai/engine/protocol.py
+++ b/vllm/entrypoints/openai/engine/protocol.py
@@ -20,7 +20,7 @@ from pydantic import (
 
 from vllm.config.utils import replace
 from vllm.entrypoints.chat_utils import make_tool_call_id
-from vllm.exceptions import VLLMValidationError
+from vllm.exceptions import VLLMServerError, VLLMValidationError
 from vllm.logger import init_logger
 from vllm.sampling_params import StructuredOutputsParams
 from vllm.utils import random_uuid
@@ -407,7 +407,7 @@ class DeltaMessage(OpenAIBaseModel):
         return data
 
 
-class GenerationError(Exception):
+class GenerationError(VLLMServerError):
     """raised when finish_reason indicates internal server error (500)"""
 
     def __init__(self, message: str = "Internal server error"):
diff --git a/vllm/entrypoints/serve/dev/rlhf/api_router.py b/vllm/entrypoints/serve/dev/rlhf/api_router.py
index 8a2494a59df..392fcf56747 100644
--- a/vllm/entrypoints/serve/dev/rlhf/api_router.py
+++ b/vllm/entrypoints/serve/dev/rlhf/api_router.py
@@ -5,7 +5,7 @@ import json
 from http import HTTPStatus
 from typing import Annotated
 
-from fastapi import APIRouter, FastAPI, HTTPException, Query, Request
+from fastapi import APIRouter, Body, FastAPI, HTTPException, Query, Request
 from fastapi.responses import JSONResponse
 
 from vllm.distributed.weight_transfer.base import (
@@ -203,11 +203,29 @@ async def update_weights(raw_request: Request):
 
 
 @router.post("/finish_weight_update")
-async def finish_weight_update(raw_request: Request):
-    await engine_client(raw_request).finish_weight_update()
+async def finish_weight_update(
+    raw_request: Request,
+    weight_version: Annotated[str | None, Body(embed=True)] = None,
+):
+    await engine_client(raw_request).finish_weight_update(weight_version)
     return JSONResponse(content={"message": "Weight update finished"})
 
 
+@router.post("/update_weight_version")
+async def update_weight_version(
+    raw_request: Request,
+    new_version: Annotated[str, Body(embed=True)],
+):
+    await engine_client(raw_request).update_weight_version(new_version)
+    return JSONResponse(content={"success": True, "new_version": new_version})
+
+
+@router.get("/weight_info")
+async def weight_info(raw_request: Request):
+    weight_version = await engine_client(raw_request).get_weight_version()
+    return JSONResponse(content={"weight_version": weight_version})
+
+
 @router.get("/get_world_size")
 async def get_world_size(
     raw_request: Request,
diff --git a/vllm/entrypoints/serve/elastic_ep/api_router.py b/vllm/entrypoints/serve/elastic_ep/api_router.py
index e711a257ddd..02a24250905 100644
--- a/vllm/entrypoints/serve/elastic_ep/api_router.py
+++ b/vllm/entrypoints/serve/elastic_ep/api_router.py
@@ -12,10 +12,7 @@ from vllm.engine.protocol import EngineClient
 from vllm.entrypoints.openai.engine.protocol import (
     ErrorResponse,
 )
-from vllm.entrypoints.serve.elastic_ep.middleware import (
-    get_scaling_elastic_ep,
-    set_scaling_elastic_ep,
-)
+from vllm.entrypoints.serve.elastic_ep.middleware import get_scaling_elastic_ep
 from vllm.entrypoints.serve.utils.api_utils import validate_json_request
 from vllm.logger import init_logger
 
@@ -64,8 +61,6 @@ async def scale_elastic_ep(raw_request: Request):
             status_code=400, detail="drain_timeout must be a positive integer"
         )
 
-    # Set scaling flag to prevent new requests
-    set_scaling_elastic_ep(True)
     client = engine_client(raw_request)
     try:
         await client.scale_elastic_ep(new_data_parallel_size, drain_timeout)
@@ -83,8 +78,6 @@ async def scale_elastic_ep(raw_request: Request):
     except Exception as e:
         logger.error("Scale failed: %s", e)
         raise HTTPException(status_code=500, detail="Scale failed") from e
-    finally:
-        set_scaling_elastic_ep(False)
 
 
 @router.post("/is_scaling_elastic_ep")
diff --git a/vllm/entrypoints/serve/utils/error_response.py b/vllm/entrypoints/serve/utils/error_response.py
index fc17a75c75a..2aa785c53bb 100644
--- a/vllm/entrypoints/serve/utils/error_response.py
+++ b/vllm/entrypoints/serve/utils/error_response.py
@@ -28,7 +28,9 @@ def create_error_response(
         )
 
         from vllm.exceptions import (
+            VLLMClientError,
             VLLMNotFoundError,
+            VLLMServerError,
             VLLMUnprocessableEntityError,
             VLLMValidationError,
         )
@@ -45,8 +47,23 @@ def create_error_response(
             err_type = "NotFoundError"
             status_code = HTTPStatus.NOT_FOUND
             param = None
+        elif isinstance(exc, VLLMClientError):
+            # Any other client-caused error defaults to 400.
+            err_type = "BadRequestError"
+            status_code = HTTPStatus.BAD_REQUEST
+            param = None
+        elif isinstance(exc, GenerationError):
+            err_type = "InternalServerError"
+            status_code = exc.status_code
+            param = None
+        elif isinstance(exc, VLLMServerError):
+            # Any other server-caused error defaults to 500.
+            err_type = "InternalServerError"
+            status_code = HTTPStatus.INTERNAL_SERVER_ERROR
+            param = None
+        # Fallback for raw exceptions not yet migrated to VLLMError.
+        # TODO(zqzten): remove these fallback handlers after migration to VLLMError
         elif isinstance(exc, (ValueError, TypeError, OverflowError)):
-            # Common validation errors from user input
             err_type = "BadRequestError"
             status_code = HTTPStatus.BAD_REQUEST
             param = None
@@ -54,10 +71,6 @@ def create_error_response(
             err_type = "NotImplementedError"
             status_code = HTTPStatus.NOT_IMPLEMENTED
             param = None
-        elif isinstance(exc, GenerationError):
-            err_type = "InternalServerError"
-            status_code = exc.status_code
-            param = None
         elif any(cls.__name__ == "TemplateError" for cls in type(exc).__mro__):
             # jinja2.TemplateError and its subclasses (avoid importing jinja2)
             err_type = "BadRequestError"
diff --git a/vllm/entrypoints/serve/utils/server_utils.py b/vllm/entrypoints/serve/utils/server_utils.py
index c6520658090..c60abfae9a2 100644
--- a/vllm/entrypoints/serve/utils/server_utils.py
+++ b/vllm/entrypoints/serve/utils/server_utils.py
@@ -31,7 +31,7 @@ from vllm.entrypoints.serve.utils.error_response import (
     create_error_response,
     sanitize_message,
 )
-from vllm.exceptions import VLLMValidationError
+from vllm.exceptions import VLLMError, VLLMValidationError
 from vllm.logger import init_logger
 from vllm.utils.gc_utils import freeze_gc_heap
 from vllm.v1.engine.exceptions import EngineDeadError, EngineGenerateError
@@ -325,6 +325,16 @@ async def log_response(request: Request, call_next):
     return response
 
 
+async def vllm_error_handler(req: Request, exc: VLLMError):
+    """Dispatch a vLLM-specific error to the appropriate handler."""
+    if isinstance(exc, (EngineGenerateError, EngineDeadError)):
+        return await engine_error_handler(req, exc)
+    elif isinstance(exc, GenerationError):
+        return await generation_error_handler(req, exc)
+    else:
+        return await exception_handler(req, exc)
+
+
 async def engine_error_handler(
     req: Request, exc: EngineDeadError | EngineGenerateError
 ):
diff --git a/vllm/envs.py b/vllm/envs.py
index 69a4ceef457..87984fe31a4 100755
--- a/vllm/envs.py
+++ b/vllm/envs.py
@@ -154,6 +154,7 @@ if TYPE_CHECKING:
     K_SCALE_CONSTANT: int = 200
     V_SCALE_CONSTANT: int = 100
     VLLM_USE_RUST_FRONTEND: bool = False
+    VLLM_USE_RUST_BENCH: bool = False
     VLLM_RUST_FRONTEND_PATH: str | None = "auto"
     VLLM_SERVER_DEV_MODE: bool = False
     VLLM_V1_OUTPUT_PROC_CHUNK_SIZE: int = 128
@@ -552,22 +553,24 @@ def _deprecated_triton_attn_use_td() -> None:
     return None
 
 
-def _resolve_rust_frontend_path() -> str | None:
-    """Resolve the Rust frontend binary path.
+def _resolve_rust_cli_path() -> str | None:
+    """Resolve the vllm-rs binary path.
 
-    Returns None if VLLM_USE_RUST_FRONTEND is not enabled.
+    Returns None unless VLLM_USE_RUST_FRONTEND or VLLM_USE_RUST_BENCH is enabled.
     When enabled, resolves VLLM_RUST_FRONTEND_PATH ("auto" by default)
     to the actual binary path.
     """
-    use_rust = bool(int(os.environ.get("VLLM_USE_RUST_FRONTEND", "0")))
+    use_rust = bool(int(os.environ.get("VLLM_USE_RUST_FRONTEND", "0"))) or bool(
+        int(os.environ.get("VLLM_USE_RUST_BENCH", "0"))
+    )
     raw = os.environ.get("VLLM_RUST_FRONTEND_PATH", "auto")
 
     if not use_rust:
         if os.environ.get("VLLM_RUST_FRONTEND_PATH") is not None:
             logger.warning(
-                "VLLM_RUST_FRONTEND_PATH is set but VLLM_USE_RUST_FRONTEND "
-                "is not enabled. The Rust frontend will not be used. "
-                "Set VLLM_USE_RUST_FRONTEND=1 to enable it."
+                "VLLM_RUST_FRONTEND_PATH is set without enabling "
+                "VLLM_USE_RUST_FRONTEND or VLLM_USE_RUST_BENCH. "
+                "Set one of them to 1 to use the vllm-rs binary."
             )
         return None
 
@@ -1349,10 +1352,12 @@ environment_variables: dict[str, Callable[[], Any]] = {
     "VLLM_USE_RUST_FRONTEND": lambda: bool(
         int(os.getenv("VLLM_USE_RUST_FRONTEND", "0"))
     ),
-    # Path to the Rust frontend binary. Defaults to "auto" which discovers
-    # the binary installed with the vllm package. Only used when
-    # VLLM_USE_RUST_FRONTEND=1.
-    "VLLM_RUST_FRONTEND_PATH": lambda: _resolve_rust_frontend_path(),
+    # If set, use the packaged Rust client for `vllm bench serve`.
+    "VLLM_USE_RUST_BENCH": lambda: bool(int(os.getenv("VLLM_USE_RUST_BENCH", "0"))),
+    # Path to the vllm-rs binary. Defaults to "auto" which discovers the
+    # binary installed with the vllm package. Used when VLLM_USE_RUST_FRONTEND=1
+    # or VLLM_USE_RUST_BENCH=1.
+    "VLLM_RUST_FRONTEND_PATH": lambda: _resolve_rust_cli_path(),
     # If set, vllm will run in development mode, which will enable
     # some additional endpoints for developing and debugging,
     # e.g. `/reset_prefix_cache`
@@ -1561,7 +1566,7 @@ environment_variables: dict[str, Callable[[], Any]] = {
     # tensors above will instead be sent via a separate message.
     # While the sending side still actually copies the tensor
     # in all cases, on the receiving side, tensors above this
-    # limit will actually be zero-copy decoded.
+    # limit will actually be zero-copy decoded. The unit is bytes.
     "VLLM_MSGPACK_ZERO_COPY_THRESHOLD": lambda: int(
         os.getenv("VLLM_MSGPACK_ZERO_COPY_THRESHOLD", "256")
     ),
diff --git a/vllm/exceptions.py b/vllm/exceptions.py
index 4112c3de24b..4383e9a6441 100644
--- a/vllm/exceptions.py
+++ b/vllm/exceptions.py
@@ -6,7 +6,25 @@
 from typing import Any
 
 
-class VLLMValidationError(ValueError):
+class VLLMError(Exception):
+    """Base class for all vLLM-specific errors.
+
+    Subclasses are split into `VLLMClientError` (caused by the request, mapped
+    to 4xx) and `VLLMServerError` (caused by the server, mapped to 5xx).
+    Dispatching on this hierarchy lets the entrypoints decide the HTTP status
+    without relying on raw Python exception types such as `ValueError`.
+    """
+
+
+class VLLMClientError(VLLMError):
+    """Base class for errors caused by the client request (4xx)."""
+
+
+class VLLMServerError(VLLMError):
+    """Base class for errors caused by the server (5xx)."""
+
+
+class VLLMValidationError(VLLMClientError):
     """vLLM-specific validation error for request validation failures.
 
     Args:
@@ -36,7 +54,7 @@ class VLLMValidationError(ValueError):
         return f"{base} ({', '.join(extras)})" if extras else base
 
 
-class VLLMNotFoundError(Exception):
+class VLLMNotFoundError(VLLMClientError):
     """vLLM-specific NotFoundError"""
 
     pass
@@ -66,7 +84,7 @@ class LoRAAdapterNotFoundError(VLLMNotFoundError):
         return self.message
 
 
-class VLLMUnprocessableEntityError(ValueError):
+class VLLMUnprocessableEntityError(VLLMClientError):
     """vLLM-specific error for unprocessable entity requests.
 
     This exception is raised when the request content is invalid or cannot be
diff --git a/vllm/inputs/engine.py b/vllm/inputs/engine.py
index f997004d2fb..bda40bebbc8 100644
--- a/vllm/inputs/engine.py
+++ b/vllm/inputs/engine.py
@@ -7,6 +7,8 @@ from typing import TYPE_CHECKING, Literal, TypeAlias
 
 from typing_extensions import NotRequired, TypedDict, assert_never
 
+from vllm.exceptions import VLLMValidationError
+
 if TYPE_CHECKING:
     import torch
 
@@ -284,7 +286,7 @@ which can be passed to `LLMEngine.add_request` or `AsyncLLM.add_request`.
 
 def _validate_enc_input(enc_input: SingletonInput) -> EncoderInput:
     if enc_input["type"] == "embeds":
-        raise ValueError(
+        raise VLLMValidationError(
             "Embedding inputs are not supported for encoder-decoder models"
         )
 
@@ -302,7 +304,7 @@ def _validate_enc_input(enc_input: SingletonInput) -> EncoderInput:
 
 def _validate_dec_input(dec_input: SingletonInput) -> DecoderEngineInput:
     if dec_input["type"] == "embeds":
-        raise ValueError(
+        raise VLLMValidationError(
             "Embedding inputs are not supported for encoder-decoder models"
         )
 
diff --git a/vllm/model_executor/kernels/linear/scaled_mm/aiter.py b/vllm/model_executor/kernels/linear/scaled_mm/aiter.py
index 1b39491ab34..da8f69fa97b 100644
--- a/vllm/model_executor/kernels/linear/scaled_mm/aiter.py
+++ b/vllm/model_executor/kernels/linear/scaled_mm/aiter.py
@@ -9,6 +9,9 @@ from vllm._aiter_ops import (
     rocm_aiter_ops,
 )
 from vllm.logger import init_logger
+from vllm.model_executor.layers.quantization.utils.fp8_utils import (
+    _upcast_e8m0_to_fp32,
+)
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
     GroupShape,
 )
@@ -16,6 +19,7 @@ from vllm.model_executor.utils import replace_parameter
 from vllm.platforms import current_platform
 
 from .BlockScaledMMLinearKernel import (
+    FP8BlockParams,
     Fp8BlockScaledMMLinearKernel,
 )
 from .cutlass import CutlassInt8ScaledMMLinearKernel
@@ -375,6 +379,17 @@ class AiterFp8BlockScaledMMKernel(Fp8BlockScaledMMLinearKernel):
             and rocm_aiter_ops.is_triton_gemm_w8a8_tuned(n, k)
         )
 
+    def process_weights_after_loading(self, layer: torch.nn.Module) -> None:
+        super().process_weights_after_loading(layer)
+
+        params = FP8BlockParams.from_layer(layer)
+        if params.weight_scale_inv is not None:
+            ws, attr = params.weight_scale_inv, params.WEIGHT_SCALE_INV
+        else:
+            ws, attr = params.weight_scale, params.WEIGHT_SCALE
+        if ws is not None and ws.dtype == torch.float8_e8m0fnu:
+            replace_parameter(layer, attr, _upcast_e8m0_to_fp32(ws).contiguous())
+
     @classmethod
     def is_supported(cls, compute_capability=None):
         return (
@@ -406,19 +421,12 @@ class AiterFp8BlockScaledMMKernel(Fp8BlockScaledMMLinearKernel):
         Bs: torch.Tensor,
     ) -> torch.Tensor:
         if As.dtype != Bs.dtype:
-            from vllm.model_executor.layers.quantization.utils.fp8_utils import (
-                _upcast_e8m0_to_fp32,
-            )
-
             if As.dtype == torch.float8_e8m0fnu:
                 As = _upcast_e8m0_to_fp32(As).contiguous()
             else:
                 As = As.to(torch.float32)
 
-            if Bs.dtype == torch.float8_e8m0fnu:
-                Bs = _upcast_e8m0_to_fp32(Bs).contiguous()
-            else:
-                Bs = Bs.to(torch.float32)
+            Bs = Bs.to(torch.float32)
 
         out_dtype = self.config.out_dtype
         if self.use_triton:
diff --git a/vllm/model_executor/layers/attention/mla_attention.py b/vllm/model_executor/layers/attention/mla_attention.py
index 71c4def404c..4a265554b77 100644
--- a/vllm/model_executor/layers/attention/mla_attention.py
+++ b/vllm/model_executor/layers/attention/mla_attention.py
@@ -772,13 +772,8 @@ class MLAAttention(nn.Module, AttentionLayerBase):
         num_mha_tokens = q.size(0) - num_mqa_tokens
 
         if self.impl.is_sparse and num_mha_tokens > 0:
-            prefill_max_seq_len = attn_metadata.prefill_max_seq_len  # type: ignore[attr-defined]
-            use_mha = (
-                self.prefill_backend is not None
-                and prefill_max_seq_len <= attn_metadata.topk_tokens  # type: ignore[attr-defined]
-                and not self._vllm_config.attention_config.sparse_mla_force_mqa
-            )
-            if not use_mha:
+            prefill_metadata = getattr(attn_metadata, "prefill", None)
+            if not getattr(prefill_metadata, "use_dense_mha", False):
                 num_mqa_tokens = q.size(0)
                 num_mha_tokens = 0
 
@@ -1413,6 +1408,10 @@ class MLACommonPrefillMetadata:
     q_data_type: torch.dtype | None = None
     output_dtype: torch.dtype | None = None
     prefill_backend: MLAPrefillBackend | None = None
+    # Whether the prefill suffix is routed through dense MHA.
+    # Indexer scoring may be skipped only for a pure-prefill batch,
+    # since decode tokens still consume top-k indices.
+    use_dense_mha: bool = False
 
 
 @dataclass
diff --git a/vllm/model_executor/layers/attention/sparse_mla_attention.py b/vllm/model_executor/layers/attention/sparse_mla_attention.py
index 19cad7986bf..1463f9fb35b 100644
--- a/vllm/model_executor/layers/attention/sparse_mla_attention.py
+++ b/vllm/model_executor/layers/attention/sparse_mla_attention.py
@@ -203,6 +203,10 @@ class SparseMLACommonMetadataBuilder(AttentionMetadataBuilder[T]):
                 q_data_type=self.model_config.dtype,
                 output_dtype=self.model_config.dtype,
                 prefill_backend=self._prefill_backend,
+                use_dense_mha=(
+                    prefill_max_seq_len <= self.topk_tokens
+                    and not self.vllm_config.attention_config.sparse_mla_force_mqa
+                ),
             )
             self._prefill_backend.prepare_metadata(prefill)
 
diff --git a/vllm/model_executor/layers/fused_moe/activation.py b/vllm/model_executor/layers/fused_moe/activation.py
index 9f8b8971ff2..baa6136485e 100644
--- a/vllm/model_executor/layers/fused_moe/activation.py
+++ b/vllm/model_executor/layers/fused_moe/activation.py
@@ -171,11 +171,20 @@ def apply_moe_activation(
         # Fused CUDA kernel: writes straight to `output`, no fp32 temporaries.
         # (The pure-torch fallback below upcast both halves to fp32 and
         # allocated ~8 temporaries per call, blowing up MoE memory.)
-        # linear_beta <= 0 signals "unset" to the kernel (up passed through).
-        beta = 1.0 if activation_situ_beta is None else activation_situ_beta
-        linear_beta = activation_situ_linear_beta
+        # Both betas come from FusedMoEConfig; a missing beta means the caller
+        # bypassed the config plumbing, so fail rather than silently use 1.0.
+        # linear_beta is genuinely optional: <= 0 signals "unset" to the kernel
+        # (up passed through), matching SituAndMul(linear_beta=None).
+        assert activation_situ_beta is not None, (
+            "SITU requires activation_situ_beta from FusedMoEConfig"
+        )
         torch.ops._C.situ_and_mul(
-            output, input, beta, -1.0 if linear_beta is None else linear_beta
+            output,
+            input,
+            activation_situ_beta,
+            -1.0
+            if activation_situ_linear_beta is None
+            else activation_situ_linear_beta,
         )
     elif activation == MoEActivation.SWIGLUOAI:
         torch.ops._C.swigluoai_and_mul(output, input)
diff --git a/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp8_moe.py b/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp8_moe.py
index c5330f3b438..59115df13b5 100644
--- a/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp8_moe.py
+++ b/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp8_moe.py
@@ -6,10 +6,13 @@
 ``convert_to_fp8_moe_kernel_format``.
 """
 
+import math
+
 import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm.logger import init_logger
+from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.experts.mxfp8_emulation_moe import (
     Mxfp8TritonExpertsBase,
 )
@@ -17,6 +20,9 @@ from vllm.platforms import current_platform
 
 logger = init_logger(__name__)
 
+_AITER_SWIGLU_ALPHA = 1.702
+_AITER_SWIGLU_BETA = 1.0
+
 
 def is_aiter_mxfp8_moe_available() -> bool:
     """True when the FlyDSL MXFP8 MoE can run here: gfx950, the ``flydsl``
@@ -93,6 +99,27 @@ class AiterMxfp8Experts(Mxfp8TritonExpertsBase):
             return False, (
                 "kernel requires the aiter flydsl package, which is not installed"
             )
+        if (
+            is_supported
+            and moe_config.activation != MoEActivation.SWIGLUOAI_UNINTERLEAVE
+        ):
+            return False, (
+                "kernel hardcodes SwiGLU-OAI activation and requires "
+                f"activation={MoEActivation.SWIGLUOAI_UNINTERLEAVE.value}; "
+                f"got activation={moe_config.activation.value}"
+            )
+        if is_supported and (
+            moe_config.swiglu_alpha is None
+            or not math.isclose(float(moe_config.swiglu_alpha), _AITER_SWIGLU_ALPHA)
+            or moe_config.swiglu_beta is None
+            or not math.isclose(float(moe_config.swiglu_beta), _AITER_SWIGLU_BETA)
+        ):
+            return False, (
+                "kernel hardcodes SwiGLU-OAI with "
+                f"alpha={_AITER_SWIGLU_ALPHA} and beta={_AITER_SWIGLU_BETA}; "
+                f"got swiglu_alpha={moe_config.swiglu_alpha} and "
+                f"swiglu_beta={moe_config.swiglu_beta}"
+            )
         return is_supported, reason
 
     def apply(
diff --git a/vllm/model_executor/layers/fused_moe/experts/flashinfer_cutedsl_moe.py b/vllm/model_executor/layers/fused_moe/experts/flashinfer_cutedsl_moe.py
index b512d51c135..823ae43b5ef 100644
--- a/vllm/model_executor/layers/fused_moe/experts/flashinfer_cutedsl_moe.py
+++ b/vllm/model_executor/layers/fused_moe/experts/flashinfer_cutedsl_moe.py
@@ -13,6 +13,9 @@ from vllm.model_executor.layers.fused_moe.config import (
 from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
     TopKWeightAndReduceNoOP,
 )
+from vllm.model_executor.layers.quantization.utils.flashinfer_utils import (
+    activation_to_flashinfer_int,
+)
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
     QuantKey,
     kNvfp4Dynamic,
@@ -76,7 +79,7 @@ class FlashInferCuteDSLExperts(mk.FusedMoEExpertsModular):
 
     @staticmethod
     def _supports_no_act_and_mul() -> bool:
-        return False
+        return True
 
     @staticmethod
     def _supports_quant_scheme(
@@ -90,7 +93,7 @@ class FlashInferCuteDSLExperts(mk.FusedMoEExpertsModular):
 
     @staticmethod
     def _supports_activation(activation: MoEActivation) -> bool:
-        return activation == MoEActivation.SILU
+        return activation in (MoEActivation.SILU, MoEActivation.RELU2_NO_MUL)
 
     @staticmethod
     def _supports_parallel_config(
@@ -163,4 +166,5 @@ class FlashInferCuteDSLExperts(mk.FusedMoEExpertsModular):
             num_local_experts=self.local_num_experts,
             local_expert_offset=self.local_expert_offset,
             moe_output=output,
+            activation_type=activation_to_flashinfer_int(activation),
         )
diff --git a/vllm/model_executor/layers/fused_moe/experts/marlin_moe.py b/vllm/model_executor/layers/fused_moe/experts/marlin_moe.py
index cfd21a7c4e6..df596ff5d1c 100644
--- a/vllm/model_executor/layers/fused_moe/experts/marlin_moe.py
+++ b/vllm/model_executor/layers/fused_moe/experts/marlin_moe.py
@@ -157,6 +157,8 @@ def _fused_marlin_moe(
     clamp_limit: float | None = None,
     gemm1_alpha: float = 1.0,
     gemm1_beta: float = 0.0,
+    activation_situ_beta: float | None = None,
+    activation_situ_linear_beta: float | None = None,
 ) -> torch.Tensor:
     assert hidden_states.ndim == 2
     M, K = hidden_states.size()
@@ -250,6 +252,8 @@ def _fused_marlin_moe(
         beta=gemm1_beta,
         topk_ids=topk_ids,
         expert_map=expert_map,
+        activation_situ_beta=activation_situ_beta,
+        activation_situ_linear_beta=activation_situ_linear_beta,
     )
 
     if output is None:
@@ -337,6 +341,8 @@ def fused_marlin_moe(
     clamp_limit: float | None = None,
     gemm1_alpha: float = 1.0,
     gemm1_beta: float = 0.0,
+    activation_situ_beta: float | None = None,
+    activation_situ_linear_beta: float | None = None,
 ) -> torch.Tensor:
     """
     This function computes a Mixture of Experts (MoE) layer using two sets of
@@ -438,6 +444,8 @@ def fused_marlin_moe(
         num_tokens_post_padded=num_tokens_post_padded,
         activation=activation,
         activation_func=activation_func,
+        activation_situ_beta=activation_situ_beta,
+        activation_situ_linear_beta=activation_situ_linear_beta,
         input_global_scale1=input_global_scale1,
         input_global_scale2=input_global_scale2,
         global_scale1=global_scale1,
@@ -505,6 +513,8 @@ def batched_fused_marlin_moe(
     gemm1_alpha: float = 1.0,
     gemm1_beta: float = 0.0,
     activation_func: Callable[..., None] = apply_moe_activation,
+    activation_situ_beta: float | None = None,
+    activation_situ_linear_beta: float | None = None,
 ) -> torch.Tensor:
     """
     This function massages the inputs so the batched hidden_states can be
@@ -613,6 +623,8 @@ def batched_fused_marlin_moe(
         apply_router_weight_on_input=apply_router_weight_on_input,
         activation=activation,
         activation_func=activation_func,
+        activation_situ_beta=activation_situ_beta,
+        activation_situ_linear_beta=activation_situ_linear_beta,
         expert_map=expert_map,
         block_size_m=block_size_m,
         sorted_token_ids=sorted_token_ids,
@@ -877,6 +889,10 @@ class MarlinExperts(LoRAExpertsMixin, MarlinExpertsBase):
                 global_num_experts=global_num_experts,
                 activation=activation,
                 activation_func=self.activation,
+                activation_situ_beta=self.moe_config.activation_situ_beta,
+                activation_situ_linear_beta=(
+                    self.moe_config.activation_situ_linear_beta
+                ),
                 moe_sum=self.moe_sum,
                 expert_map=expert_map,
                 output=output,
@@ -917,6 +933,8 @@ class MarlinExperts(LoRAExpertsMixin, MarlinExpertsBase):
             beta: float = 0.0,
             topk_ids: torch.Tensor | None = None,
             expert_map: torch.Tensor | None = None,
+            activation_situ_beta: float | None = None,
+            activation_situ_linear_beta: float | None = None,
         ) -> None:
             # act_input  = intermediate_cache1 (M*topk, 2N for gated)
             # act_output = intermediate_cache2 (M*topk, N)
@@ -955,6 +973,8 @@ class MarlinExperts(LoRAExpertsMixin, MarlinExpertsBase):
                 beta=beta,
                 topk_ids=topk_ids,
                 expert_map=expert_map,
+                activation_situ_beta=activation_situ_beta,
+                activation_situ_linear_beta=activation_situ_linear_beta,
             )
             lora_state["cache2"] = act_output
 
@@ -1002,6 +1022,8 @@ class MarlinExperts(LoRAExpertsMixin, MarlinExpertsBase):
             global_num_experts=global_num_experts,
             activation=activation,
             activation_func=activation_with_lora,
+            activation_situ_beta=self.moe_config.activation_situ_beta,
+            activation_situ_linear_beta=self.moe_config.activation_situ_linear_beta,
             moe_sum=moe_sum_with_lora,
             expert_map=expert_map,
             output=output,
@@ -1110,20 +1132,33 @@ class BatchedMarlinExperts(MarlinExpertsBase):
             act: MoEActivation,
             act_output: torch.Tensor,
             act_input: torch.Tensor,
+            *,
+            activation_situ_beta: float | None = None,
+            activation_situ_linear_beta: float | None = None,
             **kwargs,
         ) -> None:
             if act != MoEActivation.SITU:
-                self.activation(act, act_output, act_input, **kwargs)
+                self.activation(
+                    act,
+                    act_output,
+                    act_input,
+                    activation_situ_beta=activation_situ_beta,
+                    activation_situ_linear_beta=activation_situ_linear_beta,
+                    **kwargs,
+                )
                 return
 
             num_experts, max_num_tokens = hidden_states.shape[:2]
-            beta = self.moe_config.activation_situ_beta
-            linear_beta = self.moe_config.activation_situ_linear_beta
+            beta = activation_situ_beta
+            linear_beta = activation_situ_linear_beta
+            assert beta is not None, (
+                "SITU requires activation_situ_beta from FusedMoEConfig"
+            )
             torch.ops._C.masked_situ_and_mul(
                 act_output.view(num_experts, max_num_tokens, -1),
                 act_input.view(num_experts, max_num_tokens, -1),
                 expert_tokens_meta.expert_num_tokens,
-                1.0 if beta is None else beta,
+                beta,
                 -1.0 if linear_beta is None else linear_beta,
             )
 
@@ -1158,4 +1193,6 @@ class BatchedMarlinExperts(MarlinExpertsBase):
             gemm1_alpha=self.gemm1_alpha,
             gemm1_beta=self.gemm1_beta,
             activation_func=activation_func,
+            activation_situ_beta=self.moe_config.activation_situ_beta,
+            activation_situ_linear_beta=self.moe_config.activation_situ_linear_beta,
         )
diff --git a/vllm/model_executor/layers/fused_moe/experts/mxfp8_emulation_moe.py b/vllm/model_executor/layers/fused_moe/experts/mxfp8_emulation_moe.py
index 71dd7634a69..ad6083251fb 100644
--- a/vllm/model_executor/layers/fused_moe/experts/mxfp8_emulation_moe.py
+++ b/vllm/model_executor/layers/fused_moe/experts/mxfp8_emulation_moe.py
@@ -107,17 +107,13 @@ class Mxfp8EmulationTritonExperts(Mxfp8TritonExpertsBase):
             limit = self.quant_config.gemm1_clamp_limit
             if limit is None:
                 raise ValueError("SWIGLUOAI_UNINTERLEAVE requires gemm1_clamp_limit")
-            alpha = self.quant_config.gemm1_alpha
-            alpha = 1.702 if alpha is None else float(alpha)
-            beta = self.quant_config.gemm1_beta
-            beta = 1.0 if beta is None else float(beta)
             apply_moe_activation(
                 activation,
                 output,
                 input,
                 clamp_limit=float(limit),
-                alpha=alpha,
-                beta=beta,
+                alpha=self.gemm1_alpha,
+                beta=self.gemm1_beta,
             )
             return
         super().activation(activation, output, input)
diff --git a/vllm/model_executor/layers/fused_moe/experts/mxfp8_native_moe.py b/vllm/model_executor/layers/fused_moe/experts/mxfp8_native_moe.py
index 9839756880a..e8c7dc96921 100644
--- a/vllm/model_executor/layers/fused_moe/experts/mxfp8_native_moe.py
+++ b/vllm/model_executor/layers/fused_moe/experts/mxfp8_native_moe.py
@@ -362,10 +362,7 @@ class Mxfp8NativeTritonExperts(Mxfp8TritonExpertsBase):
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
         apply_router_weight_on_input: bool,
     ):
-        alpha = self.quant_config.gemm1_alpha
-        alpha = 1.702 if alpha is None else float(alpha)
-        beta = self.quant_config.gemm1_beta
-        beta = 1.0 if beta is None else float(beta)
+        # `self.gemm1_alpha` and `self.gemm1_beta`` are set by `TritonExperts.__init__`.
         limit = self.quant_config.gemm1_clamp_limit
         limit = None if limit is None else float(limit)
         out = fused_moe_mxfp8_native(
@@ -376,8 +373,8 @@ class Mxfp8NativeTritonExperts(Mxfp8TritonExpertsBase):
             self.w2_scale_val,
             topk_weights,
             topk_ids,
-            alpha=alpha,
-            beta=beta,
+            alpha=self.gemm1_alpha,
+            beta=self.gemm1_beta,
             limit=limit,
             global_num_experts=global_num_experts,
             expert_map=expert_map,
diff --git a/vllm/model_executor/layers/fused_moe/experts/trtllm_mxfp4_moe.py b/vllm/model_executor/layers/fused_moe/experts/trtllm_mxfp4_moe.py
index 7b2a6f807c6..add3eb98818 100644
--- a/vllm/model_executor/layers/fused_moe/experts/trtllm_mxfp4_moe.py
+++ b/vllm/model_executor/layers/fused_moe/experts/trtllm_mxfp4_moe.py
@@ -14,6 +14,7 @@ from vllm.model_executor.layers.fused_moe.config import (
 from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
     TopKWeightAndReduceNoOP,
 )
+from vllm.model_executor.layers.fused_moe.utils import trtllm_moe_pack_topk_ids_weights
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
     QuantKey,
     kMxfp4Static,
@@ -78,7 +79,7 @@ class TrtLlmMxfp4ExpertsBase:
         else:
             self.gemm1_clamp_limit = None
 
-        # SITU (Kimi SituGLU) TRTLLM-Gen kernel computes
+        # SITU (SituGLU) TRTLLM-Gen kernel computes
         #   left  = alpha * tanh(x0 / alpha) * sigmoid(x0)   # gate (x0)
         #   right = beta  * tanh(x1 / beta)                  # up   (x1)
         # which matches vLLM's situ_and_mul with (beta, linear_beta), so map
@@ -261,7 +262,6 @@ class TrtLlmMxfp4ExpertsMonolithic(
             routing_method_type=self.routing_method_type,
             do_finalize=True,
             activation_type=self._flashinfer_activation_type(activation),
-            is_private=True,
             tune_max_num_tokens=max(self.max_capture_size, 1),
             output=output,
             routing_replay_out=routing_replay_out,
@@ -348,8 +348,9 @@ class TrtLlmMxfp4ExpertsModular(TrtLlmMxfp4ExpertsBase, mk.FusedMoEExpertsModula
     ) -> None:
         from flashinfer import trtllm_fp4_block_scale_routed_moe
 
+        packed_tensor = trtllm_moe_pack_topk_ids_weights(topk_ids, topk_weights)
         trtllm_fp4_block_scale_routed_moe(
-            topk_ids=(topk_ids, topk_weights),
+            topk_ids=packed_tensor,
             routing_bias=None,
             hidden_states=x_quant,
             hidden_states_scale=x_scale,
@@ -379,7 +380,6 @@ class TrtLlmMxfp4ExpertsModular(TrtLlmMxfp4ExpertsBase, mk.FusedMoEExpertsModula
             do_finalize=True,
             enable_pdl=True,
             activation_type=self._flashinfer_activation_type(activation),
-            is_private=True,
             output=output,
             tune_max_num_tokens=max(self.max_capture_size, 1),
         )
diff --git a/vllm/model_executor/layers/fused_moe/fused_moe.py b/vllm/model_executor/layers/fused_moe/fused_moe.py
index 0d357fcbf45..be4930052a9 100644
--- a/vllm/model_executor/layers/fused_moe/fused_moe.py
+++ b/vllm/model_executor/layers/fused_moe/fused_moe.py
@@ -28,9 +28,12 @@ from vllm.model_executor.layers.fused_moe.moe_align_block_size import (
 from vllm.model_executor.layers.fused_moe.utils import (
     enable_swap_ab,
     moe_kernel_quantize_input,
+    resolve_moe_use_td,
+    warn_if_moe_use_td_ineffective,
 )
 from vllm.platforms import current_platform
 from vllm.triton_utils import tl, triton
+from vllm.triton_utils.allocation import set_triton_allocator
 from vllm.utils.math_utils import next_power_of_2
 from vllm.utils.platform_utils import get_device_name_as_file_name
 from vllm.utils.torch_utils import direct_register_custom_op
@@ -347,6 +350,8 @@ def fused_moe_kernel(
     per_channel_quant: tl.constexpr,
     HAS_BIAS: tl.constexpr,
     SWAP_AB: tl.constexpr,
+    # Tensor-descriptor path for the A gather and B load in the K-loop.
+    USE_TD: tl.constexpr = False,
 ):
     """
     Implements the fused computation for a Mixture of Experts (MOE) using
@@ -436,7 +441,25 @@ def fused_moe_kernel(
 
     offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N).to(tl.int64)) % N
     offs_k = tl.arange(0, BLOCK_SIZE_K)
-    if SWAP_AB:
+    # TD gather and the SWAP_AB accumulator layout are mutually exclusive.
+    tl.static_assert(not (USE_TD and SWAP_AB))
+    if USE_TD:
+        # ``tt.descriptor_gather`` requires block_shape[0] == 1 and i32 idx.
+        m_td = num_valid_tokens // top_k
+        a_desc = tl.make_tensor_descriptor(
+            base=a_ptr,
+            shape=(m_td, K),
+            strides=(stride_am, stride_ak),
+            block_shape=(1, BLOCK_SIZE_K),
+        )
+        b_desc = tl.make_tensor_descriptor(
+            base=b_ptr + off_experts * stride_be,
+            shape=(N, K),
+            strides=(stride_bn, stride_bk),
+            block_shape=(BLOCK_SIZE_N, BLOCK_SIZE_K),
+        )
+        gather_idx = (offs_token // top_k).to(tl.int32)
+    elif SWAP_AB:
         a_ptrs = a_ptr + (
             offs_k[:, None] * stride_ak + offs_token[None, :] // top_k * stride_am
         )
@@ -454,7 +477,6 @@ def fused_moe_kernel(
             + off_experts * stride_be
             + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
         )
-
     if use_int8_w8a16:
         b_scale_ptrs = (
             b_scale_ptr + off_experts * stride_bse + offs_bn[None, :] * stride_bsn
@@ -498,18 +520,21 @@ def fused_moe_kernel(
     for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
         # Load the next block of A and B, generate a mask by checking the
         # K dimension.
-        if SWAP_AB:
+        if USE_TD:
+            a = a_desc.gather(gather_idx, k * BLOCK_SIZE_K)
+            b = b_desc.load([pid_n * BLOCK_SIZE_N, k * BLOCK_SIZE_K]).T
+        elif SWAP_AB:
             a_mask = (offs_k[:, None] < K - k * BLOCK_SIZE_K) & token_mask[None, :]
             b_mask = offs_k[None, :] < K - k * BLOCK_SIZE_K
+            a = tl.load(a_ptrs, mask=a_mask, other=0.0)
+            b = tl.load(b_ptrs, mask=b_mask, other=0.0)
         else:
-            a_mask = token_mask[:, None] & (offs_k[None, :] < K - k * BLOCK_SIZE_K)
-            b_mask = offs_k[:, None] < K - k * BLOCK_SIZE_K
-        a = tl.load(
-            a_ptrs,
-            mask=a_mask,
-            other=0.0,
-        )
-        b = tl.load(b_ptrs, mask=b_mask, other=0.0)
+            a = tl.load(
+                a_ptrs,
+                mask=token_mask[:, None] & (offs_k[None, :] < K - k * BLOCK_SIZE_K),
+                other=0.0,
+            )
+            b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)
         # We accumulate along the K dimension.
         if use_int8_w8a16:
             accumulator = tl.dot(a, b.to(compute_type), acc=accumulator)
@@ -536,9 +561,10 @@ def fused_moe_kernel(
                     accumulator += tl.dot(a, b)
         else:
             accumulator += tl.dot(a, b)
-        # Advance the ptrs to the next K block.
-        a_ptrs += BLOCK_SIZE_K * stride_ak
-        b_ptrs += BLOCK_SIZE_K * stride_bk
+        if not USE_TD:
+            # Advance the ptrs to the next K block.
+            a_ptrs += BLOCK_SIZE_K * stride_ak
+            b_ptrs += BLOCK_SIZE_K * stride_bk
 
     if SWAP_AB:
         accumulator = tl.trans(accumulator, (1, 0))
@@ -765,6 +791,19 @@ def invoke_fused_moe_triton_kernel(
     else:
         SWAP_AB = False
 
+    # Quantized weights always carry a B_scale (see the asserts below); key off
+    # that rather than enumerating quant flags, which misses w8a16-fp8/nvfp4/etc.
+    is_quantized = B_scale is not None
+    warn_if_moe_use_td_ineffective("TRITON", is_quantized=is_quantized)
+
+    # TD path is unvalidated under quantization; fall back to the pointer path.
+    use_td = resolve_moe_use_td() and not is_quantized
+    if use_td:
+        # The TD path builds a tensor descriptor inside the kernel, which
+        # requires a PyTorch-backed scratch allocator to be registered
+        # (Triton raises "no allocator was set" otherwise on CUDA).
+        set_triton_allocator(A.device)
+
     if use_fp8_w8a8 or use_int8_w8a8:
         assert B_scale is not None
         assert block_shape is None or triton.cdiv(
@@ -805,6 +844,19 @@ def invoke_fused_moe_triton_kernel(
     BLOCK_SIZE_K = config.pop("BLOCK_SIZE_K")
     if block_shape is not None:
         BLOCK_SIZE_K = min(BLOCK_SIZE_K, min(block_shape[0], block_shape[1]))
+    if use_td and A.size(1) % BLOCK_SIZE_K != 0:
+        # TD gather/load feeding tl.dot with a non-block-aligned K
+        # miscompiles (~74% of output elements wrong) on real HW;
+        # this is a compiler-codegen issue, not a Python-maskable
+        # boundary gap. Fall back to the pointer-arith path.
+        logger.warning_once(
+            "Disabling VLLM_TRITON_USE_TD for this MoE launch: K=%d is not "
+            "a multiple of BLOCK_SIZE_K=%d, which triggers a known "
+            "Triton tensor-descriptor + tl.dot miscompilation.",
+            A.size(1),
+            BLOCK_SIZE_K,
+        )
+        use_td = False
     fused_moe_kernel[grid](
         A,
         B,
@@ -847,6 +899,7 @@ def invoke_fused_moe_triton_kernel(
         HAS_BIAS=HAS_BIAS,
         BLOCK_SIZE_K=BLOCK_SIZE_K,
         SWAP_AB=SWAP_AB,
+        USE_TD=use_td,
         **config,
     )
 
diff --git a/vllm/model_executor/layers/fused_moe/modular_kernel.py b/vllm/model_executor/layers/fused_moe/modular_kernel.py
index cc5e2d0f06a..f7dc5dc2c05 100644
--- a/vllm/model_executor/layers/fused_moe/modular_kernel.py
+++ b/vllm/model_executor/layers/fused_moe/modular_kernel.py
@@ -10,7 +10,6 @@ from typing import final
 import torch
 
 import vllm.envs as envs
-from vllm.forward_context import get_forward_context, is_forward_context_available
 from vllm.logger import init_logger
 from vllm.model_executor.layers.fused_moe.activation import (
     MoEActivation,
@@ -891,6 +890,8 @@ class FusedMoEExpertsModular(FusedMoEExperts):
         beta: float = 0.0,
         topk_ids: torch.Tensor | None = None,
         expert_map: torch.Tensor | None = None,
+        activation_situ_beta: float | None = None,
+        activation_situ_linear_beta: float | None = None,
     ) -> None:
         apply_moe_activation(
             activation,
@@ -901,8 +902,16 @@ class FusedMoEExpertsModular(FusedMoEExperts):
             beta=beta,
             topk_ids=topk_ids,
             expert_map=expert_map,
-            activation_situ_beta=self.moe_config.activation_situ_beta,
-            activation_situ_linear_beta=(self.moe_config.activation_situ_linear_beta),
+            activation_situ_beta=(
+                self.moe_config.activation_situ_beta
+                if activation_situ_beta is None
+                else activation_situ_beta
+            ),
+            activation_situ_linear_beta=(
+                self.moe_config.activation_situ_linear_beta
+                if activation_situ_linear_beta is None
+                else activation_situ_linear_beta
+            ),
         )
 
     @abstractmethod
@@ -1196,20 +1205,6 @@ class FusedMoEKernelModularImpl:
         The _prepare method is a wrapper around self.prepare_finalize.prepare
         that handles DBO and async.
         """
-        # Skip cudagraph/DP padding tokens uniformly across all a2a backends:
-        # forcing padded rows' expert ids to -1 makes every prepare_finalize drop
-        # them (not dispatched / not computed by the experts). The V2 model runner
-        # marks them in forward_context.is_padding; it is None for runners that do
-        # not populate it, leaving topk_ids unchanged.
-        # This requires the experts kernel to treat topk_id == -1 as a skip
-        # sentinel.
-        is_padding = None
-        if envs.VLLM_MOE_SKIP_PADDING and is_forward_context_available():
-            is_padding = get_forward_context().is_padding
-        if is_padding is not None:
-            n = topk_ids.shape[0]
-            # TODO: Properly support DBO (padding lives at the batch tail).
-            topk_ids = torch.where(is_padding[:n].unsqueeze(1), -1, topk_ids)
 
         if not self.prepare_finalize.supports_async():
             # We shouldn't be running an a2a kernel that doesn't
diff --git a/vllm/model_executor/layers/fused_moe/oracle/mxfp4.py b/vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
index 5cb0436ae50..0aed6cbf403 100644
--- a/vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
+++ b/vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
@@ -502,6 +502,7 @@ def select_mxfp4_moe_backend(
         _get_priority_backends_for_gpt_oss(), requested_activation_key
     )
 
+    unsupported_reasons = []
     for backend in AVAILABLE_BACKENDS:
         # Use requested_activation_key if provided, otherwise use backend default
         act_key = (
@@ -518,6 +519,7 @@ def select_mxfp4_moe_backend(
                 return backend, k_cls
             else:
                 logger.debug_once(_make_log_unsupported(backend, reason))
+                unsupported_reasons.append((backend, reason))
 
     if current_platform.is_xpu():
         backend = Mxfp4MoeBackend.XPU
@@ -541,26 +543,19 @@ def select_mxfp4_moe_backend(
             activation_format,
         )
 
-    if current_platform.is_rocm():
-        backend = Mxfp4MoeBackend.TRITON_UNFUSED
-        logger.info_once(_make_log_backend(backend))
-        return _return_or_raise(
-            Mxfp4MoeBackend.TRITON_UNFUSED,
-            config,
-            kMxfp4Static,
-            None,
-            activation_format,
-        )
-
-    if current_platform.is_cuda():
-        raise NotImplementedError(
-            "No MXFP4 MoE backend supports the deployment configuration. "
-            f"weight_key=kMxfp4Static, activation_key={activation_key}. "
-            "Native backends require specific hardware. "
-            "Set `VLLM_LOGGING_LEVEL=DEBUG` to see detailed unsupported reasons. "
-        )
-
-    return Mxfp4MoeBackend.NONE, None
+    unsupported_log = "; ".join(
+        [
+            f"backend: {backend.value}, reason: {reason}"
+            for backend, reason in unsupported_reasons
+        ]
+    )
+    raise NotImplementedError(
+        "No MXFP4 MoE backend supports the deployment configuration. "
+        f"weight_key=kMxfp4Static, activation_key={activation_key}. "
+        f"Candidate backends were: "
+        f"{[backend.value for backend in AVAILABLE_BACKENDS]}. "
+        f"Unsupported reasons: {unsupported_log}. "
+    )
 
 
 def select_deepseek_v4_mxfp4_moe_backend(
diff --git a/vllm/model_executor/layers/fused_moe/oracle/mxfp8.py b/vllm/model_executor/layers/fused_moe/oracle/mxfp8.py
index b9086cfa48a..2a03eacea8e 100644
--- a/vllm/model_executor/layers/fused_moe/oracle/mxfp8.py
+++ b/vllm/model_executor/layers/fused_moe/oracle/mxfp8.py
@@ -12,10 +12,10 @@ from vllm.model_executor.layers.quantization.utils.quant_utils import (
     kMxfp8Dynamic,
     kMxfp8Static,
 )
-from vllm.platforms import current_platform
 
 logger = init_logger(__name__)
 
+# Ordered by priority.
 _SUPPORTED_BACKENDS = (
     Fp8MoeBackend.FLASHINFER_TRTLLM,
     Fp8MoeBackend.DEEPGEMM,
@@ -26,6 +26,8 @@ _SUPPORTED_BACKENDS = (
     # devices / no flydsl / EP it is skipped and native is used.
     Fp8MoeBackend.AITER_MXFP8,
     Fp8MoeBackend.HUMMING,
+    Fp8MoeBackend.TRITON_MXFP8,
+    Fp8MoeBackend.EMULATION,
 )
 
 _BACKEND_NAME_MAP: dict[str, Fp8MoeBackend] = {
@@ -61,15 +63,12 @@ def _mxfp8_backend_to_kernel_cls(
 
         return [AiterMxfp8Experts]
     if backend == Fp8MoeBackend.TRITON_MXFP8:
-        # Explicit ``--moe-backend triton``: the Triton mxfp8 path, i.e.
-        # dot_scaled on MX-capable HW (gfx950) and BF16 emulation otherwise.
-        # Mirrors the ROCm auto-fallback in ``_select_rocm_mxfp8_backend``.
-        if current_platform.supports_mx():
-            from vllm.model_executor.layers.fused_moe.experts.mxfp8_native_moe import (
-                Mxfp8NativeTritonExperts,
-            )
+        from vllm.model_executor.layers.fused_moe.experts.mxfp8_native_moe import (
+            Mxfp8NativeTritonExperts,
+        )
 
-            return [Mxfp8NativeTritonExperts]
+        return [Mxfp8NativeTritonExperts]
+    if backend == Fp8MoeBackend.EMULATION:
         from vllm.model_executor.layers.fused_moe.experts.mxfp8_emulation_moe import (
             Mxfp8EmulationTritonExperts,
         )
@@ -105,35 +104,6 @@ def _select_kernel_cls(
     )
 
 
-def _select_rocm_mxfp8_backend() -> tuple[Fp8MoeBackend, type[mk.FusedMoEExperts]]:
-    """ROCm fallback when no auto-selected MXFP8 backend is available.
-
-    The aiter FlyDSL backend (``AITER_MXFP8``) is auto-picked earlier by
-    ``select_mxfp8_moe_backend`` via ``_SUPPORTED_BACKENDS`` when usable, or
-    explicitly via ``--moe-backend aiter``; this fallback handles the rest
-    (native dot_scaled on gfx950, else BF16 emulation).
-    """
-
-    if current_platform.supports_mx():
-        from vllm.model_executor.layers.fused_moe.experts.mxfp8_native_moe import (
-            Mxfp8NativeTritonExperts,
-        )
-
-        logger.info_once("Using native CDNA4 (gfx950) MXFP8 dot_scaled MoE backend.")
-        return Fp8MoeBackend.TRITON_MXFP8, Mxfp8NativeTritonExperts
-
-    from vllm.model_executor.layers.fused_moe.experts.mxfp8_emulation_moe import (
-        Mxfp8EmulationTritonExperts,
-    )
-
-    logger.info_once(
-        "No native MXFP8 MoE backend available on this device; "
-        "MXFP8 weights will be dequantized to BF16 once at load time and the "
-        "MoE will run in BF16 (no per-step dequant)."
-    )
-    return Fp8MoeBackend.EMULATION, Mxfp8EmulationTritonExperts
-
-
 def select_mxfp8_moe_backend(
     config: FusedMoEConfig,
 ) -> tuple[Fp8MoeBackend, type[mk.FusedMoEExperts]]:
@@ -167,8 +137,5 @@ def select_mxfp8_moe_backend(
         logger.info_once("Using '%s' MxFp8 MoE backend.", backend.value)
         return backend, experts_cls
 
-    # simplify the logic for rocm, refactor later when more backends are supported
-    if current_platform.is_rocm():
-        return _select_rocm_mxfp8_backend()
-
+    # TODO: add debug log with reason.
     raise ValueError("No MXFP8 MoE backends available.")
diff --git a/vllm/model_executor/layers/fused_moe/oracle/unquantized.py b/vllm/model_executor/layers/fused_moe/oracle/unquantized.py
index 8fa1a0c265c..bf77a316bb4 100644
--- a/vllm/model_executor/layers/fused_moe/oracle/unquantized.py
+++ b/vllm/model_executor/layers/fused_moe/oracle/unquantized.py
@@ -362,6 +362,13 @@ def make_unquantized_moe_kernel(
     experts_cls: type[mk.FusedMoEExperts],
     routing_tables: tuple[torch.Tensor, torch.Tensor, torch.Tensor] | None = None,
 ) -> mk.FusedMoEKernel:
+    from vllm.model_executor.layers.fused_moe.utils import (
+        warn_if_moe_use_td_ineffective,
+    )
+
+    # Warn against the selected backend, not each probed candidate.
+    warn_if_moe_use_td_ineffective(backend.value, is_quantized=False)
+
     # Create Prepare/Finalize
     is_monolithic = issubclass(experts_cls, mk.FusedMoEExpertsMonolithic)
     prepare_finalize = maybe_make_prepare_finalize(
diff --git a/vllm/model_executor/layers/fused_moe/router/router_factory.py b/vllm/model_executor/layers/fused_moe/router/router_factory.py
index c7cfccbe64b..ff4874b20e6 100644
--- a/vllm/model_executor/layers/fused_moe/router/router_factory.py
+++ b/vllm/model_executor/layers/fused_moe/router/router_factory.py
@@ -9,6 +9,7 @@ from vllm._aiter_ops import rocm_aiter_ops
 from vllm.distributed.eplb.eplb_state import EplbLayerState
 from vllm.model_executor.layers.fused_moe.config import (
     RoutingMethodType,
+    get_routing_method_type,
 )
 from vllm.model_executor.layers.fused_moe.router.aiter_shared_routed_fused_moe_router import (  # noqa: E501
     AiterSharedRoutedFusedMoERouter,
@@ -67,7 +68,8 @@ def create_fused_moe_router(
     The selection logic follows this priority order:
     1. RoutingSimulatorRouter - if VLLM_MOE_ROUTING_SIMULATION_STRATEGY env var is set
     2. ZeroExpertRouter - if zero_expert_type is not None
-    3. GroupedTopKRouter - if use_grouped_topk is True
+    3. GroupedTopKRouter - if use_grouped_topk is True and the grouping is not
+       degenerate (at most one group, with topk_group <= 1)
     4. CustomRoutingRouter - if custom_routing_function is not None
     5. FusedTopKBiasRouter - if e_score_correction_bias is not None
     6. AiterSharedRoutedFusedMoERouter - if num_fused_shared_experts > 0
@@ -143,30 +145,48 @@ def create_fused_moe_router(
                 "num_expert_group and topk_group must be provided when "
                 "use_grouped_topk is True"
             )
-        grouped_topk_router = GroupedTopKRouter(
-            top_k=top_k,
-            global_num_experts=global_num_experts,
-            eplb_state=eplb_state,
-            num_expert_group=num_expert_group,
-            topk_group=topk_group,
-            renormalize=renormalize,
-            scoring_func=scoring_func,
-            routed_scaling_factor=routed_scaling_factor,
-            e_score_correction_bias=e_score_correction_bias,
-            num_fused_shared_experts=num_fused_shared_experts,
-        )
-        if (
-            grouped_topk_router.routing_method_type != RoutingMethodType.Unspecified
-            or num_expert_group > 1
-            or topk_group > 1
-        ):
-            return grouped_topk_router
 
-        # If routing_method for GroupedTopKRouter is Unspecified and there is only
-        # one group, fallback to standard top-k routing
-        use_grouped_topk = False
-        num_expert_group = None
-        topk_group = None
+        # For topk_group <= 1, grouped implementation is pure overhead.
+        degenerate_grouping = num_expert_group <= 1 and topk_group <= 1
+        # FusedTopKRouter cannot apply routed_scaling_factor, FusedTopKBiasRouter can.
+        scaling_handled_downstream = (
+            routed_scaling_factor == 1.0 or e_score_correction_bias is not None
+        )
+
+        # Degenerating must not change the advertised routing method, which drives
+        # kernel selection. num_expert_group only affects it for biased routing.
+        def advertised_routing_method(groups: int | None) -> RoutingMethodType:
+            return get_routing_method_type(
+                scoring_func=scoring_func,
+                top_k=top_k,
+                renormalize=renormalize,
+                num_expert_group=groups,
+                has_e_score_bias=e_score_correction_bias is not None,
+                routed_scaling_factor=routed_scaling_factor,
+            )
+
+        routing_method_preserved = advertised_routing_method(
+            num_expert_group
+        ) == advertised_routing_method(None)
+
+        if not (
+            degenerate_grouping
+            and scaling_handled_downstream
+            and routing_method_preserved
+        ):
+            return GroupedTopKRouter(
+                top_k=top_k,
+                global_num_experts=global_num_experts,
+                eplb_state=eplb_state,
+                num_expert_group=num_expert_group,
+                topk_group=topk_group,
+                renormalize=renormalize,
+                scoring_func=scoring_func,
+                routed_scaling_factor=routed_scaling_factor,
+                e_score_correction_bias=e_score_correction_bias,
+                num_fused_shared_experts=num_fused_shared_experts,
+            )
+        # Otherwise fall through to the non-grouped chain below.
 
     if custom_routing_function is not None:
         return CustomRoutingRouter(
diff --git a/vllm/model_executor/layers/fused_moe/utils.py b/vllm/model_executor/layers/fused_moe/utils.py
index e3d6493dda2..cce8ccd073f 100644
--- a/vllm/model_executor/layers/fused_moe/utils.py
+++ b/vllm/model_executor/layers/fused_moe/utils.py
@@ -7,7 +7,9 @@ from typing import TYPE_CHECKING
 import torch
 import torch.nn.functional as F
 
+import vllm.envs as envs
 from vllm import _custom_ops as ops
+from vllm.logger import init_logger
 from vllm.model_executor.layers.quantization.utils.fp8_utils import (
     per_token_group_quant_fp8,
 )
@@ -39,6 +41,8 @@ from vllm.utils.math_utils import cdiv
 if TYPE_CHECKING:
     from vllm.model_executor.layers.fused_moe.config import FusedMoEConfig
 
+logger = init_logger(__name__)
+
 
 @triton.jit
 def _count_expert_num_tokens(
@@ -585,3 +589,80 @@ def enable_swap_ab(BLOCK_SIZE_M: int, BLOCK_SIZE_N: int) -> bool:
         and BLOCK_SIZE_M < 64
         and BLOCK_SIZE_N >= 64
     )
+
+
+def moe_use_td_hw_supported() -> bool:
+    """Whether the current device can run the TD (gather) path of
+    ``fused_moe_kernel`` (ignores the ``VLLM_TRITON_USE_TD`` override).
+
+    The A-load uses ``tensor_descriptor.gather``, which lowers to the PTX
+    ``tile::gather4`` instruction. That instruction is part of the
+    ``tcgen05``/Tensor Memory (TMEM) family introduced with Blackwell and has
+    no Hopper (sm90) equivalent -- ptxas rejects it there ("Feature
+    '.tile::gather4 ...' requires .target sm_100 or higher"). Unlike
+    ``scatter4``, ``gather4`` is supported across the whole sm100+ range
+    including consumer Blackwell (sm120/sm121): see triton-lang/triton#8498,
+    which enables ``gather4`` on sm120/sm121 while leaving ``scatter4``
+    unsupported there. So this gates on a blanket ``has_device_capability(100)``
+    rather than the sm100 *family* check used for the scatter store path.
+    """
+    if current_platform.is_xpu():
+        return True
+    if current_platform.is_cuda():
+        return current_platform.has_device_capability(100)
+    return False
+
+
+def resolve_moe_use_td() -> bool:
+    """Tri-state resolver for ``VLLM_TRITON_USE_TD``.
+
+    Unset auto-selects the TD path on XPU only, mirroring the attention
+    dispatcher in ``triton_attn.py``. ``1``/``0`` force it on/off regardless
+    of hardware; forcing ``1`` where it cannot compile (see
+    ``moe_use_td_hw_supported``) fails at ptxas. Blackwell CUDA (sm100+) can
+    compile it but is opt-in only, pending validation.
+    """
+    override = envs.VLLM_TRITON_USE_TD
+    if override is None:
+        return current_platform.is_xpu()
+    return override
+
+
+_warned_moe_use_td_ineffective = False
+
+
+def warn_if_moe_use_td_ineffective(
+    active_backend: str, is_quantized: bool = False
+) -> None:
+    """One-shot warning when ``VLLM_TRITON_USE_TD`` is set but ignored.
+
+    Fires when the user set the env explicitly and either (a) the active
+    MoE backend is not the fused Triton kernel, or (b) the model is
+    quantized (the TD path falls back to the pointer path under any
+    quantization).
+    """
+    global _warned_moe_use_td_ineffective
+    if _warned_moe_use_td_ineffective:
+        return
+    if envs.VLLM_TRITON_USE_TD is None:
+        return
+    is_triton = active_backend.upper() == "TRITON"
+    if is_triton and not is_quantized:
+        return
+    if not is_triton:
+        reason = (
+            f"the active MoE backend is {active_backend!r}; pass "
+            "`--moe-backend triton` to enable the tensor-descriptor path"
+        )
+    else:
+        reason = (
+            "the model uses quantized MoE weights; the TD path is "
+            "currently restricted to non-quantized weights and falls "
+            "back to the pointer path"
+        )
+    logger.warning(
+        "VLLM_TRITON_USE_TD is set to %s but %s.",
+        envs.VLLM_TRITON_USE_TD,
+        reason,
+    )
+    _warned_moe_use_td_ineffective = True
diff --git a/vllm/model_executor/layers/linear.py b/vllm/model_executor/layers/linear.py
index df105f06634..e4662148ff7 100644
--- a/vllm/model_executor/layers/linear.py
+++ b/vllm/model_executor/layers/linear.py
@@ -50,6 +50,7 @@ WEIGHT_LOADER_V2_SUPPORTED = [
     "UnquantizedLinearMethod",
     "CompressedTensorsLinearMethod",
     "CompressedTensorsLinearTransformMethod",
+    "QutlassNvFP4LinearMethod",
     "AutoAWQMarlinLinearMethod",
     "AutoAWQLinearMethod",
     "AutoGPTQLinearMethod",
diff --git a/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py b/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py
index b41ae77a95f..bc49226c75b 100644
--- a/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py
+++ b/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py
@@ -21,9 +21,7 @@ from vllm.model_executor.model_loader.weight_utils import (
 from vllm.model_executor.parameter import BasevLLMParameter
 from vllm.model_executor.utils import set_weight_attrs
 from vllm.platforms import current_platform
-from vllm.third_party.flash_linear_attention.ops.fused_norm_gate import (
-    FusedRMSNormGated,
-)
+from vllm.third_party.flash_linear_attention.ops.kda import FusedRMSNormGated
 from vllm.transformers_utils.configs.kimi_linear import KimiLinearConfig
 from vllm.v1.attention.backends.gdn_attn import GDNAttentionMetadata
 
diff --git a/vllm/model_executor/layers/mla.py b/vllm/model_executor/layers/mla.py
index a846aa5f0e9..6c0e8e069fc 100644
--- a/vllm/model_executor/layers/mla.py
+++ b/vllm/model_executor/layers/mla.py
@@ -8,6 +8,7 @@ from vllm.config import CacheConfig
 from vllm.model_executor.custom_op import PluggableLayer
 from vllm.model_executor.layers.attention import MLAAttention
 from vllm.model_executor.layers.quantization import QuantizationConfig
+from vllm.platforms import current_platform
 
 
 @dataclass
@@ -67,6 +68,7 @@ class MultiHeadLatentAttentionWrapper(PluggableLayer):
         prefix: str = "",
         skip_topk: bool = False,
         non_causal_multi_token_decode: bool = False,
+        allow_short_prefill_indexer_scoring_skip: bool = False,
     ) -> None:
         super().__init__()
         self.hidden_size = hidden_size
@@ -123,7 +125,26 @@ class MultiHeadLatentAttentionWrapper(PluggableLayer):
             topk_indices_buffer=mla_modules.topk_indices_buffer,
             non_causal_multi_token_decode=non_causal_multi_token_decode,
         )
-
+        indexer_op = getattr(self.indexer, "indexer_op", None)
+        if indexer_op is not None and hasattr(
+            indexer_op, "dense_mha_metadata_layer_name"
+        ):
+            enable_short_prefill_scoring_skip = (
+                allow_short_prefill_indexer_scoring_skip
+                and not self.skip_topk
+                and not getattr(indexer_op, "use_pcp", False)
+                and current_platform.is_cuda()
+            )
+            # The indexer and main MLA use independent decode thresholds and
+            # may classify the same short extend differently. Bind the main
+            # MLA layer name so the eager indexer op can check whether the
+            # batch's top-k indices will be consumed.
+            # PCP is excluded because indexer cache/scoring ownership differs
+            # across ranks and the no-consumer invariant has not been
+            # established there.
+            indexer_op.dense_mha_metadata_layer_name = (
+                self.mla_attn.layer_name if enable_short_prefill_scoring_skip else ""
+            )
         self.prefix = prefix
 
     def forward(
diff --git a/vllm/model_executor/layers/quantization/compressed_tensors/transform/linear.py b/vllm/model_executor/layers/quantization/compressed_tensors/transform/linear.py
index bd1964e667d..0cde020627d 100644
--- a/vllm/model_executor/layers/quantization/compressed_tensors/transform/linear.py
+++ b/vllm/model_executor/layers/quantization/compressed_tensors/transform/linear.py
@@ -25,6 +25,7 @@ from vllm.model_executor.layers.quantization.compressed_tensors.transform.module
 from vllm.model_executor.layers.quantization.compressed_tensors.transform.utils import (  # noqa: E501
     TransformTuple,
 )
+from vllm.platforms import current_platform
 
 
 class CompressedTensorsLinearTransformMethod(LinearMethodBase):
@@ -48,11 +49,12 @@ class CompressedTensorsLinearTransformMethod(LinearMethodBase):
 
         assert input_tfms or output_tfms
 
-        if is_qutlass_fp4_scheme(quant_scheme, input_tfms):
+        if is_qutlass_fp4_scheme(
+            quant_scheme, input_tfms
+        ) and current_platform.has_device_capability(100):
             return QutlassNvFP4LinearMethod(quant_method, input_tfms, output_tfms)
 
         # hadacore or dense gemm is selected by Transform module
-
         return cls(quant_method, input_tfms, output_tfms)
 
     def __init__(
diff --git a/vllm/model_executor/layers/quantization/compressed_tensors/transform/module.py b/vllm/model_executor/layers/quantization/compressed_tensors/transform/module.py
index f5589c8c07f..75505934938 100644
--- a/vllm/model_executor/layers/quantization/compressed_tensors/transform/module.py
+++ b/vllm/model_executor/layers/quantization/compressed_tensors/transform/module.py
@@ -32,7 +32,7 @@ class HadamardTransform(torch.nn.Module):
     transforms: dict[int, TransformTuple]  # info parsed from transforms config
     weight: SharedWeightParameter  # container for shared tensors
 
-    scales: dict[int, float]  # hadamard scale, usually sqrt(matrix.size(0))
+    scaled_data_ptrs: set[int] = set()
 
     def __init__(
         self,
@@ -44,7 +44,6 @@ class HadamardTransform(torch.nn.Module):
     ):
         super().__init__()
         self.transforms = transforms
-        self.scales = {}
 
         if get_tensor_model_parallel_world_size() > 1:
             raise NotImplementedError(
@@ -64,26 +63,26 @@ class HadamardTransform(torch.nn.Module):
             )
 
             data_key = self._get_data_key(scheme, weight_size)
+            # load up in model's default precision, rather than using scheme.precision
             self.weight.add_partition(
                 part_index,
                 data_key,
                 size=(weight_size, weight_size),
-                dtype=scheme.precision,
             )
 
         # validate that shared tensors and schemes are correct
         self._validate_input_transforms()
 
     def process_weights_after_loading(self):
-        for part_id in self.weight.partitions:
-            data = self.weight.partitions[part_id].data
-
+        for part_id, partition in self.weight.partitions.items():
             # required by torch.compile
             self.weight.process_weights_after_loading()
 
-            # precompute scale as a runtime multiply, not division
-            # do not fold into weight in order to utilize FWHT
-            self.scales[part_id] = 1 / math.sqrt(data.size(0))
+            # Merge normalization scale directly into weight, must be done only once
+            data_ptr = partition.data.data_ptr()
+            if data_ptr not in HadamardTransform.scaled_data_ptrs:
+                partition.data.div_(math.sqrt(partition.data.size(0)))
+                HadamardTransform.scaled_data_ptrs.add(data_ptr)
 
             # FUTURE: avoid runtime transpose by processing weights
             # prior to apply
@@ -111,26 +110,19 @@ class HadamardTransform(torch.nn.Module):
             weight = (
                 weight if self.transforms[part_id].args.inverse else weight.T
             )  # linear := x(W.T)
-            scale = self.scales[part_id]
 
             if self.transforms[part_id].scheme.head_dim is not None:
                 value = value.unflatten(-1, (-1, weight.size(0)))
-                value = (
-                    dispatch_unquantized_gemm()(
-                        self, value.to(weight.dtype), weight, None
-                    ).to(value.dtype)
-                    * scale
-                )
+                value = dispatch_unquantized_gemm()(
+                    self, value.to(weight.dtype), weight, None
+                ).to(value.dtype)
                 value = value.flatten(-2, -1)
 
                 return value
 
-            return (
-                dispatch_unquantized_gemm()(
-                    self, value.to(weight.dtype), weight, None
-                ).to(value.dtype)
-                * scale
-            )
+            return dispatch_unquantized_gemm()(
+                self, value.to(weight.dtype), weight, None
+            ).to(value.dtype)
 
     def _get_data_key(self, scheme: TransformScheme, weight_size: int) -> Hashable:
         return (id(scheme), weight_size)
diff --git a/vllm/model_executor/layers/quantization/compressed_tensors/transform/schemes/linear_qutlass_nvfp4.py b/vllm/model_executor/layers/quantization/compressed_tensors/transform/schemes/linear_qutlass_nvfp4.py
index f0bb47a728a..b8b771e3c7a 100644
--- a/vllm/model_executor/layers/quantization/compressed_tensors/transform/schemes/linear_qutlass_nvfp4.py
+++ b/vllm/model_executor/layers/quantization/compressed_tensors/transform/schemes/linear_qutlass_nvfp4.py
@@ -2,7 +2,13 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import torch
+from torch.nn.parameter import Parameter
 
+from vllm._custom_ops import fusedQuantizeNv
+from vllm.model_executor.kernels.linear import (
+    _LINEAR_BACKEND_KERNEL_MAP,
+    NvFp4LinearKernel,
+)
 from vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors import (  # noqa: E501
     CompressedTensorsScheme,
     CompressedTensorsW4A4Fp4,
@@ -11,18 +17,32 @@ from vllm.model_executor.layers.quantization.compressed_tensors.transform.linear
     CompressedTensorsLinearTransformMethod,
     TransformTuple,
 )
+from vllm.model_executor.layers.quantization.qutlass_utils import to_blocked
+from vllm.model_executor.layers.quantization.utils.nvfp4_utils import (
+    slice_nvfp4_output,
+)
+from vllm.utils.flashinfer import (
+    flashinfer_scaled_fp4_mm,
+)
 
 __all__ = ["is_qutlass_fp4_scheme", "QutlassNvFP4LinearMethod"]
 
+NVFP4_MAX = 6.0
 
+
+# QUTLASS supports transform block sizes (16, 32, 64, 128) for NVFP4
+# https://github.com/IST-DASLab/qutlass/blob/v0.2.0/qutlass/csrc/bindings.cpp#L413-L414
 def is_qutlass_fp4_scheme(
     quant_scheme: CompressedTensorsScheme | None,
     input_tfms: dict[int, TransformTuple],
 ) -> bool:
     return (
-        isinstance(quant_scheme, (CompressedTensorsW4A4Fp4,))
-        and len(input_tfms) == 1
-        and input_tfms[0].scheme.head_dim == quant_scheme.group_size
+        isinstance(quant_scheme, CompressedTensorsW4A4Fp4)
+        and len(input_tfms) >= 1
+        and all(
+            input_tfm.scheme.head_dim in (16, 32, 64, 128)
+            for input_tfm in input_tfms.values()
+        )
     )
 
 
@@ -50,15 +70,90 @@ class QutlassNvFP4LinearMethod(CompressedTensorsLinearTransformMethod):
         )
 
         assert self.input_transform is not None
-        assert len(self.input_transform.weight) == 1
-        assert self.input_transform.weight[0].size(0) == layer.scheme.group_size
+        assert len(self.input_transform.weight.partitions) >= 1
 
         return ret
 
+    @staticmethod
+    def _get_flashinfer_gemm_backend(kernel: NvFp4LinearKernel) -> str:
+        """
+        Given a kernel, find the string that is needed to be passed into
+        `flashinfer_scaled_fp4_mm`, using
+        vllm.model_executor.kernels.linear._LINEAR_BACKEND_KERNEL_MAP as source of truth
+        """
+        kernel_type = type(kernel)
+        for key, kernels in _LINEAR_BACKEND_KERNEL_MAP.items():
+            if not key.startswith("flashinfer_") or kernel_type not in kernels:
+                continue
+            backend = key.removeprefix("flashinfer_")
+            # flashinfer GEMM backend uses "cute-dsl", not "cutedsl"
+            return backend.replace("cutedsl", "cute-dsl")
+        raise ValueError(
+            f"QutlassNvFP4 transform requires a FlashInfer kernel, "
+            f"got {kernel_type.__name__}"
+        )
+
+    def process_weights_after_loading(self, layer):
+        super().process_weights_after_loading(layer)
+
+        assert self.input_transform is not None
+        layer.hadamard_matrix = self.input_transform.weight.partitions[0].data
+
+        # fusedQuantizeNv stores raw absmax as block scales (sf = absmax),
+        # while CT weights use sf = absmax * SFScaleVal / 6.0. The GEMM
+        # computes alpha * sum(fp4_a * sf_a * fp4_w * sf_w), so alpha must
+        # compensate: alpha = weight_global_scale / 6.0
+        layer.fused_alpha = Parameter(
+            layer.weight_global_scale / NVFP4_MAX, requires_grad=False
+        )
+
+        layer.fused_global_scale = Parameter(
+            torch.tensor(
+                [NVFP4_MAX],
+                dtype=torch.float32,
+                device=layer.weight_global_scale.device,
+            ),
+            requires_grad=False,
+        )
+
+        layer.flashinfer_gemm_backend = self._get_flashinfer_gemm_backend(
+            layer.scheme.kernel
+        )
+
     def apply(
         self,
         layer: torch.nn.Module,
         x: torch.Tensor,
         bias: torch.Tensor | None = None,
     ) -> torch.Tensor:
-        raise NotImplementedError()
+        assert bias is None
+        output_size = layer.output_size_per_partition
+        output_shape = [*x.shape[:-1], output_size]
+
+        x_flat = x.contiguous().flatten(end_dim=-2)
+
+        x_fp4, x_scales = fusedQuantizeNv(
+            x_flat, layer.hadamard_matrix, layer.fused_global_scale
+        )
+
+        x_scales_blocked = to_blocked(x_scales, backend="triton").view(x_scales.shape)
+
+        out = flashinfer_scaled_fp4_mm(
+            x_fp4,
+            layer.weight,
+            x_scales_blocked,
+            layer.weight_scale,
+            layer.fused_alpha,
+            x.dtype,
+            backend=layer.flashinfer_gemm_backend,
+        )
+
+        out = slice_nvfp4_output(out, output_size)
+
+        if self.output_transform is not None:
+            for part_id, (start, length) in enumerate(self.partition_ranges):
+                out[:, start : start + length] = self.output_transform(
+                    out[:, start : start + length].clone(), part_id=part_id
+                )
+
+        return out.view(*output_shape)
diff --git a/vllm/model_executor/layers/quantization/compressed_tensors/utils.py b/vllm/model_executor/layers/quantization/compressed_tensors/utils.py
index afb899cd6d7..872771af9ab 100644
--- a/vllm/model_executor/layers/quantization/compressed_tensors/utils.py
+++ b/vllm/model_executor/layers/quantization/compressed_tensors/utils.py
@@ -145,8 +145,8 @@ def find_matched_target(
 
     matched_target = (
         _find_first_match(layer_name, targets)
-        or _find_first_match(module.__class__.__name__, targets, True)
         or _match_fused_layer(layer_name, targets, fused_mapping)
+        or _find_first_match(module.__class__.__name__, targets, True)
     )
 
     return matched_target
diff --git a/vllm/model_executor/layers/quantization/inc/config_parser.py b/vllm/model_executor/layers/quantization/inc/config_parser.py
index 603b80b7cd0..6e94cad2cc6 100644
--- a/vllm/model_executor/layers/quantization/inc/config_parser.py
+++ b/vllm/model_executor/layers/quantization/inc/config_parser.py
@@ -142,6 +142,14 @@ class INCConfigParser:
         if self._config.extra_config and layer_name in self._config.extra_config:
             return get_config(layer_name)
 
+        # Suffix match: handle cases where extra_config keys use short names
+        # (e.g. "lm_head") but the layer_name is fully qualified
+        # (e.g. "model.language_model.lm_head") due to model nesting.
+        if self._config.extra_config:
+            for cfg_key in self._config.extra_config:
+                if layer_name.endswith(f".{cfg_key}"):
+                    return get_config(cfg_key)
+
         quantized = not isinstance(layer, ParallelLMHead)
         if self._config.block_name_to_quantize:
             quantized = any(
diff --git a/vllm/model_executor/layers/quantization/input_quant_fp8.py b/vllm/model_executor/layers/quantization/input_quant_fp8.py
index e8810919c20..2eb34630aa6 100644
--- a/vllm/model_executor/layers/quantization/input_quant_fp8.py
+++ b/vllm/model_executor/layers/quantization/input_quant_fp8.py
@@ -139,11 +139,6 @@ class QuantFP8(CustomOp):
         scale_ub: torch.Tensor | None = None,
         use_triton: bool = False,
     ) -> tuple[torch.Tensor, torch.Tensor]:
-        if self.is_group_quant and use_triton:
-            assert scale is None, "Dynamic group quantization does not use scale"
-
-            return torch.ops.vllm.triton_per_token_group_quant_fp8(x, self.group_size)
-
         use_aiter_quant = self.use_aiter and scale_ub is None and x.is_contiguous()
         use_aiter_per_tensor_quant = (
             use_aiter_quant and self.group_shape.is_per_tensor()
diff --git a/vllm/model_executor/layers/quantization/qutlass_utils.py b/vllm/model_executor/layers/quantization/qutlass_utils.py
index 315ecd0c009..86b0548307a 100644
--- a/vllm/model_executor/layers/quantization/qutlass_utils.py
+++ b/vllm/model_executor/layers/quantization/qutlass_utils.py
@@ -84,6 +84,7 @@ def triton_scale_swizzle(
     )
 
 
+@torch.library.custom_op("vllm::triton_mx_block_rearrange", mutates_args=())
 def triton_mx_block_rearrange(scale_tensor: torch.Tensor) -> torch.Tensor:
     """
     Rearranges an E8M0 tensor scale from row-major format to
@@ -142,6 +143,14 @@ def triton_mx_block_rearrange(scale_tensor: torch.Tensor) -> torch.Tensor:
     return out
 
 
+@triton_mx_block_rearrange.register_fake
+def _triton_mx_block_rearrange_fake(scale_tensor: torch.Tensor) -> torch.Tensor:
+    rows, cols = scale_tensor.shape
+    padded_rows = cdiv(rows, 128) * 128
+    padded_cols = cdiv(cols, 4) * 4
+    return scale_tensor.new_empty((padded_rows, padded_cols))
+
+
 def to_blocked(
     input_matrix: torch.Tensor, backend: Literal["torch", "triton"] = "triton"
 ) -> torch.Tensor:
@@ -157,7 +166,7 @@ def to_blocked(
         backend: "torch" (PyTorch path) or "triton" (Triton kernel)
 
     Returns:
-        Rearranged tensor of shape (32*cdiv(H,128), 16*cdiv(W,4))
+        Rearranged flattened tensor of size (32*cdiv(H,128) * 16*cdiv(W,4))
     """
     if backend == "triton":
         return triton_mx_block_rearrange(input_matrix).flatten()
diff --git a/vllm/model_executor/layers/quantization/utils/flashinfer_fp4_moe.py b/vllm/model_executor/layers/quantization/utils/flashinfer_fp4_moe.py
index bab3dee649b..1b513e9c0ce 100644
--- a/vllm/model_executor/layers/quantization/utils/flashinfer_fp4_moe.py
+++ b/vllm/model_executor/layers/quantization/utils/flashinfer_fp4_moe.py
@@ -103,7 +103,8 @@ def prepare_nvfp4_moe_layer_for_flashinfer_cutedsl(
     """Prepare weights for the CuteDSL wrapper-based NvFP4 MoE backend.
 
     Converts weight scale factors to MMA layout expected by CuteDslMoEWrapper,
-    and interleaves w13 gate/linear rows.
+    and interleaves w13 gate/linear rows for gated activations. Non-gated
+    activations use a single w13 projection and keep its row order unchanged.
     """
     from flashinfer.cute_dsl.utils import convert_sf_to_mma_layout
 
@@ -112,13 +113,14 @@ def prepare_nvfp4_moe_layer_for_flashinfer_cutedsl(
     a13_scale = a13_scale.max().to(torch.float32).repeat(num_experts)
     a2_scale = a2_scale.max().to(torch.float32).repeat(num_experts)
 
-    half = w13.shape[1] // 2
-    w13 = torch.cat([w13[:, half:], w13[:, :half]], dim=1)
-    w13_scale = torch.cat([w13_scale[:, half:], w13_scale[:, :half]], dim=1)
+    if layer.activation.is_gated:
+        half = w13.shape[1] // 2
+        w13 = torch.cat([w13[:, half:], w13[:, :half]], dim=1)
+        w13_scale = torch.cat([w13_scale[:, half:], w13_scale[:, :half]], dim=1)
 
-    # Interleave up/gate rows for w13 weights and scales.
-    w13 = interleave_linear_and_gate(w13, group_size=64, dim=1)
-    w13_scale = interleave_linear_and_gate(w13_scale, group_size=64, dim=1)
+        # Interleave up/gate rows for w13 weights and scales.
+        w13 = interleave_linear_and_gate(w13, group_size=64, dim=1)
+        w13_scale = interleave_linear_and_gate(w13_scale, group_size=64, dim=1)
 
     # Convert w13 scale factors: linear → swizzled → MMA layout.
     w13_scale = swizzle_blockscale(w13_scale)
diff --git a/vllm/model_executor/layers/quantization/utils/fp8_utils.py b/vllm/model_executor/layers/quantization/utils/fp8_utils.py
index 83e56a4567b..2e4fbdf4c64 100644
--- a/vllm/model_executor/layers/quantization/utils/fp8_utils.py
+++ b/vllm/model_executor/layers/quantization/utils/fp8_utils.py
@@ -34,7 +34,6 @@ from vllm.utils.deep_gemm import (
     transform_sf_into_required_layout,
 )
 from vllm.utils.platform_utils import get_device_name_as_file_name
-from vllm.utils.torch_utils import direct_register_custom_op
 
 logger = init_logger(__name__)
 
@@ -45,39 +44,6 @@ def is_fp8(x: torch.dtype | torch.Tensor) -> bool:
     return x == torch.float8_e4m3fn or x == torch.float8_e4m3fnuz
 
 
-def _triton_per_token_group_quant_fp8_impl(
-    x: torch.Tensor,
-    group_size: int,
-) -> tuple[torch.Tensor, torch.Tensor]:
-    return per_token_group_quant_fp8(
-        x, group_size, column_major_scales=False, use_ue8m0=False
-    )
-
-
-def _triton_per_token_group_quant_fp8_fake(
-    x: torch.Tensor,
-    group_size: int,
-) -> tuple[torch.Tensor, torch.Tensor]:
-    M, N = x.shape
-    x_fp8 = torch.empty((M, N), dtype=current_platform.fp8_dtype(), device=x.device)
-    out_bs = torch.empty(
-        (
-            M,
-            (N + group_size - 1) // group_size,
-        ),
-        dtype=torch.float32,
-        device=x.device,
-    )
-    return x_fp8, out_bs
-
-
-direct_register_custom_op(
-    "triton_per_token_group_quant_fp8",
-    _triton_per_token_group_quant_fp8_impl,
-    fake_impl=_triton_per_token_group_quant_fp8_fake,
-)
-
-
 def input_to_float8(
     x: torch.Tensor, dtype: torch.dtype | None = None
 ) -> tuple[torch.Tensor, torch.Tensor]:
diff --git a/vllm/model_executor/layers/rotary_embedding/mrope.py b/vllm/model_executor/layers/rotary_embedding/mrope.py
index 3c946dd130c..29ce9e5000d 100644
--- a/vllm/model_executor/layers/rotary_embedding/mrope.py
+++ b/vllm/model_executor/layers/rotary_embedding/mrope.py
@@ -5,6 +5,7 @@
 import numpy as np
 import torch
 
+from vllm.platforms import current_platform
 from vllm.triton_utils import tl, triton
 
 from .base import RotaryEmbeddingBase
@@ -24,16 +25,17 @@ def _triton_mrope_forward(
     rd: tl.constexpr,
     pad_n_qh: tl.constexpr,
     pad_n_kh: tl.constexpr,
-    pad_hd: tl.constexpr,
+    pad_rd: tl.constexpr,
     mrope_section_t: tl.constexpr,
     mrope_section_h: tl.constexpr,
     mrope_section_w: tl.constexpr,
     is_interleaved: tl.constexpr,
+    is_neox_style: tl.constexpr,
 ):
     # Adapted from
     # https://github.com/linkedin/Liger-Kernel/blob/main/src/liger_kernel/ops/qwen2vl_mrope.py
     # This version supports flatten input tensors from vllm
-    # and supports cos and sin cache with shape (3, num_tokens, head_dim // 2)
+    # and supports cos and sin cache with shape (3, num_tokens, rotary_dim // 2)
     # instead of (3, bsz, seq_len, head_dim), also supports interleaved rotary
     pid = tl.program_id(0)
     # locate start address
@@ -44,9 +46,9 @@ def _triton_mrope_forward(
     # get the cos(mθ_{i...d/2}) and sin(mθ_{i...d/2}) for token position
     # m of this program instance
     # ####################################################################
-    # Note: cos and sin now have shape (3, num_tokens, head_dim // 2)
+    # Note: cos and sin now have shape (3, num_tokens, rotary_dim // 2)
 
-    # Updated stride calculation for half head_dim
+    # Updated stride calculation for half rotary_dim
     half_rd = rd // 2
     t_cos = cos + pid * half_rd
     h_cos = t_cos + num_tokens * half_rd
@@ -55,12 +57,17 @@ def _triton_mrope_forward(
     h_sin = t_sin + num_tokens * half_rd
     w_sin = h_sin + num_tokens * half_rd
 
-    # Updated offsets for half head_dim
-    cos_offsets = tl.arange(0, pad_hd // 2)
+    # Updated offsets for half rotary_dim
+    cos_offsets = tl.arange(0, pad_rd // 2)
     if is_interleaved:
-        h_mask = ((cos_offsets % 3) == 1) & (cos_offsets <= 3 * mrope_section_h)
-        w_mask = ((cos_offsets % 3) == 2) & (cos_offsets <= 3 * mrope_section_w)
-        t_mask = ~(h_mask | w_mask)
+        valid_mask = cos_offsets < half_rd
+        h_mask = (
+            valid_mask & ((cos_offsets % 3) == 1) & (cos_offsets <= 3 * mrope_section_h)
+        )
+        w_mask = (
+            valid_mask & ((cos_offsets % 3) == 2) & (cos_offsets <= 3 * mrope_section_w)
+        )
+        t_mask = valid_mask & ~(h_mask | w_mask)
     else:
         t_end = mrope_section_t
         h_end = t_end + mrope_section_h
@@ -79,55 +86,74 @@ def _triton_mrope_forward(
     sin_row = t_sin_row + h_sin_row + w_sin_row
 
     # ####################################################################
-    # Load the left and right half of q and k for the current
-    # program instance (i.e. for the current token) separately
+    # Load the two values in each rotary pair for the current token.
+    # NeoX pairs the first and second halves, while GPT-J pairs
+    # adjacent values.
     # ####################################################################
-    # left half of the head
-    first_half_q_offsets = (
-        tl.arange(0, pad_n_qh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    )
-    first_half_k_offsets = (
-        tl.arange(0, pad_n_kh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    )
-    first_q_mask = (tl.arange(0, pad_n_qh)[:, None] < n_qh) & (
-        tl.arange(0, pad_hd // 2)[None, :] < rd // 2
-    )
-    first_k_mask = (tl.arange(0, pad_n_kh)[:, None] < n_kh) & (
-        tl.arange(0, pad_hd // 2)[None, :] < rd // 2
-    )
+    if is_neox_style:
+        rotary_offsets = tl.arange(0, pad_rd // 2)
+        first_q_offsets = tl.arange(0, pad_n_qh)[:, None] * hd + rotary_offsets[None, :]
+        first_k_offsets = tl.arange(0, pad_n_kh)[:, None] * hd + rotary_offsets[None, :]
+        first_q_mask = (tl.arange(0, pad_n_qh)[:, None] < n_qh) & (
+            rotary_offsets[None, :] < rd // 2
+        )
+        first_k_mask = (tl.arange(0, pad_n_kh)[:, None] < n_kh) & (
+            rotary_offsets[None, :] < rd // 2
+        )
 
-    q_tile_1 = tl.load(q_ptr + first_half_q_offsets, mask=first_q_mask, other=0).to(
-        sin_row.dtype
-    )
-    k_tile_1 = tl.load(k_ptr + first_half_k_offsets, mask=first_k_mask, other=0).to(
-        sin_row.dtype
-    )
+        q_tile_1 = tl.load(q_ptr + first_q_offsets, mask=first_q_mask, other=0).to(
+            sin_row.dtype
+        )
+        k_tile_1 = tl.load(k_ptr + first_k_offsets, mask=first_k_mask, other=0).to(
+            sin_row.dtype
+        )
 
-    # right half of the head
-    second_half_q_offsets = first_half_q_offsets + (rd // 2)
-    second_half_k_offsets = first_half_k_offsets + (rd // 2)
-    second_q_mask = first_q_mask
-    second_k_mask = first_k_mask
+        second_q_offsets = first_q_offsets + (rd // 2)
+        second_k_offsets = first_k_offsets + (rd // 2)
+        q_tile_2 = tl.load(q_ptr + second_q_offsets, mask=first_q_mask, other=0).to(
+            sin_row.dtype
+        )
+        k_tile_2 = tl.load(k_ptr + second_k_offsets, mask=first_k_mask, other=0).to(
+            sin_row.dtype
+        )
 
-    q_tile_2 = tl.load(q_ptr + second_half_q_offsets, mask=second_q_mask, other=0).to(
-        sin_row.dtype
-    )
-    k_tile_2 = tl.load(k_ptr + second_half_k_offsets, mask=second_k_mask, other=0).to(
-        sin_row.dtype
-    )
+        new_q_tile_1 = q_tile_1 * cos_row - q_tile_2 * sin_row
+        tl.store(q_ptr + first_q_offsets, new_q_tile_1, mask=first_q_mask)
+        new_q_tile_2 = q_tile_2 * cos_row + q_tile_1 * sin_row
+        tl.store(q_ptr + second_q_offsets, new_q_tile_2, mask=first_q_mask)
 
-    # y = [x1, x2] * [cos, cos] + [-x2, x1] * [sin, sin]
-    # Since cos and sin are now half-size,
-    # we use the same cos_row and sin_row for both halves
-    new_q_tile_1 = q_tile_1 * cos_row - q_tile_2 * sin_row
-    tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-    new_q_tile_2 = q_tile_2 * cos_row + q_tile_1 * sin_row
-    tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
+        new_k_tile_1 = k_tile_1 * cos_row - k_tile_2 * sin_row
+        tl.store(k_ptr + first_k_offsets, new_k_tile_1, mask=first_k_mask)
+        new_k_tile_2 = k_tile_2 * cos_row + k_tile_1 * sin_row
+        tl.store(k_ptr + second_k_offsets, new_k_tile_2, mask=first_k_mask)
+    else:
+        # Load and store adjacent rotary pairs contiguously. Using stride-two
+        # even/odd offsets makes Triton emit scalar 16-bit memory operations on
+        # AMD, while split/interleave only rearranges values in registers.
+        rotary_offsets = tl.arange(0, pad_rd)
+        q_offsets = tl.arange(0, pad_n_qh)[:, None] * hd + rotary_offsets[None, :]
+        k_offsets = tl.arange(0, pad_n_kh)[:, None] * hd + rotary_offsets[None, :]
+        q_mask = (tl.arange(0, pad_n_qh)[:, None] < n_qh) & (
+            rotary_offsets[None, :] < rd
+        )
+        k_mask = (tl.arange(0, pad_n_kh)[:, None] < n_kh) & (
+            rotary_offsets[None, :] < rd
+        )
 
-    new_k_tile_1 = k_tile_1 * cos_row - k_tile_2 * sin_row
-    tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-    new_k_tile_2 = k_tile_2 * cos_row + k_tile_1 * sin_row
-    tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
+        q_tile = tl.load(q_ptr + q_offsets, mask=q_mask, other=0).to(sin_row.dtype)
+        k_tile = tl.load(k_ptr + k_offsets, mask=k_mask, other=0).to(sin_row.dtype)
+        q_tile_1, q_tile_2 = tl.split(tl.reshape(q_tile, (pad_n_qh, pad_rd // 2, 2)))
+        k_tile_1, k_tile_2 = tl.split(tl.reshape(k_tile, (pad_n_kh, pad_rd // 2, 2)))
+
+        new_q_tile_1 = q_tile_1 * cos_row - q_tile_2 * sin_row
+        new_q_tile_2 = q_tile_2 * cos_row + q_tile_1 * sin_row
+        new_q_tile = tl.interleave(new_q_tile_1, new_q_tile_2)
+        tl.store(q_ptr + q_offsets, new_q_tile, mask=q_mask)
+
+        new_k_tile_1 = k_tile_1 * cos_row - k_tile_2 * sin_row
+        new_k_tile_2 = k_tile_2 * cos_row + k_tile_1 * sin_row
+        new_k_tile = tl.interleave(new_k_tile_1, new_k_tile_2)
+        tl.store(k_ptr + k_offsets, new_k_tile, mask=k_mask)
 
 
 def triton_mrope(
@@ -139,23 +165,26 @@ def triton_mrope(
     head_size: int,
     rotary_dim: int,
     mrope_interleaved: bool,
+    is_neox_style: bool,
 ) -> tuple[torch.Tensor, torch.Tensor]:
     """Qwen2VL mrope kernel.
 
     Args:
         q: [num_tokens, num_heads * head_size]
         k: [num_tokens, num_kv_heads * head_size]
-        cos: [3, num_tokens, head_size //2 ]
+        cos: [3, num_tokens, rotary_dim // 2]
             (T/H/W positions with multimodal inputs)
-        sin: [3, num_tokens, head_size //2 ]
+        sin: [3, num_tokens, rotary_dim // 2]
             (T/H/W positions with multimodal inputs)
         mrope_section: [t, h, w]
         head_size: int
+        is_neox_style: Whether rotary pairs use split-half (NeoX) or
+            adjacent (GPT-J) layout.
     """
     n_row, n_q_head_head_dim = q.shape
     n_q_head = n_q_head_head_dim // head_size
     n_kv_head = k.shape[1] // head_size
-    pad_hd = triton.next_power_of_2(head_size)
+    pad_rd = triton.next_power_of_2(rotary_dim)
     pad_n_q_head = triton.next_power_of_2(n_q_head)
     pad_n_kv_head = triton.next_power_of_2(n_kv_head)
 
@@ -166,6 +195,11 @@ def triton_mrope(
     cos = cos.contiguous()
     sin = sin.contiguous()
 
+    # Small adjacent-pair tiles perform best with one wave per program on
+    # ROCm. Keep the existing launch shape for larger rotary dimensions,
+    # NeoX, and other backends.
+    use_single_wave = current_platform.is_rocm() and not is_neox_style and pad_rd <= 64
+    num_warps = 1 if use_single_wave else 4
     _triton_mrope_forward[(n_row,)](
         q,
         k,
@@ -178,11 +212,13 @@ def triton_mrope(
         rotary_dim,
         pad_n_q_head,
         pad_n_kv_head,
-        pad_hd,
+        pad_rd,
         mrope_section[0],
         mrope_section[1],
         mrope_section[2],
         mrope_interleaved,
+        is_neox_style,
+        num_warps=num_warps,
     )
     return q, k
 
@@ -349,6 +385,7 @@ class MRotaryEmbedding(RotaryEmbeddingBase):
                 self.head_size,
                 self.rotary_dim,
                 self.mrope_interleaved,
+                self.is_neox_style,
             )
 
             return q.reshape(query_shape), k.reshape(key_shape)
diff --git a/vllm/model_executor/layers/sparse_attn_indexer.py b/vllm/model_executor/layers/sparse_attn_indexer.py
index 5b8e2bf008e..9a671be5563 100644
--- a/vllm/model_executor/layers/sparse_attn_indexer.py
+++ b/vllm/model_executor/layers/sparse_attn_indexer.py
@@ -8,7 +8,7 @@ import vllm.envs as envs
 from vllm import _custom_ops as ops
 from vllm._aiter_ops import rocm_aiter_ops
 from vllm.compilation.breakable_cudagraph import eager_break_during_capture
-from vllm.config import get_current_vllm_config
+from vllm.config import CUDAGraphMode, get_current_vllm_config
 from vllm.distributed import get_dcp_group, get_pcp_group
 from vllm.forward_context import get_forward_context
 from vllm.logger import init_logger
@@ -310,6 +310,7 @@ def sparse_attn_indexer(
     topk_indices_buffer: torch.Tensor,
     skip_k_cache_insert: bool,
     use_pcp: bool,
+    dense_mha_metadata_layer_name: LayerNameType,
     use_fp4_cache: bool = False,
     dcp_rank: int = 0,
     dcp_world_size: int = 1,
@@ -317,7 +318,8 @@ def sparse_attn_indexer(
     skip_topk_buffer_clear: bool = False,
 ) -> torch.Tensor:
     # careful! this will be None in dummy run
-    attn_metadata = get_forward_context().attn_metadata
+    forward_context = get_forward_context()
+    attn_metadata = forward_context.attn_metadata
     fp8_dtype = current_platform.fp8_dtype()
     k_cache_prefix = _resolve_layer_name(k_cache_prefix)
 
@@ -357,6 +359,7 @@ def sparse_attn_indexer(
             topk_indices_buffer,
             skip_k_cache_insert,
             use_pcp,
+            dense_mha_metadata_layer_name,
             use_fp4_cache,
         )
     attn_metadata_narrowed = attn_metadata[k_cache_prefix]
@@ -402,6 +405,24 @@ def sparse_attn_indexer(
             scale_fmt,
         )
 
+    # The indexer and main MLA may classify the same short extend differently
+    # because they use independent decode thresholds. Only the main MLA route
+    # can determine whether the top-k indices will be consumed.
+    if forward_context.cudagraph_runtime_mode != CUDAGraphMode.FULL:
+        dense_mha_layer = _resolve_layer_name(dense_mha_metadata_layer_name)
+        if dense_mha_layer:
+            mla_metadata = attn_metadata.get(dense_mha_layer)
+            prefill_metadata = getattr(mla_metadata, "prefill", None)
+            if (
+                getattr(prefill_metadata, "use_dense_mha", False)
+                and getattr(mla_metadata, "num_decode_tokens", -1) == 0
+                and not torch.cuda.is_current_stream_capturing()
+            ):
+                # Deliberately leave the buffer untouched. Dense MHA does not
+                # consume top-k indices for this batch; clearing it would be
+                # unnecessary work.
+                return topk_indices_buffer
+
     # The buffer must be pre-filled with -1 (the "no token" sentinel) before the
     # top-k kernels scatter valid indices into it. On the fused deepseek_v32
     # nvidia path, _fused_norm_rope_kernel already cleared the same
@@ -684,6 +705,7 @@ def sparse_attn_indexer_fake(
     topk_indices_buffer: torch.Tensor | None,
     skip_k_cache_insert: bool,
     use_pcp: bool,
+    dense_mha_metadata_layer_name: LayerNameType,
     use_fp4_cache: bool = False,
     dcp_rank: int = 0,
     dcp_world_size: int = 1,
@@ -739,6 +761,7 @@ class SparseAttnIndexer(CustomOp):
         self.topk_indices_buffer = topk_indices_buffer
         self.skip_k_cache_insert = skip_k_cache_insert
         self.use_fp4_cache = use_fp4_cache
+        self.dense_mha_metadata_layer_name = ""
         # DCP scalars are constant for the run; resolve them here (config is set
         # during model construction) and pass them into the custom op, rather
         # than threading them through per-step metadata.
@@ -800,6 +823,7 @@ class SparseAttnIndexer(CustomOp):
             self.topk_indices_buffer,
             self.skip_k_cache_insert,
             self.use_pcp,
+            _encode_layer_name(self.dense_mha_metadata_layer_name),
             self.use_fp4_cache,
             self.dcp_rank,
             self.dcp_world_size,
diff --git a/vllm/model_executor/model_loader/weight_utils.py b/vllm/model_executor/model_loader/weight_utils.py
index db161a58988..e1897a77f10 100644
--- a/vllm/model_executor/model_loader/weight_utils.py
+++ b/vllm/model_executor/model_loader/weight_utils.py
@@ -694,12 +694,10 @@ def _get_checkpoints_size_bytes(files: list[str]) -> int:
 
 
 def _get_available_ram_bytes() -> int:
-    """Return available RAM, honoring cgroup limits on ROCm."""
+    """Return available RAM, honoring cgroup limits."""
     import psutil
 
     host_available = psutil.virtual_memory().available
-    if not current_platform.is_rocm():
-        return host_available
 
     from vllm.utils.cpu_resource_utils import get_cgroup_memory_limit
 
diff --git a/vllm/model_executor/models/deepseek_v2.py b/vllm/model_executor/models/deepseek_v2.py
index ec36c222184..accd305f4a0 100644
--- a/vllm/model_executor/models/deepseek_v2.py
+++ b/vllm/model_executor/models/deepseek_v2.py
@@ -1180,6 +1180,9 @@ class DeepseekV2MLAAttention(nn.Module):
             # never-written topk buffer.
             skip_topk=_skip_topk and not is_mtp_layer,
             non_causal_multi_token_decode=non_causal_multi_token_decode,
+            # Do not skip scoring for MTP layers: their top-k buffer may be
+            # reused by later draft iterations through index sharing.
+            allow_short_prefill_indexer_scoring_skip=not is_mtp_layer,
         )
 
     def forward(
diff --git a/vllm/model_executor/models/glm4_1v.py b/vllm/model_executor/models/glm4_1v.py
index 787e53e8df2..19389d84463 100644
--- a/vllm/model_executor/models/glm4_1v.py
+++ b/vllm/model_executor/models/glm4_1v.py
@@ -1082,8 +1082,8 @@ class Glm4vProcessingInfo(BaseProcessingInfo):
             preprocessed_size = ImageSize(width=image_width, height=image_height)
 
         # NOTE: Frames are padded to be divisible by `temporal_patch_size`
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/qwen2_vl/image_processing_qwen2_vl.py#L294
-        padded_num_frames = num_frames + num_frames % temporal_patch_size
+        # https://github.com/huggingface/transformers/blob/v5.13.0/src/transformers/models/qwen2_vl/video_processing_qwen2_vl.py#L249-L252
+        padded_num_frames = num_frames + (-num_frames % temporal_patch_size)
 
         grid_t = max(padded_num_frames // temporal_patch_size, 1)
         grid_h = preprocessed_size.height // patch_size
diff --git a/vllm/model_executor/models/kanana_v.py b/vllm/model_executor/models/kanana_v.py
index 125d7e71c7b..b1a5f78b1b3 100644
--- a/vllm/model_executor/models/kanana_v.py
+++ b/vllm/model_executor/models/kanana_v.py
@@ -409,8 +409,8 @@ class KananaVProcessingInfo(BaseProcessingInfo):
             preprocessed_size = ImageSize(width=image_width, height=image_height)
 
         # NOTE: Frames are padded to be divisible by `temporal_patch_size`
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/qwen2_vl/image_processing_qwen2_vl.py#L294
-        padded_num_frames = num_frames + num_frames % temporal_patch_size
+        # https://github.com/huggingface/transformers/blob/v5.13.0/src/transformers/models/qwen2_vl/video_processing_qwen2_vl.py#L249-L252
+        padded_num_frames = num_frames + (-num_frames % temporal_patch_size)
 
         grid_t = max(padded_num_frames // temporal_patch_size, 1)
         grid_h = preprocessed_size.height // patch_size
diff --git a/vllm/model_executor/models/keye.py b/vllm/model_executor/models/keye.py
index dd1fb892ad1..c3d69836a79 100644
--- a/vllm/model_executor/models/keye.py
+++ b/vllm/model_executor/models/keye.py
@@ -983,7 +983,7 @@ class KeyeProcessingInfo(BaseProcessingInfo):
         else:
             preprocessed_size = ImageSize(width=image_width, height=image_height)
 
-        padded_num_frames = num_frames + num_frames % temporal_patch_size
+        padded_num_frames = num_frames + (-num_frames % temporal_patch_size)
 
         grid_t = max(padded_num_frames // temporal_patch_size, 1)
         grid_h = preprocessed_size.height // patch_size
diff --git a/vllm/model_executor/models/llava_onevision2.py b/vllm/model_executor/models/llava_onevision2.py
index 58179ec00de..9dc11d493a0 100644
--- a/vllm/model_executor/models/llava_onevision2.py
+++ b/vllm/model_executor/models/llava_onevision2.py
@@ -1345,7 +1345,7 @@ class LlavaOnevision2ProcessingInfo(BaseProcessingInfo):
             preprocessed = ImageSize(width=rw, height=rh)
         else:
             preprocessed = ImageSize(width=image_width, height=image_height)
-        padded_frames = num_frames + num_frames % temporal_patch_size
+        padded_frames = num_frames + (-num_frames % temporal_patch_size)
         grid_t = max(padded_frames // temporal_patch_size, 1)
         grid_h = preprocessed.height // patch_size
         grid_w = preprocessed.width // patch_size
diff --git a/vllm/model_executor/models/mimo_v2_omni.py b/vllm/model_executor/models/mimo_v2_omni.py
index d0d9589ae1d..747cb0e88b2 100644
--- a/vllm/model_executor/models/mimo_v2_omni.py
+++ b/vllm/model_executor/models/mimo_v2_omni.py
@@ -715,7 +715,7 @@ class MiMoV2OmniProcessingInfo(BaseProcessingInfo):
             effective_frames = num_frames * tokens_per_second
         else:
             effective_frames = num_frames
-        padded_num_frames = effective_frames + effective_frames % temporal_patch_size
+        padded_num_frames = effective_frames + (-effective_frames % temporal_patch_size)
         grid_t = max(padded_num_frames // temporal_patch_size, 1)
         grid_h = preprocessed_size.height // patch_size
         grid_w = preprocessed_size.width // patch_size
diff --git a/vllm/model_executor/models/qwen2_vl.py b/vllm/model_executor/models/qwen2_vl.py
index 539f141cbaa..e2e9f245248 100644
--- a/vllm/model_executor/models/qwen2_vl.py
+++ b/vllm/model_executor/models/qwen2_vl.py
@@ -898,8 +898,8 @@ class Qwen2VLProcessingInfo(BaseProcessingInfo):
             preprocessed_size = ImageSize(width=image_width, height=image_height)
 
         # NOTE: Frames are padded to be divisible by `temporal_patch_size`
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/qwen2_vl/image_processing_qwen2_vl.py#L294
-        padded_num_frames = num_frames + num_frames % temporal_patch_size
+        # https://github.com/huggingface/transformers/blob/v5.13.0/src/transformers/models/qwen2_vl/video_processing_qwen2_vl.py#L249-L252
+        padded_num_frames = num_frames + (-num_frames % temporal_patch_size)
 
         grid_t = max(padded_num_frames // temporal_patch_size, 1)
         grid_h = preprocessed_size.height // patch_size
diff --git a/vllm/model_executor/models/transformers/__init__.py b/vllm/model_executor/models/transformers/__init__.py
index 9dbfe4b5031..ff4bb6cd433 100644
--- a/vllm/model_executor/models/transformers/__init__.py
+++ b/vllm/model_executor/models/transformers/__init__.py
@@ -64,12 +64,15 @@ def vllm_attention_forward(
     head_dim_v = value.shape[-1]
     query, key, value = (x.transpose(1, 2) for x in (query, key, value))
     query, key, value = (x.reshape(hidden, -1) for x in (query, key, value))
-    # Pad `value` up to the query/key head size when they differ (expanded MLA).
-    if head_dim_v != head_dim_qk:
+    # Pad `value` up to the query/key head size when it is smaller (expanded
+    # MLA). A larger last dim just means `value` isn't split per head, e.g.
+    # packed grouped/multi-query projections, and needs no padding.
+    pad_value = head_dim_v < head_dim_qk
+    if pad_value:
         value = F.pad(value.view(-1, head_dim_v), (0, head_dim_qk - head_dim_v))
         value = value.reshape(hidden, -1)
     attn_output = self_attn.forward(query, key, value)
-    if head_dim_v != head_dim_qk:
+    if pad_value:
         attn_output = attn_output.view(-1, head_dim_qk)[..., :head_dim_v]
         attn_output = attn_output.reshape(hidden, -1)
     return attn_output, None
diff --git a/vllm/model_executor/warmup/kernel_warmup.py b/vllm/model_executor/warmup/kernel_warmup.py
index ba7f4cae393..537a2311fba 100644
--- a/vllm/model_executor/warmup/kernel_warmup.py
+++ b/vllm/model_executor/warmup/kernel_warmup.py
@@ -96,7 +96,7 @@ def _warmup_ll_bf16_router_gemm(model: torch.nn.Module) -> None:
     )
 
 
-def kernel_warmup(worker: "Worker"):
+def kernel_warmup(worker: "Worker", *, process_local_only: bool = False):
     from vllm.model_executor.warmup.minimax_m3_msa_warmup import (
         minimax_m3_msa_warmup,
     )
@@ -105,7 +105,8 @@ def kernel_warmup(worker: "Worker"):
         # Pooling models do not use the generation slot-mapping path.
         if not worker.model_runner.is_pooling_model:
             warm_v1_block_table_kernels(worker.model_runner)
-        # No dummy run reaches the scheduler-driven KV-block zeroing path.
+        # The KV-block zeroing kernel is driven by the scheduler's
+        # `new_block_ids_to_zero`, so no dummy run ever reaches it.
         zeroer = getattr(worker.model_runner, "_kv_block_zeroer", None)
         if zeroer is not None:
             zeroer.warmup(worker.model_runner.kv_cache_config.num_blocks)
@@ -124,6 +125,23 @@ def kernel_warmup(worker: "Worker"):
     )
 
     # Run next so input-prep kernels JIT against pristine runner state.
+    if worker.vllm_config.kernel_config.enable_jit_warmup:
+        kimi_k3_triton_warmup(worker)
+        fa4_cutedsl_warmup(worker)
+        sparse_mla_triton_warmup(worker)
+
+    if current_platform.has_device_capability(90):
+        _warmup_ll_bf16_router_gemm(worker.get_model())
+
+    if worker.vllm_config.kernel_config.enable_cutedsl_warmup:
+        # TODO(roberto): Remove after registered CuTeDSL warmups are migrated
+        # to the shared JIT warmup infrastructure.
+        # https://github.com/vllm-project/vllm/pull/47451
+        cutedsl_warmup()
+
+    if process_local_only:
+        return
+
     flashinfer_sparse_mla_decode_autotune_warmup(worker)
     deepseek_v4_sparse_mla_attention_warmup(worker)
 
@@ -149,9 +167,6 @@ def kernel_warmup(worker: "Worker"):
     elif has_flashinfer() and current_platform.has_device_capability(90):
         flashinfer_autotune(worker.model_runner)
 
-    if current_platform.has_device_capability(90):
-        _warmup_ll_bf16_router_gemm(worker.get_model())
-
     # FlashInfer attention warmup
     # Only warmup if the model has FlashInfer attention groups
     # and is not a pooling model
@@ -184,17 +199,6 @@ def kernel_warmup(worker: "Worker"):
             create_mixed_batch=True,
         )
 
-    if worker.vllm_config.kernel_config.enable_cutedsl_warmup:
-        # TODO(roberto): Remove after registered CuTeDSL warmups are migrated
-        # to the shared JIT warmup infrastructure.
-        # https://github.com/vllm-project/vllm/pull/47451
-        cutedsl_warmup()
-
-    if worker.vllm_config.kernel_config.enable_jit_warmup:
-        kimi_k3_triton_warmup(worker)
-        fa4_cutedsl_warmup(worker)
-        sparse_mla_triton_warmup(worker)
-
 
 def _flashinfer_autotune_skip_ops(runner: "GPUModelRunner") -> set[str] | None:
     if envs.VLLM_FLASHINFER_AUTOTUNE_SKIP_OPS is not None:
diff --git a/vllm/models/deepseek_v32/nvidia/attention.py b/vllm/models/deepseek_v32/nvidia/attention.py
index dcf955ad59b..0e604da87ba 100644
--- a/vllm/models/deepseek_v32/nvidia/attention.py
+++ b/vllm/models/deepseek_v32/nvidia/attention.py
@@ -494,8 +494,10 @@ class DeepseekV32Attention(MLAAttention):
                 self.indexer.max_model_len,
                 self.indexer.max_total_seq_len,
                 self.topk_indices_buffer,
-                True,  # skip_k_cache_insert
-                False,  # use_fp4_cache
+                skip_k_cache_insert=True,
+                use_pcp=False,
+                dense_mha_metadata_layer_name="",
+                use_fp4_cache=False,
                 # fused_norm_rope already cleared the topk buffer this forward.
                 skip_topk_buffer_clear=True,
             )
diff --git a/vllm/models/deepseek_v4/quant_config.py b/vllm/models/deepseek_v4/quant_config.py
index 89cf695baf0..293d71f2f41 100644
--- a/vllm/models/deepseek_v4/quant_config.py
+++ b/vllm/models/deepseek_v4/quant_config.py
@@ -120,7 +120,11 @@ class DeepseekV4FP8Config(Fp8Config):
     @staticmethod
     def _is_quark_mxfp4_ocp(hf_quant_cfg: dict) -> bool:
         """True for AMD-Quark exports whose global scheme is MXFP4."""
-        weight = (hf_quant_cfg.get("global_quant_config") or {}).get("weight") or {}
+        weight = (hf_quant_cfg.get("global_quant_config") or {}).get("weight")
+        # A non-dict weight (e.g. a list of multiple specs) means not an OCP
+        # MXFP4 scheme (e.g. NVFP4 with 2-level scale).
+        if not isinstance(weight, dict):
+            return False
         return (
             weight.get("dtype") == "fp4"
             and weight.get("qscheme") == "per_group"
diff --git a/vllm/models/inkling/nvidia/moe.py b/vllm/models/inkling/nvidia/moe.py
index 32f7489a4c4..9ff3acaa14d 100644
--- a/vllm/models/inkling/nvidia/moe.py
+++ b/vllm/models/inkling/nvidia/moe.py
@@ -589,6 +589,9 @@ class InklingMoE(nn.Module):
             param.data[lids] = vals.reshape(len(gids), *param.shape[1:]).to(
                 param.device
             )
+        elif key == "w2_weight_scale" and weight.shape[-1] == 1:
+            # Per-output-channel scales are replicated across TP ranks.
+            param.data[lids] = weight[gids].to(device=param.device, dtype=param.dtype)
         elif key.startswith("w13"):
             # Checkpoint w13 rows are interleaved [g0, u0, g1, u1, ...]; the
             # fused param layout is [w1(gate); w3(up)]. The TP-local rows form
diff --git a/vllm/models/kimi_k3/amd/linear.py b/vllm/models/kimi_k3/amd/linear.py
index c833750512e..f27c81a24f1 100644
--- a/vllm/models/kimi_k3/amd/linear.py
+++ b/vllm/models/kimi_k3/amd/linear.py
@@ -407,7 +407,8 @@ class KimiMLAAttention(nn.Module):
                 prefix=f"{prefix}.g_proj",
             )
 
-        mla_modules = MLAModules(
+        # TODO: Remove this mypy workaround once the K3 PR is fully merged.
+        mla_modules = MLAModules(  # type: ignore[call-arg]
             kv_a_layernorm=self.kv_a_layernorm,
             kv_b_proj=self.kv_b_proj,
             rotary_emb=None,
diff --git a/vllm/models/kimi_k3/amd/ops/__init__.py b/vllm/models/kimi_k3/amd/ops/__init__.py
index e69de29bb2d..208f01a7cb5 100644
--- a/vllm/models/kimi_k3/amd/ops/__init__.py
+++ b/vllm/models/kimi_k3/amd/ops/__init__.py
@@ -0,0 +1,2 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
diff --git a/vllm/models/kimi_k3/common/mm_preprocess.py b/vllm/models/kimi_k3/common/mm_preprocess.py
index 81db5faa2be..fa70ccdce0a 100644
--- a/vllm/models/kimi_k3/common/mm_preprocess.py
+++ b/vllm/models/kimi_k3/common/mm_preprocess.py
@@ -328,102 +328,3 @@ class KimiK3MultiModalProcessor(BaseMultiModalProcessor[KimiK3ProcessingInfo]):
                 replacement=get_replacement,
             ),
         ]
-
-    @staticmethod
-    def preprocess_messages(
-        messages: list[dict],
-    ) -> list[dict]:
-        """Reorder tool-result messages to match assistant tool_call order.
-
-        Supports matching by ``tool_call_id``, by synthetic
-        ``"{tool}:{zero_based_index}"`` alias, or by an explicit
-        ``(tool, index)`` pair carried on the message.  When any tool
-        message in a block cannot be resolved the whole block is left
-        in its original order (graceful fallback).
-
-        Returns a new list; caller-owned message dicts are not mutated.
-        """
-        normalized: list[dict] = []
-        i = 0
-
-        while i < len(messages):
-            message = messages[i]
-            normalized.append(message)
-            i += 1
-
-            if message.get("role") != "assistant":
-                continue
-
-            tool_calls = message.get("tool_calls")
-            if not isinstance(tool_calls, list) or not tool_calls:
-                continue
-
-            # Build lookup tables from the assistant's tool_calls.
-            targets_by_id: dict[str, tuple[str, int, int]] = {}
-            targets_by_pair: dict[tuple[str, int], tuple[str, int, int]] = {}
-            for index, tool_call in enumerate(tool_calls, start=1):
-                function = tool_call.get("function")
-                if not isinstance(function, dict):
-                    continue
-                tool = function.get("name")
-                if not isinstance(tool, str) or not tool:
-                    continue
-
-                call_target = (tool, index, index - 1)
-                targets_by_pair[(tool, index)] = call_target
-
-                aliases = [f"{tool}:{index - 1}"]
-                tool_call_id = tool_call.get("id")
-                if tool_call_id:
-                    aliases.insert(0, str(tool_call_id))
-                for alias in aliases:
-                    targets_by_id.setdefault(alias, call_target)
-
-            if not targets_by_pair:
-                continue
-
-            # Collect consecutive tool messages.
-            block_start = i
-            while i < len(messages) and messages[i].get("role") == "tool":
-                i += 1
-            if i == block_start:
-                continue
-
-            # Try to resolve every tool message in the block.
-            keyed: list[tuple[int, int, dict]] = []
-            unresolved = False
-            for order, tool_message in enumerate(messages[block_start:i]):
-                resolved: tuple[str, int, int] | None = None
-
-                # 1. Match by tool_call_id
-                tool_call_id = tool_message.get("tool_call_id")
-                if tool_call_id is not None:
-                    resolved = targets_by_id.get(str(tool_call_id))
-
-                # 2. Match by (tool, index) pair
-                if resolved is None:
-                    tool = tool_message.get("tool") or tool_message.get("name")
-                    raw_index = tool_message.get("index")
-                    if tool is not None and raw_index is not None:
-                        try:
-                            resolved = targets_by_pair.get((str(tool), int(raw_index)))
-                        except (TypeError, ValueError):
-                            resolved = None
-
-                if resolved is None:
-                    unresolved = True
-                    break
-
-                tool_name, xtml_index, sort_order = resolved
-                enriched = dict(tool_message)
-                enriched["tool"] = tool_name
-                enriched["index"] = xtml_index
-                keyed.append((sort_order, order, enriched))
-
-            if unresolved:
-                normalized.extend(messages[block_start:i])
-            else:
-                keyed.sort(key=lambda item: (item[0], item[1]))
-                normalized.extend(msg for _, _, msg in keyed)
-
-        return normalized
diff --git a/vllm/models/kimi_k3/nvidia/dspark_mla.py b/vllm/models/kimi_k3/nvidia/dspark_mla.py
index 14ef4b8aa2a..f9692d5d472 100644
--- a/vllm/models/kimi_k3/nvidia/dspark_mla.py
+++ b/vllm/models/kimi_k3/nvidia/dspark_mla.py
@@ -35,7 +35,8 @@ class ReplicatedDSparkMarkovHead(DSparkMarkovHead):
         markov_rank: int,
         prefix: str,
     ) -> None:
-        super().__init__(
+        # TODO: Remove this mypy workaround once the K3 PR is fully merged.
+        super().__init__(  # type: ignore[call-arg]
             vocab_size,
             draft_vocab_size,
             markov_rank,
diff --git a/vllm/models/kimi_k3/nvidia/kda.py b/vllm/models/kimi_k3/nvidia/kda.py
index a4b06156a89..d4fec1e0084 100644
--- a/vllm/models/kimi_k3/nvidia/kda.py
+++ b/vllm/models/kimi_k3/nvidia/kda.py
@@ -6,22 +6,20 @@ from collections.abc import Callable
 import torch
 from einops import rearrange
 from torch import nn
+from torch.nn.parameter import Parameter
 
 from vllm import _custom_ops as ops
 from vllm.compilation.breakable_cudagraph import eager_break_during_capture
 from vllm.config import VllmConfig
-from vllm.distributed import divide
+from vllm.distributed import divide, get_tensor_model_parallel_rank
 from vllm.forward_context import get_forward_context
 from vllm.logger import init_logger
 from vllm.model_executor.layers.linear import (
     ColumnParallelLinear,
+    MergedColumnParallelLinear,
     RowParallelLinear,
 )
 from vllm.model_executor.layers.mamba.gdn.base import GatedDeltaNetAttention
-from vllm.model_executor.layers.mamba.gdn.kimi_gdn_linear_attn import (
-    _KimiGDNMergedColumnParallelLinear,
-    a_log_weight_loader,
-)
 from vllm.model_executor.layers.mamba.mamba_utils import (
     MambaStateDtypeCalculator,
     MambaStateShapeCalculator,
@@ -38,15 +36,14 @@ from vllm.model_executor.model_loader.weight_utils import (
     default_weight_loader,
     sharded_weight_loader,
 )
+from vllm.model_executor.parameter import BasevLLMParameter
 from vllm.model_executor.utils import set_weight_attrs
 from vllm.models.kimi_k3.nvidia.kda_metadata import (
     KimiK3KDAAttentionBackend,
     KimiK3KDAMetadata,
 )
 from vllm.platforms import current_platform
-from vllm.third_party.flash_linear_attention.ops.fused_norm_gate import (
-    FusedRMSNormGated,
-)
+from vllm.third_party.flash_linear_attention.ops.kda import FusedRMSNormGated
 from vllm.transformers_utils.configs.kimi_linear import KimiLinearConfig
 from vllm.v1.attention.backend import AttentionBackend
 
@@ -55,6 +52,86 @@ logger = init_logger(__name__)
 _KDA_GATE_LOGBOUND_MIN = -5.0
 
 
+def a_log_weight_loader(
+    shard_axis: int,
+) -> Callable[[torch.Tensor, torch.Tensor], None]:
+    """Load KDA A_log stored as either old 4D or current 1D weights."""
+
+    def loader(param: torch.Tensor, loaded_weight: torch.Tensor) -> None:
+        tp_rank = get_tensor_model_parallel_rank()
+        shard_size = param.data.shape[shard_axis]
+        start_idx = tp_rank * shard_size
+
+        if loaded_weight.dim() == 4:
+            assert loaded_weight.shape[:2] == (1, 1), (
+                f"Expected old A_log shape (1, 1, H, 1), got {loaded_weight.shape}"
+            )
+            assert loaded_weight.shape[-1] == 1, (
+                f"Expected old A_log last dim to be 1, got {loaded_weight.shape}"
+            )
+            loaded_weight = loaded_weight.view(loaded_weight.shape[2])
+
+        loaded_weight = loaded_weight.narrow(shard_axis, start_idx, shard_size)
+        return default_weight_loader(param, loaded_weight)
+
+    return loader
+
+
+class _KimiGDNMergedColumnParallelLinear(MergedColumnParallelLinear):
+    """Merged projection with one output replicated across TP ranks."""
+
+    def __init__(
+        self,
+        input_size: int,
+        output_sizes: list[int],
+        replicated_shard_id: int,
+        tp_size: int,
+        **kwargs,
+    ) -> None:
+        self.replicated_shard_id = replicated_shard_id
+        output_sizes = output_sizes.copy()
+        output_sizes[replicated_shard_id] *= tp_size
+        super().__init__(input_size, output_sizes, **kwargs)
+
+    def weight_loader(
+        self,
+        param: Parameter,
+        loaded_weight: torch.Tensor,
+        loaded_shard_id: tuple[int, ...] | int | None = None,
+    ) -> None:
+        tp_rank = self.tp_rank
+        param_tp_rank = getattr(param, "tp_rank", None)
+        if loaded_shard_id == self.replicated_shard_id:
+            self.tp_rank = 0
+            if param_tp_rank is not None:
+                param.tp_rank = 0
+        try:
+            super().weight_loader(param, loaded_weight, loaded_shard_id)
+        finally:
+            self.tp_rank = tp_rank
+            if param_tp_rank is not None:
+                param.tp_rank = param_tp_rank
+
+    def weight_loader_v2(
+        self,
+        param: BasevLLMParameter,
+        loaded_weight: torch.Tensor,
+        loaded_shard_id: tuple[int, ...] | int | None = None,
+    ) -> None:
+        tp_rank = self.tp_rank
+        param_tp_rank = getattr(param, "tp_rank", None)
+        if loaded_shard_id == self.replicated_shard_id:
+            self.tp_rank = 0
+            if param_tp_rank is not None:
+                param.tp_rank = 0
+        try:
+            super().weight_loader_v2(param, loaded_weight, loaded_shard_id)
+        finally:
+            self.tp_rank = tp_rank
+            if param_tp_rank is not None:
+                param.tp_rank = param_tp_rank
+
+
 def is_fused_kda_decode_supported(
     num_heads: int,
     head_dim: int,
diff --git a/vllm/models/kimi_k3/nvidia/mla.py b/vllm/models/kimi_k3/nvidia/mla.py
index c6eee08e12d..3334f3bbbb4 100644
--- a/vllm/models/kimi_k3/nvidia/mla.py
+++ b/vllm/models/kimi_k3/nvidia/mla.py
@@ -335,7 +335,8 @@ class MultiHeadLatentAttention(nn.Module, AttentionLayerBase):
         kv_cache_dtype = kv_cache_dtype_str_to_dtype(
             self.kv_cache_dtype, vllm_config.model_config
         )
-        return MLAAttentionSpec(
+        # TODO: Remove this mypy workaround once the K3 PR is fully merged.
+        return MLAAttentionSpec(  # type: ignore[call-arg]
             block_size=vllm_config.cache_config.block_size,
             num_kv_heads=1,
             head_size=self.head_size,
diff --git a/vllm/multimodal/inputs.py b/vllm/multimodal/inputs.py
index 8da41afa765..62450d4e6c7 100644
--- a/vllm/multimodal/inputs.py
+++ b/vllm/multimodal/inputs.py
@@ -460,6 +460,8 @@ class BaseMultiModalField(ABC):
             device = "cpu"
         if pin_memory and self.keep_on_cpu:
             pin_memory = False
+        if device == "cpu" or device == torch.device("cpu"):
+            pin_memory = False
 
         batch = [elem.data for elem in elems]
         out = self._reduce_data(batch, pin_memory=pin_memory)
diff --git a/vllm/pooling_params.py b/vllm/pooling_params.py
index 6cb130fdbbb..5280e997c9c 100644
--- a/vllm/pooling_params.py
+++ b/vllm/pooling_params.py
@@ -7,6 +7,7 @@ from typing import Any
 import msgspec
 
 from vllm.config import ModelConfig, PoolerConfig
+from vllm.exceptions import VLLMValidationError
 from vllm.logger import init_logger
 from vllm.sampling_params import RequestOutputKind
 from vllm.tasks import PoolingTask, check_removed_pooling_task
@@ -145,7 +146,7 @@ class PoolingParams(
                     invalid_parameters.append(k)
 
             if invalid_parameters:
-                raise ValueError(
+                raise VLLMValidationError(
                     f"Task {self.task} only supports {valid_parameters} "
                     f"parameters, does not support "
                     f"{invalid_parameters} parameters"
@@ -170,21 +171,21 @@ class PoolingParams(
                 valid_range = f"[1, {embedding_size}]"
                 dimensions_in_range = 1 <= dimensions <= embedding_size
                 if not model_config.is_matryoshka:
-                    raise ValueError(
+                    raise VLLMValidationError(
                         f"Model {model_name!r} does not support Matryoshka "
                         f"embeddings; dimensions must be unset "
                         f"(received dimensions={dimensions})."
                     )
 
                 if not dimensions_in_range:
-                    raise ValueError(
+                    raise VLLMValidationError(
                         f"Model {model_name!r} only supports dimensions in "
                         f"range {valid_range}, got {dimensions}."
                     )
 
                 mds = model_config.matryoshka_dimensions
                 if mds is not None and dimensions not in mds:
-                    raise ValueError(
+                    raise VLLMValidationError(
                         f"Model {model_name!r} only supports Matryoshka "
                         f"dimensions {str(mds)}, got {dimensions}."
                     )
@@ -208,7 +209,7 @@ class PoolingParams(
                 invalid_parameters.append(k)
 
         if invalid_parameters:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"Task {self.task!r} only supports {valid_parameters} "
                 f"parameters, does not support "
                 f"{invalid_parameters} parameters"
@@ -231,7 +232,7 @@ class PoolingParams(
     def __post_init__(self) -> None:
         check_removed_pooling_task(self.task)
         if self.output_kind != RequestOutputKind.FINAL_ONLY:
-            raise ValueError(
+            raise VLLMValidationError(
                 "For pooling output_kind has to be FINAL_ONLY, "
                 f"got {self.output_kind!r}"
             )
diff --git a/vllm/sampling_params.py b/vllm/sampling_params.py
index 25e36ceb568..5dedbde372e 100644
--- a/vllm/sampling_params.py
+++ b/vllm/sampling_params.py
@@ -100,12 +100,12 @@ class StructuredOutputsParams:
             ]
         )
         if count > 1:
-            raise ValueError(
+            raise VLLMValidationError(
                 "You can only use one kind of structured outputs constraint "
                 f"but multiple are specified: {self.__dict__}"
             )
         if count < 1:
-            raise ValueError(
+            raise VLLMValidationError(
                 "You must use one kind of structured outputs constraint "
                 f"but none are specified: {self.__dict__}"
             )
@@ -166,13 +166,13 @@ class RepetitionDetectionParams:
             or self.min_pattern_size < 0
             or self.min_pattern_size > self.max_pattern_size
         ):
-            raise ValueError(
+            raise VLLMValidationError(
                 "max_pattern_size, min_pattern_size must be >=0, "
                 "with min_pattern_size <= max_pattern_size. "
                 "Set both to 0 to disable repetitive pattern detection."
             )
         if self.max_pattern_size > 0 and self.min_count < 2:
-            raise ValueError(
+            raise VLLMValidationError(
                 "min_count must be >= 2 to detect repetitive patterns "
                 "in engine output. If you do not wish to detect repetitive "
                 "patterns, set max_pattern_size to 0."
@@ -514,31 +514,33 @@ class SamplingParams(
 
     def _verify_args(self) -> None:
         if not isinstance(self.n, int):
-            raise ValueError(f"n must be an int, but is of type {type(self.n)}")
+            raise VLLMValidationError(
+                f"n must be an int, but is of type {type(self.n)}"
+            )
         if self.n < 1:
-            raise ValueError(f"n must be at least 1, got {self.n}.")
+            raise VLLMValidationError(f"n must be at least 1, got {self.n}.")
         max_n = envs.VLLM_MAX_N_SEQUENCES
         if self.n > max_n:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"n must be at most {max_n}, got {self.n}. "
                 "To increase this limit, set the VLLM_MAX_N_SEQUENCES "
                 "environment variable."
             )
         if not -2.0 <= self.presence_penalty <= 2.0:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"presence_penalty must be in [-2, 2], got {self.presence_penalty}."
             )
         if not -2.0 <= self.frequency_penalty <= 2.0:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"frequency_penalty must be in [-2, 2], got {self.frequency_penalty}."
             )
         if not math.isfinite(self.repetition_penalty):
-            raise ValueError(
+            raise VLLMValidationError(
                 "repetition_penalty must be a finite number, "
                 f"got {self.repetition_penalty}."
             )
         if self.repetition_penalty <= 0.0:
-            raise ValueError(
+            raise VLLMValidationError(
                 "repetition_penalty must be greater than zero, got "
                 f"{self.repetition_penalty}."
             )
@@ -568,15 +570,15 @@ class SamplingParams(
             )
         # quietly accept -1 as disabled, but prefer 0
         if self.top_k < -1:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"top_k must be 0 (disable), or at least 1, got {self.top_k}."
             )
         if not isinstance(self.top_k, int):
-            raise TypeError(
+            raise VLLMValidationError(
                 f"top_k must be an integer, got {type(self.top_k).__name__}"
             )
         if not 0.0 <= self.min_p <= 1.0:
-            raise ValueError(f"min_p must be in [0, 1], got {self.min_p}.")
+            raise VLLMValidationError(f"min_p must be in [0, 1], got {self.min_p}.")
         if self.max_tokens is not None and self.max_tokens < 1:
             raise VLLMValidationError(
                 f"max_tokens must be at least 1, got {self.max_tokens}.",
@@ -584,11 +586,11 @@ class SamplingParams(
                 value=self.max_tokens,
             )
         if self.min_tokens < 0:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"min_tokens must be greater than or equal to 0, got {self.min_tokens}."
             )
         if self.max_tokens is not None and self.min_tokens > self.max_tokens:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"min_tokens must be less than or equal to "
                 f"max_tokens={self.max_tokens}, got {self.min_tokens}."
             )
@@ -617,27 +619,29 @@ class SamplingParams(
             )
         assert isinstance(self.stop_token_ids, list)
         if not all(isinstance(st_id, int) for st_id in self.stop_token_ids):
-            raise ValueError(
+            raise VLLMValidationError(
                 f"stop_token_ids must contain only integers, got {self.stop_token_ids}."
             )
         assert isinstance(self.stop, list)
         if any(not stop_str for stop_str in self.stop):
-            raise ValueError("stop cannot contain an empty string.")
+            raise VLLMValidationError("stop cannot contain an empty string.")
         if self.stop and not self.detokenize:
-            raise ValueError(
+            raise VLLMValidationError(
                 "stop strings are only supported when detokenize is True. "
                 "Set detokenize=True to use stop."
             )
         assert isinstance(self.bad_words, list)
         if any(not bad_word for bad_word in self.bad_words):
-            raise ValueError(
+            raise VLLMValidationError(
                 f"bad_words cannot contain an empty string. "
                 f"Got bad_words={self.bad_words}"
             )
 
     def _verify_greedy_sampling(self) -> None:
         if self.n > 1:
-            raise ValueError(f"n must be 1 when using greedy sampling, got {self.n}.")
+            raise VLLMValidationError(
+                f"n must be 1 when using greedy sampling, got {self.n}."
+            )
 
     def update_from_generation_config(
         self,
@@ -889,7 +893,7 @@ class SamplingParams(
 
         # Some sampling parameters are not yet compatible with spec decoding.
         if self.min_p > _SAMPLING_EPS or self.logit_bias:
-            raise ValueError(
+            raise VLLMValidationError(
                 "The min_p and logit_bias sampling parameters "
                 "are not yet supported with speculative decoding."
             )
@@ -910,7 +914,7 @@ class SamplingParams(
             or self.bad_words
             or self.allowed_token_ids
         ):
-            raise ValueError(
+            raise VLLMValidationError(
                 "The temperature, min_p, seed, min_tokens, logit_bias, "
                 "bad_words, and allowed_token_ids sampling parameters "
                 "are not yet supported with diffusion models."
@@ -930,7 +934,7 @@ class SamplingParams(
             # rather than sampling left-to-right, which the grammar FSM
             # requires. Without this check, requests fail mid-generation
             # with an FSM rejection (HTTP 500). See issue #45436.
-            raise ValueError(
+            raise VLLMValidationError(
                 "Structured outputs are not yet supported for diffusion "
                 "language models. Remove the structured output constraint "
                 "(e.g. `response_format`, `structured_outputs`) from the "
@@ -938,7 +942,7 @@ class SamplingParams(
             )
 
         if tokenizer is None:
-            raise ValueError(
+            raise VLLMValidationError(
                 "Structured outputs requires a tokenizer so it can't be used with 'skip_tokenizer_init'"  # noqa: E501
             )
 
@@ -952,7 +956,7 @@ class SamplingParams(
             if backend != _backend and not (
                 backend == "auto" and self.structured_outputs._backend_was_auto
             ):
-                raise ValueError(
+                raise VLLMValidationError(
                     "Request-level structured output backend selection is not "
                     f"supported. The request specified '{_backend}', but vLLM "
                     f"was initialised with '{backend}'. This error can be "
@@ -967,7 +971,7 @@ class SamplingParams(
             and not self.structured_outputs.choice
         ):
             # It is invalid for choice to be an empty list
-            raise ValueError(
+            raise VLLMValidationError(
                 f"Choice '{self.structured_outputs.choice}' cannot be an empty list"  # noqa: E501
             )
         # Reject empty string grammar early to avoid engine-side crashes
@@ -975,16 +979,20 @@ class SamplingParams(
             isinstance(self.structured_outputs.grammar, str)
             and self.structured_outputs.grammar.strip() == ""
         ):
-            raise ValueError("structured_outputs.grammar cannot be an empty string")
+            raise VLLMValidationError(
+                "structured_outputs.grammar cannot be an empty string"
+            )
         # Reject empty string json schema early to avoid engine-side crashes
         if (
             isinstance(self.structured_outputs.json, str)
             and self.structured_outputs.json.strip() == ""
         ):
-            raise ValueError("structured_outputs.json cannot be an empty string")
+            raise VLLMValidationError(
+                "structured_outputs.json cannot be an empty string"
+            )
         # Reject json_object=False early to avoid engine-side crashes
         if self.structured_outputs.json_object is False:
-            raise ValueError(
+            raise VLLMValidationError(
                 "structured_outputs.json_object must be True if set; omit "
                 "structured_outputs to disable structured outputs"
             )
@@ -1006,7 +1014,7 @@ class SamplingParams(
             validate_xgrammar_grammar(self)
         elif backend.startswith("guidance"):
             if _is_non_tekken_mistral(tokenizer=tokenizer):
-                raise ValueError(
+                raise VLLMValidationError(
                     "Non-tekken Mistral tokenizers are not supported for the 'guidance'"
                     " structured output backend. Please either use a more recent "
                     "Mistral model, the ['xgrammar', 'outlines'] "
@@ -1026,7 +1034,7 @@ class SamplingParams(
         elif backend == "lm-format-enforcer":
             # lm format enforcer backend
             if is_mistral_tokenizer(tokenizer):
-                raise ValueError(
+                raise VLLMValidationError(
                     "Mistral tokenizer is not supported for the 'lm-format-enforcer' "
                     "structured output backend. Please use ['xgrammar', 'outlines'] "
                     "backends or tokenizer_mode='hf' instead."
diff --git a/vllm/v1/attention/backends/flex_attention.py b/vllm/v1/attention/backends/flex_attention.py
index e7fe5852194..83144751aeb 100644
--- a/vllm/v1/attention/backends/flex_attention.py
+++ b/vllm/v1/attention/backends/flex_attention.py
@@ -12,6 +12,7 @@ import torch
 import torch._dynamo.decorators
 import torch.nn.functional as F
 from torch.nn.attention.flex_attention import (
+    AuxRequest,
     BlockMask,
     _mask_mod_signature,
     _score_mod_signature,
@@ -1228,6 +1229,9 @@ class FlexAttentionImpl(AttentionImpl):
         if block_n is not None:
             self.block_n = block_n
 
+        # Optional post-attention epilogue transform
+        self.out_transform = kwargs.get("out_transform")
+
     @staticmethod
     def view_as_4d(tensor: torch.Tensor) -> torch.Tensor:
         """View a 3d tensor as 4D."""
@@ -1392,8 +1396,13 @@ class FlexAttentionImpl(AttentionImpl):
             self.scale,
             enable_gqa=enable_gqa,
             kernel_options=kernel_options,
+            return_aux=AuxRequest(lse=True) if self.out_transform is not None else None,
         )
 
+        if self.out_transform is not None:
+            out, aux = out
+            out = self.out_transform(out, aux.lse)
+
         # Flex doesn't have an out variant today, rely on epilogue fusion
         out = out.permute(0, 2, 1, 3).squeeze(0)
         output[:num_actual_tokens, :, :].copy_(out)
diff --git a/vllm/v1/attention/ops/rocm_aiter_mla_sparse.py b/vllm/v1/attention/ops/rocm_aiter_mla_sparse.py
index 6a28324208f..63bad3cfad0 100644
--- a/vllm/v1/attention/ops/rocm_aiter_mla_sparse.py
+++ b/vllm/v1/attention/ops/rocm_aiter_mla_sparse.py
@@ -442,6 +442,7 @@ def rocm_fp8_paged_mqa_logits(
                 KVBlockSize=block_size,
                 WavePerEU=2,
             )
+            out_logits.nan_to_num_(float("-inf"))
             return out_logits
         deepgemm_fp8_paged_mqa_logits_stage1 = (
             aiter_paged_mqa_logits_module.deepgemm_fp8_paged_mqa_logits_stage1
diff --git a/vllm/v1/engine/__init__.py b/vllm/v1/engine/__init__.py
index e80be0e45d7..83033ecf81d 100644
--- a/vllm/v1/engine/__init__.py
+++ b/vllm/v1/engine/__init__.py
@@ -35,8 +35,6 @@ FT_STATUS_CALL_ID = -2
 
 
 class EEPNotificationType(enum.Enum):
-    NEW_CORE_ENGINES_INIT_READY = "NEW_CORE_ENGINES_INIT_READY"
-    NEW_CORE_ENGINES_WEIGHTS_INIT_READY = "NEW_CORE_ENGINES_WEIGHTS_INIT_READY"
     RECONFIGURE_FINISHED = "RECONFIGURE_FINISHED"
     SHUTDOWN_COMPLETE = "SHUTDOWN_COMPLETE"
 
diff --git a/vllm/v1/engine/async_llm.py b/vllm/v1/engine/async_llm.py
index 922a8aa5982..1718a5adc39 100644
--- a/vllm/v1/engine/async_llm.py
+++ b/vllm/v1/engine/async_llm.py
@@ -21,6 +21,7 @@ from vllm.distributed.weight_transfer.base import (
 from vllm.engine.arg_utils import AsyncEngineArgs
 from vllm.engine.protocol import EngineClient, StreamingInput
 from vllm.entrypoints.serve.elastic_ep.middleware import set_scaling_elastic_ep
+from vllm.exceptions import VLLMClientError, VLLMValidationError
 from vllm.inputs import EngineInput, PromptType
 from vllm.logger import init_logger
 from vllm.lora.request import LoRARequest
@@ -109,6 +110,7 @@ class AsyncLLM(EngineClient):
         maybe_register_config_serialize_by_value()
 
         self.vllm_config = vllm_config
+        self._elastic_ep_lock = asyncio.Lock()
         self.model_config = vllm_config.model_config
         self.observability_config = vllm_config.observability_config
 
@@ -308,7 +310,7 @@ class AsyncLLM(EngineClient):
             and not is_pooling
             and params.prompt_logprobs
         ):
-            raise ValueError(
+            raise VLLMValidationError(
                 "--kv-sharing-fast-prefill produces incorrect logprobs for "
                 "prompt tokens, please disable it when the requests need "
                 "prompt logprobs"
@@ -475,7 +477,7 @@ class AsyncLLM(EngineClient):
                     )
                     req.external_req_id = request_id
                     if req.prompt_embeds is not None:
-                        raise ValueError(
+                        raise VLLMValidationError(
                             "prompt_embeds not supported for streaming inputs"
                         )
                     prompt_text, _, _ = extract_prompt_components(
@@ -511,7 +513,7 @@ class AsyncLLM(EngineClient):
             or params.output_kind == RequestOutputKind.FINAL_ONLY
             or params.stop
         ):
-            raise ValueError(
+            raise VLLMValidationError(
                 "Input streaming not currently supported "
                 "for pooling models, n > 1, request_kind = FINAL_ONLY "
                 "or with stop strings."
@@ -603,7 +605,7 @@ class AsyncLLM(EngineClient):
             raise
 
         # Request validation error.
-        except ValueError as e:
+        except VLLMClientError as e:
             if self.log_requests:
                 logger.info("Request %s failed (bad request): %s.", request_id, e)
             raise
@@ -868,7 +870,7 @@ class AsyncLLM(EngineClient):
             raise
 
         # Request validation error.
-        except ValueError:
+        except VLLMClientError:
             if self.log_requests:
                 logger.info("Request %s failed (bad request).", request_id)
             raise
@@ -998,17 +1000,27 @@ class AsyncLLM(EngineClient):
             "waiting for requests to drain."
         )
 
+    async def _drain_requests_for_elastic_ep(self, drain_timeout: int) -> None:
+        try:
+            logger.info(
+                "VLLM_ELASTIC_EP_DRAIN_REQUESTS is set, "
+                "waiting for requests to drain before scaling"
+            )
+            await self.wait_for_requests_to_drain(drain_timeout)
+        except BaseException:
+            set_scaling_elastic_ep(False)
+            raise
+
     async def scale_elastic_ep(
         self, new_data_parallel_size: int, drain_timeout: int = 300
     ):
-        """
-        Scale up or down the data parallel size by adding or removing
-        engine cores.
-        Args:
-            new_data_parallel_size: The new number of data parallel workers
-            drain_timeout:
-                Maximum time to wait for requests to drain (seconds)
-        """
+        """Scale the elastic EP data parallel size."""
+        async with self._elastic_ep_lock:
+            await self._scale_elastic_ep(new_data_parallel_size, drain_timeout)
+
+    async def _scale_elastic_ep(
+        self, new_data_parallel_size: int, drain_timeout: int
+    ) -> None:
         old_data_parallel_size = self.vllm_config.parallel_config.data_parallel_size
         if old_data_parallel_size == new_data_parallel_size:
             logger.info(
@@ -1017,12 +1029,7 @@ class AsyncLLM(EngineClient):
             )
             return
 
-        if envs.VLLM_ELASTIC_EP_DRAIN_REQUESTS:
-            logger.info(
-                "VLLM_ELASTIC_EP_DRAIN_REQUESTS is set, "
-                "waiting for requests to drain before scaling"
-            )
-            await self.wait_for_requests_to_drain(drain_timeout)
+        await self.engine_core.prepare_elastic_ep(new_data_parallel_size)
 
         # recreate stat loggers
         if new_data_parallel_size > old_data_parallel_size and self.log_stats:
@@ -1042,11 +1049,12 @@ class AsyncLLM(EngineClient):
             self.logger_manager.log_engine_initialized()
 
         set_scaling_elastic_ep(True)
-        try:
-            await self.engine_core.scale_elastic_ep(new_data_parallel_size)
-            self.vllm_config.parallel_config.data_parallel_size = new_data_parallel_size
-        finally:
-            set_scaling_elastic_ep(False)
+        if envs.VLLM_ELASTIC_EP_DRAIN_REQUESTS:
+            await self._drain_requests_for_elastic_ep(drain_timeout)
+
+        await self.engine_core.commit_elastic_ep()
+        self.vllm_config.parallel_config.data_parallel_size = new_data_parallel_size
+        set_scaling_elastic_ep(False)
 
     async def handle_fault(
         self, fault_tolerance_request: FaultToleranceRequest
@@ -1106,6 +1114,16 @@ class AsyncLLM(EngineClient):
             "update_weights", kwargs={"update_info": request.update_info}
         )
 
-    async def finish_weight_update(self) -> None:
-        """Finish the current weight update."""
+    async def finish_weight_update(self, weight_version: str | None = None) -> None:
+        """Finish the weight update and set its version if provided."""
         await self.collective_rpc("finish_weight_update")
+        if weight_version is not None:
+            await self.update_weight_version(weight_version)
+
+    async def update_weight_version(self, new_version: str) -> None:
+        """Set the weight version without updating weights."""
+        await self.engine_core.set_weight_version_async(new_version)
+
+    async def get_weight_version(self) -> str:
+        """Return the latest committed weight version."""
+        return await self.engine_core.get_weight_version_async()
diff --git a/vllm/v1/engine/coordinator.py b/vllm/v1/engine/coordinator.py
index 2f3b03636d7..d7f05cffc8a 100644
--- a/vllm/v1/engine/coordinator.py
+++ b/vllm/v1/engine/coordinator.py
@@ -376,6 +376,9 @@ class DPCoordinatorProc:
                     eng_index = outputs.engine_index
                     scheduler_stats = outputs.scheduler_stats
                     if scheduler_stats:
+                        # Elastic EP stats may arrive while the engine list changes.
+                        if eng_index >= len(self.engines):
+                            continue
                         # 1. Updated request load stats - update our local
                         # state with these.
                         stats = self.engines[eng_index].request_counts
diff --git a/vllm/v1/engine/core.py b/vllm/v1/engine/core.py
index 9817c474343..66135273002 100644
--- a/vllm/v1/engine/core.py
+++ b/vllm/v1/engine/core.py
@@ -7,7 +7,7 @@ import signal
 import threading
 import time
 from collections import defaultdict, deque
-from collections.abc import Callable, Generator
+from collections.abc import Callable, Generator, Sequence
 from concurrent.futures import Future
 from contextlib import ExitStack, contextmanager
 from enum import IntEnum
@@ -87,7 +87,7 @@ from vllm.v1.kv_cache_interface import KVCacheConfig, get_kv_cache_spec_kind
 from vllm.v1.metrics.stats import SchedulerIterationDetails, SchedulerStats
 from vllm.v1.outputs import ModelRunnerOutput
 from vllm.v1.request import Request, RequestStatus
-from vllm.v1.serial_utils import MsgpackDecoder, MsgpackEncoder
+from vllm.v1.serial_utils import MsgpackDecoder, MsgpackEncoder, bytestr
 from vllm.v1.structured_output import StructuredOutputManager
 from vllm.v1.utils import compute_iteration_details
 from vllm.version import __version__ as VLLM_VERSION
@@ -125,6 +125,8 @@ class EngineCore:
             )
 
         self.log_stats = log_stats
+        # Opaque weight version supplied by the caller.
+        self._weight_version = "default"
 
         # Setup Model.
         self.model_executor = executor_class(vllm_config)
@@ -325,8 +327,9 @@ class EngineCore:
 
         vllm_config.validate_block_size()
 
-        # Initialize kv cache and warmup the execution
         self.model_executor.initialize_from_config(kv_cache_configs)
+        if not envs.VLLM_ELASTIC_EP_SCALE_UP_LAUNCH:
+            self.model_executor.compile_or_warm_up_model()
 
         elapsed = time.time() - start
         compile_time = vllm_config.compilation_config.compilation_time
@@ -956,6 +959,13 @@ class EngineCore:
     ) -> list[_R]:
         return self.model_executor.collective_rpc(method, timeout, args, kwargs)
 
+    def set_weight_version(self, weight_version: str) -> None:
+        self._weight_version = weight_version
+
+    def get_weight_version(self) -> str:
+        """Return the latest committed weight version."""
+        return self._weight_version
+
     def preprocess_add_request(self, request: EngineCoreRequest) -> tuple[Request, int]:
         """Preprocess the request.
 
@@ -984,9 +994,7 @@ class EngineCore:
         raise NotImplementedError
 
     def _eep_send_engine_core_notification(
-        self,
-        notification_type: EEPNotificationType,
-        vllm_config: VllmConfig | None = None,
+        self, notification_type: EEPNotificationType
     ):
         raise NotImplementedError
 
@@ -1061,11 +1069,6 @@ class EngineCoreProc(EngineCore):
 
             self.addresses = addresses
             self.process_input_queue_block = True
-            if envs.VLLM_ELASTIC_EP_SCALE_UP_LAUNCH:
-                self._eep_send_engine_core_notification(
-                    EEPNotificationType.NEW_CORE_ENGINES_INIT_READY,
-                    vllm_config=vllm_config,
-                )
             self._init_data_parallel(vllm_config)
 
             super().__init__(
@@ -1745,10 +1748,11 @@ class EngineCoreProc(EngineCore):
         encoder = MsgpackEncoder()
         # Send buffers to reuse.
         reuse_buffers: list[bytearray] = []
-        # Keep references to outputs and buffers until zmq is finished
-        # with them (outputs may contain tensors/np arrays whose
-        # backing buffers were extracted for zero-copy send).
-        pending = deque[tuple[zmq.MessageTracker, Any, bytearray]]()
+        # Payload buffers that can't be reused yet because zmq may still be
+        # sending them.
+        # Buffers of the zero-copy tensor/ndarray frames don't need tracking
+        # here: zmq itself holds a reference to each until it's done with it.
+        pending = deque[tuple[zmq.MessageTracker, bytearray]]()
 
         # We must set linger to ensure the ENGINE_CORE_DEAD
         # message is sent prior to closing the socket.
@@ -1789,20 +1793,38 @@ class EngineCoreProc(EngineCore):
 
                 # Reclaim buffers that zmq is finished with.
                 while pending and pending[-1][0].done:
-                    reuse_buffers.append(pending.pop()[2])
+                    reclaimed = pending.pop()[1]
+                    if len(reuse_buffers) < max_reuse_bufs:
+                        reuse_buffers.append(reclaimed)
 
                 buffer = reuse_buffers.pop() if reuse_buffers else bytearray()
                 buffers = encoder.encode_into(outputs, buffer)
-                tracker = sockets[client_index].send_multipart(
-                    buffers, copy=False, track=True
+                tracker = self._send_msg_tracking_payload(
+                    sockets[client_index], buffers
                 )
                 if not tracker.done:
-                    ref = outputs if len(buffers) > 1 else None
-                    pending.appendleft((tracker, ref, buffer))
+                    pending.appendleft((tracker, buffer))
                 elif len(reuse_buffers) < max_reuse_bufs:
                     # Limit the number of buffers to reuse.
                     reuse_buffers.append(buffer)
 
+    @staticmethod
+    def _send_msg_tracking_payload(
+        socket: zmq.Socket, buffers: Sequence[bytestr]
+    ) -> zmq.MessageTracker:
+        """Send `buffers` as a zero-copy multipart message, returning a tracker
+        for the *first* frame.
+
+        Used instead of `Socket.send_multipart()` because we reuse the buffer
+        passed to `MsgpackEncoder.encode_into()`: `send_multipart()` returns a
+        tracker for the last frame only.
+        """
+        more_flag = zmq.SNDMORE if len(buffers) > 1 else 0
+        tracker = socket.send(buffers[0], more_flag, copy=False, track=True)
+        if more_flag:
+            socket.send_multipart(buffers[1:], copy=False)
+        return tracker
+
     def _handle_request_preproc_error(self, request: EngineCoreRequest) -> None:
         """Log and return a request-scoped error response for exceptions raised
         from the add request preprocessing in the input socket processing thread.
@@ -2087,12 +2109,16 @@ class DPEngineCoreProc(EngineCoreProc):
             self._maybe_publish_request_counts()
 
             if self.eep_scaling_state is not None:
-                _ = self.eep_scaling_state.progress()
-                if self.eep_scaling_state.is_complete():
-                    if self.eep_scaling_state.worker_type == "removing":
+                state = self.eep_scaling_state
+                if state.commit_requested or not state.is_ready_for_switch():
+                    state.progress()
+                if state.is_complete():
+                    if state.worker_type == "removing":
                         raise SystemExit
                     self.process_input_queue_block = True
                     self.eep_scaling_state = None
+                elif not state.commit_requested and state.is_ready_for_switch():
+                    self.process_input_queue_block = True
 
             executed = self._process_engine_step()
             self._maybe_publish_request_counts()
@@ -2162,7 +2188,7 @@ class DPEngineCoreProc(EngineCoreProc):
 
     def reinitialize_distributed(
         self, reconfig_request: ReconfigureDistributedRequest
-    ) -> None:
+    ) -> str:
         from copy import deepcopy
 
         from vllm.distributed.elastic_ep.elastic_state import ElasticEPScalingState
@@ -2194,7 +2220,10 @@ class DPEngineCoreProc(EngineCoreProc):
             == ReconfigureRankType.SHUTDOWN_CURRENT_RANK
         )
 
-        self.eep_scaling_state = ElasticEPScalingState(
+        if self.eep_scaling_state is not None:
+            raise RuntimeError("Elastic EP reconfiguration is already active")
+
+        state = ElasticEPScalingState(
             model_executor=self.model_executor,
             engine_core=self,
             vllm_config=self.vllm_config,
@@ -2203,30 +2232,34 @@ class DPEngineCoreProc(EngineCoreProc):
             scale_type="scale_down" if is_scale_down else "scale_up",
             reconfig_request=reconfig_request,
         )
+        self.eep_scaling_state = state
+
         self.process_input_queue_block = False
         logger.info(
             "[Elastic EP] Received reconfiguration request and starting scaling up/down"
         )
+        return state.ready_key
+
+    def commit_prepared_elastic_ep(self) -> None:
+        state = self.eep_scaling_state
+        if state is None or state.commit_requested or not state.is_ready_for_switch():
+            raise RuntimeError("No prepared Elastic EP reconfiguration is ready")
+        state.commit_requested = True
+        self.process_input_queue_block = False
+        logger.info("[Elastic EP] Committing prepared reconfiguration")
 
     def _eep_send_engine_core_notification(
-        self,
-        notification_type: EEPNotificationType,
-        vllm_config: VllmConfig | None = None,
+        self, notification_type: EEPNotificationType
     ):
         """
         Send notifications to EngineCoreClient, which can then forward
         the notifications to other engine core processes. It is used for:
-        1) In scale up: new core engines to notify existing core engines
-           that they are ready;
-        2) In scale down: removing core engines to notify EngineCoreClient
+        1) In scale down: removing core engines to notify EngineCoreClient
            so EngineCoreClient can release their ray placement groups;
-        3) Both scale up/down: to notify EngineCoreClient that existing
+        2) Both scale up/down: to notify EngineCoreClient that existing
            core engines have already switched to the new parallel setup.
         """
-        if vllm_config is None:
-            dp_rank = self.vllm_config.parallel_config.data_parallel_rank
-        else:
-            dp_rank = vllm_config.parallel_config.data_parallel_rank
+        dp_rank = self.vllm_config.parallel_config.data_parallel_rank
         notification_data = (notification_type.value, dp_rank)
         outputs = EngineCoreOutputs(
             utility_output=UtilityOutput(
@@ -2248,22 +2281,11 @@ class DPEngineCoreProc(EngineCoreProc):
             ):
                 socket.send_multipart(encoder.encode(outputs))
 
-    def eep_handle_engine_core_notification(
-        self, notification_type: str | EEPNotificationType
-    ):
-        """
-        Handle notification received from EngineCoreClient
-        (forwarded from new core engines).
-        """
-        assert self.eep_scaling_state is not None
-        if isinstance(notification_type, str):
-            notification_type = EEPNotificationType(notification_type)
-        self.eep_scaling_state.handle_notification(notification_type)
-
     def _eep_scale_up_before_kv_init(self):
         from vllm.distributed.elastic_ep.elastic_state import ElasticEPScalingState
 
-        self.eep_scaling_state = ElasticEPScalingState(
+        self.ignore_start_dp_wave = True
+        state = ElasticEPScalingState(
             model_executor=self.model_executor,
             engine_core=self,
             vllm_config=self.vllm_config,
@@ -2272,7 +2294,10 @@ class DPEngineCoreProc(EngineCoreProc):
             scale_type="scale_up",
             reconfig_request=None,
         )
-        self.eep_scaling_state.run_pre_kv_init_states()
+        if self.eep_scaling_state is not None:
+            raise RuntimeError("Elastic EP reconfiguration is already active")
+        self.eep_scaling_state = state
+        state.run_pre_kv_init_states()
         self.process_input_queue_block = False
 
 
diff --git a/vllm/v1/engine/core_client.py b/vllm/v1/engine/core_client.py
index 0aa4b6f3312..febaa10ce61 100644
--- a/vllm/v1/engine/core_client.py
+++ b/vllm/v1/engine/core_client.py
@@ -7,7 +7,7 @@ import sys
 import uuid
 import weakref
 from abc import ABC, abstractmethod
-from collections import Counter, defaultdict, deque
+from collections import Counter, defaultdict
 from collections.abc import Awaitable, Callable, Sequence
 from concurrent.futures import Future
 from dataclasses import dataclass
@@ -176,9 +176,21 @@ class EngineCoreClient(ABC):
     def execute_dummy_batch(self) -> None:
         raise NotImplementedError
 
+    def set_weight_version(self, weight_version: str) -> None:
+        raise NotImplementedError
+
+    def get_weight_version(self) -> str:
+        raise NotImplementedError
+
     async def execute_dummy_batch_async(self) -> None:
         raise NotImplementedError
 
+    async def set_weight_version_async(self, weight_version: str) -> None:
+        raise NotImplementedError
+
+    async def get_weight_version_async(self) -> str:
+        raise NotImplementedError
+
     def abort_requests(self, request_ids: list[str]) -> None:
         raise NotImplementedError
 
@@ -213,7 +225,10 @@ class EngineCoreClient(ABC):
         running state."""
         raise NotImplementedError
 
-    async def scale_elastic_ep(self, new_data_parallel_size: int) -> None:
+    async def commit_elastic_ep(self) -> None:
+        raise NotImplementedError
+
+    async def prepare_elastic_ep(self, new_data_parallel_size: int) -> None:
         raise NotImplementedError
 
     async def get_output_async(self) -> EngineCoreOutputs:
@@ -351,6 +366,12 @@ class InprocClient(EngineCoreClient):
     def execute_dummy_batch(self) -> None:
         self.engine_core.execute_dummy_batch()
 
+    def set_weight_version(self, weight_version: str) -> None:
+        self.engine_core.set_weight_version(weight_version)
+
+    def get_weight_version(self) -> str:
+        return self.engine_core.get_weight_version()
+
     def add_lora(self, lora_request: LoRARequest) -> bool:
         return self.engine_core.add_lora(lora_request)
 
@@ -650,11 +671,6 @@ class MPClient(EngineCoreClient):
             self.core_engine: EngineIdentity = self.core_engines[0]
             self.utility_results: dict[int, AnyFuture] = {}
 
-            # Request objects which may contain pytorch-allocated tensors
-            # that we need to keep references to until zmq is done with the
-            # underlying data.
-            self.pending_messages = deque[tuple[zmq.MessageTracker, Any]]()
-
             # Start monitoring engine core processes for unexpected failures
             self.start_engine_core_monitor()
 
@@ -686,14 +702,6 @@ class MPClient(EngineCoreClient):
         if self.resources.engine_dead:
             raise EngineDeadError()
 
-    def add_pending_message(self, tracker: zmq.MessageTracker, msg: Any):
-        if not tracker.done:
-            self.pending_messages.appendleft((tracker, msg))
-
-    def free_pending_messages(self):
-        while self.pending_messages and self.pending_messages[-1][0].done:
-            self.pending_messages.pop()
-
     def dp_engines_running(self) -> bool:
         return self.engines_running
 
@@ -875,17 +883,12 @@ class SyncMPClient(MPClient):
 
     def _send_input(self, request_type: EngineCoreRequestType, request: Any):
         self.ensure_alive()
-        self.free_pending_messages()
         # (Identity, RequestType, SerializedRequest)
         msg = (self.core_engine, request_type.value, *self.encoder.encode(request))
-
-        if len(msg) <= 3:
-            # No auxiliary buffers => no tensor backing buffers in request.
-            self.input_socket.send_multipart(msg, copy=False)
-            return
-
-        tracker = self.input_socket.send_multipart(msg, copy=False, track=True)
-        self.add_pending_message(tracker, request)
+        # Any zero-copy tensor/ndarray frames are kept alive by zmq itself
+        # until it's finished sending them (there is a ref chain from the underlying
+        # memoryview back to the original owning tensor/ndarray).
+        self.input_socket.send_multipart(msg, copy=False)
 
     def call_utility(self, method: str, *args) -> Any:
         call_id = uuid.uuid1().int >> 64
@@ -947,6 +950,12 @@ class SyncMPClient(MPClient):
     def execute_dummy_batch(self) -> None:
         self.call_utility("execute_dummy_batch")
 
+    def set_weight_version(self, weight_version: str) -> None:
+        self.call_utility("set_weight_version", weight_version)
+
+    def get_weight_version(self) -> str:
+        return self.call_utility("get_weight_version")
+
     def collective_rpc(
         self,
         method: str | Callable[..., _R],
@@ -1102,32 +1111,16 @@ class AsyncMPClient(MPClient):
             engine = self.core_engine
 
         message = (request_type.value, *self.encoder.encode(request))
-        return self._send_input_message(message, engine, request)
+        return self._send_input_message(message, engine)
 
     def _send_input_message(
-        self, message: tuple[bytestr, ...], engine: EngineIdentity, objects: Any
+        self, message: tuple[bytestr, ...], engine: EngineIdentity
     ) -> Awaitable[Any]:
-        """
-        objects is a reference to retain until zmq is finished with the
-        buffers, in case they were extracted from tensors in the request.
-        """
         self.ensure_alive()
-        self.free_pending_messages()
-
-        msg = (engine,) + message
-        if not objects or len(msg) <= 3:
-            # No auxiliary buffers => no tensor backing buffers in request.
-            return self.input_socket.send_multipart(msg, copy=False)
-
-        future: asyncio.Future[zmq.MessageTracker]
-        future = self.input_socket.send_multipart(msg, copy=False, track=True)
-
-        def add_pending(f: asyncio.Future[zmq.MessageTracker]):
-            with contextlib.suppress(BaseException):
-                self.add_pending_message(f.result(), objects)
-
-        future.add_done_callback(add_pending)
-        return future
+        # Any zero-copy tensor/ndarray frames are kept alive by zmq itself
+        # until it's finished sending them (there is a ref chain from the underlying
+        # memoryview back to the original owning tensor/ndarray).
+        return self.input_socket.send_multipart((engine,) + message, copy=False)
 
     async def call_utility_async(self, method: str, *args) -> Any:
         return await self._call_utility_async(method, *args, engine=self.core_engine)
@@ -1142,7 +1135,7 @@ class AsyncMPClient(MPClient):
             EngineCoreRequestType.UTILITY.value,
             *self.encoder.encode((self.client_index, call_id, method, args)),
         )
-        await self._send_input_message(message, engine, args)
+        await self._send_input_message(message, engine)
         self._ensure_output_queue_task()
         return await future
 
@@ -1199,6 +1192,12 @@ class AsyncMPClient(MPClient):
     async def execute_dummy_batch_async(self) -> None:
         await self.call_utility_async("execute_dummy_batch")
 
+    async def set_weight_version_async(self, weight_version: str) -> None:
+        await self.call_utility_async("set_weight_version", weight_version)
+
+    async def get_weight_version_async(self) -> str:
+        return await self.call_utility_async("get_weight_version")
+
     async def add_lora_async(self, lora_request: LoRARequest) -> bool:
         return await self.call_utility_async("add_lora", lora_request)
 
@@ -1460,6 +1459,7 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
         )
 
         assert len(self.core_engines) > 1
+        self._prepared_elastic_ep: tuple[int, int] | None = None
 
         self.eng_start_index = (
             len(self.core_engines) * self.client_index
@@ -1573,31 +1573,16 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
         if len(cache.pending_notifications[notification_type]) >= abs(
             cache.num_new_core_engines
         ):
-            if notification_type == EEPNotificationType.SHUTDOWN_COMPLETE:
-                assert isinstance(self.resources.engine_manager, CoreEngineActorManager)
-                assert cache.num_new_core_engines < 0
-                old_dp_size = len(cache.existing_core_engines)
-                new_dp_size = old_dp_size + cache.num_new_core_engines
-                self.resources.engine_manager.scale_down_elastic_ep(
-                    old_dp_size, new_dp_size
-                )
-            else:
-                await asyncio.gather(
-                    *[
-                        self._call_utility_async(
-                            "eep_handle_engine_core_notification",
-                            notification_type,
-                            engine=engine,
-                        )
-                        for engine in cache.existing_core_engines
-                    ]
-                )
-            cache.pending_notifications[notification_type] = set()
-            if notification_type in [
-                EEPNotificationType.SHUTDOWN_COMPLETE,
-                EEPNotificationType.NEW_CORE_ENGINES_WEIGHTS_INIT_READY,
-            ]:
-                self.eep_scaling_cache = None
+            engine_manager = self.resources.engine_manager
+            assert isinstance(engine_manager, CoreEngineActorManager)
+            assert cache.num_new_core_engines < 0
+            old_dp_size = len(cache.existing_core_engines)
+            new_dp_size = old_dp_size + cache.num_new_core_engines
+            engine_manager.scale_down_elastic_ep(old_dp_size, new_dp_size)
+            self.vllm_config.parallel_config.data_parallel_size_local = len(
+                engine_manager.local_engine_actors
+            )
+            self.eep_scaling_cache = None
 
     async def abort_requests_async(self, request_ids: list[str]) -> None:
         if not request_ids or self.resources.engine_dead:
@@ -1621,31 +1606,46 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
     ) -> None:
         await self._send_input(EngineCoreRequestType.ABORT, request_ids, engine)
 
-    async def scale_elastic_ep(self, new_data_parallel_size: int) -> None:
-        """Scale elastic EP data parallel size"""
+    async def commit_elastic_ep(self) -> None:
+        """Commit prepared elastic EP scaling."""
+        prepared = self._prepared_elastic_ep
+        if prepared is None:
+            raise RuntimeError("Elastic EP scaling has not been prepared")
+        new_data_parallel_size, num_redundant_experts = prepared
         cur_data_parallel_size = len(self.core_engines)
-
-        assert new_data_parallel_size != cur_data_parallel_size, (
-            f"new_data_parallel_size {new_data_parallel_size} must be "
-            f"different from cur_data_parallel_size {cur_data_parallel_size}"
+        if new_data_parallel_size > cur_data_parallel_size:
+            await self._commit_scale_up_elastic_ep(new_data_parallel_size)
+        else:
+            await self._commit_scale_down_elastic_ep(new_data_parallel_size)
+        self.vllm_config.parallel_config.eplb_config.num_redundant_experts = (
+            num_redundant_experts
         )
+        self._prepared_elastic_ep = None
 
+    async def prepare_elastic_ep(self, new_data_parallel_size: int) -> None:
+        """Prepare elastic EP scaling without routing requests to new engines."""
+        if (prepared := self._prepared_elastic_ep) is not None:
+            if prepared[0] == new_data_parallel_size:
+                return
+            raise RuntimeError("Elastic EP scaling is already prepared")
+        cur_data_parallel_size = len(self.core_engines)
         assert self.vllm_config.parallel_config.data_parallel_backend == "ray", (
             "Only ray DP backend supports scaling elastic EP"
         )
-
-        scale_up = new_data_parallel_size > cur_data_parallel_size
-
-        if scale_up:
-            await self._scale_up_elastic_ep(
-                cur_data_parallel_size, new_data_parallel_size
-            )
+        parallel_config = self.vllm_config.parallel_config
+        num_experts = self.vllm_config.model_config.get_num_experts()
+        num_redundant_experts = (
+            num_experts + parallel_config.eplb_config.num_redundant_experts
+        ) * new_data_parallel_size // cur_data_parallel_size - num_experts
+        if new_data_parallel_size < cur_data_parallel_size:
+            await self._prepare_scale_down_elastic_ep(new_data_parallel_size)
         else:
-            await self._scale_down_elastic_ep(
-                cur_data_parallel_size, new_data_parallel_size
+            await self._prepare_scale_up_elastic_ep(
+                new_data_parallel_size, num_redundant_experts
             )
+        self._prepared_elastic_ep = new_data_parallel_size, num_redundant_experts
 
-    async def _eep_wait_for_setup_switch_complete(self) -> None:
+    def _eep_wait_for_setup_switch_complete(self) -> asyncio.Future:
         """
         Wait for core engines to switch to the new setup.
 
@@ -1657,9 +1657,26 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
         future = asyncio.get_running_loop().create_future()
         self.utility_results[EEP_NOTIFICATION_CALL_ID] = future
         self._ensure_output_queue_task()
-        await future
+        return future
 
-    def _setup_elastic_ep_reconfig_bootstrap(self) -> tuple[str, int]:
+    def _wait_for_new_engine_ready(self, new_core_engines: list[bytes]) -> None:
+        new_engine_identities = set(new_core_engines)
+        sync_input_socket = zmq.Socket.shadow(self.input_socket)
+        while new_engine_identities:
+            if not sync_input_socket.poll(timeout=VLLM_ENGINE_READY_TIMEOUT_S * 1000):
+                raise TimeoutError(
+                    f"Timed out waiting for new engine core processes to "
+                    f"start. Waited "
+                    f"{VLLM_ENGINE_READY_TIMEOUT_S}s (configured by "
+                    f"VLLM_ENGINE_READY_TIMEOUT_S). To increase the "
+                    f"timeout, set the environment variable: "
+                    f"VLLM_ENGINE_READY_TIMEOUT_S=<seconds>"
+                )
+            identity, payload = sync_input_socket.recv_multipart()
+            new_engine_identities.discard(identity)
+            self._apply_ready_response(payload)
+
+    def _setup_elastic_ep_reconfig_bootstrap(self) -> None:
         from vllm.distributed.utils import create_tcp_store
         from vllm.utils.network_utils import get_open_ports_list
 
@@ -1679,36 +1696,36 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
         )
         parallel_config._coord_store_port = store.port
         self._coord_store = store
-        return ip, store.port
 
-    async def _scale_up_elastic_ep(
-        self, cur_data_parallel_size: int, new_data_parallel_size: int
-    ) -> None:
-        """Scale up the data parallel size by creating new engine cores
-        and reconfiguring existing ones."""
-        cur_data_parallel_size = len(self.core_engines)
-
-        self.eep_scaling_cache = ElasticScalingCache(
-            existing_core_engines=self.core_engines.copy(),
-            num_new_core_engines=new_data_parallel_size - cur_data_parallel_size,
-            pending_notifications=dict(),
+    def _make_reconfig_request(
+        self,
+        new_data_parallel_size: int,
+        rank_type: ReconfigureRankType = ReconfigureRankType.KEEP_CURRENT_RANK,
+    ) -> ReconfigureDistributedRequest:
+        parallel_config = self.vllm_config.parallel_config
+        return ReconfigureDistributedRequest(
+            new_data_parallel_size=new_data_parallel_size,
+            new_data_parallel_rank=rank_type,
+            new_data_parallel_rank_local=ReconfigureRankType.KEEP_CURRENT_RANK,
+            new_data_parallel_master_ip=parallel_config.data_parallel_master_ip,
+            new_data_parallel_master_port=parallel_config.data_parallel_master_port,
+            new_data_parallel_master_port_list=parallel_config._data_parallel_master_port_list,
+            coord_store_port=parallel_config._coord_store_port,
         )
 
-        parallel_config = self.vllm_config.parallel_config
-        ip, coord_store_port = self._setup_elastic_ep_reconfig_bootstrap()
+    async def _prepare_scale_up_elastic_ep(
+        self,
+        new_data_parallel_size: int,
+        num_redundant_experts: int,
+    ) -> None:
+        """Prepare scale up by creating new engine cores and reconfiguring
+        existing ones."""
+        self._setup_elastic_ep_reconfig_bootstrap()
 
         # Phase 1: Send reconfig messages to existing engines
         reconfig_futures = []
         for engine in self.core_engines:
-            reconfig_request = ReconfigureDistributedRequest(
-                new_data_parallel_size=new_data_parallel_size,
-                new_data_parallel_rank=ReconfigureRankType.KEEP_CURRENT_RANK,
-                new_data_parallel_rank_local=ReconfigureRankType.KEEP_CURRENT_RANK,
-                new_data_parallel_master_ip=ip,
-                new_data_parallel_master_port=parallel_config.data_parallel_master_port,
-                new_data_parallel_master_port_list=parallel_config._data_parallel_master_port_list,
-                coord_store_port=coord_store_port,
-            )
+            reconfig_request = self._make_reconfig_request(new_data_parallel_size)
             coro = self._call_utility_async(
                 "reinitialize_distributed", reconfig_request, engine=engine
             )
@@ -1716,51 +1733,54 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
 
         # Phase 2: Create new engines
         assert isinstance(self.resources.engine_manager, CoreEngineActorManager)
-        parallel_config.eplb_config.num_redundant_experts = 0
         start_new_worker_future = asyncio.to_thread(
             self.resources.engine_manager.scale_up_elastic_ep,
             self.vllm_config,
             new_data_parallel_size,
+            num_redundant_experts,
         )
-        wait_future = self._eep_wait_for_setup_switch_complete()
 
         # Phase 3: Wait for new engines to be created
         # and reconfig messages to be received
         await asyncio.gather(start_new_worker_future, *reconfig_futures)
+        ready_keys = [future.result() for future in reconfig_futures]
+        ready_keys.extend(
+            f"eep_ready/{rank}"
+            for rank in range(len(self.core_engines), new_data_parallel_size)
+        )
+        await asyncio.to_thread(self._coord_store.wait, ready_keys)
         logger.info("[Elastic EP] Successfully started new engines")
 
-        # Create new CoreEngine objects for the new engines
-        new_engine_identities = set()
-        for i in range(cur_data_parallel_size, new_data_parallel_size):
-            new_engine = i.to_bytes(2, "little")
-            self.core_engines.append(new_engine)
-            # NOTE(yongji): we don't update lb_engines here,
-            # we let run_engine_stats_update_task to update it.
-            new_engine_identities.add(new_engine)
+    async def _commit_scale_up_elastic_ep(self, new_data_parallel_size: int) -> None:
+        new_core_engines = [
+            rank.to_bytes(2, "little")
+            for rank in range(len(self.core_engines), new_data_parallel_size)
+        ]
 
-        # Wait for ready messages from new engines on the input socket
-        sync_input_socket = zmq.Socket.shadow(self.input_socket)
-        while new_engine_identities:
-            if not sync_input_socket.poll(
-                timeout=VLLM_ENGINE_READY_TIMEOUT_S * 1000  # convert to ms
-            ):
-                raise TimeoutError(
-                    f"Timed out waiting for new engine core processes to "
-                    f"start. Waited "
-                    f"{VLLM_ENGINE_READY_TIMEOUT_S}s (configured by "
-                    f"VLLM_ENGINE_READY_TIMEOUT_S). To increase the "
-                    f"timeout, set the environment variable: "
-                    f"VLLM_ENGINE_READY_TIMEOUT_S=<seconds>"
-                )
-            identity, payload = sync_input_socket.recv_multipart()
-            new_engine_identities.discard(identity)
-            self._apply_ready_response(payload)
+        await self.pause_scheduler_async(mode="keep", clear_cache=False)
+        wait_future = self._eep_wait_for_setup_switch_complete()
+        finish_futures = [
+            asyncio.create_task(
+                self._call_utility_async("commit_prepared_elastic_ep", engine=engine)
+            )
+            for engine in self.core_engines
+        ]
+        try:
+            await asyncio.gather(*finish_futures)
+            await wait_future
+            self._wait_for_new_engine_ready(new_core_engines)
+        except Exception:
+            wait_future.cancel()
+            raise
 
-        # NOTE(yongji): Before we schedule any requests on the new workers,
-        # we should wait for them to switch to the new setup.
-        await wait_future
+        self.core_engines.extend(new_core_engines)
         # Update the parallel config
-        self.vllm_config.parallel_config.data_parallel_size = new_data_parallel_size
+        parallel_config = self.vllm_config.parallel_config
+        parallel_config.data_parallel_size = new_data_parallel_size
+        if isinstance(self.resources.engine_manager, CoreEngineActorManager):
+            parallel_config.data_parallel_size_local = len(
+                self.resources.engine_manager.local_engine_actors
+            )
         # Notify coordinator about scale up through existing
         # stats_update_task connection
         self._ensure_stats_update_task()
@@ -1773,10 +1793,23 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
             "[Elastic EP] Scale up completed, new data parallel size: %s",
             new_data_parallel_size,
         )
+        await self.resume_scheduler_async()
 
-    async def _scale_down_elastic_ep(
-        self, cur_data_parallel_size: int, new_data_parallel_size: int
-    ) -> None:
+    async def _prepare_scale_down_elastic_ep(self, new_data_parallel_size: int) -> None:
+        self._setup_elastic_ep_reconfig_bootstrap()
+
+        reconfig_futures = []
+        for engine in self.core_engines[:new_data_parallel_size]:
+            reconfig_request = self._make_reconfig_request(new_data_parallel_size)
+            coro = self._call_utility_async(
+                "reinitialize_distributed", reconfig_request, engine=engine
+            )
+            reconfig_futures.append(asyncio.create_task(coro))
+
+        ready_keys = await asyncio.gather(*reconfig_futures)
+        await asyncio.to_thread(self._coord_store.wait, ready_keys)
+
+    async def _commit_scale_down_elastic_ep(self, new_data_parallel_size: int) -> None:
         """Scale down the data parallel size by shutting down and
         reconfiguring existing engine cores."""
         cur_data_parallel_size = len(self.core_engines)
@@ -1787,50 +1820,51 @@ class DPLBAsyncMPClient(DPAsyncMPClient):
             pending_notifications=dict(),
         )
 
-        parallel_config = self.vllm_config.parallel_config
-        ip, coord_store_port = self._setup_elastic_ep_reconfig_bootstrap()
-
+        old_core_engines = self.core_engines
+        # NOTE(yongji): Immediately stop sending requests to the removing engines.
+        self.core_engines = old_core_engines[:new_data_parallel_size]
+        self.lb_engines = self.lb_engines[:new_data_parallel_size]
         removed_dp_size = cur_data_parallel_size - new_data_parallel_size
+        pause_modes = ["keep"] * new_data_parallel_size + ["abort"] * removed_dp_size
+        pause_futures = [
+            self._call_utility_async("pause_scheduler", mode, False, engine=engine)
+            for mode, engine in zip(pause_modes, old_core_engines)
+        ]
+        await asyncio.gather(*pause_futures)
         assert isinstance(self.resources.engine_manager, CoreEngineActorManager)
         self.resources.engine_manager.remove_run_refs_for_scale_down(removed_dp_size)
+        wait_future = self._eep_wait_for_setup_switch_complete()
         reconfig_futures = []
-        for cur_dp_rank, engine in enumerate(self.core_engines):
-            reconfig_request = ReconfigureDistributedRequest(
-                new_data_parallel_size=new_data_parallel_size,
-                new_data_parallel_rank=ReconfigureRankType.KEEP_CURRENT_RANK,
-                new_data_parallel_rank_local=ReconfigureRankType.KEEP_CURRENT_RANK,
-                new_data_parallel_master_ip=ip,
-                new_data_parallel_master_port=parallel_config.data_parallel_master_port,
-                new_data_parallel_master_port_list=parallel_config._data_parallel_master_port_list,
-                coord_store_port=coord_store_port,
-            )
-            if cur_dp_rank >= new_data_parallel_size:
-                reconfig_request.new_data_parallel_rank = (
-                    ReconfigureRankType.SHUTDOWN_CURRENT_RANK
+        for cur_dp_rank, engine in enumerate(old_core_engines):
+            if cur_dp_rank < new_data_parallel_size:
+                coro = self._call_utility_async(
+                    "commit_prepared_elastic_ep", engine=engine
+                )
+            else:
+                reconfig_request = self._make_reconfig_request(
+                    new_data_parallel_size,
+                    ReconfigureRankType.SHUTDOWN_CURRENT_RANK,
+                )
+                coro = self._call_utility_async(
+                    "reinitialize_distributed", reconfig_request, engine=engine
                 )
-            coro = self._call_utility_async(
-                "reinitialize_distributed", reconfig_request, engine=engine
-            )
             reconfig_futures.append(asyncio.create_task(coro))
 
-        # NOTE(yongji): Immediately stop sending requests to the removing engines.
-        self.core_engines = self.core_engines[:new_data_parallel_size]
-        self.lb_engines = self.lb_engines[:new_data_parallel_size]
-        wait_future = self._eep_wait_for_setup_switch_complete()
+        try:
+            await asyncio.gather(*reconfig_futures)
 
-        await asyncio.gather(*reconfig_futures)
+            self.vllm_config.parallel_config.data_parallel_size = new_data_parallel_size
+            self._ensure_stats_update_task()
+            scale_down_marker = msgspec.msgpack.encode(
+                ("SCALE_ELASTIC_EP", new_data_parallel_size)
+            )
+            await self.first_req_send_socket.send(scale_down_marker)
+            await wait_future
+            await self.resume_scheduler_async()
+        except Exception:
+            wait_future.cancel()
+            raise
 
-        self.vllm_config.parallel_config.data_parallel_size = new_data_parallel_size
-        self._ensure_stats_update_task()
-        scale_down_marker = msgspec.msgpack.encode(
-            ("SCALE_ELASTIC_EP", new_data_parallel_size)
-        )
-        await self.first_req_send_socket.send(scale_down_marker)
-
-        # NOTE(yongji): Unlike scaling up,
-        # here we don't actually need to wait for the setup switch to complete.
-        # We may want to remove it in the future.
-        await wait_future
         logger.info(
             "[Elastic EP] Scale down completed, new data parallel size: %s",
             new_data_parallel_size,
diff --git a/vllm/v1/engine/exceptions.py b/vllm/v1/engine/exceptions.py
index d9f79a019e2..edb0fe4261b 100644
--- a/vllm/v1/engine/exceptions.py
+++ b/vllm/v1/engine/exceptions.py
@@ -1,12 +1,15 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-class EngineGenerateError(Exception):
+from vllm.exceptions import VLLMServerError
+
+
+class EngineGenerateError(VLLMServerError):
     """Raised when a AsyncLLM.generate() fails. Recoverable."""
 
     pass
 
 
-class EngineDeadError(Exception):
+class EngineDeadError(VLLMServerError):
     """Raised when the EngineCore dies. Unrecoverable."""
 
     def __init__(self, *args, suppress_context: bool = False, **kwargs):
diff --git a/vllm/v1/engine/input_processor.py b/vllm/v1/engine/input_processor.py
index 10a6329f5a0..735593672da 100644
--- a/vllm/v1/engine/input_processor.py
+++ b/vllm/v1/engine/input_processor.py
@@ -7,6 +7,7 @@ from typing import Any, Literal
 
 import vllm.envs as envs
 from vllm.config import VllmConfig
+from vllm.exceptions import VLLMValidationError
 from vllm.inputs import (
     EngineInput,
     PromptType,
@@ -90,7 +91,7 @@ class InputProcessor:
                 task for task in supported_tasks if task in GENERATION_TASKS
             ]
             if not supported_generation_tasks:
-                raise ValueError("This model does not support generation")
+                raise VLLMValidationError("This model does not support generation")
 
             params.verify(
                 self.model_config,
@@ -104,13 +105,13 @@ class InputProcessor:
                     self.vllm_config.reasoning_config is None
                     or not self.vllm_config.reasoning_config.enabled
                 ):
-                    raise ValueError(
+                    raise VLLMValidationError(
                         "thinking_token_budget is set but reasoning_config is "
                         "not configured. Please set --reasoning-parser "
                         "and/or --reasoning-config to use thinking_token_budget."
                     )
                 if self.use_v2_model_runner:
-                    raise ValueError(
+                    raise VLLMValidationError(
                         "thinking_token_budget is not yet supported by the V2 "
                         "model runner. Run vLLM with VLLM_USE_V2_MODEL_RUNNER=0 "
                         "to use thinking_token_budget."
@@ -120,7 +121,7 @@ class InputProcessor:
                 task for task in supported_tasks if task in POOLING_TASKS
             ]
             if not supported_pooling_tasks:
-                raise ValueError("This model does not support pooling")
+                raise VLLMValidationError("This model does not support pooling")
 
             if params.task is None:
                 if "token_embed" in supported_pooling_tasks:
@@ -131,7 +132,7 @@ class InputProcessor:
                     params.task = "plugin"
 
             if params.task not in supported_pooling_tasks:
-                raise ValueError(
+                raise VLLMValidationError(
                     f"Unsupported task: {params.task!r} "
                     f"Supported tasks: {supported_pooling_tasks}"
                 )
@@ -149,7 +150,7 @@ class InputProcessor:
 
         # LoRA request passed in while LoRA is not enabled
         if not self.lora_config:
-            raise ValueError(
+            raise VLLMValidationError(
                 f"Got lora_request {lora_request} but LoRA is not enabled!"
             )
 
@@ -261,7 +262,7 @@ class InputProcessor:
         dp_local_size = parallel_config.data_parallel_size_local
         num_ranks = dp_local_size if parallel_config.local_engines_only else dp_size
         if data_parallel_rank is not None and not (0 <= data_parallel_rank < num_ranks):
-            raise ValueError(
+            raise VLLMValidationError(
                 f"data_parallel_rank {data_parallel_rank} "
                 f"is out of range [0, {num_ranks})."
             )
@@ -393,7 +394,7 @@ class InputProcessor:
             return
 
         if prompt_len == 0 and prompt_type == "decoder":
-            raise ValueError(f"The {prompt_type} prompt cannot be empty")
+            raise VLLMValidationError(f"The {prompt_type} prompt cannot be empty")
 
         model_config = self.model_config
         max_prompt_len = (
@@ -415,7 +416,7 @@ class InputProcessor:
                     "number of text tokens."
                 )
 
-            raise ValueError(
+            raise VLLMValidationError(
                 f"The {prompt_type} prompt (length {prompt_len}) is "
                 f"longer than the maximum model length of {max_prompt_len}. "
                 f"{suggestion}"
@@ -425,7 +426,7 @@ class InputProcessor:
                 "Make sure that `max_model_len` is no smaller than the "
                 "number of text tokens (prompt + requested output tokens)."
             )
-            raise ValueError(
+            raise VLLMValidationError(
                 f"The {prompt_type} prompt (length {prompt_len}) plus the number of "
                 f"requested output tokens (at least 1) is longer than the maximum "
                 f"model length of {max_prompt_len}. {suggestion}"
@@ -457,7 +458,7 @@ class InputProcessor:
                 for mm_position in mm_positions:
                     num_embeds = mm_position.get_num_embeds()
                     if num_embeds > self.mm_encoder_cache_size:
-                        raise ValueError(
+                        raise VLLMValidationError(
                             f"The {prompt_type} prompt contains a(n) {modality} item "
                             f"with {num_embeds} embedding tokens, which exceeds the "
                             f"pre-allocated encoder cache size "
@@ -481,7 +482,9 @@ class InputProcessor:
             # truly out-of-vocabulary.
             model_vocab_size = model_config.get_vocab_size()
             if max_input_id > max(tokenizer.max_token_id, model_vocab_size - 1):
-                raise ValueError(f"Token id {max_input_id} is out of vocabulary")
+                raise VLLMValidationError(
+                    f"Token id {max_input_id} is out of vocabulary"
+                )
 
     def _validate_model_inputs(
         self,
diff --git a/vllm/v1/engine/llm_engine.py b/vllm/v1/engine/llm_engine.py
index ff86a1dffd9..17e40630859 100644
--- a/vllm/v1/engine/llm_engine.py
+++ b/vllm/v1/engine/llm_engine.py
@@ -425,6 +425,13 @@ class LLMEngine:
     ) -> list[_R]:
         return self.engine_core.collective_rpc(method, timeout, args, kwargs)
 
+    def set_weight_version(self, weight_version: str) -> None:
+        self.engine_core.set_weight_version(weight_version)
+
+    def get_weight_version(self) -> str:
+        """Return the latest committed weight version."""
+        return self.engine_core.get_weight_version()
+
     def apply_model(self, func: Callable[[nn.Module], _R]) -> list[_R]:
         return self.collective_rpc("apply_model", args=(func,))
 
diff --git a/vllm/v1/engine/utils.py b/vllm/v1/engine/utils.py
index db1896b0946..9b3bea0db9c 100644
--- a/vllm/v1/engine/utils.py
+++ b/vllm/v1/engine/utils.py
@@ -821,7 +821,10 @@ class CoreEngineActorManager:
         return placement_groups, local_dp_ranks
 
     def scale_up_elastic_ep(
-        self, cur_vllm_config: VllmConfig, new_data_parallel_size: int
+        self,
+        cur_vllm_config: VllmConfig,
+        new_data_parallel_size: int,
+        num_redundant_experts: int,
     ) -> None:
         import copy
 
@@ -864,6 +867,9 @@ class CoreEngineActorManager:
             if new_data_parallel_size > 1:
                 _apply_dp_identity_suffix(dp_vllm_config, rank)
             dp_vllm_config.parallel_config.data_parallel_size = new_data_parallel_size
+            dp_vllm_config.parallel_config.eplb_config.num_redundant_experts = (
+                num_redundant_experts
+            )
             dp_vllm_config.parallel_config.placement_group = pg
 
             # Check if this placement group is on the head node
@@ -906,39 +912,18 @@ class CoreEngineActorManager:
             self.created_placement_groups.append(pg)
             self.placement_group_is_local.append(local_client)
 
-        ray.get(
-            [
-                actor.wait_for_init.remote()
-                for actor in (
-                    self.local_engine_actors[-new_local_engines:]
-                    if new_local_engines > 0
-                    else []
-                )
-                + self.remote_engine_actors[
-                    -(len(placement_groups) - new_local_engines) :
-                ]
-            ]
-        )
-
         actors = (
             self.local_engine_actors[-new_local_engines:]
             if new_local_engines > 0
             else []
         ) + self.remote_engine_actors[-(len(placement_groups) - new_local_engines) :]
 
+        ray.get([actor.wait_for_init.remote() for actor in actors])
         for actor in actors:
             ref = actor.run.remote()
             self.run_refs.append(ref)
             self.actor_run_ref_dict[actor] = ref
 
-        cur_vllm_config.parallel_config.data_parallel_size = new_data_parallel_size
-        # Update old_vllm_config with new data_parallel_size_local if any new
-        # local engines were added
-        if new_local_engines > 0:
-            cur_vllm_config.parallel_config.data_parallel_size_local += (
-                new_local_engines
-            )
-
     def scale_down_elastic_ep(
         self, cur_data_parallel_size: int, new_data_parallel_size: int
     ) -> None:
diff --git a/vllm/v1/executor/abstract.py b/vllm/v1/executor/abstract.py
index 4063844d469..404acd50de9 100644
--- a/vllm/v1/executor/abstract.py
+++ b/vllm/v1/executor/abstract.py
@@ -116,11 +116,11 @@ class Executor(ABC):
         raise NotImplementedError
 
     def initialize_from_config(self, kv_cache_configs: list[KVCacheConfig]) -> None:
-        """
-        Initialize the KV caches and begin the model execution loop of the
-        underlying workers.
-        """
+        """Initialize the KV caches on the underlying workers."""
         self.collective_rpc("initialize_from_config", args=(kv_cache_configs,))
+
+    def compile_or_warm_up_model(self) -> None:
+        """Compile/warm up the model and capture cudagraphs on workers."""
         compilation_times: list[CompilationTimes] = self.collective_rpc(
             "compile_or_warm_up_model"
         )
diff --git a/vllm/v1/kv_offload/base.py b/vllm/v1/kv_offload/base.py
index 6ff09f23d37..b0db86cf9bf 100644
--- a/vllm/v1/kv_offload/base.py
+++ b/vllm/v1/kv_offload/base.py
@@ -13,8 +13,6 @@ from typing import TYPE_CHECKING, Any, ClassVar, NamedTuple, NewType, TypeVar
 import numpy as np
 import torch
 
-from vllm.logger import init_logger
-
 if TYPE_CHECKING:
     from vllm.distributed.kv_transfer.kv_connector.v1.offloading.metrics import (
         OffloadingConnectorStats,
@@ -27,8 +25,6 @@ from vllm.v1.kv_offload.config import OffloadingConfig
 # Use the helper functions below to construct / decompose keys.
 OffloadKey = NewType("OffloadKey", bytes)
 
-logger = init_logger(__name__)
-
 
 def make_offload_key(block_hash: bytes, group_idx: int) -> OffloadKey:
     """Pack a block hash and group index into an `OffloadKey`."""
@@ -540,10 +536,6 @@ class OffloadingSpec(ABC):
         return {}
 
     def __init__(self, config: OffloadingConfig):
-        logger.warning(
-            "Initializing OffloadingSpec. This API is experimental and "
-            "subject to change in the future as we iterate the design."
-        )
         self.config = config
         self.extra_config = config.extra_config
         self.replicated_layout: bool = False
diff --git a/vllm/v1/kv_offload/cpu/manager.py b/vllm/v1/kv_offload/cpu/manager.py
index c2ec4170b8e..0eef5cf7e79 100644
--- a/vllm/v1/kv_offload/cpu/manager.py
+++ b/vllm/v1/kv_offload/cpu/manager.py
@@ -2,7 +2,6 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 from collections import OrderedDict
 from collections.abc import Collection, Iterable
-from typing import Literal
 
 from typing_extensions import override
 
@@ -24,19 +23,15 @@ from vllm.v1.kv_offload.cpu.common import (
     CPULoadStoreSpec,
     CPUOffloadingMetrics,
 )
-from vllm.v1.kv_offload.cpu.policies.arc import ARCCachePolicy
 from vllm.v1.kv_offload.cpu.policies.base import BlockStatus, CachePolicy
-from vllm.v1.kv_offload.cpu.policies.lru import LRUCachePolicy
-
-_CACHE_POLICIES: dict[str, type[CachePolicy]] = {
-    "lru": LRUCachePolicy,
-    "arc": ARCCachePolicy,
-}
+from vllm.v1.kv_offload.cpu.policies.factory import CachePolicyFactory
 
 
 class CPUOffloadingManager(OffloadingManager):
     """
-    An OffloadingManager with a pluggable CachePolicy (LRU or ARC).
+    An OffloadingManager with a pluggable CachePolicy, resolved by name via
+    CachePolicyFactory (built in: "lru", "arc"; external policies can either
+    register their own or be loaded out-of-tree via cache_policy_module_path).
 
     The manager owns all shared logic: ref-counting, event emission,
     block pool management, and the prepare_store/complete_store skeletons.
@@ -47,7 +42,8 @@ class CPUOffloadingManager(OffloadingManager):
     def __init__(
         self,
         num_blocks: int,
-        cache_policy: Literal["lru", "arc"] = "lru",
+        cache_policy: str = "lru",
+        cache_policy_module_path: str | None = None,
         enable_events: bool = False,
         store_threshold: int = 1,
         max_tracker_size: int = 64_000,
@@ -57,12 +53,9 @@ class CPUOffloadingManager(OffloadingManager):
         self._num_allocated_blocks: int = 0
         self._free_list: list[int] = []
         self.events: list[OffloadingEvent] | None = [] if enable_events else None
-        policy_cls = _CACHE_POLICIES.get(cache_policy)
-        if policy_cls is None:
-            raise ValueError(
-                f"Unknown cache policy: {cache_policy!r}. "
-                f"Supported: {list(_CACHE_POLICIES)}"
-            )
+        policy_cls = CachePolicyFactory.get_cache_policy_cls(
+            cache_policy, cache_policy_module_path
+        )
         self._policy: CachePolicy = policy_cls(cache_capacity=num_blocks)
         # Track the number of blocks in the cache that are evictable. i.e. ref_cnt 0.
         self._num_evictable_cache_blocks: int = 0
diff --git a/vllm/v1/kv_offload/cpu/policies/arc.py b/vllm/v1/kv_offload/cpu/policies/arc.py
index f682a47e45f..d6569cbcfd2 100644
--- a/vllm/v1/kv_offload/cpu/policies/arc.py
+++ b/vllm/v1/kv_offload/cpu/policies/arc.py
@@ -48,7 +48,7 @@ class ARCCachePolicy(CachePolicy):
     """
 
     def __init__(self, cache_capacity: int):
-        self.cache_capacity: int = cache_capacity
+        super().__init__(cache_capacity)
         self.target_t1_size: float = 0.0
         self.t1: OrderedDict[OffloadKey, BlockStatus] = OrderedDict()
         self.t2: OrderedDict[OffloadKey, BlockStatus] = OrderedDict()
diff --git a/vllm/v1/kv_offload/cpu/policies/base.py b/vllm/v1/kv_offload/cpu/policies/base.py
index 2b6681e4992..908907d326f 100644
--- a/vllm/v1/kv_offload/cpu/policies/base.py
+++ b/vllm/v1/kv_offload/cpu/policies/base.py
@@ -41,8 +41,8 @@ class CachePolicy(ABC):
     and eviction, so they cannot be separated cleanly.
     """
 
-    @abstractmethod
-    def __init__(self, cache_capacity: int) -> None: ...
+    def __init__(self, cache_capacity: int) -> None:
+        self.cache_capacity = cache_capacity
 
     @abstractmethod
     def get(self, key: OffloadKey) -> BlockStatus | None:
diff --git a/vllm/v1/kv_offload/cpu/policies/factory.py b/vllm/v1/kv_offload/cpu/policies/factory.py
new file mode 100644
index 00000000000..88d11b270c9
--- /dev/null
+++ b/vllm/v1/kv_offload/cpu/policies/factory.py
@@ -0,0 +1,87 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import importlib
+from collections.abc import Callable
+
+from vllm.logger import init_logger
+from vllm.v1.kv_offload.cpu.policies.base import CachePolicy
+
+logger = init_logger(__name__)
+
+
+class CachePolicyFactory:
+    """Registry for CachePolicy implementations, resolved by name.
+
+    Mirrors OffloadingSpecFactory (vllm/v1/kv_offload/factory.py): built-in
+    policies are pre-registered below. External policies can either
+    register_cache_policy() a friendly short name up front, or skip
+    registration entirely and pass a module path at lookup time (out-of-tree,
+    no vLLM fork/patch required) -- see get_cache_policy_cls.
+    """
+
+    _registry: dict[str, Callable[[], type[CachePolicy]]] = {}
+
+    @classmethod
+    def register_cache_policy(
+        cls, name: str, module_path: str, class_name: str
+    ) -> None:
+        """Register a cache policy with a lazy-loading module and class name."""
+        if name in cls._registry:
+            raise ValueError(f"Cache policy '{name}' is already registered.")
+
+        def loader() -> type[CachePolicy]:
+            module = importlib.import_module(module_path)
+            return getattr(module, class_name)
+
+        cls._registry[name] = loader
+
+    @classmethod
+    def get_cache_policy_cls(
+        cls, name: str, module_path: str | None = None
+    ) -> type[CachePolicy]:
+        """Get a cache policy class by name.
+
+        Args:
+            name: Name of the cache policy. Checked against the registry
+                first; if it's not registered and `module_path` is given,
+                `name` is imported from there instead -- an out-of-tree
+                policy needs no register_cache_policy() call at all, just
+                this module path passed through config (mirrors
+                OffloadingSpecFactory.get_spec_cls's spec_module_path
+                fallback).
+            module_path: Python import path to load `name` from when it is
+                not a registered policy.
+
+        Returns:
+            The cache policy class.
+
+        Raises ValueError if the cache policy is neither registered nor
+        resolvable via `module_path`.
+        """
+        if name in cls._registry:
+            return cls._registry[name]()
+        if module_path is None:
+            raise ValueError(
+                f"Unknown cache policy: {name!r}. Supported: {list(cls._registry)}. "
+                "For an out-of-tree policy, also set cache_policy_module_path."
+            )
+        logger.warning(
+            "Loading out-of-tree cache policy '%s' from '%s'. This API is "
+            "experimental and subject to change in the future as we "
+            "iterate the design.",
+            name,
+            module_path,
+        )
+        module = importlib.import_module(module_path)
+        policy_cls = getattr(module, name)
+        assert issubclass(policy_cls, CachePolicy)
+        return policy_cls
+
+
+# Register built-in policies here.
+CachePolicyFactory.register_cache_policy(
+    "lru", "vllm.v1.kv_offload.cpu.policies.lru", "LRUCachePolicy"
+)
+CachePolicyFactory.register_cache_policy(
+    "arc", "vllm.v1.kv_offload.cpu.policies.arc", "ARCCachePolicy"
+)
diff --git a/vllm/v1/kv_offload/cpu/policies/lru.py b/vllm/v1/kv_offload/cpu/policies/lru.py
index efa24fe9033..e8ccf0bdef5 100644
--- a/vllm/v1/kv_offload/cpu/policies/lru.py
+++ b/vllm/v1/kv_offload/cpu/policies/lru.py
@@ -19,6 +19,7 @@ class LRUCachePolicy(CachePolicy):
     """
 
     def __init__(self, cache_capacity: int):
+        super().__init__(cache_capacity)
         # Blocks with ref_cnt 0 (not participating in any loads/stores) ordered in LRU
         self.evictable_blocks: OrderedDict[OffloadKey, None] = OrderedDict()
         self.blocks: dict[OffloadKey, BlockStatus] = {}
diff --git a/vllm/v1/kv_offload/cpu/spec.py b/vllm/v1/kv_offload/cpu/spec.py
index d755bfecbc4..9162c7b18b4 100644
--- a/vllm/v1/kv_offload/cpu/spec.py
+++ b/vllm/v1/kv_offload/cpu/spec.py
@@ -117,6 +117,9 @@ class CPUOffloadingSpec(OffloadingSpec):
         self._worker: CPUOffloadingWorker | None = None
 
         self.eviction_policy: str = self.extra_config.get("eviction_policy", "lru")
+        self.cache_policy_module_path: str | None = self.extra_config.get(
+            "cache_policy_module_path"
+        )
 
     @override
     def get_manager(self) -> OffloadingManager:
@@ -131,7 +134,8 @@ class CPUOffloadingSpec(OffloadingSpec):
 
             self._manager = CPUOffloadingManager(
                 num_blocks=self.num_blocks,
-                cache_policy=self.eviction_policy,  # type: ignore[arg-type]
+                cache_policy=self.eviction_policy,
+                cache_policy_module_path=self.cache_policy_module_path,
                 enable_events=self.kv_events_config.enable_kv_cache_events,
                 store_threshold=store_threshold,
                 max_tracker_size=max_tracker_size,
diff --git a/vllm/v1/kv_offload/factory.py b/vllm/v1/kv_offload/factory.py
index 931fda8308f..19bc401277d 100644
--- a/vllm/v1/kv_offload/factory.py
+++ b/vllm/v1/kv_offload/factory.py
@@ -35,6 +35,13 @@ class OffloadingSpecFactory:
             spec_module_path = extra_config.get("spec_module_path")
             if spec_module_path is None:
                 raise ValueError(f"Unsupported spec type: {spec_name}")
+            logger.warning(
+                "Loading out-of-tree offloading spec '%s' from '%s'. This "
+                "API is experimental and subject to change in the future "
+                "as we iterate the design.",
+                spec_name,
+                spec_module_path,
+            )
             spec_module = importlib.import_module(spec_module_path)
             spec_cls = getattr(spec_module, spec_name)
         assert issubclass(spec_cls, OffloadingSpec)
diff --git a/vllm/v1/kv_offload/file_mapper.py b/vllm/v1/kv_offload/file_mapper.py
index 8e8c19d53d6..4b12dba913d 100644
--- a/vllm/v1/kv_offload/file_mapper.py
+++ b/vllm/v1/kv_offload/file_mapper.py
@@ -35,6 +35,7 @@ class FileMapper:
         kv_cache_groups: list[dict] | None = None,
         inference_engine: str = "vllm",
         parallel_agnostic: bool = False,
+        replicated_layout: bool = False,
     ):
         """
         Initialize the file mapper. Each worker constructs its own, but
@@ -60,6 +61,10 @@ class FileMapper:
         }
         if not parallel_agnostic:
             self.fields["parallel_agnostic"] = False
+        # Only written when True so existing deployments' hashed fields are
+        # unchanged (False is the historical default and must not appear).
+        if replicated_layout:
+            self.fields["replicated_layout"] = True
         self.base_path: str = self._compute_base_path(root_dir, self.fields)
 
     @classmethod
@@ -92,7 +97,11 @@ class FileMapper:
             rank=parallel.rank,
             dtype=config.model.dtype,
             kv_cache_groups=kv_cache_groups,
-            parallel_agnostic=(parallel_agnostic and parallel.is_parallelism_agnostic),
+            parallel_agnostic=(
+                parallel_agnostic
+                and (parallel.is_parallelism_agnostic or config.replicated_layout)
+            ),
+            replicated_layout=(parallel_agnostic and config.replicated_layout),
         )
 
     def get_file_name(self, key: OffloadKey) -> str:
diff --git a/vllm/v1/kv_offload/tiering/manager.py b/vllm/v1/kv_offload/tiering/manager.py
index c6738096135..d1d3c421159 100644
--- a/vllm/v1/kv_offload/tiering/manager.py
+++ b/vllm/v1/kv_offload/tiering/manager.py
@@ -93,11 +93,13 @@ class CPUPrimaryTierOffloadingManager(CPUOffloadingManager):
         num_blocks: int,
         mmap_region: SharedOffloadRegion,
         cache_policy: str = "lru",
+        cache_policy_module_path: str | None = None,
         enable_events: bool = False,
     ):
         super().__init__(
             num_blocks=num_blocks,
-            cache_policy=cache_policy,  # type: ignore[arg-type]
+            cache_policy=cache_policy,
+            cache_policy_module_path=cache_policy_module_path,
             enable_events=enable_events,
         )
         self._mmap_region = mmap_region
diff --git a/vllm/v1/kv_offload/tiering/obj/manager.py b/vllm/v1/kv_offload/tiering/obj/manager.py
index c7e0d4c4beb..2dfc2d30fa2 100644
--- a/vllm/v1/kv_offload/tiering/obj/manager.py
+++ b/vllm/v1/kv_offload/tiering/obj/manager.py
@@ -307,15 +307,36 @@ class ObjectStoreSecondaryTierManager(SecondaryTierManager):
             else:
                 if state == NIXL_PROC:
                     continue
-                elif state == NIXL_DONE:
+                if state == NIXL_DONE:
                     success = True
                 else:
                     success = False
                     logger.warning("transfer failed job=%d state=%s", job_id, state)
+
+            try:
+                self._agent.release_xfer_handle(entry.xfer_handle)
+            except Exception as exc:
+                # Keep the entry until NIXL confirms that the transfer handle
+                # can be released. The transfer may still access primary-tier
+                # memory, so publishing its result would allow unsafe reuse.
+                logger.warning("release_xfer_handle failed for job %d: %s", job_id, exc)
+                continue
+
+            # Once the transfer handle is released, these remaining cleanup
+            # failures must not suppress the job completion. They can leak
+            # NIXL metadata, but cannot leave an active data transfer behind.
+            try:
+                self._agent.release_dlist_handle(entry.obj_handle)
+            except Exception as exc:
+                logger.warning(
+                    "release_dlist_handle failed for job %d: %s", job_id, exc
+                )
+            try:
+                self._agent.deregister_memory(entry.files_desc)
+            except Exception as exc:
+                logger.warning("deregister_memory failed for job %d: %s", job_id, exc)
+
             del self._transfers[job_id]
-            self._agent.release_xfer_handle(entry.xfer_handle)
-            self._agent.release_dlist_handle(entry.obj_handle)
-            self._agent.deregister_memory(entry.files_desc)
             self._pending_results.append(JobResult(job_id=job_id, success=success))
 
     def get_finished_jobs(self) -> Iterable[JobResult]:
diff --git a/vllm/v1/kv_offload/tiering/p2p/session/client.py b/vllm/v1/kv_offload/tiering/p2p/session/client.py
index ed91f7f9b00..07a4620780e 100644
--- a/vllm/v1/kv_offload/tiering/p2p/session/client.py
+++ b/vllm/v1/kv_offload/tiering/p2p/session/client.py
@@ -11,7 +11,6 @@ callback injected by the coordinator (which gates on ConnectAck).
 
 from __future__ import annotations
 
-import enum
 import time
 from collections.abc import Callable, Sequence
 from dataclasses import dataclass, field
@@ -35,26 +34,12 @@ _LOAD_TIMEOUT_S = 30.0
 _ABORT_ACK_TIMEOUT_S = 10.0
 
 
-class ClientPhase(enum.Enum):
-    """Lifecycle of a request's client-side lookup/fetch signalling.
-
-    Advances monotonically. Only ``finish`` reads it, to decide
-    whether a terminal empty FetchMsg is owed to release the peer's
-    lookup state (owed only from ``PROBING``: a LookupMsg went out but no
-    FetchMsg has since closed the peer's lookup phase).
-    """
-
-    REGISTERED = enum.auto()  # keys in probes/unsent, nothing sent yet
-    PROBING = enum.auto()  # LookupMsg flushed, awaiting responses
-    FETCH_SENT = enum.auto()  # FetchMsg sent (real or terminal empty)
-
-
 @dataclass
 class _InboundLoadState:
     """Client-role state for a single in-flight load request.
 
-    Lives on ``_ClientRequestState.load`` for the duration of a fetch;
-    the owning kv_request_id is the dict key, so it isn't stored here.
+    Lives in ``_ClientRequestState.loads`` keyed by round_seq for the
+    duration of a fetch; the owning kv_request_id is the outer dict key.
     """
 
     job_id: int  # opaque ID assigned by the manager to this load request
@@ -68,7 +53,7 @@ class _ClientRequestState:
 
     One entry per kv_request_id we're driving. Lookup-phase fields are
     used only by symmetric P2P (``do_p2p_fetch``); PD-only loads leave
-    ``probes``/``unsent`` empty and drive just ``phase`` and ``load``. An
+    ``probes``/``unsent`` empty and drive just ``phase`` and ``loads``. An
     entry is dropped once every field is idle — see ``ClientRole._maybe_prune``.
     """
 
@@ -82,11 +67,23 @@ class _ClientRequestState:
     # OffloadKeys registered but not yet flushed onto the wire. Drained and
     # cleared by the next flush_pending_lookups.
     unsent: list[OffloadKey] = field(default_factory=list)
+    # Current lookup round. LookupMsgs carry it, each fetch closes it and
+    # advances it, so every round's supply/demand/completion is isolated
+    # on the wire. PD clients never probe and stay on round 0.
+    round_seq: int = 0
+    # This id ran the symmetric lookup phase (register_lookup); a fetch
+    # with keys then requires every key to be a confirmed probe. PD
+    # loads never probe.
+    probed: bool = False
 
-    # Monotonic lookup/fetch signalling phase; see ``ClientPhase``.
-    phase: ClientPhase = ClientPhase.REGISTERED
-    # Set while a fetch is in flight; cleared on completion/abort/timeout.
-    load: _InboundLoadState | None = None
+    # The peer holds lookup state no FetchMsg has closed: a LookupMsg
+    # was flushed since the last fetch. finish owes a terminal empty
+    # FetchMsg while set, so the peer releases parked supply.
+    peer_lookup_open: bool = False
+    # In-flight loads keyed by the round their fetch carried. The
+    # scheduler submits loads incrementally as chunks resolve, so several
+    # can be in flight at once; TransferDone/AbortAck match by round.
+    loads: dict[int, _InboundLoadState] = field(default_factory=dict)
 
 
 class LoadResult(NamedTuple):
@@ -133,11 +130,10 @@ class ClientRole:
         # _serve_pending. Populated by register_lookup, drained by
         # flush_pending_lookups, and discarded on finish/close.
         self._flush_pending: set[str] = set()
-        # kv_request_ids with a fetch in flight (``st.load is not None``) —
-        # the work-list collect_results walks for timeouts, and the
-        # has_active_loads predicate, instead of scanning every request.
-        # Kept in exact sync with ``st.load``: armed in request_blocks,
-        # discarded wherever load is cleared, and cleared on close.
+        # kv_request_ids with at least one fetch in flight — the work-list
+        # collect_results walks for timeouts, and the has_active_loads
+        # predicate, instead of scanning every request. Kept in exact sync
+        # with ``st.loads``.
         self._active_loads: set[str] = set()
         self._completed_loads: list[LoadResult] = []
 
@@ -156,17 +152,21 @@ class ClientRole:
     def _maybe_prune(self, kv_request_id: str) -> None:
         """Drop the entry once it holds no live load or lookup state.
 
-        The sticky ``phase`` is only read by ``finish``. A probe
-        clears when its fetch is issued (``request_blocks``) or when the
-        request finishes (``finish``/``close``); in the former case
-        ``load`` is set and keeps the entry alive, in the latter the phase
-        is no longer needed — so dropping on emptiness never loses a phase
-        still in use.
+        ``peer_lookup_open`` is only read by ``finish``, and every path
+        that clears the last probe (fetch / finish / close) also settles
+        it, so dropping on emptiness never loses a flag still in use.
         """
         st = self._requests.get(kv_request_id)
-        if st is not None and st.load is None and not st.probes and not st.unsent:
+        if st is not None and not st.loads and not st.probes and not st.unsent:
             del self._requests[kv_request_id]
 
+    def _on_load_terminal(self, kv_request_id: str, st: _ClientRequestState) -> None:
+        """Wind down id-level state once no load remains in flight."""
+        if st.loads:
+            return
+        self._active_loads.discard(kv_request_id)
+        self._maybe_prune(kv_request_id)
+
     @property
     def has_active_loads(self) -> bool:
         """True if any kv_request_id has a fetch in flight."""
@@ -184,7 +184,12 @@ class ClientRole:
         block_ids: Sequence[int],
         send_ready: bool,
     ) -> None:
-        """Register a load request and send the FetchMsg."""
+        """Send the FetchMsg closing the current lookup round.
+
+        The scheduler may submit several loads per kv_request_id as its
+        matched prefix resolves incrementally; each fetch carries the
+        round it closes so the loads stay independent on the wire.
+        """
         logger.debug(
             "P2PSession %s: request_blocks job_id=%d kv_request_id=%s "
             "blocks=%d ready=%s",
@@ -195,77 +200,74 @@ class ClientRole:
             send_ready,
         )
         st = self._get_or_create_request(kv_request_id)
-        st.load = _InboundLoadState(
+        round_seq = st.round_seq
+        st.round_seq += 1
+        st.loads[round_seq] = _InboundLoadState(
             job_id=job_id,
             submitted_at=time.monotonic(),
         )
         self._active_loads.add(kv_request_id)
-        st.phase = ClientPhase.FETCH_SENT
+        st.peer_lookup_open = False
         self._send(
             {
                 TYPE_KEY: FetchMsg.TYPE,
                 FetchMsg.KV_REQUEST_ID: kv_request_id,
                 FetchMsg.KEYS: list(keys),
                 FetchMsg.BLOCK_INDEXES: [int(idx) for idx in block_ids],
+                FetchMsg.ROUND_SEQ: round_seq,
             }
         )
-        # Issuing the fetch ends this request's lookup phase, so drop all
-        # probe state. Once the peer serves this fetch both sides unpin, so
-        # the producer may evict the block; a stale cached True would
-        # otherwise let a re-scheduled lookup() return HIT without
-        # re-probing, pointing at a block the producer no longer holds.
-        # Clearing forces a fresh LookupMsg on re-schedule so the producer
-        # answers from current state. For a symmetric-P2P request (probes
-        # populated) every fetched block was a confirmed HIT; a PD-only load
-        # never probes, so probes is empty and the clear is a no-op.
-        if st.probes:
+        # Issuing the fetch closes this lookup round, so drop all probe
+        # state. Once the peer serves the fetch both sides unpin, so a
+        # stale cached True would let a later lookup() return HIT for a
+        # block the producer may have evicted; clearing forces a fresh
+        # probe under the next round.
+        if st.probed and keys:
+            assert st.probes, (
+                f"symmetric fetch for {kv_request_id} has keys but no probes"
+            )
             assert all(st.probes.get(key) is True for key in keys)
         st.probes.clear()
 
     def finish(self, kv_request_id: str) -> None:
-        """Finish a request: abort any in-flight load and release lookup state.
+        """Finish a request: abort in-flight loads and release lookup state.
 
-        Called from the session's ``finish_request``. The two branches are
-        mutually exclusive: ``load`` is set only by ``request_blocks``, which
-        also advances ``phase`` to ``FETCH_SENT``, and nothing moves it back
-        to ``PROBING`` — so a fetch in flight never coexists with the
-        ``PROBING`` phase.
-
-        - Fetch in flight (``load`` set, phase ``FETCH_SENT``): send an
-          AbortFetchMsg unless the load is already aborting, then drop it.
-        - Outstanding lookups (``PROBING``): a LookupMsg was flushed but no
-          FetchMsg has closed the peer's lookup phase. Every FetchMsg the
-          server receives in p2p mode is its "request finished" signal (it
-          releases lookup state and fires ``cb.finish_request``); when the
-          client's lookups all missed no FetchMsg is otherwise sent, so emit
-          a terminal empty one purely to trigger those semantics. In
-          ``REGISTERED`` the peer never received a LookupMsg and in
-          ``FETCH_SENT`` a FetchMsg already closed the phase, so neither owes
-          a terminal FetchMsg.
+        Called from the session's ``finish_request``. Sends an
+        AbortFetchMsg per load not already aborting, and — independently —
+        the terminal empty FetchMsg when the peer still holds lookup
+        state no fetch has closed (its "request finished" signal: it
+        releases lookup state, drains parked supply, and fires
+        ``cb.finish_request``). A later round's supply can be parked
+        while an earlier round's load is still in flight, so both can be
+        owed at once.
 
         Then drop all probe/lookup state and prune the entry.
         """
         st = self._requests.get(kv_request_id)
         if st is None:
             return
-        if st.load is not None:
-            if st.load.aborted_at is None:
+        if st.loads:
+            for round_seq, load in st.loads.items():
+                if load.aborted_at is not None:
+                    continue
                 self._send(
                     {
                         TYPE_KEY: AbortFetchMsg.TYPE,
                         AbortFetchMsg.KV_REQUEST_ID: kv_request_id,
+                        AbortFetchMsg.ROUND_SEQ: round_seq,
                     }
                 )
-            st.load = None
+            st.loads.clear()
             self._active_loads.discard(kv_request_id)
-        elif st.phase is ClientPhase.PROBING:
-            st.phase = ClientPhase.FETCH_SENT
+        if st.peer_lookup_open:
+            st.peer_lookup_open = False
             self._send(
                 {
                     TYPE_KEY: FetchMsg.TYPE,
                     FetchMsg.KV_REQUEST_ID: kv_request_id,
                     FetchMsg.KEYS: [],
                     FetchMsg.BLOCK_INDEXES: [],
+                    FetchMsg.ROUND_SEQ: st.round_seq,
                 }
             )
         st.probes.clear()
@@ -273,20 +275,21 @@ class ClientRole:
         self._flush_pending.discard(kv_request_id)
         self._maybe_prune(kv_request_id)
 
-    def on_transfer_done(self, kv_request_id: str, success: bool) -> None:
+    def on_transfer_done(
+        self, kv_request_id: str, success: bool, round_seq: int
+    ) -> None:
         """Handle a TransferDoneMsg from the peer."""
         st = self._requests.get(kv_request_id)
-        if st is not None and st.load is not None:
+        load = st.loads.pop(round_seq, None) if st is not None else None
+        if st is not None and load is not None:
             self._completed_loads.append(
                 LoadResult(
-                    job_id=st.load.job_id,
+                    job_id=load.job_id,
                     kv_request_id=kv_request_id,
                     success=success,
                 )
             )
-            st.load = None
-            self._active_loads.discard(kv_request_id)
-            self._maybe_prune(kv_request_id)
+            self._on_load_terminal(kv_request_id, st)
         else:
             # No matching in-flight load: either a duplicate
             # transfer_done from the peer (protocol violation) or a
@@ -295,41 +298,44 @@ class ClientRole:
             # so we can't tell — log so it's findable.
             logger.warning(
                 "P2PSession %s: transfer_done for unknown kv_request_id=%s "
-                "(duplicate from peer, or raced with local cancel/timeout)",
+                "round=%s (duplicate from peer, or raced with local "
+                "cancel/timeout)",
                 self._peer_id,
                 kv_request_id,
+                round_seq,
             )
 
-    def on_abort_ack(self, kv_request_id: str) -> None:
+    def on_abort_ack(self, kv_request_id: str, round_seq: int) -> None:
         """Handle an AbortAckMsg from the peer."""
         st = self._requests.get(kv_request_id)
-        if st is not None and st.load is not None:
+        load = st.loads.pop(round_seq, None) if st is not None else None
+        if st is not None and load is not None:
             logger.warning(
                 "P2PSession %s: load request %s (job_id=%d) timed out; "
                 "load job completed with failure. If this recurs, ensure "
                 "PYTHONHASHSEED is set to the same value on all nodes.",
                 self._peer_id,
                 kv_request_id,
-                st.load.job_id,
+                load.job_id,
             )
             self._completed_loads.append(
                 LoadResult(
-                    job_id=st.load.job_id,
+                    job_id=load.job_id,
                     kv_request_id=kv_request_id,
                     success=False,
                 )
             )
-            st.load = None
-            self._active_loads.discard(kv_request_id)
-            self._maybe_prune(kv_request_id)
+            self._on_load_terminal(kv_request_id, st)
         else:
             # See on_transfer_done: same ambiguity (duplicate ack
             # vs. raced with local cancel/timeout that already popped).
             logger.warning(
                 "P2PSession %s: abort_ack for unknown kv_request_id=%s "
-                "(duplicate from peer, or raced with local cancel/timeout)",
+                "round=%s (duplicate from peer, or raced with local "
+                "cancel/timeout)",
                 self._peer_id,
                 kv_request_id,
+                round_seq,
             )
 
     # ------------------------------------------------------------------
@@ -345,17 +351,18 @@ class ClientRole:
         - Once a LookupRespMsg has resolved the entry: returns the cached
           bool result on every call without popping it.
 
-        A resolved entry is retained until its fetch is issued
-        (``request_blocks`` pops it) or the request finishes
+        A resolved entry is retained until a fetch closes the round
+        (``request_blocks`` clears all probes) or the request finishes
         (``finish`` clears all entries for the id). A request's
         block set can be re-probed across steps, so popping on read would
         make a repeat probe of an already-resolved key look brand-new and
         re-queue it, emitting a redundant LookupMsg for an answer we
         already hold. Keeping the entry until fetch makes repeat probes
-        free; clearing it at fetch forces a fresh probe if the request is
-        re-scheduled, since the block is unpinned once served.
+        free; clearing at fetch forces a fresh probe under the next
+        round, since the block is unpinned once served.
         """
         st = self._get_or_create_request(kv_request_id)
+        st.probed = True
         okey = OffloadKey(key)
         if okey in st.probes:
             return st.probes[okey]
@@ -378,14 +385,11 @@ class ClientRole:
         ``on_schedule_end()``. A request's block set may be discovered
         across several scheduler steps, so more than one LookupMsg can
         go out per kv_request_id — one per step that registered new
-        keys. register_lookup() de-dups in-flight and already-resolved
-        (req_id, key) pairs, so each LookupMsg carries only the keys
-        first probed in that step. The peer's lookup phase for the id is
-        still closed by exactly one FetchMsg, which the client contract
-        guarantees is sent after every lookup for the id has resolved
-        (see request_blocks / finish). Send-gating is handled by
-        the injected ``_send`` callback (queues until ConnectAckMsg if
-        needed).
+        keys, all tagged with the current round. register_lookup()
+        de-dups in-flight and already-resolved (req_id, key) pairs, so
+        each LookupMsg carries only the keys first probed in that step.
+        Send-gating is handled by the injected ``_send`` callback
+        (queues until ConnectAckMsg if needed).
 
         Only requests that registered new keys since the last flush are
         visited — the ``_flush_pending`` work-list avoids scanning every
@@ -395,13 +399,9 @@ class ClientRole:
             st = self._requests.get(req_id)
             if st is None or not st.unsent:
                 continue
-            # Record that the peer now holds lookup state for this id so
-            # finish knows a terminal empty FetchMsg may be owed.
-            # Only promote from REGISTERED: once a fetch has gone out
-            # (FETCH_SENT) a later LookupMsg must not regress the phase, as
-            # no terminal FetchMsg is owed for an already-fetched request.
-            if st.phase is ClientPhase.REGISTERED:
-                st.phase = ClientPhase.PROBING
+            # The peer now holds lookup state for this id; finish owes a
+            # terminal empty FetchMsg until a fetch closes it.
+            st.peer_lookup_open = True
             logger.debug(
                 "P2P LOOKUP client %s: SEND LookupMsg kv_request_id=%s keys=%d",
                 self._peer_id,
@@ -413,6 +413,7 @@ class ClientRole:
                     TYPE_KEY: LookupMsg.TYPE,
                     LookupMsg.KV_REQUEST_ID: req_id,
                     LookupMsg.KEYS: list(st.unsent),
+                    LookupMsg.ROUND_SEQ: st.round_seq,
                 }
             )
             st.unsent = []
@@ -451,52 +452,56 @@ class ClientRole:
     def collect_results(self) -> list[LoadResult]:
         """Walk load timeouts and drain completed loads.
 
-        Active requests past ``_LOAD_TIMEOUT_S`` get an AbortFetchMsg
-        sent and enter the aborting phase. Aborting requests past
-        ``_ABORT_ACK_TIMEOUT_S`` are surfaced as failed loads.
+        Loads past ``_LOAD_TIMEOUT_S`` get an AbortFetchMsg sent and
+        enter the aborting phase. Aborting loads past
+        ``_ABORT_ACK_TIMEOUT_S`` are surfaced as failed.
 
         Lookups have no timeout: an unanswered probe stays None (RETRY)
         until finish_request clears it — see ``_ClientRequestState.probes``.
         """
         now = time.monotonic()
-        to_remove: list[str] = []
+        to_remove: list[tuple[str, int]] = []
         for req_id in self._active_loads:
             st = self._requests[req_id]
-            assert st.load is not None
-            load = st.load
-            if load.aborted_at is None:
-                if now - load.submitted_at >= _LOAD_TIMEOUT_S:
-                    load.aborted_at = now
-                    logger.warning(
-                        "P2PSession %s: %s timed out, sending abort",
-                        self._peer_id,
-                        req_id,
-                    )
-                    self._send(
-                        {
-                            TYPE_KEY: AbortFetchMsg.TYPE,
-                            AbortFetchMsg.KV_REQUEST_ID: req_id,
-                        }
-                    )
-            else:
-                if now - load.aborted_at >= _ABORT_ACK_TIMEOUT_S:
-                    to_remove.append(req_id)
-                    self._completed_loads.append(
-                        LoadResult(
-                            job_id=load.job_id,
-                            kv_request_id=req_id,
-                            success=False,
+            assert st.loads
+            for round_seq, load in st.loads.items():
+                if load.aborted_at is None:
+                    if now - load.submitted_at >= _LOAD_TIMEOUT_S:
+                        load.aborted_at = now
+                        logger.warning(
+                            "P2PSession %s: %s round=%s timed out, sending abort",
+                            self._peer_id,
+                            req_id,
+                            round_seq,
                         )
-                    )
-                    logger.warning(
-                        "P2PSession %s: abort_ack timed out for kv_request_id=%s",
-                        self._peer_id,
-                        req_id,
-                    )
-        for req_id in to_remove:
-            self._requests[req_id].load = None
-            self._active_loads.discard(req_id)
-            self._maybe_prune(req_id)
+                        self._send(
+                            {
+                                TYPE_KEY: AbortFetchMsg.TYPE,
+                                AbortFetchMsg.KV_REQUEST_ID: req_id,
+                                AbortFetchMsg.ROUND_SEQ: round_seq,
+                            }
+                        )
+                else:
+                    if now - load.aborted_at >= _ABORT_ACK_TIMEOUT_S:
+                        to_remove.append((req_id, round_seq))
+                        self._completed_loads.append(
+                            LoadResult(
+                                job_id=load.job_id,
+                                kv_request_id=req_id,
+                                success=False,
+                            )
+                        )
+                        logger.warning(
+                            "P2PSession %s: abort_ack timed out for "
+                            "kv_request_id=%s round=%s",
+                            self._peer_id,
+                            req_id,
+                            round_seq,
+                        )
+        for req_id, round_seq in to_remove:
+            st = self._requests[req_id]
+            st.loads.pop(round_seq)
+            self._on_load_terminal(req_id, st)
 
         results = self._completed_loads
         self._completed_loads = []
@@ -510,12 +515,12 @@ class ClientRole:
         forever on an answer that can never arrive). See ``ClientCloseResult``.
         """
         failed_jobs = [
-            st.load.job_id for st in self._requests.values() if st.load is not None
+            load.job_id for st in self._requests.values() for load in st.loads.values()
         ]
         failed_req_ids = [
             req_id
             for req_id, st in self._requests.items()
-            if st.load is not None or any(hit is None for hit in st.probes.values())
+            if st.loads or any(hit is None for hit in st.probes.values())
         ]
         self._requests.clear()
         self._flush_pending.clear()
diff --git a/vllm/v1/kv_offload/tiering/p2p/session/protocol.py b/vllm/v1/kv_offload/tiering/p2p/session/protocol.py
index 8988c7e79f3..33ae5bf2f79 100644
--- a/vllm/v1/kv_offload/tiering/p2p/session/protocol.py
+++ b/vllm/v1/kv_offload/tiering/p2p/session/protocol.py
@@ -26,14 +26,12 @@ Block Transfer Flow (happy path)
 
 1. Client sends FetchMsg with a kv_request_id and lists of
    block keys + remote indexes where it wants the data written.
-   In p2p mode FetchMsg is also the server-side "request finished"
-   signal for the id: no further ``cb.create_store_job`` will fire
-   (parked LookupMsg batches are popped, so pending-key resolution
-   cannot promote a HIT after this point), all server-side lookup
-   state for the id is released, and ``cb.finish_request`` fires on
-   each dropped batch. The client emits exactly one FetchMsg per
-   lookup-touched request, including an empty one when no blocks
-   end up being fetched.
+   A request may run several lookup→fetch rounds; symmetric-P2P
+   messages carry ROUND_SEQ so each round's supply, demand, and
+   completion stay isolated. The terminal empty FetchMsg is the
+   server-side "request finished" signal for the id: parked
+   LookupMsg batches are popped and ``cb.finish_request`` fires on
+   each.
 2. Server matches requested blocks against locally stored blocks:
    - Blocks already available are transferred immediately via RDMA.
    - Blocks not yet available are recorded as "demanded" and
@@ -178,33 +176,32 @@ class DisconnectMsg:
 
 
 class FetchMsg:
-    """Client → Server: request blocks by key and close the lookup phase.
+    """Client → Server: request blocks for one lookup round.
 
-    In p2p mode FetchMsg is also the server-side "request finished"
-    signal for ``kv_request_id``: on receipt the server (a) fires no
-    further ``cb.create_store_job`` for this id — parked LookupMsg
-    batches are popped, so ``_resolve_pending_lookups`` cannot promote
-    a HIT_PENDING / RETRY key into a fresh pin after this point — and
-    (b) calls ``cb.finish_request(batch.ctx)`` on each dropped batch
-    so the TieringManager can release per-batch bookkeeping. In the
-    all-miss case the client emits an empty FetchMsg (``KEYS``
-    and ``BLOCK_INDEXES`` both empty) purely to fire this signal.
+    A non-empty fetch closes only its round. The terminal empty FetchMsg
+    (``KEYS`` and ``BLOCK_INDEXES`` both empty) is the "request
+    finished" signal: the server pops parked LookupMsg batches, calls
+    ``cb.finish_request`` on each, and drains any leftover supply.
 
     Fields:
         KV_REQUEST_ID: Identifies this block transfer request.
         KEYS: List of block keys (OffloadKey bytes). May be empty.
         BLOCK_INDEXES: List of remote block indexes (same length as KEYS).
+        ROUND_SEQ: Lookup round this fetch closes. PD clients never probe
+            and stay on their single round 0.
     """
 
     TYPE = "fetch"
     KV_REQUEST_ID = "kv_request_id"
     KEYS = "keys"
     BLOCK_INDEXES = "block_indexes"
+    ROUND_SEQ = "round_seq"
 
     @staticmethod
     def validate(msg: dict) -> None:
         """Raise ValueError if any field has an invalid type or value."""
         _require(msg, FetchMsg.KV_REQUEST_ID, str)
+        _require_non_neg_int(msg, FetchMsg.ROUND_SEQ)
         _require_list(msg, FetchMsg.KEYS)
         _require_list(msg, FetchMsg.BLOCK_INDEXES)
         keys = msg[FetchMsg.KEYS]
@@ -229,17 +226,21 @@ class LookupMsg:
     Fields:
         KV_REQUEST_ID: Identifies this lookup transaction.
         KEYS: List of block keys (OffloadKey bytes) to probe.
+        ROUND_SEQ: Lookup round these probes belong to; pinned supply is
+            parked under it for that round's fetch.
     """
 
     TYPE = "lookup"
     KV_REQUEST_ID = "kv_request_id"
     KEYS = "keys"
+    ROUND_SEQ = "round_seq"
 
     @staticmethod
     def validate(msg: dict) -> None:
         """Raise ValueError if any field has an invalid type or value."""
         _require(msg, LookupMsg.KV_REQUEST_ID, str)
         _require_list(msg, LookupMsg.KEYS)
+        _require_non_neg_int(msg, LookupMsg.ROUND_SEQ)
 
 
 class LookupRespMsg:
@@ -284,17 +285,22 @@ class TransferDoneMsg:
     Fields:
         KV_REQUEST_ID: The request that completed.
         SUCCESS: Whether the transfer completed successfully.
+        ROUND_SEQ: The fetch round that completed. Several loads can be
+            in flight per id (the scheduler submits loads incrementally),
+            so completions are matched by round.
     """
 
     TYPE = "transfer_done"
     KV_REQUEST_ID = "kv_request_id"
     SUCCESS = "success"
+    ROUND_SEQ = "round_seq"
 
     @staticmethod
     def validate(msg: dict) -> None:
         """Raise ValueError if any field has an invalid type or value."""
         _require(msg, TransferDoneMsg.KV_REQUEST_ID, str)
         _require(msg, TransferDoneMsg.SUCCESS, bool)
+        _require_non_neg_int(msg, TransferDoneMsg.ROUND_SEQ)
 
 
 class AbortFetchMsg:
@@ -302,15 +308,18 @@ class AbortFetchMsg:
 
     Fields:
         KV_REQUEST_ID: The request to cancel.
+        ROUND_SEQ: The fetch round to cancel.
     """
 
     TYPE = "abort_fetch"
     KV_REQUEST_ID = "kv_request_id"
+    ROUND_SEQ = "round_seq"
 
     @staticmethod
     def validate(msg: dict) -> None:
         """Raise ValueError if any field has an invalid type or value."""
         _require(msg, AbortFetchMsg.KV_REQUEST_ID, str)
+        _require_non_neg_int(msg, AbortFetchMsg.ROUND_SEQ)
 
 
 class AbortAckMsg:
@@ -318,12 +327,15 @@ class AbortAckMsg:
 
     Fields:
         KV_REQUEST_ID: The request that was cancelled.
+        ROUND_SEQ: The round that was cancelled; echoes AbortFetchMsg.
     """
 
     TYPE = "abort_ack"
     KV_REQUEST_ID = "kv_request_id"
+    ROUND_SEQ = "round_seq"
 
     @staticmethod
     def validate(msg: dict) -> None:
         """Raise ValueError if any field has an invalid type or value."""
         _require(msg, AbortAckMsg.KV_REQUEST_ID, str)
+        _require_non_neg_int(msg, AbortAckMsg.ROUND_SEQ)
diff --git a/vllm/v1/kv_offload/tiering/p2p/session/server.py b/vllm/v1/kv_offload/tiering/p2p/session/server.py
index 3709ecd9c53..8558acdf919 100644
--- a/vllm/v1/kv_offload/tiering/p2p/session/server.py
+++ b/vllm/v1/kv_offload/tiering/p2p/session/server.py
@@ -53,15 +53,6 @@ class StoreResult(NamedTuple):
     success: bool
 
 
-class _InflightXfer(NamedTuple):
-    """Metadata for a single inflight RDMA transfer, keyed by transfer_id."""
-
-    kv_request_id: str
-    block_count: int
-    # The set of store job IDs that contributed blocks to this transfer.
-    job_ids: set[int]
-
-
 class _MatchResult(NamedTuple):
     """Result of block matching: pairs ready for transfer."""
 
@@ -73,12 +64,17 @@ class _MatchResult(NamedTuple):
 
 @dataclass
 class _OutboundRequestState:
-    """Server-role state for a single peer fetch request.
+    """Server-role state for a single fetch round of a peer request.
 
-    The owning ``kv_request_id`` is the ``ServerRole._requests`` dict key
-    and is not duplicated on the value.
+    A kv_request_id may run several lookup→fetch rounds; rounds live in
+    ``_ServerRequestState.outbound`` keyed by the wire ``round_seq``, so
+    terminals touch only their own round.
     """
 
+    # Supply came from inbound lookup pins (symmetric): no late
+    # submit_store can arrive, so unmatched fetch demand fails fast. PD
+    # rounds park demand for stores instead.
+    lookup_supplied: bool = False
     demand_received: bool = False
     available: dict[OffloadKey, tuple[int, int]] = field(
         default_factory=dict
@@ -88,7 +84,8 @@ class _OutboundRequestState:
     )  # key → remote_block_idx: blocks peer wants, awaiting supply
     remaining: int = 0  # blocks that need to be transferred to client
     finishing: bool = False  # Signal finish request ASAP
-    # Job IDs that submit_store'd blocks for this request and have not
+    inflight: int = 0  # transfers submitted for this round, not yet polled
+    # Job IDs that submit_store'd blocks for this round and have not
     # yet emitted a StoreResult. The terminal-finalize helper drains
     # this set; poll-done and poll-failed discard entries as their
     # StoreResults fire.
@@ -145,6 +142,21 @@ class _OutboundRequestState:
         )
 
 
+@dataclass
+class _InflightXfer:
+    """Metadata for a single inflight RDMA transfer, keyed by transfer_id."""
+
+    kv_request_id: str
+    block_count: int
+    # The set of store job IDs that contributed blocks to this transfer.
+    job_ids: set[int]
+    # Round this transfer serves and its key in ``st.outbound``;
+    # remaining/finalize apply only while the round is still registered.
+    # Dummy default for test-seeded entries.
+    round: _OutboundRequestState = field(default_factory=_OutboundRequestState)
+    round_key: int = 0
+
+
 @dataclass
 class _ActiveLookup:
     """In-flight state for one inbound LookupMsg.
@@ -159,6 +171,8 @@ class _ActiveLookup:
     lookup_id: int
     kv_request_id: str
     ctx: ReqContext
+    # Wire round these probes belong to; pins park under it.
+    round_seq: int = 0
     # Keys from the inbound LookupMsg, preserved in wire order so
     # the aggregated response goes back in the same order.
     keys: list[OffloadKey] = field(default_factory=list)
@@ -185,6 +199,7 @@ class _PendingLookup(NamedTuple):
 
     keys: list[OffloadKey]
     enqueued_at: float
+    round_seq: int = 0
 
 
 @dataclass
@@ -198,17 +213,15 @@ class _ServerRequestState:
     is idle — see ``ServerRole._maybe_prune``.
     """
 
-    # Outbound serve state (PD + symmetric producer side). None until the
-    # first add_stored_blocks / on_fetch; reset to None on finalize/abort.
-    outbound: _OutboundRequestState | None = None
+    # Fetch rounds keyed by wire round_seq. A round is created by its
+    # first supply or its fetch and removed at its terminal (finalize /
+    # failure / abort).
+    outbound: dict[int, _OutboundRequestState] = field(default_factory=dict)
     # Raw inbound LookupMsgs not yet processed against the ParentManager.
     pending_lookups: list[_PendingLookup] = field(default_factory=list)
     # Per-LookupMsg state parked with HIT_PENDING / RETRY keys, keyed by
     # the (globally unique) lookup_id and re-polled each serve.
     lookups: dict[int, _ActiveLookup] = field(default_factory=dict)
-    # Start time of a pending abort drain (``time.monotonic``); None when
-    # no abort is in progress.
-    abort_started_at: float | None = None
     # Transfer ids in ``ServerRole._inflight`` for this id. Kept in sync
     # via _inflight_add / _inflight_pop so a non-empty set is an exact
     # "has any inflight transfer" predicate and the abort drain can
@@ -256,9 +269,9 @@ class ServerRole:
         # ``parent.on_request_finished`` in ``serve_external_requests``.
         self._finished_lookup_ctxs: list[ReqContext] = []
         self._lookup_id_counter: int = 0
-        # kv_request_ids with a parked abort awaiting drain — work-list so
-        # drain_pending_aborts doesn't scan every request each poll tick.
-        self._parked_aborts: set[str] = set()
+        # Parked aborts awaiting drain, keyed by (kv_request_id, round)
+        # with the abort start time.
+        self._pending_aborts: dict[tuple[str, int], float] = {}
 
     # ------------------------------------------------------------------
     # State helpers
@@ -277,11 +290,11 @@ class ServerRole:
         st = self._requests.get(kv_request_id)
         if (
             st is not None
-            and st.outbound is None
+            and not st.outbound
             and not st.inflight_tids
             and not st.lookups
             and not st.pending_lookups
-            and st.abort_started_at is None
+            and not any(kv == kv_request_id for kv, _ in self._pending_aborts)
         ):
             del self._requests[kv_request_id]
 
@@ -295,87 +308,100 @@ class ServerRole:
         keys: Sequence[OffloadKey],
         block_ids: Sequence[int],
         job_id: JobId,
+        round_seq: int = 0,
+        *,
+        from_lookup: bool = False,
     ) -> None:
-        """New blocks stored locally — match against pending fetch demand."""
+        """New blocks stored locally — match within their fetch round.
+
+        Lookup pins carry the round they were probed under; PD
+        submit_store batches share PD's single round 0.
+        """
         self._store_jobs[job_id] = time.monotonic()
         st = self._get_or_create_request(kv_request_id)
-        if st.outbound is None:
-            st.outbound = _OutboundRequestState()
-        result = st.outbound.add_stored_blocks(keys, block_ids, job_id)
-        if result.local_idxs and st.outbound.demand_received:
-            self._submit_transfer(kv_request_id, result)
+        rnd = st.outbound.get(round_seq)
+        if rnd is None:
+            rnd = st.outbound[round_seq] = _OutboundRequestState()
+        if from_lookup:
+            rnd.lookup_supplied = True
+        result = rnd.add_stored_blocks(keys, block_ids, job_id)
+        if result.local_idxs and rnd.demand_received:
+            self._submit_transfer(kv_request_id, result, rnd, round_seq)
 
     def on_fetch(
         self,
         kv_request_id: str,
         keys: Sequence[OffloadKey],
         block_indexes: Sequence[int],
+        round_seq: int = 0,
     ) -> None:
         """Handle a FetchMsg from the peer.
 
-        In p2p mode FetchMsg is the server-side "request finished"
-        signal for ``kv_request_id``. Three consequences flow from that:
-
-        - No further ``parent.create_store_job`` will fire for this id:
-          the request's parked ``lookups`` are popped here (via
-          ``_finish_inbound_lookups``) before the next
-          ``serve_external_requests`` runs ``_resolve_pending_lookups``,
-          so any HIT_PENDING / RETRY key that would otherwise later
-          promote to HIT and pin a slot is dropped instead. Any raw
-          not-yet-processed LookupMsg for this id is dropped too.
-        - All server-side lookup state for the id is cleaned up (the
-          ``lookups`` entries themselves).
-        - The synthetic ctx is queued for ``parent.on_request_finished``
-          (fired by the next ``serve_external_requests``) so the
-          TieringManager can release per-lookup bookkeeping.
-
-        The client contract guarantees exactly one FetchMsg per
-        lookup-touched request — including an empty one when no
-        blocks end up being fetched.
-
-        Raises ``ValueError`` on a duplicate fetch for the same
-        ``kv_request_id``; the coordinator's dispatch loop turns that
-        into a protocol-error disconnect.
+        A non-empty fetch binds and closes its round, leaving lookup
+        state alone (the next round's LookupMsg may already be in
+        flight). The terminal empty fetch closes the id: parked lookups
+        are popped and every remaining round drained. A second fetch for
+        a round already holding demand raises ValueError
+        (protocol-error disconnect).
         """
         logger.debug(
-            "P2PSession %s: fetch RECEIVED kv_request_id=%s blocks=%d",
+            "P2PSession %s: fetch RECEIVED kv_request_id=%s round=%s blocks=%d",
             self._peer_id,
             kv_request_id,
+            round_seq,
             len(keys),
         )
         st = self._requests.get(kv_request_id)
-        existing = st.outbound if st is not None else None
+        existing = st.outbound.get(round_seq) if st is not None else None
         if existing is not None and existing.demand_received:
-            # A second fetch for the same kv_request_id would overwrite
-            # `remaining` and leak inflight bookkeeping. Treat as a
-            # protocol violation.
-            raise ValueError(f"duplicate fetch for kv_request_id={kv_request_id}")
+            raise ValueError(
+                f"duplicate fetch for kv_request_id={kv_request_id} round={round_seq}"
+            )
         st = self._get_or_create_request(kv_request_id)
-        if st.outbound is None:
-            st.outbound = _OutboundRequestState()
-        req = st.outbound
+        req = st.outbound.get(round_seq)
+        if req is None:
+            req = st.outbound[round_seq] = _OutboundRequestState()
         result = req.add_fetch_demand(keys, block_indexes)
+        if not keys:
+            # Terminal empty fetch: close the lookup phase and drain
+            # every round with no TransferDoneMsg (nothing waits on it).
+            self._finish_inbound_lookups(kv_request_id)
+            for key in list(st.outbound):
+                self._finalize_outbound(kv_request_id, key, send_done=False)
+            return
+        if req.lookup_supplied and req.demanded:
+            # A symmetric round's supply always precedes its fetch, so
+            # unmatched demand is unservable — fail now, not at the load
+            # timeout. PD rounds keep parking demand for stores that
+            # arrive later.
+            logger.warning(
+                "P2PSession %s: fetch kv_request_id=%s round=%s demanded %d "
+                "blocks but %d have no pinned supply; failing fetch "
+                "immediately",
+                self._peer_id,
+                kv_request_id,
+                round_seq,
+                len(keys),
+                len(req.demanded),
+            )
+            self._finalize_outbound(kv_request_id, round_seq, success=False)
+            return
         if result.local_idxs:
-            self._submit_transfer(kv_request_id, result)
-        # Close the peer's request as far as the server's lookup phase
-        # is concerned: pop parked lookups so no further
-        # ``parent.create_store_job`` fires for this id, and queue their
-        # ctxs for ``parent.on_request_finished``. Done before the
-        # finalize path below so bookkeeping releases in-order.
-        self._finish_inbound_lookups(kv_request_id)
+            self._submit_transfer(kv_request_id, result, req, round_seq)
         # Prefiller-first mode: finish_request may have run before
         # fetch arrived. If so, finalize once we know what was
         # demanded — fully satisfied → success, else early-fail.
-        if req.finishing and not self._has_inflight_for(kv_request_id):
-            self._finalize_outbound(kv_request_id)
+        if req.finishing and req.inflight == 0:
+            self._finalize_outbound(kv_request_id, round_seq)
 
-    def on_abort_fetch(self, kv_request_id: str) -> None:
-        """Handle an AbortFetchMsg from the peer."""
+    def on_abort_fetch(self, kv_request_id: str, round_seq: int = 0) -> None:
+        """Handle an AbortFetchMsg from the peer, cancelling one round."""
         # Abort for an unknown id may be a benign race/duplicate or a
         # real protocol violation; we don't track completed ids, so warn.
         st = self._requests.get(kv_request_id)
-        has_outbound = st is not None and st.outbound is not None
-        if not has_outbound and not self._has_inflight_for(kv_request_id):
+        if (st is None or not st.outbound) and not self._has_inflight_for(
+            kv_request_id
+        ):
             logger.warning(
                 "P2PSession %s: abort_fetch for unknown kv_request_id=%s "
                 "(no outbound or inflight state); benign race or stale",
@@ -385,16 +411,15 @@ class ServerRole:
         # Idempotent: receiving AbortFetchMsg again before we've sent the
         # ack just triggers another drain attempt without resetting the
         # deadline.
-        st = self._get_or_create_request(kv_request_id)
-        if st.abort_started_at is None:
-            st.abort_started_at = time.monotonic()
-            self._parked_aborts.add(kv_request_id)
-        self._drain_abort(kv_request_id)
+        self._get_or_create_request(kv_request_id)
+        self._pending_aborts.setdefault((kv_request_id, round_seq), time.monotonic())
+        self._drain_abort(kv_request_id, round_seq)
 
     def on_lookup(
         self,
         kv_request_id: str,
         keys: Sequence[OffloadKey],
+        round_seq: int = 0,
     ) -> None:
         """Enqueue a LookupMsg from a symmetric-P2P consumer.
 
@@ -406,13 +431,18 @@ class ServerRole:
         parent calls are valid.
         """
         logger.debug(
-            "P2P LOOKUP server %s: RECV LookupMsg kv_request_id=%s keys=%d",
+            "P2P LOOKUP server %s: RECV LookupMsg kv_request_id=%s round=%s keys=%d",
             self._peer_id,
             kv_request_id,
+            round_seq,
             len(keys),
         )
         self._get_or_create_request(kv_request_id).pending_lookups.append(
-            _PendingLookup(keys=list(keys), enqueued_at=time.monotonic())
+            _PendingLookup(
+                keys=list(keys),
+                enqueued_at=time.monotonic(),
+                round_seq=round_seq,
+            )
         )
         self._serve_pending.add(kv_request_id)
 
@@ -434,7 +464,7 @@ class ServerRole:
                 st.pending_lookups = []
                 for pl in pending:
                     self._process_inbound_lookup(
-                        kv_request_id, pl.keys, pl.enqueued_at, parent
+                        kv_request_id, pl.keys, pl.enqueued_at, pl.round_seq, parent
                     )
             self._resolve_pending_lookups(kv_request_id, parent)
             st = self._requests.get(kv_request_id)
@@ -481,9 +511,7 @@ class ServerRole:
             else:
                 lookup.pending.add(h)
         if new_hits:
-            self._pin_and_register_hits(
-                lookup.kv_request_id, new_hits, lookup.ctx, parent
-            )
+            self._pin_and_register_hits(lookup, new_hits, parent)
         return new_hits
 
     def _process_inbound_lookup(
@@ -491,6 +519,7 @@ class ServerRole:
         kv_request_id: str,
         keys: list[OffloadKey],
         enqueued_at: float,
+        round_seq: int,
         parent: ParentManager,
     ) -> None:
         """Resolve one enqueued LookupMsg against ``parent``.
@@ -515,6 +544,7 @@ class ServerRole:
             lookup_id=lookup_id,
             kv_request_id=kv_request_id,
             ctx=ctx,
+            round_seq=round_seq,
             keys=list(keys),
             deadline=enqueued_at + _LOOKUP_PENDING_TIMEOUT_S,
         )
@@ -547,25 +577,26 @@ class ServerRole:
 
     def _pin_and_register_hits(
         self,
-        kv_request_id: str,
+        lookup: _ActiveLookup,
         keys: list[OffloadKey],
-        ctx: ReqContext,
         parent: ParentManager,
     ) -> None:
-        """Pin primary slots for HIT keys and feed them into the
-        existing ``add_stored_blocks`` matching path.
+        """Pin primary slots for HIT keys and park them as the lookup's
+        round supply via ``add_stored_blocks``.
 
         Caller has already confirmed every key is HIT (single-threaded
         scheduler ⇒ no eviction race), so the JobMetadata returned by
         ``parent.create_store_job`` carries parallel ``keys``/``block_ids``
         of length ``len(keys)``.
         """
-        meta = parent.create_store_job(keys, ctx)
+        meta = parent.create_store_job(keys, lookup.ctx)
         self.add_stored_blocks(
-            kv_request_id,
+            lookup.kv_request_id,
             list(meta.keys),
             list(meta.block_ids),
             meta.job_id,
+            round_seq=lookup.round_seq,
+            from_lookup=True,
         )
 
     def _resolve_pending_lookups(
@@ -650,9 +681,8 @@ class ServerRole:
 
         Called on the two events that mean "no more lookup traffic for
         ``kv_request_id`` is expected on this session": the terminal
-        FetchMsg from the peer (client contract: exactly one FetchMsg
-        per lookup-touched request, even if empty), and a local
-        ``finish``. Whichever fires second is a no-op.
+        empty FetchMsg from the peer and a local ``finish``. Whichever
+        fires second is a no-op.
         """
         st = self._requests.get(kv_request_id)
         if st is None:
@@ -688,20 +718,15 @@ class ServerRole:
         self._finish_inbound_lookups(kv_request_id)
 
         st = self._requests.get(kv_request_id)
-        req = st.outbound if st is not None else None
-        if req is None:
+        if st is None:
             return
-        req.finishing = True
-        if not req.demand_received:
-            return
-        if self._has_inflight_for(kv_request_id):
-            return
-        # Remaining > 0 here: if it had hit 0, the poll-done success
-        # branch would have already cleared outbound and we'd have
-        # returned at `req is None` above. Helper derives success from
-        # remaining and emits StoreResult(success=False) for any
-        # leftover pending jobs.
-        self._finalize_outbound(kv_request_id)
+        for key, req in list(st.outbound.items()):
+            req.finishing = True
+            if not req.demand_received or req.inflight:
+                # No demand yet (prefiller-first): on_fetch finalizes via
+                # `finishing`. Inflight: the last completion finalizes.
+                continue
+            self._finalize_outbound(kv_request_id, key)
 
     def collect_results(self) -> list[StoreResult]:
         """Drain timeouts, deferred results, and transport completions.
@@ -741,20 +766,24 @@ class ServerRole:
                 )
                 continue
             results.extend(self._settle_xfer_jobs(xfer, success=True))
+            rnd = xfer.round
             st = self._requests.get(xfer.kv_request_id)
-            req = st.outbound if st is not None else None
-            if req is not None and req.demand_received:
-                req.remaining -= xfer.block_count
-                assert req.remaining >= 0, (
+            if st is not None and st.outbound.get(xfer.round_key) is rnd:
+                rnd.remaining -= xfer.block_count
+                assert rnd.remaining >= 0, (
                     f"remaining went negative for kv_request_id={xfer.kv_request_id}"
                 )
-                if req.remaining == 0:
-                    self._finalize_outbound(xfer.kv_request_id, success=True)
-                elif req.finishing and not self._has_inflight_for(xfer.kv_request_id):
-                    self._finalize_outbound(xfer.kv_request_id, success=False)
+                if rnd.remaining == 0:
+                    self._finalize_outbound(
+                        xfer.kv_request_id, xfer.round_key, success=True
+                    )
+                elif rnd.finishing and rnd.inflight == 0:
+                    self._finalize_outbound(
+                        xfer.kv_request_id, xfer.round_key, success=False
+                    )
             self._maybe_prune(xfer.kv_request_id)
 
-        failed_kv_request_ids: set[str] | None = None
+        failed_rounds: list[tuple[str, _OutboundRequestState]] | None = None
         for tid in poll_result.failed:
             xfer = self._inflight_pop(tid)
             if xfer is None:
@@ -767,35 +796,36 @@ class ServerRole:
                     tid,
                 )
                 continue
-            if failed_kv_request_ids is None:
-                failed_kv_request_ids = set()
-            failed_kv_request_ids.add(xfer.kv_request_id)
             results.extend(self._settle_xfer_jobs(xfer, success=False))
+            rnd = xfer.round
             st = self._requests.get(xfer.kv_request_id)
-            req = st.outbound if st is not None else None
-            if st is not None:
-                st.outbound = None
-            if req is not None and req.demand_received:
+            if st is not None and st.outbound.get(xfer.round_key) is rnd:
+                del st.outbound[xfer.round_key]
+                if failed_rounds is None:
+                    failed_rounds = []
+                failed_rounds.append((xfer.kv_request_id, rnd))
                 self._send(
                     {
                         TYPE_KEY: TransferDoneMsg.TYPE,
                         TransferDoneMsg.KV_REQUEST_ID: xfer.kv_request_id,
                         TransferDoneMsg.SUCCESS: False,
+                        TransferDoneMsg.ROUND_SEQ: xfer.round_key,
                     }
                 )
             self._maybe_prune(xfer.kv_request_id)
 
-        # Cancel other inflight for the same failed kv_request_ids
-        if failed_kv_request_ids:
-            ids_to_cancel = [
-                tid
-                for tid, xfer in self._inflight.items()
-                if xfer.kv_request_id in failed_kv_request_ids
-            ]
-            for tid in ids_to_cancel:
-                self._inflight_pop(tid)
-            self._transport.cancel(ids_to_cancel)
-            for kv_request_id in failed_kv_request_ids:
+        # Cancel each failed round's other inflight and fail its
+        # remaining store jobs — nothing else will settle them.
+        if failed_rounds:
+            for kv_request_id, rnd in failed_rounds:
+                ids_to_cancel = [
+                    tid for tid, x in self._inflight.items() if x.round is rnd
+                ]
+                for tid in ids_to_cancel:
+                    self._inflight_pop(tid)
+                if ids_to_cancel:
+                    self._transport.cancel(ids_to_cancel)
+                results.extend(self._fail_round_jobs(rnd))
                 self._maybe_prune(kv_request_id)
 
         return results
@@ -811,8 +841,8 @@ class ServerRole:
 
     def drain_pending_aborts(self) -> None:
         """Re-attempt every parked abort once per poll tick."""
-        for kv_request_id in list(self._parked_aborts):
-            self._drain_abort(kv_request_id)
+        for kv_request_id, round_seq in list(self._pending_aborts):
+            self._drain_abort(kv_request_id, round_seq)
 
     def close(self) -> tuple[list[int], list[ReqContext]]:
         """Tear down. Cancels inflight.
@@ -839,7 +869,7 @@ class ServerRole:
         failed_serves.extend(self._finished_lookup_ctxs)
         self._requests.clear()
         self._serve_pending.clear()
-        self._parked_aborts.clear()
+        self._pending_aborts.clear()
         self._finished_lookup_ctxs.clear()
         return failed_stores, failed_serves
 
@@ -870,6 +900,8 @@ class ServerRole:
         xfer = self._inflight.pop(tid, None)
         if xfer is None:
             return None
+        xfer.round.inflight -= 1
+        assert xfer.round.inflight >= 0
         st = self._requests.get(xfer.kv_request_id)
         if st is not None:
             st.inflight_tids.discard(tid)
@@ -880,80 +912,108 @@ class ServerRole:
     ) -> list[StoreResult]:
         """Emit StoreResults for a completed transfer's store jobs.
 
-        Pops each attached job from ``_store_jobs`` and clears it from the
-        request's pending set. A job already popped (via timeout, cancel,
-        etc.) is skipped so we never double-emit a contradictory result.
+        Pops each attached job from ``_store_jobs`` and clears it from
+        its round's pending set. A job already popped (via timeout,
+        cancel, etc.) is skipped so we never double-emit a contradictory
+        result.
         """
         results: list[StoreResult] = []
-        st = self._requests.get(xfer.kv_request_id)
-        req = st.outbound if st is not None else None
         for job_id in xfer.job_ids:
             if self._store_jobs.pop(job_id, None) is None:
                 continue
             results.append(StoreResult(job_id=job_id, success=success))
-            if req is not None:
-                req.pending_job_ids.discard(job_id)
+            xfer.round.pending_job_ids.discard(job_id)
         return results
 
     # ------------------------------------------------------------------
     # Internal — finalize / abort drain
     # ------------------------------------------------------------------
 
+    def _fail_round_jobs(self, rnd: _OutboundRequestState) -> list[StoreResult]:
+        """Fail a terminated round's still-pending store jobs (idempotent)."""
+        results: list[StoreResult] = []
+        for job_id in rnd.pending_job_ids:
+            if self._store_jobs.pop(job_id, None) is None:
+                continue
+            results.append(StoreResult(job_id=job_id, success=False))
+        rnd.pending_job_ids.clear()
+        return results
+
     def _finalize_outbound(
         self,
         kv_request_id: str,
+        round_key: int,
         success: bool | None = None,
+        send_done: bool = True,
     ) -> None:
-        """Pop the outbound state and emit terminal results.
-
-        Called when no further work will happen for this kv_request_id
-        on the server side: either request_finish has fired and there
-        are no inflight transfers, or the last inflight just completed
-        while finishing.
+        """Pop one round and emit its terminal results.
 
         If ``success`` is None, derive it from ``req.remaining == 0``.
-        The same flag is used for both the peer's TransferDoneMsg and
-        the StoreResult(s) emitted for any leftover pending job_ids.
+        ``send_done=False`` skips the TransferDoneMsg (terminal empty
+        fetch). Other rounds of the id are untouched.
         """
         st = self._requests[kv_request_id]
-        assert st.outbound is not None
-        req = st.outbound
-        st.outbound = None
+        req = st.outbound.pop(round_key)
         if success is None:
-            success = req.remaining == 0
-        for job_id in req.pending_job_ids:
-            self._store_jobs.pop(job_id, None)
-            self._pending_store_results.append(
-                StoreResult(job_id=job_id, success=success)
-            )
-        self._send(
-            {
-                TYPE_KEY: TransferDoneMsg.TYPE,
-                TransferDoneMsg.KV_REQUEST_ID: kv_request_id,
-                TransferDoneMsg.SUCCESS: success,
-            }
+            success = req.demand_received and req.remaining == 0
+        settled = self._fail_round_jobs(req) if not success else None
+        if settled is not None:
+            self._pending_store_results.extend(settled)
+        else:
+            for job_id in req.pending_job_ids:
+                if self._store_jobs.pop(job_id, None) is None:
+                    continue
+                self._pending_store_results.append(
+                    StoreResult(job_id=job_id, success=True)
+                )
+            req.pending_job_ids.clear()
+        logger.debug(
+            "P2PSession %s: finalize kv_request_id=%s round=%s success=%s "
+            "remaining=%d leftover_available=%d send_done=%s",
+            self._peer_id,
+            kv_request_id,
+            round_key,
+            success,
+            req.remaining,
+            len(req.available),
+            send_done,
         )
+        if send_done and req.demand_received:
+            self._send(
+                {
+                    TYPE_KEY: TransferDoneMsg.TYPE,
+                    TransferDoneMsg.KV_REQUEST_ID: kv_request_id,
+                    TransferDoneMsg.SUCCESS: success,
+                    TransferDoneMsg.ROUND_SEQ: round_key,
+                }
+            )
         self._maybe_prune(kv_request_id)
 
-    def _drain_abort(self, kv_request_id: str) -> None:
+    def _drain_abort(self, kv_request_id: str, round_seq: int) -> None:
         """One drain attempt for a pending abort.
 
-        Stops accepting more blocks for ``kv_request_id``, then asks the
-        transport to cancel any matching inflight transfers in
-        ``mode="wait"``. Sends ``AbortAckMsg`` once nothing remains
-        inflight, or after ``_CANCEL_DRAIN_TIMEOUT_S`` falls back to
-        ``mode="immediate"`` and acks anyway.
+        Detaches the aborted round, then asks the transport to cancel its
+        inflight transfers in ``mode="wait"``. Sends ``AbortAckMsg`` once
+        nothing remains inflight, or after ``_CANCEL_DRAIN_TIMEOUT_S``
+        falls back to ``mode="immediate"`` and acks anyway.
         """
         st = self._requests[kv_request_id]
-        st.outbound = None
-        ids = list(st.inflight_tids)
+        rnd = st.outbound.pop(round_seq, None)
+        if rnd is not None:
+            # Its transfers are being cancelled; fail its jobs now
+            # instead of leaking them to the store timeout.
+            self._pending_store_results.extend(self._fail_round_jobs(rnd))
+        ids = [
+            tid
+            for tid, x in self._inflight.items()
+            if x.kv_request_id == kv_request_id and x.round_key == round_seq
+        ]
         if not ids:
-            self._finalize_abort(kv_request_id)
+            self._finalize_abort(kv_request_id, round_seq)
             return
 
-        assert st.abort_started_at is not None
-        expired = time.monotonic() - st.abort_started_at >= _CANCEL_DRAIN_TIMEOUT_S
-        if expired:
+        started_at = self._pending_aborts[(kv_request_id, round_seq)]
+        if time.monotonic() - started_at >= _CANCEL_DRAIN_TIMEOUT_S:
             for tid in ids:
                 self._inflight_pop(tid)
             self._transport.cancel(ids, mode="immediate")
@@ -964,7 +1024,7 @@ class ServerRole:
                 kv_request_id,
                 len(ids),
             )
-            self._finalize_abort(kv_request_id)
+            self._finalize_abort(kv_request_id, round_seq)
             return
 
         still = self._transport.cancel(ids, mode="wait")
@@ -977,16 +1037,15 @@ class ServerRole:
             if tid not in still_set:
                 self._inflight_pop(tid)
         if not still:
-            self._finalize_abort(kv_request_id)
+            self._finalize_abort(kv_request_id, round_seq)
 
-    def _finalize_abort(self, kv_request_id: str) -> None:
-        st = self._requests[kv_request_id]
-        st.abort_started_at = None
-        self._parked_aborts.discard(kv_request_id)
+    def _finalize_abort(self, kv_request_id: str, round_seq: int) -> None:
+        self._pending_aborts.pop((kv_request_id, round_seq), None)
         self._send(
             {
                 TYPE_KEY: AbortAckMsg.TYPE,
                 AbortAckMsg.KV_REQUEST_ID: kv_request_id,
+                AbortAckMsg.ROUND_SEQ: round_seq,
             }
         )
         self._maybe_prune(kv_request_id)
@@ -995,7 +1054,13 @@ class ServerRole:
     # Internal — transfers and store-job timeouts
     # ------------------------------------------------------------------
 
-    def _submit_transfer(self, kv_request_id: str, result: _MatchResult) -> None:
+    def _submit_transfer(
+        self,
+        kv_request_id: str,
+        result: _MatchResult,
+        rnd: _OutboundRequestState,
+        round_key: int,
+    ) -> None:
         logger.debug(
             "P2PSession %s: NIXL write_blocks CALL kv_request_id=%s "
             "local_idxs=%d remote_idxs=%d",
@@ -1016,12 +1081,15 @@ class ServerRole:
                 transfer_id,
                 len(result.local_idxs),
             )
+            rnd.inflight += 1
             self._inflight_add(
                 transfer_id,
                 _InflightXfer(
                     kv_request_id=kv_request_id,
                     block_count=len(result.local_idxs),
                     job_ids=result.job_ids,
+                    round=rnd,
+                    round_key=round_key,
                 ),
             )
         else:
@@ -1031,22 +1099,24 @@ class ServerRole:
                 kv_request_id,
                 len(result.local_idxs),
             )
-            # The matched blocks were popped from req.demanded /
-            # req.available, but no inflight will satisfy them, so
-            # remaining will never reach 0 on its own. Mark the
-            # request as finishing so the existing terminal paths
-            # clean up: if other inflight is in flight, the last one
-            # to drain will fire _finalize_outbound(success=False)
-            # via the elif branch in collect_results. If
-            # nothing else is in flight, finalize now so the peer
-            # and the local store jobs don't wait for finish_request
-            # or for _STORE_TIMEOUT_S / _LOAD_TIMEOUT_S.
+            # The matched blocks were popped from rnd.demanded /
+            # rnd.available, but no inflight will satisfy them, so
+            # remaining will never reach 0 on its own. Mark the round
+            # as finishing so the existing terminal paths clean up: if
+            # other transfers of this round are in flight, the last one
+            # to drain will fire _finalize_outbound(success=False) via
+            # the elif branch in collect_results. If nothing else is in
+            # flight, finalize now so the peer and the local store jobs
+            # don't wait for finish_request or for _STORE_TIMEOUT_S /
+            # _LOAD_TIMEOUT_S.
+            rnd.finishing = True
             st = self._requests.get(kv_request_id)
-            req = st.outbound if st is not None else None
-            if req is not None:
-                req.finishing = True
-                if not self._has_inflight_for(kv_request_id):
-                    self._finalize_outbound(kv_request_id, success=False)
+            if (
+                st is not None
+                and st.outbound.get(round_key) is rnd
+                and rnd.inflight == 0
+            ):
+                self._finalize_outbound(kv_request_id, round_key, success=False)
 
     def _timeout_pending_store_jobs(self) -> list[StoreResult]:
         if not self._store_jobs:
diff --git a/vllm/v1/kv_offload/tiering/p2p/session/session.py b/vllm/v1/kv_offload/tiering/p2p/session/session.py
index 7d19913b0a2..4bd780159a4 100644
--- a/vllm/v1/kv_offload/tiering/p2p/session/session.py
+++ b/vllm/v1/kv_offload/tiering/p2p/session/session.py
@@ -374,26 +374,34 @@ class P2PSession:
                 for bh in msg[FetchMsg.KEYS]
             ]
             block_indexes = msg[FetchMsg.BLOCK_INDEXES]
+            round_seq = msg[FetchMsg.ROUND_SEQ]
             # Run the server-role state machine inline as today —
             # add_fetch_demand records demand against any blocks we've
             # already seen in `available`. Report the kv_request_id so
             # the manager (after poll() returns) can replay any parked
             # submit_store batches; their add_stored_blocks calls hit
             # the demand recorded here and submit transfers immediately.
-            self._server.on_fetch(kv_request_id, keys, block_indexes)
+            self._server.on_fetch(kv_request_id, keys, block_indexes, round_seq)
             self._new_fetch_ids.append(kv_request_id)
         elif msg_type == AbortFetchMsg.TYPE:
             AbortFetchMsg.validate(msg)
-            self._server.on_abort_fetch(msg[AbortFetchMsg.KV_REQUEST_ID])
+            self._server.on_abort_fetch(
+                msg[AbortFetchMsg.KV_REQUEST_ID],
+                msg[AbortFetchMsg.ROUND_SEQ],
+            )
         elif msg_type == TransferDoneMsg.TYPE:
             TransferDoneMsg.validate(msg)
             self._client.on_transfer_done(
                 msg[TransferDoneMsg.KV_REQUEST_ID],
                 msg[TransferDoneMsg.SUCCESS],
+                msg[TransferDoneMsg.ROUND_SEQ],
             )
         elif msg_type == AbortAckMsg.TYPE:
             AbortAckMsg.validate(msg)
-            self._client.on_abort_ack(msg[AbortAckMsg.KV_REQUEST_ID])
+            self._client.on_abort_ack(
+                msg[AbortAckMsg.KV_REQUEST_ID],
+                msg[AbortAckMsg.ROUND_SEQ],
+            )
         elif msg_type == LookupMsg.TYPE:
             LookupMsg.validate(msg)
             kv_request_id = msg[LookupMsg.KV_REQUEST_ID]
@@ -401,7 +409,7 @@ class P2PSession:
                 OffloadKey(bh if isinstance(bh, bytes) else bytes(bh))
                 for bh in msg[LookupMsg.KEYS]
             ]
-            self._server.on_lookup(kv_request_id, keys)
+            self._server.on_lookup(kv_request_id, keys, msg[LookupMsg.ROUND_SEQ])
         elif msg_type == LookupRespMsg.TYPE:
             LookupRespMsg.validate(msg)
             kv_request_id = msg[LookupRespMsg.KV_REQUEST_ID]
diff --git a/vllm/v1/kv_offload/tiering/spec.py b/vllm/v1/kv_offload/tiering/spec.py
index 0d200a42390..e6ec192415e 100644
--- a/vllm/v1/kv_offload/tiering/spec.py
+++ b/vllm/v1/kv_offload/tiering/spec.py
@@ -9,8 +9,13 @@ and configurable secondary tiers (e.g., Storage, Network).
 Configuration via kv_connector_extra_config:
   - cpu_bytes_to_use: (required) Bytes to allocate for CPU primary tier
   - block_size: (optional) Block size for offloaded blocks (default: GPU block size)
-  - eviction_policy: (optional) Primary tier eviction policy: "lru" or
-    "arc" (default: "lru")
+  - eviction_policy: (optional) Primary tier eviction policy: built-in "lru"/
+    "arc", or the name of a policy registered via CachePolicyFactory, or an
+    out-of-tree CachePolicy class name paired with cache_policy_module_path
+    (default: "lru")
+  - cache_policy_module_path: (optional) Python import path to load
+    eviction_policy from when it names an out-of-tree CachePolicy not
+    registered via CachePolicyFactory
   - secondary_tiers: (optional) List of secondary tier configurations
     Each secondary tier config is a dict with:
       - type: (required) Type of secondary tier (e.g., "example", "storage", "network")
@@ -178,7 +183,8 @@ class TieringOffloadingSpec(CPUOffloadingSpec):
             # Create primary tier (CPU-based)
             primary_tier = CPUPrimaryTierOffloadingManager(
                 num_blocks=self.num_blocks,
-                cache_policy=self.eviction_policy,  # type: ignore[arg-type]
+                cache_policy=self.eviction_policy,
+                cache_policy_module_path=self.cache_policy_module_path,
                 enable_events=self.kv_events_config.enable_kv_cache_events,
                 mmap_region=scheduler_mmap,
             )
diff --git a/vllm/v1/sample/logits_processor/__init__.py b/vllm/v1/sample/logits_processor/__init__.py
index 2cb89e1ea95..f319097495d 100644
--- a/vllm/v1/sample/logits_processor/__init__.py
+++ b/vllm/v1/sample/logits_processor/__init__.py
@@ -10,6 +10,7 @@ from typing import TYPE_CHECKING
 
 import torch
 
+from vllm.exceptions import VLLMValidationError
 from vllm.logger import init_logger
 from vllm.logits_process import LogitsProcessor as RequestLogitsProcessor
 from vllm.sampling_params import SamplingParams
@@ -228,7 +229,12 @@ def validate_logits_processors_parameters(
         tuple(logits_processors) if logits_processors is not None else None
     )
     for logits_procs in cached_load_custom_logitsprocs(logits_processors):
-        logits_procs.validate_params(sampling_params)
+        try:
+            logits_procs.validate_params(sampling_params)
+        except ValueError as e:
+            # Legacy custom logitsprocs may still raise ValueError from
+            # validate_params; convert for backward compatibility.
+            raise VLLMValidationError(str(e)) from e
 
 
 class AdapterLogitsProcessor(LogitsProcessor):
diff --git a/vllm/v1/sample/logits_processor/interface.py b/vllm/v1/sample/logits_processor/interface.py
index 3e426e321b3..dce0706335e 100644
--- a/vllm/v1/sample/logits_processor/interface.py
+++ b/vllm/v1/sample/logits_processor/interface.py
@@ -62,7 +62,9 @@ class LogitsProcessor(ABC):
     def validate_params(cls, sampling_params: SamplingParams):
         """Validate sampling params for this logits processor.
 
-        Raise ValueError for invalid ones.
+        Raise ``VLLMValidationError`` (preferred) / ``ValueError`` (backward compatible)
+        for invalid params. Bare ``ValueError`` is converted to ``VLLMValidationError``
+        at the engine boundary so online serving returns HTTP 400.
         """
         return None
 
diff --git a/vllm/v1/spec_decode/dflash.py b/vllm/v1/spec_decode/dflash.py
index 5d1c0629218..626dd36dea9 100644
--- a/vllm/v1/spec_decode/dflash.py
+++ b/vllm/v1/spec_decode/dflash.py
@@ -38,8 +38,12 @@ class DFlashProposer(SpecDecodeBaseProposer):
 
         # Only next_token_ids and mask tokens are query tokens, all other context is K/V
         self.max_query_tokens = self.max_batch_size * (1 + self.num_speculative_tokens)
+        self.max_padded_query_tokens = max(
+            self.max_query_tokens,
+            vllm_config.compilation_config.max_cudagraph_capture_size or 0,
+        )
         # Positions covers both context states + query states
-        self.max_positions = self.max_num_tokens + self.max_query_tokens
+        self.max_positions = self.max_num_tokens + self.max_padded_query_tokens
 
         # Separate context buffers to keep query buffer addresses stable for CUDA graphs
         self._context_slot_mapping_buffer = torch.zeros(
@@ -48,7 +52,7 @@ class DFlashProposer(SpecDecodeBaseProposer):
             device=device,
         )
         self._slot_mapping_buffer = torch.zeros(
-            self.max_query_tokens,
+            self.max_padded_query_tokens,
             dtype=torch.int64,
             device=device,
         )
@@ -58,7 +62,7 @@ class DFlashProposer(SpecDecodeBaseProposer):
             device=device,
         )
         self.positions = torch.zeros(
-            self.max_query_tokens,
+            self.max_padded_query_tokens,
             dtype=torch.int64,
             device=device,
         )
diff --git a/vllm/v1/worker/cpu/shm.py b/vllm/v1/worker/cpu/shm.py
index bd1f96c71ed..970e8a2d414 100644
--- a/vllm/v1/worker/cpu/shm.py
+++ b/vllm/v1/worker/cpu/shm.py
@@ -24,12 +24,17 @@ def fake_pin_memory(self: torch.Tensor, *args: Any, **kwargs: Any) -> torch.Tens
 class _EventPlaceholder:
     def __init__(self, *args, **kwargs) -> None:
         self.record = noop
+        self.wait = noop
         self.synchronize = noop
 
 
 class _StreamPlaceholder:
     def __init__(self, *args, **kwargs) -> None:
         self.wait_stream = noop
+        self.wait_event = noop
+        self.record_event = noop
+        self.synchronize = noop
+        self.query = lambda: True
         self.device = torch.device("cpu")
 
     def __enter__(self, *args, **kwargs):
@@ -55,6 +60,7 @@ torch.cuda.current_stream = lambda *args, **kwargs: _StreamPlaceholder()
 torch.accelerator.synchronize = noop
 torch.accelerator.empty_cache = noop
 torch.Tensor.pin_memory = fake_pin_memory
+torch.Tensor.record_stream = noop
 torch.accelerator.get_memory_info = get_memory_info
 
 # Patch vLLM torch utils
@@ -80,3 +86,9 @@ import vllm.v1.worker.gpu.buffer_utils as gpu_buffer_utils
 import vllm.v1.worker.cpu.buffer_utils as cpu_buffer_utils
 
 gpu_buffer_utils.UvaBuffer = cpu_buffer_utils.UvaBuffer
+
+# Patch Triton
+from vllm.triton_utils import HAS_TRITON, tl
+
+if HAS_TRITON:
+    tl.debug_barrier = noop
diff --git a/vllm/v1/worker/gpu/model_states/encoder_decoder.py b/vllm/v1/worker/gpu/model_states/encoder_decoder.py
index f759c0b1e15..618984c97f3 100644
--- a/vllm/v1/worker/gpu/model_states/encoder_decoder.py
+++ b/vllm/v1/worker/gpu/model_states/encoder_decoder.py
@@ -9,6 +9,7 @@ import torch.nn as nn
 
 from vllm.config import VllmConfig
 from vllm.config.compilation import CUDAGraphMode
+from vllm.utils.torch_utils import PIN_MEMORY
 from vllm.v1.kv_cache_interface import CrossAttentionSpec, KVCacheConfig
 from vllm.v1.worker.gpu.attn_utils import build_attn_metadata
 from vllm.v1.worker.gpu.input_batch import InputBatch
@@ -157,7 +158,9 @@ class EncoderDecoderModelState(ModelState):
         for_capture: bool,
         num_reqs: int,
     ) -> dict[int, tuple[torch.Tensor, np.ndarray]]:
-        encoder_seq_lens = torch.zeros(num_reqs, dtype=torch.int32, pin_memory=True)
+        encoder_seq_lens = torch.zeros(
+            num_reqs, dtype=torch.int32, pin_memory=PIN_MEMORY
+        )
         encoder_seq_lens_np = encoder_seq_lens.numpy()
         if not for_capture:
             # During normal execution, use actual encoder lengths.
diff --git a/vllm/v1/worker/gpu/warmup.py b/vllm/v1/worker/gpu/warmup.py
index 6f08bbf08c7..260ca7a5f17 100644
--- a/vllm/v1/worker/gpu/warmup.py
+++ b/vllm/v1/worker/gpu/warmup.py
@@ -169,9 +169,12 @@ def warmup_kernels(
     # a uniform decode batch.
     prompt_len = decode_query_len + 1
     prompt_token_ids = list(range(prompt_len))
+    # Upper bound on the decode steps built in `decode_steps` below.
     num_decode_steps = 1
     if not model_runner.is_pooling_model:
         num_decode_steps = 5 if num_spec_steps > 0 else 3
+    # Size the block allocation for the worst case: every request advancing
+    # decode_query_len tokens on every decode step.
     decode_len = prompt_len + num_decode_steps * decode_query_len
 
     kv_cache_groups = model_runner.kv_cache_config.kv_cache_groups
@@ -239,6 +242,8 @@ def warmup_kernels(
         nonlocal next_block_id
         return list(range(next_block_id, next_block_id := next_block_id + num_blocks))
 
+    # The KV-block zeroing kernel is driven by the scheduler's
+    # new_block_ids_to_zero, so none of the steps below reach it.
     if model_runner.kv_block_zeroer is not None:
         model_runner.kv_block_zeroer.warmup(model_runner.kv_cache_config.num_blocks)
 
@@ -286,10 +291,12 @@ def warmup_kernels(
 
         worker_sample_tokens(grammar_output)
 
+        # Per-request state carried across the decode steps.
         req_computed = [prompt_len] * num_reqs
         req_blocks = [list(prefill_block_counts) for _ in range(num_reqs)]
 
         def _run_decode_step(indices: list[int], spec_flags: list[bool]) -> None:
+            """Decode `indices`, spec-decoding the ones flagged in `spec_flags`."""
             cached_req_data = CachedRequestData.make_empty()
             cached_req_data.req_ids = [req_ids[i] for i in indices]
             cached_req_data.num_computed_tokens = [req_computed[i] for i in indices]
@@ -326,13 +333,22 @@ def warmup_kernels(
 
             worker_execute_model(decode_output)
             worker_sample_tokens(None)
+
             for i, use_spec in zip(indices, spec_flags):
                 req_computed[i] += decode_query_len if use_spec else 1
 
         all_indices = list(range(num_reqs))
         use_spec_decode = num_spec_steps > 0
-        decode_steps = [(all_indices, [use_spec_decode] * num_reqs)]
+
+        # Decode steps to warm, as (request indices, per-request spec flag).
+        # Under spec decoding the scheduler drops requests the drafter proposed
+        # nothing for, so warm each batch shape with and without draft tokens.
+        decode_steps: list[tuple[list[int], list[bool]]] = [
+            (all_indices, [use_spec_decode] * num_reqs),
+        ]
         if num_reqs >= 2:
+            # Mixed spec / non-spec: GDN and KDA reclassify the non-spec decode
+            # as a prefill and split the batch into spec/non-spec token indices.
             decode_steps.append(([0, 1], [use_spec_decode, False]))
             if use_spec_decode:
                 # Exercise the model paths that split a batch by whether each
@@ -345,8 +361,8 @@ def warmup_kernels(
         elif use_spec_decode:
             decode_steps.append(([0], [False]))
 
-        for indices, spec_flags in decode_steps:
-            _run_decode_step(indices, spec_flags)
+        for step_indices, step_spec_flags in decode_steps:
+            _run_decode_step(step_indices, step_spec_flags)
 
     # Clean up - process finish_req_ids.
     cleanup_output = SchedulerOutput.make_empty()
diff --git a/vllm/v1/worker/gpu_model_runner.py b/vllm/v1/worker/gpu_model_runner.py
index 21e18167c55..a38fa7c6a64 100644
--- a/vllm/v1/worker/gpu_model_runner.py
+++ b/vllm/v1/worker/gpu_model_runner.py
@@ -991,9 +991,16 @@ class GPUModelRunner(
             return
 
         kv_caches = getattr(self, "kv_caches", [])
-        for cache_tensor in kv_caches:
-            if cache_tensor is not None:
-                cache_tensor.zero_()
+        for cache_entry in kv_caches:
+            if cache_entry is None:
+                continue
+            # Hybrid models (Mamba, DeltaNet) store per-layer state as a
+            # list of tensors rather than a single tensor.
+            if isinstance(cache_entry, list):
+                for t in cache_entry:
+                    t.zero_()
+            else:
+                cache_entry.zero_()
 
         k_attr_names = ("_k_scale", "k_scale")
         v_attr_names = ("_v_scale", "v_scale")
diff --git a/vllm/v1/worker/gpu_worker.py b/vllm/v1/worker/gpu_worker.py
index 556b1e6c7d9..bbab98dbd7b 100644
--- a/vllm/v1/worker/gpu_worker.py
+++ b/vllm/v1/worker/gpu_worker.py
@@ -154,7 +154,7 @@ class Worker(WorkerBase):
             self.worker_sentinel = WorkerSentinel(worker=self)
         # Buffers saved before sleep
         self._sleep_saved_buffers: dict[str, torch.Tensor] = {}
-        self._sleep_rebuild_draft_metadata_buffers = False
+        self._sleep_saved_draft_buffers: dict[str, torch.Tensor] = {}
 
         # Weight transfer engine is created in `load_model` once the model
         # is available, since the engine needs a reference to the model.
@@ -200,10 +200,10 @@ class Worker(WorkerBase):
                 name: buffer.cpu().clone() for name, buffer in model.named_buffers()
             }
             draft = self.get_draft_model()
-            inner = getattr(draft, "model", None) if draft is not None else None
-            self._sleep_rebuild_draft_metadata_buffers = inner is not None and hasattr(
-                inner, "_build_fused_kv_buffers"
-            )
+            if draft is not None:
+                self._sleep_saved_draft_buffers = {
+                    name: buffer.cpu().clone() for name, buffer in draft.named_buffers()
+                }
 
         self._get_sleep_mode_backend().suspend(level)
 
@@ -228,20 +228,21 @@ class Worker(WorkerBase):
         self._get_sleep_mode_backend().resume(tags)
 
         # Restore the buffers after level 2 sleep
-        if len(self._sleep_saved_buffers):
+        wake_weights = tags is None or "weights" in tags
+        if wake_weights and len(self._sleep_saved_buffers):
             model = self.model_runner.model
             for name, buffer in model.named_buffers():
                 if name in self._sleep_saved_buffers:
                     buffer.data.copy_(self._sleep_saved_buffers[name].data)
             self._sleep_saved_buffers = {}
 
-        if self._sleep_rebuild_draft_metadata_buffers:
+        if wake_weights and len(self._sleep_saved_draft_buffers):
             draft = self.get_draft_model()
             if draft is not None:
-                inner = getattr(draft, "model", None)
-                if inner is not None and hasattr(inner, "_build_fused_kv_buffers"):
-                    inner._build_fused_kv_buffers()
-            self._sleep_rebuild_draft_metadata_buffers = False
+                for name, buffer in draft.named_buffers():
+                    if name in self._sleep_saved_draft_buffers:
+                        buffer.data.copy_(self._sleep_saved_draft_buffers[name].data)
+            self._sleep_saved_draft_buffers = {}
 
         if tags is None or "kv_cache" in tags:
             self.model_runner.post_kv_cache_wake_up()