p

Signed-off-by: khluu <khluu000@gmail.com>

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

@@ -40,16 +40,16 @@ docker run \
     python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager -tp 2 --distributed-executor-backend mp
     python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --attention-backend=TRITON_ATTN
     python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --quantization fp8
-    python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4  --block-size 64 --enforce-eager
+    python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4  --block-size 64 --enforce-eager --max-model-len 8192
     python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b  --block-size 64 --enforce-eager -tp 2
     python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b  --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel
     cd tests
     pytest -v -s v1/core --ignore=v1/core/test_reset_prefix_cache_e2e.py --ignore=v1/core/test_scheduler_e2e.py
     pytest -v -s v1/engine
     pytest -v -s v1/sample --ignore=v1/sample/test_logprobs.py --ignore=v1/sample/test_logprobs_e2e.py
-    pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py
+    pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py
     pytest -v -s v1/structured_output
     pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_tree_attention.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py
-    pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_nixl_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py
+    pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_nixl_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py -k "not (test_register_kv_caches and FLASH_ATTN and True)"
     pytest -v -s v1/test_serial_utils.py
 '

.buildkite/test_areas/distributed.yaml

@@ -15,8 +15,29 @@ steps:
   - pytest -v -s distributed/test_shm_buffer.py
   - pytest -v -s distributed/test_shm_storage.py
 
-- label: Distributed (2 GPUs)
-  timeout_in_minutes: 60
+- label: Distributed DP Tests (2 GPUs)
+  timeout_in_minutes: 20
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 2
+  source_file_dependencies:
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/executor/
+  - vllm/worker/worker_base.py
+  - vllm/v1/engine/
+  - vllm/v1/worker/
+  - tests/v1/distributed
+  - tests/v1/entrypoints/openai/test_multi_api_servers.py
+  commands:
+  # https://github.com/NVIDIA/nccl/issues/1838
+  - export NCCL_CUMEM_HOST_ENABLE=0
+  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
+  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
+  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
+  - DP_SIZE=2 pytest -v -s v1/entrypoints/openai/test_multi_api_servers.py
+
+- label: Distributed Compile + RPC Tests (2 GPUs)
+  timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -29,22 +50,31 @@ steps:
   - vllm/v1/worker/
   - tests/compile/fullgraph/test_basic_correctness.py
   - tests/compile/test_wrapper.py
-  - tests/distributed/
   - tests/entrypoints/llm/test_collective_rpc.py
-  - tests/v1/distributed
-  - tests/v1/entrypoints/openai/test_multi_api_servers.py
+  commands:
+  # https://github.com/NVIDIA/nccl/issues/1838
+  - export NCCL_CUMEM_HOST_ENABLE=0
+  - pytest -v -s entrypoints/llm/test_collective_rpc.py
+  - pytest -v -s ./compile/fullgraph/test_basic_correctness.py
+  - pytest -v -s ./compile/test_wrapper.py
+
+- label: Distributed Torchrun + Shutdown Tests (2 GPUs)
+  timeout_in_minutes: 20
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 2
+  source_file_dependencies:
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/executor/
+  - vllm/worker/worker_base.py
+  - vllm/v1/engine/
+  - vllm/v1/worker/
+  - tests/distributed/
   - tests/v1/shutdown
   - tests/v1/worker/test_worker_memory_snapshot.py
   commands:
   # https://github.com/NVIDIA/nccl/issues/1838
   - export NCCL_CUMEM_HOST_ENABLE=0
-  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
-  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
-  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
-  - DP_SIZE=2 pytest -v -s v1/entrypoints/openai/test_multi_api_servers.py
-  - pytest -v -s entrypoints/llm/test_collective_rpc.py
-  - pytest -v -s ./compile/fullgraph/test_basic_correctness.py
-  - pytest -v -s ./compile/test_wrapper.py
   - VLLM_TEST_SAME_HOST=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - VLLM_TEST_SAME_HOST=1 VLLM_TEST_WITH_DEFAULT_DEVICE_SET=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s v1/shutdown

.buildkite/test_areas/kernels.yaml

@@ -35,7 +35,7 @@ steps:
   parallelism: 2
 
 - label: Kernels MoE Test %N
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/
   - csrc/moe/
@@ -47,7 +47,7 @@ steps:
   commands:
     - pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
     - pytest -v -s kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-  parallelism: 2
+  parallelism: 5
 
 - label: Kernels Mamba Test
   timeout_in_minutes: 45

.buildkite/test_areas/pytorch.yaml

@@ -35,7 +35,7 @@ steps:
   # as it is a heavy test that is covered in other steps.
   # Use `find` to launch multiple instances of pytest so that
   # they do not suffer from https://github.com/vllm-project/vllm/issues/28965
-  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -exec pytest -s -v {} \\;"
+  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
 
 - label: PyTorch Fullgraph
   timeout_in_minutes: 30

benchmarks/kernels/cpu/benchmark_cpu_attn.py

@@ -27,7 +27,7 @@ def get_attn_isa(
     else:
         if current_platform.get_cpu_architecture() == CpuArchEnum.ARM:
             return "neon"
-        elif torch._C._cpu._is_amx_tile_supported():
+        elif torch.cpu._is_amx_tile_supported():
             return "amx"
         else:
             return "vec"

benchmarks/kernels/cpu/benchmark_cpu_fused_moe.py

@@ -24,7 +24,7 @@ except (ImportError, AttributeError) as e:
     sys.exit(1)
 
 # ISA selection following test_cpu_fused_moe.py pattern
-ISA_CHOICES = ["amx", "vec"] if torch._C._cpu._is_amx_tile_supported() else ["vec"]
+ISA_CHOICES = ["amx", "vec"] if torch.cpu._is_amx_tile_supported() else ["vec"]
 
 
 @torch.inference_mode()

cmake/external_projects/vllm_flash_attn.cmake

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

csrc/ops.h

@@ -262,7 +262,8 @@ void get_cutlass_moe_mm_data(
     torch::Tensor& problem_sizes1, torch::Tensor& problem_sizes2,
     torch::Tensor& input_permutation, torch::Tensor& output_permutation,
     const int64_t num_experts, const int64_t n, const int64_t k,
-    const std::optional<torch::Tensor>& blockscale_offsets);
+    const std::optional<torch::Tensor>& blockscale_offsets,
+    const bool is_gated);
 
 void get_cutlass_moe_mm_problem_sizes_from_expert_offsets(
     const torch::Tensor& expert_first_token_offset,

csrc/quantization/w8a8/cutlass/moe/moe_data.cu

@@ -17,8 +17,11 @@ __global__ void compute_problem_sizes(const int32_t* __restrict__ topk_ids,
                                       int32_t* problem_sizes2,
                                       int32_t* atomic_buffer,
                                       const int topk_length, const int n,
-                                      const int k) {
+                                      const int k, const bool is_gated) {
   int expert_id = blockIdx.x;
+  // For gated activations (gate + up), first GEMM output is 2*n.
+  // For non-gated activations (up only), first GEMM output is n.
+  int const n1 = is_gated ? 2 * n : n;
 
   int occurrences = 0;
   for (int i = threadIdx.x; i < topk_length; i += THREADS_PER_EXPERT) {
@@ -31,13 +34,13 @@ __global__ void compute_problem_sizes(const int32_t* __restrict__ topk_ids,
     int final_occurrences = atomic_buffer[expert_id];
     if constexpr (!SWAP_AB) {
       problem_sizes1[expert_id * 3] = final_occurrences;
-      problem_sizes1[expert_id * 3 + 1] = 2 * n;
+      problem_sizes1[expert_id * 3 + 1] = n1;
       problem_sizes1[expert_id * 3 + 2] = k;
       problem_sizes2[expert_id * 3] = final_occurrences;
       problem_sizes2[expert_id * 3 + 1] = k;
       problem_sizes2[expert_id * 3 + 2] = n;
     } else {
-      problem_sizes1[expert_id * 3] = 2 * n;
+      problem_sizes1[expert_id * 3] = n1;
       problem_sizes1[expert_id * 3 + 1] = final_occurrences;
       problem_sizes1[expert_id * 3 + 2] = k;
       problem_sizes2[expert_id * 3] = k;
@@ -107,13 +110,11 @@ __global__ void compute_arg_sorts(const int32_t* __restrict__ topk_ids,
 }
 
 namespace {
-inline void launch_compute_problem_sizes(const torch::Tensor& topk_ids,
-                                         torch::Tensor& problem_sizes1,
-                                         torch::Tensor& problem_sizes2,
-                                         torch::Tensor& atomic_buffer,
-                                         int64_t num_experts, int64_t n,
-                                         int64_t k, cudaStream_t stream,
-                                         const bool swap_ab) {
+inline void launch_compute_problem_sizes(
+    const torch::Tensor& topk_ids, torch::Tensor& problem_sizes1,
+    torch::Tensor& problem_sizes2, torch::Tensor& atomic_buffer,
+    int64_t num_experts, int64_t n, int64_t k, cudaStream_t stream,
+    const bool swap_ab, const bool is_gated) {
   int num_threads = min(THREADS_PER_EXPERT, topk_ids.numel());
 
   auto const* topk_ptr = topk_ids.data_ptr<int32_t>();
@@ -125,7 +126,7 @@ inline void launch_compute_problem_sizes(const torch::Tensor& topk_ids,
     compute_problem_sizes<SwapAB><<<num_experts, num_threads, 0, stream>>>(
         topk_ptr, ps1_ptr, ps2_ptr, atomic_ptr,
         static_cast<int>(topk_ids.numel()), static_cast<int>(n),
-        static_cast<int>(k));
+        static_cast<int>(k), is_gated);
   });
 }
 }  // namespace
@@ -222,7 +223,8 @@ void get_cutlass_moe_mm_data_caller(
     torch::Tensor& problem_sizes1, torch::Tensor& problem_sizes2,
     torch::Tensor& input_permutation, torch::Tensor& output_permutation,
     const int64_t num_experts, const int64_t n, const int64_t k,
-    const std::optional<torch::Tensor>& blockscale_offsets) {
+    const std::optional<torch::Tensor>& blockscale_offsets,
+    const bool is_gated) {
   auto stream = at::cuda::getCurrentCUDAStream(topk_ids.device().index());
   auto options_int32 =
       torch::TensorOptions().dtype(torch::kInt32).device(topk_ids.device());
@@ -236,7 +238,7 @@ void get_cutlass_moe_mm_data_caller(
 
   launch_compute_problem_sizes(topk_ids, problem_sizes1, problem_sizes2,
                                atomic_buffer, num_experts, n, k, stream,
-                               may_swap_ab);
+                               may_swap_ab, is_gated);
 
   if (blockscale_offsets.has_value()) {
     // fp4 path

csrc/quantization/w8a8/cutlass/scaled_mm_entry.cu

@@ -75,7 +75,8 @@ void get_cutlass_moe_mm_data_caller(
     torch::Tensor& problem_sizes1, torch::Tensor& problem_sizes2,
     torch::Tensor& input_permutation, torch::Tensor& output_permutation,
     const int64_t num_experts, const int64_t n, const int64_t k,
-    const std::optional<torch::Tensor>& blockscale_offsets);
+    const std::optional<torch::Tensor>& blockscale_offsets,
+    const bool is_gated);
 
 void get_cutlass_moe_mm_problem_sizes_from_expert_offsets_caller(
     const torch::Tensor& expert_first_token_offset,
@@ -278,7 +279,8 @@ void get_cutlass_moe_mm_data(
     torch::Tensor& problem_sizes1, torch::Tensor& problem_sizes2,
     torch::Tensor& input_permutation, torch::Tensor& output_permutation,
     const int64_t num_experts, const int64_t n, const int64_t k,
-    const std::optional<torch::Tensor>& blockscale_offsets) {
+    const std::optional<torch::Tensor>& blockscale_offsets,
+    const bool is_gated) {
   // This function currently gets compiled only if we have a valid cutlass moe
   // mm to run it for.
   int32_t version_num = get_sm_version_num();
@@ -288,7 +290,7 @@ void get_cutlass_moe_mm_data(
   get_cutlass_moe_mm_data_caller(topk_ids, expert_offsets, problem_sizes1,
                                  problem_sizes2, input_permutation,
                                  output_permutation, num_experts, n, k,
-                                 blockscale_offsets);
+                                 blockscale_offsets, is_gated);
   return;
 #endif
   TORCH_CHECK_NOT_IMPLEMENTED(

csrc/torch_bindings.cpp

@@ -489,8 +489,8 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
       "                        Tensor! problem_sizes1, Tensor! problem_sizes2, "
       "                        Tensor! input_permutation, "
       "                        Tensor! output_permutation, int num_experts, "
-      "                        int n, int k, Tensor? blockscale_offsets) -> "
-      "()");
+      "                        int n, int k, Tensor? blockscale_offsets, "
+      "                        bool is_gated) -> ()");
   ops.impl("get_cutlass_moe_mm_data", torch::kCUDA, &get_cutlass_moe_mm_data);
 
   // compute per-expert problem sizes from expert_first_token_offset

docs/design/attention_backends.md

@@ -127,8 +127,8 @@ Priority is **1 = highest** (tried first).
 | 3 | `FLASH_ATTN_MLA` |
 | 4 | `FLASHMLA` |
 | 5 | `TRITON_MLA` |
-| 6 | `FLASHMLA_SPARSE` |
-| 7 | `FLASHINFER_MLA_SPARSE` |
+| 6 | `FLASHINFER_MLA_SPARSE`**\*** |
+| 7 | `FLASHMLA_SPARSE` |
 
 **Ampere/Hopper (SM 8.x-9.x):**
 
@@ -140,6 +140,8 @@ Priority is **1 = highest** (tried first).
 | 4 | `TRITON_MLA` |
 | 5 | `FLASHMLA_SPARSE` |
 
+> **\*** For sparse MLA, FP8 KV cache always prefers `FLASHINFER_MLA_SPARSE`. With BF16 KV cache, `FLASHINFER_MLA_SPARSE` is preferred for low query-head counts (<= 16), while `FLASHMLA_SPARSE` is preferred otherwise.
+>
 > **Note:** ROCm and CPU platforms have their own selection logic. See the platform-specific documentation for details.
 
 ## Legend

docs/features/lora.md

@@ -389,3 +389,17 @@ vllm serve model --enable-lora --max-lora-rank 64
 # Bad: unnecessarily high, wastes memory
 vllm serve model --enable-lora --max-lora-rank 256
 ```
+
+### Restricting LoRA to Specific Modules
+
+The `--lora-target-modules` parameter allows you to restrict which model modules have LoRA applied at deployment time. This is useful for performance tuning when you only need LoRA on specific layers:
+
+```bash
+# Apply LoRA only to output projection layers
+vllm serve model --enable-lora --lora-target-modules o_proj
+
+# Apply LoRA to multiple specific modules
+vllm serve model --enable-lora --lora-target-modules o_proj qkv_proj down_proj
+```
+
+When `--lora-target-modules` is not specified, LoRA will be applied to all supported modules in the model. This parameter accepts module suffixes (the last component of the module name), such as `o_proj`, `qkv_proj`, `gate_proj`, etc.

docs/models/pooling_models.md

@@ -625,6 +625,46 @@ curl -s http://localhost:8000/rerank -H "Content-Type: application/json" -d '{
 }'
 ```
 
+### ColQwen3.5 Multi-Modal Late Interaction Models
+
+ColQwen3.5 is based on [ColPali](https://arxiv.org/abs/2407.01449), extending ColBERT's late interaction approach to **multi-modal** inputs. It uses the Qwen3.5 hybrid backbone (linear + full attention) and produces per-token L2-normalized vectors for MaxSim scoring.
+
+| Architecture | Backbone | Example HF Models |
+| - | - | - |
+| `ColQwen3_5` | Qwen3.5 | `athrael-soju/colqwen3.5-4.5B` |
+
+Start the server:
+
+```shell
+vllm serve athrael-soju/colqwen3.5-4.5B --max-model-len 4096
+```
+
+Then you can use the rerank endpoint:
+
+```shell
+curl -s http://localhost:8000/rerank -H "Content-Type: application/json" -d '{
+    "model": "athrael-soju/colqwen3.5-4.5B",
+    "query": "What is machine learning?",
+    "documents": [
+        "Machine learning is a subset of artificial intelligence.",
+        "Python is a programming language.",
+        "Deep learning uses neural networks."
+    ]
+}'
+```
+
+Or the score endpoint:
+
+```shell
+curl -s http://localhost:8000/score -H "Content-Type: application/json" -d '{
+    "model": "athrael-soju/colqwen3.5-4.5B",
+    "text_1": "What is the capital of France?",
+    "text_2": ["The capital of France is Paris.", "Python is a programming language."]
+}'
+```
+
+An example can be found here: [examples/pooling/score/colqwen3_5_rerank_online.py](../../examples/pooling/score/colqwen3_5_rerank_online.py)
+
 ### BAAI/bge-m3
 
 The `BAAI/bge-m3` model comes with extra weights for sparse and colbert embeddings but unfortunately in its `config.json`

docs/models/supported_models.md

@@ -834,6 +834,7 @@ The following table lists those that are tested in vLLM.
 | `CLIPModel` | CLIP | T / I | `openai/clip-vit-base-patch32`, `openai/clip-vit-large-patch14`, etc. | | |
 | `ColModernVBertForRetrieval` | ColModernVBERT | T / I | `ModernVBERT/colmodernvbert-merged` | | |
 | `ColPaliForRetrieval` | ColPali | T / I | `vidore/colpali-v1.3-hf` | | |
+| `ColQwen3_5` | ColQwen3.5 | T + I + V | `athrael-soju/colqwen3.5-4.5B-v3` | | |
 | `LlamaNemotronVLModel` | Llama Nemotron Embedding + SigLIP | T + I | `nvidia/llama-nemotron-embed-vl-1b-v2` | | |
 | `LlavaNextForConditionalGeneration`<sup>C</sup> | LLaVA-NeXT-based | T / I | `royokong/e5-v` | | ✅︎ |
 | `Phi3VForCausalLM`<sup>C</sup> | Phi-3-Vision-based | T + I | `TIGER-Lab/VLM2Vec-Full` | | ✅︎ |

examples/offline_inference/audio_language.py

@@ -70,6 +70,29 @@ def run_audioflamingo3(question: str, audio_count: int) -> ModelRequestData:
     )
 
 
+# CohereASR
+def run_cohere_asr(question: str, audio_count: int) -> ModelRequestData:
+    assert audio_count == 1, "CohereASR only support single audio input per prompt"
+    # TODO (ekagra): add HF ckpt after asr release
+    model_name = "/host/engines/vllm/audio/2b-release"
+
+    prompt = (
+        "<|startofcontext|><|startoftranscript|>"
+        "<|emo:undefined|><|en|><|en|><|pnc|><|noitn|>"
+        "<|notimestamp|><|nodiarize|>"
+    )
+    engine_args = EngineArgs(
+        model=model_name,
+        limit_mm_per_prompt={"audio": audio_count},
+        trust_remote_code=True,
+    )
+
+    return ModelRequestData(
+        engine_args=engine_args,
+        prompt=prompt,
+    )
+
+
 # MusicFlamingo
 def run_musicflamingo(question: str, audio_count: int) -> ModelRequestData:
     model_name = "nvidia/music-flamingo-2601-hf"
@@ -508,14 +531,15 @@ def run_whisper(question: str, audio_count: int) -> ModelRequestData:
 
 model_example_map = {
     "audioflamingo3": run_audioflamingo3,
-    "musicflamingo": run_musicflamingo,
+    "cohere_asr": run_cohere_asr,
+    "funaudiochat": run_funaudiochat,
     "gemma3n": run_gemma3n,
     "glmasr": run_glmasr,
-    "funaudiochat": run_funaudiochat,
     "granite_speech": run_granite_speech,
     "kimi_audio": run_kimi_audio,
     "midashenglm": run_midashenglm,
     "minicpmo": run_minicpmo,
+    "musicflamingo": run_musicflamingo,
     "phi4_mm": run_phi4mm,
     "qwen2_audio": run_qwen2_audio,
     "qwen2_5_omni": run_qwen2_5_omni,

examples/online_serving/openai_chat_completion_client_for_multimodal.py

@@ -20,9 +20,9 @@ run the script with
 python openai_chat_completion_client_for_multimodal.py --chat-type audio
 """
 
-import base64
 import os
 
+import pybase64 as base64
 import requests
 from openai import OpenAI
 from utils import get_first_model

examples/online_serving/openai_realtime_client.py

@@ -24,11 +24,11 @@ The script:
 
 import argparse
 import asyncio
-import base64
 import json
 
 import librosa
 import numpy as np
+import pybase64 as base64
 import websockets
 
 from vllm.assets.audio import AudioAsset

examples/online_serving/openai_realtime_microphone_client.py

@@ -18,13 +18,13 @@ Requirements: websockets, numpy, gradio
 
 import argparse
 import asyncio
-import base64
 import json
 import queue
 import threading
 
 import gradio as gr
 import numpy as np
+import pybase64 as base64
 import websockets
 
 SAMPLE_RATE = 16_000

examples/pooling/embed/embedding_requests_base64_online.py

@@ -7,8 +7,8 @@ NOTE:
 """
 
 import argparse
-import base64
 
+import pybase64 as base64
 import requests
 import torch
 

examples/pooling/embed/vision_embedding_online.py

@@ -7,10 +7,10 @@ Refer to each `run_*` function for the command to run the server for that model.
 """
 
 import argparse
-import base64
 import io
 from typing import Literal
 
+import pybase64 as base64
 from openai import OpenAI
 from openai._types import NOT_GIVEN, NotGiven
 from openai.types.chat import ChatCompletionMessageParam

examples/pooling/plugin/prithvi_geospatial_mae_io_processor.py

@@ -1,8 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import base64
 import os
 
+import pybase64 as base64
 import torch
 
 from vllm import LLM

examples/pooling/plugin/prithvi_geospatial_mae_online.py

@@ -1,9 +1,9 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
 import os
 
+import pybase64 as base64
 import requests
 
 # This example shows how to perform an online inference that generates

examples/pooling/score/colqwen3_5_rerank_online.py

@@ -0,0 +1,130 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Example of using ColQwen3.5 late interaction model for reranking.
+
+ColQwen3.5 is a multi-modal ColBERT-style model based on Qwen3.5.
+It produces per-token embeddings and uses MaxSim scoring for retrieval
+and reranking. Supports both text and image inputs.
+
+Start the server with:
+    vllm serve athrael-soju/colqwen3.5-4.5B --max-model-len 4096
+
+Then run this script:
+    python colqwen3_5_rerank_online.py
+"""
+
+import requests
+
+MODEL = "athrael-soju/colqwen3.5-4.5B"
+BASE_URL = "http://127.0.0.1:8000"
+
+headers = {"accept": "application/json", "Content-Type": "application/json"}
+
+
+def rerank_text():
+    """Text-only reranking via /rerank endpoint."""
+    print("=" * 60)
+    print("1. Text reranking (/rerank)")
+    print("=" * 60)
+
+    data = {
+        "model": MODEL,
+        "query": "What is machine learning?",
+        "documents": [
+            "Machine learning is a subset of artificial intelligence.",
+            "Python is a programming language.",
+            "Deep learning uses neural networks for complex tasks.",
+            "The weather today is sunny.",
+        ],
+    }
+
+    response = requests.post(f"{BASE_URL}/rerank", headers=headers, json=data)
+
+    if response.status_code == 200:
+        result = response.json()
+        print("\n  Ranked documents (most relevant first):")
+        for item in result["results"]:
+            doc_idx = item["index"]
+            score = item["relevance_score"]
+            print(f"    [{score:.4f}] {data['documents'][doc_idx]}")
+    else:
+        print(f"  Request failed: {response.status_code}")
+        print(f"  {response.text[:300]}")
+
+
+def score_text():
+    """Text-only scoring via /score endpoint."""
+    print()
+    print("=" * 60)
+    print("2. Text scoring (/score)")
+    print("=" * 60)
+
+    query = "What is the capital of France?"
+    documents = [
+        "The capital of France is Paris.",
+        "Berlin is the capital of Germany.",
+        "Python is a programming language.",
+    ]
+
+    data = {
+        "model": MODEL,
+        "text_1": query,
+        "text_2": documents,
+    }
+
+    response = requests.post(f"{BASE_URL}/score", headers=headers, json=data)
+
+    if response.status_code == 200:
+        result = response.json()
+        print(f"\n  Query: {query}\n")
+        for item in result["data"]:
+            idx = item["index"]
+            score = item["score"]
+            print(f"    Doc {idx} (score={score:.4f}): {documents[idx]}")
+    else:
+        print(f"  Request failed: {response.status_code}")
+        print(f"  {response.text[:300]}")
+
+
+def score_text_top_n():
+    """Text reranking with top_n filtering via /rerank endpoint."""
+    print()
+    print("=" * 60)
+    print("3. Text reranking with top_n=2 (/rerank)")
+    print("=" * 60)
+
+    data = {
+        "model": MODEL,
+        "query": "What is the capital of France?",
+        "documents": [
+            "The capital of France is Paris.",
+            "Berlin is the capital of Germany.",
+            "Python is a programming language.",
+            "The Eiffel Tower is in Paris.",
+        ],
+        "top_n": 2,
+    }
+
+    response = requests.post(f"{BASE_URL}/rerank", headers=headers, json=data)
+
+    if response.status_code == 200:
+        result = response.json()
+        print(f"\n  Top {data['top_n']} results:")
+        for item in result["results"]:
+            doc_idx = item["index"]
+            score = item["relevance_score"]
+            print(f"    [{score:.4f}] {data['documents'][doc_idx]}")
+    else:
+        print(f"  Request failed: {response.status_code}")
+        print(f"  {response.text[:300]}")
+
+
+def main():
+    rerank_text()
+    score_text()
+    score_text_top_n()
+
+
+if __name__ == "__main__":
+    main()

examples/pooling/score/colqwen3_rerank_online.py

@@ -15,9 +15,9 @@ Then run this script:
     python colqwen3_rerank_online.py
 """
 
-import base64
 from io import BytesIO
 
+import pybase64 as base64
 import requests
 from PIL import Image
 

examples/pooling/token_embed/colqwen3_token_embed_online.py

@@ -21,10 +21,10 @@ Then run this script:
 """
 
 import argparse
-import base64
 from io import BytesIO
 
 import numpy as np
+import pybase64 as base64
 import requests
 from PIL import Image
 

requirements/common.txt

@@ -37,7 +37,7 @@ pyyaml
 six>=1.16.0; python_version > '3.11' # transitive dependency of pandas that needs to be the latest version for python 3.12
 setuptools>=77.0.3,<81.0.0; python_version > '3.11' # Setuptools is used by triton, we need to ensure a modern version is installed for 3.12+ so that it does not try to import distutils, which was removed in 3.12
 einops # Required for Qwen2-VL.
-compressed-tensors == 0.13.0 # required for compressed-tensors
+compressed-tensors == 0.14.0.1 # required for compressed-tensors
 depyf==0.20.0 # required for profiling and debugging with compilation config
 cloudpickle # allows pickling lambda functions in model_executor/models/registry.py
 watchfiles # required for http server to monitor the updates of TLS files

tests/benchmarks/test_random_multimodal_dataset_video.py

@@ -1,12 +1,12 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
 import os
 from tempfile import NamedTemporaryFile
 from typing import Any, cast
 
 import cv2
+import pybase64 as base64
 import pytest
 from transformers import AutoTokenizer, PreTrainedTokenizerBase
 

tests/distributed/test_weight_transfer.py

@@ -6,10 +6,10 @@ Unit tests for engine classes (parsing, validation, registry).
 Integration tests for NCCL and IPC weight transfer between processes using Ray.
 """
 
-import base64
 import pickle
 from unittest.mock import MagicMock
 
+import pybase64 as base64
 import pytest
 import ray
 import torch

tests/entrypoints/openai/chat_completion/test_audio_in_video.py

@@ -1,15 +1,15 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
 import json
 
 import openai
+import pybase64 as base64
 import pytest
 import pytest_asyncio
 
 from tests.conftest import VideoTestAssets
-from tests.utils import RemoteOpenAIServer
+from tests.utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer
 
 MODEL_NAME = "Qwen/Qwen2.5-Omni-3B"
 
@@ -22,6 +22,7 @@ def server():
         "--enforce-eager",
         "--limit-mm-per-prompt",
         json.dumps({"audio": 3, "video": 3}),
+        *ROCM_EXTRA_ARGS,
     ]
 
     with RemoteOpenAIServer(

tests/entrypoints/openai/chat_completion/test_chat_error.py

@@ -111,7 +111,7 @@ def _build_serving_chat(engine: AsyncLLM) -> OpenAIServingChat:
             [{"prompt_token_ids": [1, 2, 3]}],
         )
 
-    serving_chat.openai_serving_render._preprocess_chat = AsyncMock(
+    serving_chat.openai_serving_render.preprocess_chat = AsyncMock(
         side_effect=_fake_preprocess_chat
     )
     return serving_chat

tests/entrypoints/openai/chat_completion/test_vision_embeds.py

@@ -1,9 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
-
 import numpy as np
+import pybase64 as base64
 import pytest
 import requests
 import torch

tests/entrypoints/openai/completion/test_completion_with_prompt_embeds.py

@@ -1,11 +1,11 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
 import io
 import json
 
 import openai  # use the official client for correctness check
+import pybase64 as base64
 import pytest
 import pytest_asyncio
 import torch

tests/entrypoints/openai/correctness/test_transcription_api_correctness.py

@@ -19,8 +19,10 @@ import soundfile
 import torch
 from datasets import load_dataset
 from evaluate import load
-from transformers import AutoTokenizer
 
+from vllm.tokenizers import get_tokenizer
+
+from ....models.registry import HF_EXAMPLE_MODELS
 from ....utils import RemoteOpenAIServer
 
 
@@ -64,8 +66,12 @@ async def bound_transcribe(sem, client, tokenizer, audio, reference):
 async def process_dataset(model, client, data, concurrent_request):
     sem = asyncio.Semaphore(concurrent_request)
 
-    # Load tokenizer once outside the loop
-    tokenizer = AutoTokenizer.from_pretrained(model)
+    model_info = HF_EXAMPLE_MODELS.find_hf_info(model)
+    tokenizer = get_tokenizer(
+        model,
+        tokenizer_mode=model_info.tokenizer_mode,
+        trust_remote_code=model_info.trust_remote_code,
+    )
 
     # Warmup call as the first `librosa.load` server-side is quite slow.
     audio, sr = data[0]["audio"]["array"], data[0]["audio"]["sampling_rate"]
@@ -144,20 +150,35 @@ def run_evaluation(
 
 
 # alternatives "openai/whisper-large-v2", "openai/whisper-large-v3-turbo"..
-@pytest.mark.parametrize("model_name", ["openai/whisper-large-v3"])
+# NOTE: Expected WER measured with equivalent hf.transformers args:
+# whisper-large-v3 + esb-datasets-earnings22-validation-tiny-filtered.
+@pytest.mark.parametrize(
+    "model_config",
+    [
+        ("openai/whisper-large-v3", 12.744980),
+        # TODO (ekagra): add HF ckpt after asr release
+        # ("/host/engines/vllm/audio/2b-release", 11.73),
+    ],
+)
 # Original dataset is 20GB+ in size, hence we use a pre-filtered slice.
 @pytest.mark.parametrize(
     "dataset_repo", ["D4nt3/esb-datasets-earnings22-validation-tiny-filtered"]
 )
-# NOTE: Expected WER measured with equivalent hf.transformers args:
-# whisper-large-v3 + esb-datasets-earnings22-validation-tiny-filtered.
-@pytest.mark.parametrize("expected_wer", [12.744980])
 def test_wer_correctness(
-    model_name, dataset_repo, expected_wer, n_examples=-1, max_concurrent_request=None
+    model_config, dataset_repo, n_examples=-1, max_concurrent_request=None
 ):
+    model_name, expected_wer = model_config
+    model_info = HF_EXAMPLE_MODELS.find_hf_info(model_name)
     # TODO refactor to use `ASRDataset`
+    server_args = [
+        "--enforce-eager",
+        f"--tokenizer_mode={model_info.tokenizer_mode}",
+    ]
+    if model_info.trust_remote_code:
+        server_args.append("--trust-remote-code")
     with RemoteOpenAIServer(
-        model_name, ["--enforce-eager"], max_wait_seconds=480
+        model_name,
+        server_args,
     ) as remote_server:
         dataset = load_hf_dataset(dataset_repo)
 
@@ -167,7 +188,14 @@ def test_wer_correctness(
 
         client = remote_server.get_async_client()
         wer = run_evaluation(
-            model_name, client, dataset, max_concurrent_request, n_examples
+            model_name,
+            client,
+            dataset,
+            max_concurrent_request,
+            n_examples,
         )
+
+        print(f"Expected WER: {expected_wer}, Actual WER: {wer}")
+
         if expected_wer:
             torch.testing.assert_close(wer, expected_wer, atol=1e-1, rtol=1e-2)

tests/entrypoints/openai/responses/conftest.py

@@ -8,6 +8,9 @@ from collections.abc import Callable
 from typing import Any
 
 import pytest
+import pytest_asyncio
+
+from tests.utils import RemoteOpenAIServer
 
 logger = logging.getLogger(__name__)
 
@@ -361,3 +364,38 @@ def log_response_diagnostics(
     )
 
     return diagnostics
+
+
+@pytest.fixture(scope="module")
+def default_server_args():
+    return [
+        "--max-model-len",
+        "8192",
+        "--enforce-eager",  # For faster startup.
+        "--enable-auto-tool-choice",
+        "--structured-outputs-config.backend",
+        "xgrammar",
+        "--tool-call-parser",
+        "hermes",
+        "--reasoning-parser",
+        "qwen3",
+    ]
+
+
+@pytest.fixture(scope="module")
+def server_with_store(default_server_args):
+    with RemoteOpenAIServer(
+        "Qwen/Qwen3-1.7B",
+        default_server_args,
+        env_dict={
+            "VLLM_ENABLE_RESPONSES_API_STORE": "1",
+            "VLLM_SERVER_DEV_MODE": "1",
+        },
+    ) as remote_server:
+        yield remote_server
+
+
+@pytest_asyncio.fixture
+async def client(server_with_store):
+    async with server_with_store.get_async_client() as async_client:
+        yield async_client

tests/entrypoints/openai/responses/test_function_call.py

@@ -134,6 +134,34 @@ async def test_function_tool_use(
     assert reasoning.type == "reasoning"
 
 
+@pytest.mark.asyncio
+@pytest.mark.parametrize("model_name", [MODEL_NAME])
+async def test_max_tokens_with_tool_choice_required(
+    client: openai.AsyncOpenAI, model_name: str
+):
+    prompt = [
+        {
+            "role": "user",
+            "content": "Can you tell me what the current weather is in Berlin and the "
+            "forecast for the next 5 days, in fahrenheit?",
+        },
+    ]
+    response = await client.responses.create(
+        model=model_name,
+        input=prompt,
+        tools=tools,
+        tool_choice="required",
+        max_output_tokens=10,
+    )
+    assert len(response.output) >= 1
+    for out in response.output:
+        # When `tool_choice="required"` and the tokens of `tools`
+        # exceed `max_output_tokens`,`function_call` should be empty.
+        # This behavior should be consistent with OpenAI
+        assert out.type != "function_call"
+    assert response.incomplete_details.reason == "max_output_tokens"
+
+
 @pytest.mark.asyncio
 async def test_named_tool_use(client: openai.AsyncOpenAI):
     def get_weather(latitude: float, longitude: float) -> str:

tests/entrypoints/openai/responses/test_harmony.py

@@ -16,7 +16,8 @@ import requests
 from openai import InternalServerError, NotFoundError, OpenAI
 from openai_harmony import Message
 
-from ....utils import RemoteOpenAIServer
+from tests.utils import RemoteOpenAIServer
+
 from .conftest import (
     BASE_TEST_ENV,
     events_contain_type,

tests/entrypoints/openai/responses/test_mcp_tools.py

@@ -9,9 +9,9 @@ import pytest_asyncio
 from openai import OpenAI
 from openai_harmony import ToolDescription, ToolNamespaceConfig
 
+from tests.utils import RemoteOpenAIServer
 from vllm.entrypoints.mcp.tool_server import MCPToolServer
 
-from ....utils import RemoteOpenAIServer
 from .conftest import (
     BASE_TEST_ENV,
     events_contain_type,

tests/entrypoints/openai/responses/test_parsable_context.py

@@ -9,7 +9,8 @@ import pytest
 import pytest_asyncio
 from openai import OpenAI
 
-from ....utils import RemoteOpenAIServer
+from tests.utils import RemoteOpenAIServer
+
 from .conftest import (
     BASE_TEST_ENV,
     has_output_type,

tests/entrypoints/openai/responses/test_simple.py

@@ -5,7 +5,8 @@ import pytest
 import pytest_asyncio
 from openai import OpenAI
 
-from ....utils import RemoteOpenAIServer
+from tests.utils import RemoteOpenAIServer
+
 from .conftest import validate_streaming_event_stack
 
 MODEL_NAME = "Qwen/Qwen3-8B"

tests/entrypoints/openai/test_cli_args.py

@@ -291,3 +291,32 @@ def test_served_model_name_parsing(tmp_path, vllm_parser, args, raises):
     else:
         with pytest.raises(raises):
             vllm_parser.parse_args(args=args)
+
+
+### Tests for LoRA target modules parsing
+def test_lora_target_modules_single(serve_parser):
+    """Test parsing single lora-target-modules argument"""
+    args = serve_parser.parse_args(
+        args=["--enable-lora", "--lora-target-modules", "o_proj"]
+    )
+    assert args.lora_target_modules == ["o_proj"]
+
+
+def test_lora_target_modules_multiple(serve_parser):
+    """Test parsing multiple lora-target-modules arguments"""
+    args = serve_parser.parse_args(
+        args=[
+            "--enable-lora",
+            "--lora-target-modules",
+            "o_proj",
+            "qkv_proj",
+            "down_proj",
+        ]
+    )
+    assert args.lora_target_modules == ["o_proj", "qkv_proj", "down_proj"]
+
+
+def test_lora_target_modules_default_none(serve_parser):
+    """Test that lora-target-modules defaults to None"""
+    args = serve_parser.parse_args(args=[])
+    assert args.lora_target_modules is None

tests/entrypoints/openai/test_realtime_validation.py

@@ -2,12 +2,12 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import asyncio
-import base64
 import json
 import warnings
 
 import librosa
 import numpy as np
+import pybase64 as base64
 import pytest
 import websockets
 

tests/entrypoints/openai/test_serving_responses.py

@@ -159,6 +159,7 @@ class TestInitializeToolSessions:
         instance = OpenAIServingResponses(
             engine_client=engine_client,
             models=models,
+            openai_serving_render=MagicMock(),
             request_logger=None,
             chat_template=None,
             chat_template_content_format="auto",
@@ -245,6 +246,7 @@ class TestValidateGeneratorInput:
         instance = OpenAIServingResponses(
             engine_client=engine_client,
             models=models,
+            openai_serving_render=MagicMock(),
             request_logger=None,
             chat_template=None,
             chat_template_content_format="auto",
@@ -308,6 +310,7 @@ async def test_reasoning_tokens_counted_for_text_reasoning_model(monkeypatch):
     serving = OpenAIServingResponses(
         engine_client=engine_client,
         models=models,
+        openai_serving_render=MagicMock(),
         request_logger=None,
         chat_template=None,
         chat_template_content_format="auto",
@@ -607,6 +610,7 @@ def _make_serving_instance_with_reasoning():
     serving = OpenAIServingResponses(
         engine_client=engine_client,
         models=models,
+        openai_serving_render=MagicMock(),
         request_logger=None,
         chat_template=None,
         chat_template_content_format="auto",

tests/entrypoints/openai/tool_parsers/test_gigachat3_tool_parser.py

@@ -13,6 +13,13 @@ from vllm.entrypoints.openai.engine.protocol import FunctionCall
 from vllm.tokenizers import TokenizerLike
 from vllm.tool_parsers import ToolParser, ToolParserManager
 
+MSG_SEP_TOKEN = "<|message_sep|>\n\n"
+ROLE_SEP_TOKEN = "<|role_sep|>\n"
+EOS_TOKEN = "</s>"
+TOOL_HEADER_GIGACHAT3 = f"function call{ROLE_SEP_TOKEN}"
+TOOL_HEADER_GIGACHAT31 = "<|function_call|>"
+
+
 SIMPLE_ARGS_DICT = {
     "action": "create",
     "id": "preferences",
@@ -24,7 +31,10 @@ SIMPLE_FUNCTION_JSON = json.dumps(
     },
     ensure_ascii=False,
 )
-SIMPLE_FUNCTION_OUTPUT = "function call" + SIMPLE_FUNCTION_JSON
+SIMPLE_FUNCTION_OUTPUT_GIGACHAT3 = (
+    f"{MSG_SEP_TOKEN}{TOOL_HEADER_GIGACHAT3}{SIMPLE_FUNCTION_JSON}"
+)
+SIMPLE_FUNCTION_OUTPUT_GIGACHAT31 = f"{TOOL_HEADER_GIGACHAT31}{SIMPLE_FUNCTION_JSON}"
 SIMPLE_FUNCTION_CALL = FunctionCall(
     name="manage_user_memory",
     arguments=json.dumps(SIMPLE_ARGS_DICT, ensure_ascii=False),
@@ -38,7 +48,12 @@ PARAMETERLESS_FUNCTION_JSON = json.dumps(
     },
     ensure_ascii=False,
 )
-PARAMETERLESS_FUNCTION_OUTPUT = "function call" + PARAMETERLESS_FUNCTION_JSON
+PARAMETERLESS_FUNCTION_OUTPUT_GIGACHAT3 = (
+    f"{MSG_SEP_TOKEN}{TOOL_HEADER_GIGACHAT3}{PARAMETERLESS_FUNCTION_JSON}"
+)
+PARAMETERLESS_FUNCTION_OUTPUT_GIGACHAT31 = (
+    f"{TOOL_HEADER_GIGACHAT31}{PARAMETERLESS_FUNCTION_JSON}"
+)
 PARAMETERLESS_FUNCTION_CALL = FunctionCall(
     name="manage_user_memory",
     arguments=json.dumps({}, ensure_ascii=False),
@@ -62,17 +77,38 @@ COMPLEX_FUNCTION_JSON = json.dumps(
     },
     ensure_ascii=False,
 )
-COMPLEX_FUNCTION_OUTPUT = "function call" + COMPLEX_FUNCTION_JSON
+COMPLEX_FUNCTION_OUTPUT_GIGACHAT3 = (
+    f"{MSG_SEP_TOKEN}{TOOL_HEADER_GIGACHAT3}{COMPLEX_FUNCTION_JSON}"
+)
+COMPLEX_FUNCTION_OUTPUT_GIGACHAT31 = f"{TOOL_HEADER_GIGACHAT31}{COMPLEX_FUNCTION_JSON}"
 COMPLEX_FUNCTION_CALL = FunctionCall(
     name="manage_user_memory",
     arguments=json.dumps(COMPLEX_ARGS_DICT, ensure_ascii=False),
 )
 
 
+CONTENT_TEXT = "I'll check that for you."
+MIXED_OUTPUT_GIGACHAT3 = f"{CONTENT_TEXT}{SIMPLE_FUNCTION_OUTPUT_GIGACHAT3}"
+MIXED_OUTPUT_GIGACHAT31 = f"{CONTENT_TEXT}{SIMPLE_FUNCTION_OUTPUT_GIGACHAT31}"
+
+
+@pytest.fixture(name="gigachat_tokenizer")
+def fixture_gigachat_tokenizer(default_tokenizer: TokenizerLike):
+    default_tokenizer.add_tokens(
+        [
+            MSG_SEP_TOKEN,
+            ROLE_SEP_TOKEN,
+            TOOL_HEADER_GIGACHAT31,
+            EOS_TOKEN,
+        ]
+    )
+    return default_tokenizer
+
+
 @pytest.mark.parametrize("streaming", [True, False])
-def test_no_tool_call(streaming: bool, default_tokenizer: TokenizerLike):
+def test_no_tool_call(streaming: bool, gigachat_tokenizer: TokenizerLike):
     tool_parser: ToolParser = ToolParserManager.get_tool_parser("gigachat3")(
-        default_tokenizer
+        gigachat_tokenizer
     )
     model_output = "How can I help you today?"
     content, tool_calls = run_tool_extraction(
@@ -85,45 +121,143 @@ def test_no_tool_call(streaming: bool, default_tokenizer: TokenizerLike):
 TEST_CASES = [
     pytest.param(
         True,
-        SIMPLE_FUNCTION_OUTPUT,
+        SIMPLE_FUNCTION_OUTPUT_GIGACHAT3,
         [SIMPLE_FUNCTION_CALL],
         None,
-        id="simple_streaming",
+        id="simple_streaming_gigachat3",
     ),
     pytest.param(
         False,
-        SIMPLE_FUNCTION_OUTPUT,
+        SIMPLE_FUNCTION_OUTPUT_GIGACHAT3,
         [SIMPLE_FUNCTION_CALL],
         None,
-        id="simple_nonstreaming",
+        id="simple_nonstreaming_gigachat3",
     ),
     pytest.param(
         True,
-        PARAMETERLESS_FUNCTION_OUTPUT,
+        PARAMETERLESS_FUNCTION_OUTPUT_GIGACHAT3,
         [PARAMETERLESS_FUNCTION_CALL],
         None,
-        id="parameterless_streaming",
+        id="parameterless_streaming_gigachat3",
     ),
     pytest.param(
         False,
-        PARAMETERLESS_FUNCTION_OUTPUT,
+        PARAMETERLESS_FUNCTION_OUTPUT_GIGACHAT3,
         [PARAMETERLESS_FUNCTION_CALL],
         None,
-        id="parameterless_nonstreaming",
+        id="parameterless_nonstreaming_gigachat3",
     ),
     pytest.param(
         True,
-        COMPLEX_FUNCTION_OUTPUT,
+        COMPLEX_FUNCTION_OUTPUT_GIGACHAT3,
         [COMPLEX_FUNCTION_CALL],
         None,
-        id="complex_streaming",
+        id="complex_streaming_gigachat3",
     ),
     pytest.param(
         False,
-        COMPLEX_FUNCTION_OUTPUT,
+        COMPLEX_FUNCTION_OUTPUT_GIGACHAT3,
         [COMPLEX_FUNCTION_CALL],
         None,
-        id="complex_nonstreaming",
+        id="complex_nonstreaming_gigachat3",
+    ),
+    pytest.param(
+        True,
+        MIXED_OUTPUT_GIGACHAT3,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_streaming_gigachat3",
+    ),
+    pytest.param(
+        False,
+        MIXED_OUTPUT_GIGACHAT3,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_nonstreaming_gigachat3",
+    ),
+    pytest.param(
+        True,
+        MIXED_OUTPUT_GIGACHAT3 + EOS_TOKEN,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_streaming_with_eos_gigachat3",
+    ),
+    pytest.param(
+        False,
+        MIXED_OUTPUT_GIGACHAT3 + EOS_TOKEN,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_nonstreaming_with_eos_gigachat3",
+    ),
+    pytest.param(
+        True,
+        SIMPLE_FUNCTION_OUTPUT_GIGACHAT31,
+        [SIMPLE_FUNCTION_CALL],
+        None,
+        id="simple_streaming_gigachat31",
+    ),
+    pytest.param(
+        False,
+        SIMPLE_FUNCTION_OUTPUT_GIGACHAT31,
+        [SIMPLE_FUNCTION_CALL],
+        None,
+        id="simple_nonstreaming_gigachat31",
+    ),
+    pytest.param(
+        True,
+        PARAMETERLESS_FUNCTION_OUTPUT_GIGACHAT31,
+        [PARAMETERLESS_FUNCTION_CALL],
+        None,
+        id="parameterless_streaming_gigachat31",
+    ),
+    pytest.param(
+        False,
+        PARAMETERLESS_FUNCTION_OUTPUT_GIGACHAT31,
+        [PARAMETERLESS_FUNCTION_CALL],
+        None,
+        id="parameterless_nonstreaming_gigachat31",
+    ),
+    pytest.param(
+        True,
+        COMPLEX_FUNCTION_OUTPUT_GIGACHAT31,
+        [COMPLEX_FUNCTION_CALL],
+        None,
+        id="complex_streaming_gigachat31",
+    ),
+    pytest.param(
+        False,
+        COMPLEX_FUNCTION_OUTPUT_GIGACHAT31,
+        [COMPLEX_FUNCTION_CALL],
+        None,
+        id="complex_nonstreaming_gigachat31",
+    ),
+    pytest.param(
+        True,
+        MIXED_OUTPUT_GIGACHAT31,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_streaming_gigachat31",
+    ),
+    pytest.param(
+        False,
+        MIXED_OUTPUT_GIGACHAT31,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_nonstreaming_gigachat31",
+    ),
+    pytest.param(
+        True,
+        MIXED_OUTPUT_GIGACHAT31 + EOS_TOKEN,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_streaming_with_eos_gigachat31",
+    ),
+    pytest.param(
+        False,
+        MIXED_OUTPUT_GIGACHAT31 + EOS_TOKEN,
+        [SIMPLE_FUNCTION_CALL],
+        CONTENT_TEXT,
+        id="mixed_content_nonstreaming_with_eos_gigachat31",
     ),
 ]
 
@@ -136,14 +270,16 @@ def test_tool_call(
     model_output: str,
     expected_tool_calls: list[FunctionCall],
     expected_content: str | None,
-    default_tokenizer: TokenizerLike,
+    gigachat_tokenizer: TokenizerLike,
 ):
     tool_parser: ToolParser = ToolParserManager.get_tool_parser("gigachat3")(
-        default_tokenizer
+        gigachat_tokenizer
     )
     content, tool_calls = run_tool_extraction(
         tool_parser, model_output, streaming=streaming
     )
+    if content == "":
+        content = None
     assert content == expected_content
     assert len(tool_calls) == len(expected_tool_calls)
     for actual, expected in zip(tool_calls, expected_tool_calls):
@@ -154,15 +290,46 @@ def test_tool_call(
         assert actual_args == expected_args
 
 
-def test_streaming_tool_call_with_large_steps(default_tokenizer: TokenizerLike):
+@pytest.mark.parametrize(
+    "model_output_deltas",
+    [
+        pytest.param(
+            [
+                CONTENT_TEXT[:3],
+                CONTENT_TEXT[3:5],
+                CONTENT_TEXT[5:],
+                MSG_SEP_TOKEN,
+                TOOL_HEADER_GIGACHAT3,
+                COMPLEX_FUNCTION_JSON[:40],
+                COMPLEX_FUNCTION_JSON[40:-1],
+                COMPLEX_FUNCTION_JSON[-1],
+            ],
+            id="gigachat3",
+        ),
+        pytest.param(
+            [
+                CONTENT_TEXT[:3],
+                CONTENT_TEXT[3:5],
+                CONTENT_TEXT[5:],
+                TOOL_HEADER_GIGACHAT31,
+                COMPLEX_FUNCTION_JSON[:40],
+                COMPLEX_FUNCTION_JSON[40:-1],
+                COMPLEX_FUNCTION_JSON[-1],
+            ],
+            id="gigachat31",
+        ),
+    ],
+)
+def test_streaming_tool_call_with_large_steps(
+    model_output_deltas: list[str],
+    gigachat_tokenizer: TokenizerLike,
+):
+    """
+    Test that the closing braces are streamed correctly.
+    """
     tool_parser: ToolParser = ToolParserManager.get_tool_parser("gigachat3")(
-        default_tokenizer
+        gigachat_tokenizer
     )
-    model_output_deltas = [
-        "function call",
-        COMPLEX_FUNCTION_JSON[:40],
-        COMPLEX_FUNCTION_JSON[40:],
-    ]
     reconstructor = run_tool_extraction_streaming(
         tool_parser,
         model_output_deltas,

tests/entrypoints/pooling/embed/test_cohere_online.py

@@ -7,10 +7,10 @@ embedding models, covering text embedding, embedding type conversions,
 response structure, batching, normalisation, and semantic similarity.
 """
 
-import base64
 import struct
 
 import numpy as np
+import pybase64 as base64
 import pytest
 import requests
 

tests/entrypoints/pooling/embed/test_cohere_online_vision.py

@@ -6,11 +6,11 @@ Validates image embedding, batching, normalisation, and embedding type
 conversions through the /v2/embed endpoint.
 """
 
-import base64
 import struct
 import zlib
 
 import numpy as np
+import pybase64 as base64
 import pytest
 import requests
 

tests/entrypoints/pooling/embed/test_online.py

@@ -1,11 +1,11 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
 import json
 
 import numpy as np
 import openai
+import pybase64 as base64
 import pytest
 import pytest_asyncio
 import requests

tests/entrypoints/pooling/embed/test_protocol.py

@@ -3,10 +3,10 @@
 """Unit tests for Cohere embed protocol: build_typed_embeddings and its
 underlying packing helpers, plus Cohere-specific serving helpers."""
 
-import base64
 import struct
 
 import numpy as np
+import pybase64 as base64
 import pytest
 
 from vllm.entrypoints.pooling.embed.protocol import (

tests/entrypoints/pooling/pooling/test_online.py

@@ -1,10 +1,10 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
 import json
 
 import numpy as np
+import pybase64 as base64
 import pytest
 import requests
 import torch

tests/evals/gsm8k/configs/moe-refactor/Nemotron-Nano-30B-NvFp4-ModelOpt-vllm-cutlass.yaml

@@ -0,0 +1,5 @@
+model_name: "nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-NVFP4"
+accuracy_threshold: 0.29
+num_questions: 1319
+num_fewshot: 5
+server_args: "--enforce-eager --max-model-len 8192 --tensor-parallel-size 2 --moe-backend=cutlass"

tests/evals/gsm8k/configs/moe-refactor/config-b200.txt

@@ -15,3 +15,4 @@ Mixtral-8x7B-BF16-fi-cutlass.yaml
 Mixtral-8x7B-BF16-triton.yaml
 Nemotron-Nano-30B-Fp8-ModelOpt-fi-trtllm.yaml
 Nemotron-Nano-30B-NvFp4-ModelOpt-fi-cutlass.yaml
+Nemotron-Nano-30B-NvFp4-ModelOpt-vllm-cutlass.yaml

tests/kernels/attention/test_cpu_attn.py

@@ -48,7 +48,7 @@ def get_attn_isa(
     else:
         if current_platform.get_cpu_architecture() == CpuArchEnum.ARM:
             return "neon"
-        elif torch._C._cpu._is_amx_tile_supported():
+        elif torch.cpu._is_amx_tile_supported():
             return "amx"
         else:
             return "vec"
@@ -400,9 +400,7 @@ def test_varlen_with_paged_kv_normal_vec(
 @pytest.mark.parametrize("use_alibi", [False])
 @pytest.mark.parametrize("use_sink", [False])
 @pytest.mark.parametrize("isa", ["amx"])
-@pytest.mark.skipif(
-    not torch._C._cpu._is_amx_tile_supported(), reason="no AMX support."
-)
+@pytest.mark.skipif(not torch.cpu._is_amx_tile_supported(), reason="no AMX support.")
 def test_varlen_with_paged_kv_normal_amx(
     seq_lens: list[tuple[int, int]],
     num_heads: tuple[int, int],

tests/kernels/attention/test_trtllm_kvfp8_dequant.py

@@ -12,6 +12,12 @@ import torch
 
 from vllm.platforms import current_platform
 
+if current_platform.is_rocm():
+    pytest.skip(
+        "trtllm kvfp8 dequant is not supported on ROCm.",
+        allow_module_level=True,
+    )
+
 FP8_DTYPE = current_platform.fp8_dtype()
 
 NUM_BLOCKS = 128

tests/kernels/helion/helpers.py

@@ -0,0 +1,67 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import json
+import tempfile
+from collections.abc import Callable
+from contextlib import contextmanager
+from pathlib import Path
+from unittest.mock import patch
+
+import helion
+
+from vllm.kernels.helion.config_manager import ConfigManager
+from vllm.kernels.helion.register import register_kernel
+from vllm.kernels.helion.utils import get_canonical_gpu_name
+
+GPU_PLATFORM = get_canonical_gpu_name()
+
+DEFAULT_CONFIGS: dict[str, helion.Config] = {
+    "default": helion.Config(block_sizes=[32]),
+}
+
+
+@contextmanager
+def dummy_kernel_registry(
+    configs: dict[str, helion.Config] | None = None,
+):
+    """Context manager providing a register function with automatic config setup.
+
+    Yields a ``register`` callable with the same signature as
+    ``register_kernel``.  Before applying the real decorator it writes a
+    config JSON for the kernel name (from ``op_name`` or ``fn.__name__``)
+    into a temporary directory backed by a fresh ``ConfigManager``.
+    """
+    if configs is None:
+        configs = DEFAULT_CONFIGS
+    config_data = {k: v.__dict__["config"] for k, v in configs.items()}
+
+    with tempfile.TemporaryDirectory() as tmpdir:
+        config_dir = Path(tmpdir)
+        ConfigManager.reset_instance()
+        cm = ConfigManager(base_dir=config_dir)
+
+        with patch(
+            "vllm.kernels.helion.config_manager.ConfigManager",
+            return_value=cm,
+        ):
+
+            def register(
+                op_name: str | None = None,
+                **kwargs,
+            ) -> Callable:
+                def decorator(fn: Callable) -> Callable:
+                    name = op_name or fn.__name__
+                    kernel_dir = config_dir / name
+                    kernel_dir.mkdir(parents=True, exist_ok=True)
+                    (kernel_dir / f"{GPU_PLATFORM}.json").write_text(
+                        json.dumps(config_data)
+                    )
+                    return register_kernel(op_name, **kwargs)(fn)
+
+                return decorator
+
+            try:
+                yield register
+            finally:
+                ConfigManager.reset_instance()

tests/kernels/helion/test_autotune.py

@@ -0,0 +1,91 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for autotuning Helion kernels, including disabled kernels with no configs."""
+
+import pytest
+import torch
+
+from vllm.utils.import_utils import has_helion
+
+if not has_helion():
+    pytest.skip(
+        "Helion is not installed. Install with: pip install vllm[helion]",
+        allow_module_level=True,
+    )
+
+import helion
+import helion.language as hl
+from helion.autotuner.base_search import BaseSearch
+
+from tests.kernels.helion.helpers import dummy_kernel_registry
+from vllm.kernels.helion.register import create_helion_decorated_kernel
+
+
+def _add_kernel(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+    out = torch.empty_like(x)
+    for tile in hl.tile(x.size()):
+        out[tile] = x[tile] + y[tile]
+    return out
+
+
+class NoCompileSearch(BaseSearch):
+    """Autotuner that returns the default config without GPU compilation.
+
+    Modeled after helion's test BasicSearch (pytorch/helion#1649).
+    """
+
+    def autotune(self, *, skip_cache: bool = False):
+        return self.config_spec.default_config()
+
+
+def _no_compile_autotuner_fn(bound_kernel, args, **kwargs):
+    return NoCompileSearch(bound_kernel, args, **kwargs)
+
+
+class TestAutotuneDisabledKernel:
+    """Test autotuning flow on disabled kernels (no platform configs)."""
+
+    def setup_method(self):
+        from vllm.kernels.helion.register import _REGISTERED_KERNELS
+
+        self._saved_registry = dict(_REGISTERED_KERNELS)
+        _REGISTERED_KERNELS.clear()
+
+    def teardown_method(self):
+        from vllm.kernels.helion.register import _REGISTERED_KERNELS
+
+        _REGISTERED_KERNELS.clear()
+        _REGISTERED_KERNELS.update(self._saved_registry)
+
+    def test_autotune_disabled_kernel_produces_valid_config(self):
+        """Register a kernel with no configs (disabled), run autotune,
+        verify it produces a valid helion.Config."""
+        with dummy_kernel_registry(configs={}) as register:
+            wrapper = register(
+                "autotune_test_kernel",
+                config_picker=lambda args, keys: "default",
+                fake_impl=lambda *a, **kw: None,
+                input_generator=lambda: {
+                    "small": (
+                        torch.randn(4, 4, device="cuda"),
+                        torch.randn(4, 4, device="cuda"),
+                    ),
+                },
+            )(_add_kernel)
+
+        assert wrapper._disabled is True
+
+        inputs = wrapper.get_inputs()
+        assert "small" in inputs
+
+        settings = helion.Settings()
+        settings.autotuner_fn = _no_compile_autotuner_fn
+        wrapper.helion_settings = settings
+
+        config = wrapper.run_autotune(inputs["small"])
+        expected_default = (
+            create_helion_decorated_kernel(_add_kernel, helion_settings=settings)
+            .bind(inputs["small"])
+            .config_spec.default_config()
+        )
+        assert config == expected_default

tests/kernels/helion/test_pattern_matching.py

@@ -52,7 +52,7 @@ def _helion_mock_context():
 
     with (
         patch(
-            "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+            "vllm.kernels.helion.config_manager.ConfigManager",
             return_value=mock_config_manager,
         ),
         patch(
@@ -87,8 +87,8 @@ class TestMakeFxHop:
                 raw_kernel_func=raw_add_scale,
                 op_name="test_make_fx",
                 fake_impl=lambda *a, **kw: None,
+                config_picker=lambda args, keys: "default",
             )
-            wrapper.register_config_picker(lambda args, keys: "default")
 
             def fn(x, y):
                 return wrapper(x, y, scale)
@@ -143,8 +143,8 @@ class TestMakeFxHop:
                 raw_kernel_func=raw_silu_mul,
                 op_name="test_pm_silu_mul",
                 fake_impl=lambda *a, **kw: None,
+                config_picker=lambda args, keys: "default",
             )
-            wrapper.register_config_picker(lambda args, keys: "default")
 
             def pattern(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
                 return torch.nn.functional.silu(x) * y

tests/kernels/helion/test_register.py

@@ -21,7 +21,9 @@ if not has_helion():
     )
 
 import helion
+import helion.language as hl
 
+from tests.kernels.helion.helpers import dummy_kernel_registry
 from vllm.kernels.helion.config_manager import ConfigManager
 from vllm.kernels.helion.register import (
     _HOP_AVAILABLE,
@@ -34,6 +36,13 @@ from vllm.kernels.helion.register import (
 )
 
 
+def _add_kernel(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+    out = torch.empty_like(x)
+    for tile in hl.tile(x.size()):
+        out[tile] = x[tile] + y[tile]
+    return out
+
+
 @pytest.fixture
 def sample_configs():
     """Create real Helion config objects for testing."""
@@ -90,7 +99,7 @@ def configured_kernel(sample_kernel, sample_configs, config_manager_with_test_co
 
     with (
         patch(
-            "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+            "vllm.kernels.helion.config_manager.ConfigManager",
             return_value=config_manager_with_test_configs,
         ),
         patch(
@@ -158,7 +167,7 @@ def create_configured_kernel_with_configs(
 
     with (
         patch(
-            "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+            "vllm.kernels.helion.config_manager.ConfigManager",
             return_value=mock_config_manager,
         ),
         patch(
@@ -189,7 +198,7 @@ class TestConfiguredHelionKernel:
 
         with (
             patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+                "vllm.kernels.helion.config_manager.ConfigManager",
                 return_value=mock_config_manager,
             ),
             patch(
@@ -266,7 +275,7 @@ class TestConfiguredHelionKernel:
         with (
             patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
             patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+                "vllm.kernels.helion.config_manager.ConfigManager",
                 return_value=mock_config_manager,
             ),
             patch(
@@ -310,7 +319,7 @@ class TestConfiguredHelionKernel:
         with (
             patch("vllm.kernels.helion.register.helion.kernel") as mock_helion_kernel,
             patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+                "vllm.kernels.helion.config_manager.ConfigManager",
                 return_value=mock_config_manager,
             ),
             patch(
@@ -346,23 +355,15 @@ class TestConfiguredHelionKernel:
 class TestHelionKernelWrapper:
     """Test suite for HelionKernelWrapper."""
 
-    def test_get_configured_op_validates_configs_available(self, sample_kernel):
-        """Test get_configured_op validates configs are available."""
+    def test_init_disables_on_missing_configs(self, sample_kernel):
+        """Test __init__ marks wrapper as disabled when configs are missing."""
 
         def fake_impl(*args, **kwargs):
             return torch.zeros_like(args[0])
 
-        wrapper = HelionKernelWrapper(
-            raw_kernel_func=sample_kernel,
-            op_name="test_kernel",
-            fake_impl=fake_impl,
-        )
-
         def default_picker(args, config_keys):
             return "default"
 
-        wrapper._config_picker = default_picker
-
         mock_config_manager = Mock(spec=ConfigManager)
         mock_config_manager.get_platform_configs = Mock(
             return_value={}
@@ -370,72 +371,7 @@ class TestHelionKernelWrapper:
 
         with (
             patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
-                return_value=mock_config_manager,
-            ),
-            patch(
-                "vllm.kernels.helion.utils.get_canonical_gpu_name",
-                return_value="nvidia_h200",
-            ),
-            pytest.raises(ValueError, match="No configs available"),
-        ):
-            wrapper.get_configured_op()
-
-    def test_get_configured_op_validates_config_picker(
-        self, sample_kernel, sample_configs
-    ):
-        """Test get_configured_op validates config picker."""
-
-        def fake_impl(*args, **kwargs):
-            return torch.zeros_like(args[0])
-
-        wrapper = HelionKernelWrapper(
-            raw_kernel_func=sample_kernel,
-            op_name="test_kernel",
-            fake_impl=fake_impl,
-        )
-        # Don't set config picker - should raise assertion error
-
-        mock_config_manager = Mock(spec=ConfigManager)
-        mock_config_manager.get_platform_configs = Mock(return_value=sample_configs)
-
-        with (
-            patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
-                return_value=mock_config_manager,
-            ),
-            patch(
-                "vllm.kernels.helion.utils.get_canonical_gpu_name",
-                return_value="nvidia_h200",
-            ),
-            pytest.raises(AssertionError, match="No config picker registered"),
-        ):
-            wrapper.get_configured_op()
-
-    def test_get_configured_op_returns_cached_kernel(
-        self, sample_kernel, sample_configs
-    ):
-        """Test get_configured_op returns cached ConfiguredHelionKernel."""
-
-        def fake_impl(*args, **kwargs):
-            return torch.zeros_like(args[0])
-
-        def default_picker(args, config_keys):
-            return "default"
-
-        wrapper = HelionKernelWrapper(
-            raw_kernel_func=sample_kernel,
-            op_name="test_kernel",
-            fake_impl=fake_impl,
-        )
-        wrapper._config_picker = default_picker
-
-        mock_config_manager = Mock(spec=ConfigManager)
-        mock_config_manager.get_platform_configs = Mock(return_value=sample_configs)
-
-        with (
-            patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+                "vllm.kernels.helion.config_manager.ConfigManager",
                 return_value=mock_config_manager,
             ),
             patch(
@@ -444,13 +380,269 @@ class TestHelionKernelWrapper:
             ),
             patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
         ):
-            mock_decorated = Mock()
-            mock_kernel.return_value = Mock(return_value=mock_decorated)
+            mock_kernel.return_value = Mock(return_value=sample_kernel)
 
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+            )
+
+            assert wrapper._disabled is True
+            assert "No configs available" in wrapper._disabled_reason
+
+    def test_disabled_wrapper_raises_on_call(self, sample_kernel):
+        """Test __call__ raises RuntimeError on a disabled wrapper."""
+
+        def fake_impl(*args, **kwargs):
+            return torch.zeros_like(args[0])
+
+        def default_picker(args, config_keys):
+            return "default"
+
+        mock_config_manager = Mock(spec=ConfigManager)
+        mock_config_manager.get_platform_configs = Mock(return_value={})
+
+        with (
+            patch(
+                "vllm.kernels.helion.config_manager.ConfigManager",
+                return_value=mock_config_manager,
+            ),
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                return_value="nvidia_h200",
+            ),
+            patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
+        ):
+            mock_kernel.return_value = Mock(return_value=sample_kernel)
+
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+            )
+
+        with pytest.raises(RuntimeError, match="is disabled"):
+            wrapper(torch.randn(4, 4), torch.randn(4, 4))
+
+    def test_disabled_wrapper_get_configured_op_raises(self, sample_kernel):
+        """Test get_configured_op raises RuntimeError on a disabled wrapper."""
+
+        def fake_impl(*args, **kwargs):
+            return torch.zeros_like(args[0])
+
+        def default_picker(args, config_keys):
+            return "default"
+
+        mock_config_manager = Mock(spec=ConfigManager)
+        mock_config_manager.get_platform_configs = Mock(return_value={})
+
+        with (
+            patch(
+                "vllm.kernels.helion.config_manager.ConfigManager",
+                return_value=mock_config_manager,
+            ),
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                return_value="nvidia_h200",
+            ),
+            patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
+        ):
+            mock_kernel.return_value = Mock(return_value=sample_kernel)
+
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+            )
+
+        with pytest.raises(RuntimeError, match="is disabled"):
+            wrapper.get_configured_op()
+
+    def test_disabled_wrapper_supports_get_inputs(self, sample_kernel):
+        """Test get_inputs works on a disabled wrapper."""
+
+        def fake_impl(*args, **kwargs):
+            return torch.zeros_like(args[0])
+
+        def default_picker(args, config_keys):
+            return "default"
+
+        expected_inputs = {"key1": (torch.randn(4),)}
+        input_gen = Mock(return_value=expected_inputs)
+
+        mock_config_manager = Mock(spec=ConfigManager)
+        mock_config_manager.get_platform_configs = Mock(return_value={})
+
+        with (
+            patch(
+                "vllm.kernels.helion.config_manager.ConfigManager",
+                return_value=mock_config_manager,
+            ),
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                return_value="nvidia_h200",
+            ),
+            patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
+        ):
+            mock_kernel.return_value = Mock(return_value=sample_kernel)
+
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+                input_generator=input_gen,
+            )
+
+        assert wrapper._disabled is True
+        result = wrapper.get_inputs()
+        assert result is expected_inputs
+
+    def test_disabled_wrapper_supports_run_autotune(self, sample_kernel):
+        """Test run_autotune works on a disabled wrapper."""
+
+        def fake_impl(*args, **kwargs):
+            return torch.zeros_like(args[0])
+
+        def default_picker(args, config_keys):
+            return "default"
+
+        mock_config_manager = Mock(spec=ConfigManager)
+        mock_config_manager.get_platform_configs = Mock(return_value={})
+
+        mock_config = Mock()
+
+        with (
+            patch(
+                "vllm.kernels.helion.config_manager.ConfigManager",
+                return_value=mock_config_manager,
+            ),
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                return_value="nvidia_h200",
+            ),
+            patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
+        ):
+            mock_kernel.return_value = Mock(return_value=sample_kernel)
+
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+            )
+
+        assert wrapper._disabled is True
+
+        with patch(
+            "vllm.kernels.helion.register.create_helion_decorated_kernel"
+        ) as mock_create:
+            mock_autotune_kernel = Mock()
+            mock_autotune_kernel.autotune.return_value = mock_config
+            mock_create.return_value = mock_autotune_kernel
+
+            inputs = (torch.randn(4, 4),)
+            result = wrapper.run_autotune(inputs)
+            assert result is mock_config
+
+    def test_init_caches_configured_kernel(self, sample_kernel, sample_configs):
+        """Test __init__ eagerly builds and caches ConfiguredHelionKernel."""
+
+        def fake_impl(*args, **kwargs):
+            return torch.zeros_like(args[0])
+
+        def default_picker(args, config_keys):
+            return "default"
+
+        mock_config_manager = Mock(spec=ConfigManager)
+        mock_config_manager.get_platform_configs = Mock(return_value=sample_configs)
+
+        with (
+            patch(
+                "vllm.kernels.helion.config_manager.ConfigManager",
+                return_value=mock_config_manager,
+            ),
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                return_value="nvidia_h200",
+            ),
+            patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
+        ):
+            mock_kernel.return_value = Mock(return_value=sample_kernel)
+
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+            )
+
+            assert wrapper._configured_kernel is not None
             result1 = wrapper.get_configured_op()
             result2 = wrapper.get_configured_op()
             assert result1 is result2
 
+    @pytest.mark.skipif(
+        not _HOP_AVAILABLE, reason="HOP path only used when HOP available"
+    )
+    def test_init_eagerly_initializes_hop_path(self):
+        """Test that register_kernel eagerly builds the configured kernel
+        on the HOP path (no custom op registration needed)."""
+        from vllm.kernels.helion.utils import get_canonical_gpu_name
+
+        configs = {"default": helion.Config(block_sizes=[4, 4])}
+        with (
+            dummy_kernel_registry(configs=configs) as register,
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                wraps=get_canonical_gpu_name,
+            ) as mock_gpu,
+        ):
+            wrapper = register(
+                config_picker=lambda args, keys: "default",
+            )(_add_kernel)
+
+            mock_gpu.assert_called_once()
+            assert wrapper._configured_kernel is not None
+
+        with patch(
+            "vllm.kernels.helion.utils.get_canonical_gpu_name",
+            side_effect=AssertionError("get_canonical_gpu_name called during __call__"),
+        ):
+            x = torch.randn(4, 4, device="cuda")
+            y = torch.randn(4, 4, device="cuda")
+            result = wrapper(x, y)
+            expected = x + y
+            assert torch.allclose(result, expected)
+
+    @pytest.mark.skipif(
+        _HOP_AVAILABLE, reason="CustomOp path not used when HOP available"
+    )
+    def test_init_eagerly_initializes(self):
+        """Test that register_kernel eagerly loads configs and detects GPU
+        during construction so __call__ needs no further initialization."""
+        from vllm.kernels.helion.utils import get_canonical_gpu_name
+
+        with (
+            dummy_kernel_registry() as register,
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                wraps=get_canonical_gpu_name,
+            ) as mock_gpu,
+        ):
+            wrapper = register(
+                config_picker=lambda args, keys: "default",
+            )(_add_kernel)
+
+            # Init must have detected GPU and built the kernel
+            mock_gpu.assert_called_once()
+            assert wrapper._configured_kernel is not None
+            assert hasattr(torch.ops.vllm_helion, wrapper.op_name)
+
     @pytest.mark.skipif(
         _HOP_AVAILABLE, reason="CustomOp path not used when HOP available"
     )
@@ -463,13 +655,6 @@ class TestHelionKernelWrapper:
         def default_picker(args, config_keys):
             return "default"
 
-        wrapper = HelionKernelWrapper(
-            raw_kernel_func=sample_kernel,
-            op_name="test_kernel",
-            fake_impl=fake_impl,
-        )
-        wrapper._config_picker = default_picker
-
         mock_config_manager = Mock(spec=ConfigManager)
         mock_config_manager.get_platform_configs = Mock(return_value=sample_configs)
 
@@ -479,7 +664,7 @@ class TestHelionKernelWrapper:
 
         with (
             patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+                "vllm.kernels.helion.config_manager.ConfigManager",
                 return_value=mock_config_manager,
             ),
             patch(
@@ -491,6 +676,13 @@ class TestHelionKernelWrapper:
         ):
             mock_decorated = Mock()
             mock_kernel.return_value = Mock(return_value=mock_decorated)
+
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+            )
             result = wrapper._get_or_register_custom_op()
             assert result is existing_op
 
@@ -506,13 +698,6 @@ class TestHelionKernelWrapper:
         def default_picker(args, config_keys):
             return "default"
 
-        wrapper = HelionKernelWrapper(
-            raw_kernel_func=sample_kernel,
-            op_name="test_kernel",
-            fake_impl=fake_impl,
-        )
-        wrapper._config_picker = default_picker
-
         mock_config_manager = Mock(spec=ConfigManager)
         mock_config_manager.get_platform_configs = Mock(return_value=sample_configs)
 
@@ -532,7 +717,7 @@ class TestHelionKernelWrapper:
 
         with (
             patch(
-                "vllm.kernels.helion.config_manager.ConfigManager.get_instance",
+                "vllm.kernels.helion.config_manager.ConfigManager",
                 return_value=mock_config_manager,
             ),
             patch(
@@ -548,6 +733,13 @@ class TestHelionKernelWrapper:
         ):
             mock_decorated = Mock()
             mock_kernel.return_value = Mock(return_value=mock_decorated)
+
+            wrapper = HelionKernelWrapper(
+                raw_kernel_func=sample_kernel,
+                op_name="test_kernel",
+                fake_impl=fake_impl,
+                config_picker=default_picker,
+            )
             result = wrapper._get_or_register_custom_op()
 
             mock_register.assert_called_once()
@@ -584,11 +776,10 @@ class TestKernelRegistry:
 
     def test_get_kernel_by_name_returns_kernel(self):
         """Test get_kernel_by_name returns registered kernel."""
-        wrapper = HelionKernelWrapper(
-            raw_kernel_func=Mock(),
-            op_name="test_kernel",
-            fake_impl=Mock(),
-        )
+        with dummy_kernel_registry() as register:
+            wrapper = register(
+                "test_kernel", config_picker=lambda args, keys: "default"
+            )(_add_kernel)
 
         from vllm.kernels.helion.register import _REGISTERED_KERNELS
 
@@ -604,112 +795,87 @@ class TestKernelRegistry:
 
     def test_register_kernel_auto_generates_fake_impl(self):
         """Test register_kernel auto-generates fake_impl when not provided."""
-        with patch("vllm.kernels.helion.register.infer_fake_impl") as mock_infer:
+        with (
+            dummy_kernel_registry() as register,
+            patch("vllm.kernels.helion.register.infer_fake_impl") as mock_infer,
+        ):
             mock_fake = Mock()
             mock_infer.return_value = mock_fake
+            wrapper = register(
+                config_picker=lambda args, keys: "default",
+            )(_add_kernel)
 
-            def original_kernel(x):
-                return x
-
-            wrapper = register_kernel(original_kernel)
-
-            mock_infer.assert_called_once_with(original_kernel, None)
-            assert wrapper._fake_impl is mock_fake
+        mock_infer.assert_called_once_with(_add_kernel, None)
+        assert wrapper._fake_impl is mock_fake
 
     def test_register_kernel_creates_wrapper(self):
         """Test register_kernel creates HelionKernelWrapper."""
-
-        def test_kernel(x):
-            return x
-
-        result = register_kernel("test_name")(test_kernel)
+        with dummy_kernel_registry() as register:
+            result = register("test_name", config_picker=lambda args, keys: "default")(
+                _add_kernel
+            )
 
         assert isinstance(result, HelionKernelWrapper)
         assert result.op_name == "test_name"
-        assert result.raw_kernel_func is test_kernel
+        assert result.raw_kernel_func is _add_kernel
 
     def test_register_kernel_auto_detects_name(self):
         """Test register_kernel uses function name when no name provided."""
+        with dummy_kernel_registry() as register:
+            wrapper = register(config_picker=lambda args, keys: "default")(_add_kernel)
 
-        @register_kernel
-        def my_test_kernel(x):
-            return x
-
-        assert my_test_kernel.op_name == "my_test_kernel"
+        assert wrapper.op_name == "_add_kernel"
 
     def test_register_kernel_registers_in_global_registry(self):
         """Test register_kernel adds wrapper to global registry."""
-
-        @register_kernel
-        def test_kernel(x):
-            return x
+        with dummy_kernel_registry() as register:
+            wrapper = register(
+                "test_kernel", config_picker=lambda args, keys: "default"
+            )(_add_kernel)
 
         registered_kernels = get_registered_kernels()
         assert "test_kernel" in registered_kernels
-        assert registered_kernels["test_kernel"] is test_kernel
+        assert registered_kernels["test_kernel"] is wrapper
 
     def test_register_kernel_passes_helion_settings(self):
         """Test register_kernel passes helion_settings to wrapper."""
-        mock_settings = Mock()
-        mock_settings.to_dict.return_value = {"debug": True}
+        settings = helion.Settings()
+        settings.print_output_code = True
 
-        @register_kernel("test_name", helion_settings=mock_settings)
-        def test_kernel(x):
-            return x
+        with dummy_kernel_registry() as register:
+            result = register(
+                "test_name",
+                config_picker=lambda args, keys: "default",
+                helion_settings=settings,
+            )(_add_kernel)
 
-        assert test_kernel.helion_settings is mock_settings
+        assert result.helion_settings is settings
 
     def test_register_kernel_supports_decorator_syntax(self):
         """Test register_kernel works with decorator arguments."""
         mock_fake = Mock()
 
-        wrapper = register_kernel("custom_name", fake_impl=mock_fake)
-
-        def test_kernel(x):
-            return x
-
-        result = wrapper(test_kernel)
+        with dummy_kernel_registry() as register:
+            result = register(
+                "custom_name",
+                config_picker=lambda args, keys: "default",
+                fake_impl=mock_fake,
+            )(_add_kernel)
 
         assert result.op_name == "custom_name"
         assert result._fake_impl is mock_fake
 
-    def test_register_kernel_bare_decorator(self):
-        """Test register_kernel works as bare decorator."""
-
-        @register_kernel
-        def test_kernel(x):
-            return x
-
-        assert isinstance(test_kernel, HelionKernelWrapper)
-        assert test_kernel.op_name == "test_kernel"
-
-    def test_registered_wrapper_can_register_config_picker(self):
-        """Test that registered wrapper can register config picker."""
-
-        @register_kernel
-        def test_kernel(x):
-            return x
-
-        def my_picker(args, config_keys):
-            return "default"
-
-        result = test_kernel.register_config_picker(my_picker)
-
-        assert result is my_picker
-        assert test_kernel._config_picker is my_picker
-
     def test_register_kernel_raises_on_duplicate_registration(self):
         """Test register_kernel raises error on duplicate names."""
+        with dummy_kernel_registry() as register:
+            register("duplicate_name", config_picker=lambda args, keys: "default")(
+                _add_kernel
+            )
 
-        @register_kernel("duplicate_name")
-        def kernel1(x):
-            return x
-
-        with pytest.raises(ValueError, match="already registered"):
-
-            @register_kernel("duplicate_name")
-            def kernel2(x):
-                return x
+            with pytest.raises(ValueError, match="already registered"):
+                register("duplicate_name", config_picker=lambda args, keys: "default")(
+                    _add_kernel
+                )
 
     def test_register_kernel_rejects_autotuner_fn_in_settings(self):
         """Test register_kernel rejects conflicting autotuner_fn."""
@@ -718,7 +884,11 @@ class TestKernelRegistry:
 
         with pytest.raises(ValueError, match="uses a custom autotuner"):
 
-            @register_kernel("test", helion_settings=mock_settings)
+            @register_kernel(
+                "test",
+                config_picker=lambda args, keys: "default",
+                helion_settings=mock_settings,
+            )
             def test_kernel(x):
                 return x
 
@@ -727,11 +897,47 @@ class TestKernelRegistry:
         mock_settings = Mock()
         mock_settings.to_dict.return_value = {"static_shapes": False}
 
-        with patch("vllm.kernels.helion.register.logger") as mock_logger:
+        with (
+            dummy_kernel_registry() as register,
+            patch("vllm.kernels.helion.register.logger") as mock_logger,
+        ):
+            register(
+                "test",
+                config_picker=lambda args, keys: "default",
+                helion_settings=mock_settings,
+            )(_add_kernel)
 
-            @register_kernel("test", helion_settings=mock_settings)
-            def test_kernel(x):
-                return x
+        mock_logger.warning.assert_not_called()
 
-            # Should not call warning
-            mock_logger.warning.assert_not_called()
+    def test_disabled_kernel_appears_in_registry(self):
+        """Test that a disabled wrapper is still in the global registry."""
+
+        def fake_impl(*args, **kwargs):
+            return torch.zeros_like(args[0])
+
+        mock_config_manager = Mock(spec=ConfigManager)
+        mock_config_manager.get_platform_configs = Mock(return_value={})
+
+        with (
+            patch(
+                "vllm.kernels.helion.config_manager.ConfigManager",
+                return_value=mock_config_manager,
+            ),
+            patch(
+                "vllm.kernels.helion.utils.get_canonical_gpu_name",
+                return_value="nvidia_h200",
+            ),
+            patch("vllm.kernels.helion.register.helion.kernel") as mock_kernel,
+        ):
+            mock_kernel.return_value = Mock(return_value=_add_kernel)
+
+            wrapper = register_kernel(
+                "disabled_kernel",
+                config_picker=lambda args, keys: "default",
+                fake_impl=fake_impl,
+            )(_add_kernel)
+
+        assert wrapper._disabled is True
+        registered = get_registered_kernels()
+        assert "disabled_kernel" in registered
+        assert registered["disabled_kernel"] is wrapper

tests/kernels/moe/test_cpu_fused_moe.py

@@ -22,7 +22,7 @@ INTERMEDIATE_DIM = [128, 2880]
 BATCH_SIZE = [1, 64, 256]
 ACT = [MoEActivation.SILU, MoEActivation.SWIGLUOAI]
 USE_BIAS = [True, False]
-ISA = ["amx", "vec"] if torch._C._cpu._is_amx_tile_supported() else ["vec"]
+ISA = ["amx", "vec"] if torch.cpu._is_amx_tile_supported() else ["vec"]
 DTYPE = [torch.bfloat16]
 
 

tests/lora/test_lora_manager.py

@@ -711,3 +711,192 @@ def test_packed_loras(default_vllm_config, dist_init, dummy_model_gate_up, devic
     torch.testing.assert_close(
         packed_lora1.lora_b[1], model_lora_clone1.get_lora("up_proj").lora_b
     )
+
+
+def _test_target_modules(
+    model,
+    target_modules: list[str] | None,
+    device: str,
+    expected_lora: list[tuple[str, type]],
+    expected_no_lora: list[tuple[str, type]],
+):
+    """Create a LoRAModelManager and assert which modules have LoRA applied."""
+    LoRAModelManager(
+        model,
+        2,
+        2,
+        2,
+        LoRAConfig(
+            max_lora_rank=8,
+            max_cpu_loras=2,
+            max_loras=2,
+            lora_dtype=DEFAULT_DTYPE,
+            target_modules=target_modules,
+        ),
+        device=device,
+    )
+    for module_path, lora_cls in expected_lora:
+        assert isinstance(model.get_submodule(module_path), lora_cls)
+    for module_path, lora_cls in expected_no_lora:
+        assert not isinstance(model.get_submodule(module_path), lora_cls)
+
+
+@pytest.mark.parametrize("device", DEVICES)
+def test_target_modules_config(default_vllm_config, dist_init, dummy_model, device):
+    """Test that target_modules config restricts which modules get LoRA applied."""
+    _test_target_modules(
+        dummy_model,
+        ["dense1"],
+        device,
+        expected_lora=[
+            ("dense1", ColumnParallelLinearWithLoRA),
+            ("layer1.dense1", ColumnParallelLinearWithLoRA),
+        ],
+        expected_no_lora=[
+            ("dense2", RowParallelLinearWithLoRA),
+            ("layer1.dense2", RowParallelLinearWithLoRA),
+        ],
+    )
+
+
+@pytest.mark.parametrize("device", DEVICES)
+def test_target_modules_multiple(default_vllm_config, dist_init, dummy_model, device):
+    """Test that multiple target_modules work correctly."""
+    _test_target_modules(
+        dummy_model,
+        ["dense1", "dense2"],
+        device,
+        expected_lora=[
+            ("dense1", ColumnParallelLinearWithLoRA),
+            ("layer1.dense1", ColumnParallelLinearWithLoRA),
+            ("dense2", RowParallelLinearWithLoRA),
+            ("layer1.dense2", RowParallelLinearWithLoRA),
+        ],
+        expected_no_lora=[],
+    )
+
+
+@pytest.mark.parametrize("device", DEVICES)
+def test_target_modules_none_uses_all(
+    default_vllm_config, dist_init, dummy_model, device
+):
+    """Test that target_modules=None uses all supported modules."""
+    _test_target_modules(
+        dummy_model,
+        None,
+        device,
+        expected_lora=[
+            ("dense1", ColumnParallelLinearWithLoRA),
+            ("layer1.dense1", ColumnParallelLinearWithLoRA),
+            ("dense2", RowParallelLinearWithLoRA),
+            ("layer1.dense2", RowParallelLinearWithLoRA),
+        ],
+        expected_no_lora=[],
+    )
+
+
+@pytest.mark.parametrize("device", DEVICES)
+def test_load_adapter_warns_on_unsupported_modules(
+    default_vllm_config, dist_init, dummy_model_gate_up, device, tmp_path
+):
+    """Test that _load_adapter warns when a LoRA adapter contains modules
+    not in the model's supported LoRA target modules."""
+    from unittest.mock import patch
+
+    import vllm.lora.worker_manager as wm_module
+
+    lora_config = LoRAConfig(
+        max_lora_rank=8, max_cpu_loras=4, max_loras=4, lora_dtype=DEFAULT_DTYPE
+    )
+
+    dummy_lora_files = f"{tmp_path}/lora_adapter"
+    os.makedirs(dummy_lora_files, exist_ok=True)
+    create_peft_lora(
+        dummy_model_gate_up,
+        save_dir=dummy_lora_files,
+        target_modules=["layer1.dense1", "dense2"],
+        lora_dtype=DEFAULT_DTYPE,
+    )
+
+    model_config = ModelConfig(max_model_len=16)
+    vllm_config = VllmConfig(model_config=model_config, lora_config=lora_config)
+    vllm_config.scheduler_config.max_num_seqs = 4
+    vllm_config.scheduler_config.max_num_batched_tokens = 2
+
+    worker_manager = WorkerLoRAManager(vllm_config, device, EMBEDDING_MODULES)
+    worker_manager.vocab_size = dummy_model_gate_up.unpadded_vocab_size
+    worker_manager.create_lora_manager(dummy_model_gate_up)
+
+    # Patch from_local_checkpoint to inject an unsupported module
+    original_from_checkpoint = LoRAModel.from_local_checkpoint
+
+    def patched_from_checkpoint(*args, **kwargs):
+        lora = original_from_checkpoint(*args, **kwargs)
+        lora.loras["unsupported_module"] = LoRALayerWeights(
+            module_name="unsupported_module",
+            rank=8,
+            lora_alpha=16,
+            lora_a=torch.randn(8, 10),
+            lora_b=torch.randn(10, 8),
+        )
+        return lora
+
+    lora_request = LoRARequest("test", 1, dummy_lora_files)
+    with (
+        patch.object(LoRAModel, "from_local_checkpoint", patched_from_checkpoint),
+        patch.object(wm_module.logger, "warning_once") as mock_warning,
+    ):
+        worker_manager._load_adapter(lora_request)
+        warning_args = mock_warning.call_args_list
+        found = any("unsupported_module" in str(call) for call in warning_args)
+        assert found, (
+            f"Expected warning about 'unsupported_module', got: {warning_args}"
+        )
+
+
+@pytest.mark.parametrize("device", DEVICES)
+def test_load_adapter_warns_on_target_modules_restriction(
+    default_vllm_config, dist_init, dummy_model_gate_up, device, tmp_path
+):
+    """Test that _load_adapter warns when a LoRA adapter contains modules
+    excluded by the deployment-time target_modules restriction."""
+    from unittest.mock import patch
+
+    import vllm.lora.worker_manager as wm_module
+
+    # Restrict to only dense2 — adapter has dense1 which will be excluded
+    lora_config = LoRAConfig(
+        max_lora_rank=8,
+        max_cpu_loras=4,
+        max_loras=4,
+        lora_dtype=DEFAULT_DTYPE,
+        target_modules=["dense2"],
+    )
+
+    dummy_lora_files = f"{tmp_path}/lora_adapter"
+    os.makedirs(dummy_lora_files, exist_ok=True)
+    create_peft_lora(
+        dummy_model_gate_up,
+        save_dir=dummy_lora_files,
+        target_modules=["layer1.dense1", "dense2"],
+        lora_dtype=DEFAULT_DTYPE,
+    )
+
+    model_config = ModelConfig(max_model_len=16)
+    vllm_config = VllmConfig(model_config=model_config, lora_config=lora_config)
+    vllm_config.scheduler_config.max_num_seqs = 4
+    vllm_config.scheduler_config.max_num_batched_tokens = 2
+
+    worker_manager = WorkerLoRAManager(vllm_config, device, EMBEDDING_MODULES)
+    worker_manager.vocab_size = dummy_model_gate_up.unpadded_vocab_size
+    worker_manager.create_lora_manager(dummy_model_gate_up)
+
+    lora_request = LoRARequest("test", 1, dummy_lora_files)
+    with patch.object(wm_module.logger, "warning_once") as mock_warning:
+        worker_manager._load_adapter(lora_request)
+        warning_args = mock_warning.call_args_list
+        # dense1 is supported by the model but excluded by target_modules
+        found = any("target_modules" in str(call) for call in warning_args)
+        assert found, (
+            f"Expected warning about target_modules restriction, got: {warning_args}"
+        )

tests/lora/test_lora_utils.py

@@ -0,0 +1,60 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+
+from vllm.lora.utils import is_in_target_modules, is_supported_lora_module
+
+
+class TestIsSupportedLoraModule:
+    """Tests for is_supported_lora_module (model-definition check)."""
+
+    def test_suffix_match(self):
+        assert is_supported_lora_module(
+            "model.layers.0.self_attn.o_proj", ["o_proj", "q_proj"]
+        )
+
+    def test_no_match(self):
+        assert not is_supported_lora_module(
+            "model.layers.0.self_attn.o_proj", ["q_proj", "k_proj"]
+        )
+
+    def test_exact_match(self):
+        assert is_supported_lora_module("o_proj", ["o_proj"])
+
+    def test_regex_suffix_matching(self):
+        """Regex anchors to end — partial suffix should not match."""
+        assert not is_supported_lora_module("model.layers.0.self_attn.o_proj", ["proj"])
+
+    def test_empty_supported_modules(self):
+        assert not is_supported_lora_module("model.layers.0.self_attn.o_proj", [])
+
+    def test_multiple_supported_modules(self):
+        supported = ["q_proj", "k_proj", "v_proj", "o_proj"]
+        assert is_supported_lora_module("model.layers.0.self_attn.v_proj", supported)
+        assert not is_supported_lora_module("model.layers.0.mlp.gate_proj", supported)
+
+
+class TestIsInTargetModules:
+    """Tests for is_in_target_modules (deployment-time filter)."""
+
+    def test_none_allows_all(self):
+        assert is_in_target_modules("model.layers.0.self_attn.o_proj", None)
+
+    def test_suffix_in_target(self):
+        assert is_in_target_modules(
+            "model.layers.0.self_attn.o_proj", ["o_proj", "q_proj"]
+        )
+
+    def test_suffix_not_in_target(self):
+        assert not is_in_target_modules(
+            "model.layers.0.self_attn.o_proj", ["q_proj", "k_proj"]
+        )
+
+    def test_empty_target_modules(self):
+        assert not is_in_target_modules("model.layers.0.self_attn.o_proj", [])
+
+    def test_exact_name_match(self):
+        assert is_in_target_modules("dense1", ["dense1", "dense2"])
+
+    def test_exact_name_no_match(self):
+        assert not is_in_target_modules("dense3", ["dense1", "dense2"])

tests/models/multimodal/generation/vlm_utils/model_utils.py

@@ -1235,9 +1235,9 @@ def voxtral_patch_hf_runner(hf_model: "HfRunner") -> "HfRunner":
     generated).
     """
 
-    import base64
     import io
 
+    import pybase64 as base64
     import soundfile as sf
 
     processor = hf_model.processor

tests/models/multimodal/pooling/test_colpali.py

@@ -7,9 +7,9 @@ ColPali is a multi-vector retrieval model based on PaliGemma backbone
 It produces per-token embeddings for both text and image inputs.
 """
 
-import base64
 from io import BytesIO
 
+import pybase64 as base64
 import pytest
 import torch
 from PIL import Image

tests/models/multimodal/pooling/test_colqwen3.py

@@ -7,9 +7,9 @@ ColBERT-style late interaction scoring (MaxSim). It produces per-token
 embeddings for both text and image inputs.
 """
 
-import base64
 from io import BytesIO
 
+import pybase64 as base64
 import pytest
 import torch
 from PIL import Image

tests/models/multimodal/pooling/test_colqwen3_5.py

@@ -0,0 +1,154 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for ColQwen3.5 late interaction model for multi-modal retrieval.
+
+ColQwen3.5 is a multi-vector retrieval model based on Qwen3.5 backbone with
+ColBERT-style late interaction scoring (MaxSim). It produces per-token
+embeddings for both text and image inputs.
+"""
+
+import pytest
+import torch
+
+from ....conftest import VllmRunner
+
+MODELS = [
+    "athrael-soju/colqwen3.5-4.5B-v3",
+]
+
+EMBED_DIMS = {
+    "athrael-soju/colqwen3.5-4.5B-v3": 320,
+}
+
+TEXT_QUERIES = [
+    "What is the capital of France?",
+    "Describe the contents of the document.",
+]
+
+TEXT_DOCUMENTS = [
+    "The capital of France is Paris.",
+    "This document contains important financial data.",
+]
+
+DTYPE = "half"
+
+
+def _run_token_embed_test(
+    vllm_runner: type[VllmRunner],
+    model: str,
+    *,
+    dtype: str,
+) -> None:
+    """Verify per-token embedding shape and L2 normalization."""
+    with vllm_runner(
+        model,
+        runner="pooling",
+        dtype=dtype,
+        max_model_len=4096,
+        enforce_eager=True,
+    ) as vllm_model:
+        outputs = vllm_model.token_embed([TEXT_QUERIES[0]])
+
+        assert len(outputs) == 1
+        emb = torch.tensor(outputs[0])
+        # Token embeddings should be 2D: [num_tokens, embed_dim]
+        assert emb.dim() == 2
+        assert emb.shape[1] == EMBED_DIMS[model]
+        assert emb.shape[0] > 1
+
+        # Verify L2 normalization
+        norms = torch.norm(emb, p=2, dim=-1)
+        torch.testing.assert_close(
+            norms,
+            torch.ones_like(norms),
+            rtol=1e-2,
+            atol=1e-2,
+        )
+
+
+def _run_late_interaction_test(
+    vllm_runner: type[VllmRunner],
+    model: str,
+    *,
+    dtype: str,
+) -> None:
+    """Verify MaxSim scoring matches manual computation."""
+    from vllm.entrypoints.pooling.score.utils import compute_maxsim_score
+
+    with vllm_runner(
+        model,
+        runner="pooling",
+        dtype=dtype,
+        max_model_len=4096,
+        enforce_eager=True,
+    ) as vllm_model:
+        q_outputs = vllm_model.token_embed([TEXT_QUERIES[0]])
+        d_outputs = vllm_model.token_embed([TEXT_DOCUMENTS[0]])
+
+        q_emb = torch.tensor(q_outputs[0])
+        d_emb = torch.tensor(d_outputs[0])
+
+        manual_score = compute_maxsim_score(q_emb, d_emb).item()
+
+        vllm_scores = vllm_model.score(TEXT_QUERIES[0], TEXT_DOCUMENTS[0])
+
+        assert len(vllm_scores) == 1
+        assert vllm_scores[0] == pytest.approx(manual_score, rel=0.01)
+
+
+def _run_relevance_test(
+    vllm_runner: type[VllmRunner],
+    model: str,
+    *,
+    dtype: str,
+) -> None:
+    """Verify that relevant documents score higher than irrelevant ones."""
+    query = "What is machine learning?"
+    documents = [
+        "Machine learning is a subset of artificial intelligence.",
+        "The weather forecast shows rain tomorrow.",
+        "Deep learning uses neural networks for complex tasks.",
+    ]
+
+    with vllm_runner(
+        model,
+        runner="pooling",
+        dtype=dtype,
+        max_model_len=4096,
+        enforce_eager=True,
+    ) as vllm_model:
+        scores = vllm_model.score(query, documents)
+
+        assert len(scores) == 3
+        assert scores[0] > scores[1], "ML doc should score higher than weather doc"
+        assert scores[2] > scores[1], "DL doc should score higher than weather doc"
+
+
+@pytest.mark.parametrize("model", MODELS)
+@pytest.mark.parametrize("dtype", [DTYPE])
+def test_colqwen3_5_token_embed(
+    vllm_runner,
+    model: str,
+    dtype: str,
+) -> None:
+    _run_token_embed_test(vllm_runner, model, dtype=dtype)
+
+
+@pytest.mark.parametrize("model", MODELS)
+@pytest.mark.parametrize("dtype", [DTYPE])
+def test_colqwen3_5_late_interaction_scoring(
+    vllm_runner,
+    model: str,
+    dtype: str,
+) -> None:
+    _run_late_interaction_test(vllm_runner, model, dtype=dtype)
+
+
+@pytest.mark.parametrize("model", MODELS)
+@pytest.mark.parametrize("dtype", [DTYPE])
+def test_colqwen3_5_relevance_ordering(
+    vllm_runner,
+    model: str,
+    dtype: str,
+) -> None:
+    _run_relevance_test(vllm_runner, model, dtype=dtype)

tests/models/multimodal/pooling/test_llama_nemotron_vl.py

@@ -9,10 +9,10 @@ Tests for the LlamaNemotronVL model family:
 Both variants share a SigLIP vision encoder with a bidirectional LLaMA backbone.
 """
 
-import base64
 from io import BytesIO
 from pathlib import Path
 
+import pybase64 as base64
 import pytest
 import torch
 from transformers import AutoModel, AutoModelForSequenceClassification, AutoProcessor

tests/models/multimodal/pooling/test_phi3v.py

@@ -3,6 +3,7 @@
 
 import pytest
 import torch.nn.functional as F
+import transformers.utils
 from PIL import Image
 
 from vllm.assets.base import get_vllm_public_assets
@@ -12,6 +13,12 @@ from ....conftest import IMAGE_ASSETS, HfRunner, PromptImageInput, VllmRunner
 from ....utils import large_gpu_test
 from ...utils import check_embeddings_close
 
+# BC for method that was deleted in Transformers v5.
+# Only needed for generating the HF reference.
+transformers.utils.is_flash_attn_greater_or_equal_2_10 = (
+    lambda: transformers.utils.is_flash_attn_greater_or_equal("2.1.0")
+)
+
 HF_TEXT_PROMPTS = [
     # T -> X
     "Find me an everyday image that matches the given caption: The label of the object is stop sign",  # noqa: E501

tests/models/multimodal/processing/test_h2ovl.py

@@ -23,7 +23,7 @@ def _get_expected_num_patches(
     min_num: int,
     max_num: int,
 ):
-    from vllm.model_executor.models.h2ovl import (
+    from vllm.transformers_utils.processors.h2ovl import (
         calculate_h2ovl_targets,
         get_h2ovl_target_ratios,
     )

tests/models/multimodal/processing/test_internvl.py

@@ -23,7 +23,7 @@ def _get_expected_num_patches(
     min_num: int,
     max_num: int,
 ):
-    from vllm.model_executor.models.internvl import (
+    from vllm.transformers_utils.processors.internvl import (
         calculate_internvl_targets,
         get_internvl_target_ratios,
     )

tests/models/multimodal/processing/test_nemotron_vl.py

@@ -23,7 +23,7 @@ def _get_expected_num_patches(
     min_num: int,
     max_num: int,
 ):
-    from vllm.model_executor.models.nemotron_vl import (
+    from vllm.transformers_utils.processors.nemotron_vl import (
         calculate_nemotron_vl_targets,
         get_nemotron_vl_target_ratios,
     )

tests/models/multimodal/processing/test_qwen2_5_omni_embed.py

@@ -185,14 +185,16 @@ def make_mock_model(hidden: int = 8):
 
     # super().embed_input_ids → use SupportsMultiModal.embed_input_ids
     def fake_super_embed(
-        ids, mm_embs=None, *, is_multimodal=None, handle_oov_mm_token=False
+        ids,
+        mm_embs=None,
+        *,
+        is_multimodal=None,
     ):
         return SupportsMultiModal.embed_input_ids(
             model,
             ids,
             mm_embs,
             is_multimodal=is_multimodal,
-            handle_oov_mm_token=handle_oov_mm_token,
         )
 
     # Bind embed_input_ids as the real method

tests/models/registry.py

@@ -639,6 +639,11 @@ _LATE_INTERACTION_EXAMPLE_MODELS = {
     "OpsColQwen3Model": _HfExamplesInfo(
         "OpenSearch-AI/Ops-Colqwen3-4B", trust_remote_code=True
     ),
+    "ColQwen3_5": _HfExamplesInfo(
+        "athrael-soju/colqwen3.5-4.5B-v3",
+        trust_remote_code=True,
+        max_model_len=4096,
+    ),
     "Qwen3VLNemotronEmbedModel": _HfExamplesInfo(
         "nvidia/nemotron-colembed-vl-4b-v2",
     ),
@@ -1116,6 +1121,11 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         tokenizer_mode="mistral",
     ),
     # [Encoder-decoder]
+    "CohereASRForConditionalGeneration": _HfExamplesInfo(
+        "/host/engines/vllm/audio/2b-release",
+        trust_remote_code=True,
+        is_available_online=False,  # TODO (ekagra): revert after asr release
+    ),
     "NemotronParseForConditionalGeneration": _HfExamplesInfo(
         "nvidia/NVIDIA-Nemotron-Parse-v1.1", trust_remote_code=True
     ),

tests/multimodal/media/test_audio.py

@@ -1,11 +1,11 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import base64
 from pathlib import Path
 from unittest.mock import patch
 
 import librosa
 import numpy as np
+import pybase64 as base64
 import pytest
 
 from vllm.multimodal.media import AudioMediaIO

tests/multimodal/media/test_connector.py

@@ -2,13 +2,13 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import asyncio
-import base64
 import mimetypes
 import os
 from tempfile import NamedTemporaryFile, TemporaryDirectory
 
 import aiohttp
 import numpy as np
+import pybase64 as base64
 import pytest
 import requests
 import torch

tests/plugins/bge_m3_sparse_plugin/bge_m3_sparse_processor/sparse_embeddings_processor.py

@@ -3,10 +3,10 @@
 
 from collections.abc import Sequence
 
-from vllm.config import VllmConfig
+from vllm.config import ModelConfig, PoolerConfig, VllmConfig
 from vllm.entrypoints.openai.engine.protocol import UsageInfo
+from vllm.entrypoints.pooling.base.protocol import EmbedRequestMixin
 from vllm.inputs.data import PromptType
-from vllm.logger import init_logger
 from vllm.outputs import PoolingRequestOutput
 from vllm.plugins.io_processors.interface import (
     IOProcessor,
@@ -16,14 +16,13 @@ from vllm.renderers import BaseRenderer
 from vllm.tokenizers.detokenizer_utils import convert_ids_list_to_tokens
 
 from .types import (
+    EMBED_TASKS,
     SparseEmbeddingCompletionRequestMixin,
     SparseEmbeddingResponse,
     SparseEmbeddingResponseData,
     SparseEmbeddingTokenWeight,
 )
 
-logger = init_logger(__name__)
-
 
 class BgeM3SparseEmbeddingsProcessor(
     IOProcessor[SparseEmbeddingCompletionRequestMixin, SparseEmbeddingResponse]
@@ -33,6 +32,22 @@ class BgeM3SparseEmbeddingsProcessor(
         self.offline_requests: list[SparseEmbeddingCompletionRequestMixin] = []
         self.online_requests: dict[str, SparseEmbeddingCompletionRequestMixin] = {}
         self.renderer: BaseRenderer = renderer
+        self.default_pooling_params = {}
+        pooler_config: PoolerConfig = vllm_config.model_config.pooler_config
+        if pooler_config is not None:
+            for param in ["use_activation", "dimensions"]:
+                if getattr(pooler_config, param, None) is None:
+                    continue
+                self.default_pooling_params[param] = getattr(pooler_config, param)
+        self.embed_dimensions = vllm_config.model_config.embedding_size
+        self.embed_request_queue: list[EmbedRequestMixin] = []
+
+    def __repr__(self) -> str:
+        return (
+            f"BgeM3SparseEmbeddingsProcessor("
+            f"embed_dimensions={self.embed_dimensions}, "
+            f"default_pooling_params={self.default_pooling_params})"
+        )
 
     def merge_pooling_params(
         self,
@@ -41,7 +56,57 @@ class BgeM3SparseEmbeddingsProcessor(
         if params is None:
             params = PoolingParams()
         # refer to PoolingCompletionRequest.to_pooling_params
-        params.task = "token_classify"
+        # set and verify pooling params
+        params.skip_reading_prefix_cache = True
+
+        raw_embed_request = self.embed_request_queue.pop(0)
+        if raw_embed_request.embed_task not in EMBED_TASKS:
+            raise ValueError(
+                f"Unsupported task {raw_embed_request}, "
+                f"Supported tasks are {EMBED_TASKS}"
+            )
+        has_dense_embed = True
+        if raw_embed_request.embed_task == "dense":
+            params.task = "embed"
+            params.skip_reading_prefix_cache = False
+        elif raw_embed_request.embed_task == "sparse":
+            params.task = "token_classify"
+            has_dense_embed = False
+        else:
+            params.task = "embed&token_classify"
+        params.use_activation = raw_embed_request.use_activation
+        if params.use_activation is None:
+            params.use_activation = True
+        if not has_dense_embed:
+            params.dimensions = None
+            return params
+
+        params.dimensions = raw_embed_request.dimensions
+
+        model_config: ModelConfig = self.vllm_config.model_config
+        for param in self.default_pooling_params:
+            if getattr(params, param, None) is None:
+                setattr(params, param, self.default_pooling_params[param])
+
+        if params.dimensions is not None:
+            if not model_config.is_matryoshka:
+                raise ValueError(
+                    f'Model "{model_config.served_model_name}" does not '
+                    f"support matryoshka representation, "
+                    f"changing output dimensions will lead to poor results."
+                )
+
+            mds = model_config.matryoshka_dimensions
+            if mds is not None:
+                if params.dimensions not in mds:
+                    raise ValueError(
+                        f"Model {model_config.served_model_name!r} "
+                        f"only supports {str(mds)} matryoshka dimensions, "
+                        f"use other output dimensions will "
+                        f"lead to poor results."
+                    )
+            elif params.dimensions < 1:
+                raise ValueError("Dimensions must be greater than 0")
         return params
 
     def parse_request(
@@ -61,14 +126,16 @@ class BgeM3SparseEmbeddingsProcessor(
         if request_id is not None:
             assert request_id not in self.online_requests, "request_id duplicated"
             self.online_requests[request_id] = prompt
+            self.embed_request_queue.extend(prompt.to_embed_requests_online())
         else:
             self.offline_requests.append(prompt)
+            self.embed_request_queue.extend(prompt.to_embed_requests_offline())
         return prompt.input
 
     def _get_sparse_embedding_request(self, request_id: str | None = None):
         if request_id:
             return self.online_requests.pop(request_id, None)
-        return self.offline_requests.pop()
+        return self.offline_requests.pop(0)
 
     def _build_sparse_embedding_token_weights(
         self,
@@ -100,26 +167,45 @@ class BgeM3SparseEmbeddingsProcessor(
     ) -> SparseEmbeddingResponse:
         num_prompt_tokens = 0
         response_data = []
-        return_tokens = self._get_sparse_embedding_request(request_id).return_tokens
+        raw_request = self._get_sparse_embedding_request(request_id)
+        has_dense_embed = raw_request.embed_task in ["dense", "dense&sparse"]
+        has_sparse_embed = raw_request.embed_task in ["sparse", "dense&sparse"]
+        embed_dimensions = 0
+        if has_dense_embed:
+            embed_dimensions = (
+                self.embed_dimensions
+                if raw_request.dimensions is None
+                else raw_request.dimensions
+            )
         for idx in range(len(model_output)):
             mo = model_output[idx]
-            sparse_embedding: dict[int, float] = {}
+            sparse_embedding_dict: dict[int, float] = {}
             num_prompt_tokens += len(mo.prompt_token_ids)
-            if len(mo.prompt_token_ids) != len(mo.outputs.data):
-                # this is the case that add_special_tokens is True,
-                # which means first token and last token are special tokens
-                mo.prompt_token_ids = mo.prompt_token_ids[1:]
-            for token_id, weight in zip(mo.prompt_token_ids, mo.outputs.data.tolist()):
-                sparse_embedding[token_id] = max(
-                    weight, sparse_embedding.get(token_id, 0.0)
+            dense_embedding: list[float] | None = None
+            sparse_embedding: list[SparseEmbeddingTokenWeight] | None = None
+            if has_dense_embed:
+                dense_embedding = mo.outputs.data[:embed_dimensions].tolist()
+            if has_sparse_embed:
+                sparse_weights = mo.outputs.data[embed_dimensions:].tolist()
+                if len(mo.prompt_token_ids) != len(sparse_weights):
+                    # this is the case that add_special_tokens is True,
+                    # which means first token and last token are special tokens
+                    mo.prompt_token_ids = mo.prompt_token_ids[1:]
+                for token_id, weight in zip(mo.prompt_token_ids, sparse_weights):
+                    sparse_embedding_dict[token_id] = max(
+                        weight, sparse_embedding_dict.get(token_id, 0.0)
+                    )
+                sparse_embedding = self._build_sparse_embedding_token_weights(
+                    sparse_embedding_dict,
+                    raw_request.return_tokens,
                 )
+
             response_data.append(
                 SparseEmbeddingResponseData(
                     index=idx,
-                    sparse_embedding=self._build_sparse_embedding_token_weights(
-                        sparse_embedding,
-                        return_tokens,
-                    ),
+                    object=raw_request.embed_task,
+                    sparse_embedding=sparse_embedding,
+                    dense_embedding=dense_embedding,
                 )
             )
 

tests/plugins/bge_m3_sparse_plugin/bge_m3_sparse_processor/types.py

@@ -1,18 +1,44 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+from typing import Literal, get_args
+
 from pydantic import BaseModel, Field
 
 from vllm.entrypoints.openai.engine.protocol import UsageInfo
-from vllm.entrypoints.pooling.base.protocol import CompletionRequestMixin
+from vllm.entrypoints.pooling.base.protocol import (
+    CompletionRequestMixin,
+    EmbedRequestMixin,
+)
+
+EmbedTask = Literal[
+    "sparse",
+    "dense",
+    "dense&sparse",
+]
+
+EMBED_TASKS: tuple[EmbedTask, ...] = get_args(EmbedTask)
 
 
-class SparseEmbeddingCompletionRequestMixin(CompletionRequestMixin):
+class SparseEmbeddingCompletionRequestMixin(CompletionRequestMixin, EmbedRequestMixin):
     return_tokens: bool | None = Field(
         default=None,
         description="Whether to return dict shows the mapping of token_id to text."
         "`None` or False means not return.",
     )
+    embed_task: EmbedTask = Field(
+        default="dense&sparse",
+        description="embed task, can be one of 'sparse', 'dense' , 'dense&sparse', "
+        "default to 'dense&sparse'",
+    )
+
+    def to_embed_requests_offline(self) -> list[EmbedRequestMixin]:
+        if isinstance(self.input, list):
+            return [self] * len(self.input)
+        return [self]
+
+    def to_embed_requests_online(self) -> list[EmbedRequestMixin]:
+        return [self]
 
 
 class SparseEmbeddingTokenWeight(BaseModel):
@@ -23,8 +49,9 @@ class SparseEmbeddingTokenWeight(BaseModel):
 
 class SparseEmbeddingResponseData(BaseModel):
     index: int
-    object: str = "sparse-embedding"
-    sparse_embedding: list[SparseEmbeddingTokenWeight]
+    object: str = "dense&sparse"
+    sparse_embedding: list[SparseEmbeddingTokenWeight] | None
+    dense_embedding: list[float] | None
 
 
 class SparseEmbeddingResponse(BaseModel):

tests/plugins/prithvi_io_processor_plugin/prithvi_io_processor/prithvi_processor.py

@@ -1,7 +1,6 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import base64
 import datetime
 import os
 import tempfile
@@ -11,6 +10,7 @@ from typing import Any
 
 import albumentations
 import numpy as np
+import pybase64 as base64
 import rasterio
 import regex as re
 import torch

tests/plugins_tests/test_bge_m3_sparse_io_processor_plugins.py

@@ -19,6 +19,12 @@ model_config = {
     ),
 }
 
+dense_embedding_sum = [
+    -0.7214539647102356,  # "What is the capital of France?"
+    -0.6926871538162231,  # "What is the capital of Germany?"
+    -0.7129564881324768,  # "What is the capital of Spain?"
+]
+
 
 def _float_close(expected: object, result: object):
     assert isinstance(expected, float) and isinstance(result, float), (
@@ -33,6 +39,12 @@ def _get_attr_or_val(obj: object | dict, key: str):
     return getattr(obj, key, None)
 
 
+def _check_dense_embedding(data, index=0):
+    assert _float_close(sum(data), dense_embedding_sum[index]), (
+        "dense-embedding result not match"
+    )
+
+
 def _check_sparse_embedding(data, check_tokens=False):
     expected_weights = [
         {"token_id": 32, "weight": 0.0552978515625, "token": "?"},
@@ -109,7 +121,7 @@ async def test_bge_m3_sparse_plugin_online(
     assert len(_get_attr_or_val(parsed_response, "data")) > 0
 
     data_entry = _get_attr_or_val(parsed_response, "data")[0]
-    assert _get_attr_or_val(data_entry, "object") == "sparse-embedding"
+    assert _get_attr_or_val(data_entry, "object") == "dense&sparse"
     assert _get_attr_or_val(data_entry, "sparse_embedding")
 
     # Verify sparse embedding format
@@ -117,6 +129,11 @@ async def test_bge_m3_sparse_plugin_online(
     assert isinstance(sparse_embedding, list)
     _check_sparse_embedding(sparse_embedding, return_tokens)
 
+    # Verify dense embedding format
+    dense_embedding = _get_attr_or_val(data_entry, "dense_embedding")
+    assert isinstance(dense_embedding, list)
+    _check_dense_embedding(dense_embedding)
+
     # Verify usage information
     usage = _get_attr_or_val(parsed_response, "usage")
     assert usage, f"usage not found for {parsed_response}"
@@ -164,6 +181,9 @@ def test_bge_m3_sparse_plugin_offline(vllm_runner, return_tokens: bool):
         sparse_embedding = output.sparse_embedding
         assert isinstance(sparse_embedding, list)
         _check_sparse_embedding(sparse_embedding, return_tokens)
+        dense_embedding = output.dense_embedding
+        assert isinstance(dense_embedding, list)
+        _check_dense_embedding(dense_embedding)
 
     # Verify usage
     assert response.usage.prompt_tokens > 0
@@ -206,6 +226,9 @@ def test_bge_m3_sparse_plugin_offline_multiple_inputs(vllm_runner):
         # Each output should have sparse embeddings
         sparse_embedding = output.sparse_embedding
         assert isinstance(sparse_embedding, list)
+        dense_embedding = output.dense_embedding
+        assert isinstance(dense_embedding, list)
+        _check_dense_embedding(dense_embedding, i)
 
     # Verify usage
     assert response.usage.prompt_tokens > 0

tests/plugins_tests/test_terratorch_io_processor_plugins.py

@@ -1,9 +1,9 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import base64
 import io
 
 import imagehash
+import pybase64 as base64
 import pytest
 import requests
 from PIL import Image

tests/renderers/test_sparse_tensor_validation.py

@@ -5,9 +5,9 @@ Tests verify that malicious sparse tensors are rejected before they can trigger
 out-of-bounds memory writes during to_dense() operations.
 """
 
-import base64
 import io
 
+import pybase64 as base64
 import pytest
 import torch
 

tests/tool_parsers/test_glm47_moe_tool_parser.py

@@ -0,0 +1,168 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# ruff: noqa: E501
+"""Tests for the GLM-4.7 tool call parser."""
+
+import json
+from unittest.mock import Mock
+
+import pytest
+
+from vllm.entrypoints.openai.chat_completion.protocol import (
+    ChatCompletionRequest,
+    ChatCompletionToolsParam,
+    FunctionDefinition,
+)
+from vllm.tokenizers import get_tokenizer
+from vllm.tool_parsers.glm47_moe_tool_parser import Glm47MoeModelToolParser
+
+MODEL = "zai-org/GLM-4.5"
+
+
+@pytest.fixture(scope="module")
+def glm47_tokenizer():
+    return get_tokenizer(tokenizer_name=MODEL)
+
+
+@pytest.fixture
+def glm47_tool_parser(glm47_tokenizer):
+    return Glm47MoeModelToolParser(glm47_tokenizer)
+
+
+@pytest.fixture
+def mock_request() -> ChatCompletionRequest:
+    request = Mock(spec=ChatCompletionRequest)
+    request.tools = [
+        ChatCompletionToolsParam(
+            function=FunctionDefinition(name="get_current_date", parameters={}),
+        ),
+        ChatCompletionToolsParam(
+            function=FunctionDefinition(
+                name="get_weather",
+                parameters={
+                    "type": "object",
+                    "properties": {
+                        "city": {"type": "string"},
+                        "date": {"type": "string"},
+                    },
+                },
+            ),
+        ),
+    ]
+    request.tool_choice = "auto"
+    return request
+
+
+class TestGlm47ExtractToolCalls:
+    def test_no_tool_call(self, glm47_tool_parser, mock_request):
+        out = "This is a plain response."
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert not r.tools_called
+        assert r.content == out
+
+    def test_zero_arg_inline(self, glm47_tool_parser, mock_request):
+        out = "<tool_call>get_current_date</tool_call>"
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert r.tools_called
+        assert r.tool_calls[0].function.name == "get_current_date"
+        assert json.loads(r.tool_calls[0].function.arguments) == {}
+        assert r.content is None
+
+    def test_zero_arg_newline(self, glm47_tool_parser, mock_request):
+        out = "<tool_call>get_current_date\n</tool_call>"
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert r.tools_called
+        assert r.tool_calls[0].function.name == "get_current_date"
+
+    def test_args_same_line(self, glm47_tool_parser, mock_request):
+        out = "<tool_call>get_weather<arg_key>city</arg_key><arg_value>Beijing</arg_value></tool_call>"
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert r.tools_called
+        assert json.loads(r.tool_calls[0].function.arguments) == {"city": "Beijing"}
+
+    def test_args_with_newlines(self, glm47_tool_parser, mock_request):
+        out = "<tool_call>get_weather\n<arg_key>city</arg_key>\n<arg_value>Beijing</arg_value>\n</tool_call>"
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert r.tools_called
+        assert json.loads(r.tool_calls[0].function.arguments) == {"city": "Beijing"}
+
+    def test_content_before(self, glm47_tool_parser, mock_request):
+        out = "Checking.<tool_call>get_current_date</tool_call>"
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert r.tools_called
+        assert r.content == "Checking."
+
+    def test_multiple(self, glm47_tool_parser, mock_request):
+        out = (
+            "<tool_call>get_weather<arg_key>city</arg_key><arg_value>Beijing</arg_value></tool_call>"
+            "<tool_call>get_weather<arg_key>city</arg_key><arg_value>Shanghai</arg_value></tool_call>"
+        )
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert len(r.tool_calls) == 2
+
+    def test_empty_content_none(self, glm47_tool_parser, mock_request):
+        out = "<tool_call>get_current_date</tool_call>"
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert r.content is None
+
+    def test_whitespace_content_none(self, glm47_tool_parser, mock_request):
+        out = "  \n  <tool_call>get_current_date</tool_call>"
+        r = glm47_tool_parser.extract_tool_calls(out, request=mock_request)
+        assert r.content is None
+
+
+def _reset(parser):
+    parser._buffer = ""
+    parser._in_tool_call = False
+    parser.current_tool_name_sent = False
+    parser._current_tool_name = None
+    parser._pending_key = None
+    parser._streaming_string_value = False
+    parser.prev_tool_call_arr = []
+    parser.current_tool_id = -1
+    parser.streamed_args_for_tool = []
+    parser._tool_call_ids = []
+    parser._args_started = []
+    parser._args_closed = []
+    parser._seen_keys = []
+
+
+class TestGlm47Streaming:
+    def test_no_args(self, glm47_tool_parser, mock_request):
+        _reset(glm47_tool_parser)
+        for chunk in ["<tool_call>", "get_current_date", "</tool_call>"]:
+            glm47_tool_parser.extract_tool_calls_streaming(
+                previous_text="",
+                current_text="",
+                delta_text=chunk,
+                previous_token_ids=[],
+                current_token_ids=[],
+                delta_token_ids=[],
+                request=mock_request,
+            )
+        assert len(glm47_tool_parser.prev_tool_call_arr) >= 1
+
+    def test_with_args(self, glm47_tool_parser, mock_request):
+        _reset(glm47_tool_parser)
+        # Split chunks so that the incremental string streaming path
+        # processes the value, its closing tag, and the tool-call closing
+        # tag in separate calls.
+        for chunk in [
+            "<tool_call>",
+            "get_weather\n",
+            "<arg_key>city</arg_key>",
+            "<arg_value>",
+            "Beijing",
+            "</arg_value>",
+            "</tool_call>",
+        ]:
+            glm47_tool_parser.extract_tool_calls_streaming(
+                previous_text="",
+                current_text="",
+                delta_text=chunk,
+                previous_token_ids=[],
+                current_token_ids=[],
+                delta_token_ids=[],
+                request=mock_request,
+            )
+        assert glm47_tool_parser.prev_tool_call_arr[0]["arguments"]["city"] == "Beijing"

tests/tool_parsers/test_glm4_moe_tool_parser.py

@@ -107,7 +107,7 @@ def test_extract_tool_calls_no_tools(glm4_moe_tool_parser, mock_request):
                     )
                 )
             ],
-            "",
+            None,
         ),
         (
             """<tool_call>get_current_weather
@@ -152,7 +152,7 @@ def test_extract_tool_calls_no_tools(glm4_moe_tool_parser, mock_request):
                     )
                 ),
             ],
-            "",
+            None,
         ),
         (
             """I'll help you check the weather. <tool_call>get_current_weather
@@ -202,7 +202,7 @@ def test_extract_tool_calls_no_tools(glm4_moe_tool_parser, mock_request):
                     )
                 )
             ],
-            "",
+            None,
         ),
         (
             """I will help you get the weather.<tool_call>get_weather

tests/utils.py

@@ -225,13 +225,31 @@ class RemoteVLLMServer:
             )
 
         self._start_server(model, vllm_serve_args, env_dict)
-        max_wait_seconds = max_wait_seconds or 360
-        self._wait_for_server(url=self.url_for("health"), timeout=max_wait_seconds)
+        max_wait_seconds = max_wait_seconds or 480
+        try:
+            self._wait_for_server(url=self.url_for("health"), timeout=max_wait_seconds)
+        except Exception:
+            # If the server never became healthy, we must still clean up
+            # the subprocess tree. Without this, a timeout in __init__
+            # leaks the server + EngineCore processes (and their GPU
+            # memory), because __exit__ is never called when __init__
+            # raises inside a ``with`` statement.
+            self._shutdown()
+            raise
 
     def __enter__(self):
         return self
 
     def __exit__(self, exc_type, exc_value, traceback):
+        self._shutdown()
+
+    def _shutdown(self) -> None:
+        """Kill the server process tree and wait for GPU memory release.
+
+        Called from both ``__exit__`` (normal path) and ``__init__``
+        (when the server fails to start). Must be safe to call even if
+        the process is already dead.
+        """
         pid = self.proc.pid
 
         # Get the process group ID. Because we used
@@ -265,33 +283,92 @@ class RemoteVLLMServer:
                 self.proc.wait(timeout=10)
                 print(f"[RemoteOpenAIServer] Server {pid} killed")
             except subprocess.TimeoutExpired:
-                # Phase 3: last resort - find and kill any orphaned children
-                self._kill_orphaned_children(pid)
+                pass
 
-        # Wait for GPU memory to actually be *freed*, not just
+        # After killing the root process, ensure all children in the
+        # process group (e.g. EngineCore workers) are also dead.
+        # On ROCm especially, surviving children hold GPU contexts and
+        # prevent VRAM from being reclaimed by the driver.
+        self._kill_process_group_survivors(pgid)
+
+        # Wait for GPU memory to actually be freed, not just
         # "stabilized at whatever level it's at".
         self._wait_for_gpu_memory_release()
 
-    def _kill_orphaned_children(self, parent_pid: int) -> None:
-        """Best-effort cleanup of any lingering child processes."""
-        try:
-            import psutil
+    def _kill_process_group_survivors(
+        self, pgid: int | None, timeout: float = 15.0
+    ) -> None:
+        """SIGKILL any processes still in the server's process group
+        and wait for them to exit.
 
-            parent = psutil.Process(parent_pid)
-            children = parent.children(recursive=True)
-            for child in children:
-                print(
-                    f"[RemoteOpenAIServer] Killing orphaned child "
-                    f"pid={child.pid} name={child.name()}"
-                )
-                child.kill()
-            psutil.wait_procs(children, timeout=5)
-        except Exception as e:
-            # psutil may not be installed, or processes already gone
-            print(f"[RemoteOpenAIServer] Orphan cleanup failed: {e}")
-            # Fallback: try to kill by pgid one more time
-            with contextlib.suppress(ProcessLookupError, OSError):
-                os.killpg(parent_pid, signal.SIGKILL)
+        Because the server is launched with ``start_new_session=True``,
+        all its children (EngineCore, workers, etc.) share the same
+        pgid. After the root process is killed, stragglers -- especially
+        on ROCm where GPU contexts linger until the *process* exits --
+        must be reaped explicitly.
+
+        Uses ``/proc`` to scan for pgid members so this works even after
+        the parent has been reaped (unlike ``psutil.Process.children``).
+        """
+        if pgid is None:
+            return
+
+        # Send SIGKILL to the entire process group one more time.
+        # This is cheap and harmless if everyone is already dead.
+        with contextlib.suppress(ProcessLookupError, OSError):
+            os.killpg(pgid, signal.SIGKILL)
+
+        # Collect surviving PIDs by scanning /proc for matching pgid.
+        # This works on Linux even after the parent has been waited on
+        # and is more reliable than psutil.Process(parent).children().
+        survivor_pids = self._find_pgid_members(pgid)
+
+        if not survivor_pids:
+            return
+
+        print(
+            f"[RemoteOpenAIServer] {len(survivor_pids)} process(es) still "
+            f"in pgid {pgid} after SIGKILL: {survivor_pids}"
+        )
+
+        # Wait for each survivor to actually exit so the GPU driver
+        # releases its VRAM.
+        deadline = time.time() + timeout
+        while survivor_pids and time.time() < deadline:
+            still_alive = []
+            for spid in survivor_pids:
+                try:
+                    os.kill(spid, 0)  # Check if still alive
+                    still_alive.append(spid)
+                except (ProcessLookupError, OSError):
+                    pass
+            survivor_pids = still_alive
+            if survivor_pids:
+                time.sleep(0.5)
+
+        if survivor_pids:
+            print(
+                f"[RemoteOpenAIServer] WARNING: processes {survivor_pids} "
+                f"in pgid {pgid} could not be killed within {timeout}s"
+            )
+
+    @staticmethod
+    def _find_pgid_members(pgid: int) -> list[int]:
+        """Return PIDs of all living processes whose pgid matches."""
+        members: list[int] = []
+        proc_path = Path("/proc")
+        if not proc_path.is_dir():
+            return members
+        for entry in proc_path.iterdir():
+            if not entry.name.isdigit():
+                continue
+            pid = int(entry.name)
+            try:
+                if os.getpgid(pid) == pgid:
+                    members.append(pid)
+            except OSError:
+                continue
+        return members
 
     def _get_gpu_memory_used(self) -> float | None:
         """Get total GPU memory used across all visible devices in bytes."""
@@ -318,13 +395,16 @@ class RemoteVLLMServer:
             return None
         return None
 
-    def _wait_for_gpu_memory_release(self, timeout: float = 60.0):
+    def _wait_for_gpu_memory_release(
+        self, timeout: float = 120.0, log_interval: float = 10.0
+    ):
         """Wait for GPU memory to drop back toward pre-server levels.
 
-        Two-phase strategy:
-          1. Try to wait for memory to return close to pre-server baseline.
-          2. If that doesn't happen, fall back to waiting for stabilization
-             and log a warning (the next server might still OOM).
+        Waits the full timeout for memory to return close to the
+        pre-server baseline. Does NOT fall back to a "stabilization"
+        heuristic -- if memory is still held when the timeout expires,
+        the test fails so the problem is surfaced immediately rather
+        than causing cascading OOM failures in every subsequent test.
         """
         baseline = self._pre_server_gpu_memory
         if baseline is None:
@@ -337,8 +417,7 @@ class RemoteVLLMServer:
         target = baseline + headroom_bytes
 
         start = time.time()
-        last_used: float | None = None
-        stable_count = 0
+        next_log_time = start + log_interval
 
         while time.time() - start < timeout:
             used = self._get_gpu_memory_used()
@@ -350,7 +429,6 @@ class RemoteVLLMServer:
             target_gb = target / 1e9
             elapsed = time.time() - start
 
-            # Phase 1: memory dropped to near baseline - we're done.
             if used <= target:
                 print(
                     f"[RemoteOpenAIServer] GPU memory released to "
@@ -359,28 +437,19 @@ class RemoteVLLMServer:
                 )
                 return
 
-            # Phase 2 (after 40s): fall back to stabilization check.
-            # This handles cases where another process is using GPU memory
-            # and we'll never reach baseline.
-            if elapsed > 40.0 and last_used is not None:
-                delta = abs(used - last_used)
-                if delta < 200 * 1024 * 1024:  # 200 MB
-                    stable_count += 1
-                    if stable_count >= 3:
-                        print(
-                            f"[RemoteOpenAIServer] WARNING: GPU memory "
-                            f"stabilized at {used_gb:.2f} GB "
-                            f"(target was {target_gb:.2f} GB). "
-                            f"Proceeding - next server may OOM."
-                        )
-                        return
-                else:
-                    stable_count = 0
+            now = time.time()
+            if now >= next_log_time:
+                print(
+                    f"[RemoteOpenAIServer] Waiting for GPU memory release: "
+                    f"{used_gb:.2f} GB (target: {target_gb:.2f} GB) "
+                    f"[{elapsed:.0f}s/{timeout:.0f}s]"
+                )
+                next_log_time = now + log_interval
 
-            last_used = used
             time.sleep(1.0)
 
-        # Timeout - log clearly so CI failures are diagnosable
+        # Timeout -- raise so the current test fails with a clear
+        # message instead of silently poisoning subsequent tests.
         final_used = self._get_gpu_memory_used()
         final_gb = final_used / 1e9 if final_used else 0.0
         raise RuntimeError(
@@ -534,7 +603,9 @@ class RemoteLaunchRenderServer(RemoteVLLMServer):
                 revision=model_config.tokenizer_revision,
             )
 
-    def _wait_for_gpu_memory_release(self, timeout: float = 30.0):
+    def _wait_for_gpu_memory_release(
+        self, timeout: float = 30.0, log_interval: float = 10.0
+    ):
         pass  # No GPU used
 
 

tests/v1/core/utils.py

@@ -47,7 +47,7 @@ def create_scheduler(
     enable_prefix_caching: bool = False,
     long_prefill_token_threshold: int = 0,
     disable_chunked_mm_input: bool = False,
-    use_kv_connector: None | bool | MockKVConfig = None,
+    use_kv_connector: None | bool | str | MockKVConfig = None,
     num_blocks: int = 10000,
     block_size: int = 16,
     max_model_len: int | None = None,
@@ -107,6 +107,11 @@ def create_scheduler(
                 "is_async": use_kv_connector.is_async,
             },
         )
+    elif isinstance(use_kv_connector, str):
+        kv_transfer_config = KVTransferConfig(
+            kv_connector=use_kv_connector,
+            kv_role="kv_both",
+        )
     elif use_kv_connector:
         kv_transfer_config = KVTransferConfig(
             kv_connector="ExampleConnector",

tests/v1/e2e/test_hybrid_chunked_prefill.py

@@ -0,0 +1,104 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+
+from vllm import SamplingParams
+from vllm.platforms import current_platform
+
+from ...utils import large_gpu_mark, multi_gpu_marks
+
+# A trivial request with a short prompt to ensure we run a mixed batch
+SMALL_MESSAGE = [
+    {
+        "role": "user",
+        "content": "The secret beta value is 64. What is the secret beta?",
+    }
+]
+
+# Sample prompt with a bunch of filler in between the critical fact and the request.
+# Both parts need to be processed properly for the model to generate the correct answer
+MESSAGES = [
+    {
+        "role": "user",
+        "content": (
+            "Important: The secret number is 42. "
+            "The sky is green in this hypothetical world. "
+            "Apples grow on trees in the forest. "
+            "Rivers flow through the valleys and mountains. "
+            "Birds sing songs in the early morning light. "
+            "The weather today is sunny with clear skies ahead. "
+            "Flowers bloom in the garden during spring season. "
+            "Now answer with ONLY the number and nothing else: "
+            "What is the secret number plus one?"
+        ),
+    }
+]
+
+
+@pytest.mark.skipif(not current_platform.is_cuda(), reason="CUDA not available")
+@pytest.mark.parametrize(
+    "model_name",
+    [
+        pytest.param("Qwen/Qwen3.5-4B", marks=[large_gpu_mark(min_gb=40)]),
+        pytest.param(
+            "nvidia/NVIDIA-Nemotron-3-Super-120B-A12B-FP8",
+            marks=[large_gpu_mark(min_gb=80)] + multi_gpu_marks(num_gpus=2),
+        ),
+    ],
+)
+@pytest.mark.parametrize("enable_prefix_caching", [False, True])
+def test_mtp_speculative_mixed_batch_short_prefill(
+    vllm_runner, model_name, enable_prefix_caching
+):
+    """Test to ensure MTP speculative decoding correctly handles
+    short prefill chunks that fall below the reorder_batch_threshold."""
+
+    # Set so large that both prefills will be classified as decodes in a mixed batch
+    # note, with prefix caching we require chunk_size >= mamba_block_size
+    chunk_size = 256 if not enable_prefix_caching else 16384
+    num_draft_tokens = 100
+
+    with vllm_runner(
+        model_name,
+        speculative_config={
+            "method": "mtp",
+            "num_speculative_tokens": num_draft_tokens,
+        },
+        max_num_batched_tokens=chunk_size,
+        max_model_len=512,
+        enforce_eager=True,
+        tensor_parallel_size=2,
+        trust_remote_code=True,
+        enable_chunked_prefill=True,
+        enable_prefix_caching=enable_prefix_caching,
+        mamba_cache_mode="align" if enable_prefix_caching else "none",
+    ) as llm:
+        sampling_params = SamplingParams(
+            temperature=0.0,
+            max_tokens=128,
+        )
+
+        # First small message gets prefilled first, under normal conditions since the
+        # batch is not yet mixed. Then the second prefill arrives as a mixed batch, but
+        # is shorter than num_speculative_tokens, so it gets misclassified as a decode
+        # and processed with the wrong state management logic,  causing the critical
+        # fact from the first chunk to be lost and the model to generate nonsense.
+        outputs = llm.get_llm().chat(
+            [SMALL_MESSAGE, MESSAGES],
+            sampling_params,
+            chat_template_kwargs={"enable_thinking": False},
+        )
+
+        responses = []
+        for output in outputs:
+            generated_text = output.outputs[0].text
+            print(f"Generated text: {generated_text!r}")
+            responses.append(generated_text)
+
+        assert "64" in responses[0], (
+            "The first response should contain the correct value of 64."
+        )
+        assert "43" in responses[1], (
+            "The second response should contain the correct value of 42+1=43."
+        )

tests/v1/entrypoints/openai/serving_responses/conftest.py

@@ -1,44 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import pytest
-import pytest_asyncio
-
-from tests.utils import RemoteOpenAIServer
-
-# Use a small reasoning model to test the responses API.
-MODEL_NAME = "Qwen/Qwen3-1.7B"
-
-
-@pytest.fixture(scope="module")
-def default_server_args():
-    return [
-        "--max-model-len",
-        "8192",
-        "--enforce-eager",  # For faster startup.
-        "--enable-auto-tool-choice",
-        "--structured-outputs-config.backend",
-        "xgrammar",
-        "--tool-call-parser",
-        "hermes",
-        "--reasoning-parser",
-        "qwen3",
-    ]
-
-
-@pytest.fixture(scope="module")
-def server_with_store(default_server_args):
-    with RemoteOpenAIServer(
-        MODEL_NAME,
-        default_server_args,
-        env_dict={
-            "VLLM_ENABLE_RESPONSES_API_STORE": "1",
-            "VLLM_SERVER_DEV_MODE": "1",
-        },
-    ) as remote_server:
-        yield remote_server
-
-
-@pytest_asyncio.fixture
-async def client(server_with_store):
-    async with server_with_store.get_async_client() as async_client:
-        yield async_client

tests/v1/executor/test_executor.py

@@ -14,12 +14,35 @@ from vllm.engine.arg_utils import AsyncEngineArgs, EngineArgs
 from vllm.sampling_params import SamplingParams
 from vllm.v1.engine.async_llm import AsyncLLM
 from vllm.v1.engine.llm_engine import LLMEngine
+from vllm.v1.executor.abstract import Executor
 from vllm.v1.executor.multiproc_executor import MultiprocExecutor
+from vllm.v1.executor.uniproc_executor import (
+    ExecutorWithExternalLauncher,
+    UniProcExecutor,
+)
 
 
 class Mock: ...
 
 
+def test_supports_async_scheduling_base_executor():
+    assert Executor.supports_async_scheduling() is False
+
+
+def test_supports_async_scheduling_uniproc_executor():
+    assert UniProcExecutor.supports_async_scheduling() is True
+
+
+def test_supports_async_scheduling_executor_with_external_launcher():
+    # ExecutorWithExternalLauncher inherits from UniProcExecutor and does not
+    # override supports_async_scheduling, so it should return True.
+    assert ExecutorWithExternalLauncher.supports_async_scheduling() is True
+
+
+def test_supports_async_scheduling_multiproc_executor():
+    assert MultiprocExecutor.supports_async_scheduling() is True
+
+
 class CustomMultiprocExecutor(MultiprocExecutor):
     def collective_rpc(
         self,

tests/v1/kv_connector/unit/test_multi_connector.py

@@ -231,10 +231,11 @@ def test_multi_example_connector_consistency():
     ]
     # First three events are from initialization (register_kv_caches,
     # set_host_xfer_buffer_ops, get_handshake_metadata), then generate() events.
-    assert events["storage1-WORKER"][:7] == [
+    assert events["storage1-WORKER"][:8] == [
         "register_kv_caches",
         "set_host_xfer_buffer_ops",
         "get_handshake_metadata",
+        "handle_preemptions",
         "bind_connector_metadata",
         "start_load_kv",
         "wait_for_layer_load",
@@ -246,10 +247,11 @@ def test_multi_example_connector_consistency():
         "update_state_after_alloc num_blocks=[0] 0",
         "build_connector_meta",
     ]
-    assert events["storage2-WORKER"][:7] == [
+    assert events["storage2-WORKER"][:8] == [
         "register_kv_caches",
         "set_host_xfer_buffer_ops",
         "get_handshake_metadata",
+        "handle_preemptions",
         "bind_connector_metadata",
         "start_load_kv",
         "wait_for_layer_load",
@@ -399,8 +401,8 @@ def test_multi_connector_handle_preemptions_integration():
         # testing the delegation behavior of MultiConnector here.
         # The connector attribute contains the KV connector.
         assert scheduler.connector is not None, "Scheduler should have a connector"
-        preempted_req_ids = {"req-1", "req-2", "req-3"}
-        scheduler.connector.handle_preemptions(preempted_req_ids)
+        connector_md = scheduler.connector.build_connector_meta(scheduler.schedule())
+        scheduler.connector.handle_preemptions(connector_md)
 
         # Verify both connectors received the handle_preemptions call
         events = get_connector_events()

tests/v1/kv_connector/unit/test_nixl_connector.py

@@ -1369,7 +1369,13 @@ def test_abort_timeout_on_prefiller(monkeypatch, distributed_executor_backend):
                     "NIXL_TELEMETRY_ENABLE": "1",
                 },
             }
-            ray.init(runtime_env=runtime_env)
+            # On XPU/ROCm, vLLM expects Ray's device key to be "GPU".
+            # Explicitly reserving GPU resources here prevents false negatives
+            # when Ray cannot auto-detect accelerator resources in test envs.
+            ray_init_kwargs: dict[str, Any] = {"runtime_env": runtime_env}
+            if not current_platform.is_cuda():
+                ray_init_kwargs["num_gpus"] = 1
+            ray.init(**ray_init_kwargs)
             try:
                 run_test_and_cleanup()
             finally:

tests/v1/kv_connector/unit/test_offloading_connector.py

@@ -363,10 +363,7 @@ class RequestRunner:
             assert kv_connector_metadata is not None
             assert isinstance(kv_connector_metadata, OffloadingConnectorMetadata)
 
-            if scheduler_output.preempted_req_ids:
-                self.worker_connector.handle_preemptions(
-                    scheduler_output.preempted_req_ids
-                )
+            self.worker_connector.handle_preemptions(kv_connector_metadata)
 
             self.worker_connector.bind_connector_metadata(kv_connector_metadata)
             self.worker_connector.start_load_kv(self._dummy_ctx)

tests/v1/kv_connector/unit/test_scheduler_kv_connector_override.py

@@ -0,0 +1,130 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+from unittest.mock import MagicMock, patch
+
+import pytest
+
+import vllm.plugins as plugins_module
+from tests.v1.core.utils import create_requests, create_scheduler
+from vllm.distributed.kv_transfer.kv_connector.factory import (
+    KVConnectorFactory,
+)
+from vllm.distributed.kv_transfer.kv_connector.v1.base import (
+    KVConnectorBase_V1,
+    KVConnectorMetadata,
+)
+from vllm.v1.core.kv_cache_manager import KVCacheBlocks
+from vllm.v1.core.kv_cache_utils import BlockHash
+from vllm.v1.core.sched.output import SchedulerOutput
+from vllm.v1.core.sched.scheduler import Scheduler
+from vllm.v1.request import Request
+
+
+class DummyConnectorMetadata(KVConnectorMetadata):
+    def __init__(self, block_hashes_by_req: dict[str, list[BlockHash]]):
+        self.block_hashes_by_req = block_hashes_by_req
+
+
+class DummyKVConnector(KVConnectorBase_V1):
+    def __init__(self, vllm_config, role, kv_cache_config=None):
+        super().__init__(vllm_config, role, kv_cache_config)
+
+    def get_num_new_matched_tokens(
+        self, request: Request, num_computed_tokens: int
+    ) -> tuple[int | None, bool]:
+        return (0, False)
+
+    def update_state_after_alloc(
+        self, request: Request, blocks: KVCacheBlocks, num_external_tokens: int
+    ):
+        pass
+
+    def build_connector_meta(
+        self, scheduler_output: SchedulerOutput
+    ) -> KVConnectorMetadata:
+        block_hashes_by_req = getattr(scheduler_output, "block_hashes_by_req", None)
+        assert block_hashes_by_req is not None, (
+            "DummyKVConnector expected 'block_hashes_by_req' on scheduler_output"
+        )
+        return DummyConnectorMetadata(
+            block_hashes_by_req=block_hashes_by_req,
+        )
+
+    def start_load_kv(self, kv_caches, finished_req_ids):
+        pass
+
+    def wait_for_layer_load(self, layer_name):
+        pass
+
+    def save_kv_layer(self, layer_name, kv_layer, attn_metadata, **kwargs):
+        pass
+
+    def wait_for_save(self):
+        pass
+
+
+def _my_plugin():
+    """Registers the dummy KV connector and overrides _build_kv_connector_meta"""
+    KVConnectorFactory.register_connector(
+        "DummyKVConnector",
+        __name__,
+        DummyKVConnector.__name__,
+    )
+
+    def _custom_build_kv_connector_meta(
+        self, connector: KVConnectorBase_V1, scheduler_output: SchedulerOutput
+    ) -> KVConnectorMetadata:
+        block_hashes_by_req: dict[str, list[BlockHash]] = {}
+        for req_id in scheduler_output.num_scheduled_tokens:
+            request = self.requests[req_id]
+            block_hashes_by_req[req_id] = request.block_hashes
+
+        scheduler_output.block_hashes_by_req = block_hashes_by_req  # type: ignore[attr-defined]
+        return connector.build_connector_meta(scheduler_output)
+
+    Scheduler._build_kv_connector_meta = _custom_build_kv_connector_meta
+
+
+@pytest.fixture
+def _load_plugin():
+    """Load the fake plugin through the real load_general_plugins() path."""
+    ep = MagicMock()
+    ep.name = "dummy_kv_connector_plugin"
+    ep.value = f"{__name__}:_my_plugin"
+    ep.load.return_value = _my_plugin
+
+    # Reset the global guard so load_general_plugins() actually runs.
+    plugins_module.plugins_loaded = False
+    with patch("importlib.metadata.entry_points", return_value=[ep]):
+        plugins_module.load_general_plugins()
+        yield
+    # Reset again so other tests are not affected.
+    plugins_module.plugins_loaded = False
+
+
+def test_connector_receives_block_hashes(_load_plugin):
+    block_size = 16
+    num_tokens = 48  # 3 full blocks worth of tokens
+    scheduler = create_scheduler(
+        use_kv_connector="DummyKVConnector", block_size=block_size
+    )
+    requests = create_requests(
+        num_requests=3, num_tokens=num_tokens, block_size=block_size
+    )
+    for req in requests:
+        scheduler.add_request(req)
+
+    output = scheduler.schedule()
+
+    # Verify the connector metadata was built with block hashes.
+    meta = output.kv_connector_metadata
+    assert isinstance(meta, DummyConnectorMetadata)
+    assert len(meta.block_hashes_by_req) == 3
+
+    for req in requests:
+        assert req.request_id in meta.block_hashes_by_req
+        # Each request has num_tokens / block_size = 3 full block hashes.
+        assert len(meta.block_hashes_by_req[req.request_id]) == (
+            num_tokens // block_size
+        )
+        assert meta.block_hashes_by_req[req.request_id] == req.block_hashes

tests/v1/kv_offload/test_cpu_offloading.py

@@ -22,6 +22,17 @@ if current_platform.is_cuda():
 elif current_platform.is_rocm():
     ATTN_BACKENDS = ["TRITON_ATTN"]
 
+# Maximum time (seconds) to wait for the async CPU offload transfer
+# to complete before giving up.
+_RESET_CACHE_TIMEOUT = 30 if current_platform.is_rocm() else 10
+
+# ZMQ poll timeout (ms) for the first event.
+_FIRST_EVENT_POLL_MS = 10_000 if current_platform.is_rocm() else 1000
+
+# Hard ceiling (seconds) on how long get_new_cpu_stored_events may loop,
+# to prevent hangs if non-CPU events keep arriving indefinitely.
+_EVENT_DRAIN_TIMEOUT = 60
+
 
 class MockSubscriber:
     """Helper class to receive and verify published events"""
@@ -47,9 +58,10 @@ class MockSubscriber:
         poller = zmq.Poller()
         poller.register(self.sub, zmq.POLLIN)
 
-        timeout = 1000  # 1 second
-        while True:
-            events = dict(poller.poll(timeout))
+        poll_ms = _FIRST_EVENT_POLL_MS
+        deadline = time.monotonic() + _EVENT_DRAIN_TIMEOUT
+        while time.monotonic() < deadline:
+            events = dict(poller.poll(poll_ms))
 
             if events.get(self.sub) != zmq.POLLIN:
                 return cpu_stored_events
@@ -63,13 +75,32 @@ class MockSubscriber:
             for event in event_batch.events:
                 if isinstance(event, BlockStored) and event.medium == "CPU":
                     cpu_stored_events.append(event)
-                    timeout = 100
+                    poll_ms = 100
+
+        return cpu_stored_events
 
     def close(self):
         """Clean up resources"""
         self.sub.close()
 
 
+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
+    are still held by the offload worker, ``reset_prefix_cache`` returns
+    ``False``.  Retry with a short sleep until it succeeds or we time out.
+    """
+    deadline = time.monotonic() + _RESET_CACHE_TIMEOUT
+    while not llm.reset_prefix_cache():
+        if time.monotonic() > deadline:
+            raise TimeoutError(
+                "reset_prefix_cache did not succeed within "
+                f"{_RESET_CACHE_TIMEOUT}s - async offload may be stuck"
+            )
+        time.sleep(0.1)
+
+
 def _latency_test(llm: LLM, subscriber: MockSubscriber):
     sampling_params = SamplingParams(max_tokens=1)
 
@@ -95,10 +126,16 @@ def _latency_test(llm: LLM, subscriber: MockSubscriber):
         gpu_hit_time = time.time() - start_time
         total_gpu_hit_time += gpu_hit_time
 
-        # reset prefix cache to avoid GPU hit.
-        llm.reset_prefix_cache()
+        # Wait for the async CPU offload to finish, then reset prefix cache
+        # so the next generate() must reload from CPU rather than GPU.
+        _wait_for_prefix_cache_reset(llm)
 
-        assert subscriber.get_new_cpu_stored_events()
+        # Verify CPU stored events arrived (offload is done before we
+        # attempt to load from CPU).
+        assert subscriber.get_new_cpu_stored_events(), (
+            f"No CPU stored events received on iteration {i}; "
+            "async offload may not have completed in time"
+        )
 
         # run generation again - this should trigger loading from CPU
         start_time = time.time()
@@ -185,6 +222,8 @@ def test_cpu_offloading(cpu_block_size: int, attn_backend: str) -> None:
         kv_events_config=kv_events_config,
         kv_transfer_config=kv_transfer_config,
         attention_config={"backend": attn_backend},
+        # ROCm: batch size 1 to reduce variability
+        **({"max_num_seqs": 1} if current_platform.is_rocm() else {}),
     )
 
     events_endpoint = events_endpoint.replace("*", "127.0.0.1")

tools/pre_commit/check_forbidden_imports.py

@@ -59,6 +59,14 @@ CHECK_IMPORTS = {
             "vllm/v1/serial_utils.py",
         },
     ),
+    "base64": ForbiddenImport(
+        pattern=r"^\s*(?:import\s+base64(?:$|\s|,)|from\s+base64\s+import)",
+        tip=(
+            "Replace 'import base64' with 'import pybase64' "
+            "or 'import pybase64 as base64'."
+        ),
+        allowed_pattern=re.compile(r"^\s*import\s+pybase64(\s*|\s+as\s+base64\s*)$"),
+    ),
     "re": ForbiddenImport(
         pattern=r"^\s*(?:import\s+re(?:$|\s|,)|from\s+re\s+import)",
         tip="Replace 'import re' with 'import regex as re' or 'import regex'.",

tools/pre_commit/generate_attention_backend_docs.py

@@ -1262,14 +1262,23 @@ When no backend is specified (the default):
 """
 
 
-def _priority_table(title: str, backends: list[str]) -> list[str]:
+def _priority_table(
+    title: str,
+    backends: list[str],
+    annotations: dict[str, str] | None = None,
+) -> list[str]:
     """Generate a priority table for a list of backends."""
+
+    def _fmt(b: str) -> str:
+        suffix = annotations.get(b, "") if annotations else ""
+        return f"`{b}`{suffix}"
+
     return [
         f"**{title}:**",
         "",
         "| Priority | Backend |",
         "| -------- | ------- |",
-        *[f"| {i} | `{b}` |" for i, b in enumerate(backends, 1)],
+        *[f"| {i} | {_fmt(b)} |" for i, b in enumerate(backends, 1)],
         "",
     ]
 
@@ -1298,11 +1307,25 @@ def generate_priority_section(priorities: dict[str, list[str]]) -> str:
 
     lines.extend(["### MLA Attention (DeepSeek-style)", ""])
 
+    mla_sm100_annotations = {
+        "FLASHINFER_MLA_SPARSE": "**\\***",
+    }
     if "mla_sm100" in priorities:
-        lines.extend(_priority_table(sm100, priorities["mla_sm100"]))
+        lines.extend(
+            _priority_table(sm100, priorities["mla_sm100"], mla_sm100_annotations)
+        )
     if "mla_default" in priorities:
         lines.extend(_priority_table(ampere, priorities["mla_default"]))
 
+    if "mla_sm100" in priorities:
+        lines.append(
+            "> **\\*** For sparse MLA, FP8 KV cache always prefers "
+            "`FLASHINFER_MLA_SPARSE`. With BF16 KV cache, `FLASHINFER_MLA_SPARSE` "
+            "is preferred for low query-head counts (<= 16), while "
+            "`FLASHMLA_SPARSE` is preferred otherwise."
+        )
+        lines.append(">")
+
     lines.append(
         "> **Note:** ROCm and CPU platforms have their own selection logic. "
         "See the platform-specific documentation for details."

vllm/_custom_ops.py

@@ -989,6 +989,7 @@ def get_cutlass_moe_mm_data(
     n: int,
     k: int,
     blockscale_offsets: torch.Tensor | None = None,
+    is_gated: bool = True,
 ):
     """
     Prepare data necessary to perform CUTLASS grouped matrix multiplications
@@ -1012,6 +1013,8 @@ def get_cutlass_moe_mm_data(
                           its computation. The number of block scale rows
                           computed with expert E is blockscale_offsets[E + 1] -
                           blockscale_offsets[E]
+    - is_gated: Whether the activation is gated (gate + up). When True, the
+                first GEMM N dimension is 2*n; when False, it is n.
     """
     return torch.ops._C.get_cutlass_moe_mm_data(
         topk_ids,
@@ -1024,6 +1027,7 @@ def get_cutlass_moe_mm_data(
         n,
         k,
         blockscale_offsets,
+        is_gated,
     )
 
 

vllm/_xpu_ops.py

@@ -426,7 +426,8 @@ class xpu_ops:
         mask = positions <= index_end_pos
         # mask: [B * N, L]
         logits = logits.masked_fill(~mask, float("-inf"))
-        topk_indices = logits.topk(topk_tokens, dim=-1)[1].to(torch.int32)  # [B * N, K]
+        real_topk = min(topk_tokens, logits.shape[-1])
+        topk_indices = logits.topk(real_topk, dim=-1)[1].to(torch.int32)  # [B * N, K]
         # ensure we don't set indices for the top k
         # that is out of range(masked already)
         # this will happen if context length is shorter than K