[CI] Discover Rust coverage artifacts from build metadata

Co-authored-by: OpenAI Codex <codex@openai.com>

.buildkite/hardware_tests/amd.yaml

@@ -18,6 +18,8 @@ steps:
       TERM: "xterm-256color"
     retry:
       automatic:
+        - exit_status: 1  # Transient Docker/BuildKit failure
+          limit: 1
         - exit_status: -1  # Agent was lost
           limit: 1
         - exit_status: -10  # Agent was lost
@@ -46,6 +48,8 @@ steps:
       VLLM_BRANCH: "$BUILDKITE_COMMIT"
     retry:
       automatic:
+        - exit_status: 1  # Transient Docker/BuildKit failure
+          limit: 1
         - exit_status: -1  # Agent was lost
           limit: 1
         - exit_status: -10  # Agent was lost
@@ -72,6 +76,8 @@ steps:
       VLLM_BRANCH: "$BUILDKITE_COMMIT"
     retry:
       automatic:
+        - exit_status: 1  # Transient Docker/BuildKit failure
+          limit: 1
         - exit_status: -1  # Agent was lost
           limit: 1
         - exit_status: -10  # Agent was lost

.buildkite/hardware_tests/cpu.yaml

@@ -18,6 +18,8 @@ steps:
   - tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
   - tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
   - tests/kernels/mamba/test_cpu_short_conv.py
+  - tests/kernels/mamba/test_causal_conv1d.py
+  - tests/kernels/mamba/test_mamba_ssm.py
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
@@ -28,7 +30,9 @@ steps:
       pytest -x -v -s tests/kernels/test_onednn.py
       pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
       pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
-      pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
+      pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
+      pytest -x -v -s tests/kernels/mamba/test_causal_conv1d.py
+      pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.py"
 
 # Note: SDE can't be downloaded from CI host because of AWS WAF
 # - label: CPU-Compatibility Tests
@@ -141,9 +145,22 @@ steps:
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
-      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
+      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py --ignore=tests/models/multimodal/generation/test_qwen2_5_vl.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
   parallelism: 4
 
+- label: CPU-Qwen2.5-VL Multimodal Tests
+  depends_on: []
+  device: intel_cpu
+  no_plugin: true
+  source_file_dependencies:
+  # - vllm/
+  - vllm/model_executor/layers/rotary_embedding
+  - tests/models/multimodal/generation/
+  commands:
+    - |
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
+      VLLM_CI_ENV=0 pytest -x -v -s tests/models/multimodal/generation/test_qwen2_5_vl.py"
+
 - label: "Arm CPU Test"
   depends_on: []
   soft_fail: false

.buildkite/hardware_tests/intel_xpu_ci/test-intel.yaml

@@ -18,7 +18,7 @@ steps:
   - label: "XPU example Test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 30
+    timeout_in_minutes: 50
     optional: true
     device: intel_gpu
     agent_tags:
@@ -39,7 +39,7 @@ steps:
   - label: "XPU V1 test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 30
+    timeout_in_minutes: 70
     optional: true
     device: intel_gpu
     agent_tags:
@@ -60,7 +60,7 @@ steps:
   - label: "XPU server test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 30
+    timeout_in_minutes: 45
     optional: true
     device: intel_gpu
     agent_tags:

.buildkite/image_build/image_build.sh

@@ -177,6 +177,18 @@ BRANCH=$4
 IMAGE_TAG=$5
 IMAGE_TAG_LATEST=${6:-} # only used for main branch, optional
 
+# When TORCH_NIGHTLY=1, build the base CI image against PyTorch nightly so the
+# entire existing pipeline runs on nightly torch (CUDA/GPU lane only). Delegate
+# to the dedicated nightly build (PYTORCH_NIGHTLY=1, CUDA 13.0) and tag it at the
+# normal IMAGE_TAG that every test step already pulls -- no separate image tag,
+# no duplicate "vLLM Against PyTorch Nightly" pipeline section.
+if [[ "${TORCH_NIGHTLY:-0}" == "1" ]]; then
+    echo "--- :warning: TORCH_NIGHTLY=1 -- building base image on PyTorch nightly"
+    SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+    exec "${SCRIPT_DIR}/image_build_torch_nightly.sh" \
+        "${REGISTRY}" "${REPO}" "${BUILDKITE_COMMIT}" "${BRANCH}" "${IMAGE_TAG}"
+fi
+
 # build config
 TARGET="test-ci"
 VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"

.buildkite/image_build/image_build_arm64.sh

@@ -19,13 +19,14 @@ if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
   echo "Image found"
 else
   echo "Image not found, proceeding with build..."
-  # build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
+  # build for arm64 GPU targets: Grace/GH200 (sm_90),
+  # Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
   # (sm_121, family-covered by 12.0 under CUDA 13)
   docker build --file docker/Dockerfile \
     --platform linux/arm64 \
     --build-arg max_jobs=16 \
     --build-arg nvcc_threads=4 \
-    --build-arg torch_cuda_arch_list="9.0 12.0" \
+    --build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \
     --build-arg USE_SCCACHE=1 \
     --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
     --tag "$IMAGE" \

.buildkite/intel_jobs/basic_correctness.yaml

@@ -3,7 +3,7 @@ depends_on:
   - image-build-xpu
 steps:
 - label: XPU Sleep Mode
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/lora_intel.yaml

@@ -86,7 +86,7 @@ steps:
       pytest -v -s lora/test_punica_ops.py::test_add_lora_fused_moe_early_exit'
 
 - label: LoRA Punica FP8/XPU Ops
-  timeout_in_minutes: 45
+  timeout_in_minutes: 60
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/misc_intel.yaml

@@ -3,7 +3,7 @@ depends_on:
   - image-build-xpu
 steps:
 - label: V1 Core + KV + Metrics
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
     label: production
@@ -33,7 +33,7 @@ steps:
       pytest -v -s v1/executor'
 
 - label: V1 Sample + Logits
-  timeout_in_minutes: 30
+  timeout_in_minutes: 90
   device: intel_gpu
   agent_tags:
     label: production
@@ -72,9 +72,7 @@ steps:
       pytest -v -s v1/test_oracle.py &&
       pytest -v -s v1/test_request.py &&
       pytest -v -s v1/test_outputs.py &&
-      pytest -v -s v1/sample/test_topk_topp_sampler.py &&
-      pytest -v -s v1/sample/test_logprobs.py &&
-      pytest -v -s v1/sample/test_logprobs_e2e.py'
+      pytest -v -s v1/sample'
 
 - label: Basic Models Tests (Initialization)
   timeout_in_minutes: 60
@@ -154,7 +152,7 @@ steps:
 
 - label: Regression
   key: regression
-  timeout_in_minutes: 30
+  timeout_in_minutes: 50
   device: intel_gpu
   agent_tags:
     label: production
@@ -188,7 +186,7 @@ steps:
 
 - label: Metrics, Tracing (2 GPUs)
   key: metrics-tracing-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   num_devices: 2
   device: intel_gpu
   agent_tags:
@@ -224,7 +222,7 @@ steps:
 
 - label: Async Engine, Inputs, Utils, Worker
   key: async-engine-inputs-utils-worker
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/models_distributed_intel.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
   key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 65
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/models_multimodal_intel.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
   key: multi-modal-models-standard-1-qwen2
-  timeout_in_minutes: 45
+  timeout_in_minutes: 70
   device: intel_gpu
   agent_tags:
     label: production
@@ -29,7 +29,7 @@ steps:
 
 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
   key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 45
+  timeout_in_minutes: 70
   device: intel_gpu
   agent_tags:
     label: production
@@ -52,7 +52,7 @@ steps:
 
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
   key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 45
+  timeout_in_minutes: 65
   device: intel_gpu
   agent_tags:
     label: production
@@ -100,7 +100,7 @@ steps:
 
 - label: Multi-Modal Processor # 44min
   key: multi-modal-processor
-  timeout_in_minutes: 45
+  timeout_in_minutes: 60
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/test-intel.yaml

@@ -17,7 +17,7 @@ steps:
   - label: "XPU example Test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 30
+    timeout_in_minutes: 50
     device: intel_gpu
     agent_tags:
       label: production
@@ -76,7 +76,7 @@ steps:
   - label: "XPU V1 test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 30
+    timeout_in_minutes: 70
     device: intel_gpu
     agent_tags:
       label: production
@@ -99,12 +99,13 @@ steps:
         pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py &&
         pytest -v -s v1/structured_output &&
         pytest -v -s v1/test_serial_utils.py &&
+        pytest -v -s v1/e2e/general/test_correctness_sliding_window.py --deselect="tests/v1/e2e/general/test_correctness_sliding_window.py::test_sliding_window_retrieval[True-1-5-google/gemma-3-1b-it]" &&
         pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py --ignore=v1/spec_decode/test_speculators_correctness.py &&
         pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py'
   - label: "XPU server test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 30
+    timeout_in_minutes: 45
     device: intel_gpu
     agent_tags:
       label: production

.buildkite/lm-eval-harness/configs/Meta-Llama-4-Maverick-17B-128E-Instruct-FP8.yaml

@@ -1,6 +1,7 @@
 # For hf script, without -t option (tensor parallel size).
 # bash .buildkite/lm-eval-harness/run-lm-eval-mmlupro-vllm-baseline.sh -m meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8 -l 250 -t 8 -f 5
 model_name: "meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8"
+rocm_safetensors_load_strategy: lazy
 required_gpu_arch:
   - gfx942
   - gfx950

.buildkite/lm-eval-harness/test_lm_eval_correctness.py

@@ -72,6 +72,11 @@ def launch_lm_eval(eval_config, tp_size):
     if moe_backend is not None:
         model_args += f"moe_backend={moe_backend},"
 
+    if current_platform.is_rocm():
+        rocm_load_strategy = eval_config.get("rocm_safetensors_load_strategy")
+        if rocm_load_strategy is not None:
+            model_args += f"safetensors_load_strategy={rocm_load_strategy},"
+
     env_vars = eval_config.get("env_vars", None)
     with scoped_env_vars(env_vars):
         results = lm_eval.simple_evaluate(

.buildkite/scripts/annotate-build-artifact.sh

@@ -3,7 +3,8 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 #
 # Append a build artifact line to the Buildkite annotation.
-# Usage: annotate-build-artifact.sh <label> <value>
+# Usage: annotate-build-artifact.sh <label> <value> <context>
 set -e
 echo "- **${1}**: \`${2}\`" | \
-  buildkite-agent annotate --append --style 'info' --context 'release-artifacts'
+  buildkite-agent annotate --append --style 'info' \
+    --context "${3:?context is required}"

.buildkite/scripts/build-macos-wheel.sh

@@ -0,0 +1,32 @@
+#!/usr/bin/env bash
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+#
+# Build the macOS arm64 CPU wheel natively on a macOS agent (the `macmini`
+# queue) into artifacts/dist/ for upload-nightly-wheels.sh.
+
+set -euo pipefail
+
+# The Rust frontend build needs protoc.
+if ! command -v protoc >/dev/null 2>&1; then
+  brew install protobuf
+fi
+
+# upload-nightly-wheels.sh expects exactly one wheel.
+rm -rf artifacts/dist
+mkdir -p artifacts/dist
+
+export VLLM_TARGET_DEVICE=cpu
+export VLLM_REQUIRE_RUST_FRONTEND=1
+export MACOSX_DEPLOYMENT_TARGET=11.0
+# uv's CPython is universal2; force an arm64-only build and tag so the wheel
+# isn't mislabelled universal2 and installed on Intel Macs where import fails.
+export ARCHFLAGS="-arch arm64"
+export _PYTHON_HOST_PLATFORM="macosx-11.0-arm64"
+export CMAKE_BUILD_PARALLEL_LEVEL="${CMAKE_BUILD_PARALLEL_LEVEL:-4}"
+
+uv venv --python 3.12
+uv pip install -r requirements/build/cpu.txt --index-strategy unsafe-best-match
+uv build --wheel --no-build-isolation -o artifacts/dist
+
+ls -l artifacts/dist/*.whl

.buildkite/scripts/check-ray-compatibility.sh

@@ -29,7 +29,11 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
         TORCH_INDEX_URL=""
     fi
 else
-    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
+        TORCH_INDEX_URL="https://download.pytorch.org/whl/nightly/cu130"
+    else
+        TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    fi
 fi
 echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"
 

.buildkite/scripts/ci-bake-rocm.sh

@@ -15,9 +15,9 @@ set -euo pipefail
 
 DEFAULT_REPO_SLUG="vllm-project/vllm"
 DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
-DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
+DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tools/install_protoc.sh rust-toolchain.toml tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
 DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
-DEFAULT_CI_BASE_DOCKERFILE_STAGES="base build_rixl build_rocshmem build_deepep mori_base ci_base"
+DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_rixl build_rocshmem build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="1"
 IMAGE_EXISTED_BEFORE_BUILD=0
 
@@ -285,7 +285,7 @@ get_content_arg_names() {
     fi | awk 'NF && !seen[$0]++'
 }
 
-compute_ci_base_content_hash() {
+compute_ci_base_content_hash_once() {
     local -a content_paths=()
     local -a content_args=()
     local dockerfile="${CI_BASE_DOCKERFILE:-}"
@@ -301,7 +301,8 @@ compute_ci_base_content_hash() {
         if [[ -n "${dockerfile}" ]]; then
             printf 'dockerfile:%s\n' "${dockerfile}"
             printf 'resolved-build-args:\n'
-            hash_dockerfile_arg_values "${dockerfile}" "${content_args[@]}"
+            hash_dockerfile_arg_values "${dockerfile}" "${content_args[@]}" \
+                || return 1
             if [[ -n "${stages}" ]]; then
                 printf 'dockerfile-stages:%s\n' "${stages}"
                 if [[ -f "${dockerfile}" ]]; then
@@ -314,6 +315,53 @@ compute_ci_base_content_hash() {
     } | sha256sum | cut -d' ' -f1
 }
 
+compute_ci_base_content_hash() {
+    local attempts="${CI_BASE_HASH_ATTEMPTS:-3}"
+    local delay_secs="${CI_BASE_HASH_RETRY_DELAY:-5}"
+    local attempt=0
+    local hash=""
+    local failed=0
+    local -a hashes=()
+
+    if [[ ! "${attempts}" =~ ^[1-9][0-9]*$ ]]; then
+        echo "Invalid CI_BASE_HASH_ATTEMPTS: ${attempts}" >&2
+        return 1
+    fi
+    if [[ ! "${delay_secs}" =~ ^[0-9]+$ ]]; then
+        echo "Invalid CI_BASE_HASH_RETRY_DELAY: ${delay_secs}" >&2
+        return 1
+    fi
+
+    for ((attempt = 1; attempt <= attempts; attempt++)); do
+        if ! hash=$(compute_ci_base_content_hash_once); then
+            echo "ci_base content hash calculation ${attempt}/${attempts} failed" >&2
+            failed=1
+        else
+            hashes+=("${hash}")
+            echo "ci_base content hash calculation ${attempt}/${attempts}: ${hash}" >&2
+        fi
+
+        if ((attempt < attempts)); then
+            sleep "${delay_secs}"
+        fi
+    done
+
+    if ((failed)) || ((${#hashes[@]} != attempts)); then
+        echo "Could not calculate a reliable ci_base content hash" >&2
+        return 1
+    fi
+
+    for hash in "${hashes[@]:1}"; do
+        if [[ "${hash}" != "${hashes[0]}" ]]; then
+            echo "ci_base content hash changed between calculations" >&2
+            printf '  observed: %s\n' "${hashes[@]}" >&2
+            return 1
+        fi
+    done
+
+    printf '%s\n' "${hashes[0]}"
+}
+
 extract_dockerfile_arg_default() {
     local dockerfile="$1"
     local arg_name="$2"
@@ -366,7 +414,11 @@ hash_dockerfile_arg_values() {
         printf 'arg:%s=%s\n' "${arg_name}" "${arg_value:-<empty>}"
         if [[ "${arg_name}" == "BASE_IMAGE" && -n "${arg_value}" ]]; then
             digest=$(resolve_image_digest "${arg_value}")
-            printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest:-unknown}"
+            if [[ -z "${digest}" ]]; then
+                echo "Failed to resolve digest for BASE_IMAGE=${arg_value}" >&2
+                return 1
+            fi
+            printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest}"
         fi
     done
 }
@@ -764,7 +816,7 @@ configure_ci_base_image_refs() {
         fi
         set_buildkite_metadata "rocm-ci-base-image" "${CI_BASE_IMAGE_TAG}"
         set_buildkite_metadata "rocm-ci-base-image-content" "${content_tag}"
-        set_buildkite_metadata "rocm-ci-base-image-commit" "${CI_BASE_IMAGE_TAG_COMMIT:-}"
+        set_buildkite_metadata "rocm-ci-base-image-commit" "${CI_BASE_IMAGE_TAG_COMMIT_REF:-}"
         set_buildkite_metadata "rocm-ci-base-image-stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
         return 0
     fi
@@ -1211,12 +1263,24 @@ uses_rocm_csrc_cache() {
     esac
 }
 
+uses_rocm_rust_cache() {
+    case "${TARGET}" in
+        rust-rocm-ci|test-rocm-ci|test-rocm-ci-with-wheel|test-rocm-ci-with-artifacts|export-wheel-rocm)
+            return 0
+            ;;
+        *)
+            return 1
+            ;;
+    esac
+}
+
 compute_rocm_csrc_content_hash() {
     local bake_dir=""
     local dockerfile_rocm=""
     local -a content_paths=(
         "requirements/common.txt"
         "requirements/rocm.txt"
+        "pyproject.toml"
         "setup.py"
         "CMakeLists.txt"
         "cmake"
@@ -1260,6 +1324,56 @@ compute_rocm_csrc_content_hash_if_needed() {
     echo "ROCm csrc content cache ref: ${ROCM_CSRC_CONTENT_CACHE_REF}"
 }
 
+compute_rocm_rust_content_hash() {
+    local bake_dir=""
+    local dockerfile_rocm=""
+    local -a content_paths=(
+        "requirements/build/rust.txt"
+        "rust/Cargo.lock"
+        "rust/Cargo.toml"
+        "rust/proto"
+        "rust/src"
+        "rust-toolchain.toml"
+        "tools/build_rust.py"
+        "tools/install_protoc.sh"
+        "build_rust.sh"
+    )
+    local -a content_args=()
+
+    bake_dir=$(dirname "${VLLM_BAKE_FILE}")
+    dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
+    mapfile -t content_args < <(
+        get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
+    )
+
+    {
+        printf 'rust-input-files-hash:%s\n' "$(compute_content_hash "${content_paths[@]}")"
+        printf 'dockerfile:%s\n' "${dockerfile_rocm}"
+        printf 'resolved-build-args:\n'
+        hash_dockerfile_arg_values "${dockerfile_rocm}" "${content_args[@]}"
+        printf 'dockerfile-stages:base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build\n'
+        if [[ -f "${dockerfile_rocm}" ]]; then
+            hash_dockerfile_stages "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build"
+        else
+            printf 'missing:%s\n' "${dockerfile_rocm}"
+        fi
+    } | sha256sum | cut -d' ' -f1
+}
+
+compute_rocm_rust_content_hash_if_needed() {
+    local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
+
+    if [[ "${ROCM_RUST_CONTENT_CACHE:-1}" == "0" ]] || ! uses_rocm_rust_cache; then
+        return 0
+    fi
+
+    ROCM_RUST_CONTENT_HASH=$(compute_rocm_rust_content_hash)
+    ROCM_RUST_CONTENT_CACHE_REF="${cache_repo}:rust-rocm-input-${ROCM_RUST_CONTENT_HASH}"
+    export ROCM_RUST_CONTENT_HASH
+    export ROCM_RUST_CONTENT_CACHE_REF
+    echo "ROCm Rust content cache ref: ${ROCM_RUST_CONTENT_CACHE_REF}"
+}
+
 write_hcl_string_list_entries() {
     local indent="$1"
     local value=""
@@ -1317,6 +1431,7 @@ write_rocm_build_arg_override() {
                 "${CI_BASE_DOCKERFILE_STAGES:-${DEFAULT_CI_BASE_DOCKERFILE_STAGES}}" \
                 "${CI_BASE_CONTENT_ARGS:-}"
             get_content_arg_names "${dockerfile_rocm}" "base csrc-build" "${ROCM_CSRC_CONTENT_ARGS:-}"
+            get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
         } | awk 'NF && !seen[$0]++'
     )
 
@@ -1365,46 +1480,133 @@ validate_cache_export_mode() {
     esac
 }
 
+validate_content_cache_export_mode() {
+    local mode="$1"
+    local env_name="$2"
+
+    case "${mode}" in
+        missing|always|never)
+            ;;
+        *)
+            echo "Error: ${env_name} must be one of: missing, always, never"
+            exit 1
+            ;;
+    esac
+}
+
+should_export_content_cache_ref() {
+    local cache_ref="$1"
+    local cache_name="$2"
+    local mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
+
+    case "${mode}" in
+        always)
+            echo "${cache_name} content cache export mode is always; exporting ${cache_ref}"
+            return 0
+            ;;
+        never)
+            echo "${cache_name} content cache export mode is never; not exporting ${cache_ref}"
+            return 1
+            ;;
+        missing|"")
+            if docker buildx imagetools inspect "${cache_ref}" >/dev/null 2>&1; then
+                echo "${cache_name} content cache exists; not re-exporting ${cache_ref}"
+                return 1
+            fi
+            echo "${cache_name} content cache missing; will export ${cache_ref}"
+            return 0
+            ;;
+        *)
+            echo "Error: ROCM_CONTENT_CACHE_EXPORT_MODE must be one of: missing, always, never"
+            exit 1
+            ;;
+    esac
+}
+
 write_rocm_cache_override() {
     local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
+    local content_cache_export_mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
     local csrc_cache_to_mode="${ROCM_CSRC_CACHE_TO_MODE:-max}"
+    local rust_cache_to_mode="${ROCM_RUST_CACHE_TO_MODE:-max}"
     local rocm_cache_to_mode="${ROCM_FINAL_CACHE_TO_MODE:-min}"
-    local -a content_cache_from=()
+    local -a csrc_content_cache_from=()
+    local -a rust_content_cache_from=()
+    local -a combined_content_cache_from=()
     local -a csrc_cache_to=()
+    local -a rust_cache_to=()
     local -a rocm_cache_to=()
     local -a export_wheel_cache_to=()
+    local export_csrc_cache=1
+    local export_rust_cache=1
 
-    if ! uses_rocm_csrc_cache; then
+    if ! uses_rocm_csrc_cache && ! uses_rocm_rust_cache; then
         return 0
     fi
 
+    validate_content_cache_export_mode \
+        "${content_cache_export_mode}" \
+        "ROCM_CONTENT_CACHE_EXPORT_MODE"
     validate_cache_export_mode "${csrc_cache_to_mode}" "ROCM_CSRC_CACHE_TO_MODE"
+    validate_cache_export_mode "${rust_cache_to_mode}" "ROCM_RUST_CACHE_TO_MODE"
     validate_cache_export_mode "${rocm_cache_to_mode}" "ROCM_FINAL_CACHE_TO_MODE"
+    echo "ROCm content cache export mode: ${content_cache_export_mode}"
     echo "ROCm csrc cache export mode: ${csrc_cache_to_mode}"
+    echo "ROCm Rust cache export mode: ${rust_cache_to_mode}"
     echo "ROCm final image cache export mode: ${rocm_cache_to_mode}"
 
     if [[ -n "${ROCM_CSRC_CONTENT_CACHE_REF:-}" ]]; then
-        content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
-        csrc_cache_to+=(
-            "type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
-        )
+        csrc_content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
+        if should_export_content_cache_ref "${ROCM_CSRC_CONTENT_CACHE_REF}" "ROCm csrc"; then
+            csrc_cache_to+=(
+                "type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
+            )
+        else
+            export_csrc_cache=0
+        fi
     fi
 
+    if [[ -n "${ROCM_RUST_CONTENT_CACHE_REF:-}" ]]; then
+        rust_content_cache_from+=("type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF}")
+        if should_export_content_cache_ref "${ROCM_RUST_CONTENT_CACHE_REF}" "ROCm Rust"; then
+            rust_cache_to+=(
+                "type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF},mode=${rust_cache_to_mode},ignore-error=true"
+            )
+        else
+            export_rust_cache=0
+        fi
+    fi
+
+    combined_content_cache_from=("${csrc_content_cache_from[@]}" "${rust_content_cache_from[@]}")
+
     # Docker Hub cache exports are best-effort. A cache-only target failure can
     # otherwise cancel the sibling image target before its manifest is pushed.
     if [[ -n "${BUILDKITE_COMMIT:-}" ]]; then
-        csrc_cache_to+=(
-            "type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
-        )
+        if [[ ${export_csrc_cache} -eq 1 ]]; then
+            csrc_cache_to+=(
+                "type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
+            )
+        fi
+        if [[ ${export_rust_cache} -eq 1 ]]; then
+            rust_cache_to+=(
+                "type=registry,ref=${cache_repo}:rust-rocm-${BUILDKITE_COMMIT},mode=${rust_cache_to_mode},ignore-error=true"
+            )
+        fi
         rocm_cache_to+=(
             "type=registry,ref=${cache_repo}:rocm-${BUILDKITE_COMMIT},mode=${rocm_cache_to_mode},ignore-error=true"
         )
     fi
 
     if [[ -n "${ROCM_CACHE_BRANCH_TAG:-}" ]]; then
-        csrc_cache_to+=(
-            "type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
-        )
+        if [[ ${export_csrc_cache} -eq 1 ]]; then
+            csrc_cache_to+=(
+                "type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
+            )
+        fi
+        if [[ ${export_rust_cache} -eq 1 ]]; then
+            rust_cache_to+=(
+                "type=registry,ref=${cache_repo}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rust_cache_to_mode},ignore-error=true"
+            )
+        fi
         rocm_cache_to+=(
             "type=registry,ref=${cache_repo}:rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rocm_cache_to_mode},ignore-error=true"
         )
@@ -1422,7 +1624,7 @@ target "csrc-rocm-ci" {
   cache-from = concat(
     get_cache_from_rocm_csrc(),
 EOF
-        write_hcl_string_list "    " "${content_cache_from[@]}"
+        write_hcl_string_list "    " "${csrc_content_cache_from[@]}"
         cat <<EOF
   )
 EOF
@@ -1430,11 +1632,23 @@ EOF
         cat <<EOF
 }
 
+target "rust-rocm-ci" {
+  cache-from = concat(
+    get_cache_from_rocm_rust(),
+EOF
+        write_hcl_string_list "    " "${rust_content_cache_from[@]}"
+        cat <<EOF
+  )
+EOF
+        write_hcl_string_list_attr "  " "cache-to" "${rust_cache_to[@]}"
+        cat <<EOF
+}
+
 target "test-rocm-ci" {
   cache-from = concat(
     get_cache_from_rocm(),
 EOF
-        write_hcl_string_list "    " "${content_cache_from[@]}"
+        write_hcl_string_list "    " "${combined_content_cache_from[@]}"
         cat <<EOF
   )
 EOF
@@ -1446,7 +1660,7 @@ target "export-wheel-rocm" {
   cache-from = concat(
     get_cache_from_rocm(),
 EOF
-        write_hcl_string_list "    " "${content_cache_from[@]}"
+        write_hcl_string_list "    " "${combined_content_cache_from[@]}"
         cat <<EOF
   )
 EOF
@@ -1736,8 +1950,8 @@ confirm_remote_image_push() {
         fi
 
         if [[ -z "${remote_revision}" \
-              && ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 \
-              && image_tag_is_commit_scoped ]]; then
+              && ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 ]] \
+            && image_tag_is_commit_scoped; then
             echo "Remote image exists under a commit-scoped tag; accepting push despite missing revision label."
             return 0
         fi
@@ -1867,36 +2081,57 @@ upload_wheel_artifacts_if_present() {
     local wheel_dir="./wheel-export"
     local artifact_dir="artifacts/vllm-rocm-install"
     local archive_name="vllm-rocm-install.tar.gz"
+    local metadata_dir="${wheel_dir}/.vllm-ci-artifact"
+    local native_base_image=""
     local whl=""
     local whl_name=""
+    local -a wheels=()
 
     if ! should_upload_wheel_artifacts; then
         return 0
     fi
 
-    if [[ ! -d "${wheel_dir}" ]] || ! ls "${wheel_dir}"/*.whl >/dev/null 2>&1; then
-        echo "No ROCm wheel artifacts found in ${wheel_dir}"
-        return 0
+    if [[ -d "${wheel_dir}" ]]; then
+        mapfile -t wheels < <(find "${wheel_dir}" -maxdepth 1 -type f -name '*.whl' -print)
+    fi
+    if [[ ${#wheels[@]} -ne 1 ]]; then
+        echo "Expected exactly one ROCm wheel in ${wheel_dir}; found ${#wheels[@]}" >&2
+        return 1
+    fi
+    whl="${wheels[0]}"
+    whl_name=$(basename "${whl}")
+    native_base_image="${CI_BASE_IMAGE_TAG_COMMIT_REF:-${CI_BASE_IMAGE:-}}"
+    if [[ -z "${native_base_image}" ]]; then
+        echo "Native ROCm artifact requires a ci_base image reference" >&2
+        return 1
     fi
 
     echo "--- :package: Uploading ROCm vLLM install artifact"
-    mkdir -p "${artifact_dir}"
+    rm -rf "${artifact_dir}" "${metadata_dir}"
+    mkdir -p "${artifact_dir}" "${metadata_dir}"
+
+    printf '%s\n' "${BUILDKITE_COMMIT:-local}" > "${metadata_dir}/commit.txt"
+    printf '%s\n' "${native_base_image}" > "${metadata_dir}/native-base-image.txt"
+    printf '%s\n' "${CI_BASE_IMAGE:-}" > "${metadata_dir}/ci-base-image.txt"
+    printf '%s\n' "${IMAGE_TAG:-}" > "${metadata_dir}/fallback-image.txt"
+    printf '%s\n' "${whl_name}" > "${metadata_dir}/wheel-filename.txt"
 
     tar -C "${wheel_dir}" -czf "${artifact_dir}/${archive_name}" .
+    (
+        cd "${artifact_dir}"
+        sha256sum "${archive_name}" > "${archive_name}.sha256"
+    )
     echo "Created ${archive_name}: $(du -sh "${artifact_dir}/${archive_name}" | cut -f1)"
-    printf '%s\n' "${CI_BASE_IMAGE:-}" > "${artifact_dir}/ci-base-image.txt"
-    printf '%s\n' "${IMAGE_TAG:-}" > "${artifact_dir}/fallback-image.txt"
-
-    for whl in "${wheel_dir}"/*.whl; do
-        [[ -f "${whl}" ]] || continue
-        whl_name=$(basename "${whl}")
-        cp "${whl}" "${artifact_dir}/${whl_name}"
-        echo "Copied ${whl_name}: $(du -sh "${artifact_dir}/${whl_name}" | cut -f1)"
-    done
+    cp "${metadata_dir}"/*.txt "${artifact_dir}/"
+    cp "${whl}" "${artifact_dir}/${whl_name}"
+    echo "Copied ${whl_name}: $(du -sh "${artifact_dir}/${whl_name}" | cut -f1)"
 
     if command -v buildkite-agent >/dev/null 2>&1; then
-        buildkite-agent artifact upload "${artifact_dir}/*"
+        buildkite-agent artifact upload "${artifact_dir}/*" || return 1
         echo "ROCm vLLM install artifacts uploaded to ${artifact_dir}/"
+    elif [[ "${BUILDKITE:-false}" == "true" ]]; then
+        echo "buildkite-agent not found; cannot upload required ROCm artifacts" >&2
+        return 1
     else
         echo "Not in Buildkite, skipping artifact upload"
     fi
@@ -1920,6 +2155,7 @@ main() {
     compute_dependency_cache_keys
     write_ci_base_label_override
     compute_rocm_csrc_content_hash_if_needed
+    compute_rocm_rust_content_hash_if_needed
     write_rocm_cache_override
     resolve_ci_base_dependency_targets
     print_bake_config
@@ -1927,6 +2163,11 @@ main() {
         echo "BAKE_PRINT_ONLY=1 set; skipping build"
         return 0
     fi
+    if should_upload_wheel_artifacts; then
+        # wheel-export is an output directory, not a BuildKit cache. Starting
+        # clean prevents a failed/retried export from packaging a stale wheel.
+        rm -rf ./wheel-export
+    fi
     seed_dependency_caches_if_needed
     run_bake
     upload_wheel_artifacts_if_present

.buildkite/scripts/generate-and-upload-nightly-index.sh

@@ -45,8 +45,10 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py --version "$SUBPATH" --curr
 echo "Uploading indices to $S3_COMMIT_PREFIX"
 aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
 
-# copy to /nightly/ only if it is on the main branch and not a PR
-if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
+# copy to /nightly/ only when enabled for a main branch build that is not a PR
+if [[ "${UPDATE_NIGHTLY_INDEX:-1}" == "1" && \
+      "$BUILDKITE_BRANCH" == "main" && \
+      "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
     echo "Uploading indices to overwrite /nightly/"
     aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/nightly/"
 fi
@@ -67,7 +69,7 @@ pure_version="${version%%+*}"
 echo "Pure version (without variant): $pure_version"
 
 # re-generate and copy to /<pure_version>/ only if it does not have "dev" in the version
-if [[ "$version" != *"dev"* ]]; then
+if [[ "${UPDATE_VERSION_INDEX:-1}" == "1" && "$version" != *"dev"* ]]; then
     echo "Re-generating indices for /$pure_version/"
     rm -rf "${INDICES_OUTPUT_DIR:?}"
     mkdir -p "$INDICES_OUTPUT_DIR"

.buildkite/scripts/hardware_ci/run-amd-test.sh

@@ -1,7 +1,7 @@
 #!/bin/bash
 
-# This script runs tests inside the corresponding ROCm docker container.
-# It handles both single-node and multi-node test configurations.
+# This script runs ROCm tests either directly in a native CI pod or inside the
+# corresponding Docker container. Multi-node tests continue to use Docker.
 #
 # Multi-node detection: Instead of matching on fragile group names, we detect
 # multi-node jobs structurally by looking for the bracket command syntax
@@ -34,10 +34,27 @@ set -o pipefail
 : "${CLICOLOR_FORCE:=1}"
 : "${PY_COLORS:=1}"
 : "${ROCM_DOCKER_TTY:=1}"
+: "${PYTHONFAULTHANDLER:=1}"
+: "${PYTEST_TIMEOUT:=2100}"
 if [[ " ${PYTEST_ADDOPTS:-} " != *" --color"* ]]; then
   PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--color=yes"
 fi
-export BUILDKIT_PROGRESS TERM FORCE_COLOR CLICOLOR_FORCE PY_COLORS PYTEST_ADDOPTS ROCM_DOCKER_TTY
+if [[ " ${PYTEST_ADDOPTS:-} " != *" --durations="* ]]; then
+  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--durations=25"
+fi
+if [[ " ${PYTEST_ADDOPTS:-} " != *" --durations-min="* ]]; then
+  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--durations-min=1.0"
+fi
+# Dump stacks after 15 minutes, then stop an individual test after 35 minutes.
+if [[ " ${PYTEST_ADDOPTS:-} " != *" faulthandler_timeout="* ]]; then
+  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }-o faulthandler_timeout=900"
+fi
+if [[ " ${PYTEST_ADDOPTS:-} " != *" --timeout-method="* &&
+  " ${PYTEST_ADDOPTS:-} " != *" --timeout-method "* ]]; then
+  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--timeout-method=thread"
+fi
+export BUILDKIT_PROGRESS TERM FORCE_COLOR CLICOLOR_FORCE PY_COLORS PYTEST_ADDOPTS PYTEST_TIMEOUT ROCM_DOCKER_TTY
+export PYTHONFAULTHANDLER
 
 # Export Python path for commands that run directly on the host. Containerized
 # tests set this to /vllm-workspace below so spawned Python processes do not
@@ -53,6 +70,28 @@ report_docker_usage() {
   docker system df || true
 }
 
+clear_ci_orchestration_env() {
+  unset -v \
+    VLLM_TEST_GROUP_NAME \
+    VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE \
+    VLLM_CI_ARTIFACT_STEP \
+    VLLM_TEST_CACHE \
+    VLLM_CI_EXECUTION_MODE \
+    VLLM_CI_WORKSPACE \
+    VLLM_CI_REQUIRE_WORKSPACE_MOUNT \
+    VLLM_TEST_COMMANDS \
+    VLLM_CI_BRANCH \
+    VLLM_CI_BASE_IMAGE \
+    VLLM_CI_FALLBACK_IMAGE \
+    VLLM_CI_DOCKER_DISABLED \
+    VLLM_CI_ARTIFACT_GLOB \
+    VLLM_CI_ARTIFACT_CHECKSUM_GLOB \
+    VLLM_CI_EXPECTED_GPU_COUNT \
+    VLLM_CI_USE_ARTIFACTS \
+    VLLM_CI_RESULTS_ROOT \
+    VLLM_ALLOW_DEPRECATED_BEAM_SEARCH
+}
+
 cleanup_network() {
   local max_nodes=${NUM_NODES:-2}
   for node in $(seq 0 $((max_nodes - 1))); do
@@ -145,7 +184,11 @@ prepare_artifact_image() {
   fi
 
   cp "${wheel_dir}"/*.whl "${context_dir}/wheels/" || return 1
-  tar -C "${wheel_dir}" --exclude='*.whl' -cf - . \
+  tar -C "${wheel_dir}" \
+    --exclude='*.whl' \
+    --exclude='.vllm-ci-artifact' \
+    --exclude='./.vllm-ci-artifact' \
+    -cf - . \
     | tar -C "${workspace_dir}" -xf - || return 1
   cat > "${context_dir}/Dockerfile" <<'EOF'
 ARG BASE_IMAGE
@@ -168,6 +211,257 @@ EOF
   return 0
 }
 
+is_native_runtime() {
+  [[ "${AMD_CI_RUNTIME:-}" == "native" || "${NATIVE_CI:-}" == "true" ]]
+}
+
+validate_native_workspace() {
+  local workspace_dir="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
+  local workspace_real=""
+  local checkout_real=""
+  local workspace_mount=""
+
+  mkdir -p "${workspace_dir}" || return 1
+  workspace_real=$(readlink -m "${workspace_dir}") || return 1
+  if [[ -n "${BUILDKITE_BUILD_CHECKOUT_PATH:-}" ]]; then
+    checkout_real=$(readlink -m "${BUILDKITE_BUILD_CHECKOUT_PATH}") || return 1
+    if [[ "${checkout_real}" == "${workspace_real}" \
+      || "${checkout_real}" == "${workspace_real}/"* \
+      || "${workspace_real}" == "${checkout_real}/"* ]]; then
+      echo "Refusing to replace ${workspace_real}; it overlaps the Buildkite checkout ${checkout_real}" >&2
+      return 1
+    fi
+  fi
+  if [[ "${VLLM_CI_REQUIRE_WORKSPACE_MOUNT:-1}" == "1" ]]; then
+    if ! command -v findmnt >/dev/null 2>&1; then
+      echo "findmnt is required to verify the native workspace mount" >&2
+      return 1
+    fi
+    workspace_mount=$(findmnt -n -T "${workspace_real}" -o TARGET 2>/dev/null || true)
+    if [[ "$(readlink -m "${workspace_mount:-/}")" != "${workspace_real}" ]]; then
+      echo "Native CI requires a dedicated volume mounted at ${workspace_real}" >&2
+      return 1
+    fi
+  fi
+}
+
+prepare_native_workspace() {
+  if [[ "${VLLM_CI_USE_ARTIFACTS:-0}" != "1" ]]; then
+    echo "Native CI requires VLLM_CI_USE_ARTIFACTS=1"
+    return 1
+  fi
+  if ! command -v buildkite-agent >/dev/null 2>&1; then
+    echo "buildkite-agent not found; cannot download ROCm wheel artifact"
+    return 1
+  fi
+  validate_native_workspace || return 1
+
+  local artifact_glob="${VLLM_CI_ARTIFACT_GLOB:-artifacts/vllm-rocm-install/vllm-rocm-install.tar.gz}"
+  local artifact_checksum_glob="${VLLM_CI_ARTIFACT_CHECKSUM_GLOB:-${artifact_glob}.sha256}"
+  local artifact_step="${VLLM_CI_ARTIFACT_STEP:-image-build-amd}"
+  local archive=""
+  local checksum=""
+  local download_dir=""
+  local metadata_dir=""
+  local recorded_base=""
+  local recorded_commit=""
+  local recorded_wheel=""
+  local workspace_dir="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
+  local wheel_dir=""
+  local attempt=0
+  local attempt_dir=""
+  local -a archives=()
+  local -a checksums=()
+  local -a wheels=()
+
+  artifact_work_dir=$(mktemp -d -t vllm-rocm-artifact.XXXXXX) || return 1
+  wheel_dir="${artifact_work_dir}/wheels"
+  mkdir -p "${wheel_dir}" || return 1
+
+  echo "--- Downloading ROCm wheel artifact from ${artifact_step} (native in-pod)"
+  for attempt in 1 2 3; do
+    attempt_dir="${artifact_work_dir}/download-${attempt}"
+    rm -rf "${attempt_dir}" || return 1
+    mkdir -p "${attempt_dir}" || return 1
+    if buildkite-agent artifact download \
+      "${artifact_glob}" "${attempt_dir}" --step "${artifact_step}" \
+      && buildkite-agent artifact download \
+        "${artifact_checksum_glob}" "${attempt_dir}" --step "${artifact_step}"; then
+      download_dir="${attempt_dir}"
+      break
+    fi
+    echo "Artifact download attempt ${attempt}/3 failed"
+    if [[ "${attempt}" -lt 3 ]]; then
+      sleep $((attempt * 2))
+    fi
+  done
+  if [[ -z "${download_dir}" ]]; then
+    echo "Failed to download ${artifact_glob} and ${artifact_checksum_glob} from ${artifact_step}"
+    return 1
+  fi
+
+  mapfile -t archives < <(
+    find "${download_dir}" -name "vllm-rocm-install.tar.gz" -type f -print
+  )
+  mapfile -t checksums < <(
+    find "${download_dir}" -name "vllm-rocm-install.tar.gz.sha256" -type f -print
+  )
+  if [[ ${#archives[@]} -ne 1 || ${#checksums[@]} -ne 1 ]]; then
+    echo "Expected exactly one ROCm archive and checksum; found ${#archives[@]} archive(s) and ${#checksums[@]} checksum(s)" >&2
+    return 1
+  fi
+  archive="${archives[0]}"
+  checksum="${checksums[0]}"
+  if [[ "$(dirname "${archive}")" != "$(dirname "${checksum}")" ]]; then
+    echo "ROCm archive and checksum were downloaded to different directories" >&2
+    return 1
+  fi
+  (
+    cd "$(dirname "${archive}")"
+    sha256sum -c "$(basename "${checksum}")"
+  ) || return 1
+
+  tar --no-same-owner -xzf "${archive}" -C "${wheel_dir}" || return 1
+  mapfile -t wheels < <(
+    find "${wheel_dir}" -maxdepth 1 -type f -name '*.whl' -print
+  )
+  if [[ ${#wheels[@]} -ne 1 ]]; then
+    echo "ROCm artifact must contain exactly one top-level wheel; found ${#wheels[@]}" >&2
+    return 1
+  fi
+  metadata_dir="${wheel_dir}/.vllm-ci-artifact"
+  for metadata_file in commit.txt native-base-image.txt wheel-filename.txt; do
+    if [[ ! -s "${metadata_dir}/${metadata_file}" ]]; then
+      echo "ROCm artifact metadata is missing ${metadata_file}" >&2
+      return 1
+    fi
+  done
+  for metadata_file in ci-base-image.txt fallback-image.txt; do
+    if [[ ! -f "${metadata_dir}/${metadata_file}" ]]; then
+      echo "ROCm artifact metadata is missing ${metadata_file}" >&2
+      return 1
+    fi
+  done
+
+  recorded_commit=$(tr -d '\r\n' < "${metadata_dir}/commit.txt")
+  recorded_base=$(tr -d '\r\n' < "${metadata_dir}/native-base-image.txt")
+  recorded_wheel=$(tr -d '\r\n' < "${metadata_dir}/wheel-filename.txt")
+  if [[ -z "${BUILDKITE_COMMIT:-}" || "${recorded_commit}" != "${BUILDKITE_COMMIT}" ]]; then
+    echo "ROCm artifact commit ${recorded_commit} does not match ${BUILDKITE_COMMIT:-unset}" >&2
+    return 1
+  fi
+  if [[ -z "${VLLM_CI_BASE_IMAGE:-}" || "${recorded_base}" != "${VLLM_CI_BASE_IMAGE}" ]]; then
+    echo "ROCm artifact base ${recorded_base} does not match ${VLLM_CI_BASE_IMAGE:-unset}" >&2
+    return 1
+  fi
+  if [[ "${recorded_wheel}" != "$(basename "${wheels[0]}")" ]]; then
+    echo "ROCm artifact wheel manifest ${recorded_wheel} does not match $(basename "${wheels[0]}")" >&2
+    return 1
+  fi
+  for required_dir in tests .buildkite requirements; do
+    if [[ ! -d "${wheel_dir}/${required_dir}" ]]; then
+      echo "ROCm wheel artifact did not contain ${required_dir}/" >&2
+      return 1
+    fi
+  done
+
+  echo "--- Installing ROCm wheel into pod environment"
+  python3 -m pip install --no-deps --force-reinstall "${wheels[0]}" || return 1
+
+  echo "--- Preparing ${workspace_dir} from artifact"
+  find "${workspace_dir}" -mindepth 1 -maxdepth 1 -exec rm -rf -- {} + || return 1
+  tar -C "${wheel_dir}" \
+    --exclude='*.whl' \
+    --exclude='.vllm-ci-artifact' \
+    --exclude='./.vllm-ci-artifact' \
+    -cf - . | tar --no-same-owner -C "${workspace_dir}" -xf - || return 1
+  if [[ ! -d "${workspace_dir}/tests" ]]; then
+    echo "Failed to stage the native test workspace" >&2
+    return 1
+  fi
+
+  return 0
+}
+
+initialize_native_environment() {
+  local job_id="${BUILDKITE_JOB_ID:-${BUILDKITE_PARALLEL_JOB:-local}}"
+  local job_id_suffix=""
+  local native_root=""
+  local hf_mount=""
+
+  if [[ "$(id -u)" -ne 0 ]]; then
+    echo "Native ROCm CI currently requires the ci_base container to run as root" >&2
+    return 1
+  fi
+
+  job_id="${job_id//[^A-Za-z0-9_.-]/_}"
+  job_id_suffix="${job_id##*-}"
+  job_id_suffix="${job_id_suffix:0:12}"
+  native_root="/tmp/vllm-native-${job_id}"
+  TMPDIR="/tmp/vllm-${job_id_suffix}/tmp"
+  VLLM_RPC_BASE_PATH="/tmp"
+  : "${TORCHINDUCTOR_CACHE_DIR:=${native_root}/cache/torchinductor}"
+  : "${TRITON_CACHE_DIR:=${native_root}/cache/triton}"
+  : "${VLLM_CACHE_ROOT:=${native_root}/cache/vllm}"
+  : "${XDG_CACHE_HOME:=${native_root}/cache/xdg}"
+  : "${HF_HOME:=/home/buildkite-agent/huggingface}"
+  : "${HF_HUB_DOWNLOAD_TIMEOUT:=300}"
+  : "${HF_HUB_ETAG_TIMEOUT:=60}"
+  export TMPDIR VLLM_RPC_BASE_PATH
+  export TORCHINDUCTOR_CACHE_DIR TRITON_CACHE_DIR VLLM_CACHE_ROOT XDG_CACHE_HOME
+  export HF_HOME HF_HUB_DOWNLOAD_TIMEOUT HF_HUB_ETAG_TIMEOUT
+  export PYTORCH_ROCM_ARCH=""
+
+  mkdir -p "${TMPDIR}" \
+    "${TORCHINDUCTOR_CACHE_DIR}" \
+    "${TRITON_CACHE_DIR}" \
+    "${VLLM_CACHE_ROOT}" \
+    "${XDG_CACHE_HOME}" \
+    "${HF_HOME}" || return 1
+
+  if [[ "${VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE:-0}" == "1" ]]; then
+    if ! command -v findmnt >/dev/null 2>&1; then
+      echo "findmnt is required to verify the native Hugging Face cache mount" >&2
+      return 1
+    fi
+    hf_mount=$(findmnt -n -T "${HF_HOME}" -o TARGET 2>/dev/null || true)
+    if [[ -z "${hf_mount}" || "${hf_mount}" == "/" ]]; then
+      echo "Native CI requires a persistent volume mounted at or above ${HF_HOME}" >&2
+      return 1
+    fi
+  fi
+}
+
+run_native_preflight() {
+  local expected_gpus="${VLLM_CI_EXPECTED_GPU_COUNT:-1}"
+
+  if [[ ! "${expected_gpus}" =~ ^[0-9]+$ ]]; then
+    echo "Invalid VLLM_CI_EXPECTED_GPU_COUNT=${expected_gpus}" >&2
+    return 1
+  fi
+
+  python3 -c "import encodings, importlib.metadata as im, importlib.util as iu; [im.version(d) for d in ('transformers', 'torch', 'ray', 'sympy', 'markupsafe', 'vllm')]; missing=[m for m in ('torch.utils.model_zoo', 'transformers.models.nomic_bert', 'ray.dag', 'sympy.physics', 'markupsafe._speedups') if iu.find_spec(m) is None]; assert not missing, missing" || return 1
+
+  if [[ "${expected_gpus}" == "0" ]]; then
+    echo "Native CPU-only AMD job: skipping ROCm device validation"
+    return 0
+  fi
+
+  echo "--- ROCm info"
+  rocminfo || return 1
+  VLLM_CI_EXPECTED_GPU_COUNT="${expected_gpus}" python3 - <<'PY'
+import os
+
+import torch
+
+expected = int(os.environ["VLLM_CI_EXPECTED_GPU_COUNT"])
+assert torch.version.hip, "PyTorch is not a ROCm build"
+assert torch.cuda.is_available(), "ROCm GPU is not available to PyTorch"
+actual = torch.cuda.device_count()
+assert actual == expected, f"Expected {expected} ROCm GPU(s), found {actual}"
+PY
+}
+
 is_multi_node() {
   local cmds="$1"
   # Primary signal: NUM_NODES environment variable set by the pipeline
@@ -350,7 +644,58 @@ re_quote_pytest_markers() {
 # Main
 ###############################################################################
 
-# --- GPU initialization ---
+if is_native_runtime; then
+  echo "--- Native in-pod ROCm CI (AMD_CI_RUNTIME=${AMD_CI_RUNTIME:-unset}, NATIVE_CI=${NATIVE_CI:-unset})"
+  artifact_work_dir=""
+
+  cleanup_native_workspace() {
+    if [[ -n "${artifact_work_dir}" ]]; then
+      rm -rf "${artifact_work_dir}"
+    fi
+  }
+  trap cleanup_native_workspace EXIT
+
+  if [[ -n "${VLLM_TEST_COMMANDS:-}" ]]; then
+    commands="${VLLM_TEST_COMMANDS}"
+    commands_source="env"
+  else
+    commands="$*"
+    commands_source="argv"
+    if [[ -z "$commands" ]]; then
+      echo "Error: No test commands provided for native CI." >&2
+      exit 1
+    fi
+  fi
+
+  if [[ "$commands_source" == "argv" ]]; then
+    commands=$(re_quote_pytest_markers "$commands")
+  fi
+
+  if is_multi_node "$commands"; then
+    echo "Native CI does not support multi-node jobs yet."
+    exit 1
+  fi
+
+  if ! initialize_native_environment; then
+    echo "Failed to initialize the native test environment"
+    exit 1
+  fi
+  if ! prepare_native_workspace; then
+    echo "Failed to prepare native test workspace"
+    exit 1
+  fi
+
+  export PYTHONPATH="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
+
+  echo "Native test commands: $commands"
+  run_native_preflight || exit 1
+  # Keep AMD CI orchestration variables out of vLLM's runtime environment.
+  clear_ci_orchestration_env
+  /bin/bash -o pipefail -c "${commands}"
+  handle_pytest_exit "$?"
+fi
+
+# --- GPU initialization for legacy Docker execution ---
 echo "--- ROCm info"
 rocminfo
 
@@ -452,25 +797,27 @@ fi
 
 echo "Final commands: $commands"
 
-# The ROCm test image often ships /vllm-workspace without .git (artifact tarball unpack).
-# tests/standalone_tests/python_only_compile.sh uses merge-base(HEAD, origin/main) for
-# wheels.vllm.ai; compute on the agent (full git checkout) and pass into the container.
-vllm_standalone_merge_base=""
-checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
-if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
-  checkout="."
+standalone_merge_base_env=()
+if [[ "$commands" == *python_only_compile.sh* ]]; then
+  # The ROCm test image often ships /vllm-workspace without .git. Resolve the
+  # wheels.vllm.ai commit from the agent checkout for this test only.
+  vllm_standalone_merge_base=""
+  checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
+  if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
+    checkout="."
+  fi
+  # Pass safe.directory per-command because Buildkite uses mixed user IDs.
+  if git -c "safe.directory=${checkout}" -C "${checkout}" rev-parse --is-inside-work-tree >/dev/null 2>&1; then
+    vllm_standalone_merge_base="$(
+      git -c "safe.directory=${checkout}" -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
+    )"
+  fi
+  if [[ -z "${vllm_standalone_merge_base}" ]]; then
+    vllm_standalone_merge_base="${BUILDKITE_COMMIT:-}"
+  fi
+  echo "INFO: passing CI_STANDALONE_MERGE_BASE into container: ${vllm_standalone_merge_base}"
+  standalone_merge_base_env=(-e "CI_STANDALONE_MERGE_BASE=${vllm_standalone_merge_base}")
 fi
-# Pass safe.directory per-command (-c) because buildkite runs will always fail
-# the next check on git 2.35.2+ due to mixed uses of root and buildkite-agent/uids.
-if git -c "safe.directory=${checkout}" -C "${checkout}" rev-parse --is-inside-work-tree >/dev/null 2>&1; then
-  vllm_standalone_merge_base="$(
-    git -c "safe.directory=${checkout}" -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
-  )"
-fi
-if [[ -z "${vllm_standalone_merge_base}" ]]; then
-  vllm_standalone_merge_base="${BUILDKITE_COMMIT:-}"
-fi
-echo "INFO: passing VLLM_STANDALONE_MERGE_BASE into container: ${vllm_standalone_merge_base}"
 
 MYPYTHONPATH="/vllm-workspace"
 
@@ -501,6 +848,7 @@ else
 fi
 
 # --- Route: multi-node vs single-node ---
+clear_ci_orchestration_env
 if is_multi_node "$commands"; then
   echo "--- Multi-node job detected"
   export DCKR_VER=$(docker --version | sed 's/Docker version \(.*\), build .*/\1/')
@@ -589,7 +937,9 @@ else
     -e FORCE_COLOR \
     -e CLICOLOR_FORCE \
     -e PY_COLORS \
+    -e PYTHONFAULTHANDLER \
     -e PYTEST_ADDOPTS \
+    -e PYTEST_TIMEOUT \
     -v "${HF_CACHE}:${HF_MOUNT}" \
     -e "HF_HOME=${HF_MOUNT}" \
     -e "PYTHONPATH=${MYPYTHONPATH}" \
@@ -599,7 +949,7 @@ else
     -e "VLLM_CACHE_ROOT=${CONTAINER_CACHE_ROOT}/vllm" \
     -e "XDG_CACHE_HOME=${CONTAINER_CACHE_ROOT}/xdg" \
     -e "PYTORCH_ROCM_ARCH=" \
-    -e "VLLM_STANDALONE_MERGE_BASE=${vllm_standalone_merge_base}" \
+    "${standalone_merge_base_env[@]}" \
     --name "${container_name}" \
     "${image_name}" \
     /bin/bash -c "${CONTAINER_PREFLIGHT} && ${commands}"

.buildkite/scripts/hardware_ci/run-cpu-test-arm.sh

@@ -40,7 +40,9 @@ function cpu_tests() {
     pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
     pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
     pytest -x -v -s tests/kernels/moe/test_cpu_int4_moe.py
-    pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py"
+    pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
+    pytest -x -v -s tests/kernels/mamba/test_causal_conv1d.py
+    pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.py"
 
   # skip tests requiring model downloads if HF_TOKEN is not set
   # due to rate-limits
@@ -97,3 +99,4 @@ function cpu_tests() {
 # All of CPU tests are expected to be finished less than 40 mins.
 export -f cpu_tests
 timeout 2h bash -c cpu_tests
+

.buildkite/scripts/hardware_ci/run-intel-ci-test.sh

@@ -35,6 +35,7 @@ case "${test_suite}" in
     pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py
     pytest -v -s v1/structured_output
     pytest -v -s v1/test_serial_utils.py
+    pytest -v -s v1/e2e/general/test_correctness_sliding_window.py --deselect="tests/v1/e2e/general/test_correctness_sliding_window.py::test_sliding_window_retrieval[True-1-5-google/gemma-3-1b-it]"
     pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py --ignore=v1/spec_decode/test_speculators_correctness.py
     pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py
     ;;

.buildkite/scripts/publish-release-images.sh

@@ -130,6 +130,22 @@ docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm
 docker push vllm/vllm-openai-rocm:latest-base
 docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
 
+# ---- XPU ----
+
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:latest-x86_64
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
+docker push vllm/vllm-openai-xpu:latest-x86_64
+docker push vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
+
+docker manifest rm vllm/vllm-openai-xpu:latest || true
+docker manifest rm vllm/vllm-openai-xpu:v${RELEASE_VERSION} || true
+docker manifest create vllm/vllm-openai-xpu:latest vllm/vllm-openai-xpu:latest-x86_64 --amend
+docker manifest create vllm/vllm-openai-xpu:v${RELEASE_VERSION} vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64 --amend
+docker manifest push vllm/vllm-openai-xpu:latest
+docker manifest push vllm/vllm-openai-xpu:v${RELEASE_VERSION}
+
 # ---- CPU ----
 # CPU images are behind separate block steps and may not have been built.
 # All-or-nothing: inspect both arches first, then either publish everything

.buildkite/scripts/rocm/build-test-image.sh

@@ -13,6 +13,18 @@ metadata_get() {
     fi
 }
 
+use_ci_base_if_present() {
+    local ci_base_image=""
+
+    ci_base_image="$(metadata_get rocm-ci-base-image)"
+    if [[ -z "${ci_base_image}" ]]; then
+        return 1
+    fi
+
+    export CI_BASE_IMAGE="${ci_base_image}"
+    echo "Using ROCm ci_base image selected by the preceding build step: ${CI_BASE_IMAGE}"
+}
+
 use_refreshed_base_if_present() {
     local base_refreshed=""
 
@@ -22,15 +34,12 @@ use_refreshed_base_if_present() {
     fi
 
     export BASE_IMAGE
-    export CI_BASE_IMAGE
     export IMAGE_TAG_LATEST
 
     BASE_IMAGE="$(metadata_get rocm-base-image)"
-    CI_BASE_IMAGE="$(metadata_get rocm-ci-base-image)"
     IMAGE_TAG_LATEST="$(metadata_get rocm-ci-image-descriptive)"
 
     echo "Using refreshed ROCm base image for test image: ${BASE_IMAGE}"
-    echo "Using refreshed ROCm ci_base image for test image: ${CI_BASE_IMAGE}"
     if [[ -n "${IMAGE_TAG_LATEST}" ]]; then
         echo "Also tagging full ROCm CI image as: ${IMAGE_TAG_LATEST}"
     fi
@@ -41,6 +50,8 @@ use_refreshed_base_if_present() {
 main() {
     local base_refreshed=0
 
+    use_ci_base_if_present || true
+
     if use_refreshed_base_if_present; then
         base_refreshed=1
     fi

.buildkite/scripts/run-rust-frontend-cargo-ci.sh

@@ -21,16 +21,21 @@ export CARGO_HOME="${CARGO_HOME:-$HOME/.cargo}"
 export RUSTUP_HOME="${RUSTUP_HOME:-$HOME/.rustup}"
 export PATH="$CARGO_HOME/bin:$PATH"
 
+PROTOC_VERSION="${PROTOC_VERSION:-31.1}"
+CARGO_BINSTALL_VERSION="${CARGO_BINSTALL_VERSION:-1.20.1}"
+UV_VERSION="${UV_VERSION:-0.11.28}"
+PYO3_PYTHON_VERSION="${PYO3_PYTHON_VERSION:-3.12}"
+
+CARGO_SORT_VERSION_REQ="${CARGO_SORT_VERSION_REQ:-2}"
+CARGO_DENY_VERSION_REQ="${CARGO_DENY_VERSION_REQ:-0.20}"
+CARGO_NEXTEST_VERSION_REQ="${CARGO_NEXTEST_VERSION_REQ:-0.9}"
+CARGO_LLVM_COV_VERSION="${CARGO_LLVM_COV_VERSION:-0.8.7}"
+
 log_section() {
   echo "--- $*"
 }
 
 install_protoc() {
-  if command -v protoc >/dev/null 2>&1; then
-    return
-  fi
-
-  local version="${PROTOC_VERSION:-31.1}"
   local arch
   case "$(uname -m)" in
     x86_64)
@@ -45,16 +50,17 @@ install_protoc() {
       ;;
   esac
 
-  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${version}/protoc-${version}-linux-${arch}.zip"
+  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${PROTOC_VERSION}/protoc-${PROTOC_VERSION}-linux-${arch}.zip"
   local tmp_dir
   tmp_dir="$(mktemp -d)"
 
-  log_section "Installing protoc ${version}"
+  log_section "Installing protoc ${PROTOC_VERSION}"
   curl -L --proto '=https' --tlsv1.2 -sSf "$url" -o "$tmp_dir/protoc.zip"
   mkdir -p "$CARGO_HOME/bin"
   unzip -q "$tmp_dir/protoc.zip" bin/protoc 'include/*' -d "$CARGO_HOME"
   chmod +x "$CARGO_HOME/bin/protoc"
   rm -rf "$tmp_dir"
+  protoc --version
 }
 
 rust_toolchain() {
@@ -75,66 +81,60 @@ install_rust_toolchain() {
 }
 
 install_cargo_binstall() {
-  if command -v cargo-binstall >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-binstall"
+  log_section "Installing cargo-binstall ${CARGO_BINSTALL_VERSION}"
   curl -L --proto '=https' --tlsv1.2 -sSf \
-    https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
-    | bash
+    "https://raw.githubusercontent.com/cargo-bins/cargo-binstall/v${CARGO_BINSTALL_VERSION}/install-from-binstall-release.sh" \
+    | env BINSTALL_VERSION="$CARGO_BINSTALL_VERSION" bash
+  cargo-binstall -V
 }
 
 install_cargo_sort() {
-  if command -v cargo-sort >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-sort"
-  install_cargo_binstall
-  cargo binstall --no-confirm cargo-sort
+  log_section "Installing cargo-sort ${CARGO_SORT_VERSION_REQ}"
+  cargo binstall --no-confirm --force "cargo-sort@${CARGO_SORT_VERSION_REQ}"
 }
 
 install_cargo_deny() {
-  if command -v cargo-deny >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-deny"
-  install_cargo_binstall
-  cargo binstall --no-confirm cargo-deny
+  log_section "Installing cargo-deny ${CARGO_DENY_VERSION_REQ}"
+  cargo binstall --no-confirm --force "cargo-deny@${CARGO_DENY_VERSION_REQ}"
 }
 
 install_cargo_nextest() {
-  if command -v cargo-nextest >/dev/null 2>&1; then
-    return
-  fi
+  log_section "Installing cargo-nextest ${CARGO_NEXTEST_VERSION_REQ}"
+  cargo binstall \
+    --no-confirm \
+    --force \
+    --secure \
+    "cargo-nextest@${CARGO_NEXTEST_VERSION_REQ}"
+}
 
-  log_section "Installing cargo-nextest"
-  install_cargo_binstall
-  cargo binstall --no-confirm --secure cargo-nextest
+install_cargo_llvm_cov() {
+  log_section "Installing cargo-llvm-cov ${CARGO_LLVM_COV_VERSION}"
+  local toolchain
+  toolchain="$(rust_toolchain)"
+  rustup component add --toolchain "$toolchain" llvm-tools-preview
+  cargo binstall \
+    --no-confirm \
+    --force \
+    --secure \
+    "cargo-llvm-cov@${CARGO_LLVM_COV_VERSION}"
 }
 
 install_uv() {
-  if command -v uv >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing uv"
-  curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.sh \
+  log_section "Installing uv ${UV_VERSION}"
+  curl -L --proto '=https' --tlsv1.2 -sSf \
+    "https://github.com/astral-sh/uv/releases/download/${UV_VERSION}/uv-installer.sh" \
     | env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
+  uv --version
 }
 
 setup_pyo3_python() {
-  local python_version="${PYO3_PYTHON_VERSION:-3.12}"
-
-  log_section "Installing Python ${python_version} for PyO3 tests"
-  uv python install "$python_version"
+  log_section "Installing Python ${PYO3_PYTHON_VERSION} for PyO3 tests"
+  uv python install "$PYO3_PYTHON_VERSION"
   PYO3_PYTHON="$(uv python find \
     --managed-python \
     --no-project \
     --resolve-links \
-    "$python_version")"
+    "$PYO3_PYTHON_VERSION")"
   export PYO3_PYTHON
 
   local python_libdir
@@ -156,6 +156,7 @@ PY
 }
 
 run_style_clippy() {
+  install_cargo_binstall
   install_cargo_sort
   install_cargo_deny
 
@@ -186,15 +187,43 @@ run_style_clippy() {
 run_tests() {
   install_uv
   setup_pyo3_python
+  install_cargo_binstall
   install_cargo_nextest
+  install_cargo_llvm_cov
 
-  log_section "Running cargo nextest"
-  cargo nextest run \
+  log_section "Running cargo nextest with Rust coverage"
+  mkdir -p artifacts
+  export LLVM_PROFILE_FILE_NAME="vllm-rust-unit-%4m.profraw"
+  cargo llvm-cov clean \
+    --manifest-path rust/Cargo.toml \
+    --profraw-only
+
+  set +e
+  cargo llvm-cov nextest \
     --manifest-path rust/Cargo.toml \
     --workspace \
     --all-features \
     --locked \
-    --no-fail-fast
+    --no-fail-fast \
+    --no-clean \
+    --lcov \
+    --output-path artifacts/rust-unit.lcov \
+    --ignore-filename-regex='/\.cargo/(registry|git)/|/rustc/|/target/'
+  local coverage_rc=$?
+
+  local upload_rc=0
+  if [[ $coverage_rc -eq 0 ]]; then
+    # shellcheck source=.buildkite/scripts/rust-coverage.sh
+    source .buildkite/scripts/rust-coverage.sh
+    rust_coverage_upload artifacts/rust-unit.lcov rust-unit
+    upload_rc=$?
+  fi
+  set -e
+
+  if [[ $coverage_rc -ne 0 ]]; then
+    return "$coverage_rc"
+  fi
+  return "$upload_rc"
 }
 
 install_protoc

.buildkite/scripts/rust-coverage.sh

@@ -0,0 +1,182 @@
+#!/bin/sh
+
+RUST_CODECOV_VERSION="v11.3.1"
+RUST_CODECOV_SHA256="ca1d64196d2d34771084afe76ea657d581bf628e31d993ff8e52ea09cc88a56d"
+
+rust_coverage_repo_root() {
+    if [ -f /vllm-workspace/.buildkite/scripts/rust-coverage.sh ]; then
+        printf '%s\n' /vllm-workspace
+    elif [ -n "${BUILDKITE_BUILD_CHECKOUT_PATH:-}" ] \
+        && [ -d "$BUILDKITE_BUILD_CHECKOUT_PATH" ]; then
+        printf '%s\n' "$BUILDKITE_BUILD_CHECKOUT_PATH"
+    else
+        git rev-parse --show-toplevel
+    fi
+}
+
+rust_coverage_start() {
+    RUST_COVERAGE_FLAG=${1:?coverage flag is required}
+    RUST_COVERAGE_DIR="/tmp/vllm-rust-coverage/${BUILDKITE_JOB_ID:-local}"
+    export RUST_COVERAGE_FLAG RUST_COVERAGE_DIR
+    mkdir -p "$RUST_COVERAGE_DIR"
+    LLVM_PROFILE_FILE="$RUST_COVERAGE_DIR/rust-%4m.profraw"
+    export LLVM_PROFILE_FILE
+    trap rust_coverage_finalize 0
+}
+
+rust_coverage_objects() {
+    rust_cov_objects_manifest="$(dirname "$(command -v llvm-cov)")/../objects"
+    python3 - "$rust_cov_objects_manifest" <<'PY'
+from pathlib import Path
+import sys
+
+for relative in Path(sys.argv[1]).read_text().splitlines():
+    for entry in sys.path:
+        path = Path(entry or ".").resolve() / relative
+        if path.is_file():
+            print(path)
+            break
+    else:
+        raise RuntimeError(f"installed Rust coverage object was not found: {relative}")
+PY
+}
+
+rust_coverage_collect() {
+    rust_cov_collect_flag=${1:?coverage flag is required}
+    rust_cov_collect_lcov="$RUST_COVERAGE_DIR/$rust_cov_collect_flag.lcov"
+
+    rust_cov_collect_objects=$(rust_coverage_objects) || return 1
+    rust_cov_collect_primary=
+    set --
+    while IFS= read -r rust_cov_collect_object; do
+        if [ -z "$rust_cov_collect_primary" ]; then
+            rust_cov_collect_primary=$rust_cov_collect_object
+        else
+            set -- "$@" "--object=$rust_cov_collect_object"
+        fi
+    done <<EOF
+$rust_cov_collect_objects
+EOF
+
+    llvm-profdata merge \
+        -sparse \
+        "$RUST_COVERAGE_DIR"/*.profraw \
+        -o "$RUST_COVERAGE_DIR/merged.profdata" || return 1
+    llvm-cov export \
+        "$rust_cov_collect_primary" \
+        "$@" \
+        --format=lcov \
+        --instr-profile="$RUST_COVERAGE_DIR/merged.profdata" \
+        --ignore-filename-regex='/\.cargo/(registry|git)/|/rustc/|/target/' \
+        > "$rust_cov_collect_lcov" || return 1
+    RUST_COVERAGE_LCOV=$rust_cov_collect_lcov
+    export RUST_COVERAGE_LCOV
+}
+
+rust_coverage_upload() {
+    rust_cov_upload_lcov=${1:?LCOV path is required}
+    rust_cov_upload_flag=${2:?coverage flag is required}
+    rust_cov_upload_repo_root=$(rust_coverage_repo_root) || return 1
+
+    if [ "$(uname -m)" != "x86_64" ]; then
+        echo "Rust coverage upload currently supports x86_64 CI agents" >&2
+        return 1
+    fi
+
+    rust_cov_upload_codecov_dir=$(mktemp -d /tmp/codecov-bin.XXXXXX) \
+        || return 1
+    curl -fsSL \
+        "https://github.com/codecov/codecov-cli/releases/download/${RUST_CODECOV_VERSION}/codecovcli_linux" \
+        -o "$rust_cov_upload_codecov_dir/codecov" || return 1
+    echo "$RUST_CODECOV_SHA256  $rust_cov_upload_codecov_dir/codecov" \
+        | sha256sum -c - || return 1
+    chmod +x "$rust_cov_upload_codecov_dir/codecov" || return 1
+
+    rust_cov_upload_slug="vllm-project/vllm"
+    if [ -n "${BUILDKITE_PULL_REQUEST:-}" ] \
+        && [ "${BUILDKITE_PULL_REQUEST}" != "false" ] \
+        && [ -n "${BUILDKITE_PULL_REQUEST_REPO:-}" ]; then
+        rust_cov_upload_slug=$(echo "$BUILDKITE_PULL_REQUEST_REPO" \
+            | sed -E 's#(git@|https?://)([^/:]+)[:/]([^/]+/[^/.]+)(\.git)?$#\3#')
+        case "$rust_cov_upload_slug" in
+            */*) ;;
+            *) rust_cov_upload_slug="vllm-project/vllm" ;;
+        esac
+    fi
+
+    rust_cov_upload_branch=${BUILDKITE_BRANCH:?BUILDKITE_BRANCH is required}
+    if [ -z "${CODECOV_TOKEN:-}" ]; then
+        # Codecov accepts tokenless public uploads on unprotected branch names.
+        # A colon-separated prefix keeps feature-branch and fork uploads from
+        # requiring a repository secret.
+        if [ -n "${BUILDKITE_PULL_REQUEST:-}" ] \
+            && [ "${BUILDKITE_PULL_REQUEST}" != "false" ]; then
+            rust_cov_upload_branch="pr${BUILDKITE_PULL_REQUEST}:$rust_cov_upload_branch"
+        else
+            rust_cov_upload_branch="buildkite:$rust_cov_upload_branch"
+        fi
+    fi
+
+    set --
+    set -- "$@" upload-process
+    set -- "$@" --file "$rust_cov_upload_lcov"
+    # LCOV paths are mapped server-side by codecov.yml. Skip the CLI's local
+    # source-line fix scanning, which is unrelated to path mapping.
+    set -- "$@" --disable-search --disable-file-fixes
+    set -- "$@" --fail-on-error --git-service github
+    set -- "$@" --build "${BUILDKITE_BUILD_NUMBER:?BUILDKITE_BUILD_NUMBER is required}"
+    set -- "$@" --branch "$rust_cov_upload_branch"
+    set -- "$@" --sha "${BUILDKITE_COMMIT:?BUILDKITE_COMMIT is required}"
+    set -- "$@" --slug "$rust_cov_upload_slug"
+    set -- "$@" --flag "$rust_cov_upload_flag"
+    set -- "$@" --name "${rust_cov_upload_flag}-${BUILDKITE_JOB_ID:?BUILDKITE_JOB_ID is required}"
+    set -- "$@" --dir "$rust_cov_upload_repo_root"
+    set -- "$@" --network-root-folder "$rust_cov_upload_repo_root"
+    if [ -n "${BUILDKITE_PULL_REQUEST:-}" ] \
+        && [ "${BUILDKITE_PULL_REQUEST}" != "false" ]; then
+        set -- "$@" --pr "$BUILDKITE_PULL_REQUEST"
+    fi
+
+    rust_cov_upload_log="$rust_cov_upload_codecov_dir/codecov.log"
+    # E2E steps run from tests/, so execute from the repository root to resolve
+    # codecov.yml and repository paths consistently.
+    (
+        cd "$rust_cov_upload_repo_root" || exit 1
+        "$rust_cov_upload_codecov_dir/codecov" "$@"
+    ) >"$rust_cov_upload_log" 2>&1
+    rust_cov_upload_rc=$?
+    cat "$rust_cov_upload_log"
+    # v11.3.1 can log API failures while returning zero even with
+    # --fail-on-error. Preserve the strict CI contract explicitly.
+    if grep -aEq 'error.* -- ' "$rust_cov_upload_log"; then
+        echo "Codecov CLI reported an upload error" >&2
+        rust_cov_upload_rc=1
+    fi
+    rm -rf "$rust_cov_upload_codecov_dir"
+    return "$rust_cov_upload_rc"
+}
+
+rust_coverage_finalize() {
+    rust_cov_finalize_test_rc=$?
+    trap - 0
+    set +e
+
+    rust_coverage_collect "$RUST_COVERAGE_FLAG"
+    rust_cov_finalize_collect_rc=$?
+
+    rust_cov_finalize_upload_rc=0
+    if [ "$rust_cov_finalize_collect_rc" -eq 0 ]; then
+        rust_coverage_upload "$RUST_COVERAGE_LCOV" "$RUST_COVERAGE_FLAG"
+        rust_cov_finalize_upload_rc=$?
+    fi
+
+    find "$RUST_COVERAGE_DIR" -type f -name '*.profraw' -delete
+
+    if [ "$rust_cov_finalize_test_rc" -ne 0 ]; then
+        exit "$rust_cov_finalize_test_rc"
+    fi
+    if [ "$rust_cov_finalize_collect_rc" -ne 0 ]; then
+        exit "$rust_cov_finalize_collect_rc"
+    fi
+    exit "$rust_cov_finalize_upload_rc"
+}

.buildkite/scripts/trigger-ci-build.sh

@@ -23,6 +23,7 @@ NC='\033[0m' # No Color
 # Default configuration
 PIPELINE="ci"
 DRY_RUN=true
+TORCH_NIGHTLY=false
 
 usage() {
     cat <<EOF
@@ -34,12 +35,14 @@ Sets RUN_ALL=1 and NIGHTLY=1 environment variables.
 SAFETY: Dry-run by default. Use --execute to actually trigger a build.
 
 Options:
-    --execute       Actually trigger the build (default: dry-run)
-    --pipeline      Buildkite pipeline slug (default: ${PIPELINE})
-    --commit        Override commit SHA (default: current HEAD)
-    --branch        Override branch name (default: current branch)
-    --message       Custom build message (default: auto-generated)
-    --help          Show this help message
+    --execute        Actually trigger the build (default: dry-run)
+    --pipeline       Buildkite pipeline slug (default: ${PIPELINE})
+    --commit         Override commit SHA (default: current HEAD)
+    --branch         Override branch name (default: current branch)
+    --message        Custom build message (default: auto-generated)
+    --torch-nightly  Also build and run the full suite against torch nightly
+                     (sets TORCH_NIGHTLY=1)
+    --help           Show this help message
 
 Prerequisites:
     - bk CLI installed: brew tap buildkite/buildkite && brew install buildkite/buildkite/bk
@@ -49,6 +52,7 @@ Examples:
     $(basename "$0")                        # Dry-run, show what would happen
     $(basename "$0") --execute              # Actually trigger the build
     $(basename "$0") --pipeline ci-shadow   # Dry-run with different pipeline
+    $(basename "$0") --torch-nightly        # Dry-run a full torch-nightly run
 EOF
     exit 1
 }
@@ -96,6 +100,10 @@ while [[ $# -gt 0 ]]; do
             MESSAGE="$2"
             shift 2
             ;;
+        --torch-nightly)
+            TORCH_NIGHTLY=true
+            shift
+            ;;
         --help|-h)
             usage
             ;;
@@ -171,11 +179,17 @@ if [[ $(echo "$REMOTE_BRANCHES" | wc -l) -gt 5 ]]; then
 fi
 echo ""
 
+# Environment variables passed to the build.
+BUILD_ENV=("RUN_ALL=1" "NIGHTLY=1")
+if [[ "$TORCH_NIGHTLY" == true ]]; then
+    BUILD_ENV+=("TORCH_NIGHTLY=1")
+fi
+
 log_info "Pipeline: ${PIPELINE}"
 log_info "Branch: ${BRANCH}"
 log_info "Commit: ${COMMIT}"
 log_info "Message: ${MESSAGE}"
-log_info "Environment: RUN_ALL=1, NIGHTLY=1"
+log_info "Environment: ${BUILD_ENV[*]}"
 echo ""
 
 # Build the command
@@ -187,9 +201,10 @@ CMD=(bk build create
     --commit "${COMMIT}"
     --branch "${BRANCH}"
     --message "${MESSAGE}"
-    --env "RUN_ALL=1"
-    --env "NIGHTLY=1"
 )
+for env_var in "${BUILD_ENV[@]}"; do
+    CMD+=(--env "${env_var}")
+done
 
 if [[ "$DRY_RUN" == true ]]; then
     echo "=========================================="
@@ -210,8 +225,14 @@ if [[ "$DRY_RUN" == true ]]; then
     echo "    --commit '$(escape_for_shell "${COMMIT}")' \\"
     echo "    --branch '$(escape_for_shell "${BRANCH}")' \\"
     echo "    --message '$(escape_for_shell "${MESSAGE}")' \\"
-    echo "    --env 'RUN_ALL=1' \\"
-    echo "    --env 'NIGHTLY=1'"
+    last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
+    for i in "${!BUILD_ENV[@]}"; do
+        if [[ $i -eq $last_idx ]]; then
+            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
+        else
+            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
+        fi
+    done
     echo ""
     echo "=========================================="
     echo -e "${YELLOW}To actually trigger this build, run:${NC}"

.buildkite/scripts/upload-nightly-wheels.sh

@@ -6,8 +6,14 @@ set -ex
 # manylinux platform tag with auditwheel.
 # Index generation is handled separately by generate-and-upload-nightly-index.sh.
 
-# shellcheck source=lib/manylinux.sh
-source .buildkite/scripts/lib/manylinux.sh
+# auditwheel is Linux-only; macOS wheels already carry a valid tag, so skip the
+# manylinux retag for them.
+WHEEL_PLATFORM="${VLLM_WHEEL_PLATFORM:-linux}"
+
+if [[ "$WHEEL_PLATFORM" == "linux" ]]; then
+  # shellcheck source=lib/manylinux.sh
+  source .buildkite/scripts/lib/manylinux.sh
+fi
 
 BUCKET="vllm-wheels"
 SUBPATH=$BUILDKITE_COMMIT
@@ -27,8 +33,10 @@ wheel="${wheel_files[0]}"
 
 # ========= detect manylinux tag and rename ==========
 
-wheel="$(apply_manylinux_tag "$wheel")"
-echo "Renamed wheel to: $wheel"
+if [[ "$WHEEL_PLATFORM" == "linux" ]]; then
+  wheel="$(apply_manylinux_tag "$wheel")"
+  echo "Renamed wheel to: $wheel"
+fi
 
 # Extract the version from the wheel
 version=$(unzip -p "$wheel" '**/METADATA' | grep '^Version: ' | cut -d' ' -f2)

.buildkite/scripts/upload-rocm-wheels.sh

@@ -113,8 +113,8 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py \
 echo "Uploading indices to $S3_COMMIT_PREFIX"
 aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
 
-# Update rocm/nightly/ if on main branch and not a PR
-if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]] || [[ "$NIGHTLY" == "1" ]]; then
+# Only scheduled nightly builds should update the moving nightly index.
+if [[ "${NIGHTLY:-0}" == "1" ]]; then
     echo "Updating rocm/nightly/ index..."
     aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/rocm/nightly/"
 fi
@@ -147,7 +147,7 @@ echo ""
 echo "Install command (by commit):"
 echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/$ROCM_SUBPATH/"
 echo ""
-if [[ "$BUILDKITE_BRANCH" == "main" ]] || [[ "$NIGHTLY" == "1" ]]; then
+if [[ "${NIGHTLY:-0}" == "1" ]]; then
     echo "Install command (nightly):"
     echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/rocm/nightly/"
 fi

.buildkite/scripts/xpu/push-nightly-builds-xpu.sh

@@ -0,0 +1,21 @@
+#!/bin/bash
+
+set -ex
+
+ORIG_TAG_NAME="$BUILDKITE_COMMIT"
+REPO="vllm/vllm-openai-xpu"
+
+echo "Pushing original XPU tag ${ORIG_TAG_NAME}-xpu to nightly tags in ${REPO}"
+
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu ${REPO}:nightly-x86_64
+docker push ${REPO}:nightly-x86_64
+
+docker manifest rm ${REPO}:nightly || true
+docker manifest rm ${REPO}:nightly-"$BUILDKITE_COMMIT" || true
+docker manifest create ${REPO}:nightly ${REPO}:nightly-x86_64 --amend
+docker manifest create ${REPO}:nightly-"$BUILDKITE_COMMIT" ${REPO}:nightly-x86_64 --amend
+docker manifest push ${REPO}:nightly
+docker manifest push ${REPO}:nightly-"$BUILDKITE_COMMIT"

.buildkite/test_areas/attention.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: V1 attention (H100-MI300)
   key: v1-attention-h100-mi300
-  timeout_in_minutes: 30
+  timeout_in_minutes: 85
   device: h100
   source_file_dependencies:
     - vllm/config/attention.py
@@ -12,11 +12,12 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention
+    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 2
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 70
+      timeout_in_minutes: 95
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -30,7 +31,7 @@ steps:
 
 - label: V1 attention (B200)
   key: v1-attention-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 80
   device: b200-k8s
   source_file_dependencies:
     - vllm/config/attention.py
@@ -38,4 +39,5 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention
+    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 2

.buildkite/test_areas/basic_correctness.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
   key: basic-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -19,6 +19,6 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 70
       depends_on:
       - image-build-amd

.buildkite/test_areas/benchmarks.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
   key: benchmarks-cli-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -13,6 +13,7 @@ steps:
   - pytest -v -s benchmarks/
   mirror:
     amd:
+      dind: false
       device: mi300_1
       depends_on:
       - image-build-amd
@@ -23,7 +24,7 @@ steps:
   num_gpus: 2
   optional: true
   working_dir: "/vllm-workspace/"
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   source_file_dependencies:
   - benchmarks/attention_benchmarks/
   - vllm/v1/attention/

.buildkite/test_areas/compile.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Sequence Parallel Correctness Tests (2 GPUs)
   key: sequence-parallel-correctness-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 80
   working_dir: "/vllm-workspace/"
   num_devices: 2
   source_file_dependencies:
@@ -19,7 +19,7 @@ steps:
 
 - label: Sequence Parallel Correctness Tests (2xH100)
   key: sequence-parallel-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
   working_dir: "/vllm-workspace/"
   device: h100
   optional: true
@@ -30,7 +30,7 @@ steps:
 
 - label: AsyncTP Correctness Tests (2xH100)
   key: asynctp-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: h100
   optional: true
@@ -41,7 +41,7 @@ steps:
 
 - label: AsyncTP Correctness Tests (B200)
   key: asynctp-correctness-tests-b200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   optional: true
@@ -52,7 +52,7 @@ steps:
 
 - label: Distributed Compile Unit Tests (2xH100)
   key: distributed-compile-unit-tests-2xh100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 2
@@ -66,7 +66,7 @@ steps:
 
 - label: Fusion and Compile Unit Tests (2xB200)
   key: fusion-and-compile-unit-tests-2xb200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   source_file_dependencies:
@@ -96,7 +96,7 @@ steps:
 
 - label: Fusion E2E Quick (H100)
   key: fusion-e2e-quick-h100
-  timeout_in_minutes: 15
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 1
@@ -115,7 +115,7 @@ steps:
 
 - label: Fusion E2E Config Sweep (H100)
   key: fusion-e2e-config-sweep-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 1
@@ -149,7 +149,7 @@ steps:
 
 - label: Fusion E2E TP2 Quick (H100)
   key: fusion-e2e-tp2-quick-h100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 2
@@ -167,7 +167,7 @@ steps:
 
 - label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
   key: fusion-e2e-tp2-ar-rms-config-sweep-h100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 2
@@ -207,7 +207,7 @@ steps:
 
 - label: Fusion E2E TP2 (B200)
   key: fusion-e2e-tp2-b200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   num_devices: 2

.buildkite/test_areas/cuda.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Platform Tests
   key: platform-tests
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
   - vllm/envs.py
@@ -18,8 +18,9 @@ steps:
     - pytest -v -s cuda/test_platform_no_cuda_init.py
 
 - label: Cudagraph
+  device: h200_35gb
   key: cudagraph
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   source_file_dependencies:
   - tests/v1/cudagraph
   - vllm/v1/cudagraph_dispatcher.py
@@ -28,4 +29,4 @@ steps:
   commands:
     - pytest -v -s v1/cudagraph/test_cudagraph_dispatch.py
     - pytest -v -s v1/cudagraph/test_cudagraph_mode.py
-    - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
+    - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py

.buildkite/test_areas/disaggregated.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
   key: distributed-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -15,8 +15,9 @@ steps:
     - bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
+      dind: false
       device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -29,7 +30,7 @@ steps:
 
 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
   key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -65,8 +66,9 @@ steps:
     - DP_EP=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
+      dind: false
       device: mi300_4
-      timeout_in_minutes: 50
+      timeout_in_minutes: 60
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -79,7 +81,7 @@ steps:
 
 - label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
   key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -90,8 +92,9 @@ steps:
     - CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
+      dind: false
       device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -104,7 +107,7 @@ steps:
 
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
   key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -115,8 +118,9 @@ steps:
     - HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
+      dind: false
       device: mi300_4
-      timeout_in_minutes: 60
+      timeout_in_minutes: 80
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -143,7 +147,7 @@ steps:
 
 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
   key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -158,7 +162,7 @@ steps:
 
 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
   key: nixlconnector-pd-spec-decode-acceptance-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: a100
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -171,8 +175,9 @@ steps:
     - bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
   mirror:
     amd:
+      dind: false
       device: mi300_2
-      timeout_in_minutes: 60
+      timeout_in_minutes: 70
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -182,11 +187,11 @@ steps:
         - vllm/platforms/rocm.py
       commands:
         - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-        - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
+        - KV_CACHE_MEMORY_BYTES=8G ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
 
 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
   key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -198,3 +203,25 @@ steps:
   commands:
     - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
+
+# P TP 4 - D DPEP 4 test case for DSv4-Flash
+- label: DSv4-Flash Disaggregated DP EP
+  key: dsv4-flash-disaggregated
+  timeout_in_minutes: 60
+  device: h200
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 8
+  env:
+    ENABLE_HMA_FLAG: "1"
+    DP_EP: "1"
+    GPU_MEMORY_UTILIZATION: "0.85"
+    PREFILLER_TP_SIZE: "4"
+    DECODER_TP_SIZE: "4"
+    PREFILL_BLOCK_SIZE: "256"
+    DECODE_BLOCK_SIZE: "256"
+    MODEL_NAMES: "deepseek-ai/DeepSeek-V4-Flash"
+    VLLM_SERVE_EXTRA_ARGS: "--trust-remote-code,--kv-cache-dtype,fp8"
+  commands:
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - bash v1/kv_connector/nixl_integration/run_accuracy_test.sh

.buildkite/test_areas/distributed.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Comm Ops
   key: distributed-comm-ops
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
 
 - label: Distributed DP Tests (2 GPUs)
   key: distributed-dp-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -39,6 +39,7 @@ steps:
   - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
   mirror:
     amd:
+      dind: false
       device: mi300_2
       depends_on:
       - image-build-amd
@@ -55,7 +56,7 @@ steps:
 
 - label: Distributed Compile + RPC Tests (2 GPUs)
   key: distributed-compile-rpc-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 65
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -78,7 +79,7 @@ steps:
 
 - label: Distributed Torchrun + Shutdown Tests (2 GPUs)
   key: distributed-torchrun-shutdown-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -133,7 +134,7 @@ steps:
 
 - label: Distributed DP Tests (4 GPUs)
   key: distributed-dp-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -154,7 +155,7 @@ steps:
 
 - label: Distributed Compile + Comm (4 GPUs)
   key: distributed-compile-comm-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 70
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -176,7 +177,7 @@ steps:
 
 - label: Distributed Tests (8xH100)
   key: distributed-tests-8xh100
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   device: h100
   num_devices: 8
   working_dir: "/vllm-workspace/tests"
@@ -212,7 +213,7 @@ steps:
 
 - label: Distributed Tests (2xH100-2xMI300)
   key: distributed-tests-2xh100-2xmi300
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
   device: h100
   optional: true
   working_dir: "/vllm-workspace/"
@@ -259,7 +260,7 @@ steps:
 
 - label: Pipeline + Context Parallelism (4 GPUs)
   key: pipeline-context-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -274,7 +275,7 @@ steps:
 
 - label: RayExecutorV2 (4 GPUs)
   key: rayexecutorv2-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:

.buildkite/test_areas/docker.yaml

@@ -3,7 +3,7 @@ depends_on:
   - image-build-cpu
 steps:
 - label: Docker Build Metadata
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   device: cpu-small
   source_file_dependencies:
     - .buildkite/release-pipeline.yaml

.buildkite/test_areas/e2e_integration.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
   key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h100
   optional: true
   num_devices: 4
@@ -14,7 +14,7 @@ steps:
 
 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
   key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h100
   optional: true
   num_devices: 4
@@ -24,7 +24,7 @@ steps:
 
 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
   key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -34,7 +34,7 @@ steps:
 
 - label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
   key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h100
   optional: true
   num_devices: 4
@@ -44,7 +44,7 @@ steps:
 
 - label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
   key: deepseek-v2-lite-prefetch-offload-accuracy-h100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
   device: h100
   optional: true
   num_devices: 1

.buildkite/test_areas/engine.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Engine
   key: engine
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/compilation/
@@ -29,13 +29,13 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 50
       depends_on:
       - image-build-amd
 
 - label: Engine (1 GPU)
   key: engine-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   source_file_dependencies:
     - vllm/v1/engine/
     - tests/v1/engine/
@@ -45,13 +45,13 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 40
+      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
 
 - label: e2e Scheduling (1 GPU)
   key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 35
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/
@@ -61,14 +61,14 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 70
       depends_on:
       - image-build-amd
 
 - label: e2e Core (1 GPU)
   device: h200_35gb
   key: e2e-core-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   source_file_dependencies:
     - vllm/v1/
     - tests/v1/e2e/general/
@@ -77,7 +77,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 35
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -87,7 +87,7 @@ steps:
 
 - label: V1 e2e (2 GPUs)
   key: v1-e2e-2-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 2
   source_file_dependencies:
@@ -114,13 +114,14 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
   mirror:
     amd:
+      dind: false
       device: mi300_2
       depends_on:
       - image-build-amd
 
 - label: V1 e2e (4 GPUs)
   key: v1-e2e-4-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 4
   source_file_dependencies:
@@ -148,7 +149,7 @@ steps:
 
 - label: V1 e2e (4xH100)
   key: v1-e2e-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
   device: h100
   num_devices: 4
   optional: true

.buildkite/test_areas/entrypoints.yaml

@@ -3,8 +3,9 @@ depends_on:
   - image-build
 steps:
 - label: Entrypoints Unit Tests
+  device: h200_35gb
   key: entrypoints-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/entrypoints
@@ -15,8 +16,9 @@ steps:
   - pytest -v -s entrypoints/weight_transfer
 
 - label: Entrypoints Integration (LLM)
+  device: h200_35gb
   key: entrypoints-integration-llm
-  timeout_in_minutes: 40
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -37,7 +39,7 @@ steps:
 - label: Entrypoints Integration (API Server)
   key: entrypoints-integration-api-server
   device: h200_35gb
-  timeout_in_minutes: 130
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -55,8 +57,9 @@ steps:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server OpenAI - Part 1)
+  device: h200_35gb
   key: entrypoints-integration-api-server-openai-part-1
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -68,13 +71,14 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 80
+      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server OpenAI - Part 2)
+  device: h200_35gb
   key: entrypoints-integration-api-server-openai-part-2
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -92,6 +96,7 @@ steps:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server Generate)
+  device: h200_35gb
   key: entrypoints-integration-api-server-generate
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
@@ -109,11 +114,12 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (Responses API)
+  device: h200_35gb
   key: entrypoints-integration-responses-api
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
@@ -126,7 +132,7 @@ steps:
 - label: Entrypoints Integration (Speech to Text)
   device: h200_35gb
   key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -138,7 +144,7 @@ steps:
 - label: Entrypoints Integration (Multimodal)
   device: h200_35gb
   key: entrypoints-integration-multimodal
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -148,6 +154,7 @@ steps:
   - pytest -v -s entrypoints/multimodal
 
 - label: Entrypoints Integration (Pooling)
+  device: h200_35gb
   key: entrypoints-integration-pooling
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
@@ -160,7 +167,7 @@ steps:
 
 - label: OpenAI API Correctness
   key: openai-api-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
   - csrc/

.buildkite/test_areas/expert_parallelism.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: EPLB Algorithm
   key: eplb-algorithm
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -16,6 +16,7 @@ steps:
   - pytest -v -s distributed/test_eplb_utils.py
   mirror:
     amd:
+      dind: false
       device: mi300_1
       depends_on:
       - image-build-amd
@@ -27,7 +28,7 @@ steps:
 
 - label: EPLB Execution # 17min
   key: eplb-execution
-  timeout_in_minutes: 27
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -39,7 +40,7 @@ steps:
 
 - label: Elastic EP Scaling Test
   key: elastic-ep-scaling-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h100
   working_dir: "/vllm-workspace/tests"
   num_devices: 4

.buildkite/test_areas/kernels.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: vLLM IR Tests
   key: vllm-ir-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 35
   device: h200_18gb
   working_dir: "/vllm-workspace/"
   source_file_dependencies:
@@ -15,19 +15,21 @@ steps:
     - pytest -v -s tests/kernels/ir
 
 - label: Kernels Core Operation Test
+  device: h200_35gb
   key: kernels-core-operation-test
-  timeout_in_minutes: 75
+  timeout_in_minutes: 120
   source_file_dependencies:
   - csrc/
   - tests/kernels/core
   - tests/kernels/test_concat_mla_q.py
   - tests/kernels/test_fused_qk_norm_rope_gate.py
   commands:
-    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py  kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py
+    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 3
 
 - label: Kernels MiniMax Reduce RMS Test (2 GPUs)
   key: kernels-minimax-reduce-rms-test-2-gpus
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   num_devices: 2
   device: h100
   source_file_dependencies:
@@ -41,7 +43,7 @@ steps:
 
 - label: Deepseek V4 Kernel Test (H100)
   key: deepseek-v4-kernel-test-h100
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -54,7 +56,7 @@ steps:
 
 - label: Deepseek V4 Kernel Test (B200)
   key: deepseek-v4-kernel-test-b200
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   device: b200-k8s
   source_file_dependencies:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -65,7 +67,7 @@ steps:
 
 - label: Kernels Attention Test %N
   key: kernels-attention-test
-  timeout_in_minutes: 35
+  timeout_in_minutes: 65
   source_file_dependencies:
   - csrc/attention/
   - vllm/v1/attention
@@ -79,7 +81,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 90
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -106,7 +108,7 @@ steps:
 
 - label: Kernels Quantization Test %N
   key: kernels-quantization-test
-  timeout_in_minutes: 90
+  timeout_in_minutes: 60
   source_file_dependencies:
   - csrc/quantization/
   - vllm/model_executor/layers/quantization
@@ -131,7 +133,7 @@ steps:
 
 - label: Kernels MoE Test %N
   key: kernels-moe-test
-  timeout_in_minutes: 25
+  timeout_in_minutes: 50
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/
   - csrc/moe/
@@ -147,7 +149,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 65
       source_file_dependencies:
       - csrc/quantization/cutlass_w8a8/moe/
       - csrc/moe/
@@ -162,8 +164,9 @@ steps:
       - image-build-amd
 
 - label: Kernels Mamba Test
+  device: h200_35gb
   key: kernels-mamba-test
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   source_file_dependencies:
   - csrc/mamba/
   - tests/kernels/mamba
@@ -172,19 +175,19 @@ steps:
     - pytest -v -s kernels/mamba
 
 - label: Kernels KDA Test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
   device: h200_18gb
   source_file_dependencies:
-  - vllm/model_executor/layers/fla/ops/kda.py
-  - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
-  - vllm/model_executor/layers/fla/ops/l2norm.py
+  - 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: 45
+  timeout_in_minutes: 35
   device: h100
   num_devices: 1
   source_file_dependencies:
@@ -211,7 +214,7 @@ steps:
 
 - label: Kernels (B200)
   key: kernels-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 80
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   # optional: true
@@ -231,6 +234,15 @@ steps:
   - vllm/v1/attention/backends/mla/flashinfer_mla.py
   - vllm/v1/attention/selector.py
   - vllm/platforms/cuda.py
+  - vllm/model_executor/kernels/linear/cute_dsl/ll_bf16.py
+  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_dotprod.py
+  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py
+  - vllm/cute_utils/
+  - vllm/model_executor/layers/mamba/ops/gdn_chunk_cutedsl/
+  - vllm/model_executor/layers/fused_moe/router/bf16x3_router_gemm_cutedsl.py
+  - tests/kernels/mamba/test_gdn_prefill_cutedsl.py
+  - tests/kernels/test_bf16x3_router_gemm_cutedsl.py
+  - tests/kernels/test_ll_bf16_gemm.py
   - tests/kernels/test_top_k_per_row.py
   commands:
     - nvidia-smi
@@ -259,24 +271,28 @@ steps:
     - pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
     - pytest -v -s tests/kernels/moe/test_trtllm_nvfp4_moe.py
     - pytest -v -s tests/kernels/moe/test_cutedsl_moe.py
+    - pytest -v -s tests/kernels/mamba/test_gdn_prefill_cutedsl.py
+    - pytest -v -s tests/kernels/test_bf16x3_router_gemm_cutedsl.py
+    - pytest -v -s tests/kernels/test_ll_bf16_gemm.py
     # e2e
     - pytest -v -s tests/models/quantization/test_nvfp4.py
 
 - label: Kernels Helion Test
   key: kernels-helion-test
-  timeout_in_minutes: 30
+  timeout_in_minutes: 115
   device: h100
   source_file_dependencies:
   - vllm/utils/import_utils.py
   - tests/kernels/helion/
   commands:
     - pip install helion==1.1.0
-    - pytest -v -s kernels/helion/
+    - pytest -v -s kernels/helion/ --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 2
 
  
 - label: Kernels FP8 MoE Test (1xH100)
   key: kernels-fp8-moe-test-1xh100
-  timeout_in_minutes: 90
+  timeout_in_minutes: 40
   device: h100
   num_devices: 1
   optional: true
@@ -293,7 +309,7 @@ steps:
 
 - label: Kernels FP8 MoE Test (2xH100)
   key: kernels-fp8-moe-test-2xh100
-  timeout_in_minutes: 90
+  timeout_in_minutes: 45
   device: h100
   num_devices: 2
   optional: true
@@ -303,7 +319,7 @@ steps:
 
 - label: Kernels Fp4 MoE Test (B200)
   key: kernels-fp4-moe-test-b200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: b200-k8s
   num_devices: 1
   optional: true
@@ -316,7 +332,7 @@ steps:
 
 - label: Kernels FusedMoE Layer Test (2 H100s)
   key: kernels-fusedmoe-layer-test-2-h100s
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
   device: h100
   num_devices: 2
   source_file_dependencies:

.buildkite/test_areas/lm_eval.yaml

@@ -5,7 +5,7 @@ steps:
 - label: LM Eval Small Models
   device: h200_35gb
   key: lm-eval-small-models
-  timeout_in_minutes: 75
+  timeout_in_minutes: 45
   source_file_dependencies:
   - csrc/
   - vllm/model_executor/layers/quantization
@@ -56,7 +56,7 @@ steps:
 
 - label: LM Eval Small Models (1xB200)
   key: lm-eval-small-models-1xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -78,9 +78,31 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small-tp.txt
 
+- label: LM Eval PCP (4xB200)
+  key: lm-eval-pcp-4xb200
+  timeout_in_minutes: 360
+  device: b200-k8s
+  num_devices: 4
+  optional: true
+  source_file_dependencies:
+  - csrc/
+  - tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP2-PCP2-EP.yaml
+  - tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP1-PCP4-EP.yaml
+  - tests/evals/gsm8k/configs/models-pcp.txt
+  - vllm/model_executor/layers/quantization
+  - vllm/config/parallel.py
+  - vllm/distributed/parallel_state.py
+  - vllm/model_executor/layers/attention/mla_attention.py
+  - vllm/model_executor/layers/attention/pcp.py
+  - vllm/v1/worker/gpu/model_runner.py
+  - vllm/v1/worker/gpu/pcp_manager.py
+  autorun_on_main: true
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-pcp.txt
+
 - label: LM Eval Large Models EP (2xB200)
   key: lm-eval-large-models-ep-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 60
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -92,7 +114,7 @@ steps:
 
 - label: LM Eval Qwen3.5 Models (2xB200)
   key: lm-eval-qwen3-5-models-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 45
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -103,13 +125,13 @@ steps:
   - vllm/transformers_utils/configs/qwen3_5_moe.py
   - vllm/model_executor/models/qwen3_next.py
   - vllm/model_executor/models/qwen3_next_mtp.py
-  - vllm/model_executor/layers/fla/ops/
+  - vllm/third_party/flash_linear_attention/ops/
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt
 
 - label: LM Eval Large Models (8xH200)
   key: lm-eval-large-models-8xh200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
   device: h200
   optional: true
   num_devices: 8
@@ -117,8 +139,9 @@ steps:
     - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200.txt
   mirror:
     amd:
+      dind: false
       device: mi300_8
-      timeout_in_minutes: 180
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       commands:
@@ -152,7 +175,7 @@ steps:
 
 - label: LM Eval Humming f16 (A100 - TEMPORARY)
   key: lm-eval-humming-f16-a100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 75
   device: a100
   optional: true
   num_devices: 1
@@ -167,7 +190,7 @@ steps:
 
 - label: LM Eval Humming Act int8 (A100 - TEMPORARY)
   key: lm-eval-humming-act-a100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 45
   device: a100
   optional: true
   num_devices: 1
@@ -182,7 +205,7 @@ steps:
 
 - label: LM Eval Humming f16 (H100 - TEMPORARY)
   key: lm-eval-humming-f16-h100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 70
   device: h100
   optional: true
   num_devices: 1
@@ -197,7 +220,7 @@ steps:
 
 - label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
   key: lm-eval-humming-act-h100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 70
   device: h100
   optional: true
   num_devices: 1
@@ -213,7 +236,7 @@ steps:
 
 - label: LM Eval Humming f16 (B200 - TEMPORARY)
   key: lm-eval-humming-f16-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
   device: b200-k8s
   optional: true
   num_devices: 1
@@ -228,7 +251,7 @@ steps:
 
 - label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
   key: lm-eval-humming-act-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
   device: b200-k8s
   optional: true
   num_devices: 1
@@ -244,7 +267,7 @@ steps:
 
 - label: LM Eval TurboQuant KV Cache
   key: lm-eval-turboquant-kv-cache
-  timeout_in_minutes: 75
+  timeout_in_minutes: 55
   device: h200_18gb
   source_file_dependencies:
   - vllm/model_executor/layers/quantization/turboquant/
@@ -256,7 +279,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (2xH100)
   key: gpqa-eval-gpt-oss-2xh100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
   device: h100
   optional: true
   num_devices: 2
@@ -270,7 +293,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (2xB200)
   key: gpqa-eval-gpt-oss-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 30
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -284,7 +307,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (DGX Spark)
   key: gpqa-eval-gpt-oss-spark
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
   device: dgx-spark
   optional: true
   num_devices: 1
@@ -313,7 +336,7 @@ steps:
 
 - label: LM Eval KV-Offload (2xH100)
   key: kv-offload-medium
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
   device: h100
   num_devices: 2
   source_file_dependencies:
@@ -327,7 +350,7 @@ steps:
 
 - label: LM Eval KV-Offload (4xH100)
   key: kv-offload-large
-  timeout_in_minutes: 60
+  timeout_in_minutes: 40
   device: h100
   num_devices: 4
   source_file_dependencies:
@@ -341,7 +364,7 @@ steps:
 
 - label: MRCR Eval Small Models
   device: h200_35gb
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   source_file_dependencies:
   - tests/evals/mrcr/
   commands:

.buildkite/test_areas/lora.yaml

@@ -5,7 +5,7 @@ steps:
 - label: LoRA %N
   device: h200_35gb
   key: lora
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/lora
   - tests/lora
@@ -16,7 +16,7 @@ steps:
     amd:
       device: mi325_1
       working_dir: "/vllm-workspace/tests"
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
       source_file_dependencies:
       - vllm/lora
       - tests/lora
@@ -27,7 +27,7 @@ steps:
 
 - label: LoRA TP (Distributed)
   key: lora-tp-distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
   num_devices: 4
   source_file_dependencies:
   - vllm/lora

.buildkite/test_areas/misc.yaml

@@ -5,7 +5,7 @@ steps:
 - label: V1 Spec Decode
   device: h200_35gb
   key: v1-spec-decode
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   source_file_dependencies:
     - vllm/config/
     - vllm/distributed/
@@ -23,14 +23,15 @@ steps:
     - pytest -v -s -m 'not slow_test' v1/spec_decode
   mirror:
     amd:
+      dind: false
       device: mi300_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
 
 - label: V1 Sample + Logits
   key: v1-sample-logits
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
     - vllm/config/
@@ -63,8 +64,9 @@ steps:
       - image-build-amd
 
 - label: V1 Core + KV + Metrics
+  device: h200_35gb
   key: v1-core-kv-metrics
-  timeout_in_minutes: 30
+  timeout_in_minutes: 80
   source_file_dependencies:
     - vllm/config/
     - vllm/distributed/
@@ -89,6 +91,7 @@ steps:
     - tests/v1/simple_kv_offload
     - tests/v1/worker
     - tests/v1/kv_connector/unit
+    - tests/v1/ec_connector/unit
     - tests/v1/metrics
     - tests/entrypoints/openai/correctness/test_lmeval.py
   commands:
@@ -101,6 +104,7 @@ steps:
     - pytest -v -s v1/simple_kv_offload
     - pytest -v -s v1/worker
     - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
+    - pytest -v -s -m 'not cpu_test' v1/ec_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/metrics
     # Integration test for streaming correctness (requires special branch).
     - pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
@@ -108,7 +112,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
 
@@ -172,7 +176,7 @@ steps:
 
 - label: Regression
   key: regression
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/config/
@@ -195,7 +199,7 @@ steps:
 - label: Examples
   device: h200_35gb
   key: examples
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
   - vllm/entrypoints
@@ -237,7 +241,7 @@ steps:
 
 - label: Metrics, Tracing (2 GPUs)
   key: metrics-tracing-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
   num_devices: 2
   source_file_dependencies:
   - vllm/config/
@@ -281,7 +285,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 20
+      timeout_in_minutes: 45
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -292,7 +296,7 @@ steps:
 - label: Async Engine, Inputs, Utils, Worker
   device: h200_35gb
   key: async-engine-inputs-utils-worker
-  timeout_in_minutes: 50
+  timeout_in_minutes: 25
   source_file_dependencies:
   - vllm/assets/
   - vllm/config/
@@ -319,7 +323,7 @@ steps:
   key: async-engine-inputs-utils-worker-config-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 65
   source_file_dependencies:
   - vllm/assets/
   - vllm/config/
@@ -381,7 +385,7 @@ steps:
 
 - label: Batch Invariance (A100)
   key: batch-invariance-a100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: a100
   source_file_dependencies:
     - vllm/v1/attention
@@ -395,7 +399,7 @@ steps:
 
 - label: Batch Invariance (H100)
   key: batch-invariance-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: h100
   source_file_dependencies:
     - vllm/v1/attention
@@ -411,7 +415,7 @@ steps:
 
 - label: Batch Invariance (B200)
   key: batch-invariance-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 35
   device: b200-k8s
   source_file_dependencies:
     - vllm/v1/attention
@@ -430,7 +434,7 @@ steps:
 - label: Acceptance Length Test (Large Models) # optional
   device: h200_35gb
   key: acceptance-length-test-large-models
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
   gpu: h100
   optional: true
   num_gpus: 1

.buildkite/test_areas/model_executor.yaml

@@ -3,13 +3,16 @@ depends_on:
   - image-build
 steps:
 - label: Model Executor
+  device: h200_35gb
   key: model-executor
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/engine/arg_utils.py
   - vllm/config/model.py
   - vllm/model_executor
+  - vllm/model_executor/warmup
   - tests/model_executor
+  - tests/model_executor/test_jit_warmup.py
   - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
   commands:
     - apt-get update && apt-get install -y curl libsodium23
@@ -25,6 +28,7 @@ steps:
     - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
   mirror:
     amd:
+      dind: false
       device: mi300_1
       depends_on:
       - image-build-amd
@@ -32,7 +36,9 @@ steps:
       - vllm/engine/arg_utils.py
       - vllm/config/model.py
       - vllm/model_executor
+      - vllm/model_executor/warmup
       - tests/model_executor
+      - tests/model_executor/test_jit_warmup.py
       - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.py

.buildkite/test_areas/model_runner_v2.yaml

@@ -5,7 +5,7 @@ steps:
 - label: Model Runner V2 Core Tests
   device: h200_35gb
   key: model-runner-v2-core-tests
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
   source_file_dependencies:
   - vllm/v1/worker/gpu/
   - vllm/v1/worker/gpu_worker.py
@@ -27,7 +27,7 @@ steps:
 - label: Model Runner V2 Examples
   device: h200_35gb
   key: model-runner-v2-examples
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
     - vllm/v1/worker/gpu/
@@ -63,7 +63,7 @@ steps:
 
 - label: Model Runner V2 Distributed (2 GPUs)
   key: model-runner-v2-distributed-2-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -84,7 +84,7 @@ steps:
 
 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
   key: model-runner-v2-pipeline-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:

.buildkite/test_areas/models_basic.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
   key: basic-models-tests-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 25
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -17,7 +17,7 @@ steps:
 - label: Basic Models Tests (Extra Initialization) %N
   device: h200_35gb
   key: basic-models-tests-extra-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 100
   source_file_dependencies:
   - vllm/model_executor/models/
   - tests/models/test_initialization.py
@@ -27,12 +27,12 @@ steps:
     # subset of supported models (the complement of the small subset in the above
     # test.) Also run if model initialization test file is modified
     - pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-  parallelism: 2
+  parallelism: 4
 
 - label: Basic Models Tests (Other)
   device: h200_35gb
   key: basic-models-tests-other
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
   source_file_dependencies:
   - vllm/
   - tests/models/test_terratorch.py
@@ -50,7 +50,7 @@ steps:
   key: basic-models-test-other-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   source_file_dependencies:
   - vllm/
   - tests/models/test_utils.py

.buildkite/test_areas/models_distributed.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
   key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:

.buildkite/test_areas/models_language.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Language Models Tests (Standard)
   key: language-models-tests-standard
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -15,13 +15,15 @@ steps:
     - pytest -v -s models/language -m 'core_model and (not slow_test)'
   mirror:
     amd:
+      dind: false
       device: mi300_1
       depends_on:
       - image-build-amd
 
 - label: Language Models Tests (Extra Standard) %N
+  device: h200_35gb
   key: language-models-tests-extra-standard
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/model_executor/models/
   - tests/models/language/pooling/test_embedding.py
@@ -35,6 +37,7 @@ steps:
   parallelism: 2
   mirror:
     amd:
+      dind: false
       device: mi300_1
       depends_on:
       - image-build-amd
@@ -49,10 +52,10 @@ steps:
       - tests/models/language/pooling/test_classification.py
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.py
-
 - label: Language Models Tests (Hybrid) %N
+  device: h200_35gb
   key: language-models-tests-hybrid
-  timeout_in_minutes: 75
+  timeout_in_minutes: 65
   source_file_dependencies:
   - vllm/
   - tests/models/language/generation
@@ -61,13 +64,13 @@ steps:
     # Note: also needed to run plamo2 model in vLLM
     - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
     - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
-    # Shard hybrid language model tests
-    - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+    # Shard the hybrid language model tests that are numerically stable on Hopper.
+    - pytest -v -s models/language/generation -m hybrid_model -k 'not granite-4.0-tiny-preview' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
   parallelism: 2
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 90
+      timeout_in_minutes: 70
       depends_on:
       - image-build-amd
       commands:
@@ -75,10 +78,24 @@ steps:
       - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
       - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
+# Granite 4 hybrid generation is sensitive to hardware-specific Triton SSD
+# autotuning (https://github.com/vllm-project/vllm/issues/25194). Keep this one
+# correctness test on L4 until its H200 output matches the Transformers reference.
+- label: Language Models Tests (Granite L4 Compatibility)
+  key: language-models-tests-granite-l4-compatibility
+  timeout_in_minutes: 65
+  source_file_dependencies:
+  - vllm/
+  - tests/models/language/generation
+  commands:
+    - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
+    - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
+    - pytest -v -s models/language/generation -m hybrid_model -k 'granite-4.0-tiny-preview'
+
 - label: Language Models Test (Extended Generation) # 80min
   device: h200_35gb
   key: language-models-test-extended-generation
-  timeout_in_minutes: 110
+  timeout_in_minutes: 65
   optional: true
   source_file_dependencies:
   - vllm/
@@ -92,7 +109,7 @@ steps:
 
 - label: Language Models Test (PPL)
   key: language-models-test-ppl
-  timeout_in_minutes: 110
+  timeout_in_minutes: 30
   device: h200_18gb
   optional: true
   source_file_dependencies:
@@ -101,10 +118,10 @@ steps:
   commands:
     - pytest -v -s models/language/generation_ppl_test
 
-- label: Language Models Test (Extended Pooling)  # 36min
+- label: Language Models Test (Extended Pooling)
   device: h200_35gb
   key: language-models-test-extended-pooling
-  timeout_in_minutes: 50
+  timeout_in_minutes: 120
   optional: true
   source_file_dependencies:
   - vllm/
@@ -114,13 +131,13 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 100
+      timeout_in_minutes: 120
       depends_on:
       - image-build-amd
 
 - label: Language Models Test (MTEB)
   key: language-models-test-mteb
-  timeout_in_minutes: 110
+  timeout_in_minutes: 45
   device: h200_18gb
   optional: true
   source_file_dependencies:

.buildkite/test_areas/models_multimodal.yaml

@@ -20,7 +20,7 @@ steps:
 
 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
   key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -38,7 +38,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
   device: h200_35gb
   key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
@@ -54,7 +54,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
   device: h200_35gb
   key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
@@ -69,22 +69,23 @@ steps:
       depends_on:
       - image-build-amd
 
-- label: Multi-Modal Processor (CPU)
+- label: Multi-Modal Processor (CPU) %N
   key: multi-modal-processor-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 60
+  timeout_in_minutes: 125
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   - tests/models/registry.py
   device: cpu-medium
   commands:
-    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
+    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+  parallelism: 4
 
 - label: Multi-Modal Processor # 44min
   key: multi-modal-processor
-  timeout_in_minutes: 60
+  timeout_in_minutes: 65
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -96,7 +97,7 @@ steps:
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
   device: h200_35gb
   key: multi-modal-accuracy-eval-small-models
-  timeout_in_minutes: 70
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
   source_file_dependencies:
   - vllm/multimodal/
@@ -106,6 +107,7 @@ steps:
   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-mm-small.txt --tp-size=1
   mirror:
     amd:
+      dind: false
       device: mi300_1
       depends_on:
       - image-build-amd
@@ -117,6 +119,7 @@ steps:
       - vllm/model_executor/model_loader/
 
 - label: Multi-Modal Models (Extended Generation 1)
+  device: h200_35gb
   key: multi-modal-models-extended-generation-1
   optional: true
   source_file_dependencies:
@@ -164,7 +167,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
       source_file_dependencies:

.buildkite/test_areas/plugins.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Plugin Tests (2 GPUs)
   key: plugin-tests-2-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:

.buildkite/test_areas/pytorch.yaml

@@ -5,7 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
   device: h200_35gb
   key: pytorch-compilation-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 90
   source_file_dependencies:
     - vllm/__init__.py
     - vllm/_aiter_ops.py
@@ -78,7 +78,7 @@ steps:
 
 - label: PyTorch Compilation Passes Unit Tests
   key: pytorch-compilation-passes-unit-tests
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
   source_file_dependencies:
     - vllm/__init__.py
     - vllm/_aiter_ops.py
@@ -109,14 +109,16 @@ steps:
   - pytest -s -v compile/passes --ignore compile/passes/distributed
   mirror:
     amd:
+      dind: false
       device: mi300_1
-      timeout_in_minutes: 180
+      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: PyTorch Fullgraph Smoke Test
+  device: h200_35gb
   key: pytorch-fullgraph-smoke-test
-  timeout_in_minutes: 35
+  timeout_in_minutes: 90
   source_file_dependencies:
   - vllm/__init__.py
   - vllm/_aiter_ops.py
@@ -148,11 +150,46 @@ steps:
   # as it is a heavy test that is covered in other steps.
   # Use `find` to launch multiple instances of pytest so that
   # they do not suffer from https://github.com/vllm-project/vllm/issues/28965
-  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
+  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_cudagraph.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
+
+# Hopper-only DeepSeek-V2-Lite cases in this file require two 29.3-GiB model
+# instances and cannot fit a 35GB MIG slice. L4 retains the original coverage:
+# those SM90 cases skip while the architecture-compatible cases still run.
+- label: PyTorch Fullgraph CUDAGraph (L4 Compatibility)
+  key: pytorch-fullgraph-cudagraph-l4-compatibility
+  timeout_in_minutes: 60
+  source_file_dependencies:
+  - vllm/__init__.py
+  - vllm/_aiter_ops.py
+  - vllm/_custom_ops.py
+  - vllm/compilation/
+  - vllm/config/
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/env_override.py
+  - vllm/envs.py
+  - vllm/forward_context.py
+  - vllm/inputs/
+  - vllm/ir/
+  - vllm/kernels/
+  - vllm/logger.py
+  - vllm/model_executor/
+  - vllm/multimodal/
+  - vllm/platforms/
+  - vllm/plugins/
+  - vllm/sampling_params.py
+  - vllm/sequence.py
+  - vllm/transformers_utils/
+  - vllm/triton_utils/
+  - vllm/utils/
+  - vllm/v1/
+  - tests/compile
+  commands:
+  - pytest -s -v compile/fullgraph/test_full_cudagraph.py
 
 - label: PyTorch Fullgraph
   key: pytorch-fullgraph
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: h200_18gb
   source_file_dependencies:
   - vllm/__init__.py
@@ -197,6 +234,7 @@ steps:
   - bash standalone_tests/pytorch_nightly_dependency.sh
   mirror:
     amd:
+      dind: false
       device: mi300_1
       depends_on:
       - image-build-amd

.buildkite/test_areas/quantization.yaml

@@ -3,8 +3,11 @@ depends_on:
   - image-build
 steps:
 - label: Quantization
+  device: h200_35gb
   key: quantization
-  timeout_in_minutes: 90
+  timeout_in_minutes: 75
+  env:
+    VLLM_USE_V2_MODEL_RUNNER: "0"
   source_file_dependencies:
   - csrc/
   - vllm/model_executor/layers/quantization
@@ -19,11 +22,15 @@ steps:
   # TODO(jerryzh168): resolve the above comment
   - uv pip install --system torchao==0.17.0 --index-url https://download.pytorch.org/whl/cu130
   - uv pip install --system conch-triton-kernels
-  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
+  # The SM90-only checkpoint currently contains a removed weight_chan_scale
+  # parameter. It was not exercised by the previous L4 job.
+  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py -k 'not test_compressed_tensors_w4a8_fp8' --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 8
 
 - label: Quantized Fusions
+  device: h200_35gb
   key: quantized-fusions
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   source_file_dependencies:
   - tests/fusion
   - vllm/model_executor/layers/fusion
@@ -35,7 +42,7 @@ steps:
 
 - label: Quantized MoE Test (B200)
   key: quantized-moe-test-b200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 120
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   source_file_dependencies:
@@ -52,10 +59,14 @@ steps:
     - pytest -s -v tests/quantization/test_blackwell_moe.py
 
 - label: Quantized Models Test
+  device: h200_35gb
   key: quantized-models-test
-  timeout_in_minutes: 60
+  timeout_in_minutes: 65
+  env:
+    VLLM_USE_V2_MODEL_RUNNER: "0"
   source_file_dependencies:
   - vllm/model_executor/layers/quantization
   - tests/models/quantization
   commands:
-    - pytest -v -s models/quantization
+    - pytest -v -s models/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 3

.buildkite/test_areas/rust_frontend.yaml

@@ -3,11 +3,16 @@ depends_on:
   - image-build
 steps:
 - label: Rust Frontend OpenAI Coverage
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - rust/
+  - build_rust.sh
+  - tools/build_rust.py
+  - rust-toolchain.toml
+  - .buildkite/scripts/rust-coverage.sh
+  - codecov.yml
   - vllm/benchmarks/
   - vllm/entrypoints/openai/
   - vllm/entrypoints/serve/
@@ -23,6 +28,7 @@ steps:
   - tests/entrypoints/openai/test_uds.py
   - tests/v1/sample/test_logprobs_e2e.py
   commands:
+  - . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
@@ -38,11 +44,16 @@ steps:
   - pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
 
 - label: Rust Frontend Serve/Admin Coverage
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - rust/
+  - build_rust.sh
+  - tools/build_rust.py
+  - rust-toolchain.toml
+  - .buildkite/scripts/rust-coverage.sh
+  - codecov.yml
   - vllm/entrypoints/openai/
   - vllm/entrypoints/serve/
   - vllm/v1/engine/
@@ -54,6 +65,7 @@ steps:
   # - tests/entrypoints/serve/dev/test_sleep.py
   - tests/entrypoints/serve/tokenize/test_tokenization.py
   commands:
+  - . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
@@ -67,39 +79,57 @@ steps:
   - pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
 
 - label: Rust Frontend Core Correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - rust/
+  - build_rust.sh
+  - tools/build_rust.py
+  - rust-toolchain.toml
+  - .buildkite/scripts/rust-coverage.sh
+  - codecov.yml
   - vllm/entrypoints/openai/
   - tests/utils.py
   - tests/entrypoints/openai/correctness/test_lmeval.py
   commands:
+  - . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
 
 - label: Rust Frontend Tool Use
-  timeout_in_minutes: 60
+  device: h200_35gb
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - rust/
+  - build_rust.sh
+  - tools/build_rust.py
+  - rust-toolchain.toml
+  - .buildkite/scripts/rust-coverage.sh
+  - codecov.yml
   - vllm/entrypoints/openai/
   - vllm/tool_parsers/
   - tests/utils.py
   - tests/tool_use/
   commands:
+  - . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
 
 - label: Rust Frontend Distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   num_devices: 4
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - rust/
+  - build_rust.sh
+  - tools/build_rust.py
+  - rust-toolchain.toml
+  - .buildkite/scripts/rust-coverage.sh
+  - codecov.yml
   - vllm/distributed/
   - vllm/engine/
   - vllm/executor/
@@ -110,6 +140,7 @@ steps:
   - tests/v1/distributed/test_hybrid_lb_dp.py
   - tests/v1/distributed/test_internal_lb_dp.py
   commands:
+  - . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - export NCCL_CUMEM_HOST_ENABLE=0

.buildkite/test_areas/rust_frontend_cargo.yaml

@@ -4,7 +4,7 @@ steps:
 - label: Rust Frontend Cargo Style + Clippy
   key: rust-frontend-cargo-style-clippy
   depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: cpu-medium
   no_plugin: true
   source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
 - label: Rust Frontend Cargo Tests
   key: rust-frontend-cargo-tests
   depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: cpu-medium
   no_plugin: true
   source_file_dependencies:
@@ -26,5 +26,7 @@ steps:
   - rust-toolchain.toml
   - .buildkite/test_areas/rust_frontend_cargo.yaml
   - .buildkite/scripts/run-rust-frontend-cargo-ci.sh
+  - .buildkite/scripts/rust-coverage.sh
+  - codecov.yml
   commands:
   - .buildkite/scripts/run-rust-frontend-cargo-ci.sh test

.buildkite/test_areas/samplers.yaml

@@ -5,7 +5,7 @@ steps:
 - label: Samplers Test
   device: h200_35gb
   key: samplers-test
-  timeout_in_minutes: 75
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/model_executor/layers
   - vllm/sampling_metadata.py

.buildkite/test_areas/spec_decode.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Spec Decode Eagle
   key: spec-decode-eagle
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -15,7 +15,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 45
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
 
 - label: Spec Decode Eagle Nightly B200
   key: spec-decode-eagle-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
 
 - label: Spec Decode Speculators + MTP
   key: spec-decode-speculators-mtp
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
   
 - label: Spec Decode Ngram + Suffix
   key: spec-decode-ngram-suffix
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 55
       # TODO(akaratza): Test after Torch >= 2.12 bump
       soft_fail: true
       depends_on:
@@ -109,7 +109,7 @@ steps:
 
 - label: Spec Decode Draft Model
   key: spec-decode-draft-model
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -134,7 +134,7 @@ steps:
 
 - label: Spec Decode Draft Model Nightly B200
   key: spec-decode-draft-model-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -146,7 +146,7 @@ steps:
 
 - label: Speculators Correctness
   key: speculators-correctness
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
   device: h100
   optional: true
   num_devices: 1
@@ -159,7 +159,7 @@ steps:
     - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
 
 - label: Spec Decode MTP hybrid (B200)
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: b200-k8s
   optional: true
   source_file_dependencies:

.buildkite/test_areas/weight_loading.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Weight Loading Multiple GPU  # 33min
   key: weight-loading-multiple-gpu
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   optional: true
@@ -15,6 +15,7 @@ steps:
     - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models.txt
   mirror:
     amd:
+      dind: false
       device: mi300_2
       depends_on:
       - image-build-amd

.github/CODEOWNERS

@@ -47,6 +47,7 @@
 
 # Rust Frontend
 /rust/ @BugenZhao @njhill
+/rust/src/bench @esmeetu
 /build_rust.sh @BugenZhao @njhill
 /rust-toolchain.toml @BugenZhao @njhill
 /.buildkite/test_areas/rust* @BugenZhao @njhill
@@ -172,7 +173,7 @@ mkdocs.yaml @hmellor
 # Kernels
 /vllm/v1/attention/ops/chunked_prefill_paged_decode.py @tdoublep
 /vllm/v1/attention/ops/triton_unified_attention.py @tdoublep
-/vllm/model_executor/layers/fla @ZJY0516 @vadiklyutiy
+/vllm/third_party/flash_linear_attention @ZJY0516 @vadiklyutiy
 
 # ROCm related: specify owner with write access to notify AMD folks for careful code review
 /vllm/**/*rocm* @tjtanaa @dllehr-amd

.github/workflows/issue_autolabel.yml

@@ -323,7 +323,7 @@ jobs:
             // {users} will be replaced with @mentions
             const ccConfig = {
               rocm: {
-                users: ['hongxiayang', 'tjtanaa', 'vllmellm'],
+                users: ['hongxiayang', 'tjtanaa', 'vllmellm', 'giuseppegrossi'],
                 message: 'CC {users} for ROCm-related issue',
               },
               mistral: {

.gitignore

@@ -18,6 +18,9 @@ vllm/third_party/deep_gemm/
 # fmha_sm100 vendored package built from source
 vllm/third_party/fmha_sm100/
 
+# tml-fa4 vendored package built from source
+vllm/third_party/tml_fa4/
+
 # triton jit
 .triton
 
@@ -254,3 +257,4 @@ vllm/grpc/vllm_engine_pb2.pyi
 
 # Ignore generated cpu headers 
 csrc/cpu/cpu_attn_dispatch_generated.h
+rust-coverage-tools/

.pre-commit-config.yaml

@@ -30,7 +30,7 @@ repos:
   - id: markdownlint-cli2
     language_version: lts
     args: [--fix]
-    exclude: ^CLAUDE\.md$
+    exclude: (^|/)CLAUDE\.md$
 - repo: https://github.com/rhysd/actionlint
   rev: v1.7.7
   hooks:
@@ -210,7 +210,7 @@ repos:
     name: Check SPDX headers
     entry: python tools/pre_commit/check_spdx_header.py
     language: python
-    types: [python]
+    types_or: [python, rust, proto]
   - id: check-root-lazy-imports
     name: Check root lazy imports
     entry: python tools/pre_commit/check_init_lazy_imports.py

.yapfignore

@@ -1,2 +0,0 @@
-collect_env.py
-vllm/model_executor/layers/fla/ops/*.py

CMakeLists.txt

@@ -70,6 +70,15 @@ endif()
 #
 set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
 set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
+# TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
+# version check would always warn. Only treat it as a nightly build when the
+# value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
+# must NOT suppress the warning for normal builds).
+if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
+  set(TORCH_NIGHTLY_BUILD TRUE)
+else()
+  set(TORCH_NIGHTLY_BUILD FALSE)
+endif()
 
 #
 # Try to find python package with an executable that exactly matches
@@ -175,7 +184,7 @@ endif()
 if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
   set(VLLM_GPU_LANG "CUDA")
 
-  if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
+  if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
     message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
       "expected for CUDA build, saw ${Torch_VERSION} instead.")
   endif()
@@ -188,7 +197,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
   enable_language(HIP)
 
   # ROCm 5.X and 6.X
-  if (ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
+  if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
       Torch_VERSION VERSION_LESS ${TORCH_SUPPORTED_VERSION_ROCM})
     message(WARNING "Pytorch version >= ${TORCH_SUPPORTED_VERSION_ROCM} "
       "expected for ROCm build, saw ${Torch_VERSION} instead.")
@@ -381,6 +390,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     "csrc/libtorch_stable/cuda_view.cu"
     "csrc/libtorch_stable/cuda_utils_kernels.cu"
     "csrc/libtorch_stable/activation_kernels.cu"
+    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
     "csrc/libtorch_stable/quantization/activation_kernels.cu"
     "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
     "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -1354,6 +1364,7 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
   set(VLLM_ROCM_EXT_SRC
     "csrc/rocm/torch_bindings.cpp"
     "csrc/rocm/skinny_gemms.cu"
+    "csrc/rocm/skinny_gemms_int4.cu"
     "csrc/rocm/attention.cu")
 
   set(VLLM_ROCM_HAS_GFX1100 OFF)
@@ -1397,6 +1408,7 @@ if (VLLM_GPU_LANG STREQUAL "CUDA")
     include(cmake/external_projects/fmha_sm100.cmake)
     include(cmake/external_projects/flashmla.cmake)
     include(cmake/external_projects/qutlass.cmake)
+    include(cmake/external_projects/tml_fa4.cmake)
 
     # vllm-flash-attn should be last as it overwrites some CMake functions
     include(cmake/external_projects/vllm_flash_attn.cmake)

README.md

@@ -48,7 +48,7 @@ vLLM is flexible and easy to use with:
 - Tool calling and reasoning parsers
 - OpenAI-compatible API server, plus Anthropic Messages API and gRPC support
 - Efficient multi-LoRA support for dense and MoE layers
-- Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
+- Support for NVIDIA GPUs, AMD GPUs, Intel GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
 
 vLLM seamlessly supports 200+ model architectures on Hugging Face, including:
 

benchmarks/attention_benchmarks/benchmark.py

@@ -75,7 +75,11 @@ def run_mla_benchmark(config: BenchmarkConfig, **kwargs) -> BenchmarkResult:
     from mla_runner import run_mla_benchmark as run_mla
 
     return run_mla(
-        config.backend, config, prefill_backend=config.prefill_backend, **kwargs
+        config.backend,
+        config,
+        prefill_backend=config.prefill_backend,
+        sparse_mla_force_mqa=config.sparse_mla_force_mqa,
+        **kwargs,
     )
 
 
@@ -592,6 +596,30 @@ def main():
         default="profile",
         help="Output file name for ncu profile (default: 'profile').",
     )
+    parser.add_argument(
+        "--torch-profile",
+        action="store_true",
+        default=False,
+        help="Collect a PyTorch profiler Chrome trace for each benchmark run.",
+    )
+    parser.add_argument(
+        "--torch-profile-dir",
+        default=None,
+        help="Directory for PyTorch profiler traces.",
+    )
+    parser.add_argument(
+        "--torch-profile-iters",
+        type=int,
+        default=3,
+        help="Number of forward passes to record per PyTorch profiler trace.",
+    )
+    parser.add_argument(
+        "--sparse-mla-mha-variants",
+        nargs="+",
+        default=None,
+        choices=["dense_mha", "mqa"],
+        help="Sparse MLA variants to run in mha_vs_mqa mode. Defaults to both.",
+    )
 
     # Parameter sweep (use YAML config for advanced sweeps)
     parser.add_argument(
@@ -641,6 +669,7 @@ def main():
 
         # Prefill backends (e.g., ["fa3", "fa4"])
         args.prefill_backends = yaml_config.get("prefill_backends", None)
+        args.prefill_backend = yaml_config.get("prefill_backend", None)
 
         # FP8 output benchmark knobs; CLI wins.
         if args.fp8_output_scale is None:
@@ -683,6 +712,9 @@ def main():
             args.num_q_heads = model.get("num_q_heads", args.num_q_heads)
             args.num_kv_heads = model.get("num_kv_heads", args.num_kv_heads)
             args.block_size = model.get("block_size", args.block_size)
+            args.max_model_len = model.get(
+                "max_model_len", getattr(args, "max_model_len", None)
+            )
             # MLA-specific dimensions
             args.kv_lora_rank = model.get("kv_lora_rank", args.kv_lora_rank)
             args.qk_nope_head_dim = model.get("qk_nope_head_dim", args.qk_nope_head_dim)
@@ -701,6 +733,21 @@ def main():
             args.cuda_graphs = yaml_config["cuda_graphs"]
         if "ncu_profile" in yaml_config:
             args.ncu_profile = yaml_config["ncu_profile"]
+        if "torch_profile" in yaml_config:
+            args.torch_profile = yaml_config["torch_profile"]
+        if "torch_profile_dir" in yaml_config:
+            args.torch_profile_dir = yaml_config["torch_profile_dir"]
+        if "torch_profile_iters" in yaml_config:
+            args.torch_profile_iters = yaml_config["torch_profile_iters"]
+        args.sparse_mla_topk_pattern = yaml_config.get(
+            "sparse_mla_topk_pattern", "random"
+        )
+        args.sparse_mla_dense_mha_max_seq_len = yaml_config.get(
+            "sparse_mla_dense_mha_max_seq_len", None
+        )
+        args.sparse_mla_mha_variants = yaml_config.get(
+            "sparse_mla_mha_variants", args.sparse_mla_mha_variants
+        )
 
         # Parameter sweep configuration
         if "parameter_sweep" in yaml_config:
@@ -842,8 +889,6 @@ def main():
                         num_kv_heads=args.num_kv_heads,
                         block_size=args.block_size,
                         device=args.device,
-                        repeats=args.repeats,
-                        warmup_iters=args.warmup_iters,
                         profile_memory=args.profile_memory,
                         kv_cache_dtype=args.kv_cache_dtype,
                         use_cuda_graphs=args.cuda_graphs,
@@ -1063,6 +1108,133 @@ def main():
                     f"\n  [yellow]Prefill always faster for batch_size={bs}[/]"
                 )
 
+    # Handle MHA vs MQA comparison mode for sparse MLA
+    elif hasattr(args, "mode") and args.mode == "mha_vs_mqa":
+        console.print("[yellow]Mode: MHA vs MQA comparison for sparse MLA[/]")
+
+        sparse_mla_topk_pattern = getattr(args, "sparse_mla_topk_pattern", "random")
+        dense_mha_max_seq_len = getattr(args, "sparse_mla_dense_mha_max_seq_len", None)
+        prefill_backend = getattr(args, "prefill_backend", None)
+        if prefill_backend:
+            console.print(f"Prefill backend: {prefill_backend}")
+        available_variants = [
+            ("dense_mha", False, "dense"),
+            ("mqa", True, "auto"),
+        ]
+        requested_variants = getattr(args, "sparse_mla_mha_variants", None)
+        if requested_variants is not None:
+            valid_variants = {label for label, _, _ in available_variants}
+            invalid_variants = sorted(set(requested_variants) - valid_variants)
+            if invalid_variants:
+                raise ValueError(
+                    "Invalid sparse_mla_mha_variants entries: "
+                    f"{invalid_variants}. Valid variants are: "
+                    f"{sorted(valid_variants)}"
+                )
+            requested_variant_set = set(requested_variants)
+            variants = [
+                variant
+                for variant in available_variants
+                if variant[0] in requested_variant_set
+            ]
+        else:
+            variants = available_variants
+        formatter = ResultsFormatter(console)
+        total = 0
+        for spec in args.batch_specs:
+            q_len = max(request.q_len for request in parse_batch_spec(spec))
+            for variant_label, _, _ in variants:
+                if (
+                    variant_label == "dense_mha"
+                    and dense_mha_max_seq_len is not None
+                    and q_len > dense_mha_max_seq_len
+                ):
+                    continue
+                total += len(backends)
+
+        with tqdm(total=total, desc="Benchmarking") as pbar:
+            for spec in args.batch_specs:
+                q_len = max(request.q_len for request in parse_batch_spec(spec))
+                for backend in backends:
+                    for variant_label, force_mqa, mha_mode in variants:
+                        if (
+                            variant_label == "dense_mha"
+                            and dense_mha_max_seq_len is not None
+                            and q_len > dense_mha_max_seq_len
+                        ):
+                            continue
+                        config = BenchmarkConfig(
+                            backend=f"{backend}_{variant_label}",
+                            batch_spec=spec,
+                            num_layers=args.num_layers,
+                            head_dim=args.head_dim,
+                            num_q_heads=args.num_q_heads,
+                            num_kv_heads=args.num_kv_heads,
+                            block_size=args.block_size,
+                            device=args.device,
+                            max_model_len=getattr(args, "max_model_len", None),
+                            kv_cache_dtype=args.kv_cache_dtype,
+                            profile_memory=args.profile_memory,
+                            use_cuda_graphs=args.cuda_graphs,
+                            ncu_profile=args.ncu_profile,
+                            torch_profile=args.torch_profile,
+                            torch_profile_dir=args.torch_profile_dir,
+                            torch_profile_iters=args.torch_profile_iters,
+                            warmup_ms=args.warmup_ms,
+                            kv_lora_rank=getattr(args, "kv_lora_rank", None),
+                            qk_nope_head_dim=getattr(args, "qk_nope_head_dim", None),
+                            qk_rope_head_dim=getattr(args, "qk_rope_head_dim", None),
+                            v_head_dim=getattr(args, "v_head_dim", None),
+                            sparse_mla_force_mqa=force_mqa,
+                            sparse_mla_mha_mode=mha_mode,
+                            sparse_mla_dense_mha_max_seq_len=dense_mha_max_seq_len,
+                            sparse_mla_topk_pattern=sparse_mla_topk_pattern,
+                            prefill_backend=prefill_backend,
+                        )
+
+                        # run_mla_benchmark needs the real backend name
+                        from mla_runner import run_mla_benchmark as run_mla
+
+                        run_label = f"{backend}_{variant_label} {spec}"
+                        pbar.set_postfix_str(run_label)
+
+                        try:
+                            result = run_mla(
+                                backend,
+                                config,
+                                prefill_backend=prefill_backend,
+                                sparse_mla_force_mqa=force_mqa,
+                            )
+                        except Exception as e:
+                            result = BenchmarkResult(
+                                config=config,
+                                mean_time=float("inf"),
+                                median_time=float("inf"),
+                                std_time=0,
+                                min_time=float("inf"),
+                                max_time=float("inf"),
+                                error=str(e),
+                            )
+
+                        all_results.append(result)
+                        if args.output_csv:
+                            formatter.save_csv(all_results, args.output_csv)
+                        if args.output_json:
+                            formatter.save_json(all_results, args.output_json)
+
+                        if not result.success:
+                            console.print(
+                                f"[red]Error {backend}_{variant_label} "
+                                f"{spec}: {result.error}[/]"
+                            )
+
+                        pbar.update(1)
+
+        # Display results with variant labels as separate "backends"
+        console.print("\n[bold green]MHA vs MQA Results:[/]")
+        variant_backends = [f"{b}_{v}" for b in backends for v, _, _ in variants]
+        formatter.print_table(all_results, variant_backends)
+
     # Handle model parameter sweep mode
     elif hasattr(args, "model_parameter_sweep") and args.model_parameter_sweep:
         # Model parameter sweep

benchmarks/attention_benchmarks/common.py

@@ -4,8 +4,10 @@
 """Common utilities for attention benchmarking."""
 
 import csv
+import gc
 import json
 import math
+from collections.abc import Sequence
 from dataclasses import asdict, dataclass
 from pathlib import Path
 from typing import Any
@@ -44,10 +46,13 @@ def run_do_bench(
     kwargs: dict[str, Any] = {"return_mode": "all"}
     if use_cuda_graphs:
         result = triton.testing.do_bench_cudagraph(benchmark_fn, **kwargs)
+        gc.collect()
+        torch.accelerator.empty_cache()
     else:
         if warmup_ms is not None:
             kwargs["warmup"] = warmup_ms
         result = triton.testing.do_bench(benchmark_fn, **kwargs)
+    torch.accelerator.synchronize()
     return result
 
 
@@ -91,42 +96,6 @@ except ImportError:
     AttentionLayerBase = object  # Fallback
 
 
-class MockKVBProj:
-    """Mock KV projection layer for MLA prefill mode.
-
-    Mimics ColumnParallelLinear behavior for kv_b_proj in MLA backends.
-    Projects kv_c_normed to [qk_nope_head_dim + v_head_dim] per head.
-    """
-
-    def __init__(self, num_heads: int, qk_nope_head_dim: int, v_head_dim: int):
-        self.num_heads = num_heads
-        self.qk_nope_head_dim = qk_nope_head_dim
-        self.v_head_dim = v_head_dim
-        self.out_dim = qk_nope_head_dim + v_head_dim
-        self.weight = torch.empty(0, dtype=torch.bfloat16)
-
-    def __call__(self, x: torch.Tensor) -> tuple[torch.Tensor]:
-        """
-        Project kv_c_normed to output space.
-
-        Args:
-            x: Input tensor [num_tokens, kv_lora_rank]
-
-        Returns:
-            Tuple containing output tensor
-                [num_tokens, num_heads, qk_nope_head_dim + v_head_dim]
-        """
-        num_tokens = x.shape[0]
-        result = torch.randn(
-            num_tokens,
-            self.num_heads,
-            self.out_dim,
-            device=x.device,
-            dtype=x.dtype,
-        )
-        return (result,)  # Return as tuple to match ColumnParallelLinear API
-
-
 class MockIndexer:
     """Mock Indexer for sparse MLA backends.
 
@@ -158,6 +127,60 @@ class MockIndexer:
         )
         self.topk_indices_buffer[:num_tokens] = indices
 
+    def fill_indices(
+        self,
+        num_tokens: int,
+        max_kv_len: int,
+        pattern: str = "random",
+        requests: Sequence[Any] | None = None,
+    ):
+        if pattern == "random":
+            self.fill_random_indices(num_tokens, max_kv_len)
+            return
+        if pattern == "prefix":
+            indices = torch.arange(
+                self.topk_tokens,
+                dtype=torch.int32,
+                device=self.topk_indices_buffer.device,
+            )
+            indices = (indices % max_kv_len).expand(num_tokens, -1)
+            self.topk_indices_buffer[:num_tokens] = indices
+            return
+        if pattern == "sliding_window":
+            if requests is None:
+                start = max(max_kv_len - self.topk_tokens, 0)
+                indices = torch.arange(
+                    start,
+                    start + self.topk_tokens,
+                    dtype=torch.int32,
+                    device=self.topk_indices_buffer.device,
+                )
+                indices = indices.clamp(max=max_kv_len - 1).expand(num_tokens, -1)
+                self.topk_indices_buffer[:num_tokens] = indices
+                return
+
+            rows = []
+            offsets = torch.arange(
+                self.topk_tokens,
+                dtype=torch.int32,
+                device=self.topk_indices_buffer.device,
+            ) - (self.topk_tokens - 1)
+            for request in requests:
+                q_len = request.q_len
+                kv_len = request.kv_len
+                context_len = kv_len - q_len
+                positions = torch.arange(
+                    context_len,
+                    kv_len,
+                    dtype=torch.int32,
+                    device=self.topk_indices_buffer.device,
+                )
+                row_indices = positions[:, None] + offsets[None, :]
+                rows.append(row_indices.clamp(min=0, max=kv_len - 1))
+            self.topk_indices_buffer[:num_tokens] = torch.cat(rows, dim=0)
+            return
+        raise ValueError(f"Unknown sparse MLA topk pattern: {pattern}")
+
 
 class MockLayer(AttentionLayerBase):
     """Mock attention layer with scale parameters and impl.
@@ -252,10 +275,14 @@ class BenchmarkConfig:
     num_kv_heads: int
     block_size: int
     device: str
+    max_model_len: int | None = None
     dtype: torch.dtype = torch.float16
     profile_memory: bool = False
-    use_cuda_graphs: bool = False
+    use_cuda_graphs: bool = True
     ncu_profile: bool = False
+    torch_profile: bool = False
+    torch_profile_dir: str | None = None
+    torch_profile_iters: int = 3
     warmup_ms: int | None = None
 
     # "auto" or "fp8"
@@ -271,6 +298,10 @@ class BenchmarkConfig:
     # Backend-specific tuning
     num_kv_splits: int | None = None  # CUTLASS MLA
     reorder_batch_threshold: int | None = None  # FlashAttn MLA, FlashMLA
+    sparse_mla_force_mqa: bool = False  # Force MQA path for sparse MLA
+    sparse_mla_mha_mode: str = "auto"  # "auto" or "dense"
+    sparse_mla_dense_mha_max_seq_len: int | None = None
+    sparse_mla_topk_pattern: str = "random"  # "random", "prefix", "sliding_window"
     num_splits: int | None = None  # FlashAttention split-K (0=auto, 1=disabled)
 
 

benchmarks/attention_benchmarks/configs/mla_sparse_mha_vs_mqa.yaml

@@ -0,0 +1,474 @@
+# Sparse MLA benchmark: forward_mha vs forward_mqa
+#
+# Usage:
+#   python benchmark.py --config configs/mla_sparse_mha_vs_mqa.yaml
+#
+# Heatmap grid:
+#   - batch_size: 1, 2, 4, 8, 16, 32
+#   - seq_len: 32, 64, 128, 256, 512, 1024, 2048
+#   - q_len: powers of two through seq_len
+#
+# Specs with q_len < seq_len include context; the q_len == seq_len diagonal
+# covers pure prefill.
+# The model shape below is the DP case. For the TP8 run, manually change
+# model.num_q_heads from 128 to 16 before rerunning this benchmark.
+
+mode: mha_vs_mqa
+
+model:
+  name: "deepseek-v3"
+  num_layers: 60
+  num_q_heads: 128
+  num_kv_heads: 1
+  head_dim: 576
+  kv_lora_rank: 512
+  qk_nope_head_dim: 128
+  qk_rope_head_dim: 64
+  v_head_dim: 128
+  block_size: 128
+  max_model_len: 2048
+
+batch_specs:
+  # Batch size 1
+  # seq_len = 32
+  - "1q1s32"
+  - "1q2s32"
+  - "1q4s32"
+  - "1q8s32"
+  - "1q16s32"
+  - "1q32"
+  # seq_len = 64
+  - "1q1s64"
+  - "1q2s64"
+  - "1q4s64"
+  - "1q8s64"
+  - "1q16s64"
+  - "1q32s64"
+  - "1q64"
+  # seq_len = 128
+  - "1q1s128"
+  - "1q2s128"
+  - "1q4s128"
+  - "1q8s128"
+  - "1q16s128"
+  - "1q32s128"
+  - "1q64s128"
+  - "1q128"
+  # seq_len = 256
+  - "1q1s256"
+  - "1q2s256"
+  - "1q4s256"
+  - "1q8s256"
+  - "1q16s256"
+  - "1q32s256"
+  - "1q64s256"
+  - "1q128s256"
+  - "1q256"
+  # seq_len = 512
+  - "1q1s512"
+  - "1q2s512"
+  - "1q4s512"
+  - "1q8s512"
+  - "1q16s512"
+  - "1q32s512"
+  - "1q64s512"
+  - "1q128s512"
+  - "1q256s512"
+  - "1q512"
+  # seq_len = 1024
+  - "1q1s1024"
+  - "1q2s1024"
+  - "1q4s1024"
+  - "1q8s1024"
+  - "1q16s1024"
+  - "1q32s1024"
+  - "1q64s1024"
+  - "1q128s1024"
+  - "1q256s1024"
+  - "1q512s1024"
+  - "1q1024"
+  # seq_len = 2048
+  - "1q1s2048"
+  - "1q2s2048"
+  - "1q4s2048"
+  - "1q8s2048"
+  - "1q16s2048"
+  - "1q32s2048"
+  - "1q64s2048"
+  - "1q128s2048"
+  - "1q256s2048"
+  - "1q512s2048"
+  - "1q1024s2048"
+  - "1q2048"
+
+  # Batch size 2
+  # seq_len = 32
+  - "2q1s32"
+  - "2q2s32"
+  - "2q4s32"
+  - "2q8s32"
+  - "2q16s32"
+  - "2q32"
+  # seq_len = 64
+  - "2q1s64"
+  - "2q2s64"
+  - "2q4s64"
+  - "2q8s64"
+  - "2q16s64"
+  - "2q32s64"
+  - "2q64"
+  # seq_len = 128
+  - "2q1s128"
+  - "2q2s128"
+  - "2q4s128"
+  - "2q8s128"
+  - "2q16s128"
+  - "2q32s128"
+  - "2q64s128"
+  - "2q128"
+  # seq_len = 256
+  - "2q1s256"
+  - "2q2s256"
+  - "2q4s256"
+  - "2q8s256"
+  - "2q16s256"
+  - "2q32s256"
+  - "2q64s256"
+  - "2q128s256"
+  - "2q256"
+  # seq_len = 512
+  - "2q1s512"
+  - "2q2s512"
+  - "2q4s512"
+  - "2q8s512"
+  - "2q16s512"
+  - "2q32s512"
+  - "2q64s512"
+  - "2q128s512"
+  - "2q256s512"
+  - "2q512"
+  # seq_len = 1024
+  - "2q1s1024"
+  - "2q2s1024"
+  - "2q4s1024"
+  - "2q8s1024"
+  - "2q16s1024"
+  - "2q32s1024"
+  - "2q64s1024"
+  - "2q128s1024"
+  - "2q256s1024"
+  - "2q512s1024"
+  - "2q1024"
+  # seq_len = 2048
+  - "2q1s2048"
+  - "2q2s2048"
+  - "2q4s2048"
+  - "2q8s2048"
+  - "2q16s2048"
+  - "2q32s2048"
+  - "2q64s2048"
+  - "2q128s2048"
+  - "2q256s2048"
+  - "2q512s2048"
+  - "2q1024s2048"
+  - "2q2048"
+
+  # Batch size 4
+  # seq_len = 32
+  - "4q1s32"
+  - "4q2s32"
+  - "4q4s32"
+  - "4q8s32"
+  - "4q16s32"
+  - "4q32"
+  # seq_len = 64
+  - "4q1s64"
+  - "4q2s64"
+  - "4q4s64"
+  - "4q8s64"
+  - "4q16s64"
+  - "4q32s64"
+  - "4q64"
+  # seq_len = 128
+  - "4q1s128"
+  - "4q2s128"
+  - "4q4s128"
+  - "4q8s128"
+  - "4q16s128"
+  - "4q32s128"
+  - "4q64s128"
+  - "4q128"
+  # seq_len = 256
+  - "4q1s256"
+  - "4q2s256"
+  - "4q4s256"
+  - "4q8s256"
+  - "4q16s256"
+  - "4q32s256"
+  - "4q64s256"
+  - "4q128s256"
+  - "4q256"
+  # seq_len = 512
+  - "4q1s512"
+  - "4q2s512"
+  - "4q4s512"
+  - "4q8s512"
+  - "4q16s512"
+  - "4q32s512"
+  - "4q64s512"
+  - "4q128s512"
+  - "4q256s512"
+  - "4q512"
+  # seq_len = 1024
+  - "4q1s1024"
+  - "4q2s1024"
+  - "4q4s1024"
+  - "4q8s1024"
+  - "4q16s1024"
+  - "4q32s1024"
+  - "4q64s1024"
+  - "4q128s1024"
+  - "4q256s1024"
+  - "4q512s1024"
+  - "4q1024"
+  # seq_len = 2048
+  - "4q1s2048"
+  - "4q2s2048"
+  - "4q4s2048"
+  - "4q8s2048"
+  - "4q16s2048"
+  - "4q32s2048"
+  - "4q64s2048"
+  - "4q128s2048"
+  - "4q256s2048"
+  - "4q512s2048"
+  - "4q1024s2048"
+  - "4q2048"
+
+  # Batch size 8
+  # seq_len = 32
+  - "8q1s32"
+  - "8q2s32"
+  - "8q4s32"
+  - "8q8s32"
+  - "8q16s32"
+  - "8q32"
+  # seq_len = 64
+  - "8q1s64"
+  - "8q2s64"
+  - "8q4s64"
+  - "8q8s64"
+  - "8q16s64"
+  - "8q32s64"
+  - "8q64"
+  # seq_len = 128
+  - "8q1s128"
+  - "8q2s128"
+  - "8q4s128"
+  - "8q8s128"
+  - "8q16s128"
+  - "8q32s128"
+  - "8q64s128"
+  - "8q128"
+  # seq_len = 256
+  - "8q1s256"
+  - "8q2s256"
+  - "8q4s256"
+  - "8q8s256"
+  - "8q16s256"
+  - "8q32s256"
+  - "8q64s256"
+  - "8q128s256"
+  - "8q256"
+  # seq_len = 512
+  - "8q1s512"
+  - "8q2s512"
+  - "8q4s512"
+  - "8q8s512"
+  - "8q16s512"
+  - "8q32s512"
+  - "8q64s512"
+  - "8q128s512"
+  - "8q256s512"
+  - "8q512"
+  # seq_len = 1024
+  - "8q1s1024"
+  - "8q2s1024"
+  - "8q4s1024"
+  - "8q8s1024"
+  - "8q16s1024"
+  - "8q32s1024"
+  - "8q64s1024"
+  - "8q128s1024"
+  - "8q256s1024"
+  - "8q512s1024"
+  - "8q1024"
+  # seq_len = 2048
+  - "8q1s2048"
+  - "8q2s2048"
+  - "8q4s2048"
+  - "8q8s2048"
+  - "8q16s2048"
+  - "8q32s2048"
+  - "8q64s2048"
+  - "8q128s2048"
+  - "8q256s2048"
+  - "8q512s2048"
+  - "8q1024s2048"
+  - "8q2048"
+
+  # Batch size 16
+  # seq_len = 32
+  - "16q1s32"
+  - "16q2s32"
+  - "16q4s32"
+  - "16q8s32"
+  - "16q16s32"
+  - "16q32"
+  # seq_len = 64
+  - "16q1s64"
+  - "16q2s64"
+  - "16q4s64"
+  - "16q8s64"
+  - "16q16s64"
+  - "16q32s64"
+  - "16q64"
+  # seq_len = 128
+  - "16q1s128"
+  - "16q2s128"
+  - "16q4s128"
+  - "16q8s128"
+  - "16q16s128"
+  - "16q32s128"
+  - "16q64s128"
+  - "16q128"
+  # seq_len = 256
+  - "16q1s256"
+  - "16q2s256"
+  - "16q4s256"
+  - "16q8s256"
+  - "16q16s256"
+  - "16q32s256"
+  - "16q64s256"
+  - "16q128s256"
+  - "16q256"
+  # seq_len = 512
+  - "16q1s512"
+  - "16q2s512"
+  - "16q4s512"
+  - "16q8s512"
+  - "16q16s512"
+  - "16q32s512"
+  - "16q64s512"
+  - "16q128s512"
+  - "16q256s512"
+  - "16q512"
+  # seq_len = 1024
+  - "16q1s1024"
+  - "16q2s1024"
+  - "16q4s1024"
+  - "16q8s1024"
+  - "16q16s1024"
+  - "16q32s1024"
+  - "16q64s1024"
+  - "16q128s1024"
+  - "16q256s1024"
+  - "16q512s1024"
+  - "16q1024"
+  # seq_len = 2048
+  - "16q1s2048"
+  - "16q2s2048"
+  - "16q4s2048"
+  - "16q8s2048"
+  - "16q16s2048"
+  - "16q32s2048"
+  - "16q64s2048"
+  - "16q128s2048"
+  - "16q256s2048"
+  - "16q512s2048"
+  - "16q1024s2048"
+  - "16q2048"
+
+  # Batch size 32
+  # seq_len = 32
+  - "32q1s32"
+  - "32q2s32"
+  - "32q4s32"
+  - "32q8s32"
+  - "32q16s32"
+  - "32q32"
+  # seq_len = 64
+  - "32q1s64"
+  - "32q2s64"
+  - "32q4s64"
+  - "32q8s64"
+  - "32q16s64"
+  - "32q32s64"
+  - "32q64"
+  # seq_len = 128
+  - "32q1s128"
+  - "32q2s128"
+  - "32q4s128"
+  - "32q8s128"
+  - "32q16s128"
+  - "32q32s128"
+  - "32q64s128"
+  - "32q128"
+  # seq_len = 256
+  - "32q1s256"
+  - "32q2s256"
+  - "32q4s256"
+  - "32q8s256"
+  - "32q16s256"
+  - "32q32s256"
+  - "32q64s256"
+  - "32q128s256"
+  - "32q256"
+  # seq_len = 512
+  - "32q1s512"
+  - "32q2s512"
+  - "32q4s512"
+  - "32q8s512"
+  - "32q16s512"
+  - "32q32s512"
+  - "32q64s512"
+  - "32q128s512"
+  - "32q256s512"
+  - "32q512"
+  # seq_len = 1024
+  - "32q1s1024"
+  - "32q2s1024"
+  - "32q4s1024"
+  - "32q8s1024"
+  - "32q16s1024"
+  - "32q32s1024"
+  - "32q64s1024"
+  - "32q128s1024"
+  - "32q256s1024"
+  - "32q512s1024"
+  - "32q1024"
+  # seq_len = 2048
+  - "32q1s2048"
+  - "32q2s2048"
+  - "32q4s2048"
+  - "32q8s2048"
+  - "32q16s2048"
+  - "32q32s2048"
+  - "32q64s2048"
+  - "32q128s2048"
+  - "32q256s2048"
+  - "32q512s2048"
+  - "32q1024s2048"
+  - "32q2048"
+
+backends:
+  - FLASHMLA_SPARSE
+
+device: "cuda:0"
+profile_memory: false
+sparse_mla_dense_mha_max_seq_len: 2048
+sparse_mla_topk_pattern: "random"
+
+output:
+  csv: "benchmark_output/mla_sparse_mha_vs_mqa.csv"
+  json: "benchmark_output/mla_sparse_mha_vs_mqa.json"

benchmarks/attention_benchmarks/mla_runner.py

@@ -9,6 +9,8 @@ needing full VllmConfig integration.
 """
 
 import statistics
+import tempfile
+from pathlib import Path
 
 import numpy as np
 import torch
@@ -17,7 +19,6 @@ from common import (
     BenchmarkResult,
     MockHfConfig,
     MockIndexer,
-    MockKVBProj,
     MockLayer,
     run_do_bench,
     run_ncu_profile,
@@ -33,8 +34,59 @@ from vllm.config import (
     VllmConfig,
     set_current_vllm_config,
 )
+from vllm.model_executor.layers.linear import ColumnParallelLinear
 from vllm.v1.attention.backends.mla.prefill.registry import MLAPrefillBackendEnum
 
+
+def _safe_profile_name(value: str) -> str:
+    return "".join(c if c.isalnum() or c in "._-" else "_" for c in value)
+
+
+def _create_kv_b_proj(
+    mla_dims: dict,
+    device: torch.device,
+):
+    kv_b_proj = ColumnParallelLinear(
+        mla_dims["kv_lora_rank"],
+        mla_dims["num_q_heads"]
+        * (mla_dims["qk_nope_head_dim"] + mla_dims["v_head_dim"]),
+        bias=False,
+        params_dtype=torch.bfloat16,
+        quant_config=None,
+        prefix="benchmark.kv_b_proj",
+    ).to(device)
+    with torch.no_grad():
+        kv_b_proj.weight.copy_(torch.randn_like(kv_b_proj.weight))
+    return kv_b_proj
+
+
+def _ensure_single_rank_model_parallel() -> None:
+    import torch.distributed as dist
+
+    from vllm.distributed import (
+        ensure_model_parallel_initialized,
+        init_distributed_environment,
+        model_parallel_is_initialized,
+    )
+
+    if not dist.is_available():
+        return
+    if not dist.is_initialized():
+        with tempfile.NamedTemporaryFile(
+            prefix="vllm_bench_dist_", delete=False
+        ) as init_file:
+            distributed_init_method = f"file://{init_file.name}"
+        init_distributed_environment(
+            world_size=1,
+            rank=0,
+            distributed_init_method=distributed_init_method,
+            local_rank=0,
+            backend="nccl",
+        )
+    if not model_parallel_is_initialized():
+        ensure_model_parallel_initialized(1, 1)
+
+
 # ============================================================================
 # VllmConfig Creation
 # ============================================================================
@@ -66,10 +118,12 @@ def create_minimal_vllm_config(
     block_size: int = 128,
     max_num_seqs: int = 256,
     max_num_batched_tokens: int = 8192,
+    max_model_len: int = 32768,
     mla_dims: dict | None = None,
     index_topk: int | None = None,
     prefill_backend: str | None = None,
     kv_cache_dtype: str = "auto",
+    sparse_mla_force_mqa: bool = False,
 ) -> VllmConfig:
     """
     Create minimal VllmConfig for MLA benchmarks.
@@ -86,6 +140,8 @@ def create_minimal_vllm_config(
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4", "flashinfer",
                         "trtllm"). Configures the attention config to force
                         the specified prefill backend.
+        sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
+                    forward_mqa (even prefill tokens).
 
     Returns:
         VllmConfig for benchmarking
@@ -131,7 +187,7 @@ def create_minimal_vllm_config(
             trust_remote_code=True,
             dtype="bfloat16",
             seed=0,
-            max_model_len=32768,
+            max_model_len=max_model_len,
             quantization=None,
             enforce_eager=False,
             max_logprobs=20,
@@ -163,7 +219,7 @@ def create_minimal_vllm_config(
     scheduler_config = SchedulerConfig(
         max_num_seqs=max_num_seqs,
         max_num_batched_tokens=max(max_num_batched_tokens, max_num_seqs),
-        max_model_len=32768,
+        max_model_len=max_model_len,
         is_encoder_decoder=False,
         enable_chunked_prefill=True,
     )
@@ -192,6 +248,9 @@ def create_minimal_vllm_config(
                 "flash_attn_version"
             ]
 
+    if sparse_mla_force_mqa:
+        vllm_config.attention_config.sparse_mla_force_mqa = True
+
     return vllm_config
 
 
@@ -548,12 +607,7 @@ def _create_backend_impl(
     # Calculate scale
     scale = 1.0 / np.sqrt(mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"])
 
-    # Create mock kv_b_proj layer for prefill mode
-    mock_kv_b_proj = MockKVBProj(
-        num_heads=mla_dims["num_q_heads"],
-        qk_nope_head_dim=mla_dims["qk_nope_head_dim"],
-        v_head_dim=mla_dims["v_head_dim"],
-    )
+    kv_b_proj = _create_kv_b_proj(mla_dims, device)
 
     # Create indexer for sparse backends
     indexer = None
@@ -584,7 +638,7 @@ def _create_backend_impl(
         "qk_rope_head_dim": mla_dims["qk_rope_head_dim"],
         "qk_head_dim": mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"],
         "v_head_dim": mla_dims["v_head_dim"],
-        "kv_b_proj": mock_kv_b_proj,
+        "kv_b_proj": kv_b_proj,
     }
 
     # Add indexer for sparse backends
@@ -785,14 +839,35 @@ def _run_single_benchmark(
     # Fill indexer with random indices for sparse backends
     is_sparse = backend_cfg.get("is_sparse", False)
     if is_sparse and indexer is not None:
-        indexer.fill_random_indices(total_q, max_kv_len)
+        indexer.fill_indices(
+            total_q,
+            max_kv_len,
+            getattr(config, "sparse_mla_topk_pattern", "random"),
+        )
 
     # Determine which forward methods to use based on metadata.
-    # Sparse MLA backends always use forward_mqa
-    has_decode = is_sparse or getattr(metadata, "decode", None) is not None
-    has_prefill = not is_sparse and getattr(metadata, "prefill", None) is not None
+    # Non-sparse backends use .decode/.prefill sub-objects.
+    # Sparse backends use num_decode_tokens/num_prefills directly.
+    #
+    # sparse_mla_force_mqa overrides: even for prefill metadata, use MQA.
+    force_mqa = getattr(config, "sparse_mla_force_mqa", False)
+    force_dense_mha = getattr(config, "sparse_mla_mha_mode", "auto") == "dense"
+    if force_mqa:
+        has_decode = True
+        has_prefill = False
+    elif is_sparse:
+        has_decode = metadata.num_decode_tokens > 0
+        has_prefill = metadata.num_prefills > 0
+    else:
+        has_decode = metadata.decode is not None
+        has_prefill = metadata.prefill is not None
     if not has_decode and not has_prefill:
         raise RuntimeError("Metadata has neither decode nor prefill metadata")
+    if is_sparse and force_dense_mha and not has_prefill:
+        raise RuntimeError(
+            "Sparse MLA dense_mha benchmark did not produce prefill metadata. "
+            "Check reorder_batch_threshold/path forcing."
+        )
 
     num_decode = (
         metadata.num_decode_tokens
@@ -871,7 +946,6 @@ def _run_single_benchmark(
                 metadata,
                 prefill_inputs["k_scale"],
                 prefill_fp8_output if fused_output else prefill_inputs["output"],
-                prefill_output_scale if fused_output else None,
             )
             if fused_output:
                 out = prefill_fp8_output
@@ -898,6 +972,48 @@ def _run_single_benchmark(
             throughput_tokens_per_sec=0.0,
         )
 
+    if config.torch_profile:
+        profile_dir = Path(
+            config.torch_profile_dir or "benchmark_outputs/torch_profiles"
+        )
+        profile_dir.mkdir(parents=True, exist_ok=True)
+        trace_name = _safe_profile_name(f"{config.backend}_{config.batch_spec}")
+        trace_path = profile_dir / f"{trace_name}.json"
+        iters = max(config.torch_profile_iters, 1)
+
+        forward_fn()
+        torch.accelerator.synchronize()
+        with torch.profiler.profile(
+            activities=[
+                torch.profiler.ProfilerActivity.CPU,
+                torch.profiler.ProfilerActivity.CUDA,
+            ],
+            record_shapes=True,
+            profile_memory=True,
+            with_stack=False,
+        ) as prof:
+            for _ in range(iters):
+                forward_fn()
+                torch.accelerator.synchronize()
+                prof.step()
+        prof.export_chrome_trace(str(trace_path))
+        print(f"Saved PyTorch profiler trace to {trace_path}")
+        print(
+            prof.key_averages().table(
+                sort_by="cuda_time_total",
+                row_limit=25,
+            )
+        )
+        return BenchmarkResult(
+            config=config,
+            mean_time=0.0,
+            median_time=0.0,
+            std_time=0.0,
+            min_time=0.0,
+            max_time=0.0,
+            throughput_tokens_per_sec=0.0,
+        )
+
     all_ms = run_do_bench(benchmark_fn, config.use_cuda_graphs, config.warmup_ms)
 
     # Convert ms to seconds per layer
@@ -920,6 +1036,7 @@ def _run_mla_benchmark_batched(
     configs_with_params: list[tuple],  # [(config, threshold, num_splits), ...]
     index_topk: int = 2048,
     prefill_backend: str | None = None,
+    sparse_mla_force_mqa: bool = False,
     output_scale: float | None = None,
     fuse_quant_op: bool = False,
 ) -> list[BenchmarkResult]:
@@ -940,6 +1057,8 @@ def _run_mla_benchmark_batched(
         index_topk: Topk value for sparse MLA backends (default 2048)
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4").
             When set, forces the specified FlashAttention version for prefill.
+        sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
+            forward_mqa (even prefill tokens).
 
     Returns:
         List of BenchmarkResult objects
@@ -980,21 +1099,41 @@ def _run_mla_benchmark_batched(
         sum(r.q_len for r in parse_batch_spec(cfg.batch_spec))
         for cfg, *_ in configs_with_params
     )
+    max_model_len = max(
+        max_total_q,
+        max(
+            getattr(cfg, "max_model_len", None) or 32768
+            for cfg, *_ in configs_with_params
+        ),
+    )
 
     # Create and set vLLM config for MLA (reused across all benchmarks)
     vllm_config = create_minimal_vllm_config(
         model_name="deepseek-v3",  # Used only for model path
         block_size=block_size,
         max_num_batched_tokens=max_total_q,
+        max_model_len=max_model_len,
         mla_dims=mla_dims,  # Use custom dims from config or default
         index_topk=index_topk if is_sparse else None,
         prefill_backend=prefill_backend,
         kv_cache_dtype=kv_cache_dtype,
+        sparse_mla_force_mqa=sparse_mla_force_mqa,
     )
 
     results = []
 
+    # Initialize workspace manager (needed by metadata builders)
+    from vllm.v1.worker.workspace import (
+        init_workspace_manager,
+        is_workspace_manager_initialized,
+    )
+
+    if not is_workspace_manager_initialized():
+        init_workspace_manager(device)
+
     with set_current_vllm_config(vllm_config):
+        _ensure_single_rank_model_parallel()
+
         # Create backend impl, layer, builder, and indexer (reused across benchmarks)
         impl, layer, builder_instance, indexer = _create_backend_impl(
             backend_cfg,
@@ -1040,9 +1179,20 @@ def _run_mla_benchmark_batched(
         for config, threshold, num_splits in configs_with_params:
             # Set threshold for this benchmark (FlashAttn/FlashMLA only)
             original_threshold = None
-            if threshold is not None and builder_instance:
+            effective_threshold = threshold
+            force_dense_mha = (
+                is_sparse
+                and getattr(config, "sparse_mla_mha_mode", "auto") == "dense"
+                and not getattr(config, "sparse_mla_force_mqa", False)
+            )
+            if force_dense_mha:
+                # Sparse MLA normally treats q_len <= 1 as decode. Use an
+                # impossible threshold so dense_mha benchmarks actually run
+                # the prefill/MHA path, including q_len=1 short extends.
+                effective_threshold = -1
+            if effective_threshold is not None and builder_instance:
                 original_threshold = builder_instance.reorder_batch_threshold
-                builder_instance.reorder_batch_threshold = threshold
+                builder_instance.reorder_batch_threshold = effective_threshold
 
             # Set num_splits for CUTLASS
             original_num_splits = None
@@ -1090,6 +1240,7 @@ def run_mla_benchmark(
     num_kv_splits: int | None = None,
     index_topk: int = 2048,
     prefill_backend: str | None = None,
+    sparse_mla_force_mqa: bool = False,
     output_scale: float | None = None,
     fuse_quant_op: bool = False,
 ) -> BenchmarkResult | list[BenchmarkResult]:
@@ -1111,6 +1262,8 @@ def run_mla_benchmark(
         index_topk: Topk value for sparse MLA backends (default 2048)
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4").
             When set, forces the specified FlashAttention version for prefill.
+        sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
+            forward_mqa (even prefill tokens).
         output_scale: Static per-tensor FP8 scale for prefill output (None = bf16).
         fuse_quant_op: With output_scale set, fuse the FP8 write into the prefill
             kernel vs a standalone post-quant kernel. See _run_single_benchmark.
@@ -1142,6 +1295,7 @@ def run_mla_benchmark(
         configs_with_params,
         index_topk,
         prefill_backend=prefill_backend,
+        sparse_mla_force_mqa=sparse_mla_force_mqa,
         output_scale=output_scale,
         fuse_quant_op=fuse_quant_op,
     )

benchmarks/kernels/bench_cp_gather_fp8.py

@@ -69,12 +69,11 @@ def make_inputs(total_tokens, num_reqs, block_size):
     # Output workspace
     dst = torch.zeros(total_tokens, HEAD_DIM, dtype=torch.bfloat16, device="cuda")
 
-    seq_lens_t = torch.tensor(seq_lens, dtype=torch.int32, device="cuda")
     workspace_starts_t = torch.tensor(
         workspace_starts, dtype=torch.int32, device="cuda"
     )
 
-    return cache, dst, block_table, seq_lens_t, workspace_starts_t
+    return cache, dst, block_table, workspace_starts_t
 
 
 def bench_scenario(label, num_reqs, total_tokens_list, save_path):
@@ -94,7 +93,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):
         )
     )
     def bench_fn(total_tokens, provider, num_reqs):
-        cache, dst, block_table, seq_lens_t, ws_starts = make_inputs(
+        cache, dst, block_table, ws_starts = make_inputs(
             total_tokens, num_reqs, BLOCK_SIZE
         )
 
@@ -102,7 +101,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):
 
         ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
             lambda: ops.cp_gather_and_upconvert_fp8_kv_cache(
-                cache, dst, block_table, seq_lens_t, ws_starts, num_reqs
+                cache, dst, block_table, ws_starts, num_reqs
             ),
             quantiles=quantiles,
             rep=500,

benchmarks/kernels/benchmark_rdna_hybrid_w4a16_gemm.py

@@ -0,0 +1,201 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Benchmark the RDNAHybridW4A16LinearKernel across decode and prefill shapes.
+
+Usage:
+    python benchmark_int4_gemm.py
+    python benchmark_int4_gemm.py --models Qwen/Qwen3-4B
+    python benchmark_int4_gemm.py --group-size 128
+"""
+
+import argparse
+import copy
+import itertools
+import os
+
+import torch
+
+from vllm.triton_utils import triton
+
+# ---------------------------------------------------------------------------
+# Weight shapes: [K, N], TP_SPLIT_DIM
+# ---------------------------------------------------------------------------
+WEIGHT_SHAPES = {
+    "Qwen/Qwen3-4B": [
+        ([2560, 3840], 1),  # qkv_proj
+        ([2560, 2560], 0),  # o_proj
+        ([2560, 19456], 1),  # gate_up_proj
+        ([9728, 2560], 0),  # down_proj
+    ],
+    "Qwen/Qwen2.5-7B-Instruct": [
+        ([3584, 4608], 1),
+        ([3584, 3584], 0),
+        ([3584, 37888], 1),
+        ([18944, 3584], 0),
+    ],
+    "trymirai/SmolLM2-1.7B-Instruct-AWQ": [
+        ([2048, 6144], 1),  # qkv_proj
+        ([2048, 2048], 0),  # o_proj
+        ([2048, 16384], 1),  # gate_up_proj
+        ([8192, 2048], 0),  # down_proj
+    ],
+    "RedHatAI/Qwen3-8B-quantized.w4a16": [
+        ([4096, 6144], 1),  # qkv_proj
+        ([4096, 4096], 0),  # o_proj
+        ([4096, 24576], 1),  # gate_up_proj
+        ([12288, 4096], 0),  # down_proj
+    ],
+}
+
+
+# ---------------------------------------------------------------------------
+# Weight packing
+# ---------------------------------------------------------------------------
+def prepare_hybrid_weights(K, N, group_size, device="cuda"):
+    """Create random weights for benchmarking.
+
+    Returns (w_q_skinny, w_s_skinny, w_fp16, w_zp). The triton path derives
+    its int32 view from w_q_skinny, so no separate int32 buffer is returned.
+    """
+    num_groups = K // group_size
+
+    # Random packed weights — actual values don't matter for throughput
+    w_q_skinny_i32 = torch.randint(
+        0, 2**31, (N, K // 8), dtype=torch.int32, device=device
+    )
+    w_q_skinny = w_q_skinny_i32.view(torch.int8).contiguous()
+    w_s_skinny = torch.randn(N, num_groups, dtype=torch.float16, device=device) * 0.01
+
+    # Raw per-group zero-points for asymmetric benchmarks
+    w_zp = torch.randint(0, 16, (N, num_groups), dtype=torch.int32, device=device).to(
+        torch.float16
+    )
+
+    # FP16 baseline for F.linear
+    w_fp16 = torch.randn(N, K, dtype=torch.float16, device=device) * 0.01
+
+    return w_q_skinny, w_s_skinny, w_fp16, w_zp
+
+
+# ---------------------------------------------------------------------------
+# Benchmark
+# ---------------------------------------------------------------------------
+PROVIDERS = ["torch-fp16", "hybrid-w4a16", "hybrid-w4a16-zp"]
+
+
+@triton.testing.perf_report(
+    triton.testing.Benchmark(
+        x_names=["batch_size"],
+        x_vals=[1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096],
+        x_log=False,
+        line_arg="provider",
+        line_vals=PROVIDERS,
+        line_names=PROVIDERS,
+        ylabel="TFLOP/s (larger is better)",
+        plot_name="FP16 vs Hybrid W4A16",
+        args={},
+    )
+)
+def benchmark(batch_size, provider, N, K, group_size, weights):
+    M = batch_size
+    device = "cuda"
+    dtype = torch.float16
+    a = torch.randn((M, K), device=device, dtype=dtype)
+
+    quantiles = [0.5, 0.2, 0.8]
+
+    if provider == "torch-fp16":
+        w_fp16 = weights["w_fp16"]
+        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+            lambda: torch.nn.functional.linear(a, w_fp16),
+            quantiles=quantiles,
+        )
+    elif provider in ("hybrid-w4a16", "hybrid-w4a16-zp"):
+        from vllm.model_executor.kernels.linear.mixed_precision import (
+            rdna_hybrid_w4a16 as _k,
+        )
+
+        _rdna_hybrid_w4a16_apply_impl = _k._rdna_hybrid_w4a16_apply_impl
+        from vllm.utils.platform_utils import num_compute_units
+
+        w = weights
+        cu_count = num_compute_units()
+        use_zp = provider == "hybrid-w4a16-zp"
+
+        def run():
+            return _rdna_hybrid_w4a16_apply_impl(
+                a,
+                w["w_q_skinny"],
+                w["w_s_skinny"],
+                w["w_zp"] if use_zp else None,
+                None,  # bias
+                cu_count,
+                group_size,
+            )
+
+        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+            run,
+            quantiles=quantiles,
+        )
+    else:
+        return 0.0, 0.0, 0.0
+
+    to_tflops = lambda t_ms: (2 * M * N * K) * 1e-12 / (t_ms * 1e-3)
+    return to_tflops(ms), to_tflops(max_ms), to_tflops(min_ms)
+
+
+def prepare_shapes(args):
+    KN_model_names = []
+    for model, tp_size in itertools.product(args.models, args.tp_sizes):
+        for KN, tp_dim in copy.deepcopy(WEIGHT_SHAPES[model]):
+            KN[tp_dim] //= tp_size
+            KN.append(model)
+            KN_model_names.append(KN)
+    return KN_model_names
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser(
+        description="Benchmark RDNAHybridW4A16LinearKernel"
+    )
+    parser.add_argument(
+        "--models",
+        nargs="+",
+        type=str,
+        default=["Qwen/Qwen3-4B"],
+        choices=list(WEIGHT_SHAPES.keys()),
+    )
+    parser.add_argument("--tp-sizes", nargs="+", type=int, default=[1])
+    parser.add_argument("--group-size", type=int, default=128)
+    parser.add_argument("--save-path", type=str, default=None)
+    args = parser.parse_args()
+
+    for K, N, model in prepare_shapes(args):
+        group_size = args.group_size
+        print(f"\n{'=' * 70}")
+        print(f"{model}, N={N} K={K}, group_size={group_size}")
+        print(f"{'=' * 70}")
+
+        w_q_skinny, w_s_skinny, w_fp16, w_zp = prepare_hybrid_weights(K, N, group_size)
+
+        weights = {
+            "w_q_skinny": w_q_skinny,
+            "w_s_skinny": w_s_skinny,
+            "w_fp16": w_fp16,
+            "w_zp": w_zp,
+        }
+
+        save_path = args.save_path or f"bench_int4_res_n{N}_k{K}"
+        os.makedirs(save_path, exist_ok=True)
+        benchmark.run(
+            print_data=True,
+            show_plots=False,
+            save_path=save_path,
+            N=N,
+            K=K,
+            group_size=group_size,
+            weights=weights,
+        )
+
+    print("\nBenchmark finished!")

benchmarks/kernels/benchmark_relu_squared.py

@@ -0,0 +1,108 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+# Benchmark ReLUSquaredActivation: custom CUDA kernel vs forward_native, both
+# eager and under torch.compile (Inductor fuses relu+square into one kernel).
+
+import itertools
+
+import torch
+import torch.nn.functional as F
+
+import vllm.model_executor.layers.activation  # noqa: F401
+from vllm.benchmarks.lib.utils import default_vllm_config
+from vllm.triton_utils import triton
+from vllm.utils.argparse_utils import FlexibleArgumentParser
+from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
+
+# Capped so the largest tensor stays under 2**31 elements: the shared activation
+# kernel computes the per-token pointer offset (blockIdx.x * d) in 32-bit, which
+# overflows for tensors with >2**32 elements. Realistic token counts are well
+# below this; the kernel-vs-native gap is already clear at these sizes.
+batch_size_range = [1, 16, 128]
+seq_len_range = [1, 16, 64, 1024]
+intermediate_size = [3072, 9728, 12288]
+configs = list(itertools.product(batch_size_range, seq_len_range, intermediate_size))
+
+
+@default_vllm_config()
+def benchmark_relu_squared(
+    batch_size: int,
+    seq_len: int,
+    intermediate_size: int,
+    provider: str,
+    dtype: torch.dtype,
+):
+    device = "cuda"
+    num_tokens = batch_size * seq_len
+    set_random_seed(42)
+    torch.set_default_device(device)
+
+    x = torch.randn(num_tokens, intermediate_size, dtype=dtype, device=device)
+    out = torch.empty_like(x)
+
+    def native(x: torch.Tensor) -> torch.Tensor:
+        return torch.square(F.relu(x))
+
+    # Verify the custom kernel matches the native implementation before timing.
+    ref = native(x)
+    torch.ops._C.relu_squared(out, x)
+    torch.testing.assert_close(out, ref)
+
+    if provider == "custom":
+        # Custom CUDA kernel — single fused kernel.
+        fn = lambda: torch.ops._C.relu_squared(out, x)
+    elif provider == "native":
+        # forward_native, eager — relu and square as separate ops.
+        fn = lambda: native(x)
+    elif provider == "native_compiled":
+        # forward_native under torch.compile — Inductor fuses relu+square.
+        # This is the real production baseline (custom ops are off when
+        # Inductor is enabled), so it is the comparison reviewers care about.
+        compiled = torch.compile(native)
+        compiled(x)  # warm up / trigger compilation before timing
+        fn = lambda: compiled(x)
+
+    ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+        fn, quantiles=[0.5, 0.2, 0.8]
+    )
+    return ms, max_ms, min_ms
+
+
+if __name__ == "__main__":
+    parser = FlexibleArgumentParser(
+        description="Benchmark ReLUSquaredActivation: custom kernel vs native."
+    )
+    parser.add_argument(
+        "--dtype",
+        type=str,
+        choices=["half", "bfloat16", "float"],
+        default="bfloat16",
+    )
+    args = parser.parse_args()
+
+    dtype = STR_DTYPE_TO_TORCH_DTYPE[args.dtype]
+
+    perf_report = triton.testing.perf_report(
+        triton.testing.Benchmark(
+            x_names=["batch_size", "seq_len", "intermediate_size"],
+            x_vals=configs,
+            line_arg="provider",
+            line_vals=["custom", "native_compiled", "native"],
+            line_names=[
+                "Custom Kernel",
+                "Native (torch.compile)",
+                "Native (eager)",
+            ],
+            styles=[("blue", "-"), ("green", "-"), ("red", "-")],
+            ylabel="ms",
+            plot_name="relu_squared-eager-performance",
+            args={},
+        )
+    )
+
+    perf_report(
+        lambda batch_size, seq_len, intermediate_size, provider: benchmark_relu_squared(
+            batch_size, seq_len, intermediate_size, provider, dtype
+        )
+    ).run(print_data=True)

build_rust.sh

@@ -8,6 +8,8 @@
 set -euo pipefail
 
 REPO_ROOT="$(cd "$(dirname "$0")" && pwd)"
+CARGO_LLVM_COV_VERSION="0.8.7"
+COVERAGE_TOOLS_DIR="$REPO_ROOT/rust-coverage-tools"
 
 # Read the required toolchain from rust-toolchain.toml.
 TOOLCHAIN=$(grep '^channel' "$REPO_ROOT/rust-toolchain.toml" | sed 's/.*= *"\(.*\)"/\1/')
@@ -30,4 +32,39 @@ else
     PROFILE_ARG="--release"
 fi
 
+rm -rf "$COVERAGE_TOOLS_DIR"
+mkdir -p "$COVERAGE_TOOLS_DIR/bin" "$COVERAGE_TOOLS_DIR/lib"
+
+if [[ "${VLLM_RUST_COVERAGE:-0}" == "1" ]]; then
+    # rustc wrapper flags are invisible to Cargo's normal fingerprinting.
+    # Keep instrumented intermediates isolated when local builds switch modes.
+    export CARGO_TARGET_DIR="$REPO_ROOT/rust/target/coverage"
+    rustup component add --toolchain "$TOOLCHAIN" llvm-tools-preview
+    cargo +"$TOOLCHAIN" install \
+        --locked \
+        --version "$CARGO_LLVM_COV_VERSION" \
+        cargo-llvm-cov
+
+    eval "$(
+        cargo +"$TOOLCHAIN" llvm-cov show-env \
+            --manifest-path "$REPO_ROOT/rust/Cargo.toml" \
+            --sh
+    )"
+
+    # Build scripts and proc macros can run during compilation. Their profiles
+    # are unrelated to runtime coverage and would otherwise pollute the tree.
+    export LLVM_PROFILE_FILE=/dev/null
+    export VLLM_RUST_COVERAGE_OBJECTS="$COVERAGE_TOOLS_DIR/objects"
+fi
+
 python3 "$REPO_ROOT/tools/build_rust.py" "$PROFILE_ARG"
+
+if [[ "${VLLM_RUST_COVERAGE:-0}" == "1" ]]; then
+    LLVM_BIN_DIR="$(dirname "$(rustup run "$TOOLCHAIN" rustc \
+        --print target-libdir)")/bin"
+
+    cp "$LLVM_BIN_DIR"/{llvm-cov,llvm-profdata} "$COVERAGE_TOOLS_DIR/bin/"
+    chmod 0755 "$COVERAGE_TOOLS_DIR/bin/"*
+    cp -L "$LLVM_BIN_DIR"/../lib/libLLVM.so* "$COVERAGE_TOOLS_DIR/lib/"
+    chmod 0644 "$COVERAGE_TOOLS_DIR/lib/"*
+fi

cmake/cpu_extension.cmake

@@ -430,6 +430,7 @@ set(VLLM_EXT_SRC
     "csrc/cpu/layernorm.cpp"
     "csrc/cpu/mla_decode.cpp"
     "csrc/cpu/pos_encoding.cpp"
+    "csrc/cpu/mamba_cpu.cpp"
     "csrc/moe/dynamic_4bit_int_moe_cpu.cpp"
     "csrc/cpu/cpu_attn.cpp"
     "csrc/cpu/torch_bindings.cpp")
@@ -489,6 +490,7 @@ if (ENABLE_X86_ISA)
         "csrc/cpu/spec_decode_utils.cpp"
         "csrc/cpu/cpu_attn.cpp"
         "csrc/cpu/dnnl_kernels.cpp"
+        "csrc/cpu/mamba_cpu.cpp"
         "csrc/cpu/torch_bindings.cpp"
         # TODO: Remove these files
         "csrc/cpu/activation.cpp"
@@ -502,6 +504,7 @@ if (ENABLE_X86_ISA)
         "csrc/cpu/utils.cpp"
         "csrc/cpu/spec_decode_utils.cpp"
         "csrc/cpu/cpu_attn.cpp"
+        "csrc/cpu/mamba_cpu.cpp"
         "csrc/cpu/dnnl_kernels.cpp"
         "csrc/cpu/torch_bindings.cpp"
         # TODO: Remove these files

cmake/external_projects/flashmla.cmake

@@ -19,7 +19,7 @@ else()
   FetchContent_Declare(
         flashmla
         GIT_REPOSITORY https://github.com/vllm-project/FlashMLA
-        GIT_TAG a6ec2ba7bd0a7dff98b3f4d3e6b52b159c48d78b
+        GIT_TAG a8f794d1251cbfd88a5011445dd5582289c727e4
         GIT_PROGRESS TRUE
         CONFIGURE_COMMAND ""
         BUILD_COMMAND ""
@@ -35,7 +35,7 @@ set(FLASHMLA_VENDOR_DIR "${CMAKE_SOURCE_DIR}/vllm/third_party/flashmla")
 file(MAKE_DIRECTORY "${FLASHMLA_VENDOR_DIR}")
 file(READ "${flashmla_SOURCE_DIR}/flash_mla/flash_mla_interface.py"
      FLASHMLA_INTERFACE_CONTENT)
-string(REPLACE "import flash_mla.cuda as flash_mla_cuda"
+string(REPLACE "flash_mla_cuda = torch.ops._flashmla_C"
                "import vllm._flashmla_C\nflash_mla_cuda = torch.ops._flashmla_C"
                FLASHMLA_INTERFACE_CONTENT
                "${FLASHMLA_INTERFACE_CONTENT}")
@@ -72,7 +72,7 @@ if(FLASH_MLA_ARCHS)
     list(APPEND VLLM_FLASHMLA_GPU_FLAGS "--expt-relaxed-constexpr" "--expt-extended-lambda" "--use_fast_math")
 
     set(FlashMLA_SOURCES
-        ${flashmla_SOURCE_DIR}/csrc/torch_api.cpp
+        ${flashmla_SOURCE_DIR}/csrc/api/api.cpp
 
         # Misc kernels for decoding
         ${flashmla_SOURCE_DIR}/csrc/smxx/decode/get_decoding_sched_meta/get_decoding_sched_meta.cu
@@ -128,6 +128,7 @@ if(FLASH_MLA_ARCHS)
 
     set(FlashMLA_Extension_INCLUDES
         ${flashmla_SOURCE_DIR}/csrc
+        ${flashmla_SOURCE_DIR}/csrc/kerutils/include
         ${flashmla_SOURCE_DIR}/csrc/extension/sm90/dense_fp8/
         ${flashmla_SOURCE_DIR}/csrc/cutlass/include
         ${flashmla_SOURCE_DIR}/csrc/cutlass/tools/util/include
@@ -152,15 +153,18 @@ if(FLASH_MLA_ARCHS)
         USE_SABI 3
         WITH_SOABI)
 
-    # Keep Stable ABI for the module, but *not* for CUDA/C++ files.
-    # This prevents Py_LIMITED_API from affecting nvcc and C++ compiles.
-    # Also enable C++20 for the FlashMLA sources (required for std::span, requires, etc.)
+    # Enable C++20 for the FlashMLA sources (required for std::span, requires, etc.)
     target_compile_options(_flashmla_C PRIVATE
-        $<$<COMPILE_LANGUAGE:CUDA>:-UPy_LIMITED_API>
-        $<$<COMPILE_LANGUAGE:CXX>:-UPy_LIMITED_API>
         $<$<COMPILE_LANGUAGE:CXX>:-std=c++20>
         $<$<COMPILE_LANGUAGE:CUDA>:-std=c++20>)
 
+    # _flashmla_C is now ABI-stable torch 2.11+
+    target_compile_definitions(_flashmla_C PRIVATE
+        TORCH_TARGET_VERSION=0x020B000000000000ULL)
+    if(VLLM_GPU_LANG STREQUAL "CUDA")
+        target_compile_definitions(_flashmla_C PRIVATE USE_CUDA)
+    endif()
+
     define_extension_target(
         _flashmla_extension_C
         DESTINATION vllm
@@ -172,11 +176,12 @@ if(FLASH_MLA_ARCHS)
         USE_SABI 3
         WITH_SOABI)
 
-    # Keep Stable ABI for the module, but *not* for CUDA/C++ files.
-    # This prevents Py_LIMITED_API from affecting nvcc and C++ compiles.
-    target_compile_options(_flashmla_extension_C PRIVATE
-        $<$<COMPILE_LANGUAGE:CUDA>:-UPy_LIMITED_API>
-        $<$<COMPILE_LANGUAGE:CXX>:-UPy_LIMITED_API>)
+    # _flashmla_extension_C is now ABI-stable w/ torch 2.11+
+    target_compile_definitions(_flashmla_extension_C PRIVATE
+        TORCH_TARGET_VERSION=0x020B000000000000ULL)
+    if(VLLM_GPU_LANG STREQUAL "CUDA")
+        target_compile_definitions(_flashmla_extension_C PRIVATE USE_CUDA)
+    endif()
 else()
     message(STATUS "FlashMLA will not compile: unsupported CUDA architecture ${CUDA_ARCHS}")
     # Create empty targets for setup.py on unsupported systems

cmake/external_projects/qutlass.cmake

@@ -22,7 +22,7 @@ if(QUTLASS_SRC_DIR)
   set(qutlass_BINARY_DIR "${CMAKE_BINARY_DIR}/qutlass-binary-dir-unused")
 else()
   set(_QUTLASS_UPSTREAM_REPO "https://github.com/IST-DASLab/qutlass.git")
-  set(_QUTLASS_UPSTREAM_TAG "830d2c4537c7396e14a02a46fbddd18b5d107c65")
+  set(_QUTLASS_UPSTREAM_TAG "e74319e3405ce6d71965732880f5dc1f52371f64")
 
   set(_qutlass_fc_root "${FETCHCONTENT_BASE_DIR}")
   if(NOT _qutlass_fc_root)
@@ -125,8 +125,6 @@ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
     CUDA_ARCHS "${QUTLASS_ARCHS}"
   )
 
-  # QuTLASS uses legacy ATen headers and cannot be built with TORCH_TARGET_VERSION.
-  # Keep it as its own extension (registers torch.ops._qutlass_C).
   define_extension_target(
     _qutlass_C
     DESTINATION vllm
@@ -139,9 +137,11 @@ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
     WITH_SOABI)
 
   target_compile_definitions(_qutlass_C PRIVATE
-    QUTLASS_DISABLE_PYBIND=1
+    QUTLASS_MINIMAL_BUILD=1
     TARGET_CUDA_ARCH=${QUTLASS_TARGET_CC}
-    CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
+    CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1
+    TORCH_TARGET_VERSION=0x020B000000000000ULL
+    USE_CUDA)
 
   set_property(SOURCE ${QUTLASS_SOURCES} APPEND PROPERTY COMPILE_OPTIONS
     $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr --use_fast_math -O3>

cmake/external_projects/tml_fa4.cmake

@@ -0,0 +1,50 @@
+include(FetchContent)
+
+if(DEFINED ENV{TML_FA4_SRC_DIR})
+  set(TML_FA4_SRC_DIR $ENV{TML_FA4_SRC_DIR})
+endif()
+
+if(TML_FA4_SRC_DIR)
+  FetchContent_Declare(
+    tml_fa4
+    SOURCE_DIR ${TML_FA4_SRC_DIR}
+    CONFIGURE_COMMAND ""
+    BUILD_COMMAND "")
+else()
+  FetchContent_Declare(
+    tml_fa4
+    GIT_REPOSITORY https://github.com/vllm-project/tml-fa4.git
+    GIT_TAG b206834606ed5b5f21f8eed6b0683f528ea9cf7d
+    GIT_PROGRESS TRUE
+    CONFIGURE_COMMAND ""
+    BUILD_COMMAND "")
+endif()
+
+FetchContent_GetProperties(tml_fa4)
+if(NOT tml_fa4_POPULATED)
+  FetchContent_Populate(tml_fa4)
+endif()
+message(STATUS "tml-fa4 is available at ${tml_fa4_SOURCE_DIR}")
+
+add_custom_target(tml_fa4)
+
+# Install into a private namespace so this implementation cannot shadow the
+# flash_attn package used by vLLM's standard attention backends.
+install(CODE "
+  file(GLOB_RECURSE TML_FA4_PY_FILES
+    \"${tml_fa4_SOURCE_DIR}/flash_attn/cute/*.py\")
+  foreach(SRC_FILE \${TML_FA4_PY_FILES})
+    file(RELATIVE_PATH REL_PATH
+      \"${tml_fa4_SOURCE_DIR}/flash_attn/cute\" \${SRC_FILE})
+    set(DST_FILE
+      \"\${CMAKE_INSTALL_PREFIX}/vllm/third_party/tml_fa4/\${REL_PATH}\")
+    get_filename_component(DST_DIR \${DST_FILE} DIRECTORY)
+    file(MAKE_DIRECTORY \${DST_DIR})
+    file(READ \${SRC_FILE} FILE_CONTENTS)
+    string(REPLACE
+      \"flash_attn.cute\"
+      \"vllm.third_party.tml_fa4\"
+      FILE_CONTENTS \"\${FILE_CONTENTS}\")
+    file(WRITE \${DST_FILE} \"\${FILE_CONTENTS}\")
+  endforeach()
+" COMPONENT tml_fa4)

cmake/external_projects/vllm_flash_attn.cmake

@@ -39,7 +39,7 @@ else()
   FetchContent_Declare(
           vllm-flash-attn
           GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG b3964b1d8b95d8e8447435668ab169a2700bab65
+          GIT_TAG 168920233059c48de6199e2cda74003b2ce3d199
           GIT_PROGRESS TRUE
           # Don't share the vllm-flash-attn build between build types
           BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn

codecov.yml

@@ -10,3 +10,16 @@ fixes:
   - "/usr/local/lib/python3.*/site-packages/vllm/::vllm/"
   - "/usr/lib/python3.*/dist-packages/vllm/::vllm/"
   - "/usr/lib/python3.*/site-packages/vllm/::vllm/"
+  # Map Rust sources built in the E2E image and on Buildkite agents.
+  - "/workspace/rust/::rust/"
+  - "/var/lib/buildkite-agent/.*/rust/::rust/"
+
+flags:
+  rust-unit:
+    paths:
+      - rust/
+    carryforward: false
+  rust-e2e:
+    paths:
+      - rust/
+    carryforward: false

csrc/cache.h

@@ -67,9 +67,8 @@ void cp_gather_and_upconvert_fp8_kv_cache(
     torch::Tensor const& src_cache,         // [NUM_BLOCKS, BLOCK_SIZE, 656]
     torch::Tensor const& dst,               // [TOT_TOKENS, 576]
     torch::Tensor const& block_table,       // [BATCH, BLOCK_INDICES]
-    torch::Tensor const& seq_lens,          // [BATCH]
     torch::Tensor const& workspace_starts,  // [BATCH]
-    int64_t batch_size);
+    int64_t batch_size, std::optional<torch::Tensor> seq_starts = std::nullopt);
 
 // Indexer K quantization and cache function
 void indexer_k_quant_and_cache(

csrc/cpu/cpu_attn_vec.hpp

@@ -102,7 +102,9 @@ class TileGemm82 {
     kv_cache_t* __restrict__ curr_b = b_tile;
 
     for (int32_t k = 0; k < dynamic_k_size; ++k) {
-      auto [fp32_b_0_reg, fp32_b_1_reg] = load_b_pair_vec(curr_b);
+      auto fp32_b_regs = load_b_pair_vec(curr_b);
+      auto fp32_b_0_reg = fp32_b_regs.first;
+      auto fp32_b_1_reg = fp32_b_regs.second;
 
       float* __restrict__ curr_m_a = curr_a;
       vec_op::unroll_loop<int32_t, M>([&](int32_t i) {

csrc/cpu/cpu_types_vsx.hpp

@@ -336,13 +336,14 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
     reg.val[1] = fp16_to_fp32_bits(raw_lo);
   }
   float reduce_sum() const {
-    AliasReg ar;
-    ar.reg = reg;
-    float result = 0;
-    unroll_loop<int, VEC_ELEM_NUM>(
-        [&result, &ar](int i) { result += ar.values[i]; });
-
-    return result;
+    // VSX horizontal reduction: 3 vector ops instead of 8 scalar adds.
+    // Step 1: pairwise sum of the two 4-wide halves
+    __vector float s = vec_add(reg.val[0], reg.val[1]);
+    // Step 2: rotate by 8 bytes (2 floats) and add
+    s = vec_add(s, vec_sld(s, s, 8));
+    // Step 3: rotate by 4 bytes (1 float) and add  => all lanes hold total
+    s = vec_add(s, vec_sld(s, s, 4));
+    return vec_extract(s, 0);
   }
   FP32Vec8 exp() const {
     f32x4x2_t out;

csrc/cpu/mamba_cpu.cpp

@@ -0,0 +1,285 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+//
+// CPU at::Tensor wrappers for Mamba decode-step kernels defined in
+// mamba_kernels.hpp.
+
+#include "cpu/mamba_kernels.hpp"
+
+#include <ATen/ATen.h>
+#include <torch/library.h>
+#include <c10/util/Optional.h>
+
+#include "cpu_types.hpp"
+
+// ---------------------------------------------------------------------------
+// causal_conv1d_update
+// ---------------------------------------------------------------------------
+at::Tensor causal_conv1d_update_cpu_impl(
+    at::Tensor& x, at::Tensor& conv_state, const at::Tensor& weight,
+    const c10::optional<at::Tensor>& bias,
+    const c10::optional<std::string>& activation,
+    const c10::optional<at::Tensor>& conv_state_indices,
+    const c10::optional<at::Tensor>& query_start_loc, int64_t pad_slot_id) {
+  bool do_silu = false;
+  if (activation.has_value()) {
+    const std::string& act = activation.value();
+    do_silu = (act == "silu" || act == "swish");
+  }
+
+  at::ScalarType dtype = x.scalar_type();
+
+  // Input x: contiguous in native dtype.
+  at::Tensor x_c = x.is_contiguous() ? x : x.contiguous();
+
+  // conv_state: NEVER copy the full paged tensor just for layout reasons.
+  // If the dtype matches we work directly on conv_state (contiguous or not)
+  // by extracting strides and passing them to the kernel.
+  // Only a dtype-conversion copy is made when types differ (rare for BF16).
+  bool state_type_ok = (conv_state.scalar_type() == dtype);
+  at::Tensor state_c = state_type_ok ? conv_state : conv_state.to(dtype);
+  // state_c and conv_state may be non-contiguous — that is intentional.
+
+  // Weight: coerce to same dtype if needed (should match in practice)
+  at::Tensor w_c =
+      (weight.scalar_type() != dtype)
+          ? weight.to(dtype).contiguous()
+          : (weight.is_contiguous() ? weight : weight.contiguous());
+
+  // Bias stays float32 (small scalar, used only for fp32 accumulation)
+  at::Tensor bias_f32;
+  if (bias.has_value() && bias.value().defined())
+    bias_f32 = bias.value().to(at::kFloat).contiguous();
+
+  int64_t batch = x_c.size(0);
+  int64_t dim = x_c.size(1);
+  int64_t seqlen = (x_c.dim() == 3) ? x_c.size(2) : 1;
+  int64_t width = w_c.size(1);
+  int64_t state_len = state_c.size(2);
+
+  // Extract strides — works for contiguous AND non-contiguous (transposed)
+  // state. stride(0): between cache slots (e.g. num_slots × dim × width-1 in
+  // contiguous) stride(1): between conv channels (dim stride) stride(2):
+  // between state elements (=1 when contiguous, =dim when transposed)
+  int64_t stride_s_slot = state_c.stride(0);
+  int64_t stride_s_dim = state_c.stride(1);
+  int64_t stride_s_state = state_c.stride(2);
+
+  at::Tensor out = x_c.clone();  // native dtype, no float32 alloc
+
+  const int32_t* cache_idx_ptr = nullptr;
+  at::Tensor cache_idx_int;
+  if (conv_state_indices.has_value()) {
+    cache_idx_int = conv_state_indices.value().to(at::kInt).contiguous();
+    cache_idx_ptr = cache_idx_int.data_ptr<int32_t>();
+  }
+
+  VLLM_DISPATCH_FLOATING_TYPES(dtype, "causal_conv1d_update", [&] {
+    mamba_cpu::causal_conv1d_update_kernel<scalar_t>(
+        x_c.data_ptr<scalar_t>(), state_c.data_ptr<scalar_t>(), stride_s_slot,
+        stride_s_dim, stride_s_state, w_c.data_ptr<scalar_t>(),
+        bias_f32.defined() ? bias_f32.data_ptr<float>() : nullptr,
+        out.data_ptr<scalar_t>(), cache_idx_ptr,
+        static_cast<int32_t>(pad_slot_id), batch, dim, seqlen, width, state_len,
+        do_silu);
+  });
+
+  // Write back only when a type-conversion copy was made.
+  // Layout-only non-contiguity is handled via strides above — no copy needed.
+  if (!state_type_ok) conv_state.copy_(state_c);
+
+  return out;
+}
+
+// ---------------------------------------------------------------------------
+// selective_state_update
+// ---------------------------------------------------------------------------
+void selective_state_update_cpu_impl(
+    at::Tensor& state,    // (nstates, nheads, dim, dstate)
+    const at::Tensor& x,  // (N, nheads, dim)
+    const at::Tensor& dt, const at::Tensor& A, const at::Tensor& B,
+    const at::Tensor& C, const c10::optional<at::Tensor>& D,
+    const c10::optional<at::Tensor>& z,
+    const c10::optional<at::Tensor>& dt_bias, bool dt_softplus,
+    const c10::optional<at::Tensor>& state_batch_indices,
+    const c10::optional<at::Tensor>& dst_state_batch_indices,
+    int64_t null_block_id, at::Tensor& out,
+    const c10::optional<at::Tensor>& num_accepted_tokens,
+    const c10::optional<at::Tensor>& cu_seqlens) {
+  at::ScalarType state_type = state.scalar_type();
+  at::ScalarType input_type = x.scalar_type();
+
+  // x, B, C must be contiguous and match input_type
+  auto ensure_input = [input_type](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = (t.scalar_type() != input_type) ? t.to(input_type) : t;
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+  at::Tensor x_in = ensure_input(x);
+  at::Tensor B_in = ensure_input(B);
+  at::Tensor C_in = ensure_input(C);
+  at::Tensor z_in;
+  if (z.has_value() && z.value().defined()) z_in = ensure_input(z.value());
+
+  // A, D, dt_bias are float32 model parameters that arrive here as expanded
+  // tensors, e.g. A is (nheads, head_dim, dstate) with strides (1, 0, 0).
+  // We need just the scalar value per head as a (nheads,) 1-D array so that
+  // A_ptr[h] in the kernel correctly reads head h's value.
+  //
+  // Strategy: peel trailing expanded (stride=0) dims via .select(), which is
+  // a zero-copy view.  For A: (nheads, head_dim, dstate) strides (1,0,0)
+  //   → .select(2,0) → (nheads, head_dim) strides (1,0)
+  //   → .select(1,0) → (nheads,)          stride  (1,)   ← contiguous, free.
+  // No allocation, no type conversion (A is already float32).
+  auto to_per_head_1d_f32 = [](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = t;
+    // Peel trailing dimensions that are broadcast (stride=0 or size=1)
+    while (r.dim() > 1) r = r.select(r.dim() - 1, 0);
+    if (r.scalar_type() != at::kFloat) r = r.to(at::kFloat);
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+
+  at::Tensor A_f32 = to_per_head_1d_f32(A);  // (nheads,) float32
+  at::Tensor D_f32, dt_bias_f32;
+  if (D.has_value() && D.value().defined())
+    D_f32 = to_per_head_1d_f32(D.value());
+  if (dt_bias.has_value() && dt_bias.value().defined())
+    dt_bias_f32 = to_per_head_1d_f32(dt_bias.value());
+
+  // dt: reduce (N, nheads, head_dim) expanded tensor → (N, nheads) BEFORE
+  // the type conversion so we convert head_dim x fewer elements.
+  at::Tensor dt_f32;
+  {
+    // If dt was expanded to (N, nheads, head_dim) with stride-0 in dim 2,
+    // take a zero-copy view of index 0 along that dim first.
+    at::Tensor t2 = (dt.dim() == 3) ? dt.select(2, 0) : dt;  // (N, nheads)
+    at::Tensor t3 = (t2.scalar_type() != at::kFloat) ? t2.to(at::kFloat) : t2;
+    dt_f32 = t3.is_contiguous() ? t3 : t3.contiguous();
+  }
+
+  int64_t nheads = state.size(1);
+  int64_t dim = state.size(2);
+  int64_t dstate = state.size(3);
+  int64_t N = (cu_seqlens.has_value() && cu_seqlens.value().defined())
+                  ? cu_seqlens.value().size(0) - 1
+                  : x_in.size(0);
+  int64_t ngroups = B_in.size(1);
+
+  // Strides
+  int64_t stride_state_n = state.stride(0);
+  int64_t stride_state_h = state.stride(1);
+  int64_t stride_state_d = state.stride(2);
+  int64_t stride_x_n = x_in.stride(0);
+  int64_t stride_x_h = x_in.stride(1);
+  int64_t stride_dt_n = dt_f32.stride(0);  // dt is (N, nheads)
+  int64_t stride_BC_n = B_in.stride(0);
+  int64_t stride_BC_g = B_in.stride(1);
+  int64_t stride_out_n = out.stride(0);
+  int64_t stride_out_h = out.stride(1);
+
+  // Optional index pointers
+  auto get_int32_ptr =
+      [](const c10::optional<at::Tensor>& opt) -> const int32_t* {
+    return (opt.has_value() && opt.value().defined())
+               ? opt.value().data_ptr<int32_t>()
+               : nullptr;
+  };
+  const int32_t* sbi_ptr = get_int32_ptr(state_batch_indices);
+  const int32_t* dsbi_ptr = get_int32_ptr(dst_state_batch_indices);
+  const int32_t* nat_ptr = get_int32_ptr(num_accepted_tokens);
+  const int32_t* csl_ptr = get_int32_ptr(cu_seqlens);
+
+  // Dispatch on (state_t, input_t, out_t): write directly into `out`
+  // without any intermediate float32 buffer.
+  VLLM_DISPATCH_FLOATING_TYPES(state_type, "ssu_state", [&] {
+    using state_t = scalar_t;
+    VLLM_DISPATCH_FLOATING_TYPES(input_type, "ssu_input", [&] {
+      using input_t = scalar_t;
+      VLLM_DISPATCH_FLOATING_TYPES(out.scalar_type(), "ssu_out", [&] {
+        using out_t = scalar_t;
+        mamba_cpu::selective_state_update_kernel<state_t, input_t, out_t>(
+            state.data_ptr<state_t>(), stride_state_n, stride_state_h,
+            stride_state_d, x_in.data_ptr<input_t>(), stride_x_n, stride_x_h,
+            dt_f32.data_ptr<float>(), stride_dt_n, A_f32.data_ptr<float>(),
+            B_in.data_ptr<input_t>(), C_in.data_ptr<input_t>(), stride_BC_n,
+            stride_BC_g, D_f32.defined() ? D_f32.data_ptr<float>() : nullptr,
+            z_in.defined() ? z_in.data_ptr<input_t>() : nullptr,
+            dt_bias_f32.defined() ? dt_bias_f32.data_ptr<float>() : nullptr,
+            out.data_ptr<out_t>(), stride_out_n, stride_out_h, sbi_ptr,
+            dsbi_ptr, static_cast<int32_t>(null_block_id), nat_ptr, csl_ptr, N,
+            nheads, ngroups, dim, dstate, dt_softplus);
+      });
+    });
+  });
+}
+
+// ---------------------------------------------------------------------------
+// mamba_chunk_scan_fwd_cpu
+// ---------------------------------------------------------------------------
+void mamba_chunk_scan_fwd_cpu_impl(
+    at::Tensor& out,  // [seqlen, nheads, headdim] — pre-allocated by caller
+    at::Tensor&
+        final_states,     // [batch, nheads, headdim, dstate] float32 contiguous
+    const at::Tensor& x,  // [seqlen, nheads, headdim]
+    const at::Tensor&
+        dt,  // [seqlen, nheads] float32 (preprocessed: bias+softplus+clamp)
+    const at::Tensor& A,                 // [nheads] float32
+    const at::Tensor& B,                 // [seqlen, ngroups, dstate]
+    const at::Tensor& C,                 // [seqlen, ngroups, dstate]
+    const c10::optional<at::Tensor>& D,  // [nheads] float32 (optional)
+    const c10::optional<at::Tensor>& z,  // [seqlen, nheads, headdim] (optional)
+    const at::Tensor& cu_seqlens         // [batch+1] int32
+) {
+  const at::ScalarType input_type = x.scalar_type();
+
+  auto ensure_contig = [input_type](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = (t.scalar_type() != input_type) ? t.to(input_type) : t;
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+  at::Tensor x_in = ensure_contig(x);
+  at::Tensor B_in = ensure_contig(B);
+  at::Tensor C_in = ensure_contig(C);
+  at::Tensor z_in;
+  if (z.has_value() && z.value().defined()) z_in = ensure_contig(z.value());
+
+  // A and D are float32 model parameters, potentially broadcast-expanded.
+  // Strip trailing broadcast dims to get a contiguous (nheads,) array.
+  auto to_per_head_f32 = [](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = t;
+    while (r.dim() > 1) r = r.select(r.dim() - 1, 0);
+    if (r.scalar_type() != at::kFloat) r = r.to(at::kFloat);
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+  at::Tensor A_f32 = to_per_head_f32(A);
+  at::Tensor D_f32;
+  if (D.has_value() && D.value().defined()) D_f32 = to_per_head_f32(D.value());
+
+  // dt: [seqlen, nheads] float32 — caller has applied bias+softplus+clamp in
+  // Python.
+  at::Tensor dt_c = dt.is_contiguous() ? dt : dt.contiguous();
+  if (dt_c.scalar_type() != at::kFloat) dt_c = dt_c.to(at::kFloat);
+
+  at::Tensor cu_int = cu_seqlens.to(at::kInt).contiguous();
+
+  const int64_t batch = final_states.size(0);
+  const int64_t nheads = final_states.size(1);
+  const int64_t headdim = final_states.size(2);
+  const int64_t dstate = final_states.size(3);
+  const int64_t ngroups = B_in.size(1);
+
+  TORCH_CHECK(final_states.is_contiguous(),
+              "mamba_chunk_scan_fwd_cpu: final_states must be contiguous");
+  TORCH_CHECK(out.is_contiguous(),
+              "mamba_chunk_scan_fwd_cpu: out must be contiguous (writes via "
+              "raw data_ptr)");
+
+  VLLM_DISPATCH_FLOATING_TYPES(input_type, "mamba_chunk_scan_fwd_cpu", [&] {
+    mamba_cpu::mamba_chunk_scan_fwd_kernel<scalar_t>(
+        final_states.data_ptr<float>(), x_in.data_ptr<scalar_t>(),
+        dt_c.data_ptr<float>(), A_f32.data_ptr<float>(),
+        B_in.data_ptr<scalar_t>(), C_in.data_ptr<scalar_t>(),
+        D_f32.defined() ? D_f32.data_ptr<float>() : nullptr,
+        z_in.defined() ? z_in.data_ptr<scalar_t>() : nullptr,
+        out.data_ptr<scalar_t>(), cu_int.data_ptr<int32_t>(), batch, nheads,
+        ngroups, headdim, dstate);
+  });
+}

csrc/cpu/mamba_kernels.hpp

@@ -0,0 +1,382 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+//
+// Fused CPU vector kernels for Mamba decode-step hotspots:
+//   - causal_conv1d_update  (depthwise 1-D conv state roll + compute)
+//   - selective_state_update (SSM recurrence, single-step)
+
+#pragma once
+
+#include "cpu_types.hpp"
+#include <cmath>
+#include <cstring>
+#include <cstdint>
+#include <algorithm>
+
+namespace mamba_cpu {
+
+// ---------------------------------------------------------------------------
+// causal_conv1d_update — templated for native BF16/FP32
+//
+// state_ptr may point to a NON-CONTIGUOUS paged KV cache tensor.
+// Explicit strides are passed so the kernel writes directly into the
+// correct memory locations without making a contiguous copy of the full
+// paged tensor (which was the source of the 34-41% direct_copy_kernel).
+//
+//   stride_s_slot  = state.stride(0)  — between cache slots
+//   stride_s_dim   = state.stride(1)  — between conv_dim channels
+//   stride_s_state = state.stride(2)  — between state elements
+//
+// When stride_s_state == 1 (contiguous), the memmove fast path is used.
+// ---------------------------------------------------------------------------
+template <typename scalar_t>
+inline void causal_conv1d_update_kernel(
+    const scalar_t* __restrict__ x_ptr, scalar_t* __restrict__ state_ptr,
+    int64_t stride_s_slot, int64_t stride_s_dim, int64_t stride_s_state,
+    const scalar_t* __restrict__ weight_ptr, const float* __restrict__ bias_ptr,
+    scalar_t* __restrict__ out_ptr, const int32_t* __restrict__ cache_idxs,
+    int32_t pad_slot_id, int64_t batch, int64_t dim, int64_t seqlen,
+    int64_t width, int64_t state_len, bool do_silu) {
+#pragma omp parallel for
+  for (int64_t b = 0; b < batch; ++b) {
+    int64_t cache_idx = (cache_idxs != nullptr) ? cache_idxs[b] : b;
+    if (cache_idx == pad_slot_id) continue;
+
+    for (int64_t t = 0; t < seqlen; ++t) {
+      const scalar_t* x_b = x_ptr + (b * dim * seqlen + t);
+      scalar_t* out_b = out_ptr + (b * dim * seqlen + t);
+      // Base of this slot in the (possibly non-contiguous) paged state
+      scalar_t* s_base = state_ptr + cache_idx * stride_s_slot;
+
+      for (int64_t d = 0; d < dim; ++d) {
+        float x_val = static_cast<float>(x_b[d * seqlen]);
+        scalar_t* sd = s_base + d * stride_s_dim;  // start of this dim's state
+        const scalar_t* w = weight_ptr + d * width;
+
+        // Accumulate in float32 for precision
+        float acc = (bias_ptr != nullptr) ? bias_ptr[d] : 0.0f;
+        for (int64_t k = 0; k < state_len; ++k) {
+          acc += static_cast<float>(w[k]) *
+                 static_cast<float>(sd[k * stride_s_state]);
+        }
+        acc += static_cast<float>(w[state_len]) * x_val;
+
+        // Shift state left and append new input.
+        // Use memmove when contiguous (stride==1); element loop otherwise.
+        if (stride_s_state == 1) {
+          if (state_len > 1)
+            std::memmove(sd, sd + 1, (state_len - 1) * sizeof(scalar_t));
+          if (state_len > 0) sd[state_len - 1] = static_cast<scalar_t>(x_val);
+        } else {
+          for (int64_t k = 0; k < state_len - 1; ++k)
+            sd[k * stride_s_state] = sd[(k + 1) * stride_s_state];
+          if (state_len > 0)
+            sd[(state_len - 1) * stride_s_state] = static_cast<scalar_t>(x_val);
+        }
+
+        if (do_silu) {
+          float sigmoid = (acc >= 0) ? 1.0f / (1.0f + std::exp(-acc))
+                                     : std::exp(acc) / (1.0f + std::exp(acc));
+          acc *= sigmoid;
+        }
+        out_b[d * seqlen] = static_cast<scalar_t>(acc);
+      }
+    }
+  }
+}
+
+// ---------------------------------------------------------------------------
+// selective_state_update
+//
+// Template parameters:
+//   state_t  - dtype of ssm_state cache (typically BFloat16)
+//   input_t  - dtype of x, B, C (typically BFloat16)
+//   out_t    - dtype of output tensor (typically BFloat16)
+//              Write directly — no float32 intermediate buffer needed.
+//
+// A, D, dt_bias are accepted as const float* (they are always float32
+// model parameters in Mamba2).  This eliminates the per-call float32→BF16
+// conversion and the .contiguous() materialisation of the broadcast-expand.
+//
+// dt is accepted as a (N, nheads) scalar-per-head tensor, not as the
+// (N, nheads, head_dim) expansion, so no .contiguous() copy is needed.
+// ---------------------------------------------------------------------------
+template <typename state_t, typename input_t, typename out_t = float>
+inline void selective_state_update_kernel(
+    state_t* __restrict__ state_ptr, int64_t stride_state_n,
+    int64_t stride_state_h, int64_t stride_state_d,
+    const input_t* __restrict__ x_ptr, int64_t stride_x_n, int64_t stride_x_h,
+    // dt: (N, nheads) — scalar per head, NOT expanded to head_dim
+    const float* __restrict__ dt_ptr, int64_t stride_dt_n,
+    // A: (nheads,) float32 — scalar per head
+    const float* __restrict__ A_ptr, const input_t* __restrict__ B_ptr,
+    const input_t* __restrict__ C_ptr, int64_t stride_BC_n, int64_t stride_BC_g,
+    // D: (nheads,) float32 — scalar per head (nullptr if not used)
+    const float* __restrict__ D_ptr,
+    // z: same shape as x (optional)
+    const input_t* __restrict__ z_ptr,
+    // dt_bias: (nheads,) float32 — scalar per head (nullptr if not used)
+    const float* __restrict__ dt_bias_ptr, out_t* __restrict__ out_ptr,
+    int64_t stride_out_n, int64_t stride_out_h,
+    const int32_t* __restrict__ state_batch_indices,
+    const int32_t* __restrict__ dst_state_batch_indices, int32_t null_block_id,
+    const int32_t* __restrict__ num_accepted_tokens,
+    const int32_t* __restrict__ cu_seqlens, int64_t N, int64_t nheads,
+    int64_t ngroups, int64_t dim, int64_t dstate, bool dt_softplus) {
+  using state_vec_t = vec_op::vec_t<state_t>;
+  using input_vec_t = vec_op::vec_t<input_t>;
+  constexpr int VEC_ELEM_NUM = 8;
+
+  int64_t nheads_per_group = nheads / ngroups;
+
+  for (int64_t seq_idx = 0; seq_idx < N; ++seq_idx) {
+    int64_t bos, seq_len;
+    if (cu_seqlens != nullptr) {
+      bos = cu_seqlens[seq_idx];
+      seq_len = cu_seqlens[seq_idx + 1] - bos;
+    } else {
+      bos = seq_idx;
+      seq_len = 1;
+    }
+
+    int64_t state_read_idx = (state_batch_indices != nullptr)
+                                 ? state_batch_indices[seq_idx]
+                                 : seq_idx;
+    if (state_read_idx == null_block_id) continue;
+
+    int64_t state_write_idx = (num_accepted_tokens == nullptr)
+                                  ? ((dst_state_batch_indices != nullptr)
+                                         ? dst_state_batch_indices[seq_idx]
+                                         : state_read_idx)
+                                  : -1;
+
+    state_t* s = state_ptr + state_read_idx * stride_state_n;
+
+    for (int64_t t = 0; t < seq_len; ++t) {
+      int64_t token_idx = bos + t;
+      const input_t* x_tok = x_ptr + token_idx * stride_x_n;
+      // dt: (N, nheads) — one float per head per token
+      const float* dt_tok = dt_ptr + token_idx * stride_dt_n;
+      const input_t* B_tok = B_ptr + token_idx * stride_BC_n;
+      const input_t* C_tok = C_ptr + token_idx * stride_BC_n;
+      out_t* out_tok = out_ptr + token_idx * stride_out_n;
+
+#pragma omp parallel for
+      for (int64_t h = 0; h < nheads; ++h) {
+        int64_t g = h / nheads_per_group;
+        const input_t* x_h = x_tok + h * stride_x_h;
+        const input_t* B_g = B_tok + g * stride_BC_g;
+        const input_t* C_g = C_tok + g * stride_BC_g;
+        out_t* out_h = out_tok + h * stride_out_h;
+        state_t* s_h = s + h * stride_state_h;
+
+        // Read scalars-per-head (A, dt, dt_bias, D) — no per-dim indexing
+        float dt_val = dt_tok[h];
+        if (dt_bias_ptr != nullptr) dt_val += dt_bias_ptr[h];
+        if (dt_softplus) {
+          dt_val = (dt_val <= 20.0f) ? std::log1p(std::exp(dt_val)) : dt_val;
+        }
+        const float A_val = A_ptr[h];  // scalar: same for all dim, dstate
+        const float D_val = (D_ptr != nullptr) ? D_ptr[h] : 0.0f;
+
+        const input_t* z_h =
+            (z_ptr != nullptr) ? z_ptr + token_idx * stride_x_n + h * stride_x_h
+                               : nullptr;
+
+        vec_op::FP32Vec8 dt_vec(dt_val);
+        // dA = exp(A * dt): A and dt are SCALARS per head, so compute once
+        // and broadcast.  This saves 7 redundant std::exp() calls that
+        // FP32Vec8::exp() would otherwise make on the broadcast vector.
+        const float dA_scalar = std::exp(A_val * dt_val);
+        vec_op::FP32Vec8 dA(dA_scalar);  // broadcast
+
+        for (int64_t d = 0; d < dim; ++d) {
+          float x_val = static_cast<float>(x_h[d]);
+
+          vec_op::FP32Vec8 out_vec(0.0f);
+          state_t* s_hd = s_h + d * stride_state_d;
+          const input_t* B_g_base = B_g;
+          const input_t* C_g_base = C_g;
+
+          vec_op::FP32Vec8 x_vec(x_val);
+          // dBx = B * x * dt — same dA for all dstate (A is scalar)
+          // s_new = s * dA + B * x * dt
+
+          int64_t n = 0;
+          for (; n <= dstate - VEC_ELEM_NUM; n += VEC_ELEM_NUM) {
+            vec_op::FP32Vec8 B_v((input_vec_t(B_g_base + n)));
+            vec_op::FP32Vec8 C_v((input_vec_t(C_g_base + n)));
+            vec_op::FP32Vec8 s_v((state_vec_t(s_hd + n)));
+
+            vec_op::FP32Vec8 dBx = B_v * x_vec * dt_vec;
+            vec_op::FP32Vec8 s_new = s_v * dA + dBx;
+
+            state_vec_t(s_new).save(s_hd + n);
+            out_vec = out_vec + s_new * C_v;
+          }
+
+          float out_val = out_vec.reduce_sum();
+          for (; n < dstate; ++n) {
+            // Reuse dA_scalar computed once per head — no exp() re-call
+            float dBx = static_cast<float>(B_g[n]) * x_val * dt_val;
+            float s_new = static_cast<float>(s_hd[n]) * dA_scalar + dBx;
+            s_hd[n] = static_cast<state_t>(s_new);
+            out_val += s_new * static_cast<float>(C_g[n]);
+          }
+
+          if (D_ptr != nullptr) out_val += x_val * D_val;
+          if (z_h != nullptr) {
+            float z_val = static_cast<float>(z_h[d]);
+            float sigmoid = (z_val >= 0)
+                                ? 1.0f / (1.0f + std::exp(-z_val))
+                                : std::exp(z_val) / (1.0f + std::exp(z_val));
+            out_val *= z_val * sigmoid;
+          }
+          out_h[d] = static_cast<out_t>(out_val);
+        }
+      }
+
+      if (num_accepted_tokens != nullptr &&
+          dst_state_batch_indices != nullptr) {
+        int64_t token_dst_idx = dst_state_batch_indices[seq_idx * seq_len + t];
+        if (token_dst_idx != null_block_id && token_dst_idx != state_read_idx) {
+          state_t* dst_s = state_ptr + token_dst_idx * stride_state_n;
+          std::memmove(dst_s, s, nheads * stride_state_h * sizeof(state_t));
+        }
+      }
+    }
+
+    if (num_accepted_tokens == nullptr && state_write_idx != null_block_id &&
+        state_write_idx != state_read_idx) {
+      state_t* dst_s = state_ptr + state_write_idx * stride_state_n;
+      std::memmove(dst_s, s, nheads * stride_state_h * sizeof(state_t));
+    }
+  }
+}
+
+// ---------------------------------------------------------------------------
+// mamba_chunk_scan_fwd
+//
+// Prefill SSM recurrence for Mamba2 / SSD models.
+//
+// Key difference from selective_state_update_kernel (decode path):
+//   - #pragma omp parallel for collapse(2) is OUTSIDE the time loop.
+//     Each thread owns a (batch, head) slice and runs the entire token
+//     sequence without any per-token OpenMP synchronisation overhead.
+//     For seqlen=256, this eliminates 256 thread-barrier launches per batch.
+//
+// `dt` arrives already processed (float32, after bias + softplus + clamp)
+// to keep this kernel simple. Preprocessing is done in the Python wrapper.
+//
+// `states_ptr` points to the [batch, nheads, headdim, dstate] float32 output
+// tensor, pre-initialised by the caller (zero or from initial_states).
+// Each (b, h) slice is private to exactly one thread via collapse(2), so
+// there are no write conflicts.
+//
+// D is treated as a scalar per head ([nheads] float32).
+// ---------------------------------------------------------------------------
+template <typename input_t>
+inline void mamba_chunk_scan_fwd_kernel(
+    float* __restrict__ states_ptr,     // [batch, nheads, headdim, dstate] f32
+    const input_t* __restrict__ x_ptr,  // [seqlen, nheads, headdim]
+    const float* __restrict__ dt_ptr,   // [seqlen, nheads] f32 (preprocessed)
+    const float* __restrict__ A_ptr,    // [nheads] f32
+    const input_t* __restrict__ B_ptr,  // [seqlen, ngroups, dstate]
+    const input_t* __restrict__ C_ptr,  // [seqlen, ngroups, dstate]
+    const float* __restrict__ D_ptr,    // [nheads] f32 (nullable)
+    const input_t* __restrict__ z_ptr,  // [seqlen, nheads, headdim] (nullable)
+    input_t* __restrict__ out_ptr,      // [seqlen, nheads, headdim]
+    const int32_t* __restrict__ cu_seqlens,  // [batch+1] int32
+    int64_t batch, int64_t nheads, int64_t ngroups, int64_t headdim,
+    int64_t dstate) {
+  using input_vec_t = vec_op::vec_t<input_t>;
+  constexpr int VEC_ELEM_NUM = 8;
+
+  const int64_t nheads_per_group = nheads / ngroups;
+  // states layout: [batch, nheads, headdim, dstate] contiguous (caller
+  // guarantee)
+  const int64_t stride_s_b = nheads * headdim * dstate;
+  const int64_t stride_s_h = headdim * dstate;
+  // stride_s_d = dstate, stride_s_n = 1
+
+#pragma omp parallel for collapse(2) schedule(static)
+  for (int64_t b = 0; b < batch; ++b) {
+    for (int64_t h = 0; h < nheads; ++h) {
+      const int64_t seq_start = cu_seqlens[b];
+      const int64_t seq_end = cu_seqlens[b + 1];
+      const int64_t g = h / nheads_per_group;
+
+      const float A_val = A_ptr[h];
+      const float D_val = (D_ptr != nullptr) ? D_ptr[h] : 0.0f;
+
+      // Working state slice: states[b, h, :, :] — float32, headdim * dstate.
+      // Fits in L1/L2 for typical dims (e.g. 64*128*4 = 32 KB).
+      float* s_bh = states_ptr + b * stride_s_b + h * stride_s_h;
+
+      for (int64_t t = seq_start; t < seq_end; ++t) {
+        const input_t* x_h = x_ptr + t * nheads * headdim + h * headdim;
+        const float* dt_h = dt_ptr + t * nheads + h;
+        const input_t* B_g = B_ptr + t * ngroups * dstate + g * dstate;
+        const input_t* C_g = C_ptr + t * ngroups * dstate + g * dstate;
+        const input_t* z_h = (z_ptr != nullptr)
+                                 ? z_ptr + t * nheads * headdim + h * headdim
+                                 : nullptr;
+        input_t* out_h = out_ptr + t * nheads * headdim + h * headdim;
+
+        const float dt_val = *dt_h;
+        const float dA_val = std::exp(A_val * dt_val);
+        const vec_op::FP32Vec8 dA_vec(dA_val);  // broadcast scalar
+        const vec_op::FP32Vec8 dt_vec(dt_val);
+
+        for (int64_t d = 0; d < headdim; ++d) {
+          const float x_val = static_cast<float>(x_h[d]);
+          float* s_bhd = s_bh + d * dstate;  // [dstate] contiguous float32
+
+          // Vectorised SSM update + readout over dstate:
+          //   s_new = s * dA + x * dt * B
+          //   y    += s_new * C
+          int64_t n = 0;
+          vec_op::FP32Vec8 y_vec(0.0f);
+          const vec_op::FP32Vec8 x_vec(x_val);
+
+          for (; n <= dstate - VEC_ELEM_NUM; n += VEC_ELEM_NUM) {
+            const vec_op::FP32Vec8 B_v((input_vec_t(B_g + n)));
+            const vec_op::FP32Vec8 C_v((input_vec_t(C_g + n)));
+            const vec_op::FP32Vec8 s_v(s_bhd + n);
+
+            const vec_op::FP32Vec8 s_new = s_v * dA_vec + x_vec * dt_vec * B_v;
+            s_new.save(s_bhd + n);
+            y_vec = y_vec + s_new * C_v;
+          }
+
+          float y_val = y_vec.reduce_sum();
+
+          // Scalar tail for remaining dstate elements
+          for (; n < dstate; ++n) {
+            const float B_n = static_cast<float>(B_g[n]);
+            const float C_n = static_cast<float>(C_g[n]);
+            const float s_new = s_bhd[n] * dA_val + x_val * dt_val * B_n;
+            s_bhd[n] = s_new;
+            y_val += s_new * C_n;
+          }
+
+          // D skip connection (scalar per head)
+          if (D_ptr != nullptr) y_val += x_val * D_val;
+
+          // z gating: out = y * z * sigmoid(z)  (SiLU)
+          if (z_h != nullptr) {
+            const float z_val = static_cast<float>(z_h[d]);
+            const float sigmoid =
+                (z_val >= 0.0f) ? 1.0f / (1.0f + std::exp(-z_val))
+                                : std::exp(z_val) / (1.0f + std::exp(z_val));
+            y_val *= z_val * sigmoid;
+          }
+
+          out_h[d] = static_cast<input_t>(y_val);
+        }
+      }
+    }
+  }
+}
+
+}  // namespace mamba_cpu

csrc/cpu/sgl-kernels/fla.cpp

@@ -1116,6 +1116,164 @@ void fused_sigmoid_gating_delta_rule_update_kernel_impl(
   });
 }
 
+// Speculative-decode variant: processes a varlen batch where each sequence has
+// ``q_len`` draft tokens, runs the recurrence sequentially over those tokens
+// (inside the kernel, so one dispatch handles the whole draft block), reads the
+// initial state from cache slot ``num_accepted-1`` and stores the state *after*
+// token ``t`` into cache slot ``t`` (multi-slot rollback, matching the GPU
+// kernel). Parallelized over (sequence, v_head); the per-sequence token loop is
+// sequential as required by the recurrence.
+template <typename scalar_t, typename param_t>
+void fused_sigmoid_gating_delta_rule_update_spec_kernel_impl(
+    const scalar_t* __restrict__ q_ptr,  // [T, HK, EK]
+    const scalar_t* __restrict__ k_ptr,  // [T, HK, EK]
+    const scalar_t* __restrict__ v_ptr,  // [T, HV, EV]
+    const param_t* __restrict__ A_log_ptr,
+    const scalar_t* __restrict__ a_ptr,  // [T, HV]
+    const scalar_t* __restrict__ dt_bias_ptr,
+    const scalar_t* __restrict__ b_ptr,  // [T, HV]
+    const int32_t* __restrict__ spec_indices_ptr,  // [N, S]
+    const int32_t* __restrict__ num_accepted_ptr,  // [N]
+    const int32_t* __restrict__ cu_seqlens_ptr,     // [N + 1]
+    float* __restrict__ state_ptr,
+    scalar_t* __restrict__ o_ptr,  // [T, HV, EV]
+    float* __restrict__ qk_scale_buf,  // [2, T, HK]
+    int64_t total_tokens,
+    int64_t batch_size,
+    int64_t spec_stride,
+    int64_t num_heads,
+    int64_t head_dim,
+    int64_t v_num_heads,
+    int64_t v_head_dim,
+    int64_t q_strideT,
+    int64_t q_strideH,
+    int64_t k_strideT,
+    int64_t k_strideH,
+    int64_t v_strideT,
+    int64_t v_strideH,
+    int64_t state_slot_stride,
+    bool use_qk_l2norm_in_kernel,
+    double softplus_threshold) {
+  using bVec = at::vec::Vectorized<scalar_t>;
+  using fVec = at::vec::Vectorized<float>;
+  constexpr int64_t VecSize = bVec::size();
+  constexpr int64_t fVecSize = fVec::size();
+  int64_t group_size = v_num_heads / num_heads;
+  double scale = 1 / std::sqrt((double)head_dim);
+  fVec scale_vec = fVec((float)scale);
+
+  if (use_qk_l2norm_in_kernel) {
+    float eps = 1e-5f;
+    at::parallel_for(0, total_tokens * num_heads, 0, [&](int64_t begin, int64_t end) {
+      for (int64_t i = begin; i < end; ++i) {
+        int64_t ti = i / num_heads;
+        int64_t ni = i % num_heads;
+        const scalar_t* qp = q_ptr + ti * q_strideT + ni * q_strideH;
+        const scalar_t* kp = k_ptr + ti * k_strideT + ni * k_strideH;
+        float sq = 0.f, sk = 0.f;
+        for (int64_t d = 0; d < head_dim; ++d) {
+          float qv = (float)qp[d];
+          sq += qv * qv;
+          float kv = (float)kp[d];
+          sk += kv * kv;
+        }
+        qk_scale_buf[ti * num_heads + ni] = 1.f / std::sqrt(sq + eps);
+        qk_scale_buf[total_tokens * num_heads + ti * num_heads + ni] = 1.f / std::sqrt(sk + eps);
+      }
+    });
+  }
+
+  at::parallel_for(0, batch_size * v_num_heads, 0, [&](int64_t begin, int64_t end) {
+    for (int64_t idx = begin; idx < end; ++idx) {
+      int64_t bi = idx / v_num_heads;
+      int64_t ni = idx % v_num_heads;
+      int64_t kh = ni / group_size;
+      int64_t q_start = cu_seqlens_ptr[bi];
+      int64_t q_len = cu_seqlens_ptr[bi + 1] - q_start;
+      if (q_len <= 0) {
+        continue;
+      }
+      int64_t acc = (int64_t)num_accepted_ptr[bi];
+      // Clamp acc-1 to >=0: when num_accepted is 0 the unclamped index reads
+      // out of bounds and yields an arbitrary prev_slot used to index the SSM
+      // state. Mirrors the GPU guard tl.maximum(num_accepted - 1, 0).
+      int64_t prev_slot =
+          (int64_t)spec_indices_ptr[bi * spec_stride + (acc > 0 ? acc - 1 : 0)];
+      for (int64_t t = 0; t < q_len; ++t) {
+        int64_t cur_slot = (int64_t)spec_indices_ptr[bi * spec_stride + t];
+        int64_t token = q_start + t;
+        const float* src = state_ptr + prev_slot * state_slot_stride + ni * head_dim * v_head_dim;
+        float* dst = state_ptr + cur_slot * state_slot_stride + ni * head_dim * v_head_dim;
+        float g_val = -std::exp((float)A_log_ptr[ni]) *
+            softplus((float)a_ptr[token * v_num_heads + ni] + (float)dt_bias_ptr[ni], softplus_threshold);
+        float g_val_exp = std::exp(g_val);
+        fVec g_val_exp_vec = fVec(g_val_exp);
+        float beta_val = 1.f / (1.f + std::exp(-(float)b_ptr[token * v_num_heads + ni]));
+        fVec beta_vec = fVec(beta_val);
+        int64_t q_offset = token * q_strideT + kh * q_strideH;
+        int64_t k_offset = token * k_strideT + kh * k_strideH;
+        float q_scale = use_qk_l2norm_in_kernel ? qk_scale_buf[token * num_heads + kh] : 1.f;
+        float k_scale =
+            use_qk_l2norm_in_kernel ? qk_scale_buf[total_tokens * num_heads + token * num_heads + kh] : 1.f;
+        int64_t v_offset = token * v_strideT + ni * v_strideH;
+        int64_t o_offset = (token * v_num_heads + ni) * v_head_dim;
+        int64_t dvi = 0;
+        for (; dvi <= v_head_dim - VecSize; dvi += VecSize) {
+          fVec kv_mem_vec0 = fVec(0.f);
+          fVec kv_mem_vec1 = fVec(0.f);
+          for (int di = 0; di < head_dim; ++di) {
+            fVec k_val_vec = fVec((float)k_ptr[k_offset + di] * k_scale);
+            fVec sv0 = fVec::loadu(src + di * v_head_dim + dvi);
+            fVec sv1 = fVec::loadu(src + di * v_head_dim + dvi + fVecSize);
+            kv_mem_vec0 = kv_mem_vec0 + sv0 * g_val_exp_vec * k_val_vec;
+            kv_mem_vec1 = kv_mem_vec1 + sv1 * g_val_exp_vec * k_val_vec;
+          }
+          bVec v_bvec = bVec::loadu(v_ptr + v_offset + dvi);
+          fVec v_vec0, v_vec1;
+          std::tie(v_vec0, v_vec1) = at::vec::convert_to_float(v_bvec);
+          fVec dt_vec0 = (v_vec0 - kv_mem_vec0) * beta_vec;
+          fVec dt_vec1 = (v_vec1 - kv_mem_vec1) * beta_vec;
+          fVec o_vec0 = fVec(0.f);
+          fVec o_vec1 = fVec(0.f);
+          for (int di = 0; di < head_dim; ++di) {
+            fVec q_vec = fVec((float)q_ptr[q_offset + di] * q_scale);
+            fVec k_vec = fVec((float)k_ptr[k_offset + di] * k_scale);
+            fVec sv0 = fVec::loadu(src + di * v_head_dim + dvi);
+            fVec sv1 = fVec::loadu(src + di * v_head_dim + dvi + fVecSize);
+            sv0 = sv0 * g_val_exp_vec + k_vec * dt_vec0;
+            sv1 = sv1 * g_val_exp_vec + k_vec * dt_vec1;
+            o_vec0 = o_vec0 + sv0 * q_vec * scale_vec;
+            o_vec1 = o_vec1 + sv1 * q_vec * scale_vec;
+            sv0.store(dst + di * v_head_dim + dvi);
+            sv1.store(dst + di * v_head_dim + dvi + fVecSize);
+          }
+          bVec o_vec = at::vec::convert_from_float<scalar_t>(o_vec0, o_vec1);
+          o_vec.store(o_ptr + o_offset + dvi);
+        }
+        for (; dvi < v_head_dim; ++dvi) {
+          float kv_mem_val = 0.f;
+          for (int di = 0; di < head_dim; ++di) {
+            float k_val = (float)k_ptr[k_offset + di] * k_scale;
+            kv_mem_val += src[di * v_head_dim + dvi] * g_val_exp * k_val;
+          }
+          float v_val = (float)v_ptr[v_offset + dvi];
+          float dt_val = (v_val - kv_mem_val) * beta_val;
+          float o_val = 0.f;
+          for (int di = 0; di < head_dim; ++di) {
+            float q_val = (float)q_ptr[q_offset + di] * q_scale;
+            float k_val = (float)k_ptr[k_offset + di] * k_scale;
+            float ns = src[di * v_head_dim + dvi] * g_val_exp + k_val * dt_val;
+            dst[di * v_head_dim + dvi] = ns;
+            o_val += ns * q_val * scale;
+          }
+          o_ptr[o_offset + dvi] = (scalar_t)o_val;
+        }
+        prev_slot = cur_slot;
+      }
+    }
+  });
+}
+
 template <typename scalar_t>
 void fused_gdn_gating_kernel_impl(
     float* __restrict__ A_log,
@@ -1500,6 +1658,103 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
   return core_attn_out;
 }
 
+// Speculative-decode update (multi-token, multi-slot rollback).
+// q: [T, HK, EK]  k: [T, HK, EK]  v: [T, HV, EV]
+// a: [T, HV]  b: [T, HV]
+// initial_state_source: [N_slots, HV, EK, EV] FP32 (updated in place)
+// spec_state_indices: [batch, S] INT32 (S = num_spec + 1)
+// num_accepted_tokens: [batch] INT32
+// cu_seqlens: [batch + 1] INT32
+// Returns output: [T, HV, EV]
+at::Tensor fused_sigmoid_gating_delta_rule_update_spec_cpu(
+    const at::Tensor& A_log,
+    const at::Tensor& dt_bias,
+    const at::Tensor& q,
+    const at::Tensor& k,
+    const at::Tensor& v,
+    const at::Tensor& a,
+    const at::Tensor& b,
+    at::Tensor& initial_state_source,
+    const at::Tensor& spec_state_indices,
+    const at::Tensor& num_accepted_tokens,
+    const at::Tensor& cu_seqlens,
+    bool use_qk_l2norm_in_kernel,
+    double softplus_beta = 1.0,
+    double softplus_threshold = 20.0) {
+  CHECK_DIM(3, q);
+  CHECK_DIM(3, v);
+  CHECK_LAST_DIM_CONTIGUOUS_INPUT(q);
+  int64_t total_tokens = q.size(0);
+  int64_t num_heads = q.size(1);
+  int64_t head_dim = q.size(2);
+  int64_t v_num_heads = v.size(1);
+  int64_t v_head_dim = v.size(2);
+  int64_t batch_size = cu_seqlens.size(0) - 1;
+  int64_t spec_stride = spec_state_indices.stride(0);
+  CHECK_INPUT_SHAPE_DTYPE<true>(k, {total_tokens, num_heads, head_dim}, q.scalar_type());
+  CHECK_INPUT_SHAPE_DTYPE<true>(v, {total_tokens, v_num_heads, v_head_dim}, q.scalar_type());
+  CHECK_INPUT_SHAPE_DTYPE<true>(a, {total_tokens, v_num_heads}, q.scalar_type());
+  CHECK_INPUT_SHAPE_DTYPE<true>(b, {total_tokens, v_num_heads}, q.scalar_type());
+  CHECK_INPUT_SHAPE_DTYPE<true>(dt_bias, {v_num_heads}, q.scalar_type());
+  CHECK_INPUT_SHAPE_DTYPE<true>(num_accepted_tokens, {batch_size}, at::kInt);
+  CHECK_INPUT_SHAPE_DTYPE<true>(cu_seqlens, {batch_size + 1}, at::kInt);
+  CHECK_EQ(v_num_heads % num_heads, 0);
+  TORCH_CHECK(A_log.sizes() == at::IntArrayRef({v_num_heads}));
+  CHECK_INPUT_SHAPE_DTYPE<true>(
+      initial_state_source,
+      {initial_state_source.size(0), v_num_heads, head_dim, v_head_dim},
+      at::kFloat);
+  TORCH_CHECK(initial_state_source.size(0) >= batch_size,
+      "initial_state_source capacity too small: size(0)=",
+      initial_state_source.size(0), ", batch_size=", batch_size);
+
+  int64_t q_strideT = q.stride(0);
+  int64_t q_strideH = q.stride(1);
+  int64_t k_strideT = k.stride(0);
+  int64_t k_strideH = k.stride(1);
+  int64_t v_strideT = v.stride(0);
+  int64_t v_strideH = v.stride(1);
+  int64_t state_slot_stride = initial_state_source.stride(0);
+
+  at::Tensor o = at::empty({total_tokens, v_num_heads, v_head_dim}, q.options());
+  at::Tensor qk_scale_buf = at::empty({2, total_tokens, num_heads}, at::kFloat);
+
+  CPU_DISPATCH_REDUCED_FLOATING_TYPES_EXT(
+      q.scalar_type(), A_log.scalar_type(), "fused_sigmoid_gating_delta_rule_update_spec_kernel_impl", [&] {
+        fused_sigmoid_gating_delta_rule_update_spec_kernel_impl<scalar_t, param_t>(
+            q.data_ptr<scalar_t>(),
+            k.data_ptr<scalar_t>(),
+            v.data_ptr<scalar_t>(),
+            A_log.data_ptr<param_t>(),
+            a.data_ptr<scalar_t>(),
+            dt_bias.data_ptr<scalar_t>(),
+            b.data_ptr<scalar_t>(),
+            spec_state_indices.data_ptr<int32_t>(),
+            num_accepted_tokens.data_ptr<int32_t>(),
+            cu_seqlens.data_ptr<int32_t>(),
+            initial_state_source.data_ptr<float>(),
+            o.data_ptr<scalar_t>(),
+            qk_scale_buf.data_ptr<float>(),
+            total_tokens,
+            batch_size,
+            spec_stride,
+            num_heads,
+            head_dim,
+            v_num_heads,
+            v_head_dim,
+            q_strideT,
+            q_strideH,
+            k_strideT,
+            k_strideH,
+            v_strideT,
+            v_strideH,
+            state_slot_stride,
+            use_qk_l2norm_in_kernel,
+            softplus_threshold);
+      });
+  return o;
+}
+
 // A_log: [num_v_heads]
 // a: [batch, num_v_heads]
 // b: [batch, num_v_heads]

csrc/cpu/torch_bindings.cpp

@@ -120,6 +120,14 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
     bool use_qk_l2norm_in_kernel, double softplus_beta = 1.0,
     double softplus_threshold = 20.0);
 
+at::Tensor fused_sigmoid_gating_delta_rule_update_spec_cpu(
+    const at::Tensor& A_log, const at::Tensor& dt_bias, const at::Tensor& q,
+    const at::Tensor& k, const at::Tensor& v, const at::Tensor& a,
+    const at::Tensor& b, at::Tensor& initial_state_source,
+    const at::Tensor& spec_state_indices, const at::Tensor& num_accepted_tokens,
+    const at::Tensor& cu_seqlens, bool use_qk_l2norm_in_kernel,
+    double softplus_beta = 1.0, double softplus_threshold = 20.0);
+
 std::tuple<at::Tensor, at::Tensor> fused_gdn_gating_cpu(
     const at::Tensor& A_log, const at::Tensor& a, const at::Tensor& b,
     const at::Tensor& dt_bias);
@@ -205,6 +213,32 @@ void compute_slot_mapping_kernel_impl(const torch::Tensor query_start_loc,
                                       torch::Tensor slot_mapping,
                                       const int64_t block_size);
 
+at::Tensor causal_conv1d_update_cpu_impl(
+    at::Tensor& x, at::Tensor& conv_state, const at::Tensor& weight,
+    const c10::optional<at::Tensor>& bias,
+    const c10::optional<std::string>& activation,
+    const c10::optional<at::Tensor>& conv_state_indices,
+    const c10::optional<at::Tensor>& query_start_loc, int64_t pad_slot_id);
+
+void selective_state_update_cpu_impl(
+    at::Tensor& state, const at::Tensor& x, const at::Tensor& dt,
+    const at::Tensor& A, const at::Tensor& B, const at::Tensor& C,
+    const c10::optional<at::Tensor>& D, const c10::optional<at::Tensor>& z,
+    const c10::optional<at::Tensor>& dt_bias, bool dt_softplus,
+    const c10::optional<at::Tensor>& state_batch_indices,
+    const c10::optional<at::Tensor>& dst_state_batch_indices,
+    int64_t null_block_id, at::Tensor& out,
+    const c10::optional<at::Tensor>& num_accepted_tokens,
+    const c10::optional<at::Tensor>& cu_seqlens);
+
+void mamba_chunk_scan_fwd_cpu_impl(at::Tensor& out, at::Tensor& final_states,
+                                   const at::Tensor& x, const at::Tensor& dt,
+                                   const at::Tensor& A, const at::Tensor& B,
+                                   const at::Tensor& C,
+                                   const c10::optional<at::Tensor>& D,
+                                   const c10::optional<at::Tensor>& z,
+                                   const at::Tensor& cu_seqlens);
+
 void init_cpu_memory_env(std::vector<int64_t> node_ids);
 
 namespace cpu_utils {
@@ -508,6 +542,15 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
       "softplus_threshold=20.0) -> Tensor");
   ops.impl("fused_sigmoid_gating_delta_rule_update_cpu", torch::kCPU,
            &fused_sigmoid_gating_delta_rule_update_cpu);
+  ops.def(
+      "fused_sigmoid_gating_delta_rule_update_spec_cpu(Tensor A_log, Tensor "
+      "dt_bias, Tensor q, Tensor k, Tensor v, Tensor a, Tensor b, "
+      "Tensor(a!) initial_state_source, Tensor spec_state_indices, "
+      "Tensor num_accepted_tokens, Tensor cu_seqlens, bool "
+      "use_qk_l2norm_in_kernel, float softplus_beta=1.0, float "
+      "softplus_threshold=20.0) -> Tensor");
+  ops.impl("fused_sigmoid_gating_delta_rule_update_spec_cpu", torch::kCPU,
+           &fused_sigmoid_gating_delta_rule_update_spec_cpu);
   ops.def(
       "fused_gdn_gating_cpu(Tensor A_log, Tensor a, Tensor b, Tensor dt_bias) "
       "-> (Tensor, Tensor)");
@@ -553,7 +596,8 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
 #endif
 
   // fused moe
-#if defined(__AVX512F__) || (defined(ARM_BF16_SUPPORT))
+#if defined(__AVX512F__) || \
+    (defined(__aarch64__) && !defined(__APPLE__) && defined(ARM_BF16_SUPPORT))
   ops.def(
       "prepack_moe_weight(Tensor weight, Tensor(a1!) packed_weight, str isa) "
       "-> ()");
@@ -564,7 +608,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
       "bool skip_weighted, "
       "str act, str isa) -> ()");
   ops.impl("cpu_fused_moe", torch::kCPU, &cpu_fused_moe);
-#endif  // #if defined(__AVX512F__) || (defined(ARM_BF16_SUPPORT))
+#endif
   ops.def(
       "mla_decode_kvcache("
       "   Tensor! out, Tensor query, Tensor kv_cache,"
@@ -577,6 +621,30 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
       "block_size) -> ()",
       &compute_slot_mapping_kernel_impl);
 
+  // Mamba CPU kernels
+  ops.def(
+      "causal_conv1d_update_cpu_vec("
+      "Tensor(a0!) x, Tensor(a1!) conv_state, Tensor weight, "
+      "Tensor? bias, str? activation, Tensor? conv_state_indices, "
+      "Tensor? query_start_loc, SymInt pad_slot_id) -> Tensor",
+      &causal_conv1d_update_cpu_impl);
+
+  ops.def(
+      "selective_state_update_cpu("
+      "Tensor(a0!) state, Tensor x, Tensor dt, Tensor A, Tensor B, Tensor C, "
+      "Tensor? D, Tensor? z, Tensor? dt_bias, bool dt_softplus, "
+      "Tensor? state_batch_indices, Tensor? dst_state_batch_indices, "
+      "SymInt null_block_id, Tensor(a13!) out, "
+      "Tensor? num_accepted_tokens, Tensor? cu_seqlens) -> ()",
+      &selective_state_update_cpu_impl);
+
+  ops.def(
+      "mamba_chunk_scan_fwd_cpu("
+      "Tensor(a0!) out, Tensor(a1!) final_states, "
+      "Tensor x, Tensor dt, Tensor A, Tensor B, Tensor C, "
+      "Tensor? D, Tensor? z, Tensor cu_seqlens) -> ()",
+      &mamba_chunk_scan_fwd_cpu_impl);
+
   ops.def("init_cpu_memory_env(SymInt[] node_ids) -> ()", &init_cpu_memory_env);
 
   // Speculative decoding kernels

csrc/libtorch_stable/activation_kernels.cu

@@ -669,6 +669,14 @@ __device__ __forceinline__ T gelu_quick_kernel(const T& x) {
   return (T)(((float)x) / (1.0f + expf(-1.702f * (float)x)));
 }
 
+template <typename T>
+__device__ __forceinline__ T relu_squared_kernel(const T& x) {
+  // relu(x)^2 — introduced in https://arxiv.org/abs/2109.08668v2
+  const float f = (float)x;
+  const float val = f > 0.0f ? f : 0.0f;
+  return (T)(val * val);
+}
+
 }  // namespace vllm
 
 void gelu_new(torch::stable::Tensor& out,    // [..., d]
@@ -688,3 +696,9 @@ void gelu_quick(torch::stable::Tensor& out,    // [..., d]
 {
   LAUNCH_ACTIVATION_KERNEL(vllm::gelu_quick_kernel);
 }
+
+void relu_squared(torch::stable::Tensor& out,    // [..., d]
+                  torch::stable::Tensor& input)  // [..., d]
+{
+  LAUNCH_ACTIVATION_KERNEL(vllm::relu_squared_kernel);
+}

csrc/libtorch_stable/cache_kernels.cu

@@ -1174,7 +1174,8 @@ __global__ void cp_gather_and_upconvert_fp8_kv_cache(
     const int32_t num_reqs, const int32_t block_size,
     const int32_t total_tokens, const int64_t block_table_stride,
     const int64_t cache_block_stride, const int64_t cache_entry_stride,
-    const int64_t dst_entry_stride) {
+    const int64_t dst_entry_stride,
+    const int32_t* __restrict__ seq_starts) {  // Optional source offsets
   const int flat_warp_id = (blockIdx.x * blockDim.x + threadIdx.x) >> 5;
   if (flat_warp_id >= total_tokens) return;
   const int lane_id = threadIdx.x & 31;
@@ -1192,7 +1193,8 @@ __global__ void cp_gather_and_upconvert_fp8_kv_cache(
 
   // Compute physical token address via block table
   const int out_token_id = flat_warp_id;
-  const int token_offset = out_token_id - workspace_starts[req_id];
+  int token_offset = out_token_id - workspace_starts[req_id];
+  if (seq_starts != nullptr) token_offset += seq_starts[req_id];
   const int cache_block_idx = token_offset / block_size;
   const int offset_in_block = token_offset % block_size;
   const int physical_block =
@@ -1383,9 +1385,9 @@ void cp_gather_and_upconvert_fp8_kv_cache(
     torch::stable::Tensor const& src_cache,    // [NUM_BLOCKS, BLOCK_SIZE, 656]
     torch::stable::Tensor const& dst,          // [TOT_TOKENS, 576]
     torch::stable::Tensor const& block_table,  // [BATCH, BLOCK_INDICES]
-    torch::stable::Tensor const& seq_lens,     // [BATCH]
     torch::stable::Tensor const& workspace_starts,  // [BATCH]
-    int64_t batch_size) {
+    int64_t batch_size,
+    std::optional<torch::stable::Tensor> seq_starts = std::nullopt) {
   torch::stable::accelerator::DeviceGuard device_guard(
       src_cache.get_device_index());
   const cudaStream_t stream = get_current_cuda_stream();
@@ -1396,20 +1398,25 @@ void cp_gather_and_upconvert_fp8_kv_cache(
   STD_TORCH_CHECK(
       block_table.scalar_type() == torch::headeronly::ScalarType::Int,
       "block_table must be int32");
-  STD_TORCH_CHECK(seq_lens.scalar_type() == torch::headeronly::ScalarType::Int,
-                  "seq_lens must be int32");
   STD_TORCH_CHECK(
       workspace_starts.scalar_type() == torch::headeronly::ScalarType::Int,
       "workspace_starts must be int32");
+  if (seq_starts.has_value()) {
+    STD_TORCH_CHECK(
+        seq_starts.value().scalar_type() == torch::headeronly::ScalarType::Int,
+        "seq_starts must be int32");
+  }
 
   STD_TORCH_CHECK(src_cache.device() == dst.device(),
                   "src_cache and dst must be on the same device");
   STD_TORCH_CHECK(src_cache.device() == block_table.device(),
                   "src_cache and block_table must be on the same device");
-  STD_TORCH_CHECK(src_cache.device() == seq_lens.device(),
-                  "src_cache and seq_lens must be on the same device");
   STD_TORCH_CHECK(src_cache.device() == workspace_starts.device(),
                   "src_cache and workspace_starts must be on the same device");
+  if (seq_starts.has_value()) {
+    STD_TORCH_CHECK(src_cache.device() == seq_starts.value().device(),
+                    "src_cache and seq_starts must be on the same device");
+  }
   auto dtype = src_cache.scalar_type();
   STD_TORCH_CHECK(
       dtype == torch::headeronly::ScalarType::Byte ||               // uint8
@@ -1438,6 +1445,9 @@ void cp_gather_and_upconvert_fp8_kv_cache(
   constexpr int warps_per_block = 8;
   const int grid_size = (total_tokens + warps_per_block - 1) / warps_per_block;
   const int block_size_threads = warps_per_block * 32;  // 256 threads
+  const int32_t* seq_starts_ptr =
+      seq_starts.has_value() ? seq_starts.value().const_data_ptr<int32_t>()
+                             : nullptr;
 
   vllm::cp_gather_and_upconvert_fp8_kv_cache<<<grid_size, block_size_threads, 0,
                                                stream>>>(
@@ -1446,7 +1456,7 @@ void cp_gather_and_upconvert_fp8_kv_cache(
       workspace_starts.const_data_ptr<int32_t>(),
       static_cast<int32_t>(batch_size), block_size, total_tokens,
       block_table_stride, cache_block_stride, cache_entry_stride,
-      dst_entry_stride);
+      dst_entry_stride, seq_starts_ptr);
 }
 
 // Macro to dispatch the kernel based on the data type.

csrc/libtorch_stable/dispatch_utils.h

@@ -21,6 +21,21 @@
 #define VLLM_STABLE_DISPATCH_FP8_CASE(enum_type, ...) \
   THO_PRIVATE_CASE_TYPE_USING_HINT(enum_type, fp8_t, __VA_ARGS__)
 
+// Same idea, for dispatching on an int32/int64 index tensor (e.g. topk_ids)
+// nested inside a value-type dispatch. Named 'idx_t' instead of 'scalar_t'.
+#define VLLM_STABLE_DISPATCH_IDX_CASE(enum_type, ...) \
+  THO_PRIVATE_CASE_TYPE_USING_HINT(enum_type, idx_t, __VA_ARGS__)
+
+#define VLLM_STABLE_DISPATCH_CASE_IDX_TYPES(...)                     \
+  VLLM_STABLE_DISPATCH_IDX_CASE(torch::headeronly::ScalarType::Int,  \
+                                __VA_ARGS__)                         \
+  VLLM_STABLE_DISPATCH_IDX_CASE(torch::headeronly::ScalarType::Long, \
+                                __VA_ARGS__)
+
+#define VLLM_STABLE_DISPATCH_IDX_TYPES(TYPE, NAME, ...) \
+  THO_DISPATCH_SWITCH(TYPE, NAME,                       \
+                      VLLM_STABLE_DISPATCH_CASE_IDX_TYPES(__VA_ARGS__))
+
 #define VLLM_STABLE_DISPATCH_CASE_FLOATING_TYPES(...)                  \
   THO_DISPATCH_CASE(torch::headeronly::ScalarType::Float, __VA_ARGS__) \
   THO_DISPATCH_CASE(torch::headeronly::ScalarType::Half, __VA_ARGS__)  \

csrc/libtorch_stable/fp32_router_gemm.cu

@@ -1,13 +1,18 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 //
-// Router GEMM: activation(T) x weight(fp32) -> fp32, H=3072, E=256, M<=32.
+// Router GEMM: activation(T) x weight(fp32) -> fp32, M<=32, for the
+// supported (E, H) pairs listed at the bottom of this file.
 // Supports bf16 or fp32 activation; weight is always fp32.
 // Adapted from dsv3_router_gemm_float_out.cu.
+// (E=256, H=6144) bf16 uses a B300-tuned wide-block geometry; see
+// invokeFp32RouterGemm.
 
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 
+#include <type_traits>
+
 // ---------------------------------------------------------------------------
 // Load helpers
 // ---------------------------------------------------------------------------
@@ -73,94 +78,113 @@ __device__ __forceinline__ void load_activation<__nv_bfloat16, 8>(
 // InputT : type of activation (float or __nv_bfloat16)
 // Weight is always fp32; output is always fp32.
 // VPT = 16 / sizeof(InputT):  4 for fp32, 8 for bf16
-template <typename InputT, int kBlockSize, int kNumTokens, int kNumExperts,
-          int kHiddenDim>
-__global__ __launch_bounds__(128, 1) void fp32_router_gemm_kernel(
-    float* out, InputT const* mat_a, float const* mat_b) {
+// Each block computes kEPB expert columns; wider blocks / kEPB > 1 are
+// selected per (shape, M) in invokeFp32RouterGemm (B300-tuned, see below).
+// kTGroups > 1 splits the tokens across groups of kBlockSize threads within
+// the block: all groups scan the same weight K-slices (group 0 misses to
+// DRAM, later groups hit L1) so weight traffic stays 1x, while per-thread
+// accumulator registers drop by kTGroups (at M=16 the 32 fp32 accumulators
+// push the kernel to 128 regs/thread and 1 block/SM).
+template <typename InputT, int kBlockSize, int kNumTokens, int kEPB,
+          int kNumExperts, int kHiddenDim, int kTGroups = 1>
+__global__ __launch_bounds__(
+    kBlockSize* kTGroups, 1) void fp32_router_gemm_kernel(float* out,
+                                                          InputT const* mat_a,
+                                                          float const* mat_b) {
   constexpr int VPT = 16 / sizeof(InputT);
   constexpr int k_elems_per_k_iteration = VPT * kBlockSize;
   constexpr int k_iterations = kHiddenDim / k_elems_per_k_iteration;
+  static_assert(kHiddenDim % k_elems_per_k_iteration == 0);
+  static_assert(kNumTokens % kTGroups == 0);
   constexpr int kWarpSize = 32;
-  constexpr int kNumWarps = kBlockSize / kWarpSize;
+  constexpr int kNumWarps = kBlockSize / kWarpSize;  // per token group
+  constexpr int kMG = kNumTokens / kTGroups;         // tokens per group
 
-  int const n_idx = blockIdx.x;
-  int const tid = threadIdx.x;
+  int const e_base = blockIdx.x * kEPB;
+  int const tid = threadIdx.x % kBlockSize;
+  int const m0 = (threadIdx.x / kBlockSize) * kMG;
   int const warpId = tid / kWarpSize;
   int const laneId = tid % kWarpSize;
 
-  float acc[kNumTokens] = {};
-  __shared__ float sm_reduction[kNumTokens][kNumWarps];
-
-  float const* b_col = mat_b + n_idx * kHiddenDim;
-
-  int k_bases[k_iterations];
-#pragma unroll
-  for (int ki = 0; ki < k_iterations; ki++) {
-    k_bases[ki] = ki * k_elems_per_k_iteration + tid * VPT;
-  }
+  float acc[kMG][kEPB] = {};
+  __shared__ float sm_reduction[kNumTokens][kEPB][kNumWarps];
 
 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
   cudaGridDependencySynchronize();
+  // Fire the PDL trigger right after our own wait instead of at kernel end:
+  // a gridsync-ing consumer is unaffected (its wait always targets full grid
+  // completion), while a consumer that reads none of our outputs (e.g. the
+  // NVFP4 activation quant, which reads the same hidden_states) can launch
+  // now and fully overlap this kernel's body.
+  cudaTriggerProgrammaticLaunchCompletion();
 #endif
 
+#pragma unroll
   for (int ki = 0; ki < k_iterations; ki++) {
-    int const k_base = k_bases[ki];
+    int const k_base = ki * k_elems_per_k_iteration + tid * VPT;
 
-    float b_float[VPT];
-    load_weight<VPT>(b_col + k_base, b_float);
+    float b_float[kEPB][VPT];
+#pragma unroll
+    for (int e = 0; e < kEPB; e++) {
+      load_weight<VPT>(mat_b + (e_base + e) * kHiddenDim + k_base, b_float[e]);
+    }
 
 #pragma unroll
-    for (int m_idx = 0; m_idx < kNumTokens; m_idx++) {
+    for (int m_idx = 0; m_idx < kMG; m_idx++) {
       float a_float[VPT];
-      load_activation<InputT, VPT>(mat_a + m_idx * kHiddenDim + k_base,
-                                   a_float);
+      load_activation<InputT, VPT>(
+          mat_a + (size_t)(m0 + m_idx) * kHiddenDim + k_base, a_float);
 #pragma unroll
-      for (int k = 0; k < VPT; k++) {
-        acc[m_idx] += a_float[k] * b_float[k];
+      for (int e = 0; e < kEPB; e++) {
+#pragma unroll
+        for (int k = 0; k < VPT; k++) {
+          acc[m_idx][e] += a_float[k] * b_float[e][k];
+        }
       }
     }
   }
 
   // Warp-level butterfly reduction
 #pragma unroll
-  for (int m = 0; m < kNumTokens; m++) {
-    float sum = acc[m];
-    sum += __shfl_xor_sync(0xffffffff, sum, 16);
-    sum += __shfl_xor_sync(0xffffffff, sum, 8);
-    sum += __shfl_xor_sync(0xffffffff, sum, 4);
-    sum += __shfl_xor_sync(0xffffffff, sum, 2);
-    sum += __shfl_xor_sync(0xffffffff, sum, 1);
-    if (laneId == 0) sm_reduction[m][warpId] = sum;
+  for (int m = 0; m < kMG; m++) {
+#pragma unroll
+    for (int e = 0; e < kEPB; e++) {
+      float sum = acc[m][e];
+      sum += __shfl_xor_sync(0xffffffff, sum, 16);
+      sum += __shfl_xor_sync(0xffffffff, sum, 8);
+      sum += __shfl_xor_sync(0xffffffff, sum, 4);
+      sum += __shfl_xor_sync(0xffffffff, sum, 2);
+      sum += __shfl_xor_sync(0xffffffff, sum, 1);
+      if (laneId == 0) sm_reduction[m0 + m][e][warpId] = sum;
+    }
   }
 
   __syncthreads();
 
-  if (tid == 0) {
+  // Parallel finalize: one thread per (m, e) output.
+  for (int idx = threadIdx.x; idx < kNumTokens * kEPB;
+       idx += kBlockSize * kTGroups) {
+    int const m = idx / kEPB;
+    int const e = idx % kEPB;
+    float final_sum = 0.0f;
 #pragma unroll
-    for (int m = 0; m < kNumTokens; m++) {
-      float final_sum = 0.0f;
-#pragma unroll
-      for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][w];
-      out[m * kNumExperts + n_idx] = final_sum;
-    }
+    for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][e][w];
+    out[m * kNumExperts + e_base + e] = final_sum;
   }
-
-#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-  cudaTriggerProgrammaticLaunchCompletion();
-#endif
 }
 
 // ---------------------------------------------------------------------------
 // Launcher
 // ---------------------------------------------------------------------------
 
-template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
-void invokeFp32RouterGemm(float* output, InputT const* mat_a,
-                          float const* mat_b, cudaStream_t stream) {
-  constexpr int kBlockSize = 128;
+template <typename InputT, int kBlockSize, int kEPB, int kNumTokens,
+          int kNumExperts, int kHiddenDim, int kTGroups = 1>
+static void launchFp32RouterGemm(float* output, InputT const* mat_a,
+                                 float const* mat_b, cudaStream_t stream) {
+  static_assert(kNumExperts % kEPB == 0);
   cudaLaunchConfig_t config;
-  config.gridDim = kNumExperts;
-  config.blockDim = kBlockSize;
+  config.gridDim = kNumExperts / kEPB;
+  config.blockDim = kBlockSize * kTGroups;
   config.dynamicSmemBytes = 0;
   config.stream = stream;
   cudaLaunchAttribute attrs[1];
@@ -168,15 +192,112 @@ void invokeFp32RouterGemm(float* output, InputT const* mat_a,
   attrs[0].val.programmaticStreamSerializationAllowed = 1;
   config.numAttrs = 1;
   config.attrs = attrs;
-  cudaLaunchKernelEx(&config,
-                     fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens,
-                                             kNumExperts, kHiddenDim>,
-                     output, mat_a, mat_b);
+  cudaLaunchKernelEx(
+      &config,
+      fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens, kEPB, kNumExperts,
+                              kHiddenDim, kTGroups>,
+      output, mat_a, mat_b);
+}
+
+static bool isBlackwellFamily() {
+  static int sm = []() {
+    int dev = 0, major = 0, minor = 0;
+    cudaGetDevice(&dev);
+    cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, dev);
+    cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, dev);
+    return major * 10 + minor;
+  }();
+  return sm >= 100;
+}
+
+template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
+void invokeFp32RouterGemm(float* output, InputT const* mat_a,
+                          float const* mat_b, cudaStream_t stream) {
+  // Geometry tuned on B300 per supported shape, bf16 activation, under a
+  // production-fidelity harness (CUDA-graph replay, per-layer cold weights).
+  // GLM-5.2 (E=256, H=6144):
+  //   M <= 4        : BS=768, EPB=1 (2.7us vs cast+cuBLAS 8.1us at M=1)
+  //   M in [5, 15]
+  //   or odd        : BS=384, EPB=2 (crossover vs BS=768 measured in (4, 8))
+  //   M >= 16, even : BS=192, EPB=2, 2 token groups (M=16 4.79us vs 5.04,
+  //                   M=24 5.71 vs 6.38, M=32 6.79 vs 7.72; M=12 loses at
+  //                   0.97x, so the boundary is 16).
+  // Only enabled on the Blackwell family where it was validated; Hopper and
+  // other shapes / fp32 activation keep the legacy geometry.
+  if constexpr (std::is_same_v<InputT, __nv_bfloat16> && kNumExperts == 256 &&
+                kHiddenDim == 6144) {
+    if (!isBlackwellFamily()) {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+      return;
+    }
+    if constexpr (kNumTokens <= 4) {
+      launchFp32RouterGemm<InputT, 768, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    } else if constexpr (kNumTokens >= 16 && kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 192, 2, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else {
+      launchFp32RouterGemm<InputT, 384, 2, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    }
+  } else if constexpr (std::is_same_v<InputT, __nv_bfloat16> &&
+                       kNumExperts == 128 && kHiddenDim == 6144) {
+    // MiniMax-M3. Legacy 128/1 only fills 128 blocks and pays the same
+    // accumulator register cliffs; B300 sweep:
+    //   even M in [6, 10] : BS=384, EPB=1, 2 token groups (1.26-1.43x)
+    //   even M >= 12      : BS=192, EPB=1, 2 token groups (1.59-1.66x at
+    //                       M >= 18; re-measured on B300+B200: 192 also wins
+    //                       M=12/14 by 5-11%% on both, ties 384 at 16)
+    //   M <= 5 / odd      : BS=384, EPB=1 (1.03-1.19x)
+    if (!isBlackwellFamily()) {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+      return;
+    }
+    if constexpr (kNumTokens >= 12 && kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 192, 1, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else if constexpr (kNumTokens >= 6 && kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 384, 1, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else {
+      launchFp32RouterGemm<InputT, 384, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    }
+  } else if constexpr (std::is_same_v<InputT, __nv_bfloat16> &&
+                       kNumExperts == 256 && kHiddenDim == 3072) {
+    // MiniMax-M2/M2.5. The 3.1MB weight is latency-floor bound at small M
+    // (legacy already optimal); token groups win only at even M >= 8
+    // (1.05-1.17x). EPB crossover measured between 12 and 16.
+    if (!isBlackwellFamily()) {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+      return;
+    }
+    if constexpr (kNumTokens >= 14 && kNumTokens % 2 == 0) {
+      // M=14 originally measured 0.91x and stayed on legacy; two fresh
+      // sweeps (B300 dev1 + B200) both put 192/2/tg2 ahead by 3.5-4%%.
+      launchFp32RouterGemm<InputT, 192, 2, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else if constexpr (kNumTokens >= 8 && kNumTokens <= 12 &&
+                         kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 192, 1, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    }
+  } else {
+    launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+        output, mat_a, mat_b, stream);
+  }
 }
 
 // ---------------------------------------------------------------------------
 // Explicit instantiations: M=1..32, for both input types, for the supported
-// (E, H) pairs:  (256, 3072) [MiniMax-M2/M2.5]  and  (128, 6144) [MiniMax-M3].
+// (E, H) pairs:  (256, 3072) [MiniMax-M2/M2.5],  (128, 6144) [MiniMax-M3]
+// and  (256, 6144) [GLM-5.2].
 // ---------------------------------------------------------------------------
 
 #define INSTANTIATE(T, M, E, H)                                    \
@@ -221,6 +342,8 @@ INSTANTIATE_ALL(float, 256, 3072)
 INSTANTIATE_ALL(__nv_bfloat16, 256, 3072)
 INSTANTIATE_ALL(float, 128, 6144)
 INSTANTIATE_ALL(__nv_bfloat16, 128, 6144)
+INSTANTIATE_ALL(float, 256, 6144)
+INSTANTIATE_ALL(__nv_bfloat16, 256, 6144)
 
 #undef INSTANTIATE_ALL
 #undef INSTANTIATE

csrc/libtorch_stable/fp32_router_gemm_entry.cu

@@ -25,10 +25,12 @@ inline int getSMVersion() {
 static constexpr int FP32_MAX_TOKENS = 32;
 
 // Supported (hidden_dim, num_experts) pairs (must match the instantiations in
-// fp32_router_gemm.cu): (3072, 256) for MiniMax-M2/M2.5, (6144, 128) for M3.
+// fp32_router_gemm.cu): (3072, 256) for MiniMax-M2/M2.5, (6144, 128) for M3,
+// (6144, 256) for GLM-5.2.
 static inline bool fp32_router_gemm_supported(int hidden_dim, int num_experts) {
   return (hidden_dim == 3072 && num_experts == 256) ||
-         (hidden_dim == 6144 && num_experts == 128);
+         (hidden_dim == 6144 && num_experts == 128) ||
+         (hidden_dim == 6144 && num_experts == 256);
 }
 
 // Forward declarations — 4 template params must match fp32_router_gemm.cu
@@ -77,6 +79,9 @@ void dispatchFp32RouterGemm(int num_experts, int hidden_dim, int num_tokens,
   } else if (num_experts == 128 && hidden_dim == 6144) {
     Fp32LoopUnroller<InputT, 128, 6144, 1, FP32_MAX_TOKENS>::unroll(
         num_tokens, output, mat_a, mat_b, stream);
+  } else if (num_experts == 256 && hidden_dim == 6144) {
+    Fp32LoopUnroller<InputT, 256, 6144, 1, FP32_MAX_TOKENS>::unroll(
+        num_tokens, output, mat_a, mat_b, stream);
   } else {
     throw std::invalid_argument(
         "fp32_router_gemm: unsupported (hidden_dim, num_experts) pair");
@@ -111,7 +116,7 @@ void fp32_router_gemm(
   STD_TORCH_CHECK(
       fp32_router_gemm_supported(hidden_dim, num_experts),
       "fp32_router_gemm: supported (hidden_dim, num_experts) pairs are "
-      "(3072, 256) and (6144, 128)");
+      "(3072, 256), (6144, 128) and (6144, 256)");
   STD_TORCH_CHECK(num_tokens <= FP32_MAX_TOKENS,
                   "fp32_router_gemm: num_tokens must be in [0, 32]");
   STD_TORCH_CHECK(

csrc/libtorch_stable/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu

@@ -39,12 +39,10 @@
  * The IQ/IK warps address the index_q/index_k sub-blocks *inside* qkv at the
  * fixed physical offsets (nq+2*nkv)*128 and (nq+2*nkv+niq)*128.
  *
- * Dense vs sparse is a compile-time choice via the ``kIsSparse``/``kInsertKV``
- * template bools (3 instantiations: dense <false,false>, sparse-profiling
- * <true,false>, sparse-serving <true,true>), so the index slots, the V slots
- * and the cache inserts fold away entirely on paths that don't use them. The
- * dense layer passes no caches/index: norm+RoPE happens in place and the
- * generic ``Attention`` layer owns the cache write.
+ * Dense vs sparse row layout and index-branch processing are separate template
+ * choices. Skip-index-topk reuse layers still have sparse rows and insert main
+ * K/V cache entries, but compile away index_q/index_k work and index-cache
+ * writes.
  *
  * Q/K and (sparse) index_q/index_k are all rewritten in place inside the fused
  * ``qkv`` tensor.  Caches (bf16) are scatter-written by slot.
@@ -223,10 +221,25 @@ __device__ __forceinline__ void storeCacheElems(
     // model dtype directly. FP8 cache dtypes use the conversion path below.
     storeElems<scalar_t>(reinterpret_cast<scalar_t*>(dst), elems);
   } else {
-#pragma unroll
+#ifdef USE_ROCM
+    // Match ROCm's model-dtype materialization before FP8 cache conversion.
+    using Converter = vllm::_typeConvert<scalar_t>;
+    using rounded_t = typename Converter::hip_type;
+    rounded_t rounded[kElemsPerLane];
+  #pragma unroll
+    for (int i = 0; i < kElemsPerLane; i++) {
+      rounded[i] = Converter::convert(elems[i]);
+    }
+  #pragma unroll
+    for (int i = 0; i < kElemsPerLane; i++) {
+      dst[i] = fp8::scaled_convert<cache_t, rounded_t, kv_dt>(rounded[i], 1.0f);
+    }
+#else
+  #pragma unroll
     for (int i = 0; i < kElemsPerLane; i++) {
       dst[i] = fp8::scaled_convert<cache_t, float, kv_dt>(elems[i], 1.0f);
     }
+#endif
   }
 }
 
@@ -262,20 +275,25 @@ __device__ __forceinline__ void storeElemsFp8(
 // Grid: 1D, ceil(num_tokens * slots_per_token / warps_per_block).
 // Each warp = one (token, slot).
 //
-// `kIsSparse` and `kInsertKV` are compile-time template bools, so all the
-// branch decisions that distinguish the dense layer from the sparse layer
-// (index slots, KV/index inserts, V slots) fold away per instantiation.
-// Three instantiations are built: dense <false,false>, sparse-profiling
-// <true,false> and sparse-serving <true,true>.  Slots per token:
+// `kHasIndex`, `kProcessIndex`, and `kInsertKV` are compile-time template
+// bools, so branch decisions that distinguish the dense layer from the sparse
+// layer (index slots, KV/index inserts, V slots) fold away per instantiation.
+// Slots per token:
 //     Q : nq                            (always — norm+RoPE)
 //     K : nkv                           (always — norm+RoPE; +K-cache insert)
 //     V : nkv  only if kInsertKV        (V-cache insert; no warps in dense)
-//     IQ: niq  only if kIsSparse        (norm+RoPE)
-//     IK: 1    only if kIsSparse        (norm+RoPE; +index-cache insert)
+//     IQ: niq  only if kProcessIndex    (norm+RoPE)
+//     IK: 1    only if kProcessIndex    (norm+RoPE; +index-cache insert)
 // cache_t/kv_dt: main attention KV-cache dtype (auto/fp8). out_idx_t/kFp8Idx:
 // indexer index-K cache + index-Q output dtype (scalar_t or e4m3 byte).
+// kHasIndex means the qkv row is laid out as sparse [q|k|v|index_q|index_k].
+// kProcessIndex controls whether this launch actually norms/ropes the index
+// branch and writes index_q/index_k outputs. Skip-index-topk reuse layers keep
+// kHasIndex=true but set kProcessIndex=false.
 template <typename scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt,
-          typename out_idx_t, bool kIsSparse, bool kInsertKV, bool kFp8Idx>
+          typename out_idx_t, bool kHasIndex, bool kInsertKV,
+          bool kProcessIndex,
+          bool kFp8Idx>
 __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     scalar_t* __restrict__ qkv,  // [N, qkv_row] in/out (packs index if sparse)
     scalar_t* __restrict__ q_out,         // [N, nq*128] contiguous, or nullptr
@@ -288,16 +306,15 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int64_t const* __restrict__ positions,       // [N] i64
     int64_t const* __restrict__ slot_mapping,    // main K/V slots or nullptr
     int64_t const* __restrict__ index_slot_mapping,  // index K slots/nullptr
-    cache_t* __restrict__ kv_cache,       // [nb,2,bs,nkv,128] or nullptr
+    cache_t* __restrict__ kv_cache,       // [nb,nkv,bs,2*128] or nullptr
     out_idx_t* __restrict__ index_cache,  // [nb*bs, 128]; scalar_t or e4m3 byte
     float const eps, int const rotary_dim, int const num_tokens, int const nq,
     int const nkv, int const niq, int const block_size,
-    // kv_cache strides (in elements) for logical shape [nb, 2, bs, nkv, 128].
-    // The head_dim (last) dim is always innermost-contiguous (stride 1), so the
-    // NHD/HND layout choice is fully captured by these four strides: NHD keeps
-    // s_token < s_head, HND swaps them. dim_base addresses head_dim directly.
-    int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
-    int64_t const kv_s_head) {
+    // kv_cache strides (in elements) for logical shape [nb, nkv, bs, 2*128].
+    // The content (last) dim is always innermost-contiguous (stride 1), so the
+    // NHD/HND layout choice is captured by the head/token strides.
+    int64_t const kv_s_block, int64_t const kv_s_head, int64_t const kv_s_token,
+    int64_t const kv_s_dim) {
 #if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
   // _typeConvert<BFloat16> is unavailable on pre-Ampere; the M3 kernel only
   // runs with bf16/fp16 inputs in practice.  Discard the bf16 body there.
@@ -309,9 +326,12 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int const laneId = threadIdx.x % 32;
     int const globalWarpIdx = blockIdx.x * warpsPerBlock + (threadIdx.x / 32);
 
+    static_assert(!kProcessIndex || kHasIndex,
+                  "index processing requires sparse row layout");
+
     // Slot layout (compile-time gated: dense has neither V nor index slots).
     int const v_slots = kInsertKV ? nkv : 0;
-    int const idx_slots = kIsSparse ? niq + 1 : 0;
+    int const idx_slots = kProcessIndex ? niq + 1 : 0;
     int const slots_per_token = nq + nkv + v_slots + idx_slots;
 
     int const tokenIdx = globalWarpIdx / slots_per_token;
@@ -322,14 +342,14 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int const k_begin = nq;
     int const v_begin = nq + nkv;             // valid only when kInsertKV
     int const iq_begin = nq + nkv + v_slots;  // index block start
-    int const ik_slot = iq_begin + niq;       // valid only when kIsSparse
+    int const ik_slot = iq_begin + niq;       // valid only when kProcessIndex
 
     bool const isQ = slot < k_begin;
     bool const isK = slot >= k_begin && slot < v_begin;
     bool isV = false;
     if constexpr (kInsertKV) isV = slot >= v_begin && slot < v_begin + nkv;
     bool isIQ = false, isIK = false;
-    if constexpr (kIsSparse) {
+    if constexpr (kProcessIndex) {
       isIQ = slot >= iq_begin && slot < ik_slot;
       isIK = slot == ik_slot;
     }
@@ -337,7 +357,7 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int const dim_base = laneId * kElemsPerLane;
     // Physical row width of qkv: the dense layer packs [q|k|v]; the sparse
     // layer additionally packs [index_q (niq heads) | index_k (1 head)].
-    int const qkv_row = (nq + 2 * nkv + (kIsSparse ? (niq + 1) : 0)) * kHeadDim;
+    int const qkv_row = (nq + 2 * nkv + (kHasIndex ? (niq + 1) : 0)) * kHeadDim;
 
     // ── Resolve source pointer + per-branch parameters. ────────────────────
     scalar_t* row_ptr = nullptr;       // in-place output location
@@ -368,10 +388,13 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
       row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
                 (nq + 2 * nkv + ih) * kHeadDim;
       norm_w = iq_norm_w;
-    } else {  // isIK -- single shared index key at (nq+2*nkv+niq)*128.
+    } else if (isIK) {
+      // Single shared index key at (nq+2*nkv+niq)*128.
       row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
                 (nq + 2 * nkv + niq) * kHeadDim;
       norm_w = ik_norm_w;
+    } else {
+      return;
     }
 
     // Store destination.  Q and index_q are gathered into dedicated contiguous
@@ -427,9 +450,12 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     // ── Cache inserts (sparse serving only). ───────────────────────────────
     if constexpr (kInsertKV) {
       // Guard (not early-return) so every thread reaches the PDL trigger below.
-      int64_t const sm = (isK || isV)
-                             ? slot_mapping[tokenIdx]
-                             : (isIK ? index_slot_mapping[tokenIdx] : -1);
+      int64_t sm = -1;
+      if (isK || isV) {
+        sm = slot_mapping[tokenIdx];
+      } else if constexpr (kProcessIndex) {
+        if (isIK) sm = index_slot_mapping[tokenIdx];
+      }
       if (sm >= 0) {  // skip padded / unscheduled tokens
         if (isIK) {
           if constexpr (kFp8Idx) {
@@ -438,16 +464,16 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
             storeElems<scalar_t>(index_cache + sm * kHeadDim + dim_base, elems);
           }
         } else if (isK || isV) {
-          // kv_cache logical shape [num_blocks, 2, block_size, nkv, head_dim].
+          // kv_cache logical shape [num_blocks, nkv, block_size, 2*head_dim].
           // Paging is logical (block = sm/block_size, token = sm%block_size);
           // the physical NHD/HND layout is honoured via the passed strides.
           int64_t const b = sm / block_size;
           int64_t const t = sm % block_size;
           int const kv = isK ? 0 : 1;
-          int64_t const off =
-              b * kv_s_block + kv * kv_s_kv + t * kv_s_token + head * kv_s_head;
-          storeCacheElems<scalar_t, cache_t, kv_dt>(kv_cache + off + dim_base,
-                                                    elems);
+          int64_t const off = b * kv_s_block + head * kv_s_head +
+                              t * kv_s_token +
+                              (kv * kHeadDim + dim_base) * kv_s_dim;
+          storeCacheElems<scalar_t, cache_t, kv_dt>(kv_cache + off, elems);
         }
       }
     }
@@ -474,14 +500,14 @@ void launchFusedMiniMaxM3(
     int64_t const* index_slot_mapping, cache_t* kv_cache, void* index_cache,
     float const eps, int const rotary_dim, int const num_tokens, int const nq,
     int const nkv, int const niq, int const block_size,
-    int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
-    int64_t const kv_s_head, bool const has_index, bool const insert_kv,
-    bool const fp8_idx, cudaStream_t stream) {
+    int64_t const kv_s_block, int64_t const kv_s_head, int64_t const kv_s_token,
+    int64_t const kv_s_dim, bool const has_index, bool const insert_kv,
+    bool const process_index, bool const fp8_idx, cudaStream_t stream) {
   // Index outputs are scalar_t (bf16) or e4m3 bytes (uint8_t); reinterpret the
   // void* pointers per instantiation in the LAUNCH macro.
   // Slot count must match the kernel's compile-time gating.
   int const v_slots = insert_kv ? nkv : 0;
-  int const idx_slots = has_index ? niq + 1 : 0;
+  int const idx_slots = process_index ? niq + 1 : 0;
   int const slots_per_token = nq + nkv + v_slots + idx_slots;
 
   constexpr int kBlockSize = 256;
@@ -508,50 +534,59 @@ void launchFusedMiniMaxM3(
   config.attrs = attrs;
   config.numAttrs = (sm_version >= 90) ? 1 : 0;
 
-  #define LAUNCH(IS_SPARSE, INSERT, FP8, OUT_T)                                \
+  #define LAUNCH(HAS_INDEX, INSERT, PROCESS_INDEX, FP8, OUT_T)                 \
     cudaLaunchKernelEx(                                                        \
         &config,                                                               \
         fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, cache_t, kv_dt, OUT_T, \
-                                              IS_SPARSE, INSERT, FP8>,         \
+                                              HAS_INDEX, INSERT,               \
+                                              PROCESS_INDEX, FP8>,             \
         qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w, k_norm_w, \
         iq_norm_w, ik_norm_w, cos_sin_cache, positions, slot_mapping,          \
         index_slot_mapping, kv_cache, reinterpret_cast<OUT_T*>(index_cache),   \
         eps, rotary_dim, num_tokens, nq, nkv, niq, block_size, kv_s_block,     \
-        kv_s_kv, kv_s_token, kv_s_head)
+        kv_s_head, kv_s_token, kv_s_dim)
 #else
   // ROCm: standard kernel launch syntax (no PDL/stream serialization).
   // clang-format off
-  #define LAUNCH(IS_SPARSE, INSERT, FP8, OUT_T)                              \
-    fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, cache_t, kv_dt, OUT_T,   \
-                                          IS_SPARSE, INSERT, FP8>            \
-        <<<grid, kBlockSize, 0, stream>>>(                                   \
-            qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w,     \
-            k_norm_w, iq_norm_w, ik_norm_w, cos_sin_cache, positions,        \
-            slot_mapping, index_slot_mapping, kv_cache,                      \
-            reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim,          \
-            num_tokens, nq, nkv, niq, block_size, kv_s_block, kv_s_kv,       \
-            kv_s_token, kv_s_head)
+  #define LAUNCH(HAS_INDEX, INSERT, PROCESS_INDEX, FP8, OUT_T)              \
+    fusedMiniMaxM3QNormRopeKVInsertKernel<                                  \
+        scalar_t, cache_t, kv_dt, OUT_T, HAS_INDEX, INSERT, PROCESS_INDEX,  \
+        FP8><<<grid, kBlockSize, 0, stream>>>(                              \
+        qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w,        \
+        k_norm_w, iq_norm_w, ik_norm_w, cos_sin_cache, positions,           \
+        slot_mapping, index_slot_mapping, kv_cache,                         \
+        reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim, num_tokens, \
+        nq, nkv, niq, block_size, kv_s_block, kv_s_head, kv_s_token,         \
+        kv_s_dim)
   // clang-format on
 #endif
 
   if (has_index) {
-    if (insert_kv) {
-      if (fp8_idx) {
-        LAUNCH(true, true, true, uint8_t);  // sparse serving, fp8 index outputs
+    if (!process_index) {
+      if (insert_kv) {
+        LAUNCH(true, true, false, false, scalar_t);
       } else {
-        LAUNCH(true, true, false, scalar_t);  // sparse serving, bf16
+        LAUNCH(true, false, false, false, scalar_t);
+      }
+    } else if (insert_kv) {
+      if (fp8_idx) {
+        LAUNCH(true, true, true, true,
+               uint8_t);  // sparse serving, fp8 index outputs
+      } else {
+        LAUNCH(true, true, true, false, scalar_t);  // sparse serving, bf16
       }
     } else {
       if (fp8_idx) {
-        LAUNCH(true, false, true, uint8_t);  // sparse profiling, fp8 index_q
+        LAUNCH(true, false, true, true,
+               uint8_t);  // sparse profiling, fp8 index_q
       } else {
-        LAUNCH(true, false, false, scalar_t);  // sparse profiling, bf16
+        LAUNCH(true, false, true, false, scalar_t);  // sparse profiling, bf16
       }
     }
   } else {
     // Dense layer: never has an index branch and never inserts here (the
     // generic Attention layer owns the KV insert).
-    LAUNCH(false, false, false, scalar_t);
+    LAUNCH(false, false, false, false, scalar_t);
   }
 #undef LAUNCH
 }
@@ -559,6 +594,7 @@ void launchFusedMiniMaxM3(
 }  // namespace minimax_m3_fused_ops
 }  // namespace vllm
 
+// clang-format off
 #define CALL_FUSED_MINIMAX_M3(_RAW_T, CACHE_T, KV_DTYPE)                       \
   vllm::minimax_m3_fused_ops::launchFusedMiniMaxM3<st, CACHE_T, KV_DTYPE>(     \
       reinterpret_cast<st*>(qkv.data_ptr()),                                   \
@@ -568,24 +604,29 @@ void launchFusedMiniMaxM3(
           : nullptr,                                                           \
       reinterpret_cast<st const*>(q_norm_weight.data_ptr()),                   \
       reinterpret_cast<st const*>(k_norm_weight.data_ptr()),                   \
-      has_index ? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr()) \
-                : nullptr,                                                     \
-      has_index ? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr()) \
-                : nullptr,                                                     \
+      process_index                                                            \
+          ? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr())       \
+          : nullptr,                                                           \
+      process_index                                                            \
+          ? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr())       \
+          : nullptr,                                                           \
       reinterpret_cast<st const*>(cos_sin_cache.data_ptr()),                   \
       reinterpret_cast<int64_t const*>(positions.data_ptr()),                  \
       insert_kv ? reinterpret_cast<int64_t const*>(slot_mapping->data_ptr())   \
                 : nullptr,                                                     \
-      insert_kv ? reinterpret_cast<int64_t const*>(                            \
-                      effective_index_slot_mapping->data_ptr())                \
-                : nullptr,                                                     \
+      (insert_kv && process_index)                                             \
+          ? reinterpret_cast<int64_t const*>(                                  \
+                effective_index_slot_mapping->data_ptr())                      \
+          : nullptr,                                                           \
       insert_kv ? reinterpret_cast<CACHE_T*>(kv_cache->data_ptr()) : nullptr,  \
-      (insert_kv && has_index)                                                 \
+      (insert_kv && process_index)                                             \
           ? reinterpret_cast<void*>(index_cache->data_ptr())                   \
           : nullptr,                                                           \
       static_cast<float>(eps), static_cast<int>(rotary_dim), num_tokens, nq,   \
-      nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_kv, kv_s_token, \
-      kv_s_head, has_index, insert_kv, fp8_idx, stream)
+      nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_head,           \
+      kv_s_token, kv_s_dim, has_index, insert_kv, process_index, fp8_idx,       \
+      stream)
+// clang-format on
 
 // ────────────────────────────────────────────────────────────────────────────
 // Torch op wrapper
@@ -602,13 +643,13 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
     int64_t num_index_heads,                                  // niq; 0 => dense
     std::optional<torch::stable::Tensor> slot_mapping,        // [N] i64
     std::optional<torch::stable::Tensor> index_slot_mapping,  // [N] i64
-    std::optional<torch::stable::Tensor> kv_cache,     // [nb,2,bs,nkv,128]
+    std::optional<torch::stable::Tensor> kv_cache,     // [nb,nkv,bs,2*128]
     std::optional<torch::stable::Tensor> index_cache,  // [nb,bs,128]
     int64_t block_size,
     std::optional<torch::stable::Tensor> q_out,  // [N, nq*128] contiguous
     std::optional<torch::stable::Tensor>
         index_q_out,  // [N, niq*128] contiguous
-    const std::string& kv_cache_dtype) {
+    const std::string& kv_cache_dtype, bool skip_index_branch) {
   STD_TORCH_CHECK(qkv.is_cuda() && qkv.is_contiguous(),
                   "qkv must be contiguous CUDA");
   STD_TORCH_CHECK(
@@ -647,6 +688,7 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
   // (1 head)]) right after [q|k|v] in the same row; the dense layer does not.
   bool const has_index = niq > 0;
   bool const insert_kv = kv_cache.has_value();
+  bool const process_index = has_index && !skip_index_branch;
   vllm::Fp8KVCacheDataType const kv_dt =
       vllm::get_fp8_kv_cache_data_type(kv_cache_dtype);
   int const kHeadDim = vllm::minimax_m3_fused_ops::kHeadDim;
@@ -662,7 +704,9 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
   STD_TORCH_CHECK(
       !insert_kv || has_index,
       "insert mode (kv_cache) requires the index branch (sparse layer)");
-  if (has_index) {
+  STD_TORCH_CHECK(has_index || !skip_index_branch,
+                  "skip_index_branch requires sparse qkv rows");
+  if (process_index) {
     STD_TORCH_CHECK(
         index_q_norm_weight.has_value() && index_k_norm_weight.has_value(),
         "index branch requires both index norm weights");
@@ -673,24 +717,26 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
                         index_k_norm_weight->numel() == kHeadDim,
                     "index norm weights must have 128 elements");
   }
-  // kv_cache strides (logical shape [nb, 2, bs, nkv, head_dim]). Read straight
+  // kv_cache strides (logical shape [nb, nkv, bs, 2*head_dim]). Read straight
   // off the tensor so the kernel honours whatever physical layout the attention
-  // backend allocated (NHD: stride order (0,1,2,3,4); HND: (0,1,3,2,4)). No new
+  // backend allocated (NHD: stride order (0,2,1,3); HND: (0,1,2,3)). No new
   // op argument is needed -- the strides ride along with the tensor itself.
-  int64_t kv_s_block = 0, kv_s_kv = 0, kv_s_token = 0, kv_s_head = 0;
+  int64_t kv_s_block = 0, kv_s_head = 0, kv_s_token = 0, kv_s_dim = 0;
   torch::stable::Tensor const* effective_index_slot_mapping = nullptr;
   if (insert_kv) {
     STD_TORCH_CHECK(
         slot_mapping.has_value() && slot_mapping->is_cuda() &&
             slot_mapping->scalar_type() == torch::headeronly::ScalarType::Long,
         "insert mode requires int64 CUDA slot_mapping");
-    STD_TORCH_CHECK(
-        !index_slot_mapping.has_value() ||
-            (index_slot_mapping->is_cuda() &&
-             index_slot_mapping->scalar_type() ==
-                 torch::headeronly::ScalarType::Long &&
-             index_slot_mapping->numel() == slot_mapping->numel()),
-        "index_slot_mapping must be int64 CUDA with slot_mapping length");
+    if (process_index) {
+      STD_TORCH_CHECK(
+          !index_slot_mapping.has_value() ||
+              (index_slot_mapping->is_cuda() &&
+               index_slot_mapping->scalar_type() ==
+                   torch::headeronly::ScalarType::Long &&
+               index_slot_mapping->numel() == slot_mapping->numel()),
+          "index_slot_mapping must be int64 CUDA with slot_mapping length");
+    }
     // Main attention KV cache: auto matches qkv, fp8 uses uint8 storage.
     if (kv_dt == vllm::Fp8KVCacheDataType::kAuto) {
       STD_TORCH_CHECK(kv_cache->scalar_type() == qkv.scalar_type(),
@@ -701,22 +747,26 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
           "fp8 kv_cache must use uint8 storage");
     }
     // Indexer index-K cache: independent dtype -- qkv dtype or fp8 e4m3.
-    STD_TORCH_CHECK(
-        index_cache.has_value() &&
-            (index_cache->scalar_type() == qkv.scalar_type() ||
-             index_cache->scalar_type() ==
-                 torch::headeronly::ScalarType::Float8_e4m3fn),
-        "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
-    STD_TORCH_CHECK(kv_cache->dim() == 5 && kv_cache->stride(4) == 1,
-                    "kv_cache must be [nb,2,bs,nkv,head_dim] with contiguous "
-                    "head_dim (stride(4)==1)");
+    if (process_index) {
+      STD_TORCH_CHECK(
+          index_cache.has_value() &&
+              (index_cache->scalar_type() == qkv.scalar_type() ||
+               index_cache->scalar_type() ==
+                   torch::headeronly::ScalarType::Float8_e4m3fn),
+          "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
+    }
+    STD_TORCH_CHECK(kv_cache->dim() == 4 && kv_cache->stride(3) == 1,
+                    "kv_cache must be [nb,nkv,bs,2*head_dim] with contiguous "
+                    "content dim (stride(3)==1)");
     kv_s_block = kv_cache->stride(0);
-    kv_s_kv = kv_cache->stride(1);
+    kv_s_head = kv_cache->stride(1);
     kv_s_token = kv_cache->stride(2);
-    kv_s_head = kv_cache->stride(3);
-    effective_index_slot_mapping = index_slot_mapping.has_value()
-                                       ? &index_slot_mapping.value()
-                                       : &slot_mapping.value();
+    kv_s_dim = kv_cache->stride(3);
+    if (process_index) {
+      effective_index_slot_mapping = index_slot_mapping.has_value()
+                                         ? &index_slot_mapping.value()
+                                         : &slot_mapping.value();
+    }
   }
   // Optional contiguous gather targets: when given, the normed/roped q (and
   // index_q) are written here instead of in place, so callers avoid a separate
@@ -731,9 +781,8 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
         "q_out must have num_tokens * num_heads * 128 elements");
   }
   if (index_q_out.has_value()) {
-    STD_TORCH_CHECK(
-        has_index,
-        "index_q_out requires the index branch (num_index_heads > 0)");
+    STD_TORCH_CHECK(process_index,
+                    "index_q_out requires index branch processing");
     STD_TORCH_CHECK(
         index_q_out->is_cuda() && index_q_out->is_contiguous() &&
             (index_q_out->scalar_type() == qkv.scalar_type() ||
@@ -750,8 +799,9 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
   // q/k/v + q_out stay qkv dtype. Both index outputs must agree.
   auto const kFp8 = torch::headeronly::ScalarType::Float8_e4m3fn;
   bool const fp8_idx =
-      (index_cache.has_value() && index_cache->scalar_type() == kFp8) ||
-      (index_q_out.has_value() && index_q_out->scalar_type() == kFp8);
+      process_index &&
+      ((index_cache.has_value() && index_cache->scalar_type() == kFp8) ||
+       (index_q_out.has_value() && index_q_out->scalar_type() == kFp8));
   if (fp8_idx) {
     STD_TORCH_CHECK(
         !index_cache.has_value() || index_cache->scalar_type() == kFp8,

csrc/libtorch_stable/moe/moe_align_sum_kernels.cu

@@ -82,6 +82,21 @@ __global__ void batched_moe_align_block_size_kernel(
 }
 }  // namespace batched_moe_align_block_size
 
+template <typename scalar_t>
+__device__ __forceinline__ int get_local_expert_id(
+    size_t idx, const scalar_t* __restrict__ topk_ids,
+    int32_t* __restrict__ expert_map, int32_t num_experts,
+    bool has_expert_map) {
+  int expert_id = topk_ids[idx];
+  if (expert_id >= num_experts || expert_id < 0) {
+    return -1;
+  }
+  if (has_expert_map) {
+    expert_id = expert_map[expert_id];
+  }
+  return expert_id;
+}
+
 template <typename scalar_t>
 __device__ void _moe_align_block_size(
     const scalar_t* __restrict__ topk_ids,
@@ -126,20 +141,15 @@ __device__ void _moe_align_block_size(
   const size_t stride = blockDim.x;
 
   for (size_t i = tid; i < numel; i += stride) {
-    int expert_id = topk_ids[i];
-    if (expert_id >= num_experts) {
-      continue;
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      int warp_idx = expert_id / experts_per_warp;
+      int expert_offset = expert_id % experts_per_warp;
+      int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
+      atomicAdd(&shared_counts[warp_idx * experts_per_warp + expert_offset],
+                mask);
     }
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid experts
-      if (expert_id == -1) continue;
-    }
-    int warp_idx = expert_id / experts_per_warp;
-    int expert_offset = expert_id % experts_per_warp;
-    int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
-    atomicAdd(&shared_counts[warp_idx * experts_per_warp + expert_offset],
-              mask);
   }
 
   __syncthreads();
@@ -227,14 +237,12 @@ __device__ void _moe_align_block_size_small_batch_expert(
   }
 
   for (size_t i = tid; i < numel; i += stride) {
-    int32_t expert_id = topk_ids[i];
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid expert
-      if (expert_id == -1) continue;
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
+      tokens_cnts[(tid + 1) * num_experts + expert_id] += mask;
     }
-    int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
-    tokens_cnts[(tid + 1) * num_experts + expert_id] += mask;
   }
 
   __syncthreads();
@@ -276,18 +284,16 @@ __device__ void _moe_align_block_size_small_batch_expert(
   }
 
   for (size_t i = tid; i < numel; i += stride) {
-    int32_t expert_id = topk_ids[i];
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid expert
-      if (expert_id == -1) continue;
-    }
-    int32_t rank_post_pad =
-        tokens_cnts[tid * num_experts + expert_id] + cumsum[expert_id];
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      int32_t rank_post_pad =
+          tokens_cnts[tid * num_experts + expert_id] + cumsum[expert_id];
 
-    if (token_mask == nullptr || token_mask[i / topk_num]) {
-      sorted_token_ids[sorted_token_ids_offset + rank_post_pad] = i;
-      ++tokens_cnts[tid * num_experts + expert_id];
+      if (token_mask == nullptr || token_mask[i / topk_num]) {
+        sorted_token_ids[sorted_token_ids_offset + rank_post_pad] = i;
+        ++tokens_cnts[tid * num_experts + expert_id];
+      }
     }
   }
 }
@@ -303,22 +309,15 @@ __device__ void _count_and_sort_expert_tokens(
   const size_t stride = blockDim.x * gridDim.y;
 
   for (size_t i = tid; i < numel; i += stride) {
-    int32_t expert_id = topk_ids[i];
-    if (expert_id >= num_experts) {
-      continue;
-    }
-
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid experts
-      if (expert_id == -1) continue;
-    }
-
-    if (token_mask == nullptr || token_mask[i / topk_num]) {
-      int32_t rank_post_pad = atomicAdd(
-          &cumsum_buffer[(model_offset * (num_experts + 1)) + expert_id], 1);
-      sorted_token_ids[max_num_tokens_padded * model_offset + rank_post_pad] =
-          i;
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      if (token_mask == nullptr || token_mask[i / topk_num]) {
+        int32_t rank_post_pad = atomicAdd(
+            &cumsum_buffer[(model_offset * (num_experts + 1)) + expert_id], 1);
+        sorted_token_ids[max_num_tokens_padded * model_offset + rank_post_pad] =
+            i;
+      }
     }
   }
 }
@@ -361,12 +360,24 @@ __global__ void count_and_sort_expert_tokens_kernel(
 template <typename scalar_t>
 constexpr int MOE_SUM_VEC = 16 / sizeof(scalar_t);
 
-template <typename scalar_t, int TOPK>
+template <typename idx_t>
+__device__ __forceinline__ bool moe_sum_pad_aware_skip(
+    const idx_t* __restrict__ topk_ids, const int32_t* __restrict__ expert_map,
+    int64_t idx) {
+  int64_t expert_id = static_cast<int64_t>(topk_ids[idx]);
+  if (expert_id < 0) return true;
+  if (expert_map != nullptr && expert_map[expert_id] < 0) return true;
+  return false;
+}
+
+template <typename scalar_t, typename idx_t, int TOPK, bool PAD_AWARE>
 __global__ void moe_sum_vec_kernel(
     scalar_t* __restrict__ out,          // [num_tokens, d], contiguous
     const scalar_t* __restrict__ input,  // [num_tokens, topk, d], d contiguous
     const int64_t num_tokens, const int d, const int64_t stride_token,
-    const int64_t stride_topk) {
+    const int64_t stride_topk, const idx_t* __restrict__ topk_ids,
+    const int32_t* __restrict__ expert_map, const int64_t stride_tk_token,
+    const int64_t stride_tk_k) {
   using vec_t = vllm::vec_n_t<scalar_t, MOE_SUM_VEC<scalar_t>>;  // 16-byte pack
   constexpr int VEC = MOE_SUM_VEC<scalar_t>;
   const int64_t n_vec = d / VEC;
@@ -376,6 +387,10 @@ __global__ void moe_sum_vec_kernel(
     const int64_t token = i / n_vec;
     const int64_t v = i % n_vec;
     const scalar_t* in_tok = input + token * stride_token + v * VEC;
+    const idx_t* tk_tok = nullptr;
+    if constexpr (PAD_AWARE) {
+      tk_tok = topk_ids + token * stride_tk_token;
+    }
 
     float acc[VEC];
 #pragma unroll
@@ -383,6 +398,11 @@ __global__ void moe_sum_vec_kernel(
 
 #pragma unroll
     for (int k = 0; k < TOPK; ++k) {
+      if constexpr (PAD_AWARE) {
+        if (moe_sum_pad_aware_skip(tk_tok, expert_map, k * stride_tk_k)) {
+          continue;
+        }
+      }
       vec_t packed = *reinterpret_cast<const vec_t*>(in_tok + k * stride_topk);
 #pragma unroll
       for (int j = 0; j < VEC; ++j) acc[j] += static_cast<float>(packed.val[j]);
@@ -395,13 +415,16 @@ __global__ void moe_sum_vec_kernel(
   }
 }
 
-// Runtime-topk variant of the above.
-template <typename scalar_t>
+// Runtime-topk variant of the above, for topk values outside the templated
+// set.
+template <typename scalar_t, typename idx_t, bool PAD_AWARE>
 __global__ void moe_sum_vec_dynamic_kernel(
     scalar_t* __restrict__ out,          // [num_tokens, d], contiguous
     const scalar_t* __restrict__ input,  // [num_tokens, topk, d], d contiguous
     const int64_t num_tokens, const int d, const int topk,
-    const int64_t stride_token, const int64_t stride_topk) {
+    const int64_t stride_token, const int64_t stride_topk,
+    const idx_t* __restrict__ topk_ids, const int32_t* __restrict__ expert_map,
+    const int64_t stride_tk_token, const int64_t stride_tk_k) {
   using vec_t = vllm::vec_n_t<scalar_t, MOE_SUM_VEC<scalar_t>>;
   constexpr int VEC = MOE_SUM_VEC<scalar_t>;
   const int64_t n_vec = d / VEC;
@@ -411,12 +434,21 @@ __global__ void moe_sum_vec_dynamic_kernel(
     const int64_t token = i / n_vec;
     const int64_t v = i % n_vec;
     const scalar_t* in_tok = input + token * stride_token + v * VEC;
+    const idx_t* tk_tok = nullptr;
+    if constexpr (PAD_AWARE) {
+      tk_tok = topk_ids + token * stride_tk_token;
+    }
 
     float acc[VEC];
 #pragma unroll
     for (int j = 0; j < VEC; ++j) acc[j] = 0.f;
 
     for (int k = 0; k < topk; ++k) {
+      if constexpr (PAD_AWARE) {
+        if (moe_sum_pad_aware_skip(tk_tok, expert_map, k * stride_tk_k)) {
+          continue;
+        }
+      }
       vec_t packed = *reinterpret_cast<const vec_t*>(in_tok + k * stride_topk);
 #pragma unroll
       for (int j = 0; j < VEC; ++j) acc[j] += static_cast<float>(packed.val[j]);
@@ -431,17 +463,28 @@ __global__ void moe_sum_vec_dynamic_kernel(
 
 // Stride-aware scalar fallback: handles unaligned/non-vectorizable hidden dims
 // (including a non-contiguous hidden stride) via per-element strided reads.
-template <typename scalar_t>
+template <typename scalar_t, typename idx_t, bool PAD_AWARE>
 __global__ void moe_sum_scalar_kernel(
     scalar_t* __restrict__ out,          // [num_tokens, d], contiguous
     const scalar_t* __restrict__ input,  // [num_tokens, topk, d]
     const int d, const int topk, const int64_t stride_token,
-    const int64_t stride_topk, const int64_t stride_hidden) {
+    const int64_t stride_topk, const int64_t stride_hidden,
+    const idx_t* __restrict__ topk_ids, const int32_t* __restrict__ expert_map,
+    const int64_t stride_tk_token, const int64_t stride_tk_k) {
   const int64_t token_idx = blockIdx.x;
   const scalar_t* in_tok = input + token_idx * stride_token;
+  const idx_t* tk_tok = nullptr;
+  if constexpr (PAD_AWARE) {
+    tk_tok = topk_ids + token_idx * stride_tk_token;
+  }
   for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
     float x = 0.f;
     for (int k = 0; k < topk; ++k) {
+      if constexpr (PAD_AWARE) {
+        if (moe_sum_pad_aware_skip(tk_tok, expert_map, k * stride_tk_k)) {
+          continue;
+        }
+      }
       x += static_cast<float>(
           VLLM_LDG(&in_tok[k * stride_topk + idx * stride_hidden]));
     }
@@ -712,8 +755,9 @@ void batched_moe_align_block_size(int64_t max_tokens_per_batch,
 }
 
 void moe_sum(torch::stable::Tensor& input,   // [num_tokens, topk, hidden_size]
-             torch::stable::Tensor& output)  // [num_tokens, hidden_size]
-{
+             torch::stable::Tensor& output,  // [num_tokens, hidden_size]
+             std::optional<torch::stable::Tensor> topk_ids,
+             std::optional<torch::stable::Tensor> expert_map) {
   // Output is dense and written in place, so it must be contiguous. The input
   // is read by its strides (no copy); only the hidden dim needs to be
   // contiguous to take the vectorized path.
@@ -732,10 +776,102 @@ void moe_sum(torch::stable::Tensor& input,   // [num_tokens, topk, hidden_size]
   const cudaStream_t stream =
       get_current_cuda_stream(output.get_device_index());
 
-#define LAUNCH_MOE_SUM_VEC(TOPK)                \
-  vllm::moe::moe_sum_vec_kernel<scalar_t, TOPK> \
-      <<<grid, dim3(block), 0, stream>>>(       \
-          out_ptr, in_ptr, num_tokens, hidden_size, stride_token, stride_topk)
+  if (topk_ids.has_value()) {
+    // Pad-aware reduce path
+    const torch::stable::Tensor& tk = topk_ids.value();
+    STD_TORCH_CHECK(tk.size(0) == num_tokens && tk.size(1) == topk,
+                    "moe_sum: topk_ids must have shape [num_tokens, topk]");
+    const int64_t stride_tk_token = tk.stride(0);
+    const int64_t stride_tk_k = tk.stride(1);
+
+    const int32_t* expert_map_ptr = nullptr;
+    if (expert_map.has_value()) {
+      STD_TORCH_CHECK(
+          expert_map->scalar_type() == torch::headeronly::ScalarType::Int,
+          "moe_sum: expert_map must be int32");
+      expert_map_ptr =
+          reinterpret_cast<const int32_t*>(expert_map->const_data_ptr());
+    }
+
+#define LAUNCH_MOE_SUM_PAD_AWARE_VEC(TOPK)                                     \
+  vllm::moe::moe_sum_vec_kernel<scalar_t, idx_t, TOPK, true>                   \
+      <<<grid, dim3(block), 0, stream>>>(                                      \
+          out_ptr, in_ptr, num_tokens, hidden_size, stride_token, stride_topk, \
+          topk_ids_ptr, expert_map_ptr, stride_tk_token, stride_tk_k)
+
+    VLLM_STABLE_DISPATCH_FLOATING_TYPES(
+        input.scalar_type(), "moe_sum_pad_aware", [&] {
+          constexpr int VEC = vllm::moe::MOE_SUM_VEC<scalar_t>;
+          constexpr int WIDTH = VEC * sizeof(scalar_t);
+          auto* out_ptr =
+              reinterpret_cast<scalar_t*>(output.mutable_data_ptr());
+          auto* in_ptr =
+              reinterpret_cast<const scalar_t*>(input.const_data_ptr());
+
+          const bool can_vec =
+              (stride_hidden == 1) && (hidden_size % VEC == 0) &&
+              (stride_token % VEC == 0) && (stride_topk % VEC == 0) &&
+              (reinterpret_cast<uintptr_t>(in_ptr) % WIDTH == 0) &&
+              (reinterpret_cast<uintptr_t>(out_ptr) % WIDTH == 0);
+
+          VLLM_STABLE_DISPATCH_IDX_TYPES(
+              tk.scalar_type(), "moe_sum_pad_aware_idx", [&] {
+                auto* topk_ids_ptr =
+                    reinterpret_cast<const idx_t*>(tk.const_data_ptr());
+                if (can_vec) {
+                  const int64_t n_vec = hidden_size / VEC;
+                  const int64_t total = num_tokens * n_vec;
+                  const int block = 256;
+                  const dim3 grid(
+                      std::min<int64_t>((total + block - 1) / block, 65535));
+                  switch (topk) {
+                    case 1:
+                      LAUNCH_MOE_SUM_PAD_AWARE_VEC(1);
+                      break;
+                    case 2:
+                      LAUNCH_MOE_SUM_PAD_AWARE_VEC(2);
+                      break;
+                    case 4:
+                      LAUNCH_MOE_SUM_PAD_AWARE_VEC(4);
+                      break;
+                    case 6:
+                      LAUNCH_MOE_SUM_PAD_AWARE_VEC(6);
+                      break;
+                    case 8:
+                      LAUNCH_MOE_SUM_PAD_AWARE_VEC(8);
+                      break;
+                    case 9:
+                      LAUNCH_MOE_SUM_PAD_AWARE_VEC(9);
+                      break;
+                    default:
+                      vllm::moe::moe_sum_vec_dynamic_kernel<scalar_t, idx_t,
+                                                            true>
+                          <<<grid, dim3(block), 0, stream>>>(
+                              out_ptr, in_ptr, num_tokens, hidden_size, topk,
+                              stride_token, stride_topk, topk_ids_ptr,
+                              expert_map_ptr, stride_tk_token, stride_tk_k);
+                      break;
+                  }
+                } else {
+                  dim3 grid(num_tokens);
+                  dim3 block(std::min(hidden_size, 1024));
+                  vllm::moe::moe_sum_scalar_kernel<scalar_t, idx_t, true>
+                      <<<grid, block, 0, stream>>>(
+                          out_ptr, in_ptr, hidden_size, topk, stride_token,
+                          stride_topk, stride_hidden, topk_ids_ptr,
+                          expert_map_ptr, stride_tk_token, stride_tk_k);
+                }
+              });
+        });
+#undef LAUNCH_MOE_SUM_PAD_AWARE_VEC
+    return;
+  }
+
+#define LAUNCH_MOE_SUM_VEC(TOPK)                                      \
+  vllm::moe::moe_sum_vec_kernel<scalar_t, int32_t, TOPK, false>       \
+      <<<grid, dim3(block), 0, stream>>>(out_ptr, in_ptr, num_tokens, \
+                                         hidden_size, stride_token,   \
+                                         stride_topk, nullptr, nullptr, 0, 0)
 
   VLLM_STABLE_DISPATCH_FLOATING_TYPES(input.scalar_type(), "moe_sum", [&] {
     constexpr int VEC = vllm::moe::MOE_SUM_VEC<scalar_t>;
@@ -775,18 +911,19 @@ void moe_sum(torch::stable::Tensor& input,   // [num_tokens, topk, hidden_size]
           LAUNCH_MOE_SUM_VEC(9);
           break;
         default:
-          vllm::moe::moe_sum_vec_dynamic_kernel<scalar_t>
-              <<<grid, dim3(block), 0, stream>>>(out_ptr, in_ptr, num_tokens,
-                                                 hidden_size, topk,
-                                                 stride_token, stride_topk);
+          vllm::moe::moe_sum_vec_dynamic_kernel<scalar_t, int32_t, false>
+              <<<grid, dim3(block), 0, stream>>>(
+                  out_ptr, in_ptr, num_tokens, hidden_size, topk, stride_token,
+                  stride_topk, nullptr, nullptr, 0, 0);
           break;
       }
     } else {
       dim3 grid(num_tokens);
       dim3 block(std::min(hidden_size, 1024));
-      vllm::moe::moe_sum_scalar_kernel<scalar_t><<<grid, block, 0, stream>>>(
-          out_ptr, in_ptr, hidden_size, topk, stride_token, stride_topk,
-          stride_hidden);
+      vllm::moe::moe_sum_scalar_kernel<scalar_t, int32_t, false>
+          <<<grid, block, 0, stream>>>(out_ptr, in_ptr, hidden_size, topk,
+                                       stride_token, stride_topk, stride_hidden,
+                                       nullptr, nullptr, 0, 0);
     }
   });
 #undef LAUNCH_MOE_SUM_VEC

csrc/libtorch_stable/moe/moe_ops.h

@@ -27,7 +27,9 @@ void topk_softplus_sqrt(
     const std::optional<torch::stable::Tensor>& input_ids,
     const std::optional<torch::stable::Tensor>& tid2eid);
 
-void moe_sum(torch::stable::Tensor& input, torch::stable::Tensor& output);
+void moe_sum(torch::stable::Tensor& input, torch::stable::Tensor& output,
+             std::optional<torch::stable::Tensor> topk_ids,
+             std::optional<torch::stable::Tensor> expert_map);
 
 void moe_align_block_size(
     torch::stable::Tensor topk_ids, int64_t num_experts, int64_t block_size,

csrc/libtorch_stable/moe/topk_softplus_sqrt_kernels.cu

@@ -44,6 +44,12 @@ typedef __hip_bfloat162 __nv_bfloat162;
 namespace vllm {
 namespace moe {
 
+template <typename HashIndType>
+__device__ __forceinline__ int64_t load_index_as_int64(const HashIndType* ptr,
+                                                       int64_t offset) {
+  return static_cast<int64_t>(ptr[offset]);
+}
+
 /// Aligned array type
 template <typename T,
           /// Number of elements in the array
@@ -65,6 +71,73 @@ __device__ __forceinline__ float toFloat(T value) {
   }
 }
 
+#ifndef USE_ROCM
+// Adapted from:
+// https://github.com/sgl-project/sglang/blob/main/python/sglang/jit_kernel/csrc/deepseek_v4/hash_topk.cuh
+template <typename OutIndType, typename HashIndType>
+__launch_bounds__(128) __global__
+    void dsv4HashTopkSoftplusSqrt(const float* input, float* output,
+                                  OutIndType* indices, int num_rows,
+                                  int num_experts, float routed_scaling_factor,
+                                  const HashIndType* input_ids,
+                                  const HashIndType* tid2eid) {
+  const int warp = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
+  const int lane = threadIdx.x % 32;
+  if (warp >= num_rows) return;
+  const int64_t token_id = load_index_as_int64(input_ids, warp);
+
+  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+  cudaGridDependencySynchronize();
+  #endif
+  int expert = 0;
+  float weight = 0.f;
+  if (lane < 6) {
+    // only load and calculate for 6 experts
+    expert = static_cast<int>(tid2eid[token_id * 6 + lane]);
+    const float x = input[warp * num_experts + expert];
+    weight = sqrtf(fmaxf(x, 0.f) + __logf(1.f + __expf(-fabsf(x))));
+  }
+  float weight_sum = weight;
+  #pragma unroll
+  for (int mask = 16; mask > 0; mask >>= 1) {
+    // sum in warp
+    weight_sum += VLLM_SHFL_XOR_SYNC(weight_sum, mask);
+  }
+
+  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+  cudaTriggerProgrammaticLaunchCompletion();
+  #endif
+  if (lane < 6) {
+    const int offset = warp * 6 + lane;
+    output[offset] =
+        weight * routed_scaling_factor / (weight_sum > 0.f ? weight_sum : 1.f);
+    indices[offset] = static_cast<OutIndType>(expert);
+  }
+}
+
+template <typename OutIndType, typename HashIndType>
+void launchDsv4HashTopk(const float* input, float* output, OutIndType* indices,
+                        int num_rows, int num_experts,
+                        double routed_scaling_factor,
+                        const HashIndType* input_ids,
+                        const HashIndType* tid2eid, cudaStream_t stream) {
+  if (num_rows == 0) return;
+  auto* kernel = &dsv4HashTopkSoftplusSqrt<OutIndType, HashIndType>;
+  cudaLaunchConfig_t config = {};
+  config.gridDim = (num_rows + 3) / 4;
+  config.blockDim = 128;
+  config.stream = stream;
+  cudaLaunchAttribute attr;
+  attr.id = cudaLaunchAttributeProgrammaticStreamSerialization;
+  attr.val.programmaticStreamSerializationAllowed = 1;
+  config.attrs = &attr;
+  config.numAttrs = 1;
+  const float scale = static_cast<float>(routed_scaling_factor);
+  cudaLaunchKernelEx(&config, kernel, input, output, indices, num_rows,
+                     num_experts, scale, input_ids, tid2eid);
+}
+#endif
+
 // ====================== TopK softplus_sqrt things
 // ===============================
 
@@ -86,14 +159,14 @@ __device__ __forceinline__ float toFloat(T value) {
 
 template <int VPT, int NUM_EXPERTS, int WARPS_PER_CTA, int BYTES_PER_LDG,
           int WARP_SIZE_PARAM, bool USE_HASH, typename IndType,
-          typename InputType = float>
+          typename HashIndType, typename InputType = float>
 __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
     void topkGatingSoftplusSqrt(
         const InputType* input, const bool* finished, float* output,
         const int num_rows, IndType* indices, int* source_rows, const int k,
         const int start_expert, const int end_expert, const bool renormalize,
         double routed_scaling_factor, const float* correction_bias,
-        const IndType* input_ids, const IndType* tid2eid) {
+        const HashIndType* input_ids, const HashIndType* tid2eid) {
   static_assert(std::is_same_v<InputType, float> ||
                     std::is_same_v<InputType, __nv_bfloat16> ||
                     std::is_same_v<InputType, __half>,
@@ -240,8 +313,8 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
 
   // Hash MoE path: indices are predetermined from lookup table
   if constexpr (USE_HASH) {
-    const IndType token_id = input_ids[thread_row];
-    const IndType* expert_indices_for_token = tid2eid + token_id * k;
+    const int64_t token_id = load_index_as_int64(input_ids, thread_row);
+    const int64_t token_expert_offset = token_id * static_cast<int64_t>(k);
 #pragma unroll
     for (int ii = 0; ii < VPT; ++ii) {
       float val = row_chunk[ii];
@@ -252,7 +325,8 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
     float selected_sum = 0.f;
 #pragma unroll
     for (int k_idx = 0; k_idx < k; ++k_idx) {
-      const int expert = expert_indices_for_token[k_idx];
+      const int expert = static_cast<int>(
+          load_index_as_int64(tid2eid, token_expert_offset + k_idx));
       const int idx = k * thread_row + k_idx;
       for (int ii = 0; ii < VPT; ++ii) {
         const int group_id = ii / ELTS_PER_LDG;
@@ -261,7 +335,7 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
                                group_id * THREADS_PER_ROW * ELTS_PER_LDG +
                                local_id;
         if (expert == expert_idx) {
-          indices[idx] = expert;
+          indices[idx] = static_cast<IndType>(expert);
           selected_sum += row_chunk[ii];
           break;
         }
@@ -285,7 +359,8 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
 
 #pragma unroll
     for (int k_idx = 0; k_idx < k; ++k_idx) {
-      const int expert = expert_indices_for_token[k_idx];
+      const int expert = static_cast<int>(
+          load_index_as_int64(tid2eid, token_expert_offset + k_idx));
       const int idx = k * thread_row + k_idx;
       for (int ii = 0; ii < VPT; ++ii) {
         const int group_id = ii / ELTS_PER_LDG;
@@ -461,14 +536,15 @@ struct TopkConstants {
   }
 
 template <int EXPERTS, int WARPS_PER_TB, int WARP_SIZE_PARAM,
-          int MAX_BYTES_PER_LDG, typename IndType, typename InputType>
+          int MAX_BYTES_PER_LDG, typename IndType, typename HashIndType,
+          typename InputType>
 void topkGatingSoftplusSqrtLauncherHelper(
     const InputType* input, const bool* finished, float* output,
     IndType* indices, int* source_row, const int num_rows, const int k,
     const int start_expert, const int end_expert, const bool renormalize,
     double routed_scaling_factor, const float* correction_bias,
-    const bool use_hash, const IndType* input_ids, const IndType* tid2eid,
-    cudaStream_t stream) {
+    const bool use_hash, const HashIndType* input_ids,
+    const HashIndType* tid2eid, cudaStream_t stream) {
   static constexpr int BYTES_PER_LDG =
       MIN(MAX_BYTES_PER_LDG, sizeof(InputType) * EXPERTS);
   using Constants =
@@ -481,7 +557,8 @@ void topkGatingSoftplusSqrtLauncherHelper(
   DISPATCH_HASH(use_hash, USE_HASH, {
     auto* kernel =
         &topkGatingSoftplusSqrt<VPT, EXPERTS, WARPS_PER_TB, BYTES_PER_LDG,
-                                WARP_SIZE_PARAM, USE_HASH, IndType, InputType>;
+                                WARP_SIZE_PARAM, USE_HASH, IndType, HashIndType,
+                                InputType>;
 #ifndef USE_ROCM
     cudaLaunchConfig_t config = {};
     config.gridDim = num_blocks;
@@ -538,13 +615,25 @@ void topkGatingSoftplusSqrtLauncherHelper(
     }
 #endif
 
-template <typename IndType, typename InputType>
+template <typename IndType, typename InputType, typename HashIndType = IndType>
 void topkGatingSoftplusSqrtKernelLauncher(
     const InputType* gating_output, float* topk_weights, IndType* topk_indices,
     int* token_expert_indices, const int num_tokens, const int num_experts,
     const int topk, const bool renormalize, double routed_scaling_factor,
-    const float* correction_bias, const bool use_hash, const IndType* input_ids,
-    const IndType* tid2eid, cudaStream_t stream) {
+    const float* correction_bias, const bool use_hash,
+    const HashIndType* input_ids, const HashIndType* tid2eid,
+    cudaStream_t stream) {
+#ifndef USE_ROCM
+  if constexpr (std::is_same_v<InputType, float>) {
+    if (use_hash && topk == 6 && renormalize &&
+        (num_experts == 256 || num_experts == 384)) {
+      launchDsv4HashTopk<IndType, HashIndType>(
+          gating_output, topk_weights, topk_indices, num_tokens, num_experts,
+          routed_scaling_factor, input_ids, tid2eid, stream);
+      return;
+    }
+  }
+#endif
   static constexpr int WARPS_PER_TB = 4;
   static constexpr int BYTES_PER_LDG_POWER_OF_2 = 16;
   // for bfloat16 dtype, we need 4 bytes loading to make sure num_experts
@@ -644,57 +733,55 @@ void dispatch_topk_softplus_sqrt_launch(
   if (correction_bias.has_value()) {
     bias_ptr = correction_bias.value().const_data_ptr<float>();
   }
-  bool use_hash = false;
-  if (tid2eid.has_value()) {
-    STD_TORCH_CHECK(input_ids.has_value(),
-                    "input_ids is required for hash MoE");
-    use_hash = true;
-  }
-  if (topk_indices.scalar_type() == torch::headeronly::ScalarType::Int) {
-    const int* input_ids_ptr = nullptr;
-    const int* tid2eid_ptr = nullptr;
-    if (tid2eid.has_value()) {
-      input_ids_ptr = input_ids.value().const_data_ptr<int>();
-      tid2eid_ptr = tid2eid.value().const_data_ptr<int>();
-    }
 
-    vllm::moe::topkGatingSoftplusSqrtKernelLauncher<int, ComputeType>(
-        gating_output, topk_weights.mutable_data_ptr<float>(),
-        topk_indices.mutable_data_ptr<int>(),
-        token_expert_indices.mutable_data_ptr<int>(), num_tokens, num_experts,
-        topk, renormalize, routed_scaling_factor, bias_ptr, use_hash,
-        input_ids_ptr, tid2eid_ptr, stream);
+  auto launch = [&](auto* topk_indices_ptr) {
+    using OutIndType =
+        typename std::remove_pointer<decltype(topk_indices_ptr)>::type;
+    if (tid2eid.has_value()) {
+      STD_TORCH_CHECK(input_ids.has_value(),
+                      "input_ids is required for hash MoE");
+      STD_TORCH_CHECK(
+          input_ids.value().scalar_type() == tid2eid.value().scalar_type(),
+          "input_ids and tid2eid must have the same dtype");
+      if (tid2eid.value().scalar_type() ==
+          torch::headeronly::ScalarType::Long) {
+        vllm::moe::topkGatingSoftplusSqrtKernelLauncher<OutIndType, ComputeType,
+                                                        int64_t>(
+            gating_output, topk_weights.mutable_data_ptr<float>(),
+            topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
+            num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
+            bias_ptr, true, input_ids.value().const_data_ptr<int64_t>(),
+            tid2eid.value().const_data_ptr<int64_t>(), stream);
+      } else {
+        STD_TORCH_CHECK(tid2eid.value().scalar_type() ==
+                        torch::headeronly::ScalarType::Int);
+        vllm::moe::topkGatingSoftplusSqrtKernelLauncher<OutIndType, ComputeType,
+                                                        int>(
+            gating_output, topk_weights.mutable_data_ptr<float>(),
+            topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
+            num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
+            bias_ptr, true, input_ids.value().const_data_ptr<int>(),
+            tid2eid.value().const_data_ptr<int>(), stream);
+      }
+    } else {
+      vllm::moe::topkGatingSoftplusSqrtKernelLauncher<OutIndType, ComputeType>(
+          gating_output, topk_weights.mutable_data_ptr<float>(),
+          topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
+          num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
+          bias_ptr, false, static_cast<const OutIndType*>(nullptr),
+          static_cast<const OutIndType*>(nullptr), stream);
+    }
+  };
+
+  if (topk_indices.scalar_type() == torch::headeronly::ScalarType::Int) {
+    launch(topk_indices.mutable_data_ptr<int>());
   } else if (topk_indices.scalar_type() ==
              torch::headeronly::ScalarType::UInt32) {
-    const uint32_t* input_ids_ptr = nullptr;
-    const uint32_t* tid2eid_ptr = nullptr;
-    if (tid2eid.has_value()) {
-      input_ids_ptr = input_ids.value().const_data_ptr<uint32_t>();
-      tid2eid_ptr = tid2eid.value().const_data_ptr<uint32_t>();
-    }
-    vllm::moe::topkGatingSoftplusSqrtKernelLauncher<uint32_t, ComputeType>(
-        gating_output, topk_weights.mutable_data_ptr<float>(),
-        topk_indices.mutable_data_ptr<uint32_t>(),
-        token_expert_indices.mutable_data_ptr<int>(), num_tokens, num_experts,
-        topk, renormalize, routed_scaling_factor, bias_ptr, use_hash,
-        input_ids_ptr, tid2eid_ptr, stream);
+    launch(topk_indices.mutable_data_ptr<uint32_t>());
   } else {
     STD_TORCH_CHECK(topk_indices.scalar_type() ==
                     torch::headeronly::ScalarType::Long);
-
-    const int64_t* input_ids_ptr = nullptr;
-    const int64_t* tid2eid_ptr = nullptr;
-    if (tid2eid.has_value()) {
-      input_ids_ptr = input_ids.value().const_data_ptr<int64_t>();
-      tid2eid_ptr = tid2eid.value().const_data_ptr<int64_t>();
-    }
-
-    vllm::moe::topkGatingSoftplusSqrtKernelLauncher<int64_t, ComputeType>(
-        gating_output, topk_weights.mutable_data_ptr<float>(),
-        topk_indices.mutable_data_ptr<int64_t>(),
-        token_expert_indices.mutable_data_ptr<int>(), num_tokens, num_experts,
-        topk, renormalize, routed_scaling_factor, bias_ptr, use_hash,
-        input_ids_ptr, tid2eid_ptr, stream);
+    launch(topk_indices.mutable_data_ptr<int64_t>());
   }
 }
 
@@ -738,4 +825,4 @@ void topk_softplus_sqrt(
     STD_TORCH_CHECK(false, "Unsupported gating_output data type: ",
                     gating_output.scalar_type());
   }
-}
+}

csrc/libtorch_stable/moe/torch_bindings.cpp

@@ -23,8 +23,13 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_moe_C, m) {
       "bias, Tensor? input_ids, Tensor? tid2eid) -> ()");
 
   // Calculate the result of moe by summing up the partial results
-  // from all selected experts.
-  m.def("moe_sum(Tensor input, Tensor! output) -> ()");
+  // from all selected experts. topk_ids/expert_map are optional and, when
+  // both given, enable pad-aware reduce that skips (token, expert)
+  // slots that were never actually computed (unrouted, or routed to an
+  // expert not owned by this rank under expert parallelism).
+  m.def(
+      "moe_sum(Tensor input, Tensor! output, Tensor? topk_ids=None, "
+      "Tensor? expert_map=None) -> ()");
 
   // Aligning the number of tokens to be processed by each expert such
   // that it is divisible by the block size.

csrc/libtorch_stable/ngram_embedding_kernels.cu

@@ -0,0 +1,96 @@
+// N-gram embedding index kernel for LongCat-Flash (n-gram embedding variant).
+//
+// Adapted from SGLang:
+// https://github.com/sgl-project/sglang/blob/main/python/sglang/jit_kernel/csrc/ngram_embedding.cuh
+//
+// For each position, computes the hashed n-gram embedding ids that index the
+// concatenated embedder table. Integer tensors are int32 except ``row_indices``
+// (int64); the token table is ``[max_running_reqs, max_context_len]`` int32,
+// where a negative entry marks an ignored token (e.g. an EOS boundary).
+
+#include "torch_utils.h"
+
+#include "ops.h"
+
+#include <cstdint>
+
+namespace vllm::ngram_embedding {
+
+constexpr int kBlockThreads = 256;
+
+__global__ void ComputeNGramIdsKernel(
+    int batch_size, int ne_n, int ne_k,
+    int* ne_weights,                       // [ne_n-1, ne_k, ne_n]
+    int* ne_mods,                          // [ne_n-1, ne_k]
+    int* exclusive_ne_embedder_size_sums,  // [(ne_n-1)*ne_k + 1]
+    int* exclusive_req_len_sums,           // [batch_size + 1]
+    int* ne_token_table,  // [max_running_reqs, max_context_len]
+    int max_context_len,
+    const int64_t* __restrict__ row_indices,  // [batch_size]
+    int* column_starts,                       // [batch_size]
+    int* n_gram_ids                           // [token_num, (ne_n-1)*ne_k]
+) {
+  const int req_id = blockIdx.x % batch_size;
+  const int config_id = (blockIdx.x - req_id) / batch_size;
+  // n and k are offset from their physical meaning: n = real_n - 2, k = real_k
+  // - 1 (they index into ne_weights / ne_mods).
+  const int k = config_id % ne_k;
+  const int n = (config_id - config_id % ne_k) / ne_k;
+  const int ne_weight_base_idx = n * ne_k * ne_n + k * ne_n;
+  const int ne_mod = ne_mods[n * ne_k + k];
+  for (int i = exclusive_req_len_sums[req_id] + threadIdx.x;
+       i < exclusive_req_len_sums[req_id + 1]; i += blockDim.x) {
+    uint64_t n_gram_id = 0;
+    const int64_t current_token_offset = i - exclusive_req_len_sums[req_id];
+    const int64_t req_token_table_index =
+        row_indices[req_id] * static_cast<int64_t>(max_context_len);
+    const int64_t current_token_table_index =
+        req_token_table_index + column_starts[req_id] + current_token_offset;
+    for (int j = 0; j < n + 2; j++) {
+      if (current_token_table_index - j < req_token_table_index) {
+        break;  // outside this request's range
+      }
+      if (ne_token_table[current_token_table_index - j] < 0) {
+        break;  // ignored token
+      }
+      const uint64_t term =
+          (uint64_t)ne_token_table[current_token_table_index - j] *
+          (uint64_t)ne_weights[ne_weight_base_idx + j];
+      n_gram_id += term % ne_mod;
+    }
+    n_gram_id %= ne_mod;
+    n_gram_id += exclusive_ne_embedder_size_sums[n * ne_k + k];
+    n_gram_ids[i * (ne_n - 1) * ne_k + n * ne_k + k] = (int)(n_gram_id);
+  }
+}
+
+}  // namespace vllm::ngram_embedding
+
+void ngram_compute_n_gram_ids(
+    int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
+    torch::stable::Tensor& ne_mods,
+    torch::stable::Tensor& exclusive_ne_embedder_size_sums,
+    torch::stable::Tensor& exclusive_req_len_sums,
+    torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
+    torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids) {
+  const int batch_size = static_cast<int>(exclusive_req_len_sums.size(0) - 1);
+  const int max_context_len = static_cast<int>(ne_token_table.size(1));
+  const int num_configs = (static_cast<int>(ne_n) - 1) * static_cast<int>(ne_k);
+  const int grid_size = num_configs * batch_size;
+  if (grid_size <= 0) return;
+
+  const torch::stable::accelerator::DeviceGuard device_guard(
+      ne_weights.get_device_index());
+  const cudaStream_t stream = get_current_cuda_stream();
+  vllm::ngram_embedding::ComputeNGramIdsKernel<<<
+      grid_size, vllm::ngram_embedding::kBlockThreads, 0, stream>>>(
+      batch_size, static_cast<int>(ne_n), static_cast<int>(ne_k),
+      ne_weights.mutable_data_ptr<int32_t>(),
+      ne_mods.mutable_data_ptr<int32_t>(),
+      exclusive_ne_embedder_size_sums.mutable_data_ptr<int32_t>(),
+      exclusive_req_len_sums.mutable_data_ptr<int32_t>(),
+      ne_token_table.mutable_data_ptr<int32_t>(), max_context_len,
+      row_indices.const_data_ptr<int64_t>(),
+      column_starts.mutable_data_ptr<int32_t>(),
+      n_gram_ids.mutable_data_ptr<int32_t>());
+}