Merge branch 'main' into wentao-fix-es-v2-bug

[Bugfix][Frontend] Flush engine reasoning parser at engine-reasoning → tool streaming boundary (#47606 )
Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com> Co-authored-by: Ben Browning <56071+bbrowning@users.noreply.github.com>
2026-07-13 14:19:23 -04:00 · 2026-07-13 14:06:10 -04:00 · 2026-07-13 17:43:59 +00:00 · 2026-07-13 17:28:50 +00:00 · 2026-07-13 20:09:41 +03:00 · 2026-07-13 12:57:36 -04:00
483 changed files with 12263 additions and 1869 deletions
@@ -141,9 +141,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
@@ -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:
@@ -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
@@ -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
@@ -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
@@ -152,7 +152,7 @@ steps:

 - label: Regression
  key: regression
-  timeout_in_minutes: 30
+  timeout_in_minutes: 50
  device: intel_gpu
  agent_tags:
    label: production
@@ -186,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:
@@ -222,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
@@ -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
@@ -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
@@ -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
@@ -104,7 +104,7 @@ steps:
  - label: "XPU server test"
    depends_on:
      - image-build-xpu
-    timeout_in_minutes: 30
+    timeout_in_minutes: 45
    device: intel_gpu
    agent_tags:
      label: production
@@ -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

@@ -764,7 +764,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 +1211,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 +1272,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 +1379,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 +1428,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 +1572,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 +1580,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 +1608,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
@@ -1920,6 +2082,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
@@ -21,16 +21,20 @@ 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}"
+
 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 +49,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 +80,48 @@ 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"
-  install_cargo_binstall
-  cargo binstall --no-confirm --secure cargo-nextest
+  log_section "Installing cargo-nextest ${CARGO_NEXTEST_VERSION_REQ}"
+  cargo binstall \
+    --no-confirm \
+    --force \
+    --secure \
+    "cargo-nextest@${CARGO_NEXTEST_VERSION_REQ}"
 }

 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 +143,7 @@ PY
 }

 run_style_clippy() {
+  install_cargo_binstall
  install_cargo_sort
  install_cargo_deny

@@ -186,6 +174,7 @@ run_style_clippy() {
 run_tests() {
  install_uv
  setup_pyo3_python
+  install_cargo_binstall
  install_cargo_nextest

  log_section "Running cargo nextest"
@@ -12,7 +12,8 @@ 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
@@ -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
@@ -23,7 +23,8 @@ steps:
  - 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
@@ -271,7 +272,8 @@ steps:
  - 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)
@@ -89,6 +89,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 +102,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
@@ -27,7 +27,7 @@ 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
@@ -69,7 +69,7 @@ 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
@@ -80,7 +80,8 @@ steps:
  - 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
@@ -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
@@ -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)
@@ -19,7 +19,7 @@ else()
  FetchContent_Declare(
        flashmla
        GIT_REPOSITORY https://github.com/vllm-project/FlashMLA
-        GIT_TAG a6ec2ba7bd0a7dff98b3f4d3e6b52b159c48d78b
+        GIT_TAG b70aff3d110a2b1a037e62eac295166b5143643a
        GIT_PROGRESS TRUE
        CONFIGURE_COMMAND ""
        BUILD_COMMAND ""
@@ -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 bb9a72e7dde0dc614ffc663e052cd6a19ce73a42
          GIT_PROGRESS TRUE
          # Don't share the vllm-flash-attn build between build types
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
@@ -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]
@@ -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);
@@ -508,6 +516,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)");
@@ -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);
+}
@@ -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
@@ -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(
@@ -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,
@@ -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;
+      }
    }
  }
 }
@@ -313,7 +313,7 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
    std::optional<torch::stable::Tensor> index_cache, int64_t block_size,
    std::optional<torch::stable::Tensor> q_out,
    std::optional<torch::stable::Tensor> index_q_out,
-    const std::string& kv_cache_dtype);
+    const std::string& kv_cache_dtype, bool skip_index_branch);

 // Sampler kernels (shared CUDA/ROCm)
 void apply_repetition_penalties_(
@@ -414,6 +414,8 @@ void gelu_new(torch::stable::Tensor& out, torch::stable::Tensor& input);
 void gelu_fast(torch::stable::Tensor& out, torch::stable::Tensor& input);
 void gelu_quick(torch::stable::Tensor& out, torch::stable::Tensor& input);

+void relu_squared(torch::stable::Tensor& out, torch::stable::Tensor& input);
+
 // INT8 quantization kernels (shared CUDA/ROCm)
 void static_scaled_int8_quant(torch::stable::Tensor& out,
                              torch::stable::Tensor const& input,
@@ -466,7 +466,7 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
      "Tensor? slot_mapping, Tensor? index_slot_mapping, "
      "Tensor!? kv_cache, Tensor!? index_cache, "
      "int block_size, Tensor!? q_out, Tensor!? index_q_out, "
-      "str kv_cache_dtype) -> ()");
+      "str kv_cache_dtype, bool skip_index_branch=False) -> ()");

  // Apply repetition penalties to logits in-place.
  ops.def(
@@ -539,6 +539,9 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
  // Quick GELU implementation.
  ops.def("gelu_quick(Tensor! out, Tensor input) -> ()");

+  // relu(x)^2 activation from https://arxiv.org/abs/2109.08668v2
+  ops.def("relu_squared(Tensor! out, Tensor input) -> ()");
+
  // Compute int8 quantized tensor for given scaling factor.
  ops.def(
      "static_scaled_int8_quant(Tensor! result, Tensor input, Tensor scale,"
@@ -715,6 +718,7 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
  ops.impl("gelu_new", TORCH_BOX(&gelu_new));
  ops.impl("gelu_fast", TORCH_BOX(&gelu_fast));
  ops.impl("gelu_quick", TORCH_BOX(&gelu_quick));
+  ops.impl("relu_squared", TORCH_BOX(&relu_squared));
  ops.impl("silu_and_mul_with_clamp", TORCH_BOX(&silu_and_mul_clamp));

  // INT8 quantization kernels
@@ -43,6 +43,8 @@ void gelu_fast(torch::Tensor& out, torch::Tensor& input);

 void gelu_quick(torch::Tensor& out, torch::Tensor& input);

+void relu_squared(torch::Tensor& out, torch::Tensor& input);
+
 void static_scaled_int8_quant(torch::Tensor& out, torch::Tensor const& input,
                              torch::Tensor const& scale,
                              std::optional<torch::Tensor> const& azp);
@@ -1,5 +1,7 @@
 # default base image
 ARG REMOTE_VLLM="0"
+ARG VLLM_REPO="https://github.com/vllm-project/vllm.git"
+ARG VLLM_BRANCH="main"
 ARG COMMON_WORKDIR=/app
 ARG BASE_IMAGE=rocm/vllm-dev:base
 ARG CI_BASE_IMAGE=rocm/vllm-dev:ci_base
@@ -110,10 +112,8 @@ WORKDIR ${COMMON_WORKDIR}
 FROM base AS fetch_vllm_0
 ONBUILD COPY ./ vllm/
 FROM base AS fetch_vllm_1
-ARG VLLM_REPO="https://github.com/vllm-project/vllm.git"
-ARG VLLM_BRANCH="main"
-ENV VLLM_REPO=${VLLM_REPO}
-ENV VLLM_BRANCH=${VLLM_BRANCH}
+ARG VLLM_REPO
+ARG VLLM_BRANCH
 ONBUILD RUN git clone ${VLLM_REPO} \
 	    && cd vllm \
 	    && git fetch -v --prune -- origin ${VLLM_BRANCH} \
@@ -124,34 +124,118 @@ ONBUILD RUN git clone ${VLLM_REPO} \
 FROM fetch_vllm_${REMOTE_VLLM} AS fetch_vllm

 # -----------------------
-# Rust build stage
-# Builds the `vllm-rs` frontend in a dedicated stage so the wheel build stages
-# don't need the rust toolchain or protoc.
-FROM fetch_vllm AS rust-build
-ARG COMMON_WORKDIR
-ARG USE_SCCACHE
+# Rust/protoc toolchain inputs
+#
+# Keep these separate from Rust source inputs so Rust source-only changes do not
+# invalidate the long-lived ci_base image that carries the toolchain.
+FROM scratch AS rust_toolchain_input_0
+COPY rust-toolchain.toml /rust-toolchain-input/vllm/rust-toolchain.toml
+COPY tools/install_protoc.sh /rust-toolchain-input/vllm/tools/install_protoc.sh
+
+FROM base AS rust_toolchain_input_1
+ARG VLLM_REPO
+ARG VLLM_BRANCH
+RUN git clone --no-checkout --filter=blob:none ${VLLM_REPO} /rust-toolchain-input/vllm \
+    && cd /rust-toolchain-input/vllm \
+    && git fetch -v --prune -- origin ${VLLM_BRANCH} \
+    && git sparse-checkout init --no-cone \
+    && git sparse-checkout set \
+        rust-toolchain.toml \
+        tools/install_protoc.sh \
+    && git checkout FETCH_HEAD
+
+FROM rust_toolchain_input_${REMOTE_VLLM} AS rust-toolchain-input
+
+# -----------------------
+# Rust build inputs
+#
+# REMOTE_VLLM=0: use only the local files that can affect Rust artifacts.
+# REMOTE_VLLM=1: fetch only those same files from the requested remote ref.
+FROM scratch AS rust_input_0
+COPY requirements/build/rust.txt /rust-input/vllm/requirements/build/rust.txt
+COPY rust/Cargo.lock rust/Cargo.toml /rust-input/vllm/rust/
+COPY rust/proto /rust-input/vllm/rust/proto
+COPY rust/src /rust-input/vllm/rust/src
+COPY rust-toolchain.toml /rust-input/vllm/rust-toolchain.toml
+COPY tools/build_rust.py /rust-input/vllm/tools/build_rust.py
+COPY build_rust.sh /rust-input/vllm/build_rust.sh
+
+FROM base AS rust_input_1
+ARG VLLM_REPO
+ARG VLLM_BRANCH
+RUN git clone --no-checkout --filter=blob:none ${VLLM_REPO} /rust-input/vllm \
+    && cd /rust-input/vllm \
+    && git fetch -v --prune -- origin ${VLLM_BRANCH} \
+    && git sparse-checkout init --no-cone \
+    && git sparse-checkout set \
+        requirements/build/rust.txt \
+        rust/Cargo.lock \
+        rust/Cargo.toml \
+        rust/proto/** \
+        rust/src/** \
+        rust-toolchain.toml \
+        tools/build_rust.py \
+        build_rust.sh \
+    && git checkout FETCH_HEAD
+
+FROM rust_input_${REMOTE_VLLM} AS rust-input
+
+# -----------------------
+# Rust/protoc toolchain
+FROM base AS rust-toolchain
+
+ENV CARGO_HOME=/root/.cargo
+ENV RUSTUP_HOME=/root/.rustup
+ENV PATH=${CARGO_HOME}/bin:${PATH}

 # protoc is used by tonic-build/prost-build.
 RUN apt-get update -q -y && apt-get install -q -y --no-install-recommends \
        ca-certificates curl unzip \
    && rm -rf /var/lib/apt/lists/*

-COPY tools/install_protoc.sh /tmp/install_protoc.sh
+COPY --from=rust-toolchain-input /rust-toolchain-input/vllm/tools/install_protoc.sh /tmp/install_protoc.sh
 RUN /tmp/install_protoc.sh && rm /tmp/install_protoc.sh

+COPY --from=rust-toolchain-input /rust-toolchain-input/vllm/rust-toolchain.toml /tmp/rust-toolchain.toml
+RUN TOOLCHAIN="$(grep '^channel' /tmp/rust-toolchain.toml | sed 's/.*= *"\(.*\)"/\1/')" \
+    && if ! command -v rustup >/dev/null 2>&1; then \
+        curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs \
+            | sh -s -- -y --default-toolchain none; \
+    fi \
+    && rustup toolchain install "${TOOLCHAIN}" \
+    && rustup default "${TOOLCHAIN}" \
+    && rustc --version \
+    && cargo --version
+
+# -----------------------
+# Rust build stage
+# Builds the `vllm-rs` frontend in a dedicated stage so the wheel build stages
+# don't need the rust toolchain or protoc.
+FROM rust-toolchain AS rust-build
+ARG COMMON_WORKDIR
+ARG USE_SCCACHE
+
 # Cap cargo parallelism to avoid exhausting the AMD CI host's open-file limit
 # (rustc spawns enough concurrent processes to hit RLIMIT_NOFILE otherwise).
 ENV CARGO_BUILD_JOBS=4
 ENV CARGO_NET_RETRY=10
 ENV RUSTUP_MAX_RETRIES=10

+WORKDIR ${COMMON_WORKDIR}/vllm
+
+COPY --from=rust-input /rust-input/vllm/requirements/build/rust.txt requirements/build/rust.txt
+
 # BuildKit can run this stage in parallel with ROCm native builds. Keep Rust on
 # a separate local sccache daemon while sharing the same remote cache backend.
 ENV SCCACHE_SERVER_PORT=4227

 RUN --mount=type=cache,id=vllm-rocm-uv,target=/root/.cache/uv \
-    cd ${COMMON_WORKDIR}/vllm \
-    && uv pip install --system -r requirements/build/rust.txt
+    uv pip install --system -r requirements/build/rust.txt
+
+COPY --from=rust-input /rust-input/vllm/rust rust
+COPY --from=rust-input /rust-input/vllm/rust-toolchain.toml rust-toolchain.toml
+COPY --from=rust-input /rust-input/vllm/tools/build_rust.py tools/build_rust.py
+COPY --from=rust-input /rust-input/vllm/build_rust.sh build_rust.sh

 # Build the release binary. Cargo's registry/git caches can be written by
 # concurrent BuildKit jobs on shared workers, so lock those cache mounts while
@@ -159,8 +243,7 @@ RUN --mount=type=cache,id=vllm-rocm-uv,target=/root/.cache/uv \
 # can outlive source updates across BuildKit cache reuse.
 RUN --mount=type=cache,id=vllm-rocm-cargo-registry,target=/root/.cargo/registry,sharing=locked \
    --mount=type=cache,id=vllm-rocm-cargo-git,target=/root/.cargo/git,sharing=locked \
-    cd ${COMMON_WORKDIR}/vllm \
-    && if [ "$USE_SCCACHE" = "1" ]; then \
+    if [ "$USE_SCCACHE" = "1" ]; then \
        export RUSTC_WRAPPER=sccache \
        && sccache --show-stats; \
    fi \
@@ -185,8 +268,8 @@ COPY requirements/common.txt requirements/common.txt
 RUN --mount=type=cache,id=vllm-rocm-uv,target=/root/.cache/uv \
    uv pip install --system -r requirements/rocm.txt

-# pyproject.toml is bind-mounted in the RUN step so metadata-only changes do
-# not invalidate the expensive native build layer.
+# pyproject.toml is bind-mounted so it is not persisted in the csrc-build image,
+# but it still participates in the BuildKit cache key for this RUN step.
 COPY setup.py CMakeLists.txt ./
 COPY tools/build_rust.py tools/build_rust.py
 COPY cmake cmake/
@@ -556,6 +639,17 @@ COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1
 FROM mori_base AS ci_base
 ARG COMMON_WORKDIR

+ENV CARGO_HOME=/root/.cargo
+ENV RUSTUP_HOME=/root/.rustup
+ENV PATH=${CARGO_HOME}/bin:${PATH}
+
+# Bake the Rust/protoc toolchain into ci_base so Rust build/test steps do not
+# need to fetch it on each per-commit image build.
+COPY --from=rust-toolchain /root/.cargo /root/.cargo
+COPY --from=rust-toolchain /root/.rustup /root/.rustup
+COPY --from=rust-toolchain /usr/local/bin/protoc /usr/local/bin/protoc
+COPY --from=rust-toolchain /usr/local/include/google /usr/local/include/google
+
 # Update rdma-core to support latest rocshmem.
 ARG DEEPEP_NIC
 RUN if [ "${DEEPEP_NIC}" = "cx7" ] || [ "${DEEPEP_NIC}" = "io" ]; then \
@@ -37,37 +37,6 @@ RUN microdnf install -y \
    python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip python${PYTHON_VERSION}-wheel  && \
    python${PYTHON_VERSION} -m venv $VIRTUAL_ENV && pip install --no-cache -U pip wheel uv && microdnf clean all

-FROM python-install AS pyarrow
-
-# Build Apache Arrow
-WORKDIR /tmp
-RUN --mount=type=cache,target=/root/.cache/uv \
-    git clone https://github.com/apache/arrow.git  -b maint-19.0.1 && \
-    cd arrow/cpp && \
-    mkdir release && cd release && \
-    cmake -DCMAKE_BUILD_TYPE=Release \
-          -DCMAKE_INSTALL_PREFIX=/usr/local \
-          -DARROW_PYTHON=ON \
-          -DARROW_PARQUET=ON \
-          -DARROW_ORC=ON \
-          -DARROW_FILESYSTEM=ON \
-          -DARROW_WITH_LZ4=ON \
-          -DARROW_WITH_ZSTD=ON \
-          -DARROW_WITH_SNAPPY=ON \
-          -DARROW_JSON=ON \
-          -DARROW_CSV=ON \
-          -DARROW_DATASET=ON \
-          -DPROTOBUF_PROTOC_EXECUTABLE=/usr/bin/protoc \
-          -DARROW_DEPENDENCY_SOURCE=BUNDLED \
-          .. && \
-    make -j$(nproc) && \
-    make install && \
-    cd ../../python && \
-    export PYARROW_PARALLEL=4 && \
-    export ARROW_BUILD_TYPE=release && \
-    uv pip install -r requirements-build.txt && \
-    python setup.py build_ext --build-type=$ARROW_BUILD_TYPE --bundle-arrow-cpp bdist_wheel
-
 FROM python-install AS rust
 ENV CARGO_HOME=/root/.cargo
 ENV RUSTUP_HOME=/root/.rustup
@@ -121,60 +90,18 @@ RUN --mount=type=cache,target=/root/.cache/uv \
 FROM python-install AS numba-builder

 ARG MAX_JOBS
-ARG NUMBA_VERSION=0.61.2
+ARG NUMBA_VERSION=0.65.0

 WORKDIR /tmp

 # Clone all required dependencies
 RUN --mount=type=cache,target=/root/.cache/uv \
    microdnf install ninja-build gcc gcc-c++ -y && \
-    git clone --recursive https://github.com/llvm/llvm-project.git -b llvmorg-15.0.7  && \
-    git clone --recursive https://github.com/numba/llvmlite.git -b v0.44.0 && \
+    git clone --recursive https://github.com/numba/llvmlite.git -b v0.47.0 && \
    git clone --recursive https://github.com/numba/numba.git -b ${NUMBA_VERSION} && \
-    cd llvm-project && mkdir build && cd  build && \
+    cd llvmlite && \
    uv pip install 'cmake<4' 'setuptools<70' numpy && \
-    export PREFIX=/usr/local && CMAKE_ARGS="${CMAKE_ARGS} -DLLVM_ENABLE_PROJECTS=lld;libunwind;compiler-rt" \
-    CFLAGS="$(echo $CFLAGS | sed 's/-fno-plt //g')" \
-    CXXFLAGS="$(echo $CXXFLAGS | sed 's/-fno-plt //g')" \
-    CMAKE_ARGS="${CMAKE_ARGS} -DFFI_INCLUDE_DIR=$PREFIX/include" \
-    CMAKE_ARGS="${CMAKE_ARGS} -DFFI_LIBRARY_DIR=$PREFIX/lib" \
-    cmake -DCMAKE_INSTALL_PREFIX="${PREFIX}"               \
-        -DCMAKE_BUILD_TYPE=Release                       \
-        -DCMAKE_LIBRARY_PATH="${PREFIX}"                 \
-        -DLLVM_ENABLE_LIBEDIT=OFF                        \
-        -DLLVM_ENABLE_LIBXML2=OFF                        \
-        -DLLVM_ENABLE_RTTI=ON                            \
-        -DLLVM_ENABLE_TERMINFO=OFF                       \
-        -DLLVM_INCLUDE_BENCHMARKS=OFF                    \
-        -DLLVM_INCLUDE_DOCS=OFF                          \
-        -DLLVM_INCLUDE_EXAMPLES=OFF                      \
-        -DLLVM_INCLUDE_GO_TESTS=OFF                      \
-        -DLLVM_INCLUDE_TESTS=OFF                         \
-        -DLLVM_INCLUDE_UTILS=ON                          \
-        -DLLVM_INSTALL_UTILS=ON                          \
-        -DLLVM_UTILS_INSTALL_DIR=libexec/llvm            \
-        -DLLVM_BUILD_LLVM_DYLIB=OFF                      \
-        -DLLVM_LINK_LLVM_DYLIB=OFF                       \
-        -DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD=WebAssembly \
-        -DLLVM_ENABLE_FFI=ON                             \
-        -DLLVM_ENABLE_Z3_SOLVER=OFF                      \
-        -DLLVM_OPTIMIZED_TABLEGEN=ON                     \
-        -DCMAKE_POLICY_DEFAULT_CMP0111=NEW               \
-        -DCOMPILER_RT_BUILD_BUILTINS=ON                  \
-        -DCOMPILER_RT_BUILTINS_HIDE_SYMBOLS=OFF          \
-        -DCOMPILER_RT_BUILD_LIBFUZZER=OFF                \
-        -DCOMPILER_RT_BUILD_CRT=OFF                      \
-        -DCOMPILER_RT_BUILD_MEMPROF=OFF                  \
-        -DCOMPILER_RT_BUILD_PROFILE=OFF                  \
-        -DCOMPILER_RT_BUILD_SANITIZERS=OFF               \
-        -DCOMPILER_RT_BUILD_XRAY=OFF                     \
-        -DCOMPILER_RT_BUILD_GWP_ASAN=OFF                 \
-        -DCOMPILER_RT_BUILD_ORC=OFF                      \
-        -DCOMPILER_RT_INCLUDE_TESTS=OFF                  \
-        ${CMAKE_ARGS} -GNinja ../llvm                    \
-    && ninja install  . && \
-    #  build llvmlite
-    cd ../../llvmlite && python setup.py bdist_wheel && \
+    python setup.py bdist_wheel && \
    cd ../numba && \
    if ! grep '#include "dynamic_annotations.h"' numba/_dispatcher.cpp; then \
       sed -i '/#include "internal\/pycore_atomic.h"/i\#include "dynamic_annotations.h"' numba/_dispatcher.cpp; \
@@ -192,21 +119,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
    cd opencv-python && \
    python -m build --wheel --installer=uv --outdir /tmp/opencv-python/dist

-## Todo(r3hankhan123): Remove guidance-builder stage once vLLM upgrades to new version of llguidance that fixes s390x issues. See https://github.com/guidance-ai/llguidance/issues/330
-FROM python-install AS guidance-builder
-WORKDIR /tmp
-ENV CARGO_HOME=/root/.cargo
-ENV RUSTUP_HOME=/root/.rustup
-ENV PATH="$CARGO_HOME/bin:$RUSTUP_HOME/bin:$PATH"
-RUN --mount=type=cache,target=/root/.cache/uv \
-    --mount=type=bind,from=rust,source=/root/.cargo,target=/root/.cargo,rw \
-    --mount=type=bind,from=rust,source=/root/.rustup,target=/root/.rustup,rw \
-    git clone https://github.com/guidance-ai/llguidance.git && \
-    cd llguidance && \
-    git checkout s390x-fix-v2  && \
-    uv pip install maturin && \
-    python -m maturin build --release --out dist --compatibility linux
-
 # # Final build stage
 FROM python-install AS vllm-cpu
 ARG PYTHON_VERSION
@@ -235,20 +147,16 @@ RUN --mount=type=bind,from=numa-build,src=/tmp/numactl-2.0.19,target=/numactl \
 RUN --mount=type=cache,target=/root/.cache/uv \
    --mount=type=bind,from=rust,source=/root/.cargo,target=/root/.cargo,rw \
    --mount=type=bind,from=rust,source=/root/.rustup,target=/root/.rustup,rw \
-    --mount=type=bind,from=pyarrow,source=/tmp/arrow/python/dist,target=/tmp/arrow-wheels \
    --mount=type=bind,from=torch-vision,source=/tmp/vision/dist,target=/tmp/vision-wheels/ \
    --mount=type=bind,from=hf-xet-builder,source=/tmp/hf-xet/dist,target=/tmp/hf-xet-wheels/ \
    --mount=type=bind,from=numba-builder,source=/tmp/llvmlite/dist,target=/tmp/llvmlite-wheels/ \
    --mount=type=bind,from=numba-builder,source=/tmp/numba/dist,target=/tmp/numba-wheels/ \
    --mount=type=bind,from=opencv-builder,source=/tmp/opencv-python/dist,target=/tmp/opencv-wheels/ \
-    --mount=type=bind,from=guidance-builder,source=/tmp/llguidance/dist,target=/tmp/guidance-wheels/ \
-     ARROW_WHL_FILE=$(ls /tmp/arrow-wheels/*.whl) && \
     VISION_WHL_FILE=$(ls /tmp/vision-wheels/*.whl) && \
     HF_XET_WHL_FILE=$(ls /tmp/hf-xet-wheels/*.whl) && \
     LLVM_WHL_FILE=$(ls /tmp/llvmlite-wheels/*.whl) && \
     NUMBA_WHL_FILE=$(ls /tmp/numba-wheels/*.whl) && \
     OPENCV_WHL_FILE=$(ls /tmp/opencv-wheels/*.whl) && \
-     GUIDANCE_WHL_FILE=$(ls /tmp/guidance-wheels/*.whl) && \
     uv pip install -v \
        $ARROW_WHL_FILE \
        $VISION_WHL_FILE \
@@ -256,7 +164,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
        $LLVM_WHL_FILE \
        $NUMBA_WHL_FILE \
        $OPENCV_WHL_FILE \
-        $GUIDANCE_WHL_FILE \
        --torch-backend cpu \
        --index-strategy unsafe-best-match \
        -r requirements/build/cpu.txt \
@@ -141,6 +141,10 @@ variable "ROCM_CSRC_CACHE_TO_MODE" {
  default = "max"
 }

+variable "ROCM_RUST_CACHE_TO_MODE" {
+  default = "max"
+}
+
 variable "ROCM_FINAL_CACHE_TO_MODE" {
  default = "min"
 }
@@ -164,6 +168,13 @@ function "get_cache_from_rocm" {
    VLLM_MERGE_BASE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:csrc-rocm-${VLLM_MERGE_BASE_COMMIT}" : "",
    ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
    ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:csrc-rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
+    # Import the source-scoped Rust frontend cache so non-Rust changes do not
+    # force a fresh cargo release build.
+    BUILDKITE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${BUILDKITE_COMMIT}" : "",
+    PARENT_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${PARENT_COMMIT}" : "",
+    VLLM_MERGE_BASE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${VLLM_MERGE_BASE_COMMIT}" : "",
+    ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
+    ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
    # Branch-scoped full image cache - fallback when parent-commit cache is evicted
    ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
    ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
@@ -204,6 +215,27 @@ function "get_cache_to_rocm_csrc" {
  ])
 }

+function "get_cache_from_rocm_rust" {
+  params = []
+  result = compact([
+    BUILDKITE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${BUILDKITE_COMMIT}" : "",
+    PARENT_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${PARENT_COMMIT}" : "",
+    VLLM_MERGE_BASE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${VLLM_MERGE_BASE_COMMIT}" : "",
+    ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
+    ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
+  ])
+}
+
+function "get_cache_to_rocm_rust" {
+  params = []
+  result = compact([
+    # Export exact-commit and branch-scoped Rust caches. A content-addressed
+    # cache ref is appended by ci-bake-rocm.sh when that wrapper is used.
+    BUILDKITE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${BUILDKITE_COMMIT},mode=${ROCM_RUST_CACHE_TO_MODE}" : "",
+    ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${ROCM_RUST_CACHE_TO_MODE}" : "",
+  ])
+}
+
 # Cache functions for upstream dependency stages (RIXL/UCX, ROCShmem, DeepEP).
 # These stages are pinned to specific upstream commit hashes, so cache keys use
 # those hashes rather than the Buildkite commit. This means the cache persists
@@ -276,6 +308,17 @@ target "csrc-rocm-ci" {
  output     = ["type=cacheonly"]
 }

+# Cache-only target for the Rust frontend build stage. Final-image cache
+# exports use mode=min and do not reliably persist intermediate cargo layers,
+# so Rust gets its own source-scoped cache target.
+target "rust-rocm-ci" {
+  inherits   = ["_common-rocm", "_ci-rocm"]
+  target     = "rust-build"
+  cache-from = get_cache_from_rocm_rust()
+  cache-to   = get_cache_to_rocm_rust()
+  output     = ["type=cacheonly"]
+}
+
 # Keep wheel export on the same CI graph as the test image build so the
 # shared build_vllm/export_vllm stages resolve identically within one bake
 # invocation. Without this, export-wheel-rocm uses the plain local target
@@ -292,13 +335,13 @@ target "export-wheel-rocm" {
 # Artifact-only vLLM build. GPU test jobs consume this artifact on top of
 # ci_base, avoiding a per-commit multi-GB image push/pull.
 group "test-rocm-ci-with-artifacts" {
-  targets = ["csrc-rocm-ci", "export-wheel-rocm"]
+  targets = ["rust-rocm-ci", "csrc-rocm-ci", "export-wheel-rocm"]
 }

 # Full test image + wheel export. Kept for fallback/debugging when a pushed
 # per-commit image is useful.
 group "test-rocm-ci-with-wheel" {
-  targets = ["csrc-rocm-ci", "test-rocm-ci", "export-wheel-rocm"]
+  targets = ["rust-rocm-ci", "csrc-rocm-ci", "test-rocm-ci", "export-wheel-rocm"]
 }

 # Image tags for the ci_base build. ci-bake-rocm.sh rewrites CI_BASE_IMAGE_TAG
@@ -358,6 +358,7 @@ th {
 | `DeepseekForCausalLM` | DeepSeek | `deepseek-ai/deepseek-llm-67b-base`, `deepseek-ai/deepseek-llm-7b-chat`, etc. | ✅︎ | ✅︎ |
 | `DeepseekV2ForCausalLM` | DeepSeek-V2 | `deepseek-ai/DeepSeek-V2`, `deepseek-ai/DeepSeek-V2-Chat`, etc. | ✅︎ | ✅︎ |
 | `DeepseekV3ForCausalLM` | DeepSeek-V3 | `deepseek-ai/DeepSeek-V3`, `deepseek-ai/DeepSeek-R1`, `deepseek-ai/DeepSeek-V3.1`, etc. | ✅︎ | ✅︎ |
+| `DeepseekV32ForCausalLM` | DeepSeek-V3.2 | `deepseek-ai/DeepSeek-V3.2`, etc. | ✅︎ | ✅︎ |
 | `DeepseekV4ForCausalLM` | DeepSeek-V4 | `deepseek-ai/DeepSeek-V4-Flash`, `deepseek-ai/DeepSeek-V4-Pro`, etc. | | ✅︎ |
 | `DotsOCRForCausalLM` | dots_ocr | `rednote-hilab/dots.ocr` | ✅︎ | ✅︎ |
 | `Ernie4_5ForCausalLM` | Ernie4.5 | `baidu/ERNIE-4.5-0.3B-PT`, etc. | ✅︎ | ✅︎ |
@@ -170,6 +170,7 @@ For further details on Weight Transfer, please refer to [this page](../../traini
 - `/pause` - Pause generation (causes denial of service)
 - `/resume` - Resume generation
 - `/is_paused` - Check if generation is paused
+- `/abort_requests` - Abort in-flight requests (all in-flight, or the given `request_ids`) without pausing the scheduler
 - `/init_weight_transfer_engine` - Initialize weight transfer engine for RLHF
 - `/start_weight_update` - Prepares the inference engine for a weight update.
 - `/update_weights` - Update model weights (can alter model behavior)
@@ -42,6 +42,7 @@ When using the vLLM HTTP server, the same functionality is available via:

 - `POST /pause?mode=keep` - Pause generation
 - `POST /resume` - Resume generation
+- `POST /abort_requests` - Abort in-flight requests without pausing the scheduler (send `{}` to abort all, or `{"request_ids": [...]}`)

 !!! note "Data Parallelism"
    When using data parallelism with vLLM's **internal load balancer** (i.e. `data_parallel_backend="ray"`), pause and resume are handled automatically across all DP ranks -- a single call is sufficient. When using an **external load balancer** (i.e. multiple independent vLLM instances behind a proxy), you must send pause and resume requests to **every** engine instance individually before and after the weight update.
@@ -49,6 +49,10 @@ update_request = WeightTransferUpdateRequest(
 )
 ```

+At the LLM/API layer, call `start_draft_weight_update()` instead of
+`start_weight_update()` to target the speculative draft model;
+`update_weights` / `finish_weight_update` are unchanged.
+
 ### WeightTransferUpdateInfo

 The base `WeightTransferUpdateInfo` is a marker class for backend-specific update info:
@@ -191,6 +191,7 @@ The following endpoints **do not require authentication** even when `--api-key`
 - `/pause` - Pause generation (causes denial of service)
 - `/resume` - Resume generation
 - `/is_paused` - Check if generation is paused
+- `/abort_requests` - Abort in-flight requests (causes loss of in-flight work)
 - `/scale_elastic_ep` - Trigger scaling operations
 - `/is_scaling_elastic_ep` - Check if scaling is in progress
 - `/init_weight_transfer_engine` - Initialize weight transfer engine for RLHF
@@ -21,7 +21,7 @@ pillow  # Required for image processing
 prometheus-fastapi-instrumentator >= 8.0.0 # v8 unblocks starlette >= 1.0
 tiktoken >= 0.6.0  # Required for DBRX tokenizer
 lm-format-enforcer == 0.11.3
-llguidance >= 1.7.0, < 1.8.0; platform_machine == "x86_64" or platform_machine == "arm64" or platform_machine == "aarch64" or platform_machine == "ppc64le"
+llguidance >= 1.7.0, < 1.8.0; platform_machine == "x86_64" or platform_machine == "arm64" or platform_machine == "aarch64" or platform_machine == "ppc64le" or platform_machine == "s390x"
 outlines_core == 0.2.14
 # required for outlines backend disk cache
 diskcache == 5.6.3
@@ -962,7 +962,7 @@ s3transfer==0.10.3
    # via boto3
 sacrebleu==2.4.3
    # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
    # via
    #   -r requirements/test/../common.txt
    #   accelerate
@@ -1159,7 +1159,7 @@ tqdm==4.67.3
    #   segmentation-models-pytorch
    #   sentence-transformers
    #   transformers
-transformers==5.10.4
+transformers==5.13.1
    # via
    #   -r requirements/test/../common.txt
    #   -r requirements/test/cuda.in
@@ -39,7 +39,7 @@ open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.10.4
+transformers==5.13.1
 tokenizers==0.22.2
 schemathesis>=4.0.0 # Required for openai schema test.
 # quantization
@@ -1053,7 +1053,7 @@ s3transfer==0.10.3
    # via boto3
 sacrebleu==2.4.3
    # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
@@ -1261,7 +1261,7 @@ tqdm==4.67.3
    #   segmentation-models-pytorch
    #   sentence-transformers
    #   transformers
-transformers==5.10.4
+transformers==5.13.1
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
@@ -29,7 +29,7 @@ opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.10.4
+transformers==5.13.1
 tokenizers==0.22.2
 schemathesis>=4.0.0 # Required for openai schema test.
 # quantization
@@ -35,7 +35,7 @@ open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.10.4
+transformers==5.13.1
 tokenizers==0.22.2
 schemathesis>=4.0.0 # Required for openai schema test
 # quantization
@@ -1035,7 +1035,7 @@ s3transfer==0.16.0
    # via boto3
 sacrebleu==2.6.0
    # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
@@ -1218,7 +1218,7 @@ tqdm==4.67.3
    #   sentence-transformers
    #   tilelang
    #   transformers
-transformers==5.10.4
+transformers==5.13.1
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
@@ -17,7 +17,7 @@ accelerate
 arctic-inference
 lm_eval[api]>=0.4.12
 modelscope<1.38
-transformers==5.10.4
+transformers==5.13.1

 # --- Audio Processing ---
 librosa
@@ -779,7 +779,7 @@ rpds-py==0.30.0
    #   referencing
 sacrebleu==2.6.0
    # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
@@ -940,7 +940,7 @@ tqdm==4.67.3
    #   pqdm
    #   sentence-transformers
    #   transformers
-transformers==5.10.4
+transformers==5.13.1
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
@@ -18,4 +18,4 @@ torchvision
 torchcodec >= 0.14 # Required for the torchcodec video decoding backend

 auto_round_lib>=0.14.0
-vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.10.1/vllm_xpu_kernels-0.1.10.1-cp38-abi3-manylinux_2_28_x86_64.whl
+vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.11/vllm_xpu_kernels-0.1.11-cp38-abi3-manylinux_2_28_x86_64.whl
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 syntax = "proto3";
 package vllm;

@@ -107,6 +110,9 @@ message KVCacheParameters {

  // KV Connector transfer parameters
  google.protobuf.Struct kv_transfer_params = 3;
+
+  // Encoder cache connector transfer parameters
+  google.protobuf.Struct ec_transfer_params = 4;
 }

 // Controls which extra candidate tokens at each position should be returned
@@ -173,6 +179,7 @@ message FinishInfo {

  google.protobuf.Struct kv_transfer_params = 6;
  //uint64 seed = 7;
+  google.protobuf.Struct ec_transfer_params = 8;
 }

 // Info for candidate tokens other than the input/sampled
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::time::Duration;

 use anyhow::{Context, Result, bail};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::sync::Arc;

 use tracing::info;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::collections::HashMap;
 use std::sync::Arc;

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use thiserror::Error;
 use thiserror_ext::{AsReport as _, Macro};

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::ops::Deref;
 use std::sync::Arc;

@@ -202,5 +205,8 @@ pub enum ChatEvent {
        finish_reason: FinishReason,
        /// Connector-specific KV transfer parameters for disaggregated serving.
        kv_transfer_params: Option<serde_json::Value>,
+        /// Connector-specific encoder cache transfer parameters for
+        /// disaggregated serving.
+        ec_transfer_params: Option<serde_json::Value>,
    },
 }
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Minimal chat facade above [`vllm_text`].
 //!
 //! This crate keeps the northbound boundary intentionally small:
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Chat-layer multimodal media preparation.
 //!
 //! This module owns the multimodal path for chat requests: it extracts media
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Prompt placeholder expansion shared across modalities.

 use std::collections::{HashMap, VecDeque};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Image-modality preparation: batch preprocessing and per-item feature
 //! build.

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::collections::HashMap;

 use half::{bf16, f16};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Video-modality preparation: per-clip preprocessing, config resolution,
 //! and per-item feature build.

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Default output processing pipeline.

 mod structural_tag;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Applies xgrammar structural-tag constraints for strict tool calling.

 use thiserror_ext::AsReport;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Adapts decoded text updates into parsed assistant deltas.
 //!
 //! This stage sits between low-level token decoding and final block assembly.
@@ -257,6 +260,7 @@ pub(crate) async fn unified_event_stream(
                        usage: finished.usage,
                        finish_reason: finished.finish_reason,
                        kv_transfer_params: finished.kv_transfer_params,
+                        ec_transfer_params: finished.ec_transfer_params,
                    })
                    .await;
                }
@@ -387,6 +391,7 @@ mod tests {
                usage: vllm_llm::TokenUsage::default(),
                finish_reason: crate::FinishReason::Stop(None),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            }),
        }
    }
@@ -628,6 +633,7 @@ mod tests {
                    usage: vllm_llm::TokenUsage::default(),
                    finish_reason: crate::FinishReason::Stop(None),
                    kv_transfer_params: None,
+                    ec_transfer_params: None,
                },
            ]
        );
@@ -671,6 +677,7 @@ mod tests {
                    usage: vllm_llm::TokenUsage::default(),
                    finish_reason: crate::FinishReason::Stop(None),
                    kv_transfer_params: None,
+                    ec_transfer_params: None,
                },
            ]
        );
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Native Harmony output processing for `gpt_oss`.
 //!
 //! Unlike the default text-first pipeline, this processor consumes
@@ -370,6 +373,7 @@ async fn harmony_assistant_event_stream(
                        usage: finished.usage,
                        finish_reason: finished.finish_reason,
                        kv_transfer_params: finished.kv_transfer_params,
+                        ec_transfer_params: finished.ec_transfer_params,
                    })
                    .await;
                }
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::sync::Arc;

 use futures::executor::block_on;
@@ -52,6 +55,7 @@ fn finished() -> Finished {
        },
        finish_reason: FinishReason::stop_eos(),
        kv_transfer_params: None,
+        ec_transfer_params: None,
    }
 }

@@ -115,6 +119,7 @@ fn interrupted_final_message_is_preserved() {
            },
            finish_reason: FinishReason::stop_eos(),
            kv_transfer_params: None,
+            ec_transfer_params: None,
        })
    );
 }
@@ -175,6 +180,7 @@ fn interrupted_analysis_message_is_preserved() {
            },
            finish_reason: FinishReason::stop_eos(),
            kv_transfer_params: None,
+            ec_transfer_params: None,
        })
    );
 }
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::pin::Pin;
 use std::sync::Arc;

@@ -48,6 +51,9 @@ pub(crate) enum AssistantEvent {
        finish_reason: FinishReason,
        /// Connector-specific KV transfer parameters for disaggregated serving.
        kv_transfer_params: Option<serde_json::Value>,
+        /// Connector-specific encoder cache transfer parameters for
+        /// disaggregated serving.
+        ec_transfer_params: Option<serde_json::Value>,
    },
 }

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Adapts parsed assistant updates into structured chat events.
 //!
 //! This module remains the final assembly stage in `vllm-chat`. Token-to-text
@@ -146,6 +149,7 @@ impl StructuredEventState {
        usage: vllm_llm::TokenUsage,
        finish_reason: FinishReason,
        kv_transfer_params: Option<serde_json::Value>,
+        ec_transfer_params: Option<serde_json::Value>,
    ) -> Result<Vec<ChatEvent>> {
        let mut events = Vec::new();
        self.close_open_text_block(&mut events);
@@ -155,6 +159,7 @@ impl StructuredEventState {
            usage,
            finish_reason,
            kv_transfer_params,
+            ec_transfer_params,
        });
        Ok(events)
    }
@@ -296,8 +301,11 @@ pub(crate) async fn structured_chat_event_stream(
                usage,
                finish_reason,
                kv_transfer_params,
+                ec_transfer_params,
            } => {
-                for next in state.finish(usage, finish_reason, kv_transfer_params)? {
+                for next in
+                    state.finish(usage, finish_reason, kv_transfer_params, ec_transfer_params)?
+                {
                    y.yield_ok(next).await;
                }
            }
@@ -334,6 +342,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            }),
        ]);

@@ -388,6 +397,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            }),
        ]);

@@ -439,6 +449,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            }),
        ]);

@@ -490,6 +501,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            }),
        ]);

@@ -557,6 +569,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            }),
        ]);

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 pub mod reasoning;
 pub mod tool;
 pub mod unified;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Reasoning parser registration and selection boundary for `vllm-chat`.

 use std::sync::{Arc, LazyLock};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::sync::Arc;

 use vllm_tokenizer::test_utils::TestTokenizer;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Tool parser registration and selection boundary for `vllm-chat`.

 use std::sync::{Arc, LazyLock};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use vllm_parser::tool::{Result, ToolParserOutput};

 use super::{ToolParser, ToolParserFactory, names};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Unified parser registration and selection boundary for `vllm-chat`.

 use std::sync::LazyLock;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! DeepSeek V3.2 prompt renderer.

 use std::collections::{HashMap, HashSet};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 mod encoding;

 use vllm_text::Prompt;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::path::PathBuf;

 use expect_test::{ExpectFile, expect, expect_file};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! DeepSeek V4 prompt renderer.
 //!
 //! Original Python implementation:
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 mod encoding;

 use vllm_text::Prompt;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::path::PathBuf;

 use expect_test::{ExpectFile, expect, expect_file};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Shared Harmony encoding helper for the GPT-OSS renderer and output parser.

 use std::sync::LazyLock;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Native Harmony chat renderer for `gpt_oss`.

 pub(crate) mod encoding;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::path::PathBuf;

 use expect_test::{ExpectFile, expect, expect_file};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use thiserror::Error as ThisError;

 #[derive(Debug, ThisError)]
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::collections::{HashSet, VecDeque};
 use std::fmt;
 use std::str::FromStr;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::collections::HashMap;

 use serde::Serialize;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Chat template support for tokenizers using Jinja2 templates.
 //!
 //! This module is inlined from SMG's tokenizer crate with local adaptations:
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use minijinja::value::{Kwargs, ViaDeserialize};
 use minijinja::{Error as MinijinjaError, ErrorKind, Value};
 use serde::Deserialize;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::sync::Arc;

 use indexmap::IndexMap;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::collections::HashMap;
 use std::sync::Arc;

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::fmt;
 use std::str::FromStr;

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::fs;
 use std::path::Path;

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::collections::HashMap;

 use llm_multimodal::ImageDetail;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::pin::Pin;
 use std::sync::Arc;
 use std::task::{Context, Poll};
@@ -22,6 +25,9 @@ pub struct CollectedAssistantMessage {
    pub finish_reason: FinishReason,
    /// Connector-specific KV transfer parameters for disaggregated serving.
    pub kv_transfer_params: Option<serde_json::Value>,
+    /// Connector-specific encoder cache transfer parameters for disaggregated
+    /// serving.
+    pub ec_transfer_params: Option<serde_json::Value>,
 }

 /// Per-request stream of chat events.
@@ -77,6 +83,7 @@ impl ChatEventStream {
                    usage,
                    finish_reason,
                    kv_transfer_params,
+                    ec_transfer_params,
                } => {
                    return Ok(CollectedAssistantMessage {
                        message: done,
@@ -89,6 +96,7 @@ impl ChatEventStream {
                        usage,
                        finish_reason,
                        kv_transfer_params,
+                        ec_transfer_params,
                    });
                }
                ChatEvent::ToolCallEnd { call, .. } => {
@@ -194,6 +202,7 @@ mod tests {
                    },
                    finish_reason: FinishReason::stop_eos(),
                    kv_transfer_params: None,
+                    ec_transfer_params: None,
                }),
            ]),
        );
@@ -234,6 +243,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            }
        );
    }
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::collections::BTreeSet;
 use std::fmt;
 use std::sync::Arc;
@@ -50,6 +53,7 @@ fn request_output(
        stop_reason,
        events: None,
        kv_transfer_params: None,
+        ec_transfer_params: None,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -75,6 +79,7 @@ fn request_output_with_logprobs(
        stop_reason,
        events: None,
        kv_transfer_params: None,
+        ec_transfer_params: None,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! Roundtrip tests for the real chat-template and output-processor pairing.
 //!
 //! The invariant under test is that a structured assistant message rendered as history can be
@@ -676,6 +679,7 @@ fn decoded_completion_stream(
                    usage: Default::default(),
                    finish_reason: FinishReason::stop_eos(),
                    kv_transfer_params: None,
+                    ec_transfer_params: None,
                }),
            }
        });
@@ -686,6 +690,7 @@ fn decoded_completion_stream(
                usage: Default::default(),
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
+                ec_transfer_params: None,
            });
            events.push(DecodedTextEvent::TextDelta {
                delta: chunk.delta,
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 //! CLI argument definitions for the `vllm-rs` binary.
 //!
 //! Python vLLM references:
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use expect_test::expect;
 use vllm_engine_core_client::TransportMode;
 use vllm_server::{Config, HttpListenerMode, ParserSelection, RendererSelection};
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 #![allow(clippy::doc_lazy_continuation)]

 use std::fmt::Display;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::{env, fmt, process};

 use time::UtcOffset;
@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 mod cli;
 mod logging;

@@ -1,3 +1,6 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
 use std::time::Duration;

 use anyhow::{Context, Result, bail};
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Wentao YeandGitHub	a9ee5833ca	Merge branch 'main' into wentao-fix-es-v2-bug	2026-07-13 14:19:23 -04:00
akii96 GitHub Ben Browning	550218b136	[Bugfix][Frontend] Flush engine reasoning parser at engine-reasoning → tool streaming boundary (#47606 ) Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com> Co-authored-by: Ben Browning <56071+bbrowning@users.noreply.github.com>	2026-07-13 14:06:10 -04:00
Gavin MorrisandGitHub	5c342876a6	[Doc] Add DeepseekV32ForCausalLM to supported_models.md (#48293 ) Signed-off-by: Gavin Morris <gmorriscs@gmail.com>	2026-07-13 17:43:59 +00:00
stefankoncarevic GitHub Harry Mellor	9427c45386	[ROCm][CI] Transformers: pass only one of input_ids/inputs_embeds (#48258 ) Signed-off-by: Stefan Koncarevic <stefan.koncarevic@amd.com> Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-13 17:28:50 +00:00
omerpaz95 GitHub Or Ozeri	43c8cbf79b	[EC Connector] CPU Offloading EC Connector (#47423 ) Signed-off-by: omerpaz95 <omerpaz95@gmail.com> Signed-off-by: Or Ozeri <oro@il.ibm.com> Co-authored-by: Or Ozeri <oro@il.ibm.com>	2026-07-13 20:09:41 +03:00
Taneem Ibrahim GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	62286308c9	[Misc] Improve Matryoshka pooling dimensions validation (#48057 ) Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-07-13 12:57:36 -04:00
Nick HillandGitHub	26587f9519	[BugFix][ModelRunner V2] Fix stale attn metadata in speculator prefill cudagraph capture (#48261 ) Signed-off-by: Nick Hill <nickhill123@gmail.com>	2026-07-13 09:39:15 -07:00
xiangdong GitHub Kunshang Ji	93e3bc8f30	[XPU][CI]Adjust timeout_in_minutes in Intel GPU CI (#48418 ) Signed-off-by: zengxian <xiangdong.zeng@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-07-13 23:11:16 +08:00
Yan MaandGitHub	c2c9f7c5e2	remove force channels_last in Idefics3MultiModalProcessor (#48467 ) Signed-off-by: Yan Ma <yan.ma@intel.com>	2026-07-13 14:18:57 +00:00
Omer Ullman ArgovandGitHub	1be6e937b2	lower memory required for capturing cudagraphs for large cudagraph sizes (#48483 ) Signed-off-by: Omer Ullman Argov <118735753+omera-nv@users.noreply.github.com>	2026-07-13 10:14:25 -04:00
Wentao YeandGitHub	b3cfca996c	[Mypy Fix] Split mypy work (#48490 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-07-13 12:42:42 +00:00
Bugen ZhaoandGitHub	487dfb3418	[CI] Add SPDX license header to Rust/Protobuf sources (#48472 ) Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-13 10:22:47 +01:00
Karthik Kothuri GitHub Claude wang.yuqi	107a03ba63	[Core] Support fp32 lm_head for generation models via head_dtype (RFC #48305 §3.6) (#48390 ) Signed-off-by: Karthik Kothuri <karthikkothuri2009@gmail.com> Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: wang.yuqi <yuqi.wang@daocloud.io>	2026-07-13 16:43:34 +08:00
Nicolò Lucchesi GitHub Claude	56a357ed33	[Bugfix][KV Cache] Don't route uniform-page-size MLA+SWA models into DeepseekV4 packing (#48256 ) Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-13 08:16:24 +00:00
Nicolò Lucchesi GitHub Claude Opus 4.8 Andreas Karatzas	bea70c7cfc	[Attention] Make sliding-window support an explicit backend capability (#48011 ) Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com>	2026-07-13 01:07:56 -07:00
Mohammad Miadh AngkadandGitHub	75fe92a316	[Distributed][Perf] Enable FlashInfer MNNVL allreduce RMS quant fusion (#48064 ) Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>	2026-07-13 15:02:59 +08:00
Andreas Karatzas GitHub Matthew Wong	b7b58d1eba	[ROCm][CI] Cache Rust builds by source inputs (#46527 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com> Co-authored-by: Matthew Wong <Matthew.Wong2@amd.com>	2026-07-13 01:14:08 -05:00
Canlin GuoandGitHub	36484e464a	[BugFix] Restore full tokens for Qwen MTP When MoE SP (#48429 ) Signed-off-by: gcanlin <canlinguosdu@gmail.com>	2026-07-13 13:29:41 +08:00
Rehan Khan GitHubLi, Jiang <jiang1.li@intel.com>	9e57de7197	[CPU] Create Proper Numa topology for s390x (#40714 ) Signed-off-by: Rehan Khan <Rehan.Khan7@ibm.com> Co-authored-by: Li, Jiang <jiang1.li@intel.com>	2026-07-13 12:58:43 +08:00
Yejing LaiandGitHub	8c5dafcd09	[Bugfix][UT]Fix EagleMiniCPMForCausalLM meet TypeError (#48452 ) Signed-off-by: Lai, Yejing <yejing.lai@intel.com>	2026-07-13 04:37:23 +00:00
guybd GitHub CursorLi, Jiang <jiang1.li@intel.com>	05fa8183a6	[CPU][Spec Decode] Support DFlash speculative decoding for GDN models on CPU (#46090 ) Signed-off-by: guybd <guy.boudoukh@intel.com> Signed-off-by: Guy Boudoukh <guy.boudoukh@intel.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Li, Jiang <jiang1.li@intel.com>	2026-07-13 04:16:18 +00:00
Rohan Potdar GitHub Claude Opus 4.8	d973cce3ca	Re-disable CUDA graph memory profiling on ROCm (#48440 ) Signed-off-by: Rohan Potdar <rohan.potdar@amd.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>	2026-07-13 03:59:20 +00:00
Fangzhou Ai GitHub Cursor	775c1589ea	[Bugfix][ROCm] Keep TP all_gather on base-class collective (#48446 ) Signed-off-by: fai <fangzhouai@gmail.com> Co-authored-by: Cursor <cursoragent@cursor.com>	2026-07-13 03:53:53 +00:00
zztandGitHub	2595d5cebc	[Model] Optimize Qwen3.5 on H20 (#48350 ) Signed-off-by: zzt <zengzetang.zzt@antgroup.com>	2026-07-13 03:30:48 +00:00
Tan Pin Siang GitHub vllmellm Hongxia Yang Jun Kang Chow tjtanaa	ee5a89f4d7	[ROCm][MiniMax-M3] Add AITER sparse paged attention (#47287 ) Signed-off-by: Tan Pin Siang <tanpinsiang@gmail.com> Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com> Co-authored-by: vllmellm <vllm.ellm@embeddedllm.com> Co-authored-by: Hongxia Yang <hongxia.yang@amd.com> Co-authored-by: Jun Kang Chow <junkangchow@gmail.com> Co-authored-by: tjtanaa <tunjian.tan@embeddedllm.com>	2026-07-12 19:27:29 -07:00
Tanish Malekar GitHub Claude Sonnet 4.6	e26264f3ef	[Kernel] Implement CUDA kernel for ReLUSquaredActivation (relu^2) (#39058 ) Signed-off-by: Tanish Malekar <tanishmalekar32@gmail.com> Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>	2026-07-12 19:18:03 -07:00
AlexHuangandGitHub	4c81772e8b	[Bugfix][KV Offloading] Fix stale transfer_jobs after reset_cache + harden job completion (#48102 ) Signed-off-by: Alex <alex.tech.lab@outlook.com>	2026-07-12 20:00:04 +03:00
Bugen ZhaoandGitHub	27c3e579f0	[CI][Rust Frontend] Pin cargo tool versions (#48222 ) Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-12 16:34:26 +01:00
omerpaz95 GitHub Or Ozeri	8df14cfc8c	[EC Connector] Add EC Transfer Params (#42433 ) Signed-off-by: omerpaz95 <omerpaz95@gmail.com> Co-authored-by: Or Ozeri <oro@il.ibm.com>	2026-07-12 14:35:33 +03:00
Jiangyun ZhuandGitHub	370b678a02	[CI][2/N] reduce CI time (#48394 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>	2026-07-12 04:16:55 -07:00
liranschour GitHub Or Ozeri	5c0c987c03	Make tiering offload region DP-replica aware (#47987 ) Signed-off-by: Liran Schour <lirans@il.ibm.com> Co-authored-by: Or Ozeri <oro@il.ibm.com>	2026-07-12 13:10:21 +03:00
Hugo CentenoandGitHub	5f8e73cb8b	[Bugfix] Guard mixed-dtype allreduce RMSNorm quant fusions (#48330 ) Signed-off-by: hcenteno <hugo.centeno@estudiantat.upc.edu>	2026-07-12 09:39:27 +00:00
aoshen02 GitHub Claude Opus 4.8mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	83762b77b0	[Frontend] Add /abort_requests to the RLHF dev API router (#47173 ) Signed-off-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-07-12 14:21:02 +08:00
Canlin Guo GitHub yewentao256	a02984ed47	[Perf][Qwen] Replace MOE all-reduce with reduce-scatter (#47006 ) Signed-off-by: gcanlin <canlinguosdu@gmail.com> Signed-off-by: yewentao256 <zhyanwentao@126.com> Co-authored-by: yewentao256 <zhyanwentao@126.com>	2026-07-12 06:14:49 +00:00
vx120 GitHub crp0128 aoshen02 Claude Opus 4.8	fc1c548093	Runtime Draft Weight Update for Speculative Decoding (#46725 ) Signed-off-by: vx120 <893600387@qq.com> Signed-off-by: vx120 <57470515+vx120@users.noreply.github.com> Signed-off-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: crp0128 <191679376@qq.com> Co-authored-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-11 22:51:53 -07:00
Jiangyun Zhu GitHub Yifan Qiao	481e481be7	[2/N][Core] support partial prefix cache hit for hybrid model (#46384 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai> Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>	2026-07-12 05:37:51 +00:00
zhao, zhenhuiandGitHub	8e981630c9	[CI][CPU] Add Qwen2-VL multimodal tests for CPU backend and fix incompatibilities (#48072 ) Signed-off-by: Zhenhui Zhao <zhenhui.zhao@intel.com>	2026-07-12 12:30:34 +08:00
Alejandro Paredes La TorreandGitHub	9a48eef89a	[Bugfix][LoRA] Support ark_linear base layer in _get_lora_device (#47690 ) Signed-off-by: AlejandroParedesLT <alejandroparedeslatorre@gmail.com>	2026-07-12 00:13:50 +00:00
Jiangyun ZhuandGitHub	1ef1c7ebba	[CI] split tests to reduce CI time (#48219 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>	2026-07-11 13:00:14 -07:00
Ievgen Bondarenko GitHub Claude Opus 4.7 Isotr0py	54503ecec0	fix(processor): route MiMo-V2-Omni media fetch through MediaConnector (#43117 ) Signed-off-by: Ievgen Bondarenko <ibondarenko@student.sierracollege.edu> Signed-off-by: Ievgen (Jack) Bondarenko <ibondarenko@student.sierracollege.edu> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com>	2026-07-11 15:52:52 +00:00
ErenAta16andGitHub	0067311536	fix(entrypoints): stop resolve_items leaking in-flight media fetch tasks on partial failure (#48333 ) Signed-off-by: ErenAta16 <erena6466@gmail.com>	2026-07-11 15:42:08 +00:00
Lucas Wilkinson GitHub Claude OpenAI Codex Matthew Bonanni Nicolò Lucchesi	51878e5b6e	[2/N][KV-Cache Layout Refactor] Pack K/V into the content dim across attention backends (#44455 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Signed-off-by: Matthew Bonanni <mbonanni@redhat.com> Signed-off-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Matthew Bonanni <mbonanni@redhat.com> Co-authored-by: Nicolò Lucchesi <nlucches@redhat.com>	2026-07-11 11:11:16 -04:00
Yejing LaiandGitHub	76fedaa2a5	[XPU][UT]Fix InternS1ProForConditionalGeneration AssertionError (#48232 ) Signed-off-by: Lai, Yejing <yejing.lai@intel.com>	2026-07-11 13:56:55 +00:00
Jee Jee Li GitHub peiyuanz Claude zhouzhou	19069bcbd5	FP32 router GEMV optimization (#48335 ) Signed-off-by: peiyuanz <peiyuanz@inferact.ai> Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai> Co-authored-by: peiyuanz <peiyuanz@inferact.ai> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: zhouzhou <zhouzhou@zhouzhoudeMacBook-Pro.local>	2026-07-11 13:07:48 +00:00
Harry MellorandGitHub	1bd8f80a64	[CI] Point CI at Transformers release rather than release branch (#48328 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-11 02:31:14 -07:00
Qiming Zhang GitHub Kunshang Ji	0b6636cbcb	[XPU]remove is_xxx from moe class and bump up kernels (#48079 ) Signed-off-by: mayuyuace <qiming1.zhang@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-07-11 09:27:13 +00:00
Harry MellorandGitHub	4a6440acef	Bump Transformers version to 5.13.0 (#47867 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-11 00:56:14 -07:00
Lucas WilkinsonandGitHub	bec0a4ede6	[Revert] [Build] Update vllm ...builds FA3 with torch stable API (#48269 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>	2026-07-11 05:20:25 +00:00
Robert Shaw GitHub Robert Shaw	3d99b0499a	[Logs] DP Supervisor Log Improvement (#48278 ) Signed-off-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal> Co-authored-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal>	2026-07-11 12:07:00 +08:00
Lucas Wilkinson GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	04d553f390	[Misc] Use meta tensor for KV cache stride calculation (#47316 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-07-10 23:24:59 -04:00
Lucas Wilkinson GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	9c18e90f6c	[BugFix] Fix packed HND KV cache reshape for FlashAttention (#47314 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-07-10 23:22:39 -04:00
Jimmy LeeandGitHub	092387963c	[BugFix] weights processing peak memory reduction for nvfp4 MoE layers (#46276 ) Signed-off-by: Jimmy Lee <hirejimmylee@gmail.com>	2026-07-11 02:05:35 +00:00
Joe Rowell GitHub mgoin	1bf3997eae	[Quantization] Bound peak memory when repacking FP4 MoE weights for Marlin (#47851 ) Signed-off-by: Joe Rowell <joerowell4@gmail.com> Signed-off-by: mgoin <mgoin64@gmail.com> Co-authored-by: mgoin <mgoin64@gmail.com>	2026-07-10 19:46:13 -06:00
Michael Goin GitHub Claude Fable 5	29fd688892	Add VLLM_FLASHINFER_AUTOTUNE_SKIP_OPS and skip CuTeDSL fp4_gemm autotuning by default (#48268 ) Signed-off-by: mgoin <mgoin64@gmail.com> Signed-off-by: Michael Goin <mgoin64@gmail.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>	2026-07-10 18:13:08 -07:00
Ashwin GiridharanandGitHub	ed908cf0a0	[Bugfix] Fix thinking_token_budget not enforced after natural </think> re-entry (#45984 ) Signed-off-by: Ashwin Giridharan <girida@amazon.com>	2026-07-10 22:47:51 +00:00
Nick Hill GitHub Claude Opus 4.8	26ff616bbf	[Bugfix][Test] Register Qwen/Qwen3.5-4B example model (#48276 ) Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>	2026-07-10 17:01:23 -04:00
gnovackandGitHub	f378f79b7c	handle topk_ids padding in align sum kernel (#47785 ) Signed-off-by: gnovack <novackgm@gmail.com>	2026-07-10 13:33:28 -07:00
Matthew Bonanni GitHub Claude	735def4fcf	[Bugfix] Fix FlashMLA dense fp8 metadata crash (num_sm_parts clamp) (#48045 ) Signed-off-by: Matthew Bonanni <mbonanni@redhat.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-10 12:24:52 -07:00
Wentao YeandGitHub	11ded87a98	Merge branch 'main' into wentao-fix-es-v2-bug	2026-06-09 14:22:55 -04:00
Wentao YeandGitHub	5e3d8ee466	Merge branch 'main' into wentao-fix-es-v2-bug	2026-06-04 14:36:12 -04:00
Wentao YeandGitHub	022e0cdf32	Merge branch 'main' into wentao-fix-es-v2-bug	2026-06-04 11:09:08 -04:00
yewentao256	16900f1f37	add test Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-04 15:08:42 +00:00
yewentao256	133ef0cbff	Merge branch 'main' into wentao-fix-es-v2-bug Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-03 18:13:22 +00:00
yewentao256	d8fe290d36	down to model runner Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-03 18:12:27 +00:00
yewentao256	555b9d8e25	fix es v2 bug Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-28 16:08:35 -04:00