Done

Signed-off-by: Jee Jee Li <pandaleefree@gmail.com>
[ROCm][Engine] Fix GPU memory leaks in engine shutdown and test workaround for async KV prefix cache reset (#38503 )
2026-04-25 13:23:05 +00:00 · 2026-04-25 05:25:20 +00:00 · 2026-04-25 03:57:41 +00:00 · 2026-04-25 02:59:31 +00:00 · 2026-04-25 11:20:59 +09:00 · 2026-04-25 00:59:20 +00:00
180 changed files with 11314 additions and 3545 deletions
@@ -46,7 +46,7 @@ echo "Image not found, proceeding with build..."

 # --- CUDA 13.0 for nightly builds ---
 # Nightly CI uses CUDA 13.0 while regular CI stays on CUDA 12.9
-NIGHTLY_CUDA_VERSION="13.0.0"
+NIGHTLY_CUDA_VERSION="13.0.2"
 NIGHTLY_BUILD_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-devel-ubuntu22.04"
 NIGHTLY_FINAL_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-base-ubuntu22.04"

@@ -37,7 +37,7 @@ steps:
        agents:
          queue: arm64_cpu_queue_release
        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04  --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04  --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
@@ -76,7 +76,7 @@ steps:
        agents:
          queue: cpu_queue_release
        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
@@ -121,7 +121,7 @@ steps:
          queue: cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"
          # re-tag to default image tag and push, just in case arm64 build fails
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
@@ -134,7 +134,7 @@ steps:
          queue: arm64_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"

      - label: "Build release image - x86_64 - CUDA 12.9"
@@ -167,7 +167,7 @@ steps:
          queue: cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
@@ -179,7 +179,7 @@ steps:
          queue: arm64_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"

      - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
@@ -25,22 +25,100 @@ export PYTHONPATH=".."
 ###############################################################################

 cleanup_docker() {
+  # Share the same lock with image pull to avoid cleanup/pull races on one node.
+  local docker_lock="/tmp/docker-pull.lock"
+  exec 9>"$docker_lock"
+  flock 9
+
  docker_root=$(docker info -f '{{.DockerRootDir}}')
  if [ -z "$docker_root" ]; then
    echo "Failed to determine Docker root directory." >&2
-    exit 1
+    flock -u 9
+    return 1
  fi
  echo "Docker root directory: $docker_root"

  disk_usage=$(df "$docker_root" | tail -1 | awk '{print $5}' | sed 's/%//')
  threshold=70
  if [ "$disk_usage" -gt "$threshold" ]; then
-    echo "Disk usage is above $threshold%. Cleaning up Docker images and volumes..."
-    docker image prune -f
-    docker volume prune -f && docker system prune --force --filter "until=72h" --all
-    echo "Docker images and volumes cleanup completed."
+    echo "Disk usage is above $threshold%. Running aggressive CI image cleanup..."
+    cleanup_old_ci_images "${REGISTRY}/${REPO}" "${image_name}" "${DOCKER_IMAGE_CLEANUP_HOURS:-72}" 1
  else
-    echo "Disk usage is below $threshold%. No cleanup needed."
+    echo "Disk usage is below $threshold%. Checking old CI images anyway."
+    cleanup_old_ci_images "${REGISTRY}/${REPO}" "${image_name}" "${DOCKER_IMAGE_CLEANUP_HOURS:-72}" 0
+  fi
+  echo "Old CI image cleanup completed."
+
+  flock -u 9
+}
+
+cleanup_old_ci_images() {
+  local repo_prefix="$1"
+  local current_image_ref="$2"
+  local ttl_hours="$3"
+  local aggressive_cleanup="$4"
+
+  if [[ -z "$repo_prefix" || "$repo_prefix" == "/" ]]; then
+    echo "Skip old-image cleanup: invalid repo prefix '${repo_prefix}'"
+    return 0
+  fi
+
+  if ! [[ "$ttl_hours" =~ ^[0-9]+$ ]]; then
+    echo "Invalid DOCKER_IMAGE_CLEANUP_HOURS='${ttl_hours}', fallback to 72"
+    ttl_hours=72
+  fi
+
+  local now_epoch cutoff_epoch
+  now_epoch=$(date +%s)
+  cutoff_epoch=$((now_epoch - ttl_hours * 3600))
+
+  local -a used_image_ids
+  mapfile -t used_image_ids < <(docker ps -aq | xargs -r docker inspect --format '{{.Image}}' | sort -u)
+
+  local removed_count=0
+  local examined_count=0
+  declare -A seen_ids=()
+
+  while read -r image_ref image_id; do
+    [[ -z "$image_ref" || -z "$image_id" ]] && continue
+    ((examined_count++))
+
+    # Keep the image this job is going to use.
+    if [[ "$image_ref" == "$current_image_ref" ]]; then
+      continue
+    fi
+
+    # Avoid duplicate deletes when multiple tags point to same image id.
+    if [[ -n "${seen_ids[$image_id]:-}" ]]; then
+      continue
+    fi
+    seen_ids[$image_id]=1
+
+    # Never delete images that are used by any container on this node.
+    if printf '%s\n' "${used_image_ids[@]}" | grep -qx "$image_id"; then
+      continue
+    fi
+
+    local created created_epoch
+    created=$(docker image inspect -f '{{.Created}}' "$image_id" 2>/dev/null || true)
+    [[ -z "$created" ]] && continue
+    created_epoch=$(date -d "$created" +%s 2>/dev/null || true)
+    [[ -z "$created_epoch" ]] && continue
+
+    if (( created_epoch < cutoff_epoch )) || [[ "$aggressive_cleanup" == "1" ]]; then
+      if docker image rm -f "$image_id" >/dev/null 2>&1; then
+        ((removed_count++))
+      fi
+    fi
+  done < <(docker image ls --no-trunc "$repo_prefix" --format '{{.Repository}}:{{.Tag}} {{.ID}}')
+
+  # Also trim old dangling layers; this is safe and does not remove referenced images.
+  docker image prune -f --filter "until=${ttl_hours}h" >/dev/null 2>&1 || true
+
+  if [[ "$aggressive_cleanup" == "1" ]]; then
+    echo "Examined ${examined_count} images under ${repo_prefix}, removed ${removed_count} unused images under disk pressure."
+  else
+    echo "Examined ${examined_count} images under ${repo_prefix}, removed ${removed_count} old images (>${ttl_hours}h)."
  fi
 }

@@ -265,8 +343,6 @@ fi

 remove_docker_container() {
  docker rm -f "${container_name}" || true
-  docker image rm -f "${image_name}" || true
-  docker system prune -f || true
 }
 trap remove_docker_container EXIT

@@ -12,9 +12,7 @@ docker build -t "${image_name}" -f docker/Dockerfile.xpu .

 # Setup cleanup
 remove_docker_container() {
-  docker rm -f "${container_name}" || true;
-  docker image rm -f "${image_name}" || true;
-  docker system prune -f || true;
+  docker rm -f "${container_name}" || true
 }
 trap remove_docker_container EXIT

@@ -28,6 +28,7 @@ steps:
    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
    - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
+    - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
  mirror:
    amd:
      device: mi325_1
@@ -34,7 +34,7 @@ install(CODE "set(CMAKE_INSTALL_LOCAL_ONLY TRUE)" ALL_COMPONENTS)
 # Supported python versions.  These versions will be searched in order, the
 # first match will be selected.  These should be kept in sync with setup.py.
 #
-set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13")
+set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13" "3.14")

 # Supported AMD GPU architectures.
 set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
@@ -14,7 +14,7 @@ Easy, fast, and cheap LLM serving for everyone
 | <a href="https://docs.vllm.ai"><b>Documentation</b></a> | <a href="https://blog.vllm.ai/"><b>Blog</b></a> | <a href="https://arxiv.org/abs/2309.06180"><b>Paper</b></a> | <a href="https://x.com/vllm_project"><b>Twitter/X</b></a> | <a href="https://discuss.vllm.ai"><b>User Forum</b></a> | <a href="https://slack.vllm.ai"><b>Developer Slack</b></a> |
 </p>

-🔥 We have built a vllm website to help you get started with vllm. Please visit [vllm.ai](https://vllm.ai) to learn more.
+🔥 We have built a vLLM website to help you get started with vLLM. Please visit [vllm.ai](https://vllm.ai) to learn more.
 For events, please visit [vllm.ai/events](https://vllm.ai/events) to join us.

 ---
@@ -50,7 +50,7 @@ vLLM is flexible and easy to use with:
 - Efficient multi-LoRA support for dense and MoE layers
 - Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.

-vLLM seamlessly supports 200+ model architectures on HuggingFace, including:
+vLLM seamlessly supports 200+ model architectures on Hugging Face, including:

 - Decoder-only LLMs (e.g., Llama, Qwen, Gemma)
 - Mixture-of-Expert LLMs (e.g., Mixtral, DeepSeek-V3, Qwen-MoE, GPT-OSS)
@@ -404,6 +404,7 @@ def _build_attention_metadata(
        query_start_loc=q_start_gpu,
        query_start_loc_cpu=q_start_cpu,
        seq_lens=seq_lens_gpu,
+        seq_lens_cpu_upper_bound=seq_lens_cpu,
        _seq_lens_cpu=seq_lens_cpu,
        _num_computed_tokens_cpu=num_computed_tokens_cpu,
        slot_mapping=slot_mapping,
@@ -16,7 +16,7 @@ from vllm.model_executor.layers.fused_moe.all2all_utils import (
    maybe_make_prepare_finalize,
 )
 from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_config
-from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp8
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import CutlassExpertsFp8
 from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts, fused_topk
 from vllm.platforms import current_platform
 from vllm.utils.argparse_utils import FlexibleArgumentParser
@@ -22,7 +22,7 @@ from vllm.model_executor.layers.fused_moe.config import (
    fp8_w8a8_moe_quant_config,
    nvfp4_moe_quant_config,
 )
-from vllm.model_executor.layers.fused_moe.cutlass_moe import (
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import (
    CutlassExpertsFp4,
 )
 from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts, fused_topk
@@ -13,7 +13,7 @@ from vllm.model_executor.layers.fused_moe.all2all_utils import (
    maybe_make_prepare_finalize,
 )
 from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_config
-from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp8
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import CutlassExpertsFp8
 from vllm.model_executor.layers.fused_moe.fused_moe import (
    fused_experts,
    fused_topk,
@@ -599,6 +599,11 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
  const int head_idx = (blockIdx.y * blockDim.x + threadIdx.x) * VEC_SIZE;
  // Find batch index within a block
  __shared__ int batch_idx[BLOCK_Y_SIZE];
+  if (threadIdx.x == 0) {
+    batch_idx[threadIdx.y] = -1;
+  }
+  __syncthreads();
+
  for (int iter = 0; iter < cuda_utils::ceil_div(batch_size, int(blockDim.x));
       iter++) {
    int tid = iter * blockDim.x + threadIdx.x;
@@ -611,16 +616,18 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
    }
  }

-#ifndef USE_ROCM
-  __syncwarp();
-#endif
+  __syncthreads();

-  if (head_idx >= head_dim || token_idx >= num_tokens) {
+  // num_tokens may be an allocation upper bound when Python avoids a D2H sync.
+  // Only tokens covered by the exact device-side cu_seq_lens are valid to
+  // gather.
+  const int batch = batch_idx[threadIdx.y];
+  if (head_idx >= head_dim || token_idx >= num_tokens || batch < 0) {
    return;
  }
-  const int inbatch_seq_idx = token_idx - cu_seq_lens[batch_idx[threadIdx.y]];
-  const int block_idx = block_table[batch_idx[threadIdx.y] * num_blocks +
-                                    inbatch_seq_idx / cache_block_size];
+  const int inbatch_seq_idx = token_idx - cu_seq_lens[batch];
+  const int block_idx =
+      block_table[batch * num_blocks + inbatch_seq_idx / cache_block_size];
  const int64_t src_block_offset = block_idx * block_stride;
  const int64_t cache_inblock_offset =
      (inbatch_seq_idx % cache_block_size) * head_dim + head_idx;
@@ -1490,6 +1497,9 @@ void concat_mla_q(torch::Tensor& ql_nope,  // [num_tokens, num_heads, nope_dim]
  TORCH_CHECK(ql_nope.stride(2) == 1, "ql_nope must have stride 1 in dim 2");
  TORCH_CHECK(q_pe.stride(2) == 1, "q_pe must have stride 1 in dim 2");
  TORCH_CHECK(q_out.stride(2) == 1, "q_out must have stride 1 in dim 2");
+  TORCH_CHECK(ql_nope.scalar_type() == at::ScalarType::Half ||
+                  ql_nope.scalar_type() == at::ScalarType::BFloat16,
+              "ql_nope must be float16 or bfloat16 dtype");

  if (num_tokens == 0) return;

@@ -1501,7 +1511,7 @@ void concat_mla_q(torch::Tensor& ql_nope,  // [num_tokens, num_heads, nope_dim]
  const at::cuda::OptionalCUDAGuard device_guard(device_of(ql_nope));
  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

-  VLLM_DISPATCH_FLOATING_TYPES(ql_nope.scalar_type(), "concat_mla_q", [&] {
+  VLLM_DISPATCH_HALF_TYPES(ql_nope.scalar_type(), "concat_mla_q", [&] {
    vllm::ConcatMLAQKernel<scalar_t, 512><<<grid_size, block_size, 0, stream>>>(
        q_out.data_ptr<scalar_t>(), ql_nope.data_ptr<scalar_t>(),
        q_pe.data_ptr<scalar_t>(), num_tokens, num_heads, q_out.stride(0),
@@ -22,7 +22,7 @@
 #   docker buildx bake -f docker/docker-bake.hcl -f docker/versions.json
 # =============================================================================

-ARG CUDA_VERSION=13.0.0
+ARG CUDA_VERSION=13.0.2
 ARG PYTHON_VERSION=3.12
 ARG UBUNTU_VERSION=22.04

@@ -77,7 +77,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
    uv pip install --system $pkgs --index-url https://download.pytorch.org/whl/nightly/cu128

 RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install --system numba==0.61.2
+    uv pip install --system numba==0.65.0

 RUN --mount=type=cache,target=/root/.cache/uv \
 uv pip install --system -r requirements/common.txt
@@ -124,10 +124,10 @@ COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1

 # RIXL/UCX build stages
 FROM base AS build_rixl
-ARG RIXL_BRANCH="f33a5599"
+ARG RIXL_BRANCH="bf4a7214"
 ARG RIXL_REPO="https://github.com/ROCm/RIXL.git"
-ARG UCX_BRANCH="da3fac2a"
-ARG UCX_REPO="https://github.com/ROCm/ucx.git"
+ARG UCX_BRANCH="7009d7a1"
+ARG UCX_REPO="https://github.com/openucx/ucx.git"
 ENV ROCM_PATH=/opt/rocm
 ENV UCX_HOME=/usr/local/ucx
 ENV RIXL_HOME=/usr/local/rixl
@@ -165,7 +165,7 @@ RUN cd /usr/local/src && \
        --disable-doxygen-doc \
        --enable-optimizations \
        --enable-devel-headers \
-        --with-rocm=/opt/rocm \
+        --with-rocm=${ROCM_PATH} \
        --with-verbs \
        --with-dm \
        --enable-mt && \
@@ -186,7 +186,12 @@ RUN git clone ${RIXL_REPO} /opt/rixl && \
    ninja install

 # Generate RIXL wheel
-RUN cd /opt/rixl && mkdir -p /app/install && \
+# Exclude libcore and libpull from auditwheel: transitive dependencies
+# that are not shipped in the wheel and vary across base images.
+RUN cd /opt/rixl && \
+    sed -i "s/--exclude 'libamdhip64\*'/--exclude 'libamdhip64*' --exclude 'libcore*' --exclude 'libpull*'/" \
+        contrib/build-wheel.sh && \
+    mkdir -p /app/install && \
    ./contrib/build-wheel.sh \
        --output-dir /app/install \
        --rocm-dir ${ROCM_PATH} \
@@ -431,6 +436,10 @@ COPY --from=export_vllm /vllm_v1 /usr/local/lib/python${PYTHON_VERSION}/dist-pac
 ENV MIOPEN_DEBUG_CONV_DIRECT=0
 ENV MIOPEN_DEBUG_CONV_GEMM=0

+# Use legacy IPC mode for HSA to avoid GPU memory pinning issues with UCX rocm_ipc
+# See: https://github.com/ROCm/rocm-libraries/issues/6266
+ENV HSA_ENABLE_IPC_MODE_LEGACY=1
+
 # Source code is used in the `python_only_compile.sh` test
 # We hide it inside `src/` so that this source code
 # will not be imported by other tests
@@ -2,7 +2,7 @@
  "_comment": "Auto-generated from Dockerfile ARGs. Do not edit manually. Run: python tools/generate_versions_json.py",
  "variable": {
    "CUDA_VERSION": {
-      "default": "13.0.0"
+      "default": "13.0.2"
    },
    "PYTHON_VERSION": {
      "default": "3.12"
@@ -11,10 +11,10 @@
      "default": "22.04"
    },
    "BUILD_BASE_IMAGE": {
-      "default": "nvidia/cuda:13.0.0-devel-ubuntu22.04"
+      "default": "nvidia/cuda:13.0.2-devel-ubuntu22.04"
    },
    "FINAL_BASE_IMAGE": {
-      "default": "nvidia/cuda:13.0.0-base-ubuntu22.04"
+      "default": "nvidia/cuda:13.0.2-base-ubuntu22.04"
    },
    "GET_PIP_URL": {
      "default": "https://bootstrap.pypa.io/get-pip.py"
@@ -108,6 +108,38 @@ P99 ITL (ms):                            8.39
 ==================================================
 ```

+#### Results Visualization
+
+The `--plot-timeline` and `--plot-dataset-stats` can be used to generate respectively the requests completion timeline and dataset prompt and output tokens statistics, which can be useful for debugging purpose or for deeper analysis.
+
+```bash
+vllm bench serve \
+    --backend vllm \
+    --model meta-llama/Llama-3.1-8B-Instruct \
+    --endpoint /v1/completions \
+    --dataset-name sharegpt \
+    --dataset-path <your data path>/ShareGPT_V3_unfiltered_cleaned_split.json \
+    --num-prompts 100 \
+    --plot-timeline \
+    --timeline-itl-thresholds 2,5 \
+    --plot-dataset-stats \
+    --save-result
+```
+
+##### Interactive Timeline
+
+The generated timeline is an interactive visualization in the form of an HTML file that can be rendered in most browsers. To customize the ITL color thresholds, one can use `--timeline-itl-thresholds` flag (default: 25ms, 50ms)
+
+Example output:
+
+<iframe src="../../assets/contributing/vllm_bench_serve_timeline.html" width="100%" height="600" frameborder="0"></iframe>
+
+##### Dataset statistics
+
+The generated figure shows the input prompt and output tokens distribution.
+
+Example output: ![Dataset Statistics](../assets/contributing/vllm_bench_serve_dataset_stats.png)
+
 #### Custom Dataset

 If the dataset you want to benchmark is not supported yet in vLLM, even then you can benchmark on it using `CustomDataset`. Your data needs to be in `.jsonl` format and needs to have "prompt" field per entry, e.g., data.jsonl
@@ -172,7 +172,7 @@ Priority is **1 = highest** (tried first).
 | `FLASHINFER` | TRTLLM† | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32, 64 | 64, 128, 256 | ✅ | ❌ | ✅ | Decoder | 10.x |
 | `FLASH_ATTN` | FA2* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ❌ | ✅ | All | ≥8.0 |
 | `FLASH_ATTN` | FA3* | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | Any | ✅ | ❌ | ✅ | All | 9.x |
-| `FLASH_ATTN` | FA4* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ❌ | ✅ | All | ≥10.0 |
+| `FLASH_ATTN` | FA4* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ✅ | ❌ | ✅ | All | ≥10.0 |
 | `FLASH_ATTN_DIFFKV` | | fp16, bf16 | `auto` | Any | Any | ❌ | ❌ | ✅ | Decoder | Any |
 | `FLEX_ATTENTION` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16` | Any | Any | ❌ | ✅ | ❌ | Decoder, Encoder Only | Any |
 | `ROCM_AITER_FA` | | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32 | 64, 128, 256 | ❌ | ❌ | ❌ | Decoder | N/A |
@@ -83,8 +83,8 @@ To be used with a particular `FusedMoEPrepareAndFinalizeModular` subclass, MoE k
 | triton | standard | all<sup>1</sup> | G,A,T | silu, gelu,</br>swigluoai,</br>silu_no_mul,</br>gelu_no_mul | Y | Y | [`fused_experts`][vllm.model_executor.layers.fused_moe.fused_moe.fused_experts],</br>[`TritonExperts`][vllm.model_executor.layers.fused_moe.fused_moe.TritonExperts] |
 | triton (batched) | batched | all<sup>1</sup> | G,A,T | silu, gelu | <sup>6</sup> | Y | [`BatchedTritonExperts`][vllm.model_executor.layers.fused_moe.fused_batched_moe.BatchedTritonExperts] |
 | deep gemm | standard,</br>batched | fp8 | G(128),A,T | silu, gelu | <sup>6</sup> | Y | </br>[`DeepGemmExperts`][vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe.DeepGemmExperts],</br>[`BatchedDeepGemmExperts`][vllm.model_executor.layers.fused_moe.experts.batched_deep_gemm_moe.BatchedDeepGemmExperts] |
-| cutlass_fp4 | standard,</br>batched | nvfp4 | A,T | silu | Y | Y | [`CutlassExpertsFp4`][vllm.model_executor.layers.fused_moe.cutlass_moe.CutlassExpertsFp4] |
-| cutlass_fp8 | standard,</br>batched | fp8 | A,T | silu, gelu | Y | Y | [`CutlassExpertsFp8`][vllm.model_executor.layers.fused_moe.cutlass_moe.CutlassExpertsFp8],</br>[`CutlasBatchedExpertsFp8`][vllm.model_executor.layers.fused_moe.cutlass_moe.CutlassBatchedExpertsFp8] |
+| cutlass_fp4 | standard,</br>batched | nvfp4 | A,T | silu | Y | Y | [`CutlassExpertsFp4`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassExpertsFp4] |
+| cutlass_fp8 | standard,</br>batched | fp8 | A,T | silu, gelu | Y | Y | [`CutlassExpertsFp8`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassExpertsFp8],</br>[`CutlasBatchedExpertsFp8`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassBatchedExpertsFp8] |
 | flashinfer | standard | nvfp4,</br>fp8 | T | <sup>5</sup> | N | Y | [`FlashInferExperts`][vllm.model_executor.layers.fused_moe.flashinfer_cutlass_moe.FlashInferExperts] |
 | gpt oss triton | standard | N/A | N/A | <sup>5</sup> | Y | Y | [`triton_kernel_fused_experts`][vllm.model_executor.layers.fused_moe.experts.gpt_oss_triton_kernels_moe.triton_kernel_fused_experts],</br>[`OAITritonExperts`][vllm.model_executor.layers.fused_moe.experts.gpt_oss_triton_kernels_moe.OAITritonExperts] |
 | marlin | standard,</br>batched | <sup>3</sup> / N/A | <sup>3</sup> / N/A | silu,</br>swigluoai | Y | Y | [`fused_marlin_moe`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.fused_marlin_moe],</br>[`MarlinExperts`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.MarlinExperts],</br>[`BatchedMarlinExperts`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.BatchedMarlinExperts] |
@@ -0,0 +1,70 @@
+# Context Extension
+
+!!! note
+    The `--rope-scaling` parameter used in older versions of vLLM is no longer supported. Please use the `--hf-overrides` method with `rope_parameters` instead.
+This directory contains examples for extending the context length of models using vLLM.
+
+## Offline Inference Example
+
+The [`context_extension.py`](../../examples/offline_inference/context_extension) script demonstrates how to extend the context length of a Qwen model using the YARN method (rope_parameters) and run a simple chat example.
+
+### Usage
+
+```bash
+python examples/offline_inference/context_extension.py
+```
+
+## OpenAI Online Method
+
+You can also use vLLM's OpenAI-compatible API to serve models with extended context length.
+
+### Usage
+
+Run the vLLM server with the following command to extend the context length using YARN:
+
+```bash
+vllm serve Qwen/Qwen3-0.6B \
+  --hf-overrides '{"rope_parameters": {"factor": 4.0, "original_max_position_embeddings": 32768, "rope_theta": 1000000, "rope_type": "yarn"}}' \
+  --max-model-len 131072
+```
+
+### Client Example
+
+After starting the server, you can use the OpenAI Python client to interact with it:
+
+```python
+from openai import OpenAI
+
+client = OpenAI(
+    base_url="http://localhost:8000/v1",
+    api_key="token-abc123"  # Dummy API key, required by the client
+)
+
+response = client.chat.completions.create(
+    model="Qwen/Qwen3-0.6B",
+    messages=[
+        {"role": "system", "content": "You are a helpful assistant"},
+        {"role": "user", "content": "Hello"}
+    ],
+    max_tokens=128,
+    temperature=0.8,
+    top_p=0.95
+)
+
+print(response.choices[0].message.content)
+```
+
+### Key Parameters
+
+The available parameters depend on the `rope_type` you choose. For detailed information about all supported RoPE types and their specific parameters, please refer to the [Hugging Face Transformers RoPE documentation](https://huggingface.co/docs/transformers/main/en/internal/rope_utils#transformers.RopeParameters).
+
+Common parameters include:
+
+- `rope_type`: The type of RoPE implementation (e.g., "yarn", "linear", "dynamic")
+- `factor`: The factor by which to extend the context length
+- `original_max_position_embeddings`: The original maximum position embeddings of the model
+
+The following parameters are specific to vLLM:
+
+- `max_model_len`: The new maximum sequence length after extension (original * factor).
+  Used for KV cache pre‑allocation and request limit at serving time.
@@ -375,8 +375,8 @@ For (G)B300, we recommend using CUDA 13, as shown in the following command.

    ```bash
    DOCKER_BUILDKIT=1 docker build \
-    --build-arg CUDA_VERSION=13.0.1 \
-    --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 \
+    --build-arg CUDA_VERSION=13.0.2 \
+    --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 \
    --build-arg max_jobs=256 \
    --build-arg nvcc_threads=2 \
    --build-arg RUN_WHEEL_CHECK=false \
@@ -78,7 +78,7 @@ The scoring models is designed to compute similarity scores between two input pr
 |-----------------------|---------------|----------------------------------------------|--------------------|--------------------------|
 | `classify` (see note) | Sequence-wise | reranker score for each sequence             | `cross-encoder`    | linear classifier        |
 | `embed`               | Sequence-wise | vector representations for each sequence     | `bi-encoder`       | cosine similarity        |
-| `token_classify`      | Token-wise    | probability vector of classes for each token | nan                | nan                      |
+| `token_classify`      | Token-wise    | probability vector of classes for each token | N/A                | N/A                      |
 | `token_embed`         | Token-wise    | vector representations for each token        | `late-interaction` | late interaction(MaxSim) |

 !!! note
@@ -86,14 +86,15 @@ The scoring models is designed to compute similarity scores between two input pr

 ### Pooling Usages

-| Pooling Usages              | Description                                                                                                                                             |
-|-----------------------------|---------------------------------------------------------------------------------------------------------------------------------------------------------|
-| Classification Usages       | Predicting which predefined category, class, or label best corresponds to a given input.                                                                |
-| Embedding Usages            | Converts unstructured data (text, images, audio, etc.) into structured numerical vectors (embeddings).                                                  |
-| Token Classification Usages | Token-wise classification                                                                                                                               |
-| Token Embedding Usages      | Token-wise embedding                                                                                                                                    |
-| Scoring Usages              | Computes similarity scores between two inputs. It supports three model types (aka `score_type`): `cross-encoder`, `late-interaction`, and `bi-encoder`. |
-| Reward Usages               | Evaluates the quality of outputs generated by a language model, acting as a proxy for human preferences.                                                |
+| Pooling Usages              | Description                                                                                                                                               |
+|-----------------------------|-----------------------------------------------------------------------------------------------------------------------------------------------------------|
+| Classification Usages       | Predicting which predefined category, class, or label best corresponds to a given input.                                                                  |
+| Embedding Usages            | Converts unstructured data (text, images, audio, etc.) into structured numerical vectors (embeddings).                                                    |
+| Token Classification Usages | Token-wise classification                                                                                                                                 |
+| Token Embedding Usages      | Token-wise embedding                                                                                                                                      |
+| Reward Usages               | Evaluates the quality of outputs generated by a language model, acting as a proxy for human preferences.                                                  |
+| Scoring Usages              | Computes similarity scores between two inputs. It supports three model types (aka `score_type`): `cross-encoder`, `late-interaction`, and `bi-encoder`.   |
+| Plugins Usages              | Allow users to customize input and output processors. For more information, please refer to [IO Processor Plugins](../../design/io_processor_plugins.md). |

 We also have some special models that support multiple pooling tasks, or have specific usage scenarios, or support special inputs and outputs.

@@ -101,9 +102,9 @@ For more detailed information, please refer to the link below.

 - [Classification Usages](classify.md)
 - [Embedding Usages](embed.md)
- [Reward Usages](reward.md)
 - [Token Classification Usages](token_classify.md)
 - [Token Embedding Usages](token_embed.md)
+- [Reward Usages](reward.md)
 - [Scoring Usages](scoring.md)
 - [Specific Model Examples](specific_models.md)

@@ -113,15 +114,17 @@ Each pooling model in vLLM supports one or more of these tasks according to
 [Pooler.get_supported_tasks][vllm.model_executor.layers.pooler.Pooler.get_supported_tasks],
 enabling the corresponding APIs.

-### Offline APIs corresponding to pooling tasks
+### Offline APIs corresponding to pooling usages

-| Task             | APIs                                                                                  |
-|------------------|---------------------------------------------------------------------------------------|
-| `embed`          | `LLM.embed(...)`, `LLM.encode(..., pooling_task="embed")`, `LLM.score(...)`(see note) |
-| `classify`       | `LLM.classify(...)`, `LLM.encode(..., pooling_task="classify")`, `LLM.score(...)`     |
-| `token_classify` | `LLM.reward(...)`, `LLM.encode(..., pooling_task="token_classify")`                   |
-| `token_embed`    | `LLM.encode(..., pooling_task="token_embed")`, `LLM.score(...)`                       |
-| `plugin`         | `LLM.encode(..., pooling_task="plugin")`                                              |
+| Pooling Usages              | Dedicated API       | Pooling task for `LLM.encode` API | Score Types                | scoring function         |
+|-----------------------------|---------------------|-----------------------------------|----------------------------|--------------------------|
+| Classification Usages       | `LLM.classify(...)` | `classify`                        | `cross-encoder` (see note) | linear classifier        |
+| Embedding Usages            | `LLM.embed(...)`    | `embed`                           | `bi-encoder`               | cosine similarity        |
+| Token Classification Usages | N/A                 | `token_classify`                  | N/A                        | N/A                      |
+| Token Embedding Usages      | N/A                 | `token_embed`                     | `late-interaction`         | late interaction(MaxSim) |
+| Reward Usages               | N/A                 | `classify` & `token_classify`     | N/A                        | N/A                      |
+| Scoring Usages              | `LLM.score(...)`    | N/A                               | N/A                        | N/A                      |
+| Plugins Usages              | N/A                 | `plugin`                          | N/A                        | N/A                      |

 !!! note
    Only when a classification model outputs num_labels equal to 1 can it be used as a scoring model and have its scoring API enabled.
@@ -147,7 +150,7 @@ It is primarily designed for [score models](scoring.md).

 The [encode][vllm.LLM.encode] method is available to all pooling models in vLLM.

-Please use one of the more specific methods or set the task directly when using `LLM.encode`, refer to the [table above](#offline-apis-corresponding-to-pooling-tasks).
+Please use one of the more specific methods or set the task directly when using `LLM.encode`, refer to the [table above](#offline-apis-corresponding-to-pooling-usages).

 ### Examples

@@ -183,9 +186,12 @@ Our Pooling API (`/pooling`) is similar to `LLM.encode`, being applicable to all

 The input format is the same as [Embeddings API](embed.md#openai-compatible-embeddings-api), but the output data can contain an arbitrary nested list, not just a 1-D list of floats.

-Please use one of the more specific APIs or set the task directly when using the Pooling API, refer to the [table above](#offline-apis-corresponding-to-pooling-tasks).
+Please use one of the more specific APIs or set the task directly when using the Pooling API, refer to the [table above](#offline-apis-corresponding-to-pooling-usages).

-Code example: [examples/pooling/pooling/pooling_online.py](../../../examples/pooling/pooling/pooling_online.py)
+Code examples:
+
+- [Online example](../../../examples/pooling/reward/token_reward_online.py)
+- [Offline example](../../../examples/pooling/reward/token_reward_offline.py)

 ### Examples

@@ -134,3 +134,13 @@ print(f"Data: {data!r}")
 ## Online Serving

 Please refer to the [pooling API](README.md#pooling-api). Pooling task corresponding to reward model types refer to the [table above](#summary).
+
+## More examples
+
+More examples can be found here: [examples/pooling/reward](../../../examples/pooling/reward)
+
+## Deprecated Features
+
+### `LLM.reward`
+
+`llm.reward` api is deprecated and will be removed in v0.23. Please use `LLM.encode` with `pooling_task="classify"` or `pooling_task="token_classify"` instead.
@@ -0,0 +1,62 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""
+Example offline usage of sequence reward models.
+
+The key distinction between sequence classification and token classification
+lies in their output granularity: sequence classification produces a single
+result for an entire input sequence, whereas token classification yields a
+result for each individual token within the sequence.
+"""
+
+from argparse import Namespace
+
+from vllm import LLM, EngineArgs
+from vllm.utils.argparse_utils import FlexibleArgumentParser
+from vllm.utils.print_utils import print_embeddings
+
+
+def parse_args():
+    parser = FlexibleArgumentParser()
+    parser = EngineArgs.add_cli_args(parser)
+    # Set example specific arguments
+    parser.set_defaults(
+        model="Skywork/Skywork-Reward-V2-Qwen3-0.6B",
+        runner="pooling",
+        enforce_eager=True,
+        max_model_len=1024,
+        trust_remote_code=True,
+    )
+    return parser.parse_args()
+
+
+def main(args: Namespace):
+    # Sample prompts.
+    prompts = [
+        "Hello, my name is",
+        "The president of the United States is",
+        "The capital of France is",
+        "The future of AI is",
+    ]
+
+    # Create an LLM.
+    # You should pass runner="pooling" for reward models
+    llm = LLM(**vars(args))
+
+    # Generate rewards. The output is a list of PoolingRequestOutput.
+    # Use pooling_task="classify" for sequence reward models.
+    outputs = llm.encode(prompts, pooling_task="classify")
+
+    # Print the outputs.
+    print("\nGenerated Outputs:\n" + "-" * 60)
+    for prompt, output in zip(prompts, outputs):
+        rewards = output.outputs.data
+        print(f"Prompt: {prompt!r}")
+        print_embeddings(rewards.tolist(), prefix="Reward")
+        print("-" * 60)
+
+
+if __name__ == "__main__":
+    args = parse_args()
+    main(args)
@@ -0,0 +1,71 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Example online usage of sequence reward models.
+
+Run `vllm serve <model> --runner pooling`
+to start up the server in vLLM. e.g.
+
+vllm serve Skywork/Skywork-Reward-V2-Qwen3-0.6B
+
+The key distinction between sequence classification and token classification
+lies in their output granularity: sequence classification produces a single
+result for an entire input sequence, whereas token classification yields a
+result for each individual token within the sequence.
+"""
+
+import argparse
+import pprint
+
+import requests
+
+
+def post_http_request(prompt: dict, api_url: str) -> requests.Response:
+    headers = {"User-Agent": "Test Client"}
+    response = requests.post(api_url, headers=headers, json=prompt)
+    return response
+
+
+def parse_args():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--host", type=str, default="localhost")
+    parser.add_argument("--port", type=int, default=8000)
+
+    return parser.parse_args()
+
+
+def main(args):
+    base_url = f"http://{args.host}:{args.port}"
+    models_url = base_url + "/v1/models"
+    pooing_url = base_url + "/pooling"
+
+    response = requests.get(models_url)
+    model = response.json()["data"][0]["id"]
+
+    # Input like Completions API
+    prompt = {"model": model, "input": "vLLM is great!"}
+    pooling_response = post_http_request(prompt=prompt, api_url=pooing_url)
+    print("-" * 50)
+    print("Pooling Response:")
+    pprint.pprint(pooling_response.json())
+    print("-" * 50)
+
+    # Input like Chat API
+    prompt = {
+        "model": model,
+        "messages": [
+            {
+                "role": "user",
+                "content": [{"type": "text", "text": "vLLM is great!"}],
+            }
+        ],
+    }
+    pooling_response = post_http_request(prompt=prompt, api_url=pooing_url)
+    print("Pooling Response:")
+    pprint.pprint(pooling_response.json())
+    print("-" * 50)
+
+
+if __name__ == "__main__":
+    args = parse_args()
+    main(args)
@@ -1,6 +1,15 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project

+"""
+Example offline usage of token reward models.
+
+The key distinction between sequence classification and token classification
+lies in their output granularity: sequence classification produces a single
+result for an entire input sequence, whereas token classification yields a
+result for each individual token within the sequence.
+"""
+
 from argparse import Namespace

 from vllm import LLM, EngineArgs
@@ -36,14 +45,14 @@ def main(args: Namespace):
    llm = LLM(**vars(args))

    # Generate rewards. The output is a list of PoolingRequestOutput.
-    outputs = llm.reward(prompts)
+    outputs = llm.encode(prompts, pooling_task="token_classify")

    # Print the outputs.
    print("\nGenerated Outputs:\n" + "-" * 60)
    for prompt, output in zip(prompts, outputs):
        rewards = output.outputs.data
        print(f"Prompt: {prompt!r}")
-        print_embeddings(rewards, prefix="Reward")
+        print_embeddings(rewards.tolist(), prefix="Reward")
        print("-" * 60)


@@ -1,12 +1,17 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 """
-Example online usage of Pooling API.
+Example online usage of token reward models.

 Run `vllm serve <model> --runner pooling`
 to start up the server in vLLM. e.g.

 vllm serve internlm/internlm2-1_8b-reward --trust-remote-code
+
+The key distinction between sequence classification and token classification
+lies in their output granularity: sequence classification produces a single
+result for an entire input sequence, whereas token classification yields a
+result for each individual token within the sequence.
 """

 import argparse
@@ -24,13 +24,14 @@ classifiers = [
    "Programming Language :: Python :: 3.11",
    "Programming Language :: Python :: 3.12",
    "Programming Language :: Python :: 3.13",
+    "Programming Language :: Python :: 3.14",
    "Intended Audience :: Developers",
    "Intended Audience :: Information Technology",
    "Intended Audience :: Science/Research",
    "Topic :: Scientific/Engineering :: Artificial Intelligence",
    "Topic :: Scientific/Engineering :: Information Analysis",
 ]
-requires-python = ">=3.10,<3.14"
+requires-python = ">=3.10,<3.15"
 dynamic = [ "version", "dependencies", "optional-dependencies"]

 [project.urls]
@@ -123,7 +124,8 @@ extend-exclude = ["tests/models/fixtures/*", "tests/prompts/*", "tests/tokenizer
    "benchmarks/sonnet.txt", "tests/lora/data/*", "build/*",
    "examples/pooling/token_embed/*", "tests/models/language/pooling/*",
    "vllm/third_party/*", "vllm/entrypoints/serve/instrumentator/static/*", "tests/entrypoints/openai/speech_to_text/test_transcription_validation.py",
-    "docs/governance/process.md", "tests/v1/engine/test_fast_incdec_prefix_err.py", ".git/*"]
+    "docs/governance/process.md", "docs/assets/contributing/vllm_bench_serve_timeline.html", 
+    "tests/v1/engine/test_fast_incdec_prefix_err.py", ".git/*"]
 ignore-hidden = false

 [tool.typos.default]
@@ -20,7 +20,7 @@ prometheus-fastapi-instrumentator >= 7.0.0
 tiktoken >= 0.6.0  # Required for DBRX tokenizer
 lm-format-enforcer == 0.11.3
 llguidance >= 1.3.0, < 1.4.0; platform_machine == "x86_64" or platform_machine == "arm64" or platform_machine == "aarch64" or platform_machine == "ppc64le"
-outlines_core == 0.2.11
+outlines_core == 0.2.14
 # required for outlines backend disk cache
 diskcache == 5.6.3
 lark == 1.2.2
@@ -4,7 +4,7 @@

 setuptools==77.0.3 # this version can reuse CMake build dir

-numba == 0.61.2; platform_machine != "s390x" # Required for N-gram speculative decoding
+numba == 0.65.0; platform_machine != "s390x" # Required for N-gram speculative decoding

 # Dependencies for CPUs
 torch==2.11.0+cpu; platform_machine == "x86_64" or platform_machine == "s390x" or platform_machine == "aarch64"
@@ -1,7 +1,7 @@
 # Common dependencies
 -r common.txt

-numba == 0.61.2 # Required for N-gram speculative decoding
+numba == 0.65.0 # Required for N-gram speculative decoding

 # Dependencies for NVIDIA GPUs
 torch==2.11.0
@@ -5,7 +5,7 @@
 grpcio==1.78.0
 grpcio-reflection==1.78.0

-numba == 0.61.2 # Required for N-gram speculative decoding
+numba == 0.65.0 # Required for N-gram speculative decoding

 # Dependencies for AMD GPUs
 datasets
@@ -20,4 +20,4 @@ conch-triton-kernels==1.2.1
 timm>=1.0.17
 # amd-quark: required for Quark quantization on ROCm 
 # To be consistent with test_quark.py
-amd-quark>=0.8.99
+amd-quark>=0.8.99
@@ -54,7 +54,7 @@ grpcio==1.78.0
 grpcio-reflection==1.78.0

 arctic-inference == 0.1.1 # Required for suffix decoding test
-numba == 0.61.2 # Required for N-gram speculative decoding
+numba == 0.65.0 # Required for N-gram speculative decoding
 numpy
 runai-model-streamer[s3,gcs,azure]==0.15.7
 fastsafetensors>=0.2.2 # 0.2.2 contains important fixes for multi-GPU mem usage
@@ -479,7 +479,7 @@ lightning-utilities==0.14.3
    #   lightning
    #   pytorch-lightning
    #   torchmetrics
-llvmlite==0.44.0
+llvmlite==0.47.0
    # via numba
 lm-eval==0.4.11
    # via -r requirements/test/cuda.in
@@ -550,7 +550,7 @@ nltk==3.9.1
    # via rouge-score
 num2words==0.5.14
    # via -r requirements/test/cuda.in
-numba==0.61.2
+numba==0.65.0
    # via
    #   -c requirements/cuda.txt
    #   -r requirements/test/cuda.in
@@ -40,7 +40,7 @@ buildkite-test-collector==0.1.9
 genai_perf>=0.0.8
 tritonclient>=2.51.0

-numba == 0.61.2 # Required for N-gram speculative decoding
+numba == 0.65.0 # Required for N-gram speculative decoding
 numpy
 runai-model-streamer[s3,gcs,azure]==0.15.7
 fastsafetensors>=0.2.2
@@ -52,7 +52,7 @@ grpcio==1.78.0
 grpcio-reflection==1.78.0

 arctic-inference==0.1.1 # Required for suffix decoding test
-numba==0.61.2 # Required for N-gram speculative decoding
+numba==0.65.0 # Required for N-gram speculative decoding
 numpy
 runai-model-streamer[s3,gcs,azure]==0.15.7
 fastsafetensors @ git+https://github.com/foundation-model-stack/fastsafetensors.git@0.2.2 # PyPI only ships CUDA wheels
@@ -559,7 +559,7 @@ llguidance==1.3.0
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
-llvmlite==0.44.0
+llvmlite==0.47.0
    # via numba
 lm-eval==0.4.11
    # via -r requirements/test/rocm.in
@@ -655,7 +655,7 @@ nltk==3.9.3
    # via rouge-score
 num2words==0.5.14
    # via -r requirements/test/rocm.in
-numba==0.61.2
+numba==0.65.0
    # via
    #   -c requirements/rocm.txt
    #   -r requirements/test/rocm.in
@@ -811,7 +811,7 @@ orjson==3.11.7
    # via
    #   genai-perf
    #   kaleido
-outlines-core==0.2.11
+outlines-core==0.2.14
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
@@ -244,7 +244,7 @@ lazy-loader==0.5
    #   scikit-image
 librosa==0.10.2.post1
    # via -r requirements/test/xpu.in
-llvmlite==0.44.0
+llvmlite==0.47.0
    # via numba
 lm-eval==0.4.11
    # via -r requirements/test/xpu.in
@@ -304,7 +304,7 @@ nltk==3.9.4
    # via rouge-score
 num2words==0.5.14
    # via -r requirements/test/xpu.in
-numba==0.61.2
+numba==0.65.0
    # via
    #   -c requirements/xpu.txt
    #   librosa
@@ -9,7 +9,7 @@ setuptools>=77.0.3,<81.0.0
 wheel
 jinja2>=3.1.6
 datasets # for benchmark scripts
-numba == 0.61.2 # Required for N-gram speculative decoding
+numba == 0.65.0 # Required for N-gram speculative decoding
 --extra-index-url=https://download.pytorch.org/whl/xpu
 torch==2.11.0+xpu
 torchaudio
@@ -927,7 +927,9 @@ def get_vllm_version() -> str:
    elif _is_tpu():
        version += f"{sep}tpu"
    elif _is_cpu():
-        if envs.VLLM_TARGET_DEVICE == "cpu":
+        # Check the local VLLM_TARGET_DEVICE (may be set by auto-detect above),
+        # not envs.VLLM_TARGET_DEVICE, so CPU-only hosts still get `+cpu`.
+        if VLLM_TARGET_DEVICE == "cpu":
            version += f"{sep}cpu"
    elif _is_xpu():
        version += f"{sep}xpu"
@@ -117,9 +117,9 @@ class TestFusedAddRMSNorm(torch.nn.Module):
        else:
            return norm_output, residual_output

-    def example_inputs(self, batch_size=8, hidden_size=16, seq_len=16):
-        hidden_states = torch.randn((batch_size * seq_len, hidden_size))
-        residual = torch.randn((batch_size * seq_len, hidden_size))
+    def example_inputs(self, batch_size=8, seq_len=16):
+        hidden_states = torch.randn((batch_size * seq_len, self.hidden_size))
+        residual = torch.randn((batch_size * seq_len, self.intermediate_size))
        return (hidden_states, residual)

    def ops_in_model(self, do_fusion):
@@ -10,7 +10,7 @@ import vllm.envs as envs
 from tests.compile.backend import TestBackend
 from tests.kernels.quantization.nvfp4_utils import quant_nvfp4_tensor
 from tests.utils import TestFP8Layer
-from vllm._aiter_ops import IS_AITER_FOUND
+from vllm._aiter_ops import IS_AITER_FOUND, rocm_aiter_ops
 from vllm._custom_ops import cutlass_scaled_fp4_mm, scaled_fp4_quant
 from vllm.compilation.passes.fusion.act_quant_fusion import (
    FUSED_OPS,
@@ -157,12 +157,13 @@ class TestSiluMulGroupFp8QuantModel(torch.nn.Module):
            activation_quant_key=self.act_quant_key,
            input_dtype=dtype,
        )
-        self.w = torch.rand(hidden_size, hidden_size).to(dtype=FP8_DTYPE).t()

-        scale_hidden_size = (hidden_size + 128 - 1) // 128
-        self.wscale = torch.rand(
-            (scale_hidden_size, scale_hidden_size), dtype=torch.float32
-        )
+        if not current_platform.is_fp8_fnuz():
+            kernel = self.w8a8_block_fp8_linear.kernel
+            orig_quant = kernel.quant_fp8
+            kernel.quant_fp8 = lambda *a, use_triton=False, **kw: orig_quant(
+                *a, use_triton=True, **kw
+            )

        self.enable_silu_mul_custom_op = self.silu_and_mul.enabled()

@@ -174,6 +175,9 @@ class TestSiluMulGroupFp8QuantModel(torch.nn.Module):
    def ops_in_model_before(self):
        return [
            SILU_MUL_OP if self.enable_silu_mul_custom_op else torch.ops.aten.mul,
+            rocm_aiter_ops.get_group_quant_op()
+            if current_platform.is_fp8_fnuz()
+            else torch.ops.vllm.triton_per_token_group_quant_fp8.default,
        ]

    def ops_in_model_after(self):
@@ -324,7 +328,6 @@ def test_fusion_silu_and_mul_quant(
    with set_current_vllm_config(config), monkeypatch.context() as m:
        fusion_passes = [ActivationQuantFusionPass(config)]
        if IS_AITER_FOUND and model_class is TestSiluMulGroupFp8QuantModel:
-            from vllm._aiter_ops import rocm_aiter_ops
            from vllm.compilation.passes.fusion.rocm_aiter_fusion import (
                RocmAiterSiluMulFp8GroupQuantFusionPass,
            )
@@ -352,10 +355,16 @@ def test_fusion_silu_and_mul_quant(
            atol, rtol = 1e-3, 1e-3
        elif isinstance(model, TestSiluMulNvfp4QuantModel):
            atol, rtol = 1e-1, 1e-1
-        elif isinstance(
-            model, (TestSiluMulGroupFp8QuantModel, TestSiluMulBlockQuantModel)
-        ):
+        elif isinstance(model, TestSiluMulGroupFp8QuantModel):
            atol, rtol = 5e-2, 5e-2
+        elif isinstance(model, TestSiluMulBlockQuantModel):
+            if current_platform.is_rocm():
+                atol, rtol = 1e-3, 1e-3
+            else:
+                # CUDA fused kernel computes silu*mul in fp32 while the reference
+                # goes through bf16/fp16 storage, so group maxima (and thus scales)
+                # can shift by one FP8-e4m3 code (~1/8 relative step).
+                atol, rtol = 5e-2, 5e-2

        torch.testing.assert_close(
            result[0].to(dtype=dtype), result2[0].to(dtype=dtype), atol=atol, rtol=rtol
@@ -205,6 +205,22 @@ def test_enforce_eager(vllm_runner, monkeypatch):
        pass


+@pytest.mark.forked
+def test_torch_compile_disable(vllm_runner, monkeypatch):
+    monkeypatch.setenv("VLLM_ENABLE_V1_MULTIPROCESSING", "0")
+    monkeypatch.setenv("TORCH_COMPILE_DISABLE", "1")
+    monkeypatch.setenv("VLLM_DISABLE_COMPILE_CACHE", "1")
+
+    with (
+        compilation_counter.expect(num_graphs_seen=0, stock_torch_compile_count=0),
+        vllm_runner(
+            "facebook/opt-125m",
+            gpu_memory_utilization=0.4,
+        ) as _,
+    ):
+        pass
+
+
 def test_splitting_ops_dynamic():
    # Default config
    config = VllmConfig()
@@ -1183,7 +1183,7 @@ class VllmRunner:
        return [req_output.outputs.data for req_output in req_outputs]

    def reward(self, prompts: list[str]) -> list[list[float]]:
-        req_outputs = self.llm.reward(prompts)
+        req_outputs = self.llm.encode(prompts, pooling_task="token_classify")
        return [req_output.outputs.data for req_output in req_outputs]

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

-import asyncio
 import random

 import pytest
@@ -361,16 +360,14 @@ def _test_async_transfer_layer_without_mtp_worker(
        communicator.set_stream(cuda_stream)

        for layer_idx in range(num_layers):
-            transfer_metadata = asyncio.run(
-                transfer_layer(
-                    old_layer_indices=old_indices_cpu[layer_idx],
-                    new_layer_indices=new_indices_cpu[layer_idx],
-                    expert_weights=expert_weights[layer_idx],
-                    expert_weights_buffer=expert_buffer,
-                    ep_group=ep_group,
-                    communicator=communicator,
-                    cuda_stream=cuda_stream,
-                )
+            transfer_metadata = transfer_layer(
+                old_layer_indices=old_indices_cpu[layer_idx],
+                new_layer_indices=new_indices_cpu[layer_idx],
+                expert_weights=expert_weights[layer_idx],
+                expert_weights_buffer=expert_buffer,
+                ep_group=ep_group,
+                communicator=communicator,
+                cuda_stream=cuda_stream,
            )
            cuda_stream.synchronize()
            move_from_buffer(
@@ -164,6 +164,58 @@ async def test_chat_error_non_stream():
        await serving_chat.create_chat_completion(request)


+@pytest.mark.asyncio
+async def test_openai_chat_keeps_mm_cache_for_engine_execution():
+    mock_engine = MagicMock(spec=AsyncLLM)
+    mock_engine.errored = False
+    mock_engine.model_config = MockModelConfig()
+    mock_engine.input_processor = MagicMock()
+    mock_engine.renderer = _build_renderer(mock_engine.model_config)
+
+    serving_chat = _build_serving_chat(mock_engine)
+
+    request = ChatCompletionRequest(
+        model=MODEL_NAME,
+        messages=[{"role": "user", "content": "Test prompt"}],
+    )
+
+    result = await serving_chat.render_chat_request(request)
+
+    assert isinstance(result, tuple)
+    assert (
+        serving_chat.openai_serving_render.preprocess_chat.call_args.kwargs[
+            "skip_mm_cache"
+        ]
+        is False
+    )
+
+
+@pytest.mark.asyncio
+async def test_renderer_only_chat_request_skips_mm_cache():
+    mock_engine = MagicMock(spec=AsyncLLM)
+    mock_engine.errored = False
+    mock_engine.model_config = MockModelConfig()
+    mock_engine.input_processor = MagicMock()
+    mock_engine.renderer = _build_renderer(mock_engine.model_config)
+
+    serving_chat = _build_serving_chat(mock_engine)
+
+    request = ChatCompletionRequest(
+        model=MODEL_NAME,
+        messages=[{"role": "user", "content": "Test prompt"}],
+    )
+
+    result = await serving_chat.openai_serving_render.render_chat_request(request)
+
+    assert result.token_ids == [1, 2, 3]
+    assert (
+        serving_chat.openai_serving_render.preprocess_chat.call_args.kwargs[
+            "skip_mm_cache"
+        ]
+        is True
+    )
+
+
@pytest.mark.asyncio
 async def test_chat_error_stream():
    """test finish_reason='error' returns 500 InternalServerError (streaming)"""
@@ -3,7 +3,7 @@

 from dataclasses import dataclass, field
 from typing import Any
-from unittest.mock import MagicMock
+from unittest.mock import AsyncMock, MagicMock

 import pytest

@@ -148,6 +148,66 @@ async def test_completion_error_non_stream():
        await serving_completion.create_completion(request)


+@pytest.mark.asyncio
+async def test_openai_completion_keeps_mm_cache_for_engine_execution():
+    mock_engine = MagicMock(spec=AsyncLLM)
+    mock_engine.errored = False
+    mock_engine.model_config = MockModelConfig()
+    mock_engine.input_processor = MagicMock()
+    mock_engine.renderer = _build_renderer(mock_engine.model_config)
+
+    serving_completion = _build_serving_completion(mock_engine)
+    serving_completion.openai_serving_render.preprocess_completion = AsyncMock(
+        return_value=[{"prompt_token_ids": [1, 2, 3]}]
+    )
+
+    request = CompletionRequest(
+        model=MODEL_NAME,
+        prompt="Test prompt",
+    )
+
+    result = await serving_completion.render_completion_request(request)
+
+    assert isinstance(result, list)
+    assert (
+        serving_completion.openai_serving_render.preprocess_completion.call_args.kwargs[
+            "skip_mm_cache"
+        ]
+        is False
+    )
+
+
+@pytest.mark.asyncio
+async def test_renderer_only_completion_request_skips_mm_cache():
+    mock_engine = MagicMock(spec=AsyncLLM)
+    mock_engine.errored = False
+    mock_engine.model_config = MockModelConfig()
+    mock_engine.input_processor = MagicMock()
+    mock_engine.renderer = _build_renderer(mock_engine.model_config)
+
+    serving_completion = _build_serving_completion(mock_engine)
+    serving_completion.openai_serving_render.preprocess_completion = AsyncMock(
+        return_value=[{"prompt_token_ids": [1, 2, 3]}]
+    )
+
+    request = CompletionRequest(
+        model=MODEL_NAME,
+        prompt="Test prompt",
+    )
+
+    result = await serving_completion.openai_serving_render.render_completion_request(
+        request
+    )
+
+    assert isinstance(result, list)
+    assert (
+        serving_completion.openai_serving_render.preprocess_completion.call_args.kwargs[
+            "skip_mm_cache"
+        ]
+        is True
+    )
+
+
@pytest.mark.asyncio
 async def test_completion_error_stream():
    """test finish_reason='error' returns 500 InternalServerError (streaming)"""
@@ -311,7 +311,7 @@ async def test_abort_timeout_exits_quickly(wait_for_engine_idle: float):
            pytest.fail("Process did not exit after SIGTERM with abort timeout")

        exit_time = time.time() - start_time
-        assert exit_time < 2, f"Default shutdown took too long: {exit_time:.1f}s"
+        assert exit_time < 2.1, f"Default shutdown took too long: {exit_time:.1f}s"
        assert proc.returncode in (0, -15, None), f"Unexpected: {proc.returncode}"

        await _assert_children_cleaned_up(child_pids)
@@ -12,7 +12,10 @@ from vllm.config.multimodal import MultiModalConfig
 from vllm.entrypoints.openai.engine.protocol import StreamOptions
 from vllm.entrypoints.openai.models.protocol import BaseModelPath
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
-from vllm.entrypoints.serve.disagg.protocol import GenerateRequest
+from vllm.entrypoints.serve.disagg.protocol import (
+    GenerateRequest,
+    GenerateResponse,
+)
 from vllm.entrypoints.serve.disagg.serving import ServingTokens
 from vllm.entrypoints.serve.render.serving import OpenAIServingRender
 from vllm.logprobs import Logprob
@@ -164,6 +167,36 @@ def _parse_sse_chunks(chunks: list[str]) -> list[Any]:
    return parsed


+@pytest.mark.asyncio
+async def test_serve_tokens_skips_mm_cache_for_remote_engine_execution():
+    engine = _mock_engine()
+
+    async def mock_generate(*args, **kwargs):
+        yield _make_request_output(
+            "req-1", token_ids=[10], finish_reason="stop", finished=True
+        )
+
+    engine.generate = MagicMock(side_effect=mock_generate)
+    serving = _build_serving_tokens(engine)
+
+    request = GenerateRequest(
+        token_ids=[1, 2, 3],
+        sampling_params=SamplingParams(max_tokens=1),
+        model=MODEL_NAME,
+        stream=False,
+    )
+
+    response = await serving.serve_tokens(request)
+
+    assert isinstance(response, GenerateResponse)
+    assert (
+        serving.openai_serving_render.preprocess_completion.call_args.kwargs[
+            "skip_mm_cache"
+        ]
+        is True
+    )
+
+
@pytest.mark.asyncio
 async def test_stream_basic():
    """Streaming returns SSE chunks with correct token_ids and ends with [DONE]."""
@@ -0,0 +1,140 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+from dataclasses import dataclass, field
+from typing import Any
+from unittest.mock import AsyncMock, MagicMock
+
+import pytest
+
+from vllm.config.multimodal import MultiModalConfig
+from vllm.entrypoints.openai.models.protocol import BaseModelPath
+from vllm.entrypoints.openai.models.serving import OpenAIServingModels
+from vllm.entrypoints.serve.render.serving import OpenAIServingRender
+from vllm.entrypoints.serve.tokenize.protocol import (
+    TokenizeChatRequest,
+    TokenizeCompletionRequest,
+)
+from vllm.entrypoints.serve.tokenize.serving import OpenAIServingTokenization
+from vllm.v1.engine.async_llm import AsyncLLM
+
+MODEL_NAME = "openai-community/gpt2"
+BASE_MODEL_PATHS = [
+    BaseModelPath(name=MODEL_NAME, model_path=MODEL_NAME),
+]
+
+
+@dataclass
+class MockHFConfig:
+    model_type: str = "any"
+
+
+@dataclass
+class MockModelConfig:
+    task = "generate"
+    runner_type = "generate"
+    model = MODEL_NAME
+    tokenizer = MODEL_NAME
+    trust_remote_code = False
+    tokenizer_mode = "auto"
+    max_model_len = 100
+    tokenizer_revision = None
+    multimodal_config = MultiModalConfig()
+    hf_config = MockHFConfig()
+    hf_text_config = MockHFConfig()
+    logits_processors: list[str] | None = None
+    diff_sampling_param: dict | None = None
+    allowed_local_media_path: str = ""
+    allowed_media_domains: list[str] | None = None
+    encoder_config = None
+    generation_config: str = "auto"
+    media_io_kwargs: dict[str, dict[str, Any]] = field(default_factory=dict)
+    skip_tokenizer_init = False
+    is_encoder_decoder: bool = False
+    is_multimodal_model: bool = False
+    renderer_num_workers: int = 1
+
+    def get_diff_sampling_param(self):
+        return self.diff_sampling_param or {}
+
+
+def _build_serving_tokenization(engine: AsyncLLM) -> OpenAIServingTokenization:
+    models = OpenAIServingModels(
+        engine_client=engine,
+        base_model_paths=BASE_MODEL_PATHS,
+    )
+    serving_render = OpenAIServingRender(
+        model_config=engine.model_config,
+        renderer=engine.renderer,
+        model_registry=models.registry,
+        request_logger=None,
+        chat_template=None,
+        chat_template_content_format="auto",
+    )
+    return OpenAIServingTokenization(
+        engine,
+        models,
+        openai_serving_render=serving_render,
+        request_logger=None,
+        chat_template=None,
+        chat_template_content_format="auto",
+    )
+
+
+@pytest.mark.asyncio
+async def test_tokenize_chat_skips_mm_cache_for_renderer_only_path():
+    mock_engine = MagicMock(spec=AsyncLLM)
+    mock_engine.errored = False
+    mock_engine.model_config = MockModelConfig()
+    mock_engine.input_processor = MagicMock()
+    mock_engine.renderer = MagicMock()
+
+    serving = _build_serving_tokenization(mock_engine)
+    serving.openai_serving_render.preprocess_chat = AsyncMock(
+        return_value=(
+            [{"role": "user", "content": "Test"}],
+            [{"prompt_token_ids": [1, 2, 3]}],
+        )
+    )
+
+    request = TokenizeChatRequest(
+        model=MODEL_NAME,
+        messages=[{"role": "user", "content": "Test prompt"}],
+    )
+
+    response = await serving.create_tokenize(request, MagicMock(headers={}))
+
+    assert response.tokens == [1, 2, 3]
+    assert (
+        serving.openai_serving_render.preprocess_chat.call_args.kwargs["skip_mm_cache"]
+        is True
+    )
+
+
+@pytest.mark.asyncio
+async def test_tokenize_completion_skips_mm_cache_for_renderer_only_path():
+    mock_engine = MagicMock(spec=AsyncLLM)
+    mock_engine.errored = False
+    mock_engine.model_config = MockModelConfig()
+    mock_engine.input_processor = MagicMock()
+    mock_engine.renderer = MagicMock()
+
+    serving = _build_serving_tokenization(mock_engine)
+    serving.openai_serving_render.preprocess_completion = AsyncMock(
+        return_value=[{"prompt_token_ids": [1, 2, 3]}]
+    )
+
+    request = TokenizeCompletionRequest(
+        model=MODEL_NAME,
+        prompt="Test prompt",
+    )
+
+    response = await serving.create_tokenize(request, MagicMock(headers={}))
+
+    assert response.tokens == [1, 2, 3]
+    assert (
+        serving.openai_serving_render.preprocess_completion.call_args.kwargs[
+            "skip_mm_cache"
+        ]
+        is True
+    )
@@ -223,7 +223,7 @@ if has_deep_ep() and not current_platform.has_device_capability(100):
    )

 if has_mori():
-    from vllm.model_executor.layers.fused_moe.mori_prepare_finalize import (
+    from vllm.model_executor.layers.fused_moe.prepare_finalize.mori import (
        MoriPrepareAndFinalize,
    )

@@ -367,7 +367,9 @@ else:
    CutlassExpertsFp8 = None

 if cutlass_fp4_supported():
-    from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp4
+    from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import (
+        CutlassExpertsFp4,
+    )

    register_experts(
        CutlassExpertsFp4,
@@ -10,10 +10,12 @@ from vllm.model_executor.layers.fused_moe.experts.batched_deep_gemm_moe import (
    BatchedDeepGemmExperts,
 )
 from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
-    BatchedPrepareAndFinalize,
    BatchedTritonExperts,
 )
 from vllm.model_executor.layers.fused_moe.modular_kernel import FusedMoEKernel
+from vllm.model_executor.layers.fused_moe.prepare_finalize.batched import (
+    BatchedPrepareAndFinalize,
+)
 from vllm.utils.deep_gemm import calc_diff, is_deep_gemm_supported

 from .test_deepgemm import make_block_quant_fp8_weights
@@ -21,7 +21,7 @@ from vllm.model_executor.layers.fused_moe.config import (
    FusedMoEQuantConfig,
    fp8_w8a8_moe_quant_config,
 )
-from vllm.model_executor.layers.fused_moe.cutlass_moe import (
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import (
    CutlassExpertsFp8,
    run_cutlass_moe_fp8,
 )
@@ -19,7 +19,7 @@ from vllm.model_executor.layers.fused_moe.all2all_utils import (
    maybe_make_prepare_finalize,
 )
 from vllm.model_executor.layers.fused_moe.config import nvfp4_moe_quant_config
-from vllm.model_executor.layers.fused_moe.cutlass_moe import (
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import (
    CutlassExpertsFp4,
 )
 from vllm.model_executor.layers.fused_moe.prepare_finalize import (
@@ -662,6 +662,52 @@ def test_eplb_map_no_redundancy(
        assert load.sum().item() == 0


+@pytest.mark.parametrize("top_k,R", [(2, 2), (4, 2), (8, 4), (8, 8)])
+def test_eplb_map_hot_expert_replica_balance(top_k, R):
+    """Hot logical expert with R replicas must be balanced across replicas
+    even when ``top_k`` is a multiple of ``R``. In that regime every top-k
+    offset for the hot expert lands on a multiple of ``top_k`` in the flat
+    ``topk_ids`` view, so per-replica assignment must not collapse onto a
+    single replica.
+    """
+    num_tokens = 8192
+    num_logical = 16
+    num_physical = R + (num_logical - 1)
+
+    l2p = torch.full((num_logical, R), -1, dtype=torch.int64, device="cuda")
+    l2p[0] = torch.arange(R, dtype=torch.int64, device="cuda")
+    for i in range(1, num_logical):
+        l2p[i, 0] = R + i - 1
+    rc = torch.tensor([R] + [1] * (num_logical - 1), dtype=torch.int64, device="cuda")
+
+    torch.manual_seed(0)
+    topk_ids = torch.randint(
+        1,
+        num_logical,
+        (num_tokens, top_k),
+        dtype=torch.int32,
+        device="cuda",
+    )
+    topk_ids[:, 0] = 0
+
+    load = torch.zeros(num_physical, dtype=torch.int32, device="cuda")
+    rec = torch.tensor(True, dtype=torch.bool, device="cuda")
+
+    eplb_map_to_physical_and_record(
+        topk_ids=topk_ids,
+        expert_load_view=load,
+        logical_to_physical_map=l2p,
+        logical_replica_count=rc,
+        record_enabled=rec,
+    )
+
+    hot_load = load[:R].float()
+    max_mean = (hot_load.max() / hot_load.mean()).item()
+    assert max_mean < 1.15, (
+        f"Hot expert replicas uneven: {hot_load.tolist()}, max/mean={max_mean:.3f}"
+    )
+
+
@pytest.mark.parametrize("record_enabled", [True, False])
@pytest.mark.parametrize(
    "l2p_map, replica_count, num_physical, topk_ids, expected_out, expected_load",
@@ -672,10 +718,12 @@ def test_eplb_map_no_redundancy(
            [2, 2, 1, 1],
            6,
            [[0, 1], [2, 3], [0, 2]],
-            # offs: 0→0%2=0→p0, 1→1%2=1→p5, 2→2%1=0→p2,
-            #        3→3%1=0→p3, 4→4%2=0→p0, 5→5%1=0→p2
-            [[0, 5], [2, 3], [0, 2]],
-            [2, 0, 2, 1, 0, 1],
+            # replica = (token_idx * KNUTH) & 0xFFFFFFFF % R.
+            # token 0 hash=0x00000000: %2=0, %1=0.
+            # token 1 hash=0x9E3779B9: %2=1, %1=0.
+            # token 2 hash=0x3C6EF372: %2=0, %1=0.
+            [[0, 1], [2, 3], [0, 2]],
+            [2, 1, 2, 1, 0, 0],
            id="partial",
        ),
        pytest.param(
@@ -684,10 +732,11 @@ def test_eplb_map_no_redundancy(
            [2, 2, 2, 2],
            8,
            [[0, 1], [2, 3], [0, 2]],
-            # offs: 0→0%2=0→p0, 1→1%2=1→p5, 2→2%2=0→p2,
-            #        3→3%2=1→p7, 4→4%2=0→p0, 5→5%2=1→p6
-            [[0, 5], [2, 7], [0, 6]],
-            [2, 0, 1, 0, 0, 1, 1, 1],
+            # token 0 hash=0x00000000: %2=0.
+            # token 1 hash=0x9E3779B9: %2=1.
+            # token 2 hash=0x3C6EF372: %2=0.
+            [[0, 1], [6, 7], [0, 2]],
+            [2, 1, 1, 0, 0, 0, 1, 1],
            id="full",
        ),
        pytest.param(
@@ -696,10 +745,11 @@ def test_eplb_map_no_redundancy(
            [4, 2, 2],
            8,
            [[0, 1], [2, 0], [1, 2]],
-            # offs: 0→0%4=0→p0, 1→1%2=1→p4, 2→2%2=0→p2,
-            #        3→3%4=3→p7, 4→4%2=0→p1, 5→5%2=1→p6
-            [[0, 4], [2, 7], [1, 6]],
-            [1, 1, 1, 0, 1, 0, 1, 1],
+            # token 0 hash=0x00000000: %4=0, %2=0.
+            # token 1 hash=0x9E3779B9: %4=1, %2=1.
+            # token 2 hash=0x3C6EF372: %4=2, %2=0.
+            [[0, 1], [6, 3], [1, 2]],
+            [1, 2, 1, 1, 0, 0, 1, 0],
            id="uneven",
        ),
    ],
@@ -18,7 +18,6 @@ from vllm.model_executor.layers.fused_moe.config import (
    RoutingMethodType,
 )
 from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
-    BatchedPrepareAndFinalize,
    BatchedTritonExperts,
    NaiveBatchedExperts,
 )
@@ -27,6 +26,9 @@ from vllm.model_executor.layers.fused_moe.fused_moe import (
    fused_experts,
 )
 from vllm.model_executor.layers.fused_moe.modular_kernel import FusedMoEKernel
+from vllm.model_executor.layers.fused_moe.prepare_finalize.batched import (
+    BatchedPrepareAndFinalize,
+)
 from vllm.model_executor.layers.fused_moe.router.fused_topk_router import fused_topk
 from vllm.model_executor.layers.fused_moe.utils import moe_kernel_quantize_input
 from vllm.utils.deep_gemm import per_block_cast_to_fp8
@@ -125,7 +125,7 @@ def test_gpt_oss_lora_tp2(
            enable_lora=True,
            max_loras=2,
            max_num_seqs=2,
-            max_num_batched_tokens=2048,
+            max_num_batched_tokens=1024,
            tensor_parallel_size=2,
            gpu_memory_utilization=0.8,
            fully_sharded_loras=fully_sharded_loras,
@@ -921,6 +921,7 @@ VLM_TEST_SETTINGS = {
        multi_image_prompt="Picture 1: <vlm_image>\nPicture 2: <vlm_image>\nDescribe these two images with one paragraph respectively.",  # noqa: E501
        max_model_len=4096,
        max_num_seqs=2,
+        num_logprobs=10,
        auto_cls=AutoModelForImageTextToText,
        vllm_output_post_proc=model_utils.qwen2_vllm_to_hf_output,
        image_size_factors=[(0.25,), (0.25, 0.25, 0.25), (0.25, 0.2, 0.15)],
@@ -0,0 +1,166 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+from dataclasses import dataclass, field
+
+import pytest
+
+from vllm.multimodal.video import sample_frames_from_video
+from vllm.platforms import current_platform
+
+from ....conftest import IMAGE_ASSETS, VIDEO_ASSETS
+from ....utils import create_new_process_for_each_test
+from .vlm_utils.builders import sample_frames_with_video_metadata
+
+
+@dataclass
+class VitCudagraphTestConfig:
+    model: str
+    modalities: list[str] = field(default_factory=lambda: ["image", "video"])
+    image_prompt: str | None = None
+    video_prompt: str | None = None
+    dtype: str = "bfloat16"
+    max_model_len: int = 4096
+    max_tokens: int = 64
+    max_num_seqs: int = 2
+    num_video_frames: int = 16
+    needs_video_metadata: bool = False
+    vllm_runner_kwargs: dict = field(default_factory=dict)
+    marks: list = field(default_factory=list)
+
+
+def params_with_marks(
+    configs: dict[str, VitCudagraphTestConfig],
+) -> list[pytest.param]:
+    return [
+        pytest.param(model_id, marks=cfg.marks) for model_id, cfg in configs.items()
+    ]
+
+
+def qwen_vl_chat_template(content: str) -> str:
+    return f"<|im_start|>user\n{content}<|im_end|>\n<|im_start|>assistant\n"
+
+
+MODEL_CONFIGS: dict[str, VitCudagraphTestConfig] = {
+    "qwen3_vl": VitCudagraphTestConfig(
+        model="Qwen/Qwen3-VL-2B-Instruct",
+        image_prompt=qwen_vl_chat_template(
+            "<|vision_start|><|image_pad|><|vision_end|>What is in this image?"
+        ),
+        video_prompt=qwen_vl_chat_template(
+            "<|vision_start|><|video_pad|><|vision_end|>"
+            "Describe this video in one sentence."
+        ),
+        needs_video_metadata=True,
+        marks=[pytest.mark.core_model],
+    ),
+    # TODO: Add more models below.
+}
+
+
+def get_compilation_config():
+    return {
+        "cudagraph_mm_encoder": True,
+        "encoder_cudagraph_max_vision_items_per_batch": 1,
+        "encoder_cudagraph_max_frames_per_batch": 16,
+    }
+
+
+# ---------------------------------------------------------------------------
+# Tests
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize("model_id", params_with_marks(MODEL_CONFIGS))
+@pytest.mark.skipif(not current_platform.is_cuda(), reason="Requires CUDA")
+@create_new_process_for_each_test()
+def test_vit_cudagraph_image(model_id, vllm_runner, image_assets):
+    config = MODEL_CONFIGS[model_id]
+
+    if "image" not in config.modalities:
+        pytest.skip(f"{model_id} does not support the image modality.")
+
+    image_prompts = IMAGE_ASSETS.prompts(
+        {
+            "stop_sign": config.image_prompt,  # type: ignore[typeddict-item]
+            "cherry_blossom": config.image_prompt,  # type: ignore[typeddict-item]
+        }
+    )
+    images = [[asset.pil_image] for asset in image_assets]
+
+    with vllm_runner(
+        config.model,
+        dtype=config.dtype,
+        max_model_len=config.max_model_len,
+        max_num_seqs=config.max_num_seqs,
+        limit_mm_per_prompt={"image": 1},
+        compilation_config=get_compilation_config(),
+        **config.vllm_runner_kwargs,
+    ) as vllm_model:
+        outputs = vllm_model.generate_greedy(
+            image_prompts, config.max_tokens, images=images
+        )
+
+        # Basic validation that we got a response
+        assert len(outputs) == 2
+        output_ids, output_text = outputs[0]
+
+        # Ensure we got some output
+        assert len(output_ids) > 0
+        assert len(output_text) > 0
+
+        # Ensure the output is a string
+        assert isinstance(output_text, str)
+
+
+@pytest.mark.parametrize("model_id", params_with_marks(MODEL_CONFIGS))
+@pytest.mark.skipif(not current_platform.is_cuda(), reason="Requires CUDA")
+@create_new_process_for_each_test()
+def test_vit_cudagraph_video(model_id, vllm_runner, video_assets):
+    config = MODEL_CONFIGS[model_id]
+
+    if "video" not in config.modalities:
+        pytest.skip(f"{model_id} does not support the video modality")
+
+    video_prompts = VIDEO_ASSETS.prompts(
+        {
+            "baby_reading": config.video_prompt,  # type: ignore[typeddict-item]
+        }
+    )
+    if config.needs_video_metadata:
+        sampled_vids = [
+            sample_frames_with_video_metadata(
+                (asset.np_ndarrays, asset.metadata), config.num_video_frames
+            )
+            for asset in video_assets
+        ]
+    else:
+        sampled_vids = [
+            sample_frames_from_video(asset.np_ndarrays, config.num_video_frames)
+            for asset in video_assets
+        ]
+    videos = [sampled_vids[0]]
+
+    with vllm_runner(
+        config.model,
+        dtype=config.dtype,
+        max_model_len=config.max_model_len,
+        max_num_seqs=config.max_num_seqs,
+        limit_mm_per_prompt={"video": 1},
+        compilation_config=get_compilation_config(),
+        **config.vllm_runner_kwargs,
+    ) as vllm_model:
+        outputs = vllm_model.generate_greedy(
+            video_prompts, config.max_tokens, videos=videos
+        )
+
+        # Basic validation that we got a response
+        assert len(outputs) == 1
+        output_ids, output_text = outputs[0]
+
+        # Ensure we got some output
+        assert len(output_ids) > 0
+        assert len(output_text) > 0
+
+        # Ensure the output is a string
+        assert isinstance(output_text, str)
@@ -4,6 +4,7 @@ import pytest

 from tests.reasoning.utils import run_reasoning_extraction
 from vllm.reasoning import ReasoningParser, ReasoningParserManager
+from vllm.reasoning.hy_v3_reasoning_parser import HYV3ReasoningParser
 from vllm.tokenizers import get_tokenizer

 parser_name = "hy_v3"
@@ -241,3 +242,33 @@ def test_is_reasoning_end_full_prompt(
    token_ids = hy_v3_tokenizer.convert_tokens_to_ids(tokens)
    check_is_reasoning_end = parser.is_reasoning_end(token_ids)
    assert check_is_reasoning_end == is_reasoning_end
+
+
+def test_constructor_does_not_mutate_shared_chat_template_kwargs(hy_v3_tokenizer):
+    parser_cls = ReasoningParserManager.get_reasoning_parser(parser_name)
+    chat_template_kwargs = {"reasoning_effort": "low"}
+
+    first_parser: ReasoningParser = parser_cls(
+        hy_v3_tokenizer,
+        chat_template_kwargs=chat_template_kwargs,
+    )
+    second_parser: ReasoningParser = parser_cls(
+        hy_v3_tokenizer,
+        chat_template_kwargs=chat_template_kwargs,
+    )
+
+    assert chat_template_kwargs == {"reasoning_effort": "low"}
+    assert isinstance(first_parser, HYV3ReasoningParser)
+    assert isinstance(second_parser, HYV3ReasoningParser)
+    assert first_parser._identity_parser is None
+    assert second_parser._identity_parser is None
+
+
+def test_constructor_falls_back_to_outer_reasoning_effort(hy_v3_tokenizer):
+    parser: ReasoningParser = ReasoningParserManager.get_reasoning_parser(parser_name)(
+        hy_v3_tokenizer,
+        reasoning_effort="low",
+    )
+
+    assert isinstance(parser, HYV3ReasoningParser)
+    assert parser._identity_parser is None
@@ -78,6 +78,25 @@ WITHOUT_THINK_STREAM = {
    "content": None,
 }

+# --- <tool_call> without </think> (implicit reasoning end) ---
+
+TOOL_CALL_BODY = (
+    "<tool_call>\n<function=bash>\n<parameter=command>"
+    "\ncat /etc/hosts\n</parameter>\n</function>\n</tool_call>"
+)
+
+TOOL_CALL_NO_THINK_END = {
+    "output": "I need to read the file.\n\n" + TOOL_CALL_BODY,
+    "reasoning": "I need to read the file.\n\n",
+    "content": TOOL_CALL_BODY,
+}
+
+TOOL_CALL_WITH_THINK_NO_END = {
+    "output": "<think>I need to read the file.\n\n" + TOOL_CALL_BODY,
+    "reasoning": "I need to read the file.\n\n",
+    "content": TOOL_CALL_BODY,
+}
+
 # --- Edge cases ---

 COMPLETE_REASONING = {
@@ -199,6 +218,26 @@ TEST_CASES = [
        TRUNCATED_NO_START_TOKEN_STREAM,
        id="truncated_no_start_token_stream",
    ),
+    pytest.param(
+        False,
+        TOOL_CALL_NO_THINK_END,
+        id="tool_call_no_think_end",
+    ),
+    pytest.param(
+        True,
+        TOOL_CALL_NO_THINK_END,
+        id="tool_call_no_think_end_stream",
+    ),
+    pytest.param(
+        False,
+        TOOL_CALL_WITH_THINK_NO_END,
+        id="tool_call_with_think_no_end",
+    ),
+    pytest.param(
+        True,
+        TOOL_CALL_WITH_THINK_NO_END,
+        id="tool_call_with_think_no_end_stream",
+    ),
 ]


@@ -255,6 +294,13 @@ MULTI_TOKEN_DELTA_CASES = [
        "content",
        id="no_start_end_grouped_with_content",
    ),
+    pytest.param(
+        # <tool_call> arrives in a separate delta after reasoning text
+        ["I need to read the file.\n\n", "<tool_call>\n<function=bash>"],
+        "I need to read the file.\n\n",
+        "<tool_call>\n<function=bash>",
+        id="tool_call_implicit_reasoning_end",
+    ),
 ]


@@ -296,6 +342,12 @@ THINKING_DISABLED_CASES = [
        "Some output without think tokens",
        id="thinking_disabled_no_think_tokens",
    ),
+    pytest.param(
+        "I need to read the file.\n\n" + TOOL_CALL_BODY,
+        None,
+        "I need to read the file.\n\n" + TOOL_CALL_BODY,
+        id="thinking_disabled_with_tool_call",
+    ),
 ]


@@ -34,6 +34,14 @@ def _make_renderer_mock(mm_limits: dict[str, int]) -> MagicMock:
    mm_processor = MagicMock()
    mm_processor.info.allowed_mm_limits = mm_limits
    renderer.mm_processor = mm_processor
+    renderer._readonly_mm_processor = None
+    renderer._warmup_mm_processor = BaseRenderer._warmup_mm_processor.__get__(
+        renderer, BaseRenderer
+    )
+    renderer._clear_processor_cache = BaseRenderer._clear_processor_cache
+    renderer.clear_mm_cache = MagicMock()
+    renderer.model_config.max_model_len = 128
+    renderer.model_config.get_multimodal_config.return_value.limit_per_prompt = {}

    return renderer

@@ -109,3 +117,19 @@ class TestMmWarmupSkippedWhenNoProcessor:
        BaseRenderer.warmup(renderer, ChatParams())

        renderer.model_config.get_multimodal_config.assert_not_called()
+
+
+class TestReadonlyMmWarmup:
+    """Readonly MM processor warmup must mirror the render path behavior."""
+
+    def test_readonly_processor_apply_called_and_cache_cleared(self):
+        renderer = _make_renderer_mock({"image": 1})
+        readonly_mm_processor = MagicMock()
+        readonly_mm_processor.info.allowed_mm_limits = {"image": 1}
+        renderer._readonly_mm_processor = readonly_mm_processor
+
+        with patch("vllm.multimodal.processing.TimingContext", autospec=True):
+            BaseRenderer.warmup(renderer, ChatParams())
+
+        readonly_mm_processor.apply.assert_called_once()
+        readonly_mm_processor.cache.clear_cache.assert_called_once()
@@ -4,15 +4,22 @@
 from unittest.mock import MagicMock, patch

 import pytest
+from transformers import AutoTokenizer

 from vllm.entrypoints.openai.engine.protocol import ExtractedToolCallInformation
-from vllm.tokenizers import TokenizerLike
 from vllm.tool_parsers.llama_tool_parser import Llama3JsonToolParser

+LLAMA_MODEL = "meta-llama/Llama-3.2-1B-Instruct"
+
+
+@pytest.fixture(scope="module")
+def llama_tokenizer():
+    return AutoTokenizer.from_pretrained(LLAMA_MODEL)
+

@pytest.fixture
-def parser(default_tokenizer: TokenizerLike):
-    return Llama3JsonToolParser(default_tokenizer)
+def parser(llama_tokenizer):
+    return Llama3JsonToolParser(llama_tokenizer)


 def test_extract_tool_calls_simple(parser):
@@ -107,6 +107,7 @@ def create_common_attn_metadata(
        query_start_loc=query_start_loc,
        query_start_loc_cpu=query_start_loc_cpu,
        seq_lens=seq_lens,
+        seq_lens_cpu_upper_bound=seq_lens_cpu,
        _seq_lens_cpu=seq_lens_cpu,
        _num_computed_tokens_cpu=num_computed_tokens_cpu,
        num_reqs=batch_spec.batch_size,
@@ -1310,6 +1310,54 @@ def test_dflash_acceptance_rates(dflash_config):
    cleanup_dist_env_and_memory()


+@single_gpu_only
+def test_synthetic_acceptance_rate():
+    """Verify that synthetic rejection sampling produces an acceptance
+    length close to the requested mean acceptance length."""
+    num_spec_tokens = 3
+    expected_acceptance_len = 1.875
+    tolerance = 0.15
+
+    spec_llm = LLM(
+        model="meta-llama/Llama-3.2-1B-Instruct",
+        trust_remote_code=True,
+        speculative_config={
+            "method": "eagle3",
+            "model": "nm-testing/Llama3_2_1B_speculator.eagle3",
+            "num_speculative_tokens": num_spec_tokens,
+            "max_model_len": 2048,
+            "rejection_sample_method": "synthetic",
+            "synthetic_acceptance_length": expected_acceptance_len,
+        },
+        max_model_len=2048,
+        enforce_eager=True,
+        disable_log_stats=False,
+    )
+
+    test_prompts = get_test_prompts(mm_enabled=False, num_prompts=50)
+    spec_llm.chat(
+        test_prompts,
+        SamplingParams(temperature=0, max_tokens=64, ignore_eos=True),
+    )
+
+    metrics = spec_llm.get_metrics()
+    acceptance_len = compute_acceptance_len(metrics)
+
+    print(
+        f"Synthetic acceptance length: {acceptance_len:.3f}"
+        f" (expected={expected_acceptance_len:.3f},"
+        f" tolerance=±{tolerance})"
+    )
+    assert abs(acceptance_len - expected_acceptance_len) <= tolerance, (
+        f"Synthetic acceptance length {acceptance_len:.3f} is not within"
+        f" ±{tolerance} of expected {expected_acceptance_len:.3f}"
+    )
+
+    del spec_llm
+    torch.accelerator.empty_cache()
+    cleanup_dist_env_and_memory()
+
+
 def test_dflash_correctness(dflash_config):
    """
    E2E test for DFlash (block diffusion) speculative decoding.
@@ -0,0 +1,222 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Verify that GPU memory is fully released after RixlConnector shutdown on ROCm.
+
+Regression test for ROCm/ucx#33: UCX rocm_ipc transport permanently pinned
+GPU memory via hsa_amd_ipc_memory_create during ucp_mem_map, causing
+GPU memory to be unrecoverable after engine shutdown.
+"""
+
+import gc
+
+import pytest
+import torch
+
+from vllm.platforms import current_platform
+
+pytestmark = pytest.mark.skipif(
+    not current_platform.is_rocm(),
+    reason="ROCm platform required",
+)
+
+
+def _mb(b: int) -> float:
+    return b / (1024 * 1024)
+
+
+def _gpu_snapshot(tag: str, prev_alloc: float = 0.0) -> dict:
+    """Print and return current GPU memory stats."""
+    torch.accelerator.synchronize()
+    alloc = torch.accelerator.memory_allocated()
+    reserved = torch.accelerator.memory_reserved()
+    # mem_get_info is not available on torch.accelerator
+    try:
+        drv_free, drv_total = torch.cuda.mem_get_info()
+        drv_used = drv_total - drv_free
+        drv_pct = drv_used / drv_total * 100
+    except Exception:
+        drv_used = drv_total = drv_pct = 0
+    alloc_mb = _mb(alloc)
+    drv_used_mb = _mb(drv_used)
+    delta = alloc_mb - prev_alloc
+    print(
+        f"  {tag:<40s} | {alloc_mb:>9.1f} alloc | "
+        f"{_mb(reserved):>9.1f} rsrvd | "
+        f"{drv_used_mb:>9.1f} driver ({drv_pct:.1f}%) | "
+        f"delta {delta:>+9.1f}"
+    )
+    return {
+        "tag": tag,
+        "alloc_mb": alloc_mb,
+        "drv_used_mb": drv_used_mb,
+        "drv_pct": drv_pct,
+    }
+
+
+def _full_gpu_cleanup():
+    """gc.collect + torch empty_cache, multiple rounds."""
+    gc.unfreeze()
+    for _ in range(3):
+        if gc.collect() == 0:
+            break
+    torch.accelerator.empty_cache()
+
+
+@pytest.mark.parametrize("model_name, sw_size", [("google/gemma-3-1b-it", 512)])
+def test_gpu_memory_rixl_hma(model_name, sw_size):
+    """Track GPU memory through NixlConnector create/infer/shutdown cycle."""
+    from vllm import LLM, SamplingParams
+    from vllm.config import KVTransferConfig
+    from vllm.distributed.parallel_state import cleanup_dist_env_and_memory
+
+    llm_kwargs = {
+        "model": model_name,
+        "enforce_eager": True,
+        "gpu_memory_utilization": 0.5,
+        "kv_transfer_config": KVTransferConfig(
+            kv_connector="NixlConnector",
+            kv_role="kv_both",
+        ),
+        "max_model_len": 2048,
+        "disable_hybrid_kv_cache_manager": False,
+        "max_num_batched_tokens": 1024,
+        "enable_prefix_caching": False,
+        "block_size": 16,
+    }
+
+    print("\n" + "=" * 90)
+    print("GPU MEMORY -- RIXL NixlConnector HMA (ROCm)")
+    print("=" * 90)
+    gc.collect()
+    torch.accelerator.empty_cache()
+    torch.accelerator.reset_peak_memory_stats()
+    snap0 = _gpu_snapshot("0. baseline", 0.0)
+
+    # create + infer
+    llm = LLM(**llm_kwargs)
+    snap1 = _gpu_snapshot("1. after LLM()", snap0["alloc_mb"])
+
+    llm.generate(
+        ["hi" * 1401],
+        SamplingParams(
+            temperature=0.0,
+            max_tokens=1,
+            extra_args={
+                "kv_transfer_params": {
+                    "do_remote_decode": True,
+                    "do_remote_prefill": False,
+                    "remote_engine_id": None,
+                    "remote_block_ids": None,
+                    "remote_host": None,
+                    "remote_port": None,
+                }
+            },
+        ),
+    )
+    snap2 = _gpu_snapshot("2. after generate()", snap1["alloc_mb"])
+
+    # shutdown + cleanup
+    print("\n--- shutdown ---")
+    llm.llm_engine.engine_core.shutdown()
+    _gpu_snapshot("3. after shutdown()", snap2["alloc_mb"])
+
+    del llm
+    _full_gpu_cleanup()
+    cleanup_dist_env_and_memory()
+    _full_gpu_cleanup()
+    torch._dynamo.reset()
+    gc.collect()
+    torch.accelerator.empty_cache()
+    snap_final = _gpu_snapshot("4. final", snap2["alloc_mb"])
+
+    # summary
+    print("\n" + "=" * 90)
+    baseline = snap0["alloc_mb"]
+    final = snap_final["alloc_mb"]
+    peak = snap2["alloc_mb"]
+    total_alloc = peak - baseline
+
+    print(
+        f"  PyTorch:  baseline={baseline:.0f}  peak={peak:.0f}  "
+        f"final={final:.0f}  "
+        f"leaked={final - baseline:.0f} MB"
+        + (
+            f" ({(final - baseline) / total_alloc * 100:.1f}%)"
+            if total_alloc > 0
+            else ""
+        )
+    )
+
+    drv_base = snap0["drv_used_mb"]
+    drv_final = snap_final["drv_used_mb"]
+    drv_leaked = drv_final - drv_base
+    print(
+        f"  Driver:   baseline={drv_base:.0f} ({snap0['drv_pct']:.1f}%)  "
+        f"peak={snap2['drv_used_mb']:.0f} ({snap2['drv_pct']:.1f}%)  "
+        f"final={drv_final:.0f} ({snap_final['drv_pct']:.1f}%)  "
+        f"leaked={drv_leaked:.0f} MB"
+    )
+    print("=" * 90)
+
+    # Peak driver memory used above baseline
+    drv_peak = snap2["drv_used_mb"] - drv_base
+    leak_pct = (drv_leaked / drv_peak * 100) if drv_peak > 0 else 0
+    max_leak_pct = 10
+    assert leak_pct <= max_leak_pct, (
+        f"{drv_leaked:.0f} MB ({leak_pct:.1f}%) of driver-level GPU memory "
+        f"not freed after NixlConnector shutdown "
+        f"(peak allocation: {drv_peak:.0f} MB, threshold: {max_leak_pct}%)"
+    )
+
+
+@pytest.mark.parametrize("model_name", ["google/gemma-3-1b-it"])
+def test_gpu_memory_no_rixl_baseline(model_name):
+    """Same workload without NixlConnector.  Comparing driver-level memory
+    between this and test_gpu_memory_rixl_hma isolates UCX/RIXL impact."""
+    from vllm import LLM, SamplingParams
+    from vllm.distributed.parallel_state import cleanup_dist_env_and_memory
+
+    print("\n" + "=" * 90)
+    print("CONTROL -- same model, no RIXL connector")
+    print("=" * 90)
+    gc.collect()
+    torch.accelerator.empty_cache()
+    snap0 = _gpu_snapshot("baseline", 0.0)
+
+    llm = LLM(
+        model=model_name,
+        enforce_eager=True,
+        gpu_memory_utilization=0.5,
+        max_model_len=2048,
+        max_num_batched_tokens=1024,
+        enable_prefix_caching=False,
+        block_size=16,
+    )
+    _gpu_snapshot("after LLM()", snap0["alloc_mb"])
+
+    llm.generate(["hi " * 500], SamplingParams(max_tokens=1))
+    snap_peak = _gpu_snapshot("after generate()", snap0["alloc_mb"])
+
+    llm.llm_engine.engine_core.shutdown()
+    del llm
+    _full_gpu_cleanup()
+    cleanup_dist_env_and_memory()
+    _full_gpu_cleanup()
+    torch._dynamo.reset()
+    gc.collect()
+    torch.accelerator.empty_cache()
+    snap_final = _gpu_snapshot("final", snap0["alloc_mb"])
+
+    drv_base = snap0["drv_used_mb"]
+    drv_leaked = snap_final["drv_used_mb"] - drv_base
+    drv_peak = snap_peak["drv_used_mb"] - drv_base
+    print(f"\n  Driver leaked (no rixl): {drv_leaked:.0f} MB")
+    print("=" * 90)
+
+    leak_pct = (drv_leaked / drv_peak * 100) if drv_peak > 0 else 0
+    max_leak_pct = 10
+    assert leak_pct <= max_leak_pct, (
+        f"{drv_leaked:.0f} MB ({leak_pct:.1f}%) of driver-level GPU memory "
+        f"not freed after baseline shutdown "
+        f"(peak allocation: {drv_peak:.0f} MB, threshold: {max_leak_pct}%)"
+    )
@@ -87,10 +87,13 @@ class MockSubscriber:
 def _wait_for_prefix_cache_reset(llm: LLM) -> None:
    """Wait for async offload transfers to finish so prefix cache can reset.

-    The GPU-to-CPU offload runs on a CUDA stream asynchronously.  While blocks
+    The GPU-to-CPU offload runs on a CUDA stream asynchronously. While blocks
    are still held by the offload worker, ``reset_prefix_cache`` returns
-    ``False``.  Retry with a short sleep until it succeeds or we time out.
+    ``False``. Between retries we send a dummy single-token prefill to force
+    the engine to step, which polls the worker for completed transfers and
+    frees GPU blocks.
    """
+    _dummy_params = SamplingParams(max_tokens=1)
    deadline = time.monotonic() + _RESET_CACHE_TIMEOUT
    while not llm.reset_prefix_cache():
        if time.monotonic() > deadline:
@@ -98,7 +101,13 @@ def _wait_for_prefix_cache_reset(llm: LLM) -> None:
                "reset_prefix_cache did not succeed within "
                f"{_RESET_CACHE_TIMEOUT}s - async offload may be stuck"
            )
-        time.sleep(0.1)
+        # Force an engine step so the scheduler polls get_finished()
+        # and releases GPU blocks held by in-flight async stores.
+        llm.generate(
+            [TokensPrompt(prompt_token_ids=[0])],
+            _dummy_params,
+            use_tqdm=False,
+        )


 def _latency_test(llm: LLM, subscriber: MockSubscriber):
@@ -933,3 +933,64 @@ def test_sample_recovered_tokens(
        device=DEVICE_TYPE,
    )
    assert torch.equal(recovered_token_ids, ref_recovered_token_ids)
+
+
+########################### Tests for Synthetic Rejection Sampling #########
+
+
+def _make_synthetic_sampler(rates: list[float]) -> RejectionSampler:
+    mock_sampler = Mock(spec=Sampler)
+    mock_sampler.logprobs_mode = "raw_logprobs"
+    spec_config = Mock()
+    spec_config.rejection_sample_method = "synthetic"
+    spec_config.synthetic_acceptance_rates = rates
+    return RejectionSampler(mock_sampler, spec_config, torch.device(DEVICE_TYPE))
+
+
+def _make_sampling_metadata(all_greedy: bool) -> SamplingMetadata:
+    temperature = None if all_greedy else torch.tensor([1.0, 1.0], device=DEVICE_TYPE)
+    return create_sampling_metadata(all_greedy=all_greedy, temperature=temperature)
+
+
+@pytest.mark.parametrize("all_greedy", [True, False])
+def test_synthetic_all_accepted(all_greedy: bool):
+    """With all rates=1.0, every draft token is accepted."""
+    sampler = _make_synthetic_sampler([1.0, 1.0])
+    spec_tokens = [[1, 2], [3]]
+    output_tokens = [[10, 20, 50], [30, 40]]
+
+    metadata = _make_sampling_metadata(all_greedy)
+    logits = create_logits_tensor(output_tokens)
+    bonus = torch.tensor([50, 40], device=DEVICE_TYPE)
+    spec_decode_metadata = create_spec_decode_metadata(spec_tokens, logits)
+
+    mock_sampler_output(sampler, bonus)
+    output = sampler(spec_decode_metadata, None, logits, metadata)
+    expected = torch.tensor(
+        [[1, 2, 50], [3, 40, PLACEHOLDER_TOKEN_ID]],
+        dtype=torch.int,
+        device=DEVICE_TYPE,
+    )
+    assert torch.equal(output.sampled_token_ids, expected)
+
+
+@pytest.mark.parametrize("all_greedy", [True, False])
+def test_synthetic_all_rejected(all_greedy: bool):
+    """With all rates=0.0, the first token is always rejected."""
+    sampler = _make_synthetic_sampler([0.0, 0.0])
+    spec_tokens = [[1, 2], [3]]
+    output_tokens = [[10, 20, 50], [30, 40]]
+
+    metadata = _make_sampling_metadata(all_greedy)
+    logits = create_logits_tensor(output_tokens)
+    bonus = torch.tensor([50, 40], device=DEVICE_TYPE)
+    spec_decode_metadata = create_spec_decode_metadata(spec_tokens, logits)
+
+    mock_sampler_output(sampler, bonus)
+    output = sampler(spec_decode_metadata, None, logits, metadata)
+    result = output.sampled_token_ids
+    # Exactly one token emitted per sequence (the rejection fallback),
+    # followed by placeholders.
+    for row in result:
+        assert row[0] != PLACEHOLDER_TOKEN_ID
+        assert (row[1:] == PLACEHOLDER_TOKEN_ID).all()
@@ -165,6 +165,7 @@ def get_mt_bench_prompts(
        no_stream=True,
        disable_shuffle=False,
        skip_chat_template=False,
+        trust_remote_code=False,
    )
    samples = get_samples(args, tokenizer)
    prompt_ids = [
@@ -210,8 +211,8 @@ def extract_acceptance_metrics(metrics, num_spec_tokens: int) -> dict:

@large_gpu_mark(min_gb=40)
@pytest.mark.skipif(
-    not current_platform.is_cuda(),
-    reason="This test is only supported on CUDA platform.",
+    not current_platform.is_cuda_alike(),
+    reason="This test is only supported on CUDA-alike platforms.",
 )
@pytest.mark.parametrize(
    "model_config",
@@ -741,9 +741,9 @@ def test_set_inputs_first_pass_parallel_drafting():
@pytest.mark.parametrize("pp_size", [1, 2])
@pytest.mark.parametrize("use_distinct_embed_tokens", [True, False])
@pytest.mark.parametrize("use_distinct_lm_head", [True, False])
-@mock.patch("vllm.v1.spec_decode.eagle.get_pp_group")
-@mock.patch("vllm.v1.spec_decode.eagle.get_layers_from_vllm_config")
-@mock.patch("vllm.v1.spec_decode.eagle.get_model")
+@mock.patch("vllm.v1.spec_decode.llm_base_proposer.get_pp_group")
+@mock.patch("vllm.v1.spec_decode.llm_base_proposer.get_layers_from_vllm_config")
+@mock.patch("vllm.v1.spec_decode.llm_base_proposer.get_model")
 def test_load_model(
    mock_get_model,
    mock_get_layers,
@@ -61,9 +61,9 @@ def _create_mtp_proposer(num_speculative_tokens: int) -> EagleProposer:
    return EagleProposer(vllm_config=vllm_config, device=DEVICE_TYPE)


-@mock.patch("vllm.v1.spec_decode.eagle.get_pp_group")
-@mock.patch("vllm.v1.spec_decode.eagle.get_layers_from_vllm_config")
-@mock.patch("vllm.v1.spec_decode.eagle.get_model")
+@mock.patch("vllm.v1.spec_decode.llm_base_proposer.get_pp_group")
+@mock.patch("vllm.v1.spec_decode.llm_base_proposer.get_layers_from_vllm_config")
+@mock.patch("vllm.v1.spec_decode.llm_base_proposer.get_model")
 def test_mtp_load_model_unified(mock_get_model, mock_get_layers, mock_get_pp_group):
    """Test MTP-specific model loading with unified model approach."""

@@ -2,33 +2,48 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 import pytest

-from vllm.v1.worker.gpu.spec_decode.synthetic_rejection_sampler_utils import (
-    compute_synthetic_rejection_sampler_params,
+from vllm.config.speculative import SpeculativeConfig
+from vllm.v1.spec_decode.utils import unconditional_to_conditional_rates
+
+
+def test_unconditional_to_conditional_rates_basic():
+    # c_0 = p_0; c_i = p_i / p_{i-1}
+    assert unconditional_to_conditional_rates([0.9, 0.5, 0.2]) == pytest.approx(
+        [0.9, 0.5 / 0.9, 0.2 / 0.5]
+    )
+
+
+def test_unconditional_to_conditional_rates_handles_zero():
+    # After a zero, subsequent conditional rates are clamped to 0 (the chain
+    # has already terminated in the kernel, so these values are unused).
+    assert unconditional_to_conditional_rates([1.0, 0.6, 0.0, 0.0]) == pytest.approx(
+        [1.0, 0.6, 0.0, 0.0]
+    )
+
+
+def test_unconditional_to_conditional_rates_all_ones():
+    assert unconditional_to_conditional_rates([1.0, 1.0, 1.0]) == pytest.approx(
+        [1.0, 1.0, 1.0]
+    )
+
+
+@pytest.mark.parametrize(
+    "length,n,expected",
+    [
+        (2.6, 3, [1.0, 0.6, 0.0]),
+        (1.0, 3, [0.0, 0.0, 0.0]),
+        (4.0, 3, [1.0, 1.0, 1.0]),
+        (2.0, 3, [1.0, 0.0, 0.0]),
+        (3.5, 4, [1.0, 1.0, 0.5, 0.0]),
+    ],
 )
-
-NUM_SPECULATIVE_STEPS = [1, 2, 3, 4, 5, 7, 10]
-ACCEPTANCE_RATES = [i / 100 for i in range(0, 100)]
+def test_acceptance_length_to_rates(length, n, expected):
+    assert SpeculativeConfig._acceptance_length_to_rates(length, n) == pytest.approx(
+        expected
+    )


-@pytest.mark.parametrize("num_speculative_steps", NUM_SPECULATIVE_STEPS)
-def test_compute_synthetic_rejection_sampler_params(num_speculative_steps: int):
-    """Test that the base acceptance rate and decay factor generated for
-    synthetic rejection sampling have a mean joint acceptance probability
-    that matches the desired acceptance rate."""
-    tol = 1e-9
-    for desired_acceptance_rate in ACCEPTANCE_RATES:
-        base_rate, decay_factor = compute_synthetic_rejection_sampler_params(
-            desired_acceptance_rate, num_speculative_steps, tol=tol
-        )
-
-        # Compute the mean of joint acceptance probabilities across
-        # all speculative positions.
-        joint_prob = 1.0
-        mean_joint = 0.0
-        for i in range(num_speculative_steps):
-            joint_prob *= base_rate * decay_factor**i
-            mean_joint += joint_prob
-        mean_joint /= num_speculative_steps
-
-        assert abs(desired_acceptance_rate - mean_joint) < 10 * tol
-        assert base_rate <= 1.0
+def test_resolve_length_produces_minvariance_schedule():
+    assert SpeculativeConfig._resolve_synthetic_acceptance_rates(
+        3, None, 2.6
+    ) == pytest.approx([1.0, 0.6, 0.0])
@@ -241,11 +241,13 @@ def forward_attention(
        )
    kv_cache_spec = create_standard_kv_cache_spec(vllm_config)
    builder = builder_cls(kv_cache_spec, [], vllm_config, q.device)
+    seq_lens_cpu = seq_lens.cpu()
    common_attn_metadata = CommonAttentionMetadata(
        query_start_loc=query_start_loc,
        query_start_loc_cpu=query_start_loc.cpu(),
        seq_lens=seq_lens,
-        _seq_lens_cpu=seq_lens.cpu(),
+        seq_lens_cpu_upper_bound=seq_lens_cpu,
+        _seq_lens_cpu=seq_lens_cpu,
        _num_computed_tokens_cpu=context_lens.cpu(),
        num_reqs=batch_size,
        num_actual_tokens=num_actual_tokens,
@@ -101,7 +101,7 @@ NVSHMEM_URL="https://developer.download.nvidia.com/compute/nvshmem/redist/libnvs

 pushd "$WORKSPACE"
 echo "Downloading NVSHMEM ${NVSHMEM_VER} for ${NVSHMEM_SUBDIR} ..."
-curl -fSL "${NVSHMEM_URL}" -o "${NVSHMEM_FILE}"
+curl -fSL --retry 3 --retry-delay 2 "${NVSHMEM_URL}" -o "${NVSHMEM_FILE}"
 tar -xf "${NVSHMEM_FILE}"
 rm -rf nvshmem
 mv "${NVSHMEM_FILE%.tar.xz}" nvshmem
@@ -634,9 +634,10 @@ def parse_flash_attn_features() -> dict[str, dict[str, Any]]:
    except Exception:
        return {}

-    # Analyze the functions to determine FA3-specific features
+    # Analyze the functions to determine FA3/FA4-specific features
    fa3_supports_fp8 = False
    fa3_supports_sinks = False
+    fa4_supports_sinks = False
    fa3_compute_cap: str | None = None
    fa4_compute_cap: str | None = None

@@ -656,17 +657,49 @@ def parse_flash_attn_features() -> dict[str, dict[str, Any]]:
                    fa3_supports_fp8 = True
                    break

-        # Check flash_attn_supports_sinks - looks for `get_flash_attn_version() == 3`
+        # Check flash_attn_supports_sinks - looks for `fa_version == 3/4`
+        # or `get_flash_attn_version() == 3/4` (also accepts `in (3, 4)`)
        if node.name == "flash_attn_supports_sinks":
            for n in ast.walk(node):
                if (
                    isinstance(n, ast.Compare)
-                    and isinstance(n.left, ast.Call)
-                    and isinstance(n.left.func, ast.Name)
-                    and n.left.func.id == "get_flash_attn_version"
+                    and len(n.ops) == 1
+                    and isinstance(n.ops[0], ast.Eq)
+                    and isinstance(n.comparators[0], ast.Constant)
                ):
-                    fa3_supports_sinks = True
-                    break
+                    is_version_compare = (
+                        isinstance(n.left, ast.Name) and n.left.id == "fa_version"
+                    ) or (
+                        isinstance(n.left, ast.Call)
+                        and isinstance(n.left.func, ast.Name)
+                        and n.left.func.id == "get_flash_attn_version"
+                    )
+                    if is_version_compare:
+                        val = n.comparators[0].value
+                        if val == 3:
+                            fa3_supports_sinks = True
+                        elif val == 4:
+                            fa4_supports_sinks = True
+                elif (
+                    isinstance(n, ast.Compare)
+                    and len(n.ops) == 1
+                    and isinstance(n.ops[0], ast.In)
+                    and isinstance(n.comparators[0], (ast.Tuple, ast.List, ast.Set))
+                ):
+                    is_version_compare = (
+                        isinstance(n.left, ast.Name) and n.left.id == "fa_version"
+                    ) or (
+                        isinstance(n.left, ast.Call)
+                        and isinstance(n.left.func, ast.Name)
+                        and n.left.func.id == "get_flash_attn_version"
+                    )
+                    if is_version_compare:
+                        for elt in n.comparators[0].elts:
+                            if isinstance(elt, ast.Constant):
+                                if elt.value == 3:
+                                    fa3_supports_sinks = True
+                                elif elt.value == 4:
+                                    fa4_supports_sinks = True

        # Check get_flash_attn_version for FA3/FA4 compute capability
        if node.name == "get_flash_attn_version":
@@ -731,7 +764,7 @@ def parse_flash_attn_features() -> dict[str, dict[str, Any]]:
        "fa4": {
            "compute_capability": fa4_compute_cap,
            "supports_fp8": False,
-            "supports_sink": False,
+            "supports_sink": fa4_supports_sinks,
        },
    }

@@ -1782,6 +1782,8 @@ class rocm_aiter_ops:
        need_renorm: bool,
        routed_scaling_factor: float = 1.0,
    ) -> None:
+        if correction_bias.dtype != gating_output.dtype:
+            correction_bias = correction_bias.to(gating_output.dtype)
        torch.ops.vllm.rocm_aiter_biased_grouped_topk(
            gating_output,
            correction_bias,
@@ -92,6 +92,72 @@ if hasattr(torch.ops._xpu_C, "int4_gemm_w4a16"):
        return torch.empty((M, N), dtype=input.dtype, device=input.device)


+def _gdn_attention_core_xpu_impl(
+    core_attn_out: torch.Tensor,
+    z: torch.Tensor,
+    projected_states_qkvz: torch.Tensor,
+    projected_states_ba: torch.Tensor,
+    layer_name: str,
+) -> None:
+    """Custom op wrapping the XPU SYCL GDN kernel for torch.compile."""
+    from vllm.forward_context import get_forward_context
+    from vllm.v1.attention.backends.gdn_attn import GDNAttentionMetadata
+
+    forward_context = get_forward_context()
+    self = forward_context.no_compile_layers[layer_name]
+    attn_metadata_raw = forward_context.attn_metadata
+
+    if attn_metadata_raw is None:
+        return
+
+    assert isinstance(attn_metadata_raw, dict)
+    attn_metadata = attn_metadata_raw[self.prefix]
+    assert isinstance(attn_metadata, GDNAttentionMetadata)
+
+    # TODO: xpu does not support speculative decoding yet
+    assert attn_metadata.spec_sequence_masks is None  # type: ignore[attr-defined]
+
+    conv_weights = self.conv1d.weight.view(
+        self.conv1d.weight.size(0), self.conv1d.weight.size(2)
+    )
+
+    torch.ops._xpu_C.gdn_attention(
+        core_attn_out,
+        z,
+        projected_states_qkvz,
+        projected_states_ba,
+        self.num_k_heads,
+        self.num_v_heads,
+        self.head_k_dim,
+        self.head_v_dim,
+        conv_state=self.kv_cache[0],
+        ssm_state=self.kv_cache[1],
+        conv_weights=conv_weights,
+        conv_bias=self.conv1d.bias,
+        activation=self.activation,
+        A_log=self.A_log,
+        dt_bias=self.dt_bias,
+        num_prefills=attn_metadata.num_prefills,  # type: ignore[attr-defined]
+        num_decodes=attn_metadata.num_decodes,  # type: ignore[attr-defined]
+        has_initial_state=attn_metadata.has_initial_state,  # type: ignore[attr-defined]
+        non_spec_query_start_loc=attn_metadata.non_spec_query_start_loc,  # type: ignore[attr-defined]
+        non_spec_state_indices_tensor=attn_metadata.non_spec_state_indices_tensor,  # type: ignore[attr-defined]
+        num_actual_tokens=attn_metadata.num_actual_tokens,  # type: ignore[attr-defined]
+        tp_size=self.tp_size,
+        reorder_input=not self.gqa_interleaved_layout,
+    )
+
+
+def _gdn_attention_core_xpu_fake(
+    core_attn_out: torch.Tensor,
+    z: torch.Tensor,
+    projected_states_qkvz: torch.Tensor,
+    projected_states_ba: torch.Tensor,
+    layer_name: str,
+) -> None:
+    return
+
+
 def _xpu_ops_deepseek_scaling_rope_impl(
    positions: torch.Tensor,
    query: torch.Tensor,
@@ -618,6 +684,13 @@ class xpu_ops:
                fake_impl=_xpu_mxfp4_quantize_fake,
            )

+            direct_register_custom_op(
+                op_name="gdn_attention_core_xpu",
+                op_func=_gdn_attention_core_xpu_impl,
+                mutates_args=["core_attn_out", "z"],
+                fake_impl=_gdn_attention_core_xpu_fake,
+            )
+
            _OPS_REGISTERED = True


@@ -1611,14 +1611,12 @@ def add_cli_args(parser: argparse.ArgumentParser):
    )
    parser.add_argument(
        "--timeline-itl-thresholds",
-        type=float,
-        nargs=2,
-        default=[25.0, 50.0],
-        metavar=("THRESHOLD1", "THRESHOLD2"),
+        type=str,
+        default="25,50",
        help="ITL thresholds in milliseconds for timeline plot coloring. "
-        "Specify two values to categorize inter-token latencies into three groups: "
-        "below first threshold (green), between thresholds (orange), "
-        "and above second threshold (red). Default: 25 50 (milliseconds).",
+        "Specify two comma-separated values to categorize inter-token "
+        "latencies into three groups: below first threshold (green), "
+        "between thresholds (orange), and above second threshold (red).",
    )
    parser.add_argument(
        "--plot-dataset-stats",
@@ -1637,6 +1635,19 @@ async def main_async(args: argparse.Namespace) -> dict[str, Any]:
    random.seed(args.seed)
    np.random.seed(args.seed)

+    # Validate timeline ITL thresholds
+    if args.plot_timeline:
+        try:
+            itl_thresholds = [
+                float(t.strip()) for t in args.timeline_itl_thresholds.split(",")
+            ]
+            if len(itl_thresholds) != 2:
+                raise ValueError(
+                    f"Expected 2 ITL threshold values, got {len(itl_thresholds)}"
+                )
+        except ValueError as e:
+            raise ValueError(f"Invalid --timeline-itl-thresholds format: {e}") from e
+
    # Validate ramp-up arguments
    if args.ramp_up_strategy is not None:
        if args.request_rate != float("inf"):
@@ -1906,7 +1917,9 @@ async def main_async(args: argparse.Namespace) -> dict[str, Any]:

                timeline_path = Path(file_name).with_suffix(".timeline.html")
                # Convert thresholds from milliseconds to seconds
-                itl_thresholds_sec = [t / 1000.0 for t in args.timeline_itl_thresholds]
+                itl_thresholds_sec = [
+                    float(t) / 1000.0 for t in args.timeline_itl_thresholds.split(",")
+                ]
                generate_timeline_plot(
                    per_request_data, timeline_path, itl_thresholds=itl_thresholds_sec
                )
@@ -744,6 +744,7 @@ class CompilationConfig:
        "vllm::linear_attention",
        "vllm::plamo2_mamba_mixer",
        "vllm::gdn_attention_core",
+        "vllm::gdn_attention_core_xpu",
        "vllm::olmo_hybrid_gdn_full_forward",
        "vllm::kda_attention",
        "vllm::sparse_attn_indexer",
@@ -953,8 +953,12 @@ class ModelConfig:
                "mxfp4",
                "gpt_oss_mxfp4",
                "cpu_awq",
+                "humming",
                "gguf",
            ]
+            # if the user specifies humming, we should always use humming
+            if self.quantization == "humming":
+                overrides = ["humming"] + overrides
            quantization_methods = [
                q for q in supported_quantization if q not in overrides
            ]
@@ -189,12 +189,64 @@ class SpeculativeConfig:
    distribution, but the latter yields a higher acceptance rate at the cost
    of more memory to cache draft logits."""

-    synthetic_acceptance_rate: float | None = None
-    """Average acceptance rate for synthetic rejection sampling. Draft
-    tokens are accepted with a position-dependent probability that decays
-    geometrically, calibrated so that the mean rate across all speculative
-    positions equals this value. Only used when rejection_sample_method
-    is 'synthetic'. Must be in [0, 1]."""
+    synthetic_acceptance_rates: list[float] | None = None
+    """Per-position *unconditional* acceptance rates for synthetic rejection
+    sampling. Position i's entry is the marginal probability that the first
+    i+1 draft tokens are all accepted; the list must have length
+    num_speculative_tokens, each entry in [0, 1], and be monotonically
+    non-increasing. Only valid when rejection_sample_method is 'synthetic'.
+    Mutually exclusive with synthetic_acceptance_length."""
+
+    synthetic_acceptance_length: float | None = None
+    """Target mean acceptance length for synthetic rejection sampling, in
+    [1, num_speculative_tokens + 1]. Resolved internally to
+    synthetic_acceptance_rates. Only valid when rejection_sample_method is 'synthetic'.
+    Mutually exclusive with synthetic_acceptance_rates."""
+
+    @staticmethod
+    def _acceptance_length_to_rates(length: float, n: int) -> list[float]:
+        """Mean acceptance length to unconditional per-position rates, using
+        the minimum-variance schedule."""
+        num_drafts = length - 1  # expected number of accepted draft tokens
+        num_full = int(num_drafts)
+        return (
+            [1.0] * num_full + [num_drafts - num_full] + [0.0] * (n - num_full - 1)
+        )[:n]
+
+    @staticmethod
+    def _resolve_synthetic_acceptance_rates(
+        n: int,
+        rates: list[float] | None,
+        length: float | None,
+    ) -> list[float]:
+        """Return per-position unconditional acceptance rates from exactly one
+        of `rates` or `length` (validates range, length, and monotonicity)."""
+        if (rates is None) == (length is None):
+            raise ValueError(
+                "rejection_sample_method='synthetic' requires exactly one of "
+                "synthetic_acceptance_rates or synthetic_acceptance_length."
+            )
+        if rates is not None:
+            if len(rates) != n:
+                raise ValueError(
+                    f"synthetic_acceptance_rates must have length {n}, got {rates}."
+                )
+            if not all(0.0 <= r <= 1.0 for r in rates):
+                raise ValueError(
+                    f"synthetic_acceptance_rates entries must be in [0, 1], "
+                    f"got {rates}."
+                )
+            if any(rates[i] > rates[i - 1] for i in range(1, n)):
+                raise ValueError(
+                    f"synthetic_acceptance_rates must be non-increasing, got {rates}."
+                )
+            return list(rates)
+        assert length is not None
+        if not 1.0 <= length <= float(n + 1):
+            raise ValueError(
+                f"synthetic_acceptance_length must be in [1, {n + 1}], got {length}."
+            )
+        return SpeculativeConfig._acceptance_length_to_rates(length, n)

    def compute_hash(self) -> str:
        """
@@ -818,6 +870,23 @@ class SpeculativeConfig:
                f"than zero ({self.num_speculative_tokens})."
            )

+        if self.rejection_sample_method == "synthetic":
+            # Consolidate to per-position rates
+            self.synthetic_acceptance_rates = self._resolve_synthetic_acceptance_rates(
+                self.num_speculative_tokens,
+                self.synthetic_acceptance_rates,
+                self.synthetic_acceptance_length,
+            )
+            self.synthetic_acceptance_length = None
+        elif (
+            self.synthetic_acceptance_rates is not None
+            or self.synthetic_acceptance_length is not None
+        ):
+            raise ValueError(
+                "synthetic_acceptance_rates / synthetic_acceptance_length "
+                "are only valid with rejection_sample_method='synthetic'."
+            )
+
        if self.draft_model_config:
            self.draft_model_config.verify_with_parallel_config(
                self.draft_parallel_config
@@ -900,6 +900,13 @@ class VllmConfig:
            self.compilation_config.mode = CompilationMode.NONE
            self.compilation_config.cudagraph_mode = CUDAGraphMode.NONE

+        if os.environ.get("TORCH_COMPILE_DISABLE") == "1":
+            logger.warning(
+                "TORCH_COMPILE_DISABLE is set, disabling torch.compile. "
+                "This is equivalent to setting -cc.mode=none"
+            )
+            self.compilation_config.mode = CompilationMode.NONE
+
        if self.compilation_config.backend == "eager" or (
            self.compilation_config.mode is not None
            and self.compilation_config.mode != CompilationMode.VLLM_COMPILE
@@ -4,7 +4,6 @@
 The async worker that transfers experts in the background.
 """

-import asyncio
 import threading
 from typing import TYPE_CHECKING

@@ -36,21 +35,15 @@ def start_async_worker(
        assert device_index is not None
        torch.accelerator.set_device_index(device_index)
        cuda_stream = torch.cuda.Stream(device=device_index)
-        loop = asyncio.new_event_loop()
-        asyncio.set_event_loop(loop)
        try:
-            loop.run_until_complete(
-                transfer_run_periodically(
-                    state=state,
-                    eplb_group=eplb_group,
-                    cuda_stream=cuda_stream,
-                    is_profile=is_profile,
-                )
+            transfer_run_periodically(
+                state=state,
+                eplb_group=eplb_group,
+                cuda_stream=cuda_stream,
+                is_profile=is_profile,
            )
        except Exception as exc:  # pragma: no cover - diagnostic path
            logger.exception("async loop error (Rank %d): %s", rank, str(exc))
-        finally:
-            loop.close()

    thread = threading.Thread(target=thread_target, daemon=True)
    thread.start()
@@ -83,7 +76,7 @@ def run_rebalance_experts(
    return new_physical_to_logical_map


-async def transfer_run_periodically(
+def transfer_run_periodically(
    state: "EplbState",
    eplb_group: ProcessGroup,
    cuda_stream: torch.cuda.Stream,
@@ -118,7 +111,7 @@ async def transfer_run_periodically(
            # model_state.expert_buffer, which will be consumed by the main thread in
            # move_to_workspace
            while model_state.rebalanced and layer_idx < num_layers:
-                transfer_metadata = await transfer_layer(
+                transfer_metadata = transfer_layer(
                    old_layer_indices=physical_to_logical_map_cpu[layer_idx],
                    new_layer_indices=new_physical_to_logical_map[layer_idx],
                    expert_weights=model_state.model.expert_weights[layer_idx],
@@ -418,7 +418,7 @@ def move_from_buffer(
            w[dst].copy_(w[src], non_blocking=True)


-async def transfer_layer(
+def transfer_layer(
    old_layer_indices: torch.Tensor,
    new_layer_indices: torch.Tensor,
    expert_weights: Sequence[torch.Tensor],
@@ -120,6 +120,13 @@ class OffloadingConnectorScheduler:
        self.config = SchedulerOffloadConfig.from_spec(spec)
        self.manager: OffloadingManager = spec.get_manager()

+        attention_groups: list[int] = []
+        for idx, _ in enumerate(spec.kv_cache_config.kv_cache_groups):
+            # currently treat all groups as full attention
+            attention_groups.append(idx)
+
+        self.lookup_groups = attention_groups
+
        self._req_status: dict[ReqId, RequestOffloadState] = {}
        # requests to load for the current scheduler step
        self._reqs_to_load: dict[ReqId, TransferSpec] = {}
@@ -204,64 +211,88 @@ class OffloadingConnectorScheduler:
                group_state.block_ids.clear()
        else:
            req_status = RequestOffloadState(config=self.config, req=request)
-            req_status.update_offload_keys()
            self._req_status[request.request_id] = req_status

+        req_status.update_offload_keys()
        req_status.num_locally_computed_tokens = num_computed_tokens

-        # Below assertions will be removed once this function supports HMA
-        assert len(self.config.kv_group_configs) == 1
-        assert len(req_status.group_states) == 1
-        group_config = self.config.kv_group_configs[0]
-        group_state = req_status.group_states[0]
+        for gs in req_status.group_states:
+            self.manager.touch(gs.offload_keys)

-        num_blocks = request.num_tokens // group_config.offloaded_block_size
+        # Start with the full request size as the maximum loadable
+        max_hit_size_tokens: int = req_status.req.num_tokens
+        num_hit_tokens: int = 0
+        defer_lookup = False
+        delay_request = False
+        for group_idx in self.lookup_groups:
+            group_config: GroupOffloadConfig = self.config.kv_group_configs[group_idx]
+            offloaded_block_size = group_config.offloaded_block_size
+            offload_keys = req_status.group_states[group_idx].offload_keys

-        assert len(request.block_hashes) // self.config.block_size_factor == num_blocks
-        offload_keys = group_state.offload_keys
+            num_blocks = max_hit_size_tokens // offloaded_block_size
+            assert len(offload_keys) >= num_blocks

-        self.manager.touch(offload_keys)
+            # Constrain to block-aligned boundary for this group
+            max_hit_size_tokens = num_blocks * offloaded_block_size
+            num_hit_tokens = max_hit_size_tokens - num_computed_tokens
+            if num_hit_tokens < offloaded_block_size:
+                # we can only load less than a block, better skip
+                return 0, False

-        full_block_tokens = group_config.offloaded_block_size * num_blocks
-        if full_block_tokens - num_computed_tokens < group_config.offloaded_block_size:
-            # we can load less than a block, skip
-            return 0, False
+            start_block_idx = num_computed_tokens // offloaded_block_size
+            offload_keys = offload_keys[start_block_idx:num_blocks]
+            # Full attention relies on all previous KV cache blocks.
+            # Thus, we search for a maximal prefix of KV cache which are all cached.
+            block_hits = self._maximal_prefix_lookup(
+                offload_keys, req_status.req_context
+            )
+            if block_hits == 0:
+                return 0, False

-        start_block_idx = num_computed_tokens // group_config.offloaded_block_size
-        # Full attention relays on all previous KV cache blocks.
-        # Thus, we search for a maximal prefix of KV cache which are all cached.
-        hits = self._maximal_prefix_lookup(
-            offload_keys[start_block_idx:], req_status.req_context
-        )
-        if hits is None:
-            # indicates a lookup that should be tried later
+            if block_hits is None:
+                defer_lookup = True
+            else:
+                # Further constrain based on what's actually available by backend
+                max_hit_size_tokens = offloaded_block_size * (
+                    start_block_idx + block_hits
+                )
+
+            num_hit_tokens = max_hit_size_tokens - num_computed_tokens
+            if num_hit_tokens < offloaded_block_size:
+                # we can only load less than a block, better skip
+                return 0, False
+
+            if (
+                block_hits
+                and self._blocks_being_loaded
+                and any(
+                    key in self._blocks_being_loaded
+                    for key in offload_keys[:block_hits]
+                )
+            ):
+                # hit blocks are being loaded, delay request
+                delay_request = True
+
+        if defer_lookup:
+            logger.debug(
+                "Offloading manager delayed request %s as backend requested",
+                req_status.req.request_id,
+            )
+            return None, False
+
+        if delay_request:
+            logger.debug(
+                "Delaying request %s since some of its blocks are already being loaded",
+                req_status.req.request_id,
+            )
            return None, False
-        if hits == 0:
-            return 0, False

-        num_hit_tokens = (
-            group_config.offloaded_block_size * (start_block_idx + hits)
-            - num_computed_tokens
-        )
        logger.debug(
            "Request %s hit %s offloaded tokens after %s GPU hit tokens",
            request.request_id,
            num_hit_tokens,
            num_computed_tokens,
        )
-        if num_hit_tokens < group_config.offloaded_block_size:
-            return 0, False
-
-        if self._blocks_being_loaded and any(
-            key in self._blocks_being_loaded
-            for key in offload_keys[start_block_idx : start_block_idx + hits]
-        ):
-            # hit blocks are being loaded, delay request
-            logger.debug(
-                "Delaying request %s since some of its blocks are already being loaded",
-                request.request_id,
-            )
-            return None, False

        return num_hit_tokens, True

@@ -7,6 +7,7 @@ to avoid certain eager import breakage."""

 import importlib.metadata
 import sys
+from importlib.util import find_spec

 from vllm.logger import init_logger

@@ -34,47 +35,63 @@ def main():

    cli_env_setup()

-    # For 'vllm bench *': use CPU instead of UnspecifiedPlatform by default
-    if len(sys.argv) > 1 and sys.argv[1] == "bench":
-        logger.debug(
-            "Bench command detected, must ensure current platform is not "
-            "UnspecifiedPlatform to avoid device type inference error"
-        )
-        from vllm import platforms
-
-        if platforms.current_platform.is_unspecified():
-            from vllm.platforms.cpu import CpuPlatform
-
-            platforms.current_platform = CpuPlatform()
-            logger.info(
-                "Unspecified platform detected, switching to CPU Platform instead."
+    # If `--omni` arg is passed to the CLI, delegate to vLLM Omni's entrypoint handling
+    if "--omni" in sys.argv:
+        # NOTE: Check the spec instead of importing directly here, since things could
+        # fail with ImportError due to mismatched versions if things are moved around.
+        spec = find_spec("vllm_omni")
+        if spec is None:
+            logger.error(
+                "--omni flag requires a valid instance of vllm-omni to be installed."
            )
+            sys.exit(1)

-    parser = FlexibleArgumentParser(
-        description="vLLM CLI",
-        epilog=VLLM_SUBCMD_PARSER_EPILOG.format(subcmd="[subcommand]"),
-    )
-    parser.add_argument(
-        "-v",
-        "--version",
-        action="version",
-        version=importlib.metadata.version("vllm"),
-    )
-    subparsers = parser.add_subparsers(required=False, dest="subparser")
-    cmds = {}
-    for cmd_module in CMD_MODULES:
-        new_cmds = cmd_module.cmd_init()
-        for cmd in new_cmds:
-            cmd.subparser_init(subparsers).set_defaults(dispatch_function=cmd.cmd)
-            cmds[cmd.name] = cmd
-    args = parser.parse_args()
-    if args.subparser in cmds:
-        cmds[args.subparser].validate(args)
+        from vllm_omni.entrypoints.cli.main import main as omni_main

-    if hasattr(args, "dispatch_function"):
-        args.dispatch_function(args)
+        logger.info("Delegating entrypoint handling to vllm-omni")
+        omni_main()
    else:
-        parser.print_help()
+        # For 'vllm bench *': use CPU instead of UnspecifiedPlatform by default
+        if len(sys.argv) > 1 and sys.argv[1] == "bench":
+            logger.debug(
+                "Bench command detected, must ensure current platform is not "
+                "UnspecifiedPlatform to avoid device type inference error"
+            )
+            from vllm import platforms
+
+            if platforms.current_platform.is_unspecified():
+                from vllm.platforms.cpu import CpuPlatform
+
+                platforms.current_platform = CpuPlatform()
+                logger.info(
+                    "Unspecified platform detected, switching to CPU Platform instead."
+                )
+
+        parser = FlexibleArgumentParser(
+            description="vLLM CLI",
+            epilog=VLLM_SUBCMD_PARSER_EPILOG.format(subcmd="[subcommand]"),
+        )
+        parser.add_argument(
+            "-v",
+            "--version",
+            action="version",
+            version=importlib.metadata.version("vllm"),
+        )
+        subparsers = parser.add_subparsers(required=False, dest="subparser")
+        cmds = {}
+        for cmd_module in CMD_MODULES:
+            new_cmds = cmd_module.cmd_init()
+            for cmd in new_cmds:
+                cmd.subparser_init(subparsers).set_defaults(dispatch_function=cmd.cmd)
+                cmds[cmd.name] = cmd
+        args = parser.parse_args()
+        if args.subparser in cmds:
+            cmds[args.subparser].validate(args)
+
+        if hasattr(args, "dispatch_function"):
+            args.dispatch_function(args)
+        else:
+            parser.print_help()


 if __name__ == "__main__":
@@ -1166,19 +1166,16 @@ class LLM:

        if pooling_task is None:
            raise ValueError(
-                "pooling_task required for `LLM.encode`\n"
-                "Please use one of the more specific methods or set the "
-                "pooling_task when using `LLM.encode`:\n"
-                "  - For embeddings, use `LLM.embed(...)` "
-                'or `pooling_task="embed"`.\n'
-                "  - For classification logits, use `LLM.classify(...)` "
-                'or `pooling_task="classify"`.\n'
-                "  - For similarity scores, use `LLM.score(...)`.\n"
-                "  - For rewards, use `LLM.reward(...)` "
-                'or `pooling_task="token_classify"`\n'
-                "  - For token classification, "
-                'use `pooling_task="token_classify"`\n'
-                '  - For multi-vector retrieval, use `pooling_task="token_embed"`'
+                """
+                pooling_task required for `LLM.encode`.
+                Please use one of the more specific methods or set the pooling_task when using `LLM.encode`:
+                  - For embeddings, use `LLM.embed(...)` or `pooling_task="embed"`.
+                  - For classification logits, use `LLM.classify(...)` or `pooling_task="classify"`.
+                  - For similarity scores, use `LLM.score(...)`.
+                  - For rewards, `pooling_task="classify"` or `pooling_task="token_classify"`.
+                  - For token classification, use `pooling_task="token_classify"`.
+                  - For multi-vector retrieval, use `pooling_task="token_embed"`.
+                """  # noqa: E501
            )

        if (
@@ -1340,6 +1337,11 @@ class LLM:
            A list of `PoolingRequestOutput` objects containing the
            pooled hidden states in the same order as the input prompts.
        """
+        logger.warning_once(
+            "`llm.reward` api is deprecated and will be removed in v0.23. "
+            'Please use `LLM.encode` with `pooling_task="classify"` or '
+            '`pooling_task="token_classify"` instead.'
+        )
        return self.encode(
            prompts,
            use_tqdm=use_tqdm,
@@ -73,13 +73,9 @@ from vllm.reasoning import ReasoningParser
 from vllm.renderers import ChatParams
 from vllm.sampling_params import BeamSearchParams, SamplingParams
 from vllm.tokenizers import TokenizerLike
-from vllm.tool_parsers.mistral_tool_parser import (
-    MistralToolCall,
-    MistralToolParser,
-)
 from vllm.tool_parsers.utils import partial_json_loads
 from vllm.utils.collection_utils import as_list
-from vllm.utils.mistral import is_mistral_tokenizer
+from vllm.utils.mistral import is_mistral_tokenizer, is_mistral_tool_parser

 if TYPE_CHECKING:
    from vllm.entrypoints.serve.render.serving import OpenAIServingRender
@@ -143,10 +139,12 @@ class OpenAIServingChat(OpenAIServing):
            enable_auto_tools=enable_auto_tools,
            model_name=self.model_config.model,
        )
-        _is_mistral_tool_parser = self.tool_parser is not None and issubclass(
-            self.tool_parser, MistralToolParser
-        )
-        if _is_mistral_tool_parser and self.reasoning_parser_cls is not None:
+        if (
+            is_mistral_tool_parser(self.tool_parser)
+            and self.reasoning_parser_cls is not None
+        ):
+            from vllm.tool_parsers.mistral_tool_parser import MistralToolParser
+
            MistralToolParser.model_can_reason = True

        self.exclude_tools_when_tool_choice_none = exclude_tools_when_tool_choice_none
@@ -823,6 +821,10 @@ class OpenAIServingChat(OpenAIServing):
                        harmony_tools_streamed[i] |= tools_streamed_flag
                    # Mistral grammar path: combined reasoning + tool streaming
                    elif is_mistral_grammar_path:
+                        from vllm.tool_parsers.mistral_tool_parser import (
+                            MistralToolParser,
+                        )
+
                        assert tool_parser is not None
                        assert isinstance(tool_parser, MistralToolParser)
                        assert reasoning_end_arr is not None
@@ -904,6 +906,10 @@ class OpenAIServingChat(OpenAIServing):
                            else:
                                # Generate ID based on tokenizer type
                                if is_mistral_tokenizer(tokenizer):
+                                    from vllm.tool_parsers.mistral_tool_parser import (
+                                        MistralToolCall,
+                                    )
+
                                    tool_call_id = MistralToolCall.generate_random_id()
                                else:
                                    tool_call_id = make_tool_call_id(
@@ -1275,8 +1281,6 @@ class OpenAIServingChat(OpenAIServing):
        request_metadata: RequestResponseMetadata,
        reasoning_parser: ReasoningParser | None = None,
    ) -> ErrorResponse | ChatCompletionResponse:
-        from vllm.tokenizers.mistral import MistralTokenizer
-
        created_time = int(time.time())
        final_res: RequestOutput | None = None

@@ -1393,12 +1397,17 @@ class OpenAIServingChat(OpenAIServing):
                enable_auto_tools=self.enable_auto_tools,
                tool_parser_cls=self.tool_parser,
            )
-            tool_call_class = (
-                MistralToolCall if is_mistral_tokenizer(tokenizer) else ToolCall
-            )
+            if is_mistral_tokenizer(tokenizer):
+                from vllm.tool_parsers.mistral_tool_parser import MistralToolCall
+
+                tool_call_class: type[ToolCall] = MistralToolCall
+            else:
+                tool_call_class = ToolCall

            use_mistral_tool_parser = request._grammar_from_tool_parser
            if use_mistral_tool_parser:
+                from vllm.tool_parsers.mistral_tool_parser import MistralToolParser
+
                tool_call_items = MistralToolParser.build_non_streaming_tool_calls(
                    tool_calls
                )
@@ -1436,7 +1445,7 @@ class OpenAIServingChat(OpenAIServing):
                        # Generate ID using the correct format (kimi_k2 or random),
                        # but leave it to the class if it's Mistral to preserve
                        # 9-char IDs
-                        if isinstance(tokenizer, MistralTokenizer):
+                        if is_mistral_tokenizer(tokenizer):
                            tool_call_class_items.append(tool_call_class(function=tc))
                        else:
                            generated_id = make_tool_call_id(
@@ -1469,7 +1478,7 @@ class OpenAIServingChat(OpenAIServing):
                        # Generate ID using the correct format (kimi_k2 or random),
                        # but leave it to the class if it's Mistral to preserve
                        # 9-char IDs
-                        if isinstance(tokenizer, MistralTokenizer):
+                        if is_mistral_tokenizer(tokenizer):
                            tool_call_class_items.append(
                                tool_call_class(function=tool_call)
                            )
@@ -1519,7 +1528,7 @@ class OpenAIServingChat(OpenAIServing):
                            # Generate ID using the correct format (kimi_k2 or random),
                            # but leave it to the class if it's Mistral to preserve
                            # 9-char IDs
-                            if isinstance(tokenizer, MistralTokenizer):
+                            if is_mistral_tokenizer(tokenizer):
                                tool_call_items.append(tool_call_class(function=tc))
                            else:
                                generated_id = make_tool_call_id(
@@ -65,7 +65,6 @@ from vllm.renderers.inputs.preprocess import (
 from vllm.sampling_params import BeamSearchParams, SamplingParams
 from vllm.tokenizers import TokenizerLike
 from vllm.tool_parsers import ToolParser
-from vllm.tool_parsers.mistral_tool_parser import MistralToolParser
 from vllm.tracing import (
    contains_trace_headers,
    extract_trace_headers,
@@ -73,6 +72,7 @@ from vllm.tracing import (
 )
 from vllm.utils import random_uuid
 from vllm.utils.async_utils import collect_from_async_generator
+from vllm.utils.mistral import is_mistral_tool_parser

 logger = init_logger(__name__)

@@ -615,8 +615,7 @@ class OpenAIServing:
        # let the parser handle the output.
        use_mistral_tool_parser = (
            isinstance(request, ChatCompletionRequest)
-            and tool_parser_cls is not None
-            and issubclass(tool_parser_cls, MistralToolParser)
+            and is_mistral_tool_parser(tool_parser_cls)
            and request._grammar_from_tool_parser
        )

@@ -55,9 +55,8 @@ from vllm.renderers.inputs.preprocess import (
    prompt_to_seq,
 )
 from vllm.tool_parsers import ToolParser
-from vllm.tool_parsers.mistral_tool_parser import MistralToolParser
 from vllm.utils import random_uuid
-from vllm.utils.mistral import is_mistral_tokenizer
+from vllm.utils.mistral import is_mistral_tokenizer, is_mistral_tool_parser
 from vllm.utils.mistral import mt as _mt

 logger = init_logger(__name__)
@@ -137,7 +136,7 @@ class OpenAIServingRender:
                "Beam search is not supported by the render endpoint"
            )

-        result = await self.render_chat(request)
+        result = await self.render_chat(request, skip_mm_cache=True)
        if isinstance(result, ErrorResponse):
            return result

@@ -185,6 +184,8 @@ class OpenAIServingRender:
    async def render_chat(
        self,
        request: ChatCompletionRequest,
+        *,
+        skip_mm_cache: bool = False,
    ) -> tuple[list[ConversationMessage], list[EngineInput]] | ErrorResponse:
        """Core preprocessing logic for chat requests (no model/engine check).

@@ -253,7 +254,7 @@ class OpenAIServingRender:
                default_template_kwargs=self.default_chat_template_kwargs,
                tool_dicts=tool_dicts,
                tool_parser=tool_parser,
-                skip_mm_cache=True,
+                skip_mm_cache=skip_mm_cache,
                reasoning_parser=self.reasoning_parser,
            )
        else:
@@ -277,7 +278,7 @@ class OpenAIServingRender:
        error_check_ret = await self._check_model(request)
        if error_check_ret is not None:
            return error_check_ret
-        result = await self.render_completion(request)
+        result = await self.render_completion(request, skip_mm_cache=True)
        if isinstance(result, ErrorResponse):
            return result
        generate_requests: list[GenerateRequest] = []
@@ -323,6 +324,8 @@ class OpenAIServingRender:
    async def render_completion(
        self,
        request: CompletionRequest,
+        *,
+        skip_mm_cache: bool = False,
    ) -> list[EngineInput] | ErrorResponse:
        """Core preprocessing logic for completion requests (no model/engine check).

@@ -345,7 +348,7 @@ class OpenAIServingRender:
            request,
            prompt_input=request.prompt,
            prompt_embeds=request.prompt_embeds,
-            skip_mm_cache=True,
+            skip_mm_cache=skip_mm_cache,
        )

        return engine_inputs
@@ -582,7 +585,7 @@ class OpenAIServingRender:
            tool_choice = getattr(request, "tool_choice", "none")
            tokenizer = renderer.get_tokenizer()
            is_mistral_grammar_eligible = (
-                issubclass(tool_parser, MistralToolParser)
+                is_mistral_tool_parser(tool_parser)
                and is_mistral_tokenizer(tokenizer)
                and tokenizer.supports_grammar
            )
@@ -152,6 +152,10 @@ if TYPE_CHECKING:
    VLLM_RAY_EXTRA_ENV_VARS_TO_COPY: str = ""
    VLLM_MARLIN_USE_ATOMIC_ADD: bool = False
    VLLM_MARLIN_INPUT_DTYPE: Literal["int8", "fp8"] | None = None
+    VLLM_HUMMING_ONLINE_QUANT_CONFIG: dict[str, Any] | None = None
+    VLLM_HUMMING_INPUT_QUANT_CONFIG: dict[str, Any] | None = None
+    VLLM_HUMMING_USE_F16_ACCUM: bool = False
+    VLLM_HUMMING_MOE_GEMM_TYPE: Literal["indexed", "grouped", "auto"] | None = None
    VLLM_MXFP4_USE_MARLIN: bool | None = None
    VLLM_DEEPEPLL_NVFP4_DISPATCH: bool = False
    VLLM_V1_USE_OUTLINES_CACHE: bool = False
@@ -285,6 +289,15 @@ def maybe_convert_bool(value: str | None) -> bool | None:
    return bool(int(value))


+def maybe_convert_json_str_or_file(value: str | None) -> dict[str, Any] | None:
+    if value is None:
+        return None
+    if os.path.exists(value):
+        with open(value) as f:
+            return json.load(f)
+    return json.loads(value)
+
+
 def disable_compile_cache() -> bool:
    return bool(int(os.getenv("VLLM_DISABLE_COMPILE_CACHE", "0")))

@@ -1193,6 +1206,25 @@ environment_variables: dict[str, Callable[[], Any]] = {
    "VLLM_MARLIN_INPUT_DTYPE": env_with_choices(
        "VLLM_MARLIN_INPUT_DTYPE", None, ["int8", "fp8"]
    ),
+    # The online quantization dtype for humming kernel
+    "VLLM_HUMMING_ONLINE_QUANT_CONFIG": lambda: maybe_convert_json_str_or_file(
+        os.environ.get("VLLM_HUMMING_ONLINE_QUANT_CONFIG", None)
+    ),
+    # The activation dtype config for humming kernel
+    "VLLM_HUMMING_INPUT_QUANT_CONFIG": lambda: maybe_convert_json_str_or_file(
+        os.environ.get("VLLM_HUMMING_INPUT_QUANT_CONFIG", None)
+    ),
+    # Whether to use fp16 accumulator mma
+    "VLLM_HUMMING_USE_F16_ACCUM": lambda: maybe_convert_bool(
+        os.environ.get("VLLM_HUMMING_USE_F16_ACCUM", "0")
+    ),
+    # Whether to use indexed gemm for humming moe
+    # if 1, force use indexed gemm
+    # if 0, force use grouped gemm
+    # if None, choose better gemm type automatically
+    "VLLM_HUMMING_MOE_GEMM_TYPE": lambda: maybe_convert_bool(
+        os.environ.get("VLLM_HUMMING_MOE_GEMM_TYPE", None)
+    ),
    # Whether to use DeepEPLL kernels for NVFP4 quantization and dispatch method
    # only supported on Blackwell GPUs and with
    # https://github.com/deepseek-ai/DeepEP/pull/341
@@ -597,6 +597,7 @@ class Attention(nn.Module, AttentionLayerBase):
                block_size=block_size,
                num_kv_heads=self.num_kv_heads,
                head_size=self.head_size,
+                head_size_v=self.head_size_v,
                dtype=self.kv_cache_torch_dtype,
                kv_quant_mode=quant_mode,
                sliding_window=self.sliding_window,
@@ -90,15 +90,23 @@ def create_cross_attention_backend(
            assert new_metadata.encoder_seq_lens_cpu is not None
            max_encoder_len = int(new_metadata.encoder_seq_lens_cpu.max())
            new_metadata.max_seq_len = max_encoder_len
-            # Any computed tokens indicated decode step>1 (no chunked prefill)
-            num_cache_decodes = (
-                (common_attn_metadata.num_computed_tokens_cpu > 0).sum().item()
+            # Any computed tokens indicates decode step>1 (no chunked prefill).
+            # The upper bound is exact for this `> 0` test - prefill rows have
+            # num_computed == 0 and decode rows have num_computed > 0.
+            query_lens_cpu = (
+                common_attn_metadata.query_start_loc_cpu[1:]
+                - common_attn_metadata.query_start_loc_cpu[:-1]
            )
+            assert common_attn_metadata.seq_lens_cpu_upper_bound is not None
+            num_computed_tokens_cpu = (
+                common_attn_metadata.seq_lens_cpu_upper_bound - query_lens_cpu
+            )
+            num_cache_decodes = (num_computed_tokens_cpu > 0).sum().item()
            if num_cache_decodes > 0:
                # CrossAttn KV cache has already been populated on first decoder step,
                # skip slot_mapping calculation for requests that do not need
                # reshape_and_cache.
-                num_tokens = common_attn_metadata.num_computed_tokens_cpu.numpy()
+                num_tokens = num_computed_tokens_cpu.numpy()
                new_metadata.encoder_seq_lens_cpu = np.where(
                    num_tokens > 0, 0, new_metadata.encoder_seq_lens_cpu
                )
@@ -1822,13 +1822,18 @@ class MLACommonMetadataBuilder(AttentionMetadataBuilder[M]):

        prefill_metadata = None
        if num_prefills > 0:
-            num_computed_tokens_cpu = (
-                common_attn_metadata.compute_num_computed_tokens().cpu()
-            )
-
            reqs_start = num_decodes  # prefill_start

-            context_lens_cpu = num_computed_tokens_cpu[reqs_start:num_reqs]
+            # Upper bound is exact for prefill rows (no D2H sync).
+            seq_lens_cpu = common_attn_metadata.seq_lens_cpu_upper_bound
+            assert seq_lens_cpu is not None
+            prefill_query_lens_cpu = (
+                query_start_loc_cpu[reqs_start + 1 : num_reqs + 1]
+                - query_start_loc_cpu[reqs_start:num_reqs]
+            )
+            context_lens_cpu = (
+                seq_lens_cpu[reqs_start:num_reqs] - prefill_query_lens_cpu
+            )
            max_context_len_cpu = context_lens_cpu.max().item()
            num_prefills_with_context_cpu = (context_lens_cpu > 0).sum().item()
            prefill_query_start_loc = (
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Jee Jee Li	931b3f6110	Done Signed-off-by: Jee Jee Li <pandaleefree@gmail.com>	2026-04-25 13:23:05 +00:00
Andreas KaratzasandGitHub	95995bbef8	[ROCm][Engine] Fix GPU memory leaks in engine shutdown and test workaround for async KV prefix cache reset (#38503 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-04-25 05:25:20 +00:00
Chenguang Zheng GitHub OpenAI Codex	07351e0883	[Feature] Warm up readonly multimodal processor during renderer startup (#40797 ) Signed-off-by: Chenguang ZHENG <645327136@qq.com> Co-authored-by: OpenAI Codex <codex@openai.com>	2026-04-25 03:57:41 +00:00
Andreas KaratzasandGitHub	428b988c98	[ROCm][CI] Fix `trust_remote_code` AttributeError in EAGLE3 acceptance length test (#40306 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-04-25 02:59:31 +00:00
Andreas KaratzasandGitHub	e54894fc85	[ROCm][CI] Fix TestSiluMulGroupFp8QuantModel after W8A8 block linear refactor (#39799 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-04-25 11:20:59 +09:00
Angela YiandGitHub	bc2ae5a3d6	[Test] Increase qwen2_vl num_logprobs to fix torch 2.12 update (#40818 ) Signed-off-by: Angela Yi <angelayi@meta.com>	2026-04-25 00:59:20 +00:00
Wentao YeandGitHub	a474da2813	[Refactor] Remove unused dead code (#40640 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-04-25 07:28:18 +08:00
Lucas KabelaandGitHub	ce6a199ecc	[BE][Bugfix] Respect TORCH_COMPILE_DISABLE env var at the vLLM config level for torch 2.12 (#40715 ) Signed-off-by: Lucas Kabela <lucaskabela@meta.com>	2026-04-24 16:25:03 -07:00
Ignacio SicaandGitHub	f88763efc3	[Bugfix] add seq_lens_cpu_upper_bound to CommonAttentionMetadata in mla_runner.py (#40844 ) Signed-off-by: ignaciosica <mignacio.sica@gmail.com>	2026-04-24 23:13:52 +00:00
Artem PerevedentsevandGitHub	333529deae	[EPLB] Fix replica selection bias in fused_moe router (#40810 ) Signed-off-by: Artem Perevedentsev <aperevedents@nvidia.com>	2026-04-24 22:06:41 +00:00
Zhang JianandGitHub	8825608205	[Bugfix][CI] Fix wrong residual shape in TestFusedAddRMSNorm.example_inputs that causes flaky test (#40629 ) Signed-off-by: Zhang Jian <jianmusings@gmail.com>	2026-04-24 16:40:07 -04:00
qli88andGitHub	095d2f87e8	[Bug] Fix GLM-5.1 running error on ROCm platform (#40763 ) Signed-off-by: Qiang Li <qiang.li2@amd.com>	2026-04-24 19:54:40 +00:00
Neil Schemenauer GitHub Simon Mo	21792520e7	[Build] Add Python 3.14 to supported version list. (#34770 ) Signed-off-by: Neil Schemenauer <nas@arctrix.com> Co-authored-by: Simon Mo <simon.mo@hey.com>	2026-04-24 10:24:05 -07:00
Alex BrooksandGitHub	5e11b40365	[Frontend] Delegate to vLLM Omni When `--omni` Passed (#40744 ) Signed-off-by: Alex Brooks <albrooks@redhat.com>	2026-04-24 12:30:00 -04:00
labAxiaoming GitHub Cyrus Leung	f768b4473e	[Docs] Add docs for context extension using the yarn method (#37430 ) Signed-off-by: xiaoming <1259730330@qq.com> Signed-off-by: labAxiaoming <34019940+labAxiaoming@users.noreply.github.com> Co-authored-by: Cyrus Leung <cyrus.tl.leung@gmail.com>	2026-04-24 08:26:09 -07:00
JartXandGitHub	914d0464c1	[Refactor] Unify 2D/3D kernels in triton_unified_attention (#40631 ) Signed-off-by: JartX <sagformas@epdcenter.es>	2026-04-24 17:18:06 +02:00
Jinzhen LinandGitHub	9f771b3ab9	[Quantization] add humming quantization kernel (#34556 )	2026-04-24 09:29:44 -04:00
Itay AlroyandGitHub	c9d3c6e6af	fused_moe: treat NIXL EP as batched experts (#40412 ) Signed-off-by: Itay Alroy <ialroy@nvidia.com>	2026-04-24 08:05:31 -05:00
Or OzeriandGitHub	51adca74e6	[kv_offload+HMA][9/N]: Support lookup with multiple KV groups (#39401 ) Signed-off-by: Or Ozeri <oro@il.ibm.com>	2026-04-24 15:32:29 +03:00
Netanel HaberandGitHub	e8eb0490ce	[Bugfix][MoE] Unpad routed output before shared expert add [Fixes #35949 ] (#40794 ) Signed-off-by: Netanel Haber <nhaber@nvidia.com>	2026-04-24 11:53:23 +00:00
Jiangyun ZhuandGitHub	e8ee2a78db	[Attention] use diff kv backend for mimo v2 flash (#40045 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>	2026-04-24 11:25:55 +00:00
Thomas GitHub thomasmaindron Claude Opus 4.6 Chauncey	2ec18f5df4	[Bugfix][Parser] Fix Mistral tool parser for HF tokenizers (#39294 ) Signed-off-by: thomasmaindron <thomasmaindron@users.noreply.github.com> Co-authored-by: thomasmaindron <thomasmaindron@users.noreply.github.com> Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com>	2026-04-24 19:01:56 +08:00
Dmitry TokarevandGitHub	6dec49f27e	[Build] Bump CUDA to 13.0.2 to match PyTorch 2.11.0 (#40669 ) Signed-off-by: Dmitry Tokarev <dtokarev@nvidia.com>	2026-04-24 10:27:11 +00:00
Shanshan ShenandGitHub	b5587e1013	[CI/Build] Add e2e test for ViT CUDA graph (#40780 ) Signed-off-by: shen-shanshan <467638484@qq.com>	2026-04-24 18:12:14 +08:00
milesialandGitHub	9ad5abe772	Fix Nano Nemotron VL static image inputs (#40724 ) Signed-off-by: Alexandre Milesi <milesial@users.noreply.github.com>	2026-04-24 09:18:55 +00:00
Woosuk KwonandGitHub	7d3195ea9f	[Bugfix] Fix IMA in DSA + MTP (#40772 ) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>	2026-04-24 01:40:20 -07:00
Luciano Martins GitHub Luciano Martins Isotr0py Isotr0py	512f522192	[Model] Gemma4: add bidirectional vision attention for sliding layers with window guard (#40534 ) Signed-off-by: Luciano Martins <lucianommartins@users.noreply.github.com> Signed-off-by: Luciano Martins <lucianomartins@google.com> Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: Luciano Martins <lucianommartins@users.noreply.github.com> Co-authored-by: Isotr0py <2037008807@qq.com> Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn>	2026-04-24 08:27:46 +00:00
Yuwen Zhou GitHub Kunshang Ji	4c34b2f6fc	[XPU] Enable torch.compile for XPU GDN attention (#39466 ) Signed-off-by: yuwenzho <yuwen.zhou@intel.com> Signed-off-by: Yuwen Zhou <yuwen.zhou@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-04-24 16:26:16 +08:00
Xin YangandGitHub	cf8a613a87	Support only half types for concat_mla_q kernel (#37892 ) Signed-off-by: Xin Yang <xyangx@amazon.com>	2026-04-23 23:51:05 -07:00
xiangdongandGitHub	01acf96c6f	[XPU][CI] Fix Docker cleanup races on Intel CI runners (#40761 ) Signed-off-by: zengxian <xiangdong.zeng@intel.com>	2026-04-24 14:08:45 +08:00
Jackmin801 GitHub Claude	079a4cf399	[MoE] Move cutlass moe to fused_moe/experts/ (#40574 ) Signed-off-by: Jackmin801 <ongjackm@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-04-24 06:05:49 +00:00
wang.yuqi GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>Cyrus Leung	9744b699ba	[Deprecate] Deprecate LLM.reward offline api, use LLM.encode instead. (#40688 ) Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io> Signed-off-by: wang.yuqi <noooop@126.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> Co-authored-by: Cyrus Leung <cyrus.tl.leung@gmail.com>	2026-04-24 05:37:50 +00:00
Bugen Zhao GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	c662b4359e	[Bugfix] Avoid mutating `chat_template_kwargs` in `HYV3ReasoningParser` initialization (#40713 ) Signed-off-by: Bugen Zhao <i@bugenzhao.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-04-24 13:08:58 +08:00
lyd1992andGitHub	100c7b65e7	[Platform] Fix RISC-V platform detection (lscpu parsing + non-NUMA meminfo) (#40427 ) Signed-off-by: liuyudong <liuyudong@iscas.ac.cn>	2026-04-24 04:33:05 +00:00
Neil SchemenauerandGitHub	56bdf85e10	[Feature] Avoid eager import of the "mistral_common" package. (#40043 ) Signed-off-by: Neil Schemenauer <nas@arctrix.com>	2026-04-24 02:49:16 +00:00
Vinayak KumarandGitHub	eba73068ea	[Doc] fix capitalization consistency in README (vLLM, Hugging Face) (#40729 ) Signed-off-by: Vinayak Mishra <vinayakmishra448@gmail.com>	2026-04-24 02:23:54 +00:00
Nick HillandGitHub	e9f331d72e	[MRV2] Ensure warmup covers prefill path (#40746 ) Signed-off-by: Nick Hill <nickhill123@gmail.com>	2026-04-24 01:33:26 +00:00
yzong-rh GitHub timon0305 sfeng33	c9bf77df92	[BUG]: fix HF tokenizer concurrent borrow in tool parsers (#40059 ) Signed-off-by: Yifan <yzong@redhat.com> Co-authored-by: timon0305 <timon0305@outlook.com> Co-authored-by: sfeng33 <4florafeng@gmail.com>	2026-04-23 18:20:30 -07:00
Dmitry Tokarev GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	3041344287	[Misc] Added curl retries in install_python_libraries.sh (#36700 ) Signed-off-by: Dmitry Tokarev <dtokarev@nvidia.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-04-24 01:19:30 +00:00
Doug CamposandGitHub	92762edc53	[Bugfix] Treat <tool_call> as implicit reasoning end in Qwen3 parser (#35687 ) Signed-off-by: Doug Campos <qmx@qmx.me>	2026-04-24 09:10:04 +08:00
Benjamin Chislett GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	626daa2076	[Feat] Unified Synthetic Acceptance Rate for V1 and V2 (#40662 ) Signed-off-by: Benjamin Chislett <bchislett@nvidia.com> Signed-off-by: Benjamin Chislett <chislett.ben@gmail.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-04-24 00:48:08 +00:00
Nick HillandGitHub	fe85a92e86	[Core] Avoid seq_lens_cpu GPU->CPU sync (#40654 ) Signed-off-by: Nick Hill <nickhill123@gmail.com>	2026-04-24 00:35:55 +00:00
Sage MooreandGitHub	62b1bbe470	[EPLB] Remove asyncio infrastructure from Async EPLB (#40730 ) Signed-off-by: Sage Moore <sage@neuralmagic.com>	2026-04-24 00:21:15 +00:00
Hemanth AcharyaandGitHub	fa4b70555b	[ROCm] Cast score correction bias tensor during model construction for DeepSeek/Kimi-K2 (#39999 ) Signed-off-by: Hemanth Acharya <heachary@amd.com>	2026-04-24 09:02:12 +09:00
Jackmin801 GitHub Robert Shaw	447c372ac5	[MoE] Move remaining PrepareAndFinalize to prepare finalize folder (#39009 ) Signed-off-by: Robert Shaw <robertgshaw2@gmail.com> Signed-off-by: Jackmin801 <ongjackm@gmail.com> Co-authored-by: Robert Shaw <robertgshaw2@gmail.com>	2026-04-23 20:00:53 -04:00
Sophie du Couédic GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	ff2c2bd80a	[Docs]Add documentation for bench serve visualization arguments (#40539 ) Signed-off-by: Sophie du Couédic <sop@zurich.ibm.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-04-23 15:48:29 -07:00
Matthew BonanniandGitHub	cde8d24710	[Spec Decode] Move `SpecDecodeBaseProposer` out of `eagle.py` (#40732 ) Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>	2026-04-23 22:28:27 +00:00
bnellnmandGitHub	4a6dd1c3cc	[Bugfix] Fix DeepSeek V2-Lite Accuracy drop (#40673 ) Signed-off-by: Bill Nell <bnell@redhat.com>	2026-04-23 18:11:37 -04:00
czhu-cohere GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	7ff65b1900	[Bugfix] Fix workspace resize leaking reserved GPU memory (#39226 ) Signed-off-by: root <conway.zhu@cohere.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-04-23 20:50:05 +00:00