Upstream sync 25/N: merge d785eb51ce (15 commits, conflict-free)

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

@@ -23,8 +23,8 @@ DEFAULT_CI_BASE_METADATA_VERSION="3"
 # local-source stages rather than unreachable remote-fetch alternatives.
 DEFAULT_ROCM_CSRC_CONTENT_FILES=".dockerignore requirements/common.txt requirements/rocm.txt pyproject.toml setup.py CMakeLists.txt cmake csrc vllm/envs.py vllm/__init__.py tools/build_rust.py"
 DEFAULT_ROCM_CSRC_DOCKERFILE_STAGES="base fetch_vllm_0 fetch_vllm build_vllm_dependencies rocm-triton-kernels csrc-build"
-DEFAULT_ROCM_RUST_CONTENT_FILES=".dockerignore 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"
-DEFAULT_ROCM_RUST_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust-toolchain-input rust_input_0 rust-input rust-toolchain rust-build"
+DEFAULT_ROCM_RUST_CONTENT_FILES=".dockerignore .git_archival.txt pyproject.toml 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"
+DEFAULT_ROCM_RUST_DOCKERFILE_STAGES="base fetch_vllm_0 fetch_vllm vllm-version rust_toolchain_input_0 rust-toolchain-input rust_input_0 rust-input rust-toolchain rust-build"
 # Docker's 128-character tag limit minus the longest cache prefix
 # ("csrc-rocm-branch-" and "rust-rocm-branch-", both 17 characters).
 ROCM_CACHE_BRANCH_TAG_MAX_LEN=111
@@ -523,8 +523,8 @@ prepare_ci_build_context() {
         return 1
     fi
 
-    # setuptools-scm understands Git's stable archive format, so full-source
-    # wheels retain their exact version without copying mutable Git history.
+    # setuptools-scm understands Git's stable archive format, so wheels and
+    # Rust artifacts retain their exact version without copying Git history.
     if ! is_ci_base_target; then
         write_ci_git_archival_metadata "${source_root}" "${context_root}" \
             || return $?

.buildkite/test-amd.yaml

@@ -2356,7 +2356,7 @@ steps:
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
-  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
+  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not server_load"
   - pytest -v -s entrypoints/scale_out/token_in_token_out/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
   - pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
   - pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"

.buildkite/test_areas/rust_frontend.yaml

@@ -58,7 +58,7 @@ steps:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
   # server_load can be flaky under the Rust frontend; keep it excluded for now.
-  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
+  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not server_load"
   # test_generate_logprobs expects Python-style top_logprobs truncation (dedup sampled + cap at max(k, 1)).
   - pytest -v -s entrypoints/scale_out/token_in_token_out/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
   - pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"

benchmarks/kernels/benchmark_moe.py

@@ -766,6 +766,7 @@ def get_model_params(config):
         "DeepseekV2ForCausalLM",
         "DeepseekV3ForCausalLM",
         "DeepseekV32ForCausalLM",
+        "DeepseekV4ForCausalLM",
         "GlmMoeDsaForCausalLM",
         "Glm4MoeForCausalLM",
         "Glm4MoeLiteForCausalLM",

csrc/libtorch_stable/gdn/fused_gdn_decode_kernel.cu

@@ -146,7 +146,7 @@ __device__ __forceinline__ Sum2 warp_reduce_sum_pair(float x, float y) {
   return {x, y};
 }
 
-template <typename StateT>
+template <typename StateT, int ValueHeadsPerKeyHead>
 __global__ __launch_bounds__(kThreads, 2) void gdn_decode_post_conv_mtp_kernel(
     const __nv_bfloat16* __restrict__ mixed_qkv,
     const __nv_bfloat16* __restrict__ a, const __nv_bfloat16* __restrict__ b,
@@ -186,7 +186,7 @@ __global__ __launch_bounds__(kThreads, 2) void gdn_decode_post_conv_mtp_kernel(
     return;
   }
 
-  const int key_head = value_head / 8;
+  const int key_head = value_head / ValueHeadsPerKeyHead;
   __shared__ StateT shared_state[2][kChunkV][kDimK];
   __shared__ float shared_q[kMaxMtpTokens][kDimK];
   __shared__ float shared_k[kMaxMtpTokens][kDimK];
@@ -370,7 +370,7 @@ __global__ __launch_bounds__(kThreads, 2) void gdn_decode_post_conv_mtp_kernel(
   }
 }
 
-template <typename StateT>
+template <typename StateT, int ValueHeadsPerKeyHead>
 void launch_gdn_decode_post_conv_mtp(
     torch::stable::Tensor const& mixed_qkv, torch::stable::Tensor const& a_log,
     torch::stable::Tensor const& dt_bias,
@@ -395,17 +395,19 @@ void launch_gdn_decode_post_conv_mtp(
       get_current_cuda_stream(mixed_qkv.get_device_index());
   const int num_requests = static_cast<int>(state_indices.size(0));
   const dim3 grid(num_requests, num_value_heads);
-  gdn_decode_post_conv_mtp_kernel<StateT><<<grid, kThreads, 0, stream>>>(
-      static_cast<const __nv_bfloat16*>(mixed_qkv.data_ptr()), a, b,
-      static_cast<const float*>(a_log.data_ptr()), dt_bias.data_ptr(),
-      static_cast<const int*>(state_indices.data_ptr()),
-      static_cast<const int*>(cu_seqlens.data_ptr()),
-      static_cast<const int*>(num_accepted_tokens.data_ptr()),
-      static_cast<StateT*>(state.data_ptr()), output_gate,
-      norm_weight.data_ptr(), static_cast<__nv_bfloat16*>(out.data_ptr()),
-      num_key_heads, num_value_heads, static_cast<int>(state_indices.size(1)),
-      dt_bias_type, norm_weight.scalar_type() == ScalarType::BFloat16,
-      static_cast<float>(scale), static_cast<float>(norm_eps), strides);
+  gdn_decode_post_conv_mtp_kernel<StateT, ValueHeadsPerKeyHead>
+      <<<grid, kThreads, 0, stream>>>(
+          static_cast<const __nv_bfloat16*>(mixed_qkv.data_ptr()), a, b,
+          static_cast<const float*>(a_log.data_ptr()), dt_bias.data_ptr(),
+          static_cast<const int*>(state_indices.data_ptr()),
+          static_cast<const int*>(cu_seqlens.data_ptr()),
+          static_cast<const int*>(num_accepted_tokens.data_ptr()),
+          static_cast<StateT*>(state.data_ptr()), output_gate,
+          norm_weight.data_ptr(), static_cast<__nv_bfloat16*>(out.data_ptr()),
+          num_key_heads, num_value_heads,
+          static_cast<int>(state_indices.size(1)), dt_bias_type,
+          norm_weight.scalar_type() == ScalarType::BFloat16,
+          static_cast<float>(scale), static_cast<float>(norm_eps), strides);
   const cudaError_t error = cudaGetLastError();
   STD_TORCH_CHECK(error == cudaSuccess,
                   "GDN decode MTP post-conv kernel launch failed: ",
@@ -479,8 +481,12 @@ void fused_gdn_decode_post_conv_mtp(
   STD_TORCH_CHECK(key_width > 0 && key_width % (2 * kDimK) == 0,
                   "mixed_qkv width is inconsistent with state");
   const int num_key_heads = static_cast<int>(key_width / (2 * kDimK));
-  STD_TORCH_CHECK(num_value_heads == 8 * num_key_heads,
-                  "GDN decode MTP fusion requires HV/H=8");
+  const int value_heads_per_key_head = num_value_heads / num_key_heads;
+  STD_TORCH_CHECK(
+      num_value_heads % num_key_heads == 0 &&
+          ((value_heads_per_key_head >= 1 && value_heads_per_key_head <= 4) ||
+           value_heads_per_key_head == 8),
+      "GDN decode MTP fusion requires HV/H in {1, 2, 3, 4, 8}");
 
   STD_TORCH_CHECK(state_indices.dim() == 2 && state_indices.size(0) > 0 &&
                       state_indices.size(1) > 0 &&
@@ -541,17 +547,35 @@ void fused_gdn_decode_post_conv_mtp(
   const auto* b_ptr = static_cast<const __nv_bfloat16*>(b.data_ptr());
   const auto* output_gate_ptr =
       static_cast<const __nv_bfloat16*>(output_gate.data_ptr());
-  if (state_scalar_type == ScalarType::Float) {
-    launch_gdn_decode_post_conv_mtp<float>(
-        mixed_qkv, a_log, dt_bias, state_indices, cu_seqlens,
-        num_accepted_tokens, state, norm_weight, out, a_ptr, b_ptr,
-        output_gate_ptr, num_key_heads, num_value_heads, scale, norm_eps,
-        strides);
-  } else {
-    launch_gdn_decode_post_conv_mtp<__nv_bfloat16>(
+  const auto launch = [&]<typename StateT, int ValueHeadsPerKeyHead>() {
+    launch_gdn_decode_post_conv_mtp<StateT, ValueHeadsPerKeyHead>(
         mixed_qkv, a_log, dt_bias, state_indices, cu_seqlens,
         num_accepted_tokens, state, norm_weight, out, a_ptr, b_ptr,
         output_gate_ptr, num_key_heads, num_value_heads, scale, norm_eps,
         strides);
+  };
+  const auto dispatch_state_type = [&]<int ValueHeadsPerKeyHead>() {
+    if (state_scalar_type == ScalarType::Float) {
+      launch.template operator()<float, ValueHeadsPerKeyHead>();
+    } else {
+      launch.template operator()<__nv_bfloat16, ValueHeadsPerKeyHead>();
+    }
+  };
+  switch (value_heads_per_key_head) {
+    case 1:
+      dispatch_state_type.template operator()<1>();
+      break;
+    case 2:
+      dispatch_state_type.template operator()<2>();
+      break;
+    case 3:
+      dispatch_state_type.template operator()<3>();
+      break;
+    case 4:
+      dispatch_state_type.template operator()<4>();
+      break;
+    default:
+      dispatch_state_type.template operator()<8>();
+      break;
   }
 }

docker/Dockerfile

@@ -279,11 +279,16 @@ ENV CARGO_BUILD_JOBS=4
 # separate local sccache daemon while sharing the same remote cache backend.
 ENV SCCACHE_SERVER_PORT=4227
 
+# Only Rust inputs are present, so suppress the artificial dirty state caused
+# by other tracked files being absent from this build stage.
+ENV SETUPTOOLS_SCM_PRETEND_METADATA={dirty=false}
+
 # Build the release artifacts. Cache cargo registry/git, but not target/,
 # because stale target metadata can outlive source updates across BuildKit
 # cache reuse.
 RUN --mount=type=cache,target=/root/.cargo/registry,sharing=locked \
     --mount=type=cache,target=/root/.cargo/git,sharing=locked \
+    --mount=type=bind,source=.git,target=.git \
     --mount=type=secret,id=aws-credentials,target=/root/.aws/credentials,required=false \
     if [ "$USE_SCCACHE" = "1" ]; then \
         if [ -n "${SCCACHE_ENDPOINT}" ]; then export SCCACHE_ENDPOINT="${SCCACHE_ENDPOINT}"; fi; \

docker/Dockerfile.cpu

@@ -166,11 +166,16 @@ COPY build_rust.sh build_rust.sh
 # (rustc spawns enough concurrent processes to hit RLIMIT_NOFILE otherwise).
 ENV CARGO_BUILD_JOBS=4
 
+# Only Rust inputs are present, so suppress the artificial dirty state caused
+# by other tracked files being absent from this build stage.
+ENV SETUPTOOLS_SCM_PRETEND_METADATA={dirty=false}
+
 # Build the release artifacts. Cache cargo registry/git, but not target/,
 # because stale target metadata can outlive source updates across BuildKit
 # cache reuse.
 RUN --mount=type=cache,target=/root/.cargo/registry,sharing=locked \
     --mount=type=cache,target=/root/.cargo/git,sharing=locked \
+    --mount=type=bind,source=.git,target=.git \
     --mount=type=secret,id=aws-credentials,target=/root/.aws/credentials,required=false \
     if [ "$USE_SCCACHE" = "1" ]; then \
         export RUSTC_WRAPPER=sccache; \

docker/Dockerfile.rocm

@@ -125,6 +125,18 @@ ONBUILD RUN git clone ${VLLM_REPO} \
                && git fetch upstream ; fi
 FROM fetch_vllm_${REMOTE_VLLM} AS fetch_vllm
 
+# Resolve the device-independent source version once for every downstream
+# artifact. setuptools-scm reads Git metadata for local/remote sources and the
+# archive metadata generated for canonical CI contexts.
+FROM fetch_vllm AS vllm-version
+ARG COMMON_WORKDIR
+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 \
+    && python3 -c \
+        "from setuptools_scm import get_version; print(get_version())" \
+        > /vllm-version.txt
+
 # -----------------------
 # Rust/protoc toolchain inputs
 #
@@ -238,6 +250,7 @@ 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
+COPY --from=vllm-version /vllm-version.txt /tmp/vllm-version.txt
 
 # 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
@@ -249,6 +262,7 @@ RUN --mount=type=cache,id=vllm-rocm-cargo-registry,target=/root/.cargo/registry,
         export RUSTC_WRAPPER=sccache \
         && sccache --show-stats; \
     fi \
+    && export SETUPTOOLS_SCM_PRETEND_VERSION="$(cat /tmp/vllm-version.txt)" \
     && bash build_rust.sh \
     && test -x vllm/vllm-rs \
     && if [ "$USE_SCCACHE" = "1" ]; then \
@@ -334,13 +348,15 @@ RUN --mount=type=bind,source=pyproject.toml,target=${COMMON_WORKDIR}/vllm/pyproj
 FROM build_vllm_dependencies AS build_vllm
 COPY --from=fetch_vllm ${COMMON_WORKDIR}/vllm ${COMMON_WORKDIR}/vllm
 COPY --from=csrc-build ${COMMON_WORKDIR}/vllm/dist /precompiled-wheels
+COPY --from=vllm-version /vllm-version.txt /tmp/vllm-version.txt
 
 # Drop the pre-built Rust artifacts into the source tree. setup.py detects
 # them and ships them as-is, skipping the local Rust build.
 COPY --from=rust-build ${COMMON_WORKDIR}/vllm/vllm/vllm-rs ${COMMON_WORKDIR}/vllm/vllm/vllm-rs
 COPY --from=rust-build ${COMMON_WORKDIR}/vllm/vllm/_rust_*.so ${COMMON_WORKDIR}/vllm/vllm/
 
 RUN cd "${COMMON_WORKDIR}/vllm" \
+    && export SETUPTOOLS_SCM_PRETEND_VERSION="$(cat /tmp/vllm-version.txt)" \
     && export VLLM_USE_PRECOMPILED=1 \
     && export VLLM_PRECOMPILED_WHEEL_LOCATION="$(ls /precompiled-wheels/*.whl)" \
     && export VLLM_DOCKER_BUILD_CONTEXT=1 \
@@ -640,6 +656,7 @@ FROM fetch_vllm AS build_vllm_wheel_release
 
 ARG COMMON_WORKDIR
 ARG ROCM_TRITON_KERNELS_COMMIT
+COPY --from=vllm-version /vllm-version.txt /tmp/vllm-version.txt
 COPY --from=rocm-triton-kernels \
     /opt/rocm-triton-kernels /opt/rocm-triton-kernels
 ENV TRITON_KERNELS_SRC_DIR=/opt/rocm-triton-kernels
@@ -678,17 +695,12 @@ RUN if [ "$GIT_REPO_CHECK" != "0" ]; then \
         cd vllm && bash tools/check_repo.sh; \
     fi
 
-# Extract version from git BEFORE any modifications (pin_rocm_dependencies.py modifies requirements/rocm.txt)
-# This ensures setuptools_scm sees clean repo state for version detection
-RUN --mount=type=bind,source=.git,target=vllm/.git \
-    --mount=type=cache,target=/root/.cache/uv \
+# Install the packaging helpers before modifying requirements/rocm.txt.
+RUN --mount=type=cache,target=/root/.cache/uv \
     grep -Fq "set(TRITON_KERNELS_TAG \"${ROCM_TRITON_KERNELS_COMMIT}\")" \
         vllm/cmake/external_projects/triton_kernels.cmake \
     && cd vllm \
-    && uv pip install --system setuptools_scm regex \
-    && VLLM_VERSION=$(python3 -c "import setuptools_scm; print(setuptools_scm.get_version())") \
-    && echo "Detected vLLM version: ${VLLM_VERSION}" \
-    && echo "${VLLM_VERSION}" > /tmp/vllm_version.txt
+    && uv pip install --system setuptools_scm regex
 
 # Fail if git-based package dependencies are found in requirements files
 # (uv doesn't handle git+ URLs well, and packages should be distributed on PyPI)
@@ -726,13 +738,12 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     && echo "Building vLLM with custom wheels from /install" \
     && uv pip install --system --find-links /install -r requirements/rocm.txt
 
-# Build wheel using pre-extracted version to avoid dirty state from modified requirements/rocm.txt
+# Build wheel using the shared source version to avoid dirty state from modified requirements/rocm.txt
 # (setup.py auto-detects ccache/sccache in PATH)
-RUN --mount=type=bind,source=.git,target=vllm/.git \
-    --mount=type=cache,id=vllm-rocm-ccache,target=/root/.cache/ccache \
+RUN --mount=type=cache,id=vllm-rocm-ccache,target=/root/.cache/ccache \
     cd vllm \
     && export CCACHE_BASEDIR="$PWD" \
-    && export SETUPTOOLS_SCM_PRETEND_VERSION=$(cat /tmp/vllm_version.txt) \
+    && export SETUPTOOLS_SCM_PRETEND_VERSION="$(cat /tmp/vllm-version.txt)" \
     && echo "Building wheel with version: ${SETUPTOOLS_SCM_PRETEND_VERSION}" \
     && MAX_JOBS="${MAX_JOBS:-$(nproc)}" python3 setup.py bdist_wheel --dist-dir=dist