[Misc] fix qwen3.5 config (#34604)

Signed-off-by: Seiji Eicher <seiji@anyscale.com>

.github/CODEOWNERS

@@ -2,7 +2,9 @@
 # for more info about CODEOWNERS file
 
 # This lists cover the "core" components of vLLM that require careful review
-/vllm/executor/executor_base.py @zhuohan123 @youkaichao @alexm-redhat @njhill @22quinn
+/vllm/compilation @zou3519 @youkaichao @ProExpertProg
+/vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery
+/vllm/lora @jeejeelee
 /vllm/model_executor/layers/attention @LucasWilkinson
 /vllm/model_executor/layers/fused_moe @mgoin @pavanimajety
 /vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety
@@ -11,18 +13,34 @@
 /vllm/model_executor/layers/batch_invariant.py @yewentao256 
 /vllm/multimodal @DarkLight1337 @ywang96 @NickLucche @tjtanaa
 /vllm/vllm_flash_attn @LucasWilkinson
-/vllm/lora @jeejeelee
-/vllm/reasoning @aarnphm @chaunceyjiang
-/vllm/entrypoints @aarnphm @chaunceyjiang
-/vllm/tool_parsers @aarnphm @chaunceyjiang
-/vllm/compilation @zou3519 @youkaichao @ProExpertProg
-/vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery
 CMakeLists.txt @tlrmchlsmth @LucasWilkinson
 
 # Any change to the VllmConfig changes can have a large user-facing impact,
 # so spam a lot of people
 /vllm/config @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @hmellor @yewentao256 @ProExpertProg
-/vllm/config/cache.py @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @hmellor @yewentao256 @ProExpertProg @heheda12345
+/vllm/config/cache.py @heheda12345
+
+# Entrypoints
+/vllm/entrypoints/anthropic @mgoin @DarkLight1337
+/vllm/entrypoints/cli @hmellor @mgoin @DarkLight1337 @russellb
+/vllm/entrypoints/mcp @heheda12345
+/vllm/entrypoints/openai @aarnphm @chaunceyjiang @DarkLight1337 @russellb
+/vllm/entrypoints/openai/realtime @njhill
+/vllm/entrypoints/openai/speech_to_text @NickLucche
+/vllm/entrypoints/pooling @noooop
+/vllm/entrypoints/sagemaker @DarkLight1337
+/vllm/entrypoints/serve @njhill
+/vllm/entrypoints/*.py @njhill
+/vllm/entrypoints/chat_utils.py @DarkLight1337
+/vllm/entrypoints/llm.py @DarkLight1337
+
+# Input/Output Processing
+/vllm/sampling_params.py @njhill @NickLucche
+/vllm/pooling_params.py @noooop @DarkLight1337
+/vllm/tokenizers @DarkLight1337 @njhill
+/vllm/renderers @DarkLight1337 @njhill
+/vllm/reasoning @aarnphm @chaunceyjiang
+/vllm/tool_parsers @aarnphm @chaunceyjiang
 
 # vLLM V1
 /vllm/v1/attention @LucasWilkinson
@@ -115,8 +133,8 @@ mkdocs.yaml @hmellor
 /vllm/model_executor/models/mixtral*.py @patrickvonplaten
 /vllm/model_executor/models/voxtral*.py @patrickvonplaten
 /vllm/model_executor/models/pixtral*.py @patrickvonplaten
+/vllm/tokenizers/mistral.py @patrickvonplaten
 /vllm/transformers_utils/configs/mistral.py @patrickvonplaten
-/vllm/transformers_utils/tokenizers/mistral.py @patrickvonplaten
 
 # Kernels
 /vllm/v1/attention/ops/chunked_prefill_paged_decode.py @tdoublep
@@ -152,9 +170,7 @@ mkdocs.yaml @hmellor
 /examples/pooling @noooop
 /tests/models/*/pooling* @noooop
 /tests/entrypoints/pooling @noooop
-/vllm/entrypoints/pooling @noooop
 /vllm/config/pooler.py @noooop
-/vllm/pooling_params.py @noooop
 /vllm/model_executor/layers/pooler @noooop
 
 # Security guide and policies

benchmarks/kernels/benchmark_moe.py

@@ -100,13 +100,38 @@ def benchmark_config(
     dtype: torch.dtype,
     use_fp8_w8a8: bool,
     use_int8_w8a16: bool,
+    use_int4_w4a16: bool = False,
     num_iters: int = 100,
     block_quant_shape: list[int] = None,
     use_deep_gemm: bool = False,
 ) -> float:
     init_dtype = torch.float16 if use_fp8_w8a8 else dtype
     x = torch.randn(num_tokens, hidden_size, dtype=dtype)
-    if use_int8_w8a16:
+    if use_int4_w4a16:
+        # Int4 packed weights: 2 int4 values per uint8 byte
+        # K dimension is packed (halved)
+        intermediate_size = shard_intermediate_size // 2  # after silu_and_mul
+        w1 = torch.randint(
+            0,
+            255,
+            (
+                num_experts,
+                shard_intermediate_size,
+                hidden_size // 2,  # int4 packing
+            ),
+            dtype=torch.uint8,
+        )
+        w2 = torch.randint(
+            0,
+            255,
+            (
+                num_experts,
+                hidden_size,
+                intermediate_size // 2,  # int4 packing
+            ),
+            dtype=torch.uint8,
+        )
+    elif use_int8_w8a16:
         w1 = torch.randint(
             -127,
             127,
@@ -140,7 +165,20 @@ def benchmark_config(
     w2_scale = None
     a1_scale = None
     a2_scale = None
-    if use_int8_w8a16:
+    if use_int4_w4a16:
+        if block_quant_shape is None:
+            raise ValueError("block_quant_shape is required for int4_w4a16")
+        group_size = block_quant_shape[1]
+        # Scales shape: (E, N, K // group_size) in fp16
+        w1_scale = torch.rand(
+            (num_experts, shard_intermediate_size, hidden_size // group_size),
+            dtype=dtype,
+        )
+        w2_scale = torch.rand(
+            (num_experts, hidden_size, intermediate_size // group_size),
+            dtype=dtype,
+        )
+    elif use_int8_w8a16:
         w1_scale = torch.randn(
             (num_experts, 2 * shard_intermediate_size), dtype=torch.float32
         )
@@ -199,6 +237,7 @@ def benchmark_config(
             a1_scale=a1_scale,
             a2_scale=a2_scale,
             block_shape=block_quant_shape,
+            weight_dtype="int4" if use_int4_w4a16 else None,
         )
 
         deep_gemm_experts = None
@@ -481,6 +520,7 @@ class BenchmarkWorker:
         dtype: torch.dtype,
         use_fp8_w8a8: bool,
         use_int8_w8a16: bool,
+        use_int4_w4a16: bool = False,
         block_quant_shape: list[int] = None,
         use_deep_gemm: bool = False,
     ) -> tuple[dict[str, int], float]:
@@ -488,7 +528,10 @@ class BenchmarkWorker:
 
         set_random_seed(self.seed)
         dtype_str = _get_config_dtype_str(
-            dtype, use_int8_w8a16=use_int8_w8a16, use_fp8_w8a8=use_fp8_w8a8
+            dtype,
+            use_int8_w8a16=use_int8_w8a16,
+            use_fp8_w8a8=use_fp8_w8a8,
+            use_int4_w4a16=use_int4_w4a16,
         )
         # NOTE(woosuk): The current naming convention uses w2.shape[2], which
         # is the intermediate size after silu_and_mul.
@@ -519,6 +562,7 @@ class BenchmarkWorker:
             dtype,
             use_fp8_w8a8,
             use_int8_w8a16,
+            use_int4_w4a16=use_int4_w4a16,
             num_iters=100,
             block_quant_shape=block_quant_shape,
             use_deep_gemm=use_deep_gemm,
@@ -535,6 +579,7 @@ class BenchmarkWorker:
         dtype: torch.dtype,
         use_fp8_w8a8: bool,
         use_int8_w8a16: bool,
+        use_int4_w4a16: bool,
         search_space: list[dict[str, int]],
         block_quant_shape: list[int],
         use_deep_gemm: bool,
@@ -545,7 +590,7 @@ class BenchmarkWorker:
         best_config = None
         best_time = float("inf")
         if current_platform.is_rocm():
-            is_fp16 = not (use_fp8_w8a8 or use_int8_w8a16)
+            is_fp16 = not (use_fp8_w8a8 or use_int8_w8a16 or use_int4_w4a16)
             search_space = prune_rocm_search_space(
                 num_tokens,
                 shard_intermediate_size,
@@ -574,6 +619,7 @@ class BenchmarkWorker:
                         dtype,
                         use_fp8_w8a8,
                         use_int8_w8a16,
+                        use_int4_w4a16,
                         num_iters=20,
                         block_quant_shape=block_quant_shape,
                         use_deep_gemm=use_deep_gemm,
@@ -621,6 +667,7 @@ def sort_config(config: BenchmarkConfig) -> BenchmarkConfig:
             else {}
         ),
         **({"kpack": config["kpack"]} if "kpack" in config else {}),
+        **({"SPLIT_K": config["SPLIT_K"]} if "SPLIT_K" in config else {}),
     }
 
 
@@ -633,11 +680,15 @@ def save_configs(
     dtype: torch.dtype,
     use_fp8_w8a8: bool,
     use_int8_w8a16: bool,
+    use_int4_w4a16: bool,
     block_quant_shape: list[int],
     save_dir: str,
 ) -> None:
     dtype_str = _get_config_dtype_str(
-        dtype, use_int8_w8a16=use_int8_w8a16, use_fp8_w8a8=use_fp8_w8a8
+        dtype,
+        use_int8_w8a16=use_int8_w8a16,
+        use_fp8_w8a8=use_fp8_w8a8,
+        use_int4_w4a16=use_int4_w4a16,
     )
 
     # NOTE(woosuk): The current naming convention uses w2.shape[2], which
@@ -739,6 +790,38 @@ def get_model_params(config):
     return E, topk, intermediate_size, hidden_size
 
 
+def get_quantization_group_size(config) -> int | None:
+    """Extract the quantization group size from the HF model config.
+
+    This reads directly from the HuggingFace config object (as returned by
+    ``get_config()``), not from vLLM's quantization config classes.
+
+    Supports AWQ/GPTQ-style configs (direct 'group_size' key) and
+    compressed-tensors configs (nested inside 'config_groups').
+    """
+    quantization_config = getattr(config, "quantization_config", {})
+    if not isinstance(quantization_config, dict):
+        return None
+    # AWQ / GPTQ style: group_size is a top-level key
+    gs = quantization_config.get("group_size")
+    if gs is not None:
+        return gs
+    # compressed-tensors style: group_size is nested in config_groups
+    config_groups = quantization_config.get("config_groups", {})
+    if not isinstance(config_groups, dict):
+        return None
+    for group_cfg in config_groups.values():
+        if not isinstance(group_cfg, dict):
+            continue
+        weights = group_cfg.get("weights", {})
+        if not isinstance(weights, dict):
+            continue
+        gs = weights.get("group_size")
+        if gs is not None:
+            return gs
+    return None
+
+
 def main(args: argparse.Namespace):
     print(args)
 
@@ -757,7 +840,20 @@ def main(args: argparse.Namespace):
     dtype = torch.float16 if current_platform.is_rocm() else config.dtype
     use_fp8_w8a8 = args.dtype == "fp8_w8a8"
     use_int8_w8a16 = args.dtype == "int8_w8a16"
+    use_int4_w4a16 = args.dtype == "int4_w4a16"
     block_quant_shape = get_weight_block_size_safety(config)
+    if use_int4_w4a16:
+        group_size = get_quantization_group_size(config)
+        if group_size is None:
+            raise ValueError(
+                "Could not determine group_size from model config. "
+                "The model's quantization_config must contain a 'group_size' "
+                "field (AWQ/GPTQ) or 'config_groups.*.weights.group_size' "
+                "(compressed-tensors)."
+            )
+        # For int4_w4a16, block_shape = [0, group_size]
+        # block_shape[0]=0 means no block quantization on N dimension
+        block_quant_shape = [0, group_size]
 
     if args.batch_size is None:
         batch_sizes = [
@@ -811,8 +907,20 @@ def main(args: argparse.Namespace):
         return ray.get(outputs)
 
     if args.tune:
-        is_fp16 = not (use_fp8_w8a8 or use_int8_w8a16)
-        search_space = get_configs_compute_bound(is_fp16, block_quant_shape)
+        # int4_w4a16 weights are uint8-packed, not fp16; treat like fp8 for
+        # search space generation (no matrix_instr_nonkdim/kpack exploration).
+        is_fp16 = not (use_fp8_w8a8 or use_int8_w8a16 or use_int4_w4a16)
+        # For int4_w4a16, the group_size constraint on BLOCK_SIZE_K does not
+        # apply: the gptq_awq kernel handles arbitrary BLOCK_SIZE_K regardless
+        # of group_size. Skip block_quant_shape filtering to keep the full
+        # search space (e.g. BLOCK_SIZE_K=64 with group_size=128).
+        tune_block_quant_shape = None if use_int4_w4a16 else block_quant_shape
+        search_space = get_configs_compute_bound(is_fp16, tune_block_quant_shape)
+        if use_int4_w4a16:
+            # SPLIT_K is a required kernel constexpr for gptq_awq kernel;
+            # only SPLIT_K=1 is used at runtime, so fix it during tuning.
+            for cfg in search_space:
+                cfg["SPLIT_K"] = 1
         print(f"Start tuning over {len(search_space)} configurations...")
         if use_deep_gemm:
             raise ValueError(
@@ -832,6 +940,7 @@ def main(args: argparse.Namespace):
                     dtype,
                     use_fp8_w8a8,
                     use_int8_w8a16,
+                    use_int4_w4a16,
                     search_space,
                     block_quant_shape,
                     use_deep_gemm,
@@ -851,6 +960,7 @@ def main(args: argparse.Namespace):
             dtype,
             use_fp8_w8a8,
             use_int8_w8a16,
+            use_int4_w4a16,
             block_quant_shape,
             args.save_dir,
         )
@@ -869,6 +979,7 @@ def main(args: argparse.Namespace):
                     dtype,
                     use_fp8_w8a8,
                     use_int8_w8a16,
+                    use_int4_w4a16,
                     block_quant_shape,
                     use_deep_gemm,
                 )
@@ -891,7 +1002,10 @@ if __name__ == "__main__":
     )
     parser.add_argument("--enable-expert-parallel", "-enable-ep", action="store_true")
     parser.add_argument(
-        "--dtype", type=str, choices=["auto", "fp8_w8a8", "int8_w8a16"], default="auto"
+        "--dtype",
+        type=str,
+        choices=["auto", "fp8_w8a8", "int8_w8a16", "int4_w4a16"],
+        default="auto",
     )
     parser.add_argument("--use-deep-gemm", action="store_true")
     parser.add_argument(

csrc/cuda_view.cu

@@ -2,33 +2,58 @@
 #include <torch/cuda.h>
 #include <cuda_runtime.h>
 
-// This function assumes that `cpu_tensor` is a CPU tensor allocated with pinned
-// memory, and that UVA (Unified Virtual Addressing) is enabled.
+// This function assumes that `cpu_tensor` is a CPU tensor,
+// and that UVA (Unified Virtual Addressing) is enabled.
 torch::Tensor get_cuda_view_from_cpu_tensor(torch::Tensor& cpu_tensor) {
   TORCH_CHECK(cpu_tensor.device().is_cpu(), "Input tensor must be on CPU");
 
-  // Get raw host pointer from CPU tensor
-  void* host_ptr = cpu_tensor.data_ptr();
+  // handle empty tensor
+  if (cpu_tensor.numel() == 0) {
+    return torch::empty(cpu_tensor.sizes(),
+                        cpu_tensor.options().device(torch::kCUDA));
+  }
+
+  if (cpu_tensor.is_pinned()) {
+    // If CPU tensor is pinned, directly get the device pointer.
+    void* host_ptr = const_cast<void*>(cpu_tensor.data_ptr());
+    void* device_ptr = nullptr;
+    cudaError_t err = cudaHostGetDevicePointer(&device_ptr, host_ptr, 0);
+    TORCH_CHECK(err == cudaSuccess,
+                "cudaHostGetDevicePointer failed: ", cudaGetErrorString(err));
+
+    return torch::from_blob(
+        device_ptr, cpu_tensor.sizes(), cpu_tensor.strides(),
+        [base = cpu_tensor](void*) {},  // keep cpu tensor alive
+        cpu_tensor.options().device(torch::kCUDA));
+  }
+
+  // If CPU tensor is not pinned, allocate a new pinned memory buffer.
+  torch::Tensor contiguous_cpu = cpu_tensor.contiguous();
+  size_t nbytes = contiguous_cpu.nbytes();
+
+  void* host_ptr = nullptr;
+  cudaError_t err = cudaHostAlloc(&host_ptr, nbytes, cudaHostAllocMapped);
+  if (err != cudaSuccess) {
+    AT_ERROR("cudaHostAlloc failed: ", cudaGetErrorString(err));
+  }
+
+  err = cudaMemcpy(host_ptr, contiguous_cpu.data_ptr(), nbytes,
+                   cudaMemcpyDefault);
+  if (err != cudaSuccess) {
+    cudaFreeHost(host_ptr);
+    AT_ERROR("cudaMemcpy failed: ", cudaGetErrorString(err));
+  }
 
-  // Get a device pointer corresponding to the pinned host memory
   void* device_ptr = nullptr;
-  cudaError_t err = cudaHostGetDevicePointer(&device_ptr, host_ptr, 0);
-  TORCH_CHECK(err == cudaSuccess,
-              "cudaHostGetDevicePointer failed: ", cudaGetErrorString(err));
+  err = cudaHostGetDevicePointer(&device_ptr, host_ptr, 0);
+  if (err != cudaSuccess) {
+    cudaFreeHost(host_ptr);
+    AT_ERROR("cudaHostGetDevicePointer failed: ", cudaGetErrorString(err));
+  }
 
-  // We'll use the same sizes, strides, and dtype as the CPU tensor.
-  // TODO: check if layout is respected.
-  auto sizes = cpu_tensor.sizes();
-  auto strides = cpu_tensor.strides();
-  auto options = cpu_tensor.options().device(torch::kCUDA);
+  auto deleter = [host_ptr](void*) { cudaFreeHost(host_ptr); };
 
-  // use default no-op deleter, since the memory is owned by the original CPU
-  // tensor
-  torch::Tensor cuda_tensor =
-      torch::from_blob(device_ptr, sizes, strides, options);
-
-  TORCH_CHECK(cuda_tensor.device().is_cuda(),
-              "Resulting tensor is not on CUDA device");
-
-  return cuda_tensor;
-}
+  return torch::from_blob(device_ptr, contiguous_cpu.sizes(),
+                          contiguous_cpu.strides(), deleter,
+                          contiguous_cpu.options().device(torch::kCUDA));
+}

csrc/rocm/skinny_gemms.cu

@@ -1568,8 +1568,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
         {
   #endif
           unsigned int kOff = k + (thrd * A_CHUNK);
-          unsigned int kOffcp =
-              k_str + kOff;  // min__(K - A_CHUNK, k_str + kOff);
+          unsigned int kOffcp = min__(K - A_CHUNK, k_str + kOff);
           for (unsigned int n = 0; n < N; n += CHUNKK * sprdN) {
             __builtin_amdgcn_global_load_lds(
                 (int*)(&A[min__(

csrc/topk.cu

@@ -349,7 +349,7 @@ void setup_kernel_smem_once() {
 void large_context_topk(
     const torch::Tensor& logits, torch::Tensor& indices,
     const torch::Tensor& seq_lens,
-    c10::optional<torch::Tensor> row_starts = c10::nullopt) {
+    std::optional<torch::Tensor> row_starts = std::nullopt) {
   TORCH_CHECK(logits.is_cuda(), "logits must be a CUDA tensor");
   TORCH_CHECK(indices.is_cuda(), "indices must be a CUDA tensor");
   TORCH_CHECK(seq_lens.is_cuda(), "seq_lens must be a CUDA tensor");

docs/models/pooling_models.md

@@ -311,20 +311,31 @@ An OpenAI client example can be found here: [examples/pooling/embed/openai_embed
 
 [ColBERT](https://arxiv.org/abs/2004.12832) (Contextualized Late Interaction over BERT) is a retrieval model that uses per-token embeddings and MaxSim scoring for document ranking. Unlike single-vector embedding models, ColBERT retains token-level representations and computes relevance scores through late interaction, providing better accuracy while being more efficient than cross-encoders.
 
-vLLM supports ColBERT models for reranking tasks, automatically applying MaxSim scoring for query-document relevance:
+vLLM supports ColBERT models with multiple encoder backbones:
+
+| Architecture | Backbone | Example HF Models |
+|---|---|---|
+| `HF_ColBERT` | BERT | `answerdotai/answerai-colbert-small-v1`, `colbert-ir/colbertv2.0` |
+| `ColBERTModernBertModel` | ModernBERT | `lightonai/GTE-ModernColBERT-v1` |
+| `ColBERTJinaRobertaModel` | Jina XLM-RoBERTa | `jinaai/jina-colbert-v2` |
+
+**BERT-based ColBERT** models work out of the box:
 
 ```shell
 vllm serve answerdotai/answerai-colbert-small-v1
 ```
 
-Currently supports ColBERT models with standard BERT encoders (e.g., `answerdotai/answerai-colbert-small-v1`, `colbert-ir/colbertv2.0`).
-
-ColBERT models with modified encoder architectures are not yet supported, including BERT variants with rotary embeddings (e.g., `jinaai/jina-colbert-v2`) or other custom encoders (e.g., `LiquidAI/LFM2-ColBERT-350M`).
-
-If your standard BERT ColBERT model's config doesn't specify the architecture as `HF_ColBERT`, override it with:
+For **non-BERT backbones**, use `--hf-overrides` to set the correct architecture:
 
 ```shell
-vllm serve your-colbert-model --hf-overrides '{"architectures": ["HF_ColBERT"]}'
+# ModernBERT backbone
+vllm serve lightonai/GTE-ModernColBERT-v1 \
+    --hf-overrides '{"architectures": ["ColBERTModernBertModel"]}'
+
+# Jina XLM-RoBERTa backbone
+vllm serve jinaai/jina-colbert-v2 \
+    --hf-overrides '{"architectures": ["ColBERTJinaRobertaModel"]}' \
+    --trust-remote-code
 ```
 
 Then you can use the rerank endpoint:
@@ -363,6 +374,77 @@ curl -s http://localhost:8000/pooling -H "Content-Type: application/json" -d '{
 
 An example can be found here: [examples/pooling/score/colbert_rerank_online.py](../../examples/pooling/score/colbert_rerank_online.py)
 
+### ColQwen3 Multi-Modal Late Interaction Models
+
+ColQwen3 is based on [ColPali](https://arxiv.org/abs/2407.01449), which extends ColBERT's late interaction approach to **multi-modal** inputs. While ColBERT operates on text-only token embeddings, ColPali/ColQwen3 can embed both **text and images** (e.g. PDF pages, screenshots, diagrams) into per-token L2-normalized vectors and compute relevance via MaxSim scoring. ColQwen3 specifically uses Qwen3-VL as its vision-language backbone.
+
+| Architecture | Backbone | Example HF Models |
+|---|---|---|
+| `ColQwen3` | Qwen3-VL | `TomoroAI/tomoro-colqwen3-embed-4b`, `TomoroAI/tomoro-colqwen3-embed-8b` |
+| `OpsColQwen3Model` | Qwen3-VL | `OpenSearch-AI/Ops-Colqwen3-4B`, `OpenSearch-AI/Ops-Colqwen3-8B` |
+
+Start the server:
+
+```shell
+vllm serve TomoroAI/tomoro-colqwen3-embed-4b --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": "TomoroAI/tomoro-colqwen3-embed-4b",
+    "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": "TomoroAI/tomoro-colqwen3-embed-4b",
+    "text_1": "What is the capital of France?",
+    "text_2": ["The capital of France is Paris.", "Python is a programming language."]
+}'
+```
+
+You can also get the raw token embeddings using the pooling endpoint with `token_embed` task:
+
+```shell
+curl -s http://localhost:8000/pooling -H "Content-Type: application/json" -d '{
+    "model": "TomoroAI/tomoro-colqwen3-embed-4b",
+    "input": "What is machine learning?",
+    "task": "token_embed"
+}'
+```
+
+For **image inputs**, use the chat-style `messages` field so that the vLLM multimodal processor handles them correctly:
+
+```shell
+curl -s http://localhost:8000/pooling -H "Content-Type: application/json" -d '{
+    "model": "TomoroAI/tomoro-colqwen3-embed-4b",
+    "messages": [
+        {
+            "role": "user",
+            "content": [
+                {"type": "image_url", "image_url": {"url": "data:image/png;base64,<BASE64>"}},
+                {"type": "text", "text": "Describe the image."}
+            ]
+        }
+    ]
+}'
+```
+
+Examples can be found here:
+
+- Multi-vector retrieval: [examples/pooling/token_embed/colqwen3_token_embed_online.py](../../examples/pooling/token_embed/colqwen3_token_embed_online.py)
+- Reranking: [examples/pooling/score/colqwen3_rerank_online.py](../../examples/pooling/score/colqwen3_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

@@ -728,6 +728,8 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `OpenPanguVLForConditionalGeneration` | openpangu-VL | T + I<sup>E+</sup> + V<sup>E+</sup> |`FreedomIntelligence/openPangu-VL-7B` | ✅︎ | ✅︎ |
 | `Ovis` | Ovis2, Ovis1.6 | T + I<sup>+</sup> | `AIDC-AI/Ovis2-1B`, `AIDC-AI/Ovis1.6-Llama3.2-3B`, etc. | | ✅︎ |
 | `Ovis2_5` | Ovis2.5 | T + I<sup>+</sup> + V | `AIDC-AI/Ovis2.5-9B`, etc. | | |
+| `Ovis2_6ForCausalLM` | Ovis2.6 | T + I<sup>+</sup> + V | `AIDC-AI/Ovis2.6-2B`, etc. | | |
+| `Ovis2_6_MoeForCausalLM` | Ovis2.6 | T + I<sup>+</sup> + V | `AIDC-AI/Ovis2.6-30B-A3B`, etc. | | |
 | `PaddleOCRVLForConditionalGeneration` | Paddle-OCR | T + I<sup>+</sup> | `PaddlePaddle/PaddleOCR-VL`, etc. | | |
 | `PaliGemmaForConditionalGeneration` | PaliGemma, PaliGemma 2 | T + I<sup>E</sup> | `google/paligemma-3b-pt-224`, `google/paligemma-3b-mix-224`, `google/paligemma2-3b-ft-docci-448`, etc. | ✅︎ | ✅︎ |
 | `Phi3VForCausalLM` | Phi-3-Vision, Phi-3.5-Vision | T + I<sup>E+</sup> | `microsoft/Phi-3-vision-128k-instruct`, `microsoft/Phi-3.5-vision-instruct`, etc. | | ✅︎ |

examples/online_serving/data_parallel_pause_resume.py

@@ -0,0 +1,135 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Test pause/resume with Data Parallel (DP) via HTTP API.
+
+This example demonstrates coordinated pause/resume across multiple DP ranks.
+The pause synchronizes across all DP engines via all-reduce.
+
+Prerequisites:
+    Start a vLLM server with data parallelism:
+
+    $ VLLM_SERVER_DEV_MODE=1 vllm serve facebook/opt-125m \
+        --enforce-eager \
+        --data-parallel-size 4 \
+        --tensor-parallel-size 1
+
+    Then run this script:
+
+    $ python data_parallel_pause_resume.py
+
+The test verifies pause works by:
+1. Starting a streaming generation request
+2. Pausing the server mid-generation
+3. Sleeping for PAUSE_DURATION seconds
+4. Resuming the server
+5. Verifying there was a gap in token generation matching the pause duration
+"""
+
+import argparse
+import threading
+import time
+
+import requests
+from openai import OpenAI
+
+BASE_URL = "http://localhost:8000"
+MODEL_NAME = "facebook/opt-125m"
+PAUSE_DURATION = 3.0
+
+
+def pause_generation(base_url: str, mode: str = "keep") -> None:
+    """Pause generation via HTTP endpoint."""
+    url = f"{base_url}/pause"
+    response = requests.post(url, params={"mode": mode}, timeout=60)
+    response.raise_for_status()
+    print("Server paused")
+
+
+def resume_generation(base_url: str) -> None:
+    """Resume generation via HTTP endpoint."""
+    url = f"{base_url}/resume"
+    response = requests.post(url, timeout=60)
+    response.raise_for_status()
+    print("Server resumed")
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--base-url", default=BASE_URL)
+    parser.add_argument("--model", default=MODEL_NAME)
+    args = parser.parse_args()
+
+    client = OpenAI(
+        base_url=f"{args.base_url}/v1",
+        api_key="EMPTY",
+    )
+
+    prompt = "Write a long story about a dragon. Once upon a time"
+    token_times: list[float] = []
+    pause_token_idx = 0
+    pause_triggered = threading.Event()
+
+    def generator_thread():
+        """Stream tokens and record timestamps."""
+        stream = client.completions.create(
+            model=args.model,
+            prompt=prompt,
+            max_tokens=50,
+            stream=True,
+        )
+        for chunk in stream:
+            if chunk.choices[0].text:
+                token_times.append(time.monotonic())
+                token_count = len(token_times)
+                print(f"Token {token_count}: {chunk.choices[0].text!r}")
+
+                # Signal controller after some tokens
+                if token_count >= 5 and not pause_triggered.is_set():
+                    pause_triggered.set()
+
+    def controller_thread():
+        """Pause and resume the server."""
+        nonlocal pause_token_idx
+
+        # Wait for some tokens
+        pause_triggered.wait()
+
+        print(f"\nPausing server (keep mode) at token {len(token_times)}...")
+        pause_generation(args.base_url, mode="keep")
+        pause_token_idx = len(token_times)
+        print(f"Sleeping for {PAUSE_DURATION}s...")
+
+        time.sleep(PAUSE_DURATION)
+
+        print("Resuming server...")
+        resume_generation(args.base_url)
+        print("Resumed!\n")
+
+    # Run both threads
+    gen_thread = threading.Thread(target=generator_thread)
+    ctrl_thread = threading.Thread(target=controller_thread)
+
+    gen_thread.start()
+    ctrl_thread.start()
+
+    gen_thread.join()
+    ctrl_thread.join()
+
+    # Check gap at the pause point
+    if pause_token_idx < len(token_times):
+        pause_gap = token_times[pause_token_idx] - token_times[pause_token_idx - 1]
+        print(
+            f"\nGap after pause (token {pause_token_idx} -> "
+            f"{pause_token_idx + 1}): {pause_gap:.3f}s"
+        )
+        if pause_gap >= PAUSE_DURATION * 0.9:
+            print("Test passed! Pause synchronized across DP ranks.")
+        else:
+            print(f"Test failed! Expected ~{PAUSE_DURATION}s gap, got {pause_gap:.3f}s")
+    else:
+        print("Test failed! No tokens were generated after resuming.")
+
+
+if __name__ == "__main__":
+    main()

examples/pooling/plugin/prithvi_geospatial_mae_io_processor.py

@@ -28,7 +28,7 @@ def main():
     )
 
     llm = LLM(
-        model="christian-pinto/Prithvi-EO-2.0-300M-TL-VLLM",
+        model="ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11",
         skip_tokenizer_init=True,
         trust_remote_code=True,
         enforce_eager=True,

examples/pooling/plugin/prithvi_geospatial_mae_offline.py

@@ -391,7 +391,7 @@ if __name__ == "__main__":
     parser.add_argument(
         "--model",
         type=str,
-        default="christian-pinto/Prithvi-EO-2.0-300M-TL-VLLM",
+        default="ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11",
         help="Path to a checkpoint file to load from.",
     )
     parser.add_argument(

examples/pooling/plugin/prithvi_geospatial_mae_online.py

@@ -14,9 +14,7 @@ import requests
 # - install TerraTorch v1.1 (or later):
 #   pip install terratorch>=v1.1
 # - start vllm in serving mode with the below args
-#   --model='christian-pinto/Prithvi-EO-2.0-300M-TL-VLLM'
-#   --model-impl terratorch
-#   --trust-remote-code
+#   --model='ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11'
 #   --skip-tokenizer-init --enforce-eager
 #   --io-processor-plugin terratorch_segmentation
 #   --enable-mm-embeds
@@ -34,7 +32,7 @@ def main():
             "out_data_format": "b64_json",
         },
         "priority": 0,
-        "model": "christian-pinto/Prithvi-EO-2.0-300M-TL-VLLM",
+        "model": "ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11",
     }
 
     ret = requests.post(server_endpoint, json=request_payload_url)

examples/pooling/score/colbert_rerank_online.py

@@ -1,15 +1,27 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 """
-Example of using ColBERT late interaction model for reranking.
+Example of using ColBERT late interaction models for reranking and scoring.
 
 ColBERT (Contextualized Late Interaction over BERT) uses per-token embeddings
 and MaxSim scoring for document reranking, providing better accuracy than
 single-vector models while being more efficient than cross-encoders.
 
-Start the server with:
+vLLM supports ColBERT with multiple encoder backbones. Start the server
+with one of the following:
+
+    # BERT backbone (works out of the box)
     vllm serve answerdotai/answerai-colbert-small-v1
 
+    # ModernBERT backbone
+    vllm serve lightonai/GTE-ModernColBERT-v1 \
+        --hf-overrides '{"architectures": ["ColBERTModernBertModel"]}'
+
+    # Jina XLM-RoBERTa backbone
+    vllm serve jinaai/jina-colbert-v2 \
+        --hf-overrides '{"architectures": ["ColBERTJinaRobertaModel"]}' \
+        --trust-remote-code
+
 Then run this script:
     python colbert_rerank_online.py
 """
@@ -18,39 +30,62 @@ import json
 
 import requests
 
-url = "http://127.0.0.1:8000/rerank"
+# Change this to match the model you started the server with
+MODEL = "answerdotai/answerai-colbert-small-v1"
+BASE_URL = "http://127.0.0.1:8000"
 
 headers = {"accept": "application/json", "Content-Type": "application/json"}
 
-data = {
-    "model": "answerdotai/answerai-colbert-small-v1",
-    "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.",
-    ],
-}
+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.",
+]
+
+
+def rerank_example():
+    """Use the /rerank endpoint to rank documents by query relevance."""
+    print("=== Rerank Example ===")
+
+    data = {
+        "model": MODEL,
+        "query": "What is machine learning?",
+        "documents": documents,
+    }
+
+    response = requests.post(f"{BASE_URL}/rerank", headers=headers, json=data)
+    result = response.json()
+    print(json.dumps(result, indent=2))
+
+    print("\nRanked documents (most relevant first):")
+    for item in result["results"]:
+        doc_idx = item["index"]
+        score = item["relevance_score"]
+        print(f"  Score {score:.4f}: {documents[doc_idx]}")
+
+
+def score_example():
+    """Use the /score endpoint for pairwise query-document scoring."""
+    print("\n=== Score Example ===")
+
+    data = {
+        "model": MODEL,
+        "text_1": "What is machine learning?",
+        "text_2": [
+            "Machine learning is a subset of AI.",
+            "The weather is sunny.",
+        ],
+    }
+
+    response = requests.post(f"{BASE_URL}/score", headers=headers, json=data)
+    result = response.json()
+    print(json.dumps(result, indent=2))
 
 
 def main():
-    response = requests.post(url, headers=headers, json=data)
-
-    if response.status_code == 200:
-        print("ColBERT Rerank Request successful!")
-        result = response.json()
-        print(json.dumps(result, indent=2))
-
-        # Show ranked results
-        print("\nRanked documents (most relevant first):")
-        for item in result["results"]:
-            doc_idx = item["index"]
-            score = item["relevance_score"]
-            print(f"  Score {score:.4f}: {data['documents'][doc_idx]}")
-    else:
-        print(f"Request failed with status code: {response.status_code}")
-        print(response.text)
+    rerank_example()
+    score_example()
 
 
 if __name__ == "__main__":

examples/pooling/score/colqwen3_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 late interaction model for reranking.
+
+ColQwen3 is a multi-modal ColBERT-style model based on Qwen3-VL.
+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 TomoroAI/tomoro-colqwen3-embed-4b --max-model-len 50000
+
+Then run this script:
+    python colqwen3_rerank_online.py
+"""
+
+import requests
+
+MODEL = "TomoroAI/tomoro-colqwen3-embed-4b"
+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/token_embed/colqwen3_token_embed_online.py

@@ -0,0 +1,198 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# ruff: noqa: E501
+
+"""
+Example online usage of Pooling API for ColQwen3 multi-vector retrieval.
+
+ColQwen3 is a multi-modal late interaction model based on Qwen3-VL that
+produces per-token embeddings (320-dim, L2-normalized) for both text and
+image inputs. Similarity is computed via MaxSim scoring.
+
+This example mirrors the official TomoroAI inference code
+(https://huggingface.co/TomoroAI/tomoro-colqwen3-embed-4b) but uses the
+vLLM serving API instead of local HuggingFace model loading.
+
+Start the server with:
+    vllm serve TomoroAI/tomoro-colqwen3-embed-4b --max-model-len 4096
+
+Then run this script:
+    python colqwen3_token_embed_online.py
+"""
+
+import argparse
+import base64
+from io import BytesIO
+
+import numpy as np
+import requests
+from PIL import Image
+
+# ── Helpers ─────────────────────────────────────────────────
+
+
+def post_http_request(payload: dict, api_url: str) -> requests.Response:
+    headers = {"User-Agent": "Test Client"}
+    return requests.post(api_url, headers=headers, json=payload)
+
+
+def load_image(url: str) -> Image.Image:
+    """Download an image from URL (handles Wikimedia 403)."""
+    for hdrs in ({}, {"User-Agent": "Mozilla/5.0 (compatible; ColQwen3-demo/1.0)"}):
+        resp = requests.get(url, headers=hdrs, timeout=10)
+        if resp.status_code == 403:
+            continue
+        resp.raise_for_status()
+        return Image.open(BytesIO(resp.content)).convert("RGB")
+    raise RuntimeError(f"Could not fetch image from {url}")
+
+
+def encode_image_base64(image: Image.Image) -> str:
+    """Encode a PIL image to a base64 data URI."""
+    buf = BytesIO()
+    image.save(buf, format="PNG")
+    return "data:image/png;base64," + base64.b64encode(buf.getvalue()).decode()
+
+
+def compute_maxsim(q_emb: np.ndarray, d_emb: np.ndarray) -> float:
+    """Compute ColBERT-style MaxSim score between query and document."""
+    sim = q_emb @ d_emb.T
+    return float(sim.max(axis=-1).sum())
+
+
+# ── Encode functions ────────────────────────────────────────
+
+
+def encode_queries(texts: list[str], model: str, api_url: str) -> list[np.ndarray]:
+    """Encode text queries → list of multi-vector embeddings."""
+    resp = post_http_request({"model": model, "input": texts}, api_url)
+    return [np.array(item["data"]) for item in resp.json()["data"]]
+
+
+def encode_images(image_urls: list[str], model: str, api_url: str) -> list[np.ndarray]:
+    """Encode image documents → list of multi-vector embeddings.
+
+    Images are sent via the chat-style `messages` field so that the
+    vLLM multimodal processor handles them correctly.
+    """
+    embeddings = []
+    for url in image_urls:
+        print(f"  Loading: {url.split('/')[-1]}...")
+        image = load_image(url)
+        image_uri = encode_image_base64(image)
+        resp = post_http_request(
+            {
+                "model": model,
+                "messages": [
+                    {
+                        "role": "user",
+                        "content": [
+                            {"type": "image_url", "image_url": {"url": image_uri}},
+                            {"type": "text", "text": "Describe the image."},
+                        ],
+                    }
+                ],
+            },
+            api_url,
+        )
+        result = resp.json()
+        if resp.status_code != 200 or "data" not in result:
+            print(f"    Error ({resp.status_code}): {str(result)[:200]}")
+            continue
+        embeddings.append(np.array(result["data"][0]["data"]))
+    return embeddings
+
+
+# ── Main ────────────────────────────────────────────────────
+
+
+def parse_args():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--host", type=str, default="localhost")
+    parser.add_argument("--port", type=int, default=8000)
+    parser.add_argument(
+        "--model",
+        type=str,
+        default="TomoroAI/tomoro-colqwen3-embed-4b",
+    )
+    return parser.parse_args()
+
+
+def main(args):
+    pooling_url = f"http://{args.host}:{args.port}/pooling"
+    score_url = f"http://{args.host}:{args.port}/score"
+    model = args.model
+
+    # Same sample data as the official TomoroAI example
+    queries = [
+        "Retrieve the city of Singapore",
+        "Retrieve the city of Beijing",
+        "Retrieve the city of London",
+    ]
+    image_urls = [
+        "https://upload.wikimedia.org/wikipedia/commons/2/27/Singapore_skyline_2022.jpg",
+        "https://upload.wikimedia.org/wikipedia/commons/6/61/Beijing_skyline_at_night.JPG",
+        "https://upload.wikimedia.org/wikipedia/commons/4/49/London_skyline.jpg",
+    ]
+
+    # ── 1) Text query embeddings ────────────────────────────
+    print("=" * 60)
+    print("1. Encode text queries (multi-vector)")
+    print("=" * 60)
+    query_embeddings = encode_queries(queries, model, pooling_url)
+    for i, emb in enumerate(query_embeddings):
+        norm = float(np.linalg.norm(emb[0]))
+        print(f'  Query {i}: {emb.shape}  (L2 norm: {norm:.4f})  "{queries[i]}"')
+
+    # ── 2) Image document embeddings ────────────────────────
+    print()
+    print("=" * 60)
+    print("2. Encode image documents (multi-vector)")
+    print("=" * 60)
+    doc_embeddings = encode_images(image_urls, model, pooling_url)
+    for i, emb in enumerate(doc_embeddings):
+        print(f"  Doc {i}:   {emb.shape}  {image_urls[i].split('/')[-1]}")
+
+    # ── 3) Cross-modal MaxSim scoring ───────────────────────
+    if doc_embeddings:
+        print()
+        print("=" * 60)
+        print("3. Cross-modal MaxSim scores (text queries × image docs)")
+        print("=" * 60)
+        # Header
+        print(f"{'':>35s}", end="")
+        for j in range(len(doc_embeddings)):
+            print(f"  Doc {j:>2d}", end="")
+        print()
+        # Score matrix
+        for i, q_emb in enumerate(query_embeddings):
+            print(f"  {queries[i]:<33s}", end="")
+            for j, d_emb in enumerate(doc_embeddings):
+                score = compute_maxsim(q_emb, d_emb)
+                print(f"  {score:6.2f}", end="")
+            print()
+
+    # ── 4) Text-only /score endpoint ────────────────────────
+    print()
+    print("=" * 60)
+    print("4. Text-only late interaction scoring (/score endpoint)")
+    print("=" * 60)
+    text_query = "What is the capital of France?"
+    text_docs = [
+        "The capital of France is Paris.",
+        "Berlin is the capital of Germany.",
+        "Python is a programming language.",
+    ]
+    resp = post_http_request(
+        {"model": model, "text_1": text_query, "text_2": text_docs},
+        score_url,
+    )
+    print(f'  Query: "{text_query}"\n')
+    for item in resp.json()["data"]:
+        idx = item["index"]
+        print(f"  Doc {idx} (score={item['score']:.4f}): {text_docs[idx]}")
+
+
+if __name__ == "__main__":
+    args = parse_args()
+    main(args)

requirements/rocm-test.txt

@@ -96,3 +96,5 @@ albumentations==1.4.6
 transformers==4.57.3
 # Pin HF Hub version
 huggingface-hub==0.36.2
+# Pin Mistral Common
+mistral-common[image,audio]==1.9.1

requirements/test.in

@@ -56,7 +56,14 @@ runai-model-streamer[s3,gcs]==0.15.3
 fastsafetensors>=0.2.2 # 0.2.2 contains important fixes for multi-GPU mem usage
 pydantic>=2.12 # 2.11 leads to error on python 3.13
 decord==0.6.0
-terratorch @ git+https://github.com/IBM/terratorch.git@1.1.rc3 # required for PrithviMAE test
+terratorch >= 1.2.2 # Required for Prithvi tests
+imagehash # Required for Prithvi tests
+segmentation-models-pytorch > 0.4.0 # Required for Prithvi tests
+
 gpt-oss >= 0.0.7; python_version > '3.11'
 
 perceptron # required for isaac test
+
+# Newer versions of datasets require torchcoded, that makes the tests fail in CI because of a missing library.
+# Older versions are in conflict with teerratorch requirements.
+datasets>=3.3.0,<=3.6.0

requirements/test.txt

@@ -1,7 +1,9 @@
 # This file was autogenerated by uv via the following command:
 #    uv pip compile requirements/test.in -o requirements/test.txt --index-strategy unsafe-best-match --torch-backend cu129 --python-platform x86_64-manylinux_2_28 --python-version 3.12
 absl-py==2.1.0
-    # via rouge-score
+    # via
+    #   rouge-score
+    #   tensorboard
 accelerate==1.0.1
     # via
     #   lm-eval
@@ -31,9 +33,7 @@ albumentations==1.4.6
     #   -r requirements/test.in
     #   terratorch
 alembic==1.16.4
-    # via
-    #   mlflow
-    #   optuna
+    # via optuna
 annotated-doc==0.0.4
     # via fastapi
 annotated-types==0.7.0
@@ -74,8 +74,6 @@ bitsandbytes==0.46.1
     #   lightning
 black==24.10.0
     # via datamodel-code-generator
-blinker==1.9.0
-    # via flask
 blobfile==3.0.0
     # via -r requirements/test.in
 bm25s==0.2.13
@@ -93,9 +91,7 @@ bounded-pool-executor==0.0.3
 buildkite-test-collector==0.1.9
     # via -r requirements/test.in
 cachetools==5.5.2
-    # via
-    #   google-auth
-    #   mlflow-skinny
+    # via google-auth
 certifi==2024.8.30
     # via
     #   fiona
@@ -106,6 +102,7 @@ certifi==2024.8.30
     #   pyproj
     #   rasterio
     #   requests
+    #   sentry-sdk
 cffi==1.17.1
     # via soundfile
 chardet==5.2.0
@@ -120,15 +117,14 @@ click==8.1.7
     #   click-plugins
     #   cligj
     #   fiona
-    #   flask
     #   jiwer
-    #   mlflow-skinny
     #   nltk
     #   rasterio
     #   ray
     #   schemathesis
     #   typer
     #   uvicorn
+    #   wandb
 click-plugins==1.1.1.2
     # via
     #   fiona
@@ -137,8 +133,6 @@ cligj==0.7.2
     # via
     #   fiona
     #   rasterio
-cloudpickle==3.1.1
-    # via mlflow-skinny
 colorama==0.4.6
     # via
     #   perceptron
@@ -163,16 +157,15 @@ cupy-cuda12x==13.6.0
     # via ray
 cycler==0.12.1
     # via matplotlib
-databricks-sdk==0.59.0
-    # via mlflow-skinny
 datamodel-code-generator==0.26.3
     # via -r requirements/test.in
 dataproperty==1.0.1
     # via
     #   pytablewriter
     #   tabledata
-datasets==3.0.2
+datasets==3.3.0
     # via
+    #   -r requirements/test.in
     #   evaluate
     #   lm-eval
     #   mteb
@@ -180,6 +173,8 @@ decorator==5.1.1
     # via librosa
 decord==0.6.0
     # via -r requirements/test.in
+diffusers==0.36.0
+    # via terratorch
 dill==0.3.8
     # via
     #   datasets
@@ -191,15 +186,11 @@ distlib==0.3.9
 dnspython==2.7.0
     # via email-validator
 docker==7.1.0
-    # via
-    #   gpt-oss
-    #   mlflow
+    # via gpt-oss
 docopt==0.6.2
     # via num2words
 docstring-parser==0.17.0
     # via jsonargparse
-efficientnet-pytorch==0.7.1
-    # via segmentation-models-pytorch
 einops==0.8.1
     # via
     #   -r requirements/test.in
@@ -217,9 +208,7 @@ encodec==0.1.1
 evaluate==0.4.3
     # via lm-eval
 fastapi==0.128.0
-    # via
-    #   gpt-oss
-    #   mlflow-skinny
+    # via gpt-oss
 fastparquet==2024.11.0
     # via genai-perf
 fastrlock==0.8.2
@@ -230,6 +219,7 @@ filelock==3.16.1
     # via
     #   blobfile
     #   datasets
+    #   diffusers
     #   huggingface-hub
     #   ray
     #   torch
@@ -237,8 +227,6 @@ filelock==3.16.1
     #   virtualenv
 fiona==1.10.1
     # via torchgeo
-flask==3.1.1
-    # via mlflow
 fonttools==4.55.0
     # via matplotlib
 fqdn==1.5.1
@@ -249,7 +237,7 @@ frozenlist==1.5.0
     # via
     #   aiohttp
     #   aiosignal
-fsspec==2024.9.0
+fsspec==2024.12.0
     # via
     #   datasets
     #   evaluate
@@ -257,6 +245,7 @@ fsspec==2024.9.0
     #   huggingface-hub
     #   lightning
     #   pytorch-lightning
+    #   tacoreader
     #   torch
 ftfy==6.3.1
     # via open-clip-torch
@@ -269,7 +258,7 @@ geopandas==1.0.1
 gitdb==4.0.12
     # via gitpython
 gitpython==3.1.44
-    # via mlflow-skinny
+    # via wandb
 google-api-core==2.24.2
     # via
     #   google-cloud-core
@@ -277,7 +266,6 @@ google-api-core==2.24.2
     #   opencensus
 google-auth==2.40.2
     # via
-    #   databricks-sdk
     #   google-api-core
     #   google-cloud-core
     #   google-cloud-storage
@@ -296,25 +284,17 @@ googleapis-common-protos==1.70.0
     # via google-api-core
 gpt-oss==0.0.8
     # via -r requirements/test.in
-graphene==3.4.3
-    # via mlflow
 graphql-core==3.2.6
-    # via
-    #   graphene
-    #   graphql-relay
-    #   hypothesis-graphql
-graphql-relay==3.2.0
-    # via graphene
+    # via hypothesis-graphql
 greenlet==3.2.3
     # via sqlalchemy
 grpcio==1.78.0
     # via
     #   grpcio-tools
     #   ray
+    #   tensorboard
 grpcio-tools==1.78.0
     # via -r requirements/test.in
-gunicorn==23.0.0
-    # via mlflow
 h11==0.14.0
     # via
     #   httpcore
@@ -338,12 +318,14 @@ httpcore==1.0.6
 httpx==0.27.2
     # via
     #   -r requirements/test.in
+    #   diffusers
     #   perceptron
     #   schemathesis
 huggingface-hub==0.36.2
     # via
     #   accelerate
     #   datasets
+    #   diffusers
     #   evaluate
     #   open-clip-torch
     #   peft
@@ -379,11 +361,13 @@ idna==3.10
     #   jsonschema
     #   requests
     #   yarl
+imagehash==4.3.2
+    # via -r requirements/test.in
 imageio==2.37.0
     # via scikit-image
 importlib-metadata==8.7.0
     # via
-    #   mlflow-skinny
+    #   diffusers
     #   opentelemetry-api
 importlib-resources==6.5.2
     # via typeshed-client
@@ -395,14 +379,10 @@ isoduration==20.11.0
     # via jsonschema
 isort==5.13.2
     # via datamodel-code-generator
-itsdangerous==2.2.0
-    # via flask
 jinja2==3.1.6
     # via
     #   datamodel-code-generator
-    #   flask
     #   genai-perf
-    #   mlflow
     #   torch
 jiwer==3.0.5
     # via -r requirements/test.in
@@ -415,12 +395,14 @@ joblib==1.4.2
     #   librosa
     #   nltk
     #   scikit-learn
-jsonargparse==4.35.0
+jsonargparse==4.46.0
     # via
     #   lightning
     #   terratorch
 jsonlines==4.0.0
     # via lm-eval
+jsonnet==0.21.0
+    # via jsonargparse
 jsonpointer==3.0.0
     # via jsonschema
 jsonschema==4.23.0
@@ -449,13 +431,13 @@ libnacl==2.1.0
     # via tensorizer
 librosa==0.10.2.post1
     # via -r requirements/test.in
-lightly==1.5.20
+lightly==1.5.22
     # via
     #   terratorch
     #   torchgeo
 lightly-utils==0.0.2
     # via lightly
-lightning==2.5.1.post0
+lightning==2.6.1
     # via
     #   terratorch
     #   torchgeo
@@ -476,12 +458,11 @@ lxml==5.3.0
 mako==1.3.10
     # via alembic
 markdown==3.8.2
-    # via mlflow
+    # via tensorboard
 markdown-it-py==3.0.0
     # via rich
 markupsafe==3.0.1
     # via
-    #   flask
     #   jinja2
     #   mako
     #   werkzeug
@@ -489,7 +470,6 @@ matplotlib==3.9.2
     # via
     #   -r requirements/test.in
     #   lightning
-    #   mlflow
     #   pycocotools
     #   torchgeo
 mbstrdecoder==1.1.3
@@ -501,10 +481,6 @@ mdurl==0.1.2
     # via markdown-it-py
 mistral-common==1.9.1
     # via -r requirements/test.in
-mlflow==2.22.0
-    # via terratorch
-mlflow-skinny==2.22.0
-    # via mlflow
 more-itertools==10.5.0
     # via lm-eval
 mpmath==1.3.0
@@ -523,8 +499,6 @@ multiprocess==0.70.16
     # via
     #   datasets
     #   evaluate
-munch==4.0.0
-    # via pretrainedmodels
 mypy-extensions==1.0.0
     # via black
 networkx==3.2.1
@@ -553,6 +527,7 @@ numpy==2.2.6
     #   cupy-cuda12x
     #   datasets
     #   decord
+    #   diffusers
     #   einx
     #   encodec
     #   evaluate
@@ -560,13 +535,13 @@ numpy==2.2.6
     #   genai-perf
     #   geopandas
     #   h5py
+    #   imagehash
     #   imageio
     #   librosa
     #   lightly
     #   lightly-utils
     #   matplotlib
     #   mistral-common
-    #   mlflow
     #   mteb
     #   numba
     #   numexpr
@@ -578,6 +553,7 @@ numpy==2.2.6
     #   perceptron
     #   pycocotools
     #   pyogrio
+    #   pywavelets
     #   rasterio
     #   rioxarray
     #   rouge-score
@@ -590,8 +566,10 @@ numpy==2.2.6
     #   shapely
     #   soxr
     #   statsmodels
+    #   tensorboard
     #   tensorboardx
     #   tensorizer
+    #   terratorch
     #   tifffile
     #   torchgeo
     #   torchmetrics
@@ -659,7 +637,6 @@ opencv-python-headless==4.13.0.90
     #   mistral-common
 opentelemetry-api==1.35.0
     # via
-    #   mlflow-skinny
     #   opentelemetry-exporter-prometheus
     #   opentelemetry-sdk
     #   opentelemetry-semantic-conventions
@@ -669,7 +646,6 @@ opentelemetry-proto==1.36.0
     # via ray
 opentelemetry-sdk==1.35.0
     # via
-    #   mlflow-skinny
     #   opentelemetry-exporter-prometheus
     #   ray
 opentelemetry-semantic-conventions==0.56b0
@@ -687,7 +663,6 @@ packaging==24.2
     #   evaluate
     #   fastparquet
     #   geopandas
-    #   gunicorn
     #   huggingface-hub
     #   hydra-core
     #   kornia
@@ -695,7 +670,6 @@ packaging==24.2
     #   lightning
     #   lightning-utilities
     #   matplotlib
-    #   mlflow-skinny
     #   optuna
     #   peft
     #   plotly
@@ -708,10 +682,12 @@ packaging==24.2
     #   rioxarray
     #   scikit-image
     #   statsmodels
+    #   tensorboard
     #   tensorboardx
     #   torchmetrics
     #   transformers
     #   typepy
+    #   wandb
     #   xarray
 pandas==2.2.3
     # via
@@ -720,8 +696,8 @@ pandas==2.2.3
     #   fastparquet
     #   genai-perf
     #   geopandas
-    #   mlflow
     #   statsmodels
+    #   tacoreader
     #   torchgeo
     #   xarray
 pathspec==0.12.1
@@ -740,7 +716,9 @@ perf-analyzer==0.1.0
     # via genai-perf
 pillow==10.4.0
     # via
+    #   diffusers
     #   genai-perf
+    #   imagehash
     #   imageio
     #   lightly-utils
     #   matplotlib
@@ -748,6 +726,7 @@ pillow==10.4.0
     #   perceptron
     #   scikit-image
     #   segmentation-models-pytorch
+    #   tensorboard
     #   torchgeo
     #   torchvision
 platformdirs==4.3.6
@@ -755,6 +734,7 @@ platformdirs==4.3.6
     #   black
     #   pooch
     #   virtualenv
+    #   wandb
 plotly==5.24.1
     # via genai-perf
 pluggy==1.5.0
@@ -769,8 +749,6 @@ portalocker==2.10.1
     # via sacrebleu
 pqdm==0.2.0
     # via -r requirements/test.in
-pretrainedmodels==0.7.4
-    # via segmentation-models-pytorch
 prometheus-client==0.22.0
     # via
     #   opentelemetry-exporter-prometheus
@@ -786,12 +764,13 @@ protobuf==6.33.2
     #   google-api-core
     #   googleapis-common-protos
     #   grpcio-tools
-    #   mlflow-skinny
     #   opentelemetry-proto
     #   proto-plus
     #   ray
+    #   tensorboard
     #   tensorboardx
     #   tensorizer
+    #   wandb
 psutil==6.1.0
     # via
     #   accelerate
@@ -801,11 +780,12 @@ py==1.11.0
     # via pytest-forked
 py-spy==0.4.0
     # via ray
-pyarrow==18.0.0
+pyarrow==23.0.0
     # via
     #   datasets
     #   genai-perf
-    #   mlflow
+    #   tacoreader
+    #   terratorch
 pyasn1==0.6.1
     # via
     #   pyasn1-modules
@@ -831,11 +811,11 @@ pydantic==2.12.0
     #   gpt-oss
     #   lightly
     #   mistral-common
-    #   mlflow-skinny
     #   mteb
     #   openai-harmony
     #   pydantic-extra-types
     #   ray
+    #   wandb
 pydantic-core==2.41.1
     # via pydantic
 pydantic-extra-types==2.10.5
@@ -873,7 +853,6 @@ pytest==8.3.5
     #   pytest-subtests
     #   pytest-timeout
     #   schemathesis
-    #   terratorch
 pytest-asyncio==0.24.0
     # via -r requirements/test.in
 pytest-cov==6.3.0
@@ -896,7 +875,6 @@ python-dateutil==2.9.0.post0
     # via
     #   arrow
     #   botocore
-    #   graphene
     #   lightly
     #   matplotlib
     #   pandas
@@ -913,6 +891,8 @@ pytz==2024.2
     # via
     #   pandas
     #   typepy
+pywavelets==1.9.0
+    # via imagehash
 pyyaml==6.0.2
     # via
     #   accelerate
@@ -923,7 +903,6 @@ pyyaml==6.0.2
     #   huggingface-hub
     #   jsonargparse
     #   lightning
-    #   mlflow-skinny
     #   omegaconf
     #   optuna
     #   peft
@@ -934,6 +913,7 @@ pyyaml==6.0.2
     #   timm
     #   transformers
     #   vocos
+    #   wandb
 rapidfuzz==3.12.1
     # via jiwer
 rasterio==1.4.3
@@ -951,6 +931,7 @@ referencing==0.35.1
     #   jsonschema-specifications
 regex==2024.9.11
     # via
+    #   diffusers
     #   nltk
     #   open-clip-torch
     #   sacrebleu
@@ -959,8 +940,8 @@ regex==2024.9.11
 requests==2.32.3
     # via
     #   buildkite-test-collector
-    #   databricks-sdk
     #   datasets
+    #   diffusers
     #   docker
     #   evaluate
     #   google-api-core
@@ -970,15 +951,16 @@ requests==2.32.3
     #   lightly
     #   lm-eval
     #   mistral-common
-    #   mlflow-skinny
     #   mteb
     #   pooch
     #   ray
     #   responses
     #   schemathesis
     #   starlette-testclient
+    #   tacoreader
     #   tiktoken
     #   transformers
+    #   wandb
 responses==0.25.3
     # via genai-perf
 rfc3339-validator==0.1.4
@@ -991,6 +973,7 @@ rich==13.9.4
     #   lightning
     #   mteb
     #   perceptron
+    #   terratorch
     #   typer
 rioxarray==0.19.0
     # via terratorch
@@ -1017,47 +1000,55 @@ sacrebleu==2.4.3
 safetensors==0.4.5
     # via
     #   accelerate
+    #   diffusers
     #   open-clip-torch
     #   peft
+    #   segmentation-models-pytorch
     #   timm
     #   transformers
 schemathesis==3.39.15
     # via -r requirements/test.in
 scikit-image==0.25.2
-    # via albumentations
+    # via
+    #   albumentations
+    #   terratorch
 scikit-learn==1.5.2
     # via
     #   albumentations
     #   librosa
     #   lm-eval
-    #   mlflow
     #   mteb
     #   sentence-transformers
+    #   terratorch
 scipy==1.13.1
     # via
     #   albumentations
     #   bm25s
+    #   imagehash
     #   librosa
-    #   mlflow
     #   mteb
     #   scikit-image
     #   scikit-learn
     #   sentence-transformers
     #   statsmodels
     #   vocos
-segmentation-models-pytorch==0.4.0
+segmentation-models-pytorch==0.5.0
     # via
+    #   -r requirements/test.in
     #   terratorch
     #   torchgeo
 sentence-transformers==5.2.0
     # via
     #   -r requirements/test.in
     #   mteb
+sentry-sdk==2.52.0
+    # via wandb
 setuptools==77.0.3
     # via
     #   grpcio-tools
     #   lightning-utilities
     #   pytablewriter
+    #   tensorboard
     #   torch
 shapely==2.1.1
     # via
@@ -1075,7 +1066,6 @@ six==1.16.0
     #   python-dateutil
     #   rfc3339-validator
     #   rouge-score
-    #   segmentation-models-pytorch
 smart-open==7.1.0
     # via ray
 smmap==5.0.2
@@ -1099,12 +1089,9 @@ soxr==0.5.0.post1
 sqlalchemy==2.0.41
     # via
     #   alembic
-    #   mlflow
     #   optuna
 sqlitedict==2.1.0
     # via lm-eval
-sqlparse==0.5.3
-    # via mlflow-skinny
 starlette==0.50.0
     # via
     #   fastapi
@@ -1124,6 +1111,8 @@ tabledata==1.3.3
     # via pytablewriter
 tabulate==0.9.0
     # via sacrebleu
+tacoreader==0.5.6
+    # via terratorch
 tblib==3.1.0
     # via -r requirements/test.in
 tcolorpy==0.1.6
@@ -1133,13 +1122,19 @@ tenacity==9.1.2
     #   gpt-oss
     #   lm-eval
     #   plotly
+tensorboard==2.20.0
+    # via terratorch
+tensorboard-data-server==0.7.2
+    # via tensorboard
 tensorboardx==2.6.4
     # via lightning
 tensorizer==2.10.1
     # via -r requirements/test.in
 termcolor==3.1.0
-    # via gpt-oss
-terratorch @ git+https://github.com/IBM/terratorch.git@07184fcf91a1324f831ff521dd238d97fe350e3e
+    # via
+    #   gpt-oss
+    #   terratorch
+terratorch==1.2.2
     # via -r requirements/test.in
 threadpoolctl==3.5.0
     # via scikit-learn
@@ -1172,7 +1167,6 @@ torch==2.10.0+cu129
     #   -r requirements/test.in
     #   accelerate
     #   bitsandbytes
-    #   efficientnet-pytorch
     #   encodec
     #   kornia
     #   lightly
@@ -1181,7 +1175,6 @@ torch==2.10.0+cu129
     #   mteb
     #   open-clip-torch
     #   peft
-    #   pretrainedmodels
     #   pytorch-lightning
     #   runai-model-streamer
     #   segmentation-models-pytorch
@@ -1213,12 +1206,11 @@ torchvision==0.25.0+cu129
     #   -r requirements/test.in
     #   lightly
     #   open-clip-torch
-    #   pretrainedmodels
     #   segmentation-models-pytorch
     #   terratorch
     #   timm
     #   torchgeo
-tqdm==4.66.6
+tqdm==4.67.3
     # via
     #   datasets
     #   evaluate
@@ -1232,10 +1224,11 @@ tqdm==4.66.6
     #   optuna
     #   peft
     #   pqdm
-    #   pretrainedmodels
     #   pytorch-lightning
     #   segmentation-models-pytorch
     #   sentence-transformers
+    #   tacoreader
+    #   terratorch
     #   tqdm-multiprocess
     #   transformers
 tqdm-multiprocess==0.0.11
@@ -1274,14 +1267,12 @@ typing-extensions==4.15.0
     #   alembic
     #   chz
     #   fastapi
-    #   graphene
     #   grpcio
     #   huggingface-hub
     #   librosa
     #   lightning
     #   lightning-utilities
     #   mistral-common
-    #   mlflow-skinny
     #   mteb
     #   opentelemetry-api
     #   opentelemetry-sdk
@@ -1299,6 +1290,7 @@ typing-extensions==4.15.0
     #   typer
     #   typeshed-client
     #   typing-inspection
+    #   wandb
 typing-inspection==0.4.2
     # via pydantic
 tzdata==2024.2
@@ -1313,25 +1305,26 @@ urllib3==2.2.3
     #   lightly
     #   requests
     #   responses
+    #   sentry-sdk
     #   tritonclient
 uvicorn==0.35.0
-    # via
-    #   gpt-oss
-    #   mlflow-skinny
+    # via gpt-oss
 vector-quantize-pytorch==1.21.2
     # via -r requirements/test.in
 virtualenv==20.31.2
     # via ray
 vocos==0.1.0
     # via -r requirements/test.in
+wandb==0.24.2
+    # via terratorch
 wcwidth==0.2.13
     # via ftfy
 webcolors==24.11.1
     # via jsonschema
 werkzeug==3.1.3
     # via
-    #   flask
     #   schemathesis
+    #   tensorboard
 word2number==1.1
     # via lm-eval
 wrapt==1.17.2

scripts/autotune_helion_kernels.py

@@ -0,0 +1,430 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Autotune registered Helion kernels for optimal configurations.
+
+Usage:
+    # Autotune all registered kernels
+    python scripts/autotune_helion_kernels.py
+
+    # Autotune specific kernel
+    python scripts/autotune_helion_kernels.py --kernels silu_mul_fp8
+
+    # Autotune multiple kernels
+    python scripts/autotune_helion_kernels.py --kernels silu_mul_fp8 rms_norm_fp8
+
+    # Force re-autotuning
+    python scripts/autotune_helion_kernels.py --force
+
+    # List available kernels
+    python scripts/autotune_helion_kernels.py --list
+"""
+
+import argparse
+import sys
+import time
+from dataclasses import dataclass
+
+import torch
+
+try:
+    import helion
+
+    from vllm.kernels.helion import (
+        ConfigManager,
+        get_kernel_by_name,
+        get_registered_kernels,
+    )
+    from vllm.kernels.helion.utils import get_canonical_gpu_name
+    from vllm.logger import init_logger
+    from vllm.utils.import_utils import has_helion
+except ImportError as e:
+    print(f"Error importing vLLM: {e}")
+    print("Please ensure vLLM is installed and in your Python path")
+    sys.exit(1)
+
+logger = init_logger("vllm.scripts.autotune_helion_kernels")
+
+
+@dataclass
+class AutotuneResult:
+    status: str  # "success" | "partial" | "error" | "skipped"
+    successful: int
+    failed: int
+    configs: dict[str, "helion.Config"]
+    message: str = ""
+
+
+def list_kernels() -> None:
+    kernels = get_registered_kernels()
+
+    if not kernels:
+        print("No Helion kernels found in registry.")
+        return
+
+    print("Available Helion kernels:")
+    print("=" * 50)
+
+    for name in sorted(kernels.keys()):
+        print(f"  {name}")
+
+    print(f"\nTotal: {len(kernels)} kernels")
+
+
+def check_requirements() -> bool:
+    if not torch.cuda.is_available():
+        logger.error("CUDA is not available. Helion autotuning requires GPU.")
+        return False
+
+    if not has_helion():
+        logger.error("Helion is not installed. Please install Helion package.")
+        return False
+
+    return True
+
+
+def autotune_kernel(
+    kernel_name: str,
+    platform: str,
+    config_manager: ConfigManager,
+    force: bool = False,
+    autotune_effort: str = "quick",
+) -> AutotuneResult:
+    logger.debug(
+        "Starting autotune for kernel '%s' with effort='%s'",
+        kernel_name,
+        autotune_effort,
+    )
+    kernel_wrapper = get_kernel_by_name(kernel_name)
+    if kernel_wrapper is None:
+        error_msg = f"Kernel '{kernel_name}' not found in registry"
+        logger.error(error_msg)
+        return AutotuneResult(
+            status="error",
+            message=error_msg,
+            successful=0,
+            failed=0,
+            configs={},
+        )
+
+    try:
+        inputs_dict = kernel_wrapper.get_inputs()
+    except NotImplementedError:
+        error_msg = f"Kernel '{kernel_name}' has no input generator registered"
+        logger.error(error_msg)
+        return AutotuneResult(
+            status="error",
+            message=error_msg,
+            successful=0,
+            failed=0,
+            configs={},
+        )
+
+    try:
+        logger.info(
+            "Autotuning kernel '%s' for platform '%s' with %d configs",
+            kernel_name,
+            platform,
+            len(inputs_dict),
+        )
+
+        configs_to_autotune = {}
+        if not force:
+            existing_configs = config_manager.get_platform_configs(
+                kernel_name, platform
+            )
+            for config_key, inputs in inputs_dict.items():
+                if config_key in existing_configs:
+                    logger.debug(
+                        "Config '%s' already exists for platform '%s', skipping",
+                        config_key,
+                        platform,
+                    )
+                else:
+                    configs_to_autotune[config_key] = inputs
+        else:
+            logger.debug("Force mode enabled, will re-autotune all configs")
+            configs_to_autotune = inputs_dict
+
+        if not configs_to_autotune:
+            logger.info(
+                "All configs already exist for kernel '%s' on platform '%s'. "
+                "Use --force to re-autotune.",
+                kernel_name,
+                platform,
+            )
+            return AutotuneResult(
+                status="skipped",
+                message="All configs already exist",
+                successful=0,
+                failed=0,
+                configs={},
+            )
+
+        total_start_time = time.time()
+        autotuned_configs = {}
+        failed_configs = []
+
+        for config_key, inputs in configs_to_autotune.items():
+            logger.info("Autotuning config: %s", config_key)
+            logger.debug(
+                "Input shapes: %s",
+                [getattr(inp, "shape", type(inp).__name__) for inp in inputs],
+            )
+
+            try:
+                config_start_time = time.time()
+                config = kernel_wrapper.run_autotune(inputs, autotune_effort)
+                config_duration = time.time() - config_start_time
+
+                # Save immediately for checkpointing
+                config_manager.save_configs(kernel_name, platform, {config_key: config})
+
+                autotuned_configs[config_key] = config
+                logger.debug("Config details: %s", config)
+
+                logger.info(
+                    "✓ Autotuned and saved config '%s' (%.2fs)",
+                    config_key,
+                    config_duration,
+                )
+
+            except (RuntimeError, ValueError, OSError) as e:
+                logger.exception(
+                    "Failed to autotune config '%s': %s",
+                    config_key,
+                    e,
+                )
+                failed_configs.append(config_key)
+
+        total_duration = time.time() - total_start_time
+        successful = len(autotuned_configs)
+        failed = len(failed_configs)
+
+        logger.info(
+            "Completed autotuning for kernel '%s': %d successful, %d failed (%.2fs)",
+            kernel_name,
+            successful,
+            failed,
+            total_duration,
+        )
+
+        status = "success" if failed == 0 else "partial"
+        return AutotuneResult(
+            status=status,
+            successful=successful,
+            failed=failed,
+            configs=autotuned_configs,
+        )
+
+    except (KeyError, RuntimeError, ValueError, OSError) as e:
+        error_msg = f"Unexpected error: {e}"
+        logger.exception("Failed to autotune kernel '%s': %s", kernel_name, e)
+        return AutotuneResult(
+            status="error",
+            message=error_msg,
+            successful=0,
+            failed=0,
+            configs={},
+        )
+
+
+def summarize_results(results: dict[str, AutotuneResult]) -> bool:
+    logger.info("=" * 50)
+    logger.info("Autotuning Results Summary")
+    logger.info("=" * 50)
+
+    total_successful = 0
+    total_failed = 0
+    success_kernels = []
+    partial_kernels = []
+    error_kernels = []
+    skipped_kernels = []
+
+    for kernel_name, result in results.items():
+        total_successful += result.successful
+        total_failed += result.failed
+
+        if result.status == "success":
+            success_kernels.append(f"{kernel_name} ({result.successful} configs)")
+            logger.info("✓ %s: %d configs successful", kernel_name, result.successful)
+        elif result.status == "partial":
+            partial_kernels.append(
+                f"{kernel_name} ({result.successful} ok, {result.failed} failed)"
+            )
+            logger.warning(
+                "⚠ %s: %d successful, %d failed",
+                kernel_name,
+                result.successful,
+                result.failed,
+            )
+        elif result.status == "error":
+            error_kernels.append(f"{kernel_name}: {result.message or 'Unknown error'}")
+            logger.error("✗ %s: %s", kernel_name, result.message or "Unknown error")
+        elif result.status == "skipped":
+            skipped_kernels.append(f"{kernel_name}: {result.message or 'Skipped'}")
+            logger.info("- %s: %s", kernel_name, result.message or "Skipped")
+
+    logger.info("=" * 50)
+    logger.info(
+        "Summary: %d total configs (%d successful, %d failed)",
+        total_successful + total_failed,
+        total_successful,
+        total_failed,
+    )
+    logger.info(
+        "Kernels: %d success, %d partial, %d error, %d skipped",
+        len(success_kernels),
+        len(partial_kernels),
+        len(error_kernels),
+        len(skipped_kernels),
+    )
+
+    has_failures = bool(error_kernels or partial_kernels)
+
+    if not has_failures:
+        if total_successful > 0:
+            logger.info("All configs autotuned successfully!")
+        else:
+            logger.info("No new configs were generated (all may already exist)")
+
+    return not has_failures
+
+
+def get_kernels_to_autotune(requested_kernels: list[str] | None) -> list[str]:
+    all_kernels = get_registered_kernels()
+    if not all_kernels:
+        logger.error("No Helion kernels found in registry")
+        sys.exit(1)
+
+    if not requested_kernels:
+        return list(all_kernels.keys())
+
+    if len(requested_kernels) != len(set(requested_kernels)):
+        duplicates = [
+            k for k in set(requested_kernels) if requested_kernels.count(k) > 1
+        ]
+        logger.error("Duplicate kernel names in --kernels flag: %s", duplicates)
+        sys.exit(1)
+
+    kernels_to_autotune = []
+    missing_kernels = []
+
+    for kernel_name in requested_kernels:
+        if kernel_name in all_kernels:
+            kernels_to_autotune.append(kernel_name)
+        else:
+            missing_kernels.append(kernel_name)
+
+    if missing_kernels:
+        logger.error("Kernel(s) not found: %s", missing_kernels)
+        logger.error("Available kernels: %s", list(all_kernels.keys()))
+        sys.exit(1)
+
+    return kernels_to_autotune
+
+
+def main():
+    parser = argparse.ArgumentParser(
+        description="Autotune Helion kernels",
+        formatter_class=argparse.RawDescriptionHelpFormatter,
+        epilog=__doc__.split("Usage:")[1] if "Usage:" in __doc__ else "",
+    )
+
+    parser.add_argument(
+        "--kernels",
+        nargs="+",
+        help="Kernel(s) to autotune (default: all kernels)",
+    )
+
+    parser.add_argument(
+        "--config-dir",
+        type=str,
+        help="Config directory for config files (default: vLLM helion configs dir)",
+    )
+
+    parser.add_argument(
+        "--list",
+        action="store_true",
+        help="List available Helion kernels and exit",
+    )
+
+    parser.add_argument(
+        "--force",
+        action="store_true",
+        help=(
+            "Force re-autotuning even if configs already exist for the "
+            "platform and config keys"
+        ),
+    )
+
+    parser.add_argument(
+        "--autotune-effort",
+        type=str,
+        default="quick",
+        help=(
+            "Helion autotune effort level: 'quick' (smaller search) or "
+            "'full' (full search budget) (default: quick)"
+        ),
+    )
+
+    parser.add_argument(
+        "--verbose",
+        action="store_true",
+        help="Enable verbose logging",
+    )
+
+    args = parser.parse_args()
+
+    import logging
+
+    if args.verbose:
+        logging.getLogger("vllm").setLevel(logging.DEBUG)
+        logger.debug("Verbose mode enabled")
+        logger.debug("Arguments: %s", vars(args))
+    else:
+        logging.getLogger("vllm").setLevel(logging.INFO)
+
+    if args.list:
+        list_kernels()
+        return
+
+    if not check_requirements():
+        sys.exit(1)
+
+    platform = get_canonical_gpu_name()
+    logger.info("Detected GPU platform: %s", platform)
+
+    config_manager = (
+        ConfigManager(args.config_dir) if args.config_dir else ConfigManager()
+    )
+
+    try:
+        config_manager.ensure_base_dir_writable()
+    except OSError as e:
+        logger.error("Failed to access config directory: %s", e)
+        sys.exit(1)
+
+    kernels_to_autotune = get_kernels_to_autotune(args.kernels)
+
+    logger.info(
+        "Will autotune %d kernel(s) for platform '%s': %s",
+        len(kernels_to_autotune),
+        platform,
+        kernels_to_autotune,
+    )
+
+    results = {}
+    for kernel_name in kernels_to_autotune:
+        result = autotune_kernel(
+            kernel_name, platform, config_manager, args.force, args.autotune_effort
+        )
+        results[kernel_name] = result
+
+    success = summarize_results(results)
+    sys.exit(0 if success else 1)
+
+
+if __name__ == "__main__":
+    main()

tests/basic_correctness/test_cpu_offload.py

@@ -1,10 +1,29 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+import pytest
+
 from ..utils import compare_two_settings
 
 
-def test_cpu_offload():
+@pytest.mark.parametrize("disable_pin_memory", [False, True])
+@pytest.mark.parametrize("disable_uva", [False, True])
+def test_cpu_offload(disable_pin_memory, disable_uva):
+    env_vars = {
+        "VLLM_WEIGHT_OFFLOADING_DISABLE_PIN_MEMORY": str(int(disable_pin_memory)),
+        "VLLM_WEIGHT_OFFLOADING_DISABLE_UVA": str(int(disable_uva)),
+    }
+
+    args = ["--cpu-offload-gb", "1"]
+
+    # cuda graph only works with UVA offloading
+    if disable_uva:
+        args.append("--enforce-eager")
+
     compare_two_settings(
-        "hmellor/tiny-random-LlamaForCausalLM", [], ["--cpu-offload-gb", "1"]
+        model="hmellor/tiny-random-LlamaForCausalLM",
+        arg1=[],
+        arg2=args,
+        env1=None,
+        env2=env_vars,
     )

tests/config/test_multimodal_config.py

@@ -3,6 +3,7 @@
 
 import pytest
 
+from vllm.config.model import ModelConfig
 from vllm.config.multimodal import MultiModalConfig
 from vllm.v1.attention.backends.registry import AttentionBackendEnum
 
@@ -23,3 +24,20 @@ def test_mm_encoder_attn_backend_hash_updates():
         mm_encoder_attn_backend=AttentionBackendEnum.FLASH_ATTN
     ).compute_hash()
     assert base_hash != overridden_hash
+
+
+def test_language_model_only_does_not_affect_mm_hash():
+    """language_model_only does not affect the ViT computation graph,
+    so it should not change the multimodal config hash."""
+    base_hash = MultiModalConfig().compute_hash()
+    lm_only_hash = MultiModalConfig(language_model_only=True).compute_hash()
+    assert base_hash == lm_only_hash
+
+
+def test_language_model_only_affects_model_hash():
+    """language_model_only affects the LM computation graph,
+    so it should change the model config hash."""
+    model = "llava-hf/llava-1.5-7b-hf"
+    base_hash = ModelConfig(model).compute_hash()
+    lm_only_hash = ModelConfig(model, language_model_only=True).compute_hash()
+    assert base_hash != lm_only_hash

tests/conftest.py

@@ -419,7 +419,6 @@ class HfRunner:
             self.tokenizer: "PreTrainedTokenizer | PreTrainedTokenizerFast" = (
                 AutoTokenizer.from_pretrained(
                     model_name,
-                    dtype=dtype,
                     trust_remote_code=trust_remote_code,
                 )
             )
@@ -430,7 +429,6 @@ class HfRunner:
 
         self.processor = AutoProcessor.from_pretrained(
             model_name,
-            dtype=dtype,
             trust_remote_code=trust_remote_code,
         )
         if skip_tokenizer_init:

tests/detokenizer/test_min_tokens.py

@@ -39,7 +39,6 @@ def test_min_tokens_with_stop(min_tokens: int, stop: str, truth: str):
         mm_features=None,
         sampling_params=params,
         pooling_params=None,
-        eos_token_id=None,
         arrival_time=0.0,
         lora_request=None,
         cache_salt=None,

tests/detokenizer/test_stop_string_while_stop_model_terminates.py

@@ -35,7 +35,6 @@ def _make_request(stop, include_stop_str_in_output: bool, min_tokens: int = 0):
         mm_features=None,
         sampling_params=params,
         pooling_params=None,
-        eos_token_id=None,
         arrival_time=0.0,
         lora_request=None,
         cache_salt=None,

tests/distributed/test_comm_ops.py

@@ -19,6 +19,8 @@ from vllm.distributed import (
     tensor_model_parallel_all_reduce,
     tensor_model_parallel_reduce_scatter,
 )
+from vllm.distributed.parallel_state import GroupCoordinator, TensorMetadata
+from vllm.v1.worker.gpu_worker import AsyncIntermediateTensors
 
 from ..utils import (
     init_test_distributed_environment,
@@ -200,6 +202,111 @@ def send_recv_tensor_dict_test_worker(
         torch.testing.assert_close(recv_dict["f"], test_dict["f"])
 
 
+class _DummyWork:
+    def __init__(self) -> None:
+        self.wait_calls = 0
+
+    def wait(self) -> None:
+        self.wait_calls += 1
+
+
+class _DummyAllGatherGroup:
+    def __init__(self, world_size: int, rank_in_group: int) -> None:
+        self.world_size = world_size
+        self.rank_in_group = rank_in_group
+
+    def all_gather(self, t: torch.Tensor, dim: int = 0) -> torch.Tensor:
+        # duplicate local slice across ranks.
+        assert dim == 0
+        return torch.cat([t for _ in range(self.world_size)], dim=0)
+
+
+def _make_group_for_unit_test(
+    rank_in_group: int = 0, world_size: int = 2
+) -> GroupCoordinator:
+    # avoid running GroupCoordinator.__init__ (it wires up real process groups).
+    g = GroupCoordinator.__new__(GroupCoordinator)
+    g.world_size = world_size
+    g.rank_in_group = rank_in_group
+    g.ranks = list(range(world_size))
+    g.use_cpu_custom_send_recv = False
+    g.device_group = None
+    g.cpu_group = None
+    return g
+
+
+def test_irecv_tensor_dict_send_allgather_postprocess_binds_keys(
+    monkeypatch: pytest.MonkeyPatch,
+) -> None:
+    def fake_irecv(t: torch.Tensor, *args: Any, **kwargs: Any) -> _DummyWork:
+        t.fill_(1)
+        return _DummyWork()
+
+    monkeypatch.setattr(torch.distributed, "is_initialized", lambda: True)
+    monkeypatch.setattr(torch.distributed, "irecv", fake_irecv)
+
+    g = _make_group_for_unit_test(rank_in_group=0, world_size=2)
+    # 2 tensors so we can catch late-binding bugs in postprocess closures.
+    metadata_list = [
+        ("a", TensorMetadata("cpu", torch.int32, torch.Size([4]))),
+        ("b", TensorMetadata("cpu", torch.int32, torch.Size([4]))),
+    ]
+    g.recv_object = lambda src=None: metadata_list  # type: ignore[method-assign]
+
+    ag = _DummyAllGatherGroup(world_size=2, rank_in_group=0)
+    td, handles, postprocess = g.irecv_tensor_dict(all_gather_group=ag)
+
+    assert td is not None
+    assert len(handles) == 2
+    assert len(postprocess) == 2
+
+    # before postprocess, dict holds the TP slice (shape 2).
+    assert td["a"].shape == torch.Size([2])
+    assert td["b"].shape == torch.Size([2])
+
+    # simulate worker-side "defer wait": wait + postprocess later.
+    for handle in handles:
+        handle.wait()
+    for fn in postprocess:
+        fn()
+
+    # after postprocess, dict values are reconstructed to full shape (shape 4),
+    # and each key should be updated independently
+    assert td["a"].shape == torch.Size([4])
+    assert td["b"].shape == torch.Size([4])
+    torch.testing.assert_close(td["a"], torch.ones(4, dtype=torch.int32))
+    torch.testing.assert_close(td["b"], torch.ones(4, dtype=torch.int32))
+
+
+def test_async_intermediate_tensors_lazy_wait() -> None:
+    work = _DummyWork()
+    post_calls = {"n": 0}
+
+    def post() -> None:
+        post_calls["n"] += 1
+
+    it = AsyncIntermediateTensors(
+        {"x": torch.tensor([1])},
+        comm_handles=[work],
+        comm_postprocess=[post],
+    )
+
+    # accessing non-tensor attributes should not trigger wait.
+    assert it.kv_connector_output is None
+    assert work.wait_calls == 0
+    assert post_calls["n"] == 0
+
+    # first access of `.tensors` triggers wait + postprocess.
+    _ = it.tensors
+    assert work.wait_calls == 1
+    assert post_calls["n"] == 1
+
+    # subsequent access should not re-wait.
+    _ = it.tensors
+    assert work.wait_calls == 1
+    assert post_calls["n"] == 1
+
+
 @ray.remote(num_gpus=1, max_calls=1)
 def send_recv_test_worker(
     monkeypatch: pytest.MonkeyPatch,

tests/entrypoints/openai/responses/test_sampling_params.py

@@ -4,8 +4,17 @@
 """Unit tests for ResponsesRequest.to_sampling_params() parameter mapping."""
 
 import pytest
+import torch
+from openai.types.responses.response_format_text_json_schema_config import (
+    ResponseFormatTextJSONSchemaConfig,
+)
+from pydantic import ValidationError
 
-from vllm.entrypoints.openai.responses.protocol import ResponsesRequest
+from vllm.entrypoints.openai.responses.protocol import (
+    ResponsesRequest,
+    ResponseTextConfig,
+)
+from vllm.sampling_params import StructuredOutputsParams
 
 
 class TestResponsesRequestSamplingParams:
@@ -76,9 +85,6 @@ class TestResponsesRequestSamplingParams:
 
     def test_seed_bounds_validation(self):
         """Test that seed values outside torch.long bounds are rejected."""
-        import torch
-        from pydantic import ValidationError
-
         # Test seed below minimum
         with pytest.raises(ValidationError) as exc_info:
             ResponsesRequest(
@@ -111,3 +117,40 @@ class TestResponsesRequestSamplingParams:
             seed=torch.iinfo(torch.long).max,
         )
         assert request_max.seed == torch.iinfo(torch.long).max
+
+    def test_structured_outputs_passed_through(self):
+        """Test that structured_outputs field is passed to SamplingParams."""
+        structured_outputs = StructuredOutputsParams(grammar="root ::= 'hello'")
+        request = ResponsesRequest(
+            model="test-model",
+            input="test input",
+            structured_outputs=structured_outputs,
+        )
+
+        sampling_params = request.to_sampling_params(default_max_tokens=1000)
+
+        assert sampling_params.structured_outputs is not None
+        assert sampling_params.structured_outputs.grammar == "root ::= 'hello'"
+
+    def test_structured_outputs_and_json_schema_conflict(self):
+        """Test that specifying both structured_outputs and json_schema raises."""
+        structured_outputs = StructuredOutputsParams(grammar="root ::= 'hello'")
+        text_config = ResponseTextConfig()
+        text_config.format = ResponseFormatTextJSONSchemaConfig(
+            type="json_schema",
+            name="test",
+            schema={"type": "object"},
+        )
+        request = ResponsesRequest(
+            model="test-model",
+            input="test input",
+            structured_outputs=structured_outputs,
+            text=text_config,
+        )
+
+        with pytest.raises(ValueError) as exc_info:
+            request.to_sampling_params(default_max_tokens=1000)
+
+        assert "Cannot specify both structured_outputs and text.format" in str(
+            exc_info.value
+        )

tests/entrypoints/openai/test_chat_error.py

@@ -54,16 +54,22 @@ class MockModelConfig:
     media_io_kwargs: dict[str, dict[str, Any]] = field(default_factory=dict)
     skip_tokenizer_init = False
     is_encoder_decoder: bool = False
+    is_multimodal_model: bool = False
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}
 
 
+@dataclass
+class MockVllmConfig:
+    model_config: MockModelConfig
+
+
 def _build_renderer(model_config: MockModelConfig):
     _, tokenizer_name, _, kwargs = tokenizer_args_from_config(model_config)
 
-    return HfRenderer(
-        model_config,
+    return HfRenderer.from_config(
+        MockVllmConfig(model_config),
         tokenizer_kwargs={**kwargs, "tokenizer_name": tokenizer_name},
     )
 

tests/entrypoints/openai/test_completion_error.py

@@ -53,11 +53,17 @@ class MockModelConfig:
     media_io_kwargs: dict[str, dict[str, Any]] = field(default_factory=dict)
     skip_tokenizer_init = False
     is_encoder_decoder: bool = False
+    is_multimodal_model: bool = False
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}
 
 
+@dataclass
+class MockVllmConfig:
+    model_config: MockModelConfig
+
+
 def _build_serving_completion(engine: AsyncLLM) -> OpenAIServingCompletion:
     models = OpenAIServingModels(
         engine_client=engine,
@@ -73,8 +79,8 @@ def _build_serving_completion(engine: AsyncLLM) -> OpenAIServingCompletion:
 def _build_renderer(model_config: MockModelConfig):
     _, tokenizer_name, _, kwargs = tokenizer_args_from_config(model_config)
 
-    return HfRenderer(
-        model_config,
+    return HfRenderer.from_config(
+        MockVllmConfig(model_config),
         tokenizer_kwargs={**kwargs, "tokenizer_name": tokenizer_name},
     )
 

tests/entrypoints/openai/test_gptoss_structural_tags_integration.py

@@ -23,6 +23,7 @@ class TestGptOssStructuralTagsIntegration:
         """Create a mock tokenizer."""
         tokenizer = Mock()
         tokenizer.encode = Mock(return_value=[1, 2, 3, 4, 5])
+        tokenizer.vocab = {"<|end|>": 6}
         return tokenizer
 
     @pytest.fixture

tests/entrypoints/openai/test_lora_resolvers.py

@@ -52,11 +52,17 @@ class MockModelConfig:
     generation_config: str = "auto"
     skip_tokenizer_init: bool = False
     is_encoder_decoder: bool = False
+    is_multimodal_model: bool = False
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}
 
 
+@dataclass
+class MockVllmConfig:
+    model_config: MockModelConfig
+
+
 class MockLoRAResolver(LoRAResolver):
     async def resolve_lora(
         self, base_model_name: str, lora_name: str
@@ -90,8 +96,8 @@ def register_mock_resolver():
 def _build_renderer(model_config: MockModelConfig):
     _, tokenizer_name, _, kwargs = tokenizer_args_from_config(model_config)
 
-    return HfRenderer(
-        model_config,
+    return HfRenderer.from_config(
+        MockVllmConfig(model_config),
         tokenizer_kwargs={**kwargs, "tokenizer_name": tokenizer_name},
     )
 

tests/entrypoints/openai/test_serving_chat.py

@@ -529,16 +529,22 @@ class MockModelConfig:
     media_io_kwargs: dict[str, dict[str, Any]] = field(default_factory=dict)
     skip_tokenizer_init: bool = False
     is_encoder_decoder: bool = False
+    is_multimodal_model: bool = False
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}
 
 
+@dataclass
+class MockVllmConfig:
+    model_config: MockModelConfig
+
+
 def _build_renderer(model_config: MockModelConfig):
     _, tokenizer_name, _, kwargs = tokenizer_args_from_config(model_config)
 
-    return HfRenderer(
-        model_config,
+    return HfRenderer.from_config(
+        MockVllmConfig(model_config),
         tokenizer_kwargs={**kwargs, "tokenizer_name": tokenizer_name},
     )
 
@@ -749,8 +755,10 @@ async def test_serving_chat_mistral_token_ids_prompt_is_validated():
     mock_engine.io_processor = MagicMock()
 
     mock_tokenizer = MagicMock(spec=MistralTokenizer)
-    mock_renderer = MistralRenderer(mock_engine.model_config, tokenizer_kwargs={})
-    mock_renderer._tokenizer = mock_tokenizer
+    mock_renderer = MistralRenderer(
+        MockVllmConfig(mock_engine.model_config),
+        tokenizer=mock_tokenizer,
+    )
     # Force the Mistral chat template renderer to return token IDs.
     # Choose a prompt length that is < max_model_len, but large enough that
     # adding max_tokens should exceed the model context window.
@@ -788,8 +796,10 @@ async def test_serving_chat_mistral_token_ids_prompt_too_long_is_rejected():
     mock_engine.io_processor = MagicMock()
 
     mock_tokenizer = MagicMock(spec=MistralTokenizer)
-    mock_renderer = MistralRenderer(mock_engine.model_config, tokenizer_kwargs={})
-    mock_renderer._tokenizer = mock_tokenizer
+    mock_renderer = MistralRenderer(
+        MockVllmConfig(mock_engine.model_config),
+        tokenizer=mock_tokenizer,
+    )
     # prompt_token_ids length == max_model_len should be rejected for
     # completion-like requests (ChatCompletionRequest).
     mock_renderer.render_messages_async = AsyncMock(

tests/entrypoints/openai/test_serving_tokens.py

@@ -1,6 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+import os
+
 import httpx
 import pytest
 import pytest_asyncio
@@ -46,6 +48,27 @@ def server(request):
         "--max-model-len",
         "1024",
         "--enforce-eager",
+        # On ROCm (e.g. MI355X/gfx950), bf16 GEMM results can differ by
+        # 1 ULP when the batch dimension (M) changes, because different M
+        # values cause the Tensile backend to select different tile
+        # configurations with different fp32 accumulation orders. With
+        # prefix caching, cache-miss prefills compute all tokens in one
+        # pass (large M) while cache-hit requests compute only the
+        # uncached suffix (small M), seeding a divergence that amplifies
+        # through the residual stream and flips argmax tokens.
+        # See: https://github.com/vllm-project/vllm/issues/33123
+        #
+        # Either disable prefix caching entirely, or enable it with
+        # --deterministic-prefix-caching which forces cache-miss prefills
+        # to split at block boundaries so the suffix GEMM shape is always
+        # identical regardless of cache state.
+        #
+        # Option A: disable prefix caching
+        "--no-enable-prefix-caching",
+        #
+        # Option B: deterministic prefix caching
+        # "--enable-prefix-caching",
+        # "--deterministic-prefix-caching",
     ]
 
     extra_args = getattr(request, "param", None)
@@ -56,7 +79,11 @@ def server(request):
             else [str(extra_args)]
         )
 
-    with RemoteOpenAIServer(MODEL_NAME, args) as remote_server:
+    envs = os.environ.copy()
+    # See: https://github.com/vllm-project/vllm/pull/33493#issuecomment-3888060787
+    envs["VLLM_ROCM_USE_SKINNY_GEMM"] = "0"
+
+    with RemoteOpenAIServer(MODEL_NAME, args, env_dict=envs) as remote_server:
         yield remote_server
 
 

tests/entrypoints/openai/test_video.py

@@ -13,13 +13,12 @@ from vllm.platforms import current_platform
 from ...utils import RemoteOpenAIServer
 
 MODEL_NAME = "llava-hf/llava-onevision-qwen2-0.5b-ov-hf"
-MAXIMUM_VIDEOS = 4
+MAXIMUM_VIDEOS = 3
 
 TEST_VIDEO_URLS = [
-    "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/BigBuckBunny.mp4",
-    "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/ElephantsDream.mp4",
-    "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/ForBiggerBlazes.mp4",
-    "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/ForBiggerFun.mp4",
+    "https://www.bogotobogo.com/python/OpenCV_Python/images/mean_shift_tracking/slow_traffic_small.mp4",
+    "https://github.com/opencv/opencv/raw/refs/tags/4.12.0/samples/data/vtest.avi",
+    "https://github.com/opencv/opencv/raw/refs/tags/4.12.0/samples/data/Megamind.avi",
 ]
 
 

tests/entrypoints/pooling/score/test_online_colbert.py

@@ -8,10 +8,8 @@ import requests
 from tests.utils import RemoteOpenAIServer
 from vllm.entrypoints.pooling.score.protocol import RerankResponse, ScoreResponse
 
-# ColBERT model - using answerai-colbert-small-v1 as it's a smaller model
 MODEL_NAME = "answerdotai/answerai-colbert-small-v1"
-COLBERT_DIM = 96  # This model uses 96-dimensional output
-DTYPE = "half"
+COLBERT_DIM = 96
 MAX_MODEL_LEN = 512
 
 
@@ -26,129 +24,119 @@ def server():
         yield remote_server
 
 
-@pytest.mark.parametrize("model_name", [MODEL_NAME])
-def test_colbert_rerank(server: RemoteOpenAIServer, model_name: str):
-    """Test ColBERT rerank endpoint."""
-    query = "What is the capital of France?"
-    documents = [
-        "The capital of Brazil is Brasilia.",
-        "The capital of France is Paris.",
-    ]
+class TestColBERTOnline:
+    def test_rerank(self, server: RemoteOpenAIServer):
+        """Test ColBERT rerank endpoint."""
+        query = "What is the capital of France?"
+        documents = [
+            "The capital of Brazil is Brasilia.",
+            "The capital of France is Paris.",
+        ]
 
-    rerank_response = requests.post(
-        server.url_for("rerank"),
-        json={
-            "model": model_name,
-            "query": query,
-            "documents": documents,
-        },
-    )
-    rerank_response.raise_for_status()
-    rerank = RerankResponse.model_validate(rerank_response.json())
+        rerank_response = requests.post(
+            server.url_for("rerank"),
+            json={
+                "model": MODEL_NAME,
+                "query": query,
+                "documents": documents,
+            },
+        )
+        rerank_response.raise_for_status()
+        rerank = RerankResponse.model_validate(rerank_response.json())
 
-    assert rerank.id is not None
-    assert rerank.results is not None
-    assert len(rerank.results) == 2
+        assert rerank.id is not None
+        assert rerank.results is not None
+        assert len(rerank.results) == 2
 
-    # The relevant document (Paris) should have higher score
-    paris_result = next(r for r in rerank.results if r.index == 1)
-    brazil_result = next(r for r in rerank.results if r.index == 0)
+        paris_result = next(r for r in rerank.results if r.index == 1)
+        brazil_result = next(r for r in rerank.results if r.index == 0)
 
-    assert paris_result.relevance_score > brazil_result.relevance_score
+        assert paris_result.relevance_score > brazil_result.relevance_score
 
+    def test_rerank_top_n(self, server: RemoteOpenAIServer):
+        """Test ColBERT rerank with top_n parameter."""
+        query = "What is the capital of France?"
+        documents = [
+            "The capital of Brazil is Brasilia.",
+            "The capital of France is Paris.",
+            "Machine learning is a field of AI.",
+        ]
 
-@pytest.mark.parametrize("model_name", [MODEL_NAME])
-def test_colbert_rerank_top_n(server: RemoteOpenAIServer, model_name: str):
-    """Test ColBERT rerank with top_n parameter."""
-    query = "What is the capital of France?"
-    documents = [
-        "The capital of Brazil is Brasilia.",
-        "The capital of France is Paris.",
-        "Machine learning is a field of AI.",
-    ]
+        rerank_response = requests.post(
+            server.url_for("rerank"),
+            json={
+                "model": MODEL_NAME,
+                "query": query,
+                "documents": documents,
+                "top_n": 2,
+            },
+        )
+        rerank_response.raise_for_status()
+        rerank = RerankResponse.model_validate(rerank_response.json())
 
-    rerank_response = requests.post(
-        server.url_for("rerank"),
-        json={
-            "model": model_name,
-            "query": query,
-            "documents": documents,
-            "top_n": 2,
-        },
-    )
-    rerank_response.raise_for_status()
-    rerank = RerankResponse.model_validate(rerank_response.json())
+        assert len(rerank.results) == 2
+        assert rerank.results[0].index == 1
 
-    assert len(rerank.results) == 2
-    # Top result should be about Paris (index 1)
-    assert rerank.results[0].index == 1
+    def test_score(self, server: RemoteOpenAIServer):
+        """Test ColBERT score endpoint."""
+        text_1 = "What is the capital of France?"
+        text_2 = ["The capital of France is Paris.", "Python is a language."]
 
+        score_response = requests.post(
+            server.url_for("score"),
+            json={
+                "model": MODEL_NAME,
+                "text_1": text_1,
+                "text_2": text_2,
+            },
+        )
+        score_response.raise_for_status()
+        score = ScoreResponse.model_validate(score_response.json())
 
-@pytest.mark.parametrize("model_name", [MODEL_NAME])
-def test_colbert_score(server: RemoteOpenAIServer, model_name: str):
-    """Test ColBERT score endpoint."""
-    text_1 = "What is the capital of France?"
-    text_2 = ["The capital of France is Paris.", "Python is a language."]
+        assert score.id is not None
+        assert score.data is not None
+        assert len(score.data) == 2
 
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": model_name,
-            "text_1": text_1,
-            "text_2": text_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
+        assert score.data[0].score > score.data[1].score
 
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
+    def test_token_embed(self, server: RemoteOpenAIServer):
+        """Test ColBERT token_embed task via pooling endpoint."""
+        text = "What is the capital of France?"
 
-    # The relevant document should have higher score
-    assert score.data[0].score > score.data[1].score
+        pooling_response = requests.post(
+            server.url_for("pooling"),
+            json={
+                "model": MODEL_NAME,
+                "input": text,
+                "task": "token_embed",
+            },
+        )
+        pooling_response.raise_for_status()
+        pooling = pooling_response.json()
 
+        assert "data" in pooling
+        assert len(pooling["data"]) == 1
 
-@pytest.mark.parametrize("model_name", [MODEL_NAME])
-def test_colbert_token_embed(server: RemoteOpenAIServer, model_name: str):
-    """Test ColBERT token_embed task via pooling endpoint."""
-    text = "What is the capital of France?"
+        embeddings = pooling["data"][0]["data"]
+        assert isinstance(embeddings, list)
+        assert len(embeddings) > 0
+        assert len(embeddings[0]) == COLBERT_DIM
 
-    pooling_response = requests.post(
-        server.url_for("pooling"),
-        json={
-            "model": model_name,
-            "input": text,
-            "task": "token_embed",
-        },
-    )
-    pooling_response.raise_for_status()
-    pooling = pooling_response.json()
+    def test_embed_not_supported(self, server: RemoteOpenAIServer):
+        """Test that ColBERT model does not support 'embed' task."""
+        task = "embed"
+        text = "What is the capital of France?"
 
-    assert "data" in pooling
-    assert len(pooling["data"]) == 1
+        response = requests.post(
+            server.url_for("pooling"),
+            json={
+                "model": MODEL_NAME,
+                "input": text,
+                "task": task,
+            },
+        )
 
-    # Token embeddings should be 2D
-    embeddings = pooling["data"][0]["data"]
-    assert isinstance(embeddings, list)
-    assert len(embeddings) > 0  # Should have tokens
-    assert len(embeddings[0]) == COLBERT_DIM
-
-
-@pytest.mark.parametrize("model_name", [MODEL_NAME])
-def test_colbert_embed_not_supported(server: RemoteOpenAIServer, model_name: str):
-    """Test that ColBERT model does not support 'embed' task."""
-    task = "embed"
-    text = "What is the capital of France?"
-
-    response = requests.post(
-        server.url_for("pooling"),
-        json={
-            "model": model_name,
-            "input": text,
-            "task": task,
-        },
-    )
-
-    assert response.json()["error"]["type"] == "BadRequestError"
-    assert response.json()["error"]["message"].startswith(f"Unsupported task: {task!r}")
+        assert response.json()["error"]["type"] == "BadRequestError"
+        assert response.json()["error"]["message"].startswith(
+            f"Unsupported task: {task!r}"
+        )

tests/kernels/core/test_mrope.py

@@ -4,8 +4,6 @@ from typing import NamedTuple
 
 import pytest
 import torch
-from packaging.version import Version
-from transformers import __version__ as TRANSFORMERS_VERSION
 
 from vllm.model_executor.layers.rotary_embedding import get_rope
 from vllm.platforms import current_platform
@@ -46,31 +44,13 @@ class MRoPETestInfo(NamedTuple):
     marks: list[pytest.MarkDecorator] = []
 
 
-TRANSFORMERS_BASE_VERSION = Version(TRANSFORMERS_VERSION).base_version
-
 MODELS_TO_TEST = [
     MRoPETestInfo(model_name="zai-org/GLM-4.1V-9B-Thinking"),
     MRoPETestInfo(model_name="Qwen/Qwen2-VL-7B-Instruct"),
     MRoPETestInfo(model_name="Qwen/Qwen2-VL-72B-Instruct"),
     MRoPETestInfo(model_name="Qwen/Qwen2.5-VL-72B-Instruct"),
-    MRoPETestInfo(
-        model_name="Qwen/Qwen3-VL-4B-Instruct",
-        marks=[
-            pytest.mark.skipif(
-                Version(TRANSFORMERS_BASE_VERSION) < Version("4.57.0"),
-                reason="Qwen3-VL only available after Transformers v4.57",
-            )
-        ],
-    ),
-    MRoPETestInfo(
-        model_name="Qwen/Qwen3-VL-30B-A3B-Instruct",
-        marks=[
-            pytest.mark.skipif(
-                Version(TRANSFORMERS_BASE_VERSION) < Version("4.57.0"),
-                reason="Qwen3-VL only available after Transformers v4.57",
-            )
-        ],
-    ),
+    MRoPETestInfo(model_name="Qwen/Qwen3-VL-4B-Instruct"),
+    MRoPETestInfo(model_name="Qwen/Qwen3-VL-30B-A3B-Instruct"),
 ]
 
 num_tokens_list = [11, 8192]

tests/kernels/helion/test_register.py

@@ -554,10 +554,18 @@ class TestKernelRegistry:
     """Test suite for kernel registry functionality."""
 
     def setup_method(self):
-        """Clear the registry before each test."""
+        """Save and clear the registry before each test."""
+        from vllm.kernels.helion.register import _REGISTERED_KERNELS
+
+        self._saved_registry = dict(_REGISTERED_KERNELS)
+        _REGISTERED_KERNELS.clear()
+
+    def teardown_method(self):
+        """Restore the registry after each test."""
         from vllm.kernels.helion.register import _REGISTERED_KERNELS
 
         _REGISTERED_KERNELS.clear()
+        _REGISTERED_KERNELS.update(self._saved_registry)
 
     def test_get_registered_kernels_returns_copy(self):
         """Test get_registered_kernels returns copy of registry."""

tests/kernels/helion/test_silu_mul_fp8.py

@@ -0,0 +1,331 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import torch
+import torch.nn.functional as F
+
+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,
+    )
+
+from vllm.kernels.helion.config_manager import ConfigManager
+from vllm.kernels.helion.ops.silu_mul_fp8 import (
+    pick_silu_mul_fp8_config,
+    silu_mul_fp8,
+    silu_mul_fp8_baseline,
+)
+
+
+def skip_if_platform_unsupported():
+    try:
+        from vllm.kernels.helion.utils import get_canonical_gpu_name
+
+        if not torch.cuda.is_available():
+            pytest.skip("CUDA not available")
+
+        platform = get_canonical_gpu_name()
+
+        try:
+            config_manager = ConfigManager.get_instance()
+        except RuntimeError:
+            config_manager = ConfigManager()
+
+        configs = config_manager.get_platform_configs("silu_mul_fp8", platform)
+        if len(configs) == 0:
+            pytest.skip("Current GPU platform not supported for silu_mul_fp8 kernel")
+
+    except (ImportError, RuntimeError, KeyError):
+        pytest.skip("Error detecting platform support for silu_mul_fp8 kernel")
+
+
+@pytest.fixture(autouse=True)
+def reset_config_manager_singleton():
+    ConfigManager.reset_instance()
+    ConfigManager()
+    yield
+    ConfigManager.reset_instance()
+
+
+class TestSiluMulFp8ConfigPicker:
+    def test_config_picker_exact_match(self):
+        config_keys = [
+            "intermediate_2048_batchsize_256",
+            "intermediate_4096_batchsize_256",
+        ]
+
+        input_tensor = torch.randn(32, 4096, dtype=torch.bfloat16, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+        args = (input_tensor, scale)
+
+        selected_key = pick_silu_mul_fp8_config(args, config_keys)
+        assert selected_key == "intermediate_2048_batchsize_256"
+
+    def test_config_picker_closest_match(self):
+        config_keys = [
+            "intermediate_2048_batchsize_256",
+            "intermediate_4096_batchsize_256",
+        ]
+        # Use 7000 (intermediate_size=3500) which is closer to 4096 than 2048
+        input_tensor = torch.randn(32, 7000, dtype=torch.bfloat16, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+        args = (input_tensor, scale)
+
+        selected_key = pick_silu_mul_fp8_config(args, config_keys)
+        assert selected_key == "intermediate_4096_batchsize_256"
+
+    def test_config_picker_fallback_to_default(self):
+        config_keys = ["default", "some_other_key"]
+
+        input_tensor = torch.randn(32, 4096, dtype=torch.bfloat16, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+        args = (input_tensor, scale)
+
+        selected_key = pick_silu_mul_fp8_config(args, config_keys)
+        assert selected_key == "default"
+
+    def test_config_picker_no_configs(self):
+        config_keys: list[str] = []
+
+        input_tensor = torch.randn(32, 4096, dtype=torch.bfloat16, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+        args = (input_tensor, scale)
+
+        selected_key = pick_silu_mul_fp8_config(args, config_keys)
+        assert selected_key is None
+
+    @pytest.mark.parametrize("intermediate_size", [2048, 4096, 5120])
+    def test_config_picker_different_sizes(self, intermediate_size):
+        config_keys = [
+            "intermediate_2048_batchsize_256",
+            "intermediate_4096_batchsize_256",
+            "intermediate_5120_batchsize_256",
+        ]
+
+        input_tensor = torch.randn(
+            32, 2 * intermediate_size, dtype=torch.bfloat16, device="cuda"
+        )
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+        args = (input_tensor, scale)
+
+        selected_key = pick_silu_mul_fp8_config(args, config_keys)
+        expected_key = f"intermediate_{intermediate_size}_batchsize_256"
+        assert selected_key == expected_key
+
+
+class TestSiluMulFp8Correctness:
+    @pytest.mark.parametrize("batch_size", [1, 8, 32, 128])
+    @pytest.mark.parametrize("intermediate_size", [2048, 3000, 3500, 4096, 5000])
+    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
+    def test_silu_mul_fp8_correctness(self, batch_size, intermediate_size, dtype):
+        skip_if_platform_unsupported()
+
+        input_size = 2 * intermediate_size
+        input_tensor = torch.randn(batch_size, input_size, dtype=dtype, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+
+        reference_output = silu_mul_fp8_baseline(input_tensor, scale)
+        helion_output = silu_mul_fp8(input_tensor, scale)
+
+        assert helion_output.shape == reference_output.shape
+        assert helion_output.dtype == torch.float8_e4m3fn
+        assert reference_output.dtype == torch.float8_e4m3fn
+
+        ref_f32 = reference_output.to(torch.float32)
+        helion_f32 = helion_output.to(torch.float32)
+        # FP8 E4M3 has limited precision. Values near quantization boundaries
+        # can round differently due to intermediate precision differences.
+        torch.testing.assert_close(
+            helion_f32,
+            ref_f32,
+            atol=0.05,
+            rtol=0.05,
+            msg=f"Mismatch at batch={batch_size}, size={intermediate_size}",
+        )
+
+    def test_silu_mul_fp8_shape_inference(self):
+        skip_if_platform_unsupported()
+        batch_size, input_size = 32, 8192
+        intermediate_size = input_size // 2
+
+        input_tensor = torch.randn(
+            batch_size, input_size, dtype=torch.bfloat16, device="cuda"
+        )
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+
+        output = silu_mul_fp8(input_tensor, scale)
+
+        expected_shape = (batch_size, intermediate_size)
+        assert output.shape == expected_shape
+        assert output.dtype == torch.float8_e4m3fn
+
+    def test_silu_mul_fp8_scale_variations(self):
+        skip_if_platform_unsupported()
+        batch_size, input_size = 16, 4096
+
+        input_tensor = torch.randn(
+            batch_size, input_size, dtype=torch.bfloat16, device="cuda"
+        )
+
+        scales = [0.1, 0.5, 1.0, 2.0, 10.0]
+
+        for scale_val in scales:
+            scale = torch.tensor([scale_val], dtype=torch.float32, device="cuda")
+
+            reference_output = silu_mul_fp8_baseline(input_tensor, scale)
+            helion_output = silu_mul_fp8(input_tensor, scale)
+            ref_f32 = reference_output.to(torch.float32)
+            helion_f32 = helion_output.to(torch.float32)
+
+            torch.testing.assert_close(
+                helion_f32,
+                ref_f32,
+                atol=0.05,
+                rtol=0.05,
+                msg=f"Mismatch for scale={scale_val}",
+            )
+
+    @pytest.mark.parametrize(
+        "shape",
+        [
+            (1, 4096),
+            (16, 4096),
+            (128, 4096),
+            (1024, 4096),
+            (1, 8192),
+            (16, 8192),
+            (128, 8192),
+        ],
+    )
+    def test_silu_mul_fp8_various_shapes(self, shape):
+        skip_if_platform_unsupported()
+
+        input_tensor = torch.randn(*shape, dtype=torch.bfloat16, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+
+        reference_output = silu_mul_fp8_baseline(input_tensor, scale)
+        helion_output = silu_mul_fp8(input_tensor, scale)
+
+        assert helion_output.shape == reference_output.shape
+
+        ref_f32 = reference_output.to(torch.float32)
+        helion_f32 = helion_output.to(torch.float32)
+
+        torch.testing.assert_close(
+            helion_f32, ref_f32, atol=0.05, rtol=0.05, msg=f"Mismatch for shape={shape}"
+        )
+
+
+def silu_mul_fp8_pytorch(input: torch.Tensor, scale: torch.Tensor) -> torch.Tensor:
+    """Pure PyTorch reference using F.silu.
+
+    This matches vLLM's SiluAndMul.forward_native exactly:
+    F.silu(x[..., :d]) * x[..., d:]
+    """
+    d = input.shape[-1] // 2
+    result = F.silu(input[..., :d]) * input[..., d:]
+    return (result.to(torch.float32) / scale).to(torch.float8_e4m3fn)
+
+
+class TestSiluMulFp8PytorchReference:
+    """Tests comparing Helion kernel against pure PyTorch implementation.
+
+    Uses tighter tolerance since both use PyTorch's FP8 conversion
+    (same rounding mode), unlike the vLLM C++ baseline which uses
+    NVIDIA's hardware FP8 conversion with different rounding.
+    """
+
+    @pytest.mark.parametrize("batch_size", [1, 8, 32, 128, 256])
+    @pytest.mark.parametrize("intermediate_size", [1024, 2048, 4096])
+    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16])
+    def test_silu_mul_fp8_vs_pytorch(self, batch_size, intermediate_size, dtype):
+        skip_if_platform_unsupported()
+
+        input_tensor = torch.randn(
+            batch_size, 2 * intermediate_size, dtype=dtype, device="cuda"
+        )
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+
+        pytorch_output = silu_mul_fp8_pytorch(input_tensor, scale)
+        helion_output = silu_mul_fp8(input_tensor, scale)
+
+        assert helion_output.shape == pytorch_output.shape
+        assert helion_output.dtype == torch.float8_e4m3fn
+
+        pytorch_f32 = pytorch_output.to(torch.float32)
+        helion_f32 = helion_output.to(torch.float32)
+
+        # Tolerance accounts for FP8 quantization boundary effects
+        torch.testing.assert_close(
+            helion_f32,
+            pytorch_f32,
+            atol=0.05,
+            rtol=0.05,
+            msg=(
+                f"Mismatch at batch={batch_size}, size={intermediate_size}, "
+                f"dtype={dtype}"
+            ),
+        )
+
+    @pytest.mark.parametrize(
+        "shape",
+        [
+            (1, 2, 4096),  # 3D input
+            (2, 4, 2048),  # 3D input
+            (1, 1, 1, 8192),  # 4D input
+        ],
+    )
+    def test_silu_mul_fp8_multidim_vs_pytorch(self, shape):
+        skip_if_platform_unsupported()
+
+        input_tensor = torch.randn(*shape, dtype=torch.bfloat16, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+
+        pytorch_output = silu_mul_fp8_pytorch(input_tensor, scale)
+        helion_output = silu_mul_fp8(input_tensor, scale)
+
+        assert helion_output.shape == pytorch_output.shape
+
+        pytorch_f32 = pytorch_output.to(torch.float32)
+        helion_f32 = helion_output.to(torch.float32)
+
+        torch.testing.assert_close(
+            helion_f32,
+            pytorch_f32,
+            atol=0.05,
+            rtol=0.05,
+            msg=f"Mismatch for shape={shape}",
+        )
+
+
+class TestSiluMulFp8Integration:
+    def test_kernel_registration_integration(self):
+        from vllm.kernels.helion.register import get_registered_kernels
+
+        registered_kernels = get_registered_kernels()
+        assert "silu_mul_fp8" in registered_kernels
+
+        kernel_wrapper = registered_kernels["silu_mul_fp8"]
+        assert kernel_wrapper.op_name == "silu_mul_fp8"
+        assert kernel_wrapper._config_picker is not None
+
+    def test_fake_impl_functionality(self):
+        skip_if_platform_unsupported()
+        from vllm.kernels.helion.register import get_registered_kernels
+
+        input_tensor = torch.randn(32, 4096, dtype=torch.bfloat16, device="cuda")
+        scale = torch.tensor([0.5], dtype=torch.float32, device="cuda")
+        registered_kernels = get_registered_kernels()
+        kernel_wrapper = registered_kernels["silu_mul_fp8"]
+        fake_impl = kernel_wrapper._fake_impl
+
+        fake_output = fake_impl(input_tensor, scale)
+
+        expected_shape = (32, 2048)
+        assert fake_output.shape == expected_shape
+        assert fake_output.dtype == torch.float8_e4m3fn
+        assert fake_output.device == input_tensor.device

tests/kernels/moe/test_flashinfer.py

@@ -71,7 +71,8 @@ def quant_fp8_per_tensor_batches(a):
 
     for i in range(num_batches):
         a_fp8, a_global_sf = input_to_float8(a[i])
-        a_global_sf = 1.0 / a_global_sf
+        if a_global_sf.numel() == 1:
+            a_global_sf = a_global_sf.view(1, 1)
         a_quant.append(a_fp8)
         a_scales.append(a_global_sf)
 
@@ -81,6 +82,20 @@ def quant_fp8_per_tensor_batches(a):
     return result_a_quant, result_a_scales
 
 
+def check_accuracy(ref_output, actual_output, atol=0.1, rtol=0.85, percent=0.925):
+    close = torch.isclose(ref_output, actual_output, atol=atol, rtol=rtol)
+    match_ratio = close.float().mean()
+    assert match_ratio >= percent, (
+        f"Match ratio {match_ratio:.4f} is below the threshold {percent:.4f}"
+    )
+
+    mismatch_percent = 1.0 - match_ratio.item()
+    assert mismatch_percent <= 1 - percent, (
+        f"Mismatch percentage {mismatch_percent:.4f} is above the threshold "
+        f"{1 - percent:.4f}"
+    )
+
+
 @dataclass
 class TestData:
     hidden_states: torch.Tensor
@@ -104,14 +119,16 @@ class TestData:
         is_gated = activation.is_gated
 
         hidden_states = torch.randn((m, k), device="cuda", dtype=torch.bfloat16) / 10
-        w13 = torch.randn(
-            (e, (2 * n) if is_gated else n, k), device="cuda", dtype=torch.bfloat16
+        w13 = (
+            torch.randn(
+                (e, (2 * n) if is_gated else n, k), device="cuda", dtype=torch.bfloat16
+            )
+            / 10
         )
-        w2 = torch.randn((e, k, n), device="cuda", dtype=torch.bfloat16)
+        w2 = torch.randn((e, k, n), device="cuda", dtype=torch.bfloat16) / 10
 
         # Scale to fp8
         _, a1_scale = input_to_float8(hidden_states)
-        a1_scale = 1.0 / a1_scale
         a2_scale = torch.scalar_tensor(1.0).to(device="cuda").to(dtype=torch.float32)
         w13_quantized, w13_weight_scale = quant_fp8_per_tensor_batches(w13)
         w2_quantized, w2_weight_scale = quant_fp8_per_tensor_batches(w2)
@@ -124,14 +141,16 @@ class TestData:
         layer.w2_input_scale = a2_scale
         layer.w13_weight_scale = w13_weight_scale
         layer.w2_weight_scale = w2_weight_scale
+        layer.activation = activation
         # Setup dummy config.
         layer.moe_parallel_config = mk.FusedMoEParallelConfig.make_no_parallel()
 
         # flashinfer expects swapped rows for w13
-        layer.w13_weight.data = swap_w13_to_w31(layer.w13_weight.data)
+        if is_gated:
+            layer.w13_weight.data = swap_w13_to_w31(layer.w13_weight.data)
         if is_trtllm:
             rotate_weights_for_fi_trtllm_fp8_per_tensor_moe(
-                layer.w13_weight, layer.w2_weight
+                layer.w13_weight, layer.w2_weight, is_gated
             )
             register_scales_for_trtllm_fp8_per_tensor_moe(
                 layer,
@@ -162,12 +181,14 @@ class TestData:
 @pytest.mark.parametrize("m,n,k", MNK_FACTORS)
 @pytest.mark.parametrize("e", NUM_EXPERTS)
 @pytest.mark.parametrize("topk", TOP_KS)
+@pytest.mark.parametrize("activation", [MoEActivation.SILU, MoEActivation.RELU2_NO_MUL])
 def test_flashinfer_per_tensor_moe_fp8_no_graph(
     m: int,
     n: int,
     k: int,
     e: int,
     topk: int,
+    activation: MoEActivation,
     monkeypatch,
 ):
     if not current_platform.has_device_capability(100):
@@ -175,7 +196,9 @@ def test_flashinfer_per_tensor_moe_fp8_no_graph(
     set_random_seed(7)
     monkeypatch.setenv("VLLM_FUSED_MOE_CHUNK_SIZE", "8192")
     with set_current_vllm_config(vllm_config):
-        td = TestData.make_moe_tensors_8bit(m, k, n, e, is_trtllm=True)
+        td = TestData.make_moe_tensors_8bit(
+            m, k, n, e, is_trtllm=True, activation=activation
+        )
 
         score = torch.randn((m, e), device="cuda", dtype=torch.bfloat16)
         topk_weights, topk_ids = Llama4MoE.custom_routing_function(
@@ -200,7 +223,7 @@ def test_flashinfer_per_tensor_moe_fp8_no_graph(
             topk_weights=topk_weights,
             topk_ids=topk_ids,
             inplace=False,
-            activation=MoEActivation.SILU,
+            activation=activation,
             global_num_experts=e,
             expert_map=None,
             apply_router_weight_on_input=True,
@@ -219,7 +242,13 @@ def test_flashinfer_per_tensor_moe_fp8_no_graph(
             apply_router_weight_on_input=True,
         )
 
-        torch.testing.assert_close(output, flashinfer_output, atol=5.5e-2, rtol=1e-2)
+        check_accuracy(
+            ref_output=output,
+            actual_output=flashinfer_output,
+            atol=0.1,
+            rtol=0.85,
+            percent=0.925,
+        )
 
 
 @pytest.mark.parametrize("m,n,k", MNK_FACTORS)
@@ -320,8 +349,13 @@ def test_flashinfer_cutlass_moe_fp8_no_graph(
             expert_map=None,
             apply_router_weight_on_input=True,
         )
-        torch.testing.assert_close(
-            output, flashinfer_cutlass_output, atol=5.5e-2, rtol=1e-2
+
+        check_accuracy(
+            ref_output=output,
+            actual_output=flashinfer_cutlass_output,
+            atol=0.1,
+            rtol=0.85,
+            percent=0.925,
         )
 
 
@@ -364,3 +398,80 @@ def test_convert_moe_weights_to_flashinfer_trtllm_block_layout(
 
     assert w13_converted.shape[0] == num_experts
     assert w2_converted.shape[0] == num_experts
+
+
+def test_flashinfer_blockscale_fp8_none_expert_group(monkeypatch):
+    """Test that flashinfer_fused_moe_blockscale_fp8 handles num_expert_group=None.
+
+    Regression test for https://github.com/vllm-project/vllm/issues/34477
+    MiniMax-M2.1 uses sigmoid scoring with e_score_correction_bias but no
+    grouped top-k, resulting in num_expert_group=None. This triggered a crash
+    in the flashinfer kernel when DeepSeekV3 routing was selected.
+    """
+    if not current_platform.has_device_capability(100):
+        pytest.skip("Test requires SM >= 100 (Blackwell)")
+
+    import vllm.model_executor.layers.fused_moe.flashinfer_trtllm_moe  # noqa: E501, F401
+    from tests.kernels.quant_utils import native_per_token_group_quant_fp8
+
+    set_random_seed(7)
+    monkeypatch.setenv("VLLM_FUSED_MOE_CHUNK_SIZE", "8192")
+
+    e = 16  # num_experts (must be divisible by 4)
+    topk = 6  # top_k > 1 triggers DeepSeekV3 routing with sigmoid
+    m, n, k = 10, 4096, 5120
+    block_shape = [128, 128]
+    block_k = block_shape[1]
+
+    with set_current_vllm_config(vllm_config):
+        # Create BF16 hidden states
+        x = torch.randn((m, k), device="cuda", dtype=torch.bfloat16) / 10
+
+        # Create FP8 block-scale quantized weights
+        w13_bf16 = torch.randn((e, 2 * n, k), device="cuda", dtype=torch.bfloat16) / 10
+        w2_bf16 = torch.randn((e, k, n), device="cuda", dtype=torch.bfloat16) / 10
+
+        # Quantize weights per-block to FP8
+        w13_fp8_list, w13_scale_list = [], []
+        w2_fp8_list, w2_scale_list = [], []
+        for i in range(e):
+            wq, ws = native_per_token_group_quant_fp8(w13_bf16[i], block_k)
+            w13_fp8_list.append(wq)
+            w13_scale_list.append(ws)
+
+            wq, ws = native_per_token_group_quant_fp8(w2_bf16[i], block_k)
+            w2_fp8_list.append(wq)
+            w2_scale_list.append(ws)
+
+        w13_fp8 = torch.stack(w13_fp8_list)
+        w13_scale = torch.stack(w13_scale_list)
+        w2_fp8 = torch.stack(w2_fp8_list)
+        w2_scale = torch.stack(w2_scale_list)
+
+        # DeepSeekV3 routing uses float32 logits + optional bias
+        routing_logits = torch.randn((m, e), device="cuda", dtype=torch.float32)
+        routing_bias = torch.randn(e, device="cuda", dtype=torch.float32)
+
+        # This should NOT crash with num_expert_group=None
+        output = torch.ops.vllm.flashinfer_fused_moe_blockscale_fp8(
+            routing_logits=routing_logits,
+            routing_bias=routing_bias,
+            x=x,
+            w13_weight=w13_fp8,
+            w13_weight_scale_inv=w13_scale,
+            w2_weight=w2_fp8,
+            w2_weight_scale_inv=w2_scale,
+            global_num_experts=e,
+            top_k=topk,
+            num_expert_group=None,
+            topk_group=None,
+            intermediate_size=n,
+            expert_offset=0,
+            local_num_experts=e,
+            block_shape=block_shape,
+            routing_method_type=RoutingMethodType.DeepSeekV3,
+            routed_scaling=1.0,
+        )
+
+        assert output is not None
+        assert output.shape == (m, k)

tests/kernels/moe/utils.py

@@ -7,11 +7,6 @@ import vllm._custom_ops as ops
 from tests.kernels.quant_utils import per_block_cast_to_int8
 from tests.kernels.quantization.nvfp4_utils import FLOAT4_E2M1_MAX, FLOAT8_E4M3_MAX
 from vllm.model_executor.layers.activation import SiluAndMul
-from vllm.model_executor.layers.fused_moe import (
-    TritonExperts,
-    fused_experts,
-    fused_topk,
-)
 from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
@@ -24,10 +19,15 @@ from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
     BatchedTritonExperts,
     NaiveBatchedExperts,
 )
+from vllm.model_executor.layers.fused_moe.fused_moe import (
+    TritonExperts,
+    fused_experts,
+)
 from vllm.model_executor.layers.fused_moe.modular_kernel import FusedMoEModularKernel
 from vllm.model_executor.layers.fused_moe.prepare_finalize import (
     MoEPrepareAndFinalizeNoEP,
 )
+from vllm.model_executor.layers.fused_moe.router.fused_topk_router import fused_topk
 from vllm.model_executor.layers.fused_moe.utils import moe_kernel_quantize_input
 from vllm.utils.deep_gemm import per_block_cast_to_fp8
 from vllm.utils.math_utils import round_up

tests/models/language/pooling/test_colbert.py

@@ -1,16 +1,47 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Tests for ColBERT late interaction scoring."""
+"""Tests for ColBERT late interaction scoring.
+
+Tests are parametrized across multiple ColBERT backbones to ensure the
+generic ColBERT support works with different encoder architectures.
+"""
 
 import pytest
 import torch
 
 from vllm.entrypoints.pooling.score.utils import compute_maxsim_score
 
-# ColBERT model - using answerai-colbert-small-v1 as it's a smaller model
-# suitable for testing (based on BERT-base)
-COLBERT_MODEL = "answerdotai/answerai-colbert-small-v1"
-COLBERT_DIM = 96  # This model uses 96-dimensional output
+# -----------------------------------------------------------------------
+# Model definitions: (model_name, colbert_dim, extra vllm_runner kwargs)
+# -----------------------------------------------------------------------
+COLBERT_MODELS = {
+    "bert": {
+        "model": "answerdotai/answerai-colbert-small-v1",
+        "colbert_dim": 96,
+        "max_model_len": 512,
+        "extra_kwargs": {},
+    },
+    "modernbert": {
+        "model": "lightonai/GTE-ModernColBERT-v1",
+        "colbert_dim": 128,
+        "max_model_len": 299,
+        "extra_kwargs": {
+            "hf_overrides": {
+                "architectures": ["ColBERTModernBertModel"],
+            },
+        },
+    },
+    "jina": {
+        "model": "jinaai/jina-colbert-v2",
+        "colbert_dim": 128,
+        "max_model_len": 8192,
+        "extra_kwargs": {
+            "hf_overrides": {
+                "architectures": ["ColBERTJinaRobertaModel"],
+            },
+        },
+    },
+}
 
 TEXTS_1 = [
     "What is the capital of France?",
@@ -25,80 +56,121 @@ TEXTS_2 = [
 DTYPE = "half"
 
 
+# -----------------------------------------------------------------------
+# Fixtures
+# -----------------------------------------------------------------------
+
+
+@pytest.fixture(params=list(COLBERT_MODELS.keys()), scope="module")
+def colbert_spec(request):
+    """Return the model spec dict for the current parametrization."""
+    return COLBERT_MODELS[request.param]
+
+
 @pytest.fixture(scope="module")
-def colbert_model_name():
-    return COLBERT_MODEL
+def colbert_model_name(colbert_spec):
+    return colbert_spec["model"]
 
 
-def test_colbert_token_embed(vllm_runner, colbert_model_name):
+@pytest.fixture(scope="module")
+def colbert_dim(colbert_spec):
+    return colbert_spec["colbert_dim"]
+
+
+@pytest.fixture(scope="module")
+def colbert_max_model_len(colbert_spec):
+    return colbert_spec["max_model_len"]
+
+
+@pytest.fixture(scope="module")
+def colbert_extra_kwargs(colbert_spec):
+    return colbert_spec["extra_kwargs"]
+
+
+# -----------------------------------------------------------------------
+# Tests
+# -----------------------------------------------------------------------
+
+
+def test_colbert_token_embed(
+    vllm_runner,
+    colbert_model_name,
+    colbert_dim,
+    colbert_max_model_len,
+    colbert_extra_kwargs,
+):
     """Test that ColBERT model produces token embeddings."""
     with vllm_runner(
         colbert_model_name,
         runner="pooling",
         dtype=DTYPE,
-        max_model_len=512,
+        max_model_len=colbert_max_model_len,
         enforce_eager=True,
+        **colbert_extra_kwargs,
     ) as vllm_model:
-        # Get token embeddings for a single text
         outputs = vllm_model.token_embed([TEXTS_1[0]])
 
         assert len(outputs) == 1
-        # Token embeddings should be 2D: [num_tokens, colbert_dim]
         emb = torch.tensor(outputs[0])
         assert emb.dim() == 2
-        assert emb.shape[1] == COLBERT_DIM
-        # Should have at least a few tokens
+        assert emb.shape[1] == colbert_dim
         assert emb.shape[0] > 1
 
 
-def test_colbert_late_interaction_1_to_1(vllm_runner, colbert_model_name):
+def test_colbert_late_interaction_1_to_1(
+    vllm_runner,
+    colbert_model_name,
+    colbert_max_model_len,
+    colbert_extra_kwargs,
+):
     """Test ColBERT late interaction scoring with 1:1 query-document pair."""
     with vllm_runner(
         colbert_model_name,
         runner="pooling",
         dtype=DTYPE,
-        max_model_len=512,
+        max_model_len=colbert_max_model_len,
         enforce_eager=True,
+        **colbert_extra_kwargs,
     ) as vllm_model:
-        # Get token embeddings
         q_outputs = vllm_model.token_embed([TEXTS_1[0]])
         d_outputs = vllm_model.token_embed([TEXTS_2[0]])
 
         q_emb = torch.tensor(q_outputs[0])
         d_emb = torch.tensor(d_outputs[0])
 
-        # Compute MaxSim manually
         manual_score = compute_maxsim_score(q_emb, d_emb).item()
 
-        # Use the score API (which should internally use _late_interaction_score)
         vllm_scores = vllm_model.score(TEXTS_1[0], TEXTS_2[0])
 
         assert len(vllm_scores) == 1
         assert vllm_scores[0] == pytest.approx(manual_score, rel=0.01)
 
 
-def test_colbert_late_interaction_1_to_N(vllm_runner, colbert_model_name):
+def test_colbert_late_interaction_1_to_N(
+    vllm_runner,
+    colbert_model_name,
+    colbert_max_model_len,
+    colbert_extra_kwargs,
+):
     """Test ColBERT late interaction scoring with 1:N query-documents."""
     with vllm_runner(
         colbert_model_name,
         runner="pooling",
         dtype=DTYPE,
-        max_model_len=512,
+        max_model_len=colbert_max_model_len,
         enforce_eager=True,
+        **colbert_extra_kwargs,
     ) as vllm_model:
-        # Get token embeddings
         q_outputs = vllm_model.token_embed([TEXTS_1[0]])
         d_outputs = vllm_model.token_embed(TEXTS_2)
 
         q_emb = torch.tensor(q_outputs[0])
 
-        # Compute MaxSim manually for each document
         manual_scores = []
         for d_out in d_outputs:
             d_emb = torch.tensor(d_out)
             manual_scores.append(compute_maxsim_score(q_emb, d_emb).item())
 
-        # Use the score API
         vllm_scores = vllm_model.score(TEXTS_1[0], TEXTS_2)
 
         assert len(vllm_scores) == 2
@@ -106,27 +178,30 @@ def test_colbert_late_interaction_1_to_N(vllm_runner, colbert_model_name):
             assert vllm_scores[i] == pytest.approx(manual_scores[i], rel=0.01)
 
 
-def test_colbert_late_interaction_N_to_N(vllm_runner, colbert_model_name):
+def test_colbert_late_interaction_N_to_N(
+    vllm_runner,
+    colbert_model_name,
+    colbert_max_model_len,
+    colbert_extra_kwargs,
+):
     """Test ColBERT late interaction scoring with N:N query-documents."""
     with vllm_runner(
         colbert_model_name,
         runner="pooling",
         dtype=DTYPE,
-        max_model_len=512,
+        max_model_len=colbert_max_model_len,
         enforce_eager=True,
+        **colbert_extra_kwargs,
     ) as vllm_model:
-        # Get token embeddings
         q_outputs = vllm_model.token_embed(TEXTS_1)
         d_outputs = vllm_model.token_embed(TEXTS_2)
 
-        # Compute MaxSim manually for each pair
         manual_scores = []
         for q_out, d_out in zip(q_outputs, d_outputs):
             q_emb = torch.tensor(q_out)
             d_emb = torch.tensor(d_out)
             manual_scores.append(compute_maxsim_score(q_emb, d_emb).item())
 
-        # Use the score API
         vllm_scores = vllm_model.score(TEXTS_1, TEXTS_2)
 
         assert len(vllm_scores) == 2
@@ -134,8 +209,13 @@ def test_colbert_late_interaction_N_to_N(vllm_runner, colbert_model_name):
             assert vllm_scores[i] == pytest.approx(manual_scores[i], rel=0.01)
 
 
-def test_colbert_relevance_ordering(vllm_runner, colbert_model_name):
-    """Test that ColBERT scores relevant documents higher than irrelevant ones."""
+def test_colbert_relevance_ordering(
+    vllm_runner,
+    colbert_model_name,
+    colbert_max_model_len,
+    colbert_extra_kwargs,
+):
+    """Test that ColBERT scores relevant documents higher than irrelevant."""
     query = "What is machine learning?"
     documents = [
         "Machine learning is a subset of artificial intelligence.",
@@ -147,48 +227,73 @@ def test_colbert_relevance_ordering(vllm_runner, colbert_model_name):
         colbert_model_name,
         runner="pooling",
         dtype=DTYPE,
-        max_model_len=512,
+        max_model_len=colbert_max_model_len,
         enforce_eager=True,
+        **colbert_extra_kwargs,
     ) as vllm_model:
         scores = vllm_model.score(query, documents)
 
         assert len(scores) == 3
-        # ML-related documents should score higher than unrelated Python doc
-        # Document 0 (ML definition) should be most relevant
-        # Document 2 (Deep learning) should also be relevant
-        # Document 1 (Python) should be least relevant
         assert scores[0] > scores[1], "ML doc should score higher than Python doc"
         assert scores[2] > scores[1], "DL doc should score higher than Python doc"
 
 
-def test_colbert_embed_not_supported(vllm_runner, colbert_model_name):
+def test_colbert_embed_not_supported(
+    vllm_runner,
+    colbert_model_name,
+    colbert_max_model_len,
+    colbert_extra_kwargs,
+):
     """Test that ColBERT model does not support 'embed' task."""
     with (
         vllm_runner(
             colbert_model_name,
             runner="pooling",
             dtype=DTYPE,
-            max_model_len=512,
+            max_model_len=colbert_max_model_len,
             enforce_eager=True,
+            **colbert_extra_kwargs,
         ) as vllm_model,
         pytest.raises(ValueError, match="Embedding API is not supported"),
     ):
         vllm_model.embed([TEXTS_1[0]])
 
 
-def test_colbert_hf_comparison(vllm_runner, colbert_model_name):
-    """Test that vLLM ColBERT produces same embeddings as HuggingFace."""
+# -----------------------------------------------------------------------
+# Per-model HuggingFace comparison tests
+# -----------------------------------------------------------------------
+
+
+def _assert_embeddings_close(vllm_outputs, hf_embeddings):
+    """Assert that vLLM and HuggingFace embeddings match."""
+    for i, (hf_emb, vllm_out) in enumerate(zip(hf_embeddings, vllm_outputs)):
+        vllm_emb = torch.tensor(vllm_out).float()
+
+        assert hf_emb.shape == vllm_emb.shape, (
+            f"Shape mismatch for text {i}: HF {hf_emb.shape} vs vLLM {vllm_emb.shape}"
+        )
+
+        torch.testing.assert_close(
+            vllm_emb,
+            hf_emb,
+            rtol=1e-2,
+            atol=1e-2,
+            msg=f"Embedding mismatch for text {i}",
+        )
+
+
+def test_colbert_hf_comparison_bert(vllm_runner):
+    """Test that vLLM ColBERT produces same embeddings as HuggingFace (BERT)."""
     import torch.nn.functional as F
     from huggingface_hub import hf_hub_download
     from safetensors.torch import load_file
     from transformers import AutoTokenizer, BertModel
 
+    model_name = COLBERT_MODELS["bert"]["model"]
     test_texts = [TEXTS_1[0], TEXTS_2[0]]
 
-    # Get vLLM embeddings first (to avoid GPU memory contention)
-    # Use fp32 to match HuggingFace default precision for fair comparison
     with vllm_runner(
-        colbert_model_name,
+        model_name,
         runner="pooling",
         dtype="float32",
         max_model_len=512,
@@ -196,14 +301,11 @@ def test_colbert_hf_comparison(vllm_runner, colbert_model_name):
     ) as vllm_model:
         vllm_outputs = vllm_model.token_embed(test_texts)
 
-    # Get HuggingFace reference embeddings on CPU
-    # Load the base BERT model and manually apply the ColBERT linear projection
-    hf_tokenizer = AutoTokenizer.from_pretrained(colbert_model_name)
-    hf_bert = BertModel.from_pretrained(colbert_model_name)
+    hf_tokenizer = AutoTokenizer.from_pretrained(model_name)
+    hf_bert = BertModel.from_pretrained(model_name)
     hf_bert.eval()
 
-    # Load the ColBERT linear weights from safetensors
-    weights_path = hf_hub_download(colbert_model_name, filename="model.safetensors")
+    weights_path = hf_hub_download(model_name, filename="model.safetensors")
     weights = load_file(weights_path)
     linear_weight = weights["linear.weight"]  # [96, 384]
 
@@ -212,36 +314,103 @@ def test_colbert_hf_comparison(vllm_runner, colbert_model_name):
         inputs = hf_tokenizer(text, return_tensors="pt")
         with torch.no_grad():
             outputs = hf_bert(**inputs)
-            # Get last hidden state: [1, seq_len, 384]
             hidden_states = outputs.last_hidden_state
-            # Apply ColBERT linear projection: [1, seq_len, 96]
             token_emb = F.linear(hidden_states, linear_weight)
-            # L2 normalize
             token_emb = F.normalize(token_emb, p=2, dim=-1)
             hf_embeddings.append(token_emb.squeeze(0).float())
 
-    # Compare embeddings
-    for i, (hf_emb, vllm_out) in enumerate(zip(hf_embeddings, vllm_outputs)):
-        vllm_emb = torch.tensor(vllm_out).float()
+    _assert_embeddings_close(vllm_outputs, hf_embeddings)
 
-        # Print first few components for debugging
-        print(f"\n=== Text {i}: '{test_texts[i][:30]}...' ===")
-        print(f"HF shape: {hf_emb.shape}, vLLM shape: {vllm_emb.shape}")
-        print(f"HF first token, first 10 dims:   {hf_emb[0, :10].tolist()}")
-        print(f"vLLM first token, first 10 dims: {vllm_emb[0, :10].tolist()}")
-        print(f"HF last token, first 10 dims:    {hf_emb[-1, :10].tolist()}")
-        print(f"vLLM last token, first 10 dims:  {vllm_emb[-1, :10].tolist()}")
 
-        # Should have same shape
-        assert hf_emb.shape == vllm_emb.shape, (
-            f"Shape mismatch for text {i}: HF {hf_emb.shape} vs vLLM {vllm_emb.shape}"
-        )
+def test_colbert_hf_comparison_modernbert(vllm_runner):
+    """Test that vLLM ColBERT produces same embeddings as HuggingFace
+    (ModernBERT)."""
+    import torch.nn.functional as F
+    from huggingface_hub import hf_hub_download
+    from safetensors.torch import load_file
+    from transformers import AutoModel, AutoTokenizer
 
-        # Should have same values (with tolerance for fp16)
-        torch.testing.assert_close(
-            vllm_emb,
-            hf_emb,
-            rtol=1e-2,
-            atol=1e-2,
-            msg=f"Embedding mismatch for text {i}",
-        )
+    spec = COLBERT_MODELS["modernbert"]
+    model_name = spec["model"]
+    test_texts = [TEXTS_1[0], TEXTS_2[0]]
+
+    with vllm_runner(
+        model_name,
+        runner="pooling",
+        dtype="float32",
+        max_model_len=spec["max_model_len"],
+        enforce_eager=True,
+        **spec["extra_kwargs"],
+    ) as vllm_model:
+        vllm_outputs = vllm_model.token_embed(test_texts)
+
+    hf_tokenizer = AutoTokenizer.from_pretrained(model_name)
+    hf_model = AutoModel.from_pretrained(model_name)
+    hf_model.eval()
+
+    # Load projection from sentence-transformers 1_Dense layer
+    dense_path = hf_hub_download(model_name, filename="1_Dense/model.safetensors")
+    dense_weights = load_file(dense_path)
+    linear_weight = dense_weights["linear.weight"]  # [128, 768]
+
+    hf_embeddings = []
+    for text in test_texts:
+        inputs = hf_tokenizer(text, return_tensors="pt")
+        with torch.no_grad():
+            outputs = hf_model(**inputs)
+            hidden_states = outputs.last_hidden_state
+            token_emb = F.linear(hidden_states, linear_weight)
+            token_emb = F.normalize(token_emb, p=2, dim=-1)
+            hf_embeddings.append(token_emb.squeeze(0).float())
+
+    _assert_embeddings_close(vllm_outputs, hf_embeddings)
+
+
+def test_colbert_hf_comparison_jina(vllm_runner):
+    """Test that vLLM ColBERT produces same embeddings as HuggingFace
+    (Jina XLM-RoBERTa)."""
+    import torch.nn.functional as F
+    from huggingface_hub import hf_hub_download
+    from safetensors.torch import load_file
+    from transformers import AutoModel, AutoTokenizer
+
+    spec = COLBERT_MODELS["jina"]
+    model_name = spec["model"]
+    test_texts = [TEXTS_1[0], TEXTS_2[0]]
+
+    with vllm_runner(
+        model_name,
+        runner="pooling",
+        dtype="float32",
+        max_model_len=spec["max_model_len"],
+        enforce_eager=True,
+        **spec["extra_kwargs"],
+    ) as vllm_model:
+        vllm_outputs = vllm_model.token_embed(test_texts)
+
+    hf_tokenizer = AutoTokenizer.from_pretrained(
+        model_name,
+        trust_remote_code=True,
+    )
+    hf_model = AutoModel.from_pretrained(
+        model_name,
+        trust_remote_code=True,
+    )
+    hf_model.eval()
+
+    # Load projection from main checkpoint
+    weights_path = hf_hub_download(model_name, filename="model.safetensors")
+    weights = load_file(weights_path)
+    linear_weight = weights["linear.weight"]  # [128, 1024]
+
+    hf_embeddings = []
+    for text in test_texts:
+        inputs = hf_tokenizer(text, return_tensors="pt")
+        with torch.no_grad():
+            outputs = hf_model(**inputs)
+            hidden_states = outputs.last_hidden_state
+            token_emb = F.linear(hidden_states.float(), linear_weight.float())
+            token_emb = F.normalize(token_emb, p=2, dim=-1)
+            hf_embeddings.append(token_emb.squeeze(0).float())
+
+    _assert_embeddings_close(vllm_outputs, hf_embeddings)

tests/models/multimodal/generation/test_common.py

@@ -961,12 +961,6 @@ VLM_TEST_SETTINGS = {
                 limit_mm_per_prompt={"image": 4},
             )
         ],
-        marks=[
-            pytest.mark.skipif(
-                Version(TRANSFORMERS_VERSION) == Version("4.57.1"),
-                reason="This model is broken in Transformers v4.57.1",
-            )
-        ],
     ),
     # regression test for https://github.com/vllm-project/vllm/issues/15122
     "qwen2_5_vl-windows-attention": VLMTestInfo(

tests/models/multimodal/generation/test_voxtral.py

@@ -4,16 +4,18 @@
 import json
 
 import pytest
-import pytest_asyncio
 from mistral_common.audio import Audio
 from mistral_common.protocol.instruct.chunk import AudioChunk, RawAudio, TextChunk
 from mistral_common.protocol.instruct.messages import UserMessage
+from transformers import VoxtralForConditionalGeneration
 
 from vllm.tokenizers.mistral import MistralTokenizer
 
 from ....conftest import AudioTestAssets
 from ....utils import RemoteOpenAIServer
+from ...utils import check_logprobs_close
 from .test_ultravox import MULTI_AUDIO_PROMPT, run_multi_audio_test
+from .vlm_utils import model_utils
 
 MODEL_NAME = "mistralai/Voxtral-Mini-3B-2507"
 MISTRAL_FORMAT_ARGS = [
@@ -26,40 +28,21 @@ MISTRAL_FORMAT_ARGS = [
 ]
 
 
-@pytest.fixture()
-def server(request, audio_assets: AudioTestAssets):
-    args = [
-        "--enforce-eager",
-        "--limit-mm-per-prompt",
-        json.dumps({"audio": len(audio_assets)}),
-    ] + MISTRAL_FORMAT_ARGS
-
-    with RemoteOpenAIServer(
-        MODEL_NAME, args, env_dict={"VLLM_AUDIO_FETCH_TIMEOUT": "30"}
-    ) as remote_server:
-        yield remote_server
-
-
-@pytest_asyncio.fixture
-async def client(server):
-    async with server.get_async_client() as async_client:
-        yield async_client
-
-
-def _get_prompt(audio_assets, question):
+def _get_prompt(audio_assets: AudioTestAssets, question: str) -> list[int]:
+    """Build a token-ID prompt via mistral_common for vLLM offline inference."""
     tokenizer = MistralTokenizer.from_pretrained(MODEL_NAME)
 
     audios = [
-        Audio.from_file(str(audio_assets[i].get_local_path()), strict=False)
-        for i in range(len(audio_assets))
+        Audio.from_file(str(asset.get_local_path()), strict=False)
+        for asset in audio_assets
     ]
     audio_chunks = [
         AudioChunk(input_audio=RawAudio.from_audio(audio)) for audio in audios
     ]
 
-    text_chunk = TextChunk(text=question)
-    messages = [UserMessage(content=[*audio_chunks, text_chunk]).to_openai()]
-
+    messages = [
+        UserMessage(content=[*audio_chunks, TextChunk(text=question)]).to_openai()
+    ]
     return tokenizer.apply_chat_template(messages=messages)
 
 
@@ -77,7 +60,7 @@ def test_models_with_multiple_audios(
     vllm_prompt = _get_prompt(audio_assets, MULTI_AUDIO_PROMPT)
     run_multi_audio_test(
         vllm_runner,
-        [(vllm_prompt, [audio.audio_and_sample_rate for audio in audio_assets])],
+        [(vllm_prompt, [a.audio_and_sample_rate for a in audio_assets])],  # type: ignore[list-item]
         MODEL_NAME,
         dtype=dtype,
         max_tokens=max_tokens,
@@ -86,30 +69,142 @@ def test_models_with_multiple_audios(
     )
 
 
-@pytest.mark.asyncio
-async def test_online_serving(client, audio_assets: AudioTestAssets):
-    """Exercises online serving with/without chunked prefill enabled."""
+def test_online_serving(vllm_runner, audio_assets: AudioTestAssets):
+    """Two-layer accuracy and serving validation using Mistral format.
 
-    def asset_to_chunk(asset):
+    1. Offline vLLM greedy output (runs first to avoid CUDA fork issues
+       with multiprocessing - see vlm_utils/core.py).
+    2. Online OpenAI-compatible API output must match offline — validates
+       that the serving path (chat template, audio encoding, tokenization)
+       does not corrupt anything.
+
+    Steps run sequentially so each releases the GPU before the next starts.
+    """
+
+    question = f"What's happening in these {len(audio_assets)} audio clips?"
+    max_tokens = 10
+    audio_data = [asset.audio_and_sample_rate for asset in audio_assets]
+
+    vllm_prompt = _get_prompt(audio_assets, question)
+    with vllm_runner(
+        MODEL_NAME,
+        dtype="half",
+        enforce_eager=True,
+        tokenizer_mode="mistral",
+        config_format="mistral",
+        load_format="mistral",
+        limit_mm_per_prompt={"audio": len(audio_assets)},
+    ) as vllm_model:
+        offline_outputs = vllm_model.generate_greedy(
+            [vllm_prompt],
+            max_tokens,
+            audios=[audio_data],
+        )
+
+    offline_text = offline_outputs[0][1]
+    assert offline_text, "Offline vLLM inference produced empty output"
+
+    def _asset_to_openai_chunk(asset):
         audio = Audio.from_file(str(asset.get_local_path()), strict=False)
         audio.format = "wav"
-        audio_dict = AudioChunk.from_audio(audio).to_openai()
-        return audio_dict
+        return AudioChunk.from_audio(audio).to_openai()
 
-    audio_chunks = [asset_to_chunk(asset) for asset in audio_assets]
-    text = f"What's happening in these {len(audio_assets)} audio clips?"
     messages = [
         {
             "role": "user",
-            "content": [*audio_chunks, {"type": "text", "text": text}],
+            "content": [
+                *[_asset_to_openai_chunk(a) for a in audio_assets],
+                {"type": "text", "text": question},
+            ],
         }
     ]
 
-    chat_completion = await client.chat.completions.create(
-        model=MODEL_NAME, messages=messages, max_tokens=10
+    server_args = [
+        "--enforce-eager",
+        "--limit-mm-per-prompt",
+        json.dumps({"audio": len(audio_assets)}),
+        *MISTRAL_FORMAT_ARGS,
+    ]
+
+    with RemoteOpenAIServer(
+        MODEL_NAME,
+        server_args,
+        env_dict={"VLLM_AUDIO_FETCH_TIMEOUT": "30"},
+    ) as remote_server:
+        client = remote_server.get_client()
+        completion = client.chat.completions.create(
+            model=MODEL_NAME,
+            messages=messages,
+            max_tokens=max_tokens,
+            temperature=0,
+        )
+
+    assert len(completion.choices) == 1
+    choice = completion.choices[0]
+    assert choice.finish_reason == "length"
+    assert choice.message.content == offline_text, (
+        f"Online serving output does not match offline inference.\n"
+        f"  Online:  {choice.message.content!r}\n"
+        f"  Offline: {offline_text!r}"
     )
 
-    assert len(chat_completion.choices) == 1
-    choice = chat_completion.choices[0]
-    assert choice.message.content == "In the first audio clip, you hear a brief"
-    assert choice.finish_reason == "length"
+
+def test_hf_reference(hf_runner, vllm_runner, audio_assets: AudioTestAssets):
+    """Compare vLLM Mistral-format output against HF Transformers reference.
+
+    Instead of requiring an exact text match (which is brittle across
+    attention backends), we compare per-token logprobs using the standard
+    check_logprobs_close helper: when tokens diverge at a position, each
+    runner's chosen token must appear in the other's top-k logprobs.
+
+    Marked xfail(strict=False) so remaining edge-case mismatches
+    don't block CI.
+    """
+    question = f"What's happening in these {len(audio_assets)} audio clips?"
+    max_tokens = 10
+    num_logprobs = 5
+    audio_data = [asset.audio_and_sample_rate for asset in audio_assets]
+
+    vllm_prompt = _get_prompt(audio_assets, question)
+    with vllm_runner(
+        MODEL_NAME,
+        dtype="half",
+        enforce_eager=True,
+        tokenizer_mode="mistral",
+        config_format="mistral",
+        load_format="mistral",
+        limit_mm_per_prompt={"audio": len(audio_assets)},
+    ) as vllm_model:
+        vllm_outputs = vllm_model.generate_greedy_logprobs(
+            [vllm_prompt],
+            max_tokens,
+            num_logprobs,
+            audios=[audio_data],
+        )
+    assert vllm_outputs[0][1], "vLLM inference produced empty output"
+
+    with hf_runner(
+        MODEL_NAME,
+        dtype="half",
+        auto_cls=VoxtralForConditionalGeneration,
+    ) as hf_model:
+        hf_model = model_utils.voxtral_patch_hf_runner(hf_model)
+        hf_outputs = hf_model.generate_greedy_logprobs_limit(
+            [question],
+            max_tokens,
+            num_logprobs,
+            audios=[audio_data],
+        )
+    assert hf_outputs[0][1], "HF Transformers produced empty output"
+
+    print(
+        f"HF Reference Comparison\n"
+        f"  vLLM: {vllm_outputs[0][1]!r}\n"
+        f"  HF:   {hf_outputs[0][1]!r}"
+    )
+    check_logprobs_close(
+        outputs_0_lst=vllm_outputs,
+        outputs_1_lst=hf_outputs,
+        name_0="vllm",
+        name_1="hf",
+    )

tests/models/multimodal/generation/test_voxtral_realtime.py

@@ -10,11 +10,11 @@ from mistral_common.protocol.transcription.request import (
     TranscriptionRequest,
 )
 from mistral_common.tokens.tokenizers.mistral import MistralTokenizer
+from mistral_common.tokens.tokenizers.tekken import SpecialTokenPolicy
 
 from vllm import LLM, EngineArgs, SamplingParams
 from vllm.assets.audio import AudioAsset
 from vllm.engine.arg_utils import AsyncEngineArgs
-from vllm.model_executor.models.voxtral_realtime import VoxtralRealtimeBuffer
 from vllm.v1.engine.async_llm import AsyncLLM
 
 MODEL_NAME = "mistralai/Voxtral-Mini-4B-Realtime-2602"
@@ -27,7 +27,7 @@ ENGINE_CONFIG = dict(
     load_format="mistral",
     tokenizer_mode="mistral",
     enforce_eager=True,
-    gpu_memory_utilization=0.4,
+    gpu_memory_utilization=0.9,
 )
 
 
@@ -114,6 +114,9 @@ def test_voxtral_realtime_forward(audio_assets, tokenizer, engine):
 
 @pytest.mark.asyncio
 async def test_voxtral_realtime_generator(audio_assets, tokenizer, async_engine):
+    # Lazy import to avoid CUDA-reinitialization error
+    from vllm.model_executor.models.voxtral_realtime import VoxtralRealtimeBuffer
+
     sampling_params = SamplingParams(temperature=0.0, max_tokens=1)
     audio_config = tokenizer.instruct_tokenizer.audio_encoder.audio_config
 
@@ -146,6 +149,9 @@ async def test_voxtral_realtime_generator(audio_assets, tokenizer, async_engine)
 
         output_tokens_list.append(output_tokens)
 
-    texts = [tokenizer.decode(output_tokens) for output_tokens in output_tokens_list]
+    texts = [
+        tokenizer.decode(output_tokens, special_token_policy=SpecialTokenPolicy.IGNORE)
+        for output_tokens in output_tokens_list
+    ]
     texts[1] = texts[1].replace("a base hit", "OBS").replace("oh my", "oh, my")
     assert texts == EXPECTED_TEXT

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

@@ -1215,3 +1215,91 @@ def tarsier_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
         hf_processor.patch_size = vision_encoder_info.get_patch_size()
 
     return hf_model
+
+
+def voxtral_patch_hf_runner(hf_model: "HfRunner") -> "HfRunner":
+    """Patch HfRunner for Voxtral's conversation-based processor.
+
+    Two issues in HfRunner require patching:
+
+    1. VoxtralProcessor requires ``apply_chat_template()`` with conversation
+       dicts (accepting ``url``, ``path``, or ``base64`` audio) rather than
+       the standard ``processor(text=, audio=, sampling_rate=)`` interface.
+    2. HfRunner.get_inputs cannot handle multi-audio per prompt because it
+       mis-unpacks ``[(arr1, sr1), (arr2, sr2)]`` via a ``len == 2`` check.
+
+    We override ``get_inputs`` to build conversation dicts and call
+    ``apply_chat_template`` directly, bypassing both issues. We also wrap
+    ``model.generate`` to strip prompt tokens before decoding, since
+    HfRunner.generate calls batch_decode on the full sequence (prompt +
+    generated).
+    """
+
+    import base64
+    import io
+
+    import soundfile as sf
+
+    processor = hf_model.processor
+
+    def _audio_to_base64(audio_array, sample_rate: int) -> str:
+        """Encode a numpy audio array as a base64 WAV string."""
+        buf = io.BytesIO()
+        sf.write(buf, audio_array, int(sample_rate), format="WAV")
+        return base64.b64encode(buf.getvalue()).decode("ascii")
+
+    def patched_get_inputs(prompts, images=None, videos=None, audios=None, **kwargs):
+        all_inputs = []
+        for i, prompt in enumerate(prompts):
+            content: list[dict] = []
+
+            if audios is not None and audios[i] is not None:
+                items = audios[i]
+                if not isinstance(items, list):
+                    items = [items]
+                for item in items:
+                    if isinstance(item, (list, tuple)) and len(item) == 2:
+                        arr, sr = item
+                    else:
+                        arr, sr = item, 16_000
+                    content.append(
+                        {
+                            "type": "audio",
+                            "base64": _audio_to_base64(arr, sr),
+                        }
+                    )
+
+            content.append({"type": "text", "text": prompt})
+
+            inputs = processor.apply_chat_template(
+                [{"role": "user", "content": content}]
+            )
+            if hasattr(inputs, "to"):
+                inputs = inputs.to(dtype=hf_model.dtype)
+            all_inputs.append(inputs)
+
+        return all_inputs
+
+    _orig_generate = hf_model.model.generate
+
+    def patched_generate(*args, **kwargs):
+        """Strip prompt tokens so only generated tokens are decoded."""
+        input_ids = kwargs.get("input_ids")
+        if input_ids is None and args:
+            input_ids = args[0]
+        prompt_len = input_ids.shape[1] if input_ids is not None else 0
+
+        output = _orig_generate(*args, **kwargs)
+        if prompt_len:
+            if isinstance(output, torch.Tensor):
+                output = output[:, prompt_len:]
+            else:
+                # GenerateDecoderOnlyOutput - trim sequences but preserve
+                # scores/logits so generate_greedy_logprobs_limit can
+                # extract per-token logprobs.
+                output.sequences = output.sequences[:, prompt_len:]
+        return output
+
+    hf_model.get_inputs = patched_get_inputs  # type: ignore[method-assign, assignment]
+    hf_model.model.generate = patched_generate  # type: ignore[method-assign]
+    return hf_model

tests/models/multimodal/pooling/test_colqwen3.py

@@ -0,0 +1,156 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for ColQwen3 late interaction model for multi-modal retrieval.
+
+ColQwen3 is a multi-vector retrieval model based on Qwen3-VL 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 = [
+    "TomoroAI/tomoro-colqwen3-embed-4b",
+    "OpenSearch-AI/Ops-Colqwen3-4B",
+]
+
+EMBED_DIMS = {
+    "TomoroAI/tomoro-colqwen3-embed-4b": 320,
+    "OpenSearch-AI/Ops-Colqwen3-4B": 2560,
+}
+
+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_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_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_relevance_ordering(
+    vllm_runner,
+    model: str,
+    dtype: str,
+) -> None:
+    _run_relevance_test(vllm_runner, model, dtype=dtype)

tests/models/multimodal/pooling/test_prithvi_mae.py

@@ -40,7 +40,7 @@ def _run_test(
         vllm_model.llm.encode(prompt, pooling_task="plugin")
 
 
-MODELS = ["mgazz/Prithvi-EO-2.0-300M-TL-Sen1Floods11"]
+MODELS = ["ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11"]
 
 
 @pytest.mark.core_model

tests/models/multimodal/processing/test_common.py

@@ -184,22 +184,42 @@ def get_text_token_prompts(
     text_prompt: str | None
     token_prompt: list[int]
     if isinstance(tokenizer, MistralTokenizer):
-        images = parsed_data.get("image", [])
-        request = ChatCompletionRequest(
-            messages=[
-                UserMessage(
-                    content=[
-                        TextChunk(text=""),
-                        *(ImageChunk(image=image) for image in images),
-                    ]
-                ),
-            ]
+        # ChatCompletionRequest only supports ImageChunk natively;
+        # for other modalities (e.g. audio), fall back to the model's
+        # own dummy inputs builder which knows the right placeholders.
+        has_non_image = any(
+            k != "image" and count > 0 for k, count in mm_counts.items()
         )
-        res = tokenizer.mistral.encode_chat_completion(request)
 
-        # Mistral does not support decode_tokens with skip_special_tokens=False
-        text_prompt = None
-        token_prompt = res.tokens
+        if has_non_image:
+            inputs = dummy_inputs.get_dummy_processor_inputs(
+                model_config.max_model_len,
+                mm_counts,
+            )
+            text_prompt = None
+            token_prompt = (
+                inputs.prompt
+                if isinstance(inputs.prompt, list)
+                else tokenizer.encode(inputs.prompt, add_special_tokens=False)
+            )
+        else:
+            images = parsed_data.get("image", [])
+            request = ChatCompletionRequest(
+                messages=[
+                    UserMessage(
+                        content=[
+                            TextChunk(text=""),
+                            *(ImageChunk(image=image) for image in images),
+                        ]
+                    ),
+                ]
+            )
+            res = tokenizer.mistral.encode_chat_completion(request)
+
+            # Mistral does not support decode_tokens with
+            # skip_special_tokens=False
+            text_prompt = None
+            token_prompt = res.tokens
     else:
         inputs = dummy_inputs.get_dummy_processor_inputs(
             model_config.max_model_len,

tests/models/multimodal/processing/test_gemma3.py

@@ -168,6 +168,7 @@ def test_get_image_size_with_most_features(
         image_width=max_image_size.width,
         image_height=max_image_size.height,
         processor=hf_processor,
+        mm_kwargs=hf_processor_mm_kwargs,
     )
 
     prompt = "<start_of_image>"

tests/models/multimodal/processing/test_idefics3.py

@@ -3,7 +3,9 @@
 """Tests for Idefics3's multimodal preprocessing kwargs."""
 
 import pytest
+from packaging.version import Version
 from transformers import Idefics3Config
+from transformers import __version__ as TRANSFORMERS_VERSION
 
 from vllm.multimodal import MULTIMODAL_REGISTRY
 
@@ -11,6 +13,10 @@ from ....conftest import ImageTestAssets
 from ...utils import build_model_context
 
 
+@pytest.mark.skipif(
+    Version(TRANSFORMERS_VERSION) < Version("5.2.0"),
+    reason="See https://github.com/huggingface/transformers/pull/43948",
+)
 @pytest.mark.parametrize("model_id", ["HuggingFaceM4/Idefics3-8B-Llama3"])
 @pytest.mark.parametrize(
     ("mm_processor_kwargs", "expected_toks_per_img"),
@@ -63,7 +69,11 @@ def test_processor_override(
 
     # Ensure the placeholders format are correct
     hf_processor = processor.info.get_hf_processor(**hf_processor_mm_kwargs)
-    hf_processed_inputs = hf_processor(text=prompt, images=mm_data["image"])
+    hf_processed_inputs = hf_processor(
+        text=prompt,
+        images=mm_data["image"],
+        **processor.info.ctx.get_merged_mm_kwargs(hf_processor_mm_kwargs),
+    )
     assert processed_inputs["prompt_token_ids"] == hf_processed_inputs["input_ids"][0]
 
     # Ensure we have the right number of placeholders per num_crops size

tests/models/multimodal/processing/test_qwen2_vl.py

@@ -82,6 +82,7 @@ def test_get_image_size_with_most_features(
         image_width=max_image_size.width,
         image_height=max_image_size.height,
         image_processor=hf_processor.image_processor,
+        mm_kwargs=hf_processor_mm_kwargs,
     )
 
     prompt = "<|vision_start|><|image_pad|><|vision_end|>"

tests/models/multimodal/processing/test_smolvlm.py

@@ -3,7 +3,9 @@
 """Tests for smolvlm's multimodal preprocessing kwargs."""
 
 import pytest
+from packaging.version import Version
 from transformers import SmolVLMConfig
+from transformers import __version__ as TRANSFORMERS_VERSION
 
 from vllm.multimodal import MULTIMODAL_REGISTRY
 
@@ -11,6 +13,10 @@ from ....conftest import ImageTestAssets
 from ...utils import build_model_context
 
 
+@pytest.mark.skipif(
+    Version(TRANSFORMERS_VERSION) < Version("5.2.0"),
+    reason="See https://github.com/huggingface/transformers/pull/43948",
+)
 @pytest.mark.parametrize("model_id", ["HuggingFaceTB/SmolVLM2-2.2B-Instruct"])
 @pytest.mark.parametrize(
     ("mm_processor_kwargs", "expected_toks_per_img"),
@@ -63,7 +69,11 @@ def test_processor_override(
 
     # Ensure the placeholders format are correct
     hf_processor = processor.info.get_hf_processor(**hf_processor_mm_kwargs)
-    hf_processed_inputs = hf_processor(text=prompt, images=mm_data["image"])
+    hf_processed_inputs = hf_processor(
+        text=prompt,
+        images=mm_data["image"],
+        **processor.info.ctx.get_merged_mm_kwargs(hf_processor_mm_kwargs),
+    )
     assert processed_inputs["prompt_token_ids"] == hf_processed_inputs["input_ids"][0]
 
     # Ensure we have the right number of placeholders per num_crops size

tests/models/registry.py

@@ -529,6 +529,15 @@ _EMBEDDING_EXAMPLE_MODELS = {
     # [Text-only]
     "BertModel": _HfExamplesInfo("BAAI/bge-base-en-v1.5"),
     "HF_ColBERT": _HfExamplesInfo("answerdotai/answerai-colbert-small-v1"),
+    "ColBERTModernBertModel": _HfExamplesInfo(
+        "lightonai/GTE-ModernColBERT-v1",
+        hf_overrides={"architectures": ["ColBERTModernBertModel"]},
+    ),
+    "ColBERTJinaRobertaModel": _HfExamplesInfo(
+        "jinaai/jina-colbert-v2",
+        trust_remote_code=True,
+        hf_overrides={"architectures": ["ColBERTJinaRobertaModel"]},
+    ),
     "BgeM3EmbeddingModel": _HfExamplesInfo("BAAI/bge-m3"),
     "Gemma2Model": _HfExamplesInfo("BAAI/bge-multilingual-gemma2"),
     "Gemma3TextModel": _HfExamplesInfo("google/embeddinggemma-300m"),
@@ -588,6 +597,12 @@ _EMBEDDING_EXAMPLE_MODELS = {
         "TIGER-Lab/VLM2Vec-Full", trust_remote_code=True
     ),
     "Qwen2VLForConditionalGeneration": _HfExamplesInfo("MrLight/dse-qwen2-2b-mrl-v1"),
+    "ColQwen3": _HfExamplesInfo(
+        "TomoroAI/tomoro-colqwen3-embed-4b", trust_remote_code=True
+    ),
+    "OpsColQwen3Model": _HfExamplesInfo(
+        "OpenSearch-AI/Ops-Colqwen3-4B", trust_remote_code=True
+    ),
     "SiglipModel": _HfExamplesInfo("google/siglip-base-patch16-224"),
     "PrithviGeoSpatialMAE": _HfExamplesInfo(
         "ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11",
@@ -915,6 +930,12 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         },
     ),
     "Ovis2_5": _HfExamplesInfo("AIDC-AI/Ovis2.5-2B", trust_remote_code=True),
+    "Ovis2_6ForCausalLM": _HfExamplesInfo(
+        "AIDC-AI/Ovis2.6-2B", is_available_online=False, trust_remote_code=True
+    ),
+    "Ovis2_6_MoeForCausalLM": _HfExamplesInfo(
+        "AIDC-AI/Ovis2.6-30B-A3B", trust_remote_code=True
+    ),
     "PaddleOCRVLForConditionalGeneration": _HfExamplesInfo(
         "PaddlePaddle/PaddleOCR-VL",
         trust_remote_code=True,

tests/models/test_terratorch.py

@@ -13,7 +13,7 @@ from tests.utils import create_new_process_for_each_test
     "model",
     [
         "ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11",
-        "mgazz/Prithvi_v2_eo_300_tl_unet_agb",
+        "ibm-nasa-geospatial/Prithvi-EO-2.0-300M-BurnScars",
     ],
 )
 def test_inference(

tests/plugins/prithvi_io_processor_plugin/prithvi_io_processor/prithvi_processor.py

@@ -44,12 +44,8 @@ datamodule_config: DataModuleConfig = {
     "no_label_replace": -1,
     "num_workers": 8,
     "test_transform": [
-        albumentations.Resize(
-            always_apply=False, height=448, interpolation=1, p=1, width=448
-        ),
-        albumentations.pytorch.ToTensorV2(
-            transpose_mask=False, always_apply=True, p=1.0
-        ),
+        albumentations.Resize(height=448, interpolation=1, p=1, width=448),
+        albumentations.pytorch.ToTensorV2(transpose_mask=False, p=1.0),
     ],
 }
 

tests/plugins_tests/test_io_processor_plugins.py

@@ -1,18 +1,39 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 import base64
+import io
 
+import imagehash
 import pytest
 import requests
+from PIL import Image
 
 from tests.utils import RemoteOpenAIServer
 from vllm.config import VllmConfig
 from vllm.entrypoints.pooling.pooling.protocol import IOProcessorResponse
 from vllm.plugins.io_processors import get_io_processor
 
-MODEL_NAME = "ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11"
+models_config = {
+    "ibm-nasa-geospatial/Prithvi-EO-2.0-300M-TL-Sen1Floods11": {
+        "image_url": "https://huggingface.co/christian-pinto/Prithvi-EO-2.0-300M-TL-VLLM/resolve/main/valencia_example_2024-10-26.tiff",  # noqa: E501
+        "out_hash": "aa6d92ad25926a5e",
+        "plugin": "prithvi_to_tiff",
+    },
+    "ibm-nasa-geospatial/Prithvi-EO-2.0-300M-BurnScars": {
+        "image_url": "https://huggingface.co/ibm-nasa-geospatial/Prithvi-EO-2.0-300M-BurnScars/resolve/main/examples/subsetted_512x512_HLS.S30.T10SEH.2018190.v1.4_merged.tif",  # noqa: E501
+        "out_hash": "c07f4f602da73552",
+        "plugin": "prithvi_to_tiff",
+    },
+}
 
-image_url = "https://huggingface.co/christian-pinto/Prithvi-EO-2.0-300M-TL-VLLM/resolve/main/valencia_example_2024-10-26.tiff"  # noqa: E501
+
+def _compute_image_hash(base64_data: str) -> str:
+    # Decode the base64 output and create image from byte stream
+    decoded_image = base64.b64decode(base64_data)
+    image = Image.open(io.BytesIO(decoded_image))
+
+    # Compute perceptual hash of the output image
+    return str(imagehash.phash(image))
 
 
 def test_loading_missing_plugin():
@@ -22,33 +43,39 @@ def test_loading_missing_plugin():
 
 
 @pytest.fixture(scope="function")
-def server():
+def server(model_name, plugin):
     args = [
         "--runner",
         "pooling",
         "--enforce-eager",
-        "--trust-remote-code",
         "--skip-tokenizer-init",
         # Limit the maximum number of parallel requests
         # to avoid the model going OOM in CI.
         "--max-num-seqs",
         "32",
         "--io-processor-plugin",
-        "prithvi_to_tiff",
-        "--model-impl",
-        "terratorch",
+        plugin,
         "--enable-mm-embeds",
     ]
 
-    with RemoteOpenAIServer(MODEL_NAME, args) as remote_server:
+    with RemoteOpenAIServer(model_name, args) as remote_server:
         yield remote_server
 
 
 @pytest.mark.asyncio
-@pytest.mark.parametrize("model_name", [MODEL_NAME])
+@pytest.mark.parametrize(
+    "model_name, image_url, plugin, expected_hash",
+    [
+        (model_name, config["image_url"], config["plugin"], config["out_hash"])
+        for model_name, config in models_config.items()
+    ],
+)
 async def test_prithvi_mae_plugin_online(
     server: RemoteOpenAIServer,
     model_name: str,
+    image_url: str | dict,
+    plugin: str,
+    expected_hash: str,
 ):
     request_payload_url = {
         "data": {
@@ -74,16 +101,25 @@ async def test_prithvi_mae_plugin_online(
 
     # verify the output is formatted as expected for this plugin
     plugin_data = parsed_response.data
-
     assert all(plugin_data.get(attr) for attr in ["type", "format", "data"])
 
-    # We just check that the output is a valid base64 string.
-    # Raises an exception and fails the test if the string is corrupted.
-    base64.b64decode(plugin_data["data"])
+    # Compute the output image hash and compare it against the expected hash
+    image_hash = _compute_image_hash(plugin_data["data"])
+    assert image_hash == expected_hash, (
+        f"Image hash mismatch: expected {expected_hash}, got {image_hash}"
+    )
 
 
-@pytest.mark.parametrize("model_name", [MODEL_NAME])
-def test_prithvi_mae_plugin_offline(vllm_runner, model_name: str):
+@pytest.mark.parametrize(
+    "model_name, image_url, plugin, expected_hash",
+    [
+        (model_name, config["image_url"], config["plugin"], config["out_hash"])
+        for model_name, config in models_config.items()
+    ],
+)
+def test_prithvi_mae_plugin_offline(
+    vllm_runner, model_name: str, image_url: str | dict, plugin: str, expected_hash: str
+):
     img_prompt = dict(
         data=image_url,
         data_format="url",
@@ -96,13 +132,12 @@ def test_prithvi_mae_plugin_offline(vllm_runner, model_name: str):
         runner="pooling",
         skip_tokenizer_init=True,
         enable_mm_embeds=True,
-        trust_remote_code=True,
         enforce_eager=True,
         # Limit the maximum number of parallel requests
         # to avoid the model going OOM in CI.
-        max_num_seqs=1,
-        model_impl="terratorch",
-        io_processor_plugin="prithvi_to_tiff",
+        max_num_seqs=32,
+        io_processor_plugin=plugin,
+        default_torch_num_threads=1,
     ) as llm_runner:
         pooler_output = llm_runner.get_llm().encode(img_prompt, pooling_task="plugin")
     output = pooler_output[0].outputs
@@ -110,6 +145,8 @@ def test_prithvi_mae_plugin_offline(vllm_runner, model_name: str):
     # verify the output is formatted as expected for this plugin
     assert all(hasattr(output, attr) for attr in ["type", "format", "data"])
 
-    # We just check that the output is a valid base64 string.
-    # Raises an exception and fails the test if the string is corrupted.
-    base64.b64decode(output.data)
+    # Compute the output image hash and compare it against the expected hash
+    image_hash = _compute_image_hash(output.data)
+    assert image_hash == expected_hash, (
+        f"Image hash mismatch: expected {expected_hash}, got {image_hash}"
+    )

tests/reasoning/test_gptoss_reasoning_parser.py

@@ -17,7 +17,9 @@ def gpt_oss_tokenizer():
 
 USER_MESSAGE_START = "<|start|>user<|message|>"
 REASONING_SECTION_START = "<|end|><|start|>assistant<|channel|>analysis<|message|>"
-ASSISTANT_CONTENT_START_PREFIX = "<|end|><|start|>assistant<|channel|>final"
+END = "<|end|>"
+ASSISTANT_START = "<|start|>assistant"
+ASSISTANT_CONTENT_START_PREFIX = END + ASSISTANT_START + "<|channel|>final"
 ASSISTANT_CONTENT_START_SUFFIX = "<|message|>"
 ASSISTANT_CONTENT_START = (
     ASSISTANT_CONTENT_START_PREFIX + ASSISTANT_CONTENT_START_SUFFIX
@@ -97,6 +99,20 @@ COMPLEX_CONTENT_2 = {
     "is_reasoning_end": True,
 }
 
+MULTI_TURN_CONTENT = {
+    "output": USER_MESSAGE_START
+    + "1st turn user message"
+    + REASONING_SECTION_START
+    + "1st turn reasoning"
+    + ASSISTANT_CONTENT_START
+    + "1st turn response"
+    + END
+    + USER_MESSAGE_START
+    + "2nd turn user message"
+    + END
+    + ASSISTANT_START,
+    "is_reasoning_end": False,
+}
 TEST_CASES = [
     BASIC_CONTENT,
     BASIC_REASONING_ONLY,
@@ -106,6 +122,7 @@ TEST_CASES = [
     COMPLEX_CONTENT_1,
     COMPLEX_CONTENT_1_WITH_CONTENT,
     COMPLEX_CONTENT_2,
+    MULTI_TURN_CONTENT,
 ]
 
 

tests/renderers/test_completions.py

@@ -38,6 +38,12 @@ class MockModelConfig:
     enable_prompt_embeds: bool = True
     skip_tokenizer_init: bool = False
     is_encoder_decoder: bool = False
+    is_multimodal_model: bool = False
+
+
+@dataclass
+class MockVllmConfig:
+    model_config: MockModelConfig
 
 
 @dataclass
@@ -72,16 +78,17 @@ def _build_renderer(
     _, tokenizer_name, _, kwargs = tokenizer_args_from_config(model_config)
 
     renderer = HfRenderer(
-        model_config,
-        tokenizer_kwargs={**kwargs, "tokenizer_name": tokenizer_name},
+        MockVllmConfig(model_config),
+        tokenizer=(
+            None
+            if model_config.skip_tokenizer_init
+            else DummyTokenizer(
+                truncation_side=truncation_side,
+                max_chars_per_token=max_chars_per_token,
+            )
+        ),
     )
 
-    if not model_config.skip_tokenizer_init:
-        renderer._tokenizer = DummyTokenizer(
-            truncation_side=truncation_side,
-            max_chars_per_token=max_chars_per_token,
-        )
-
     return renderer
 
 
@@ -104,14 +111,14 @@ class TestValidatePrompt:
         renderer = _build_renderer(MockModelConfig())
 
         with pytest.raises(ValueError, match="at least one prompt"):
-            renderer.render_prompts(_preprocess_prompt(renderer.config, []))
+            renderer.render_prompts(_preprocess_prompt(renderer.model_config, []))
 
     def test_invalid_type(self):
         renderer = _build_renderer(MockModelConfig())
 
         with pytest.raises(TypeError, match="should be a list of integers"):
             renderer.render_prompts(
-                _preprocess_prompt(renderer.config, [[1, 2], ["foo", "bar"]])  # type: ignore[arg-type]
+                _preprocess_prompt(renderer.model_config, [[1, 2], ["foo", "bar"]])  # type: ignore[arg-type]
             )
 
 
@@ -120,7 +127,9 @@ class TestRenderPrompt:
         renderer = _build_renderer(MockModelConfig())
 
         tokens = [101, 7592, 2088]
-        prompts = renderer.render_prompts(_preprocess_prompt(renderer.config, tokens))
+        prompts = renderer.render_prompts(
+            _preprocess_prompt(renderer.model_config, tokens)
+        )
         results = renderer.tokenize_prompts(
             prompts,
             TokenizeParams(max_total_tokens=100),
@@ -134,7 +143,7 @@ class TestRenderPrompt:
 
         token_lists = [[101, 7592, 2088], [102, 1234, 5678, 9012], [103, 4567]]
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, token_lists)
+            _preprocess_prompt(renderer.model_config, token_lists)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -151,7 +160,7 @@ class TestRenderPrompt:
 
         text_input = "x" * 10
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, text_input)
+            _preprocess_prompt(renderer.model_config, text_input)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -166,7 +175,7 @@ class TestRenderPrompt:
 
         text_list_input = ["x" * 10, "x" * 12, "x" * 14]
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, text_list_input)
+            _preprocess_prompt(renderer.model_config, text_list_input)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -181,7 +190,7 @@ class TestRenderPrompt:
         renderer = _build_renderer(MockModelConfig())
 
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, "x" * 200)
+            _preprocess_prompt(renderer.model_config, "x" * 200)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -195,7 +204,7 @@ class TestRenderPrompt:
         renderer = _build_renderer(MockModelConfig())
 
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, "x" * 200)
+            _preprocess_prompt(renderer.model_config, "x" * 200)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -209,7 +218,7 @@ class TestRenderPrompt:
         renderer = _build_renderer(MockModelConfig())
 
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, "x" * 200)
+            _preprocess_prompt(renderer.model_config, "x" * 200)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -224,7 +233,7 @@ class TestRenderPrompt:
 
         long_tokens = [100, 101, 102, 103, 104, 105, 106, 107, 108, 109]  # 10 tokens
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, long_tokens)
+            _preprocess_prompt(renderer.model_config, long_tokens)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -240,7 +249,7 @@ class TestRenderPrompt:
 
         long_tokens = [100, 101, 102, 103, 104, 105, 106, 107, 108, 109]  # 10 tokens
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, long_tokens)
+            _preprocess_prompt(renderer.model_config, long_tokens)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -257,7 +266,7 @@ class TestRenderPrompt:
         # Exceeds max_total_tokens and max_total_tokens * VLLM_MAX_CHARS_PER_TOKEN
         long_tokens = "x" * 150
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, long_tokens)
+            _preprocess_prompt(renderer.model_config, long_tokens)
         )
 
         with pytest.raises(
@@ -270,7 +279,7 @@ class TestRenderPrompt:
             )
 
         # Should not even attempt tokenization
-        assert renderer._tokenizer._captured_encode_kwargs == {}
+        assert renderer.tokenizer._captured_encode_kwargs == {}
 
     def test_text_max_length_exceeded_nonobvious(self):
         renderer = _build_renderer(MockModelConfig(), max_chars_per_token=2)
@@ -278,7 +287,7 @@ class TestRenderPrompt:
         # Exceeds max_total_tokens but not max_total_tokens * VLLM_MAX_CHARS_PER_TOKEN
         long_tokens = "x" * 150
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, long_tokens)
+            _preprocess_prompt(renderer.model_config, long_tokens)
         )
 
         with pytest.raises(
@@ -291,15 +300,15 @@ class TestRenderPrompt:
             )
 
         # Should only tokenize the first max_total_tokens + 1 tokens
-        assert renderer._tokenizer._captured_encode_kwargs["truncation"] is True
-        assert renderer._tokenizer._captured_encode_kwargs["max_length"] == 101
+        assert renderer.tokenizer._captured_encode_kwargs["truncation"] is True
+        assert renderer.tokenizer._captured_encode_kwargs["max_length"] == 101
 
     def test_token_max_length_exceeded(self):
         renderer = _build_renderer(MockModelConfig())
 
         long_tokens = list(range(150))  # Exceeds max_total_tokens=100
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, long_tokens)
+            _preprocess_prompt(renderer.model_config, long_tokens)
         )
 
         with pytest.raises(
@@ -315,7 +324,7 @@ class TestRenderPrompt:
         renderer = _build_renderer(MockModelConfig(skip_tokenizer_init=True))
 
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, "Hello world")
+            _preprocess_prompt(renderer.model_config, "Hello world")
         )
 
         with pytest.raises(ValueError, match="`skip_tokenizer_init=True`"):
@@ -328,7 +337,9 @@ class TestRenderPrompt:
         renderer = _build_renderer(MockModelConfig())
 
         tokens = [1, 2, 3, 4]
-        prompts = renderer.render_prompts(_preprocess_prompt(renderer.config, tokens))
+        prompts = renderer.render_prompts(
+            _preprocess_prompt(renderer.model_config, tokens)
+        )
         results = renderer.tokenize_prompts(
             prompts,
             TokenizeParams(
@@ -358,7 +369,7 @@ class TestRenderEmbedPrompt:
         embed_bytes = self._create_test_embed_bytes(tensor_input)
 
         prompts = renderer.render_prompts(
-            _preprocess_prompt(renderer.config, embed_bytes)
+            _preprocess_prompt(renderer.model_config, embed_bytes)
         )
         results = renderer.tokenize_prompts(
             prompts,
@@ -379,7 +390,7 @@ class TestRenderEmbedPrompt:
 
         prompts = renderer.render_prompts(
             _preprocess_prompt(
-                renderer.config,
+                renderer.model_config,
                 [self._create_test_embed_bytes(t) for t in tensor_inputs],
             )
         )
@@ -400,7 +411,7 @@ class TestRenderEmbedPrompt:
 
         prompts = renderer.render_prompts(
             _preprocess_prompt(
-                renderer.config, self._create_test_embed_bytes(tensor_input)
+                renderer.model_config, self._create_test_embed_bytes(tensor_input)
             )
         )
         results = renderer.tokenize_prompts(
@@ -427,7 +438,7 @@ class TestRenderEmbedPrompt:
 
             prompts = renderer.render_prompts(
                 _preprocess_prompt(
-                    renderer.config, self._create_test_embed_bytes(tensor_input)
+                    renderer.model_config, self._create_test_embed_bytes(tensor_input)
                 )
             )
             results = renderer.tokenize_prompts(
@@ -446,7 +457,7 @@ class TestRenderEmbedPrompt:
 
         prompts = renderer.render_prompts(
             _preprocess_prompt(
-                renderer.config, self._create_test_embed_bytes(tensor_input)
+                renderer.model_config, self._create_test_embed_bytes(tensor_input)
             )
         )
         results = renderer.tokenize_prompts(
@@ -466,7 +477,7 @@ class TestRenderEmbedPrompt:
 
         prompts = renderer.render_prompts(
             _preprocess_prompt(
-                renderer.config,
+                renderer.model_config,
                 [text_input, self._create_test_embed_bytes(tensor_input)],
             )
         )

tests/renderers/test_mistral.py

@@ -36,6 +36,12 @@ class MockModelConfig:
     enable_prompt_embeds: bool = True
     skip_tokenizer_init: bool = False
     is_encoder_decoder: bool = False
+    is_multimodal_model: bool = False
+
+
+@dataclass
+class MockVllmConfig:
+    model_config: MockModelConfig
 
 
 @pytest.mark.asyncio
@@ -50,8 +56,10 @@ async def test_async_mistral_tokenizer_does_not_block_event_loop():
     mock_model_config = MockModelConfig(skip_tokenizer_init=True)
     mock_tokenizer = Mock(spec=MistralTokenizer)
     mock_tokenizer.apply_chat_template = mocked_apply_chat_template
-    mock_renderer = MistralRenderer(mock_model_config, tokenizer_kwargs={})
-    mock_renderer._tokenizer = mock_tokenizer
+    mock_renderer = MistralRenderer(
+        MockVllmConfig(mock_model_config),
+        tokenizer=mock_tokenizer,
+    )
 
     task = mock_renderer.render_messages_async([], ChatParams())
 

tests/test_inputs.py

@@ -3,7 +3,7 @@
 
 import pytest
 
-from vllm.config import ModelConfig
+from vllm.config import ModelConfig, VllmConfig
 from vllm.inputs.preprocess import InputPreprocessor
 
 pytestmark = pytest.mark.cpu_test
@@ -20,7 +20,8 @@ pytestmark = pytest.mark.cpu_test
 )
 def test_preprocessor_always_mm_code_path(model_id, prompt):
     model_config = ModelConfig(model=model_id)
-    input_preprocessor = InputPreprocessor(model_config)
+    vllm_config = VllmConfig(model_config=model_config)
+    input_preprocessor = InputPreprocessor(vllm_config)
 
     # HF processor adds sep token
     tokenizer = input_preprocessor.get_tokenizer()

tests/tokenizers_/test_detokenize.py

@@ -67,7 +67,6 @@ def _run_incremental_decode(
         mm_features=None,
         sampling_params=params,
         pooling_params=None,
-        eos_token_id=None,
         arrival_time=0.0,
         lora_request=None,
         cache_salt=None,

tests/tool_parsers/test_step3p5_tool_parser.py

@@ -1123,7 +1123,7 @@ rectangle
 
     # Encode all content tokens at once
     all_token_ids = step3p5_tokenizer.encode(model_output, add_special_tokens=False)
-    eos_token_id = getattr(step3p5_tokenizer, "eos_token_id", None)
+    eos_token_id = step3p5_tokenizer.eos_token_id
 
     # Include EOS token in delta_token_ids if available
     if eos_token_id is not None:

tests/v1/attention/test_sparse_mla_backends.py

@@ -23,6 +23,17 @@ from vllm import _custom_ops as ops
 from vllm.config import set_current_vllm_config
 from vllm.model_executor.layers.linear import ColumnParallelLinear
 from vllm.platforms import current_platform
+
+# TODO: Integrate ROCMAiterMLASparseBackend for ROCm.
+# The ROCm sparse MLA backend (rocm_aiter_mla_sparse.py) has a compatible
+# forward_mqa interface but needs validation on ROCm hardware.
+if not current_platform.is_cuda():
+    pytest.skip(
+        "Sparse MLA backend tests currently only support CUDA. "
+        "ROCm support requires integrating ROCMAiterMLASparseBackend.",
+        allow_module_level=True,
+    )
+
 from vllm.utils.math_utils import cdiv
 from vllm.v1.attention.backends.mla.flashinfer_mla_sparse import (
     FlashInferMLASparseBackend,

tests/v1/core/test_kv_cache_utils.py

@@ -84,13 +84,15 @@ def make_request(
             )
             mm_features.append(mm_feature)
 
+    sampling_params = SamplingParams(max_tokens=17)
+    sampling_params.update_from_generation_config({}, eos_token_id=100)
+
     return Request(
         request_id=request_id,
         prompt_token_ids=prompt_token_ids,
         mm_features=mm_features if mm_features else None,
-        sampling_params=SamplingParams(max_tokens=17),
+        sampling_params=sampling_params,
         pooling_params=None,
-        eos_token_id=100,
         lora_request=None,
         cache_salt=cache_salt,
         block_hasher=get_request_block_hasher(block_size, hash_fn),

tests/v1/core/test_prefix_caching.py

@@ -75,13 +75,15 @@ def make_request(
             )
             mm_features.append(mm_feature)
 
+    sampling_params = SamplingParams(max_tokens=17, prompt_logprobs=prompt_logprobs)
+    sampling_params.update_from_generation_config({}, eos_token_id=100)
+
     return Request(
         request_id=request_id,
         prompt_token_ids=prompt_token_ids,
         mm_features=mm_features if mm_features else None,
-        sampling_params=SamplingParams(max_tokens=17, prompt_logprobs=prompt_logprobs),
+        sampling_params=sampling_params,
         pooling_params=None,
-        eos_token_id=100,
         lora_request=lora_request,
         cache_salt=cache_salt,
         block_hasher=get_request_block_hasher(block_size, hash_fn),
@@ -742,6 +744,12 @@ def _make_hybrid_kv_cache_config(
             shapes=(1, 1),
             dtypes=(torch.float32,),
         ),
+        "mamba_align": lambda: MambaSpec(
+            block_size=block_size,
+            shapes=(1, 1),
+            dtypes=(torch.float32,),
+            mamba_cache_mode="align",
+        ),
     }
 
     kv_cache_groups = [
@@ -960,6 +968,46 @@ def test_prefill_hybrid_model_combinations_eagle(
     manager.free(req1)
 
 
+def test_prefill_hybrid_model_mamba_align():
+    """Test that MambaManager.cache_blocks() handles null blocks in align mode.
+
+    Regression test for https://github.com/vllm-project/vllm/issues/34361.
+    In mamba_cache_mode="align", allocate_new_blocks() pads req_to_blocks with
+    null blocks. cache_full_blocks() correctly skips them, but
+    MambaManager.cache_blocks() must also skip null blocks when tracking
+    cached_blocks_this_step.
+    """
+    block_size = 16
+    num_blocks = 30
+
+    kv_cache_config = _make_hybrid_kv_cache_config(
+        block_size, num_blocks, ["full", "mamba_align"]
+    )
+    manager = KVCacheManager(
+        kv_cache_config,
+        max_model_len=8192,
+        enable_caching=True,
+        hash_block_size=block_size,
+    )
+
+    hash_fn = sha256
+
+    # 3 full blocks (48 tokens) + 7 partial tokens = 55 tokens total
+    all_token_ids = [i for i in range(3) for _ in range(block_size)] + [3] * 7
+
+    # First request: allocate_slots should not crash with the assertion error
+    # in MambaManager.cache_blocks() when null blocks are present.
+    req0 = make_request("0", all_token_ids, block_size, hash_fn)
+    computed_blocks, num_computed_tokens = manager.get_computed_blocks(req0)
+    assert num_computed_tokens == 0
+
+    blocks = manager.allocate_slots(req0, 55, num_computed_tokens, computed_blocks)
+    assert blocks is not None
+    assert len(blocks.get_block_ids()) == 2  # full_attn + mamba groups
+
+    manager.free(req0)
+
+
 def test_prefill_plp():
     """Test prefill with APC and some prompt logprobs (plp) requests.
 

tests/v1/core/test_priority_scheduler_random.py

@@ -48,10 +48,9 @@ def _create_random_request(
 
     request_id = uuid.uuid4().hex
 
-    sampling_params = SamplingParams(
-        ignore_eos=False,
-        max_tokens=max_tokens,
-    )
+    sampling_params = SamplingParams(ignore_eos=False, max_tokens=max_tokens)
+    sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
+
     mm_features = []
     for j, position in enumerate(mm_positions):
         identifier = f"{request_id}_hash_{j}"
@@ -79,7 +78,6 @@ def _create_random_request(
         sampling_params=sampling_params,
         pooling_params=None,
         mm_features=mm_features if mm_features else None,
-        eos_token_id=EOS_TOKEN_ID,
         arrival_time=arrival_time,
         priority=priority,
         block_hasher=block_hasher,

tests/v1/core/test_scheduler.py

@@ -469,8 +469,7 @@ def test_stop_via_update_from_output():
 
     # Test case 4: Ignore EOS flag
     scheduler = create_scheduler(num_speculative_tokens=2)
-    requests = create_requests(num_requests=1, max_tokens=10)
-    requests[0].sampling_params.ignore_eos = True
+    requests = create_requests(num_requests=1, max_tokens=10, ignore_eos=True)
     requests[0].num_computed_tokens = requests[0].num_tokens
     scheduler.requests[requests[0].request_id] = requests[0]
     scheduler.running.append(requests[0])
@@ -515,12 +514,12 @@ def test_check_stop_min_tokens():
         max_tokens=20,
         min_tokens=5,
     )
+    sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
     request = Request(
         request_id="0",
         prompt_token_ids=[0, 1, 2],
         sampling_params=sampling_params,
         pooling_params=None,
-        eos_token_id=EOS_TOKEN_ID,
     )
     # Simulate having generated 3 output tokens (less than min_tokens=5)
     request.append_output_token_ids([10, 11, EOS_TOKEN_ID])  # EOS token present
@@ -551,12 +550,12 @@ def test_check_stop_min_tokens():
         max_tokens=20,
         min_tokens=0,
     )
+    sampling_params_no_min.update_from_generation_config({}, EOS_TOKEN_ID)
     request_no_min = Request(
         request_id="1",
         prompt_token_ids=[0, 1, 2],
         sampling_params=sampling_params_no_min,
         pooling_params=None,
-        eos_token_id=EOS_TOKEN_ID,
     )
     request_no_min.append_output_token_ids([10, EOS_TOKEN_ID])
 
@@ -571,12 +570,12 @@ def test_check_stop_min_tokens():
         min_tokens=5,
         stop_token_ids=[42],
     )
+    sampling_params_stop.update_from_generation_config({}, EOS_TOKEN_ID)
     request_stop = Request(
         request_id="2",
         prompt_token_ids=[0, 1, 2],
         sampling_params=sampling_params_stop,
         pooling_params=None,
-        eos_token_id=EOS_TOKEN_ID,
     )
     # Only 3 output tokens, less than min_tokens=5, but has stop token
     request_stop.append_output_token_ids([10, 11, 42])
@@ -1877,6 +1876,7 @@ def create_requests_with_priority(
         stop_token_ids=stop_token_ids,
         prompt_logprobs=prompt_logprobs,
     )
+    sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
     requests = []
 
     if mm_hashes_list is not None:
@@ -1938,7 +1938,6 @@ def create_requests_with_priority(
             sampling_params=sampling_params,
             pooling_params=None,
             mm_features=mm_features if mm_features else None,
-            eos_token_id=EOS_TOKEN_ID,
             arrival_time=arrival_times[i],
             priority=priorities[i],
             block_hasher=block_hasher,
@@ -2429,13 +2428,13 @@ def test_schedule_skip_tokenizer_init_structured_output_request():
         max_tokens=16,
         structured_outputs=structured_outputs_params,
     )
+    sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
     request = Request(
         request_id="0",
         prompt_token_ids=[0, 1],
         mm_features=None,
         sampling_params=sampling_params,
         pooling_params=None,
-        eos_token_id=EOS_TOKEN_ID,
     )
     scheduler.add_request(request)
     output = scheduler.schedule()

tests/v1/core/utils.py

@@ -174,6 +174,7 @@ def create_requests(
     num_tokens: int = 10,
     mm_hashes_list: list[list[str]] | None = None,
     mm_positions: list[list[PlaceholderRange]] | None = None,
+    ignore_eos: bool = False,
     max_tokens: int = 16,
     stop_token_ids: list[int] | None = None,
     prompt_logprobs: int | None = None,
@@ -188,11 +189,12 @@ def create_requests(
 
     block_hasher = get_request_block_hasher(block_size, sha256)
     sampling_params = SamplingParams(
-        ignore_eos=False,
+        ignore_eos=ignore_eos,
         max_tokens=max_tokens,
         stop_token_ids=stop_token_ids,
         prompt_logprobs=prompt_logprobs,
     )
+    sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
     requests = []
 
     if mm_hashes_list is not None:
@@ -250,7 +252,6 @@ def create_requests(
             sampling_params=sampling_params,
             pooling_params=None,
             mm_features=mm_features if mm_features else None,
-            eos_token_id=EOS_TOKEN_ID,
             block_hasher=block_hasher,
         )
         requests.append(request)

tests/v1/distributed/test_async_llm_dp.py

@@ -12,6 +12,7 @@ from vllm import SamplingParams
 from vllm.config import VllmConfig
 from vllm.engine.arg_utils import AsyncEngineArgs
 from vllm.inputs import PromptType
+from vllm.outputs import RequestOutput
 from vllm.platforms import current_platform
 from vllm.sampling_params import RequestOutputKind
 from vllm.v1.engine.async_llm import AsyncLLM
@@ -181,3 +182,145 @@ async def test_load(
             assert slogger.finished_req_count > NUM_REQUESTS // (DP_SIZE + 1), (
                 f"requests are imbalanced: {stats_loggers}"
             )
+
+
+# =============================================================================
+# DP Pause/Resume Tests
+# =============================================================================
+
+DP_PAUSE_MODEL = "hmellor/tiny-random-LlamaForCausalLM"
+DP_PAUSE_PROMPT = "This is a test of data parallel pause"
+
+
+@pytest.mark.asyncio
+async def test_dp_pause_resume_basic():
+    """Pausing from the client (one call) pauses all DP ranks; resume clears it."""
+    if current_platform.is_rocm():
+        pytest.skip("DP pause tests use mp backend only")
+    with ExitStack() as after:
+        engine_args = AsyncEngineArgs(
+            model=DP_PAUSE_MODEL,
+            enforce_eager=True,
+            tensor_parallel_size=int(os.getenv("TP_SIZE", 1)),
+            data_parallel_size=DP_SIZE,
+            data_parallel_backend="mp",
+        )
+        engine = AsyncLLM.from_engine_args(engine_args)
+        after.callback(engine.shutdown)
+
+        assert not await engine.is_paused()
+        await engine.pause_generation(mode="abort")
+        assert await engine.is_paused()
+        await engine.resume_generation()
+        assert not await engine.is_paused()
+
+        # Engine still works after resume
+        sampling_params = SamplingParams(max_tokens=5)
+        async for out in engine.generate(
+            request_id="after-resume",
+            prompt=DP_PAUSE_PROMPT,
+            sampling_params=sampling_params,
+        ):
+            pass
+        assert out.finished
+
+
+@pytest.mark.asyncio
+async def test_dp_pause_abort():
+    """Pause with abort from one client aborts in-flight requests on all DP ranks."""
+    if current_platform.is_rocm():
+        pytest.skip("DP pause tests use mp backend only")
+    with ExitStack() as after:
+        engine_args = AsyncEngineArgs(
+            model=DP_PAUSE_MODEL,
+            enforce_eager=True,
+            tensor_parallel_size=int(os.getenv("TP_SIZE", 1)),
+            data_parallel_size=DP_SIZE,
+            data_parallel_backend="mp",
+        )
+        engine = AsyncLLM.from_engine_args(engine_args)
+        after.callback(engine.shutdown)
+
+        # Start several requests so they are distributed across ranks
+        sampling_params = SamplingParams(max_tokens=500, ignore_eos=True)
+        num_requests = 4
+        outputs_by_id: dict[str, list[RequestOutput]] = {}
+
+        async def gen(rid: str):
+            out_list: list[RequestOutput] = []
+            outputs_by_id[rid] = out_list
+            async for out in engine.generate(
+                request_id=rid,
+                prompt=DP_PAUSE_PROMPT,
+                sampling_params=sampling_params,
+            ):
+                out_list.append(out)
+            return out_list[-1] if out_list else None
+
+        tasks = [asyncio.create_task(gen(f"req-{i}")) for i in range(num_requests)]
+        # Wait for some tokens on at least one request
+        while not any(len(o) >= 2 for o in outputs_by_id.values()):
+            await asyncio.sleep(0.02)
+
+        await engine.pause_generation(mode="abort")
+
+        finals = await asyncio.gather(*tasks)
+        for i, final in enumerate(finals):
+            assert final is not None, f"req-{i} had no output"
+            assert final.finished
+            assert final.outputs[0].finish_reason == "abort"
+
+        assert await engine.is_paused()
+        await engine.resume_generation()
+        assert not await engine.is_paused()
+
+        # New request completes after resume
+        async for out in engine.generate(
+            request_id="after-abort",
+            prompt=DP_PAUSE_PROMPT,
+            sampling_params=SamplingParams(max_tokens=5),
+        ):
+            pass
+        assert out.finished
+        assert not engine.output_processor.has_unfinished_requests()
+
+
+@pytest.mark.asyncio
+async def test_dp_pause_keep_then_resume():
+    """Pause with keep queues new requests; resume allows them to run."""
+    if current_platform.is_rocm():
+        pytest.skip("DP pause tests use mp backend only")
+    with ExitStack() as after:
+        engine_args = AsyncEngineArgs(
+            model=DP_PAUSE_MODEL,
+            enforce_eager=True,
+            tensor_parallel_size=int(os.getenv("TP_SIZE", 1)),
+            data_parallel_size=DP_SIZE,
+            data_parallel_backend="mp",
+        )
+        engine = AsyncLLM.from_engine_args(engine_args)
+        after.callback(engine.shutdown)
+
+        await engine.pause_generation(mode="keep")
+        assert await engine.is_paused()
+
+        request_done = asyncio.Event()
+
+        async def gen():
+            async for out in engine.generate(
+                request_id="queued-keep",
+                prompt=DP_PAUSE_PROMPT,
+                sampling_params=SamplingParams(max_tokens=5),
+            ):
+                pass
+            request_done.set()
+            return out
+
+        task = asyncio.create_task(gen())
+        await asyncio.sleep(0.2)
+        assert not request_done.is_set()
+
+        await engine.resume_generation()
+        final = await asyncio.wait_for(task, timeout=10.0)
+        assert final.finished
+        assert not await engine.is_paused()

tests/v1/e2e/test_mamba_prefix_cache.py

@@ -11,8 +11,10 @@ import datasets
 import pytest
 import torch
 
+from tests.utils import create_new_process_for_each_test
 from vllm import LLM, SamplingParams, TokensPrompt
 from vllm.config import CacheConfig
+from vllm.distributed import cleanup_dist_env_and_memory
 from vllm.model_executor.layers.mamba.mamba_utils import MambaStateCopyFunc
 from vllm.sequence import IntermediateTensors
 from vllm.v1.attention.backends.utils import CommonAttentionMetadata
@@ -74,14 +76,11 @@ def get_fake_sample_fn() -> SamplerOutput:
                 ),
                 logprobs_tensors=None,
             )
-        num_sampled_tokens = spec_decode_metadata.cu_num_sampled_tokens[0].item() + 1
         accpeted_tokens = prompt_token_ids[
             first_token_id_index : first_token_id_index
             + min(num_accepted_tokens, logits.shape[0])
         ]
-        sampled_token_ids = accpeted_tokens + [-1] * (
-            num_sampled_tokens - len(accpeted_tokens)
-        )
+        sampled_token_ids = accpeted_tokens
         return SamplerOutput(
             sampled_token_ids=torch.tensor(
                 [sampled_token_ids], device="cuda", dtype=torch.int32
@@ -103,6 +102,7 @@ def get_fake_propose_draft_token_ids_fn():
         aux_hidden_states: list[torch.Tensor] | None,
         spec_decode_metadata: SpecDecodeMetadata | None,
         common_attn_metadata: CommonAttentionMetadata,
+        slot_mappings: dict[str, torch.Tensor] | list[dict[str, torch.Tensor]] | None,
     ) -> list[list[int]]:
         num_computed_tokens_cpu_tensor = self.input_batch.num_computed_tokens_cpu_tensor
         num_computed_tokens = num_computed_tokens_cpu_tensor[0].item()
@@ -121,7 +121,24 @@ def get_fake_propose_draft_token_ids_fn():
                 first_token_id_index : first_token_id_index + num_speculative_tokens
             ]
         ]
-        return proposed_draft_token_ids
+
+        next_token_ids = torch.tensor(
+            prompt_token_ids[
+                first_token_id_index - 1 : first_token_id_index
+                - 1
+                + num_accepted_tokens
+            ],
+            device="cuda",
+            dtype=torch.int32,
+        )
+
+        valid_sampled_tokens_count = torch.tensor(
+            [num_accepted_tokens], device="cuda", dtype=torch.int32
+        )
+
+        self._copy_valid_sampled_token_count(next_token_ids, valid_sampled_tokens_count)
+
+        return torch.tensor(proposed_draft_token_ids, device="cuda", dtype=torch.int32)
 
     return fake_propose_draft_token_ids_fn
 
@@ -181,6 +198,7 @@ mamba_kv_cache_dict = {}
 
 def get_fake_execute_model_fn(original_execute_model_fn: Callable):
     last_num_computed_tokens = 0
+    num_prompt_tokens = None
 
     def fake_execute_model_fn(
         self: GPUModelRunner,
@@ -198,10 +216,30 @@ def get_fake_execute_model_fn(original_execute_model_fn: Callable):
             mamba_group_id
         ].layer_names[0]
         nonlocal last_num_computed_tokens
+        nonlocal num_prompt_tokens
+
+        if (
+            len(scheduler_output.scheduled_new_reqs) > 0
+            and scheduler_output.scheduled_new_reqs[0].prompt_token_ids is not None
+        ):
+            # record number of prompt tokens
+            num_prompt_tokens = len(
+                scheduler_output.scheduled_new_reqs[0].prompt_token_ids
+            )
+
         if len(scheduler_output.scheduled_cached_reqs.req_ids) > 0:
             num_computed_tokens = (
                 scheduler_output.scheduled_cached_reqs.num_computed_tokens[0]
             )
+            if (
+                self.num_spec_tokens
+                and num_prompt_tokens is not None
+                and num_computed_tokens > num_prompt_tokens
+            ):
+                # NOTE (tdoublep) with async scheduling, the scheduler does not have an
+                # accurate measure of the number of computed tokens; we need to subtract
+                # the number of reject tokens from the previous timestep.
+                num_computed_tokens -= num_speculative_tokens + 1 - num_accepted_tokens
             if (
                 num_computed_tokens // BLOCK_SIZE
                 > last_num_computed_tokens // BLOCK_SIZE
@@ -401,6 +439,9 @@ def _run_ref_mamba_state_worker():
         }
         torch.save(cpu_state_ref, "mamba_kv_cache_dict_ref.pth")
         mamba_kv_cache_dict.clear()
+        del engine
+        torch.cuda.empty_cache()
+        cleanup_dist_env_and_memory()
     except Exception:
         traceback.print_exc()
         raise
@@ -473,10 +514,7 @@ def apply_patch(monkeypatch: pytest.MonkeyPatch):
     monkeypatch.setattr(mamba_utils, "do_mamba_copy_block", fake_copy_fn)
 
 
-@pytest.mark.skip(
-    reason="Skipping test_mamba_prefix_cache because it is based on spec "
-    "decode which is not allowed now."
-)
+@create_new_process_for_each_test()
 def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
     run_ref_mamba_state_in_subprocess()
     apply_patch(monkeypatch)
@@ -490,9 +528,9 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
             step_actions=[
                 StepAction(0, 554, [1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(554, 4, [], (-1, -1), (-1, -1)),
-                StepAction(555, 4, [], (-1, -1), (-1, -1)),
+                StepAction(555, 4, [1, 1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(556, 4, [], (-1, -1), (-1, -1)),
-                StepAction(557, 4, [1, 1, 1, 1, 1], (0, 1), (-1, -1)),
+                StepAction(557, 4, [], (0, 1), (-1, -1)),
                 StepAction(558, 4, [], (-1, -1), (-1, -1)),
                 StepAction(559, 4, [], (-1, -1), (1, 0)),
                 StepAction(560, 4, [], (-1, -1), (-1, -1)),
@@ -507,8 +545,8 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
             step_actions=[
                 StepAction(0, 554, [1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(554, 4, [], (-1, -1), (-1, -1)),
-                StepAction(556, 4, [], (-1, -1), (-1, -1)),
-                StepAction(558, 4, [1, 1, 1, 1, 1], (1, 1), (2, 0)),
+                StepAction(556, 4, [1, 1, 1, 1, 1], (-1, -1), (-1, -1)),
+                StepAction(558, 4, [], (1, 1), (2, 0)),
                 StepAction(560, 4, [], (-1, -1), (-1, -1)),
                 StepAction(562, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
             ],
@@ -523,7 +561,8 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
                 StepAction(555, 4, [], (-1, -1), (-1, -1)),
                 StepAction(557, 4, [1, 1, 1, 1, 1], (1, 1), (-1, -1)),
                 StepAction(559, 4, [], (-1, -1), (1, 0)),
-                StepAction(561, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
+                StepAction(561, 4, [], (-1, -1), (-1, -1)),
+                StepAction(563, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
             ],
         ),
         "accept_3_1": TestConfig(
@@ -533,9 +572,10 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
             step_actions=[
                 StepAction(0, 553, [1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(553, 4, [], (-1, -1), (-1, -1)),
-                StepAction(556, 4, [], (-1, -1), (-1, -1)),
-                StepAction(559, 4, [1, 1, 1, 1, 1], (2, 1), (1, 0)),
-                StepAction(562, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
+                StepAction(556, 4, [1, 1, 1, 1, 1], (-1, -1), (-1, -1)),
+                StepAction(559, 4, [], (2, 1), (1, 0)),
+                StepAction(562, 4, [], (-1, -1), (-1, -1)),
+                StepAction(565, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
             ],
         ),
         "accept_3_2": TestConfig(
@@ -558,7 +598,8 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
                 StepAction(0, 555, [1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(555, 4, [], (-1, -1), (-1, -1)),
                 StepAction(558, 4, [1, 1, 1, 1, 1], (2, 1), (2, 0)),
-                StepAction(561, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
+                StepAction(561, 4, [], (-1, -1), (-1, -1)),
+                StepAction(564, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
             ],
         ),
         "accept_4_1": TestConfig(
@@ -569,8 +610,8 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
                 StepAction(0, 553, [1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(553, 4, [], (-1, -1), (-1, -1)),
                 StepAction(557, 4, [1, 1, 1, 1, 1], (3, 1), (3, 0)),
-                StepAction(561, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
-                StepAction(565, 4, [], (-1, -1), (-1, -1)),
+                StepAction(561, 4, [], (-1, -1), (-1, -1)),
+                StepAction(565, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
             ],
         ),
         "accept_4_2": TestConfig(
@@ -581,8 +622,8 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
                 StepAction(0, 554, [1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(554, 4, [], (-1, -1), (-1, -1)),
                 StepAction(558, 4, [1, 1, 1, 1, 1], (3, 1), (2, 0)),
-                StepAction(562, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
-                StepAction(566, 4, [], (-1, -1), (-1, -1)),
+                StepAction(562, 4, [], (-1, -1), (-1, -1)),
+                StepAction(566, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
             ],
         ),
         "accept_4_3": TestConfig(
@@ -593,7 +634,8 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
                 StepAction(0, 555, [1, 1, 1, 1], (-1, -1), (-1, -1)),
                 StepAction(555, 4, [], (-1, -1), (-1, -1)),
                 StepAction(559, 4, [1, 1, 1, 1, 1], (3, 1), (1, 0)),
-                StepAction(563, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
+                StepAction(563, 4, [], (-1, -1), (-1, -1)),
+                StepAction(567, 4, [0, 1, 1, 1, 1], (-1, -1), (-1, -1)),
             ],
         ),
         "accept_4_4": TestConfig(
@@ -762,3 +804,6 @@ def test_mamba_prefix_cache(monkeypatch: pytest.MonkeyPatch):
         mamba_state_ref = torch.load("mamba_kv_cache_dict_ref.pth")
         check_mamba_state_equal(mamba_state_ref, mamba_kv_cache_dict, keys_to_check)
         mamba_kv_cache_dict.clear()
+    del engine
+    torch.cuda.empty_cache()
+    cleanup_dist_env_and_memory()

tests/v1/e2e/test_streaming_input.py

@@ -19,7 +19,7 @@ import pytest
 import pytest_asyncio
 
 from vllm import SamplingParams
-from vllm.inputs import StreamingInput
+from vllm.engine.protocol import StreamingInput
 from vllm.outputs import RequestOutput
 from vllm.platforms import current_platform
 from vllm.sampling_params import RequestOutputKind

tests/v1/engine/test_async_llm.py

@@ -708,9 +708,7 @@ async def test_pause_resume_basic():
         # Test all modes with no requests in flight
         for mode in ("abort", "wait", "keep"):
             await engine.pause_generation(mode=mode)
-            # "keep" only freezes the scheduler; it does not set _paused
-            if mode != "keep":
-                assert await engine.is_paused()
+            assert await engine.is_paused()
             await engine.resume_generation()
             assert not await engine.is_paused()
 
@@ -808,6 +806,53 @@ async def test_pause_abort():
         assert final_output2.finished
 
 
+@pytest.mark.asyncio
+async def test_pause_then_abort_queued_request():
+    """Test that aborting a request that was submitted while paused (in
+    _paused_adds_queue) aborts it and notifies the client; the request does
+    not run after resume.
+    """
+    with ExitStack() as after:
+        with set_default_torch_num_threads(1):
+            engine = AsyncLLM.from_engine_args(TEXT_ENGINE_ARGS)
+        after.callback(engine.shutdown)
+
+        request_id = "abort-queued-request"
+        sampling_params = SamplingParams(max_tokens=20, ignore_eos=True)
+        outputs: list[RequestOutput] = []
+
+        # Pause first so the next add goes to _paused_adds_queue
+        await engine.pause_generation(mode="keep")
+        assert await engine.is_paused()
+
+        async def gen():
+            async for out in engine.generate(
+                request_id=request_id,
+                prompt=TEXT_PROMPT,
+                sampling_params=sampling_params,
+            ):
+                outputs.append(out)
+            return outputs[-1] if outputs else None
+
+        gen_task = asyncio.create_task(gen())
+
+        # Give the request time to reach the engine and sit in _paused_adds_queue
+        await asyncio.sleep(0.2)
+
+        # Abort the queued request
+        await engine.abort(request_id, internal=False)
+
+        # Resume so the engine can process and deliver the abort output
+        await engine.resume_generation()
+
+        final_output = await asyncio.wait_for(gen_task, timeout=10.0)
+        assert final_output is not None
+        assert final_output.finished
+        assert final_output.outputs[0].finish_reason == "abort"
+        # Request was never run, so no tokens
+        assert len(final_output.outputs[0].token_ids) == 0
+
+
 @pytest.mark.asyncio
 async def test_pause_wait():
     """Test that mode='wait' waits for in-flight requests to complete."""

tests/v1/engine/test_engine_core.py

@@ -54,7 +54,6 @@ def make_request() -> EngineCoreRequest:
         mm_features=None,
         sampling_params=SamplingParams(),
         pooling_params=None,
-        eos_token_id=None,
         arrival_time=time.time(),
         lora_request=None,
         cache_salt=None,

tests/v1/engine/test_engine_core_client.py

@@ -8,6 +8,7 @@ import os
 import signal
 import time
 import uuid
+from concurrent.futures import Future
 from dataclasses import dataclass
 from threading import Thread
 from types import SimpleNamespace
@@ -69,7 +70,6 @@ def make_request(
         mm_features=None,
         sampling_params=params,
         pooling_params=None,
-        eos_token_id=None,
         arrival_time=time.time(),
         lora_request=None,
         cache_salt=None,
@@ -279,6 +279,24 @@ def echo_dc_nested(
     return structures.get(structure_type, val)
 
 
+def future_echo(self, value: Any, num_wait_loops: int = 2) -> Future:
+    """Utility that returns a Future completed by a per_step_hook after
+    num_wait_loops engine steps (tests deferred utility path).
+    """
+    future: Future = Future()
+    remaining = [num_wait_loops]
+
+    def _step(engine: EngineCore) -> bool:
+        remaining[0] -= 1
+        if remaining[0] <= 0:
+            future.set_result(value)
+            return True  # remove hook
+        return False
+
+    self.per_step_hooks.add(_step)
+    return future
+
+
 # --- Fixtures for subprocess patching ---
 # These create sitecustomize.py files that patch EngineCore in spawned
 # subprocesses. This is necessary because ROCm requires 'spawn' multiprocessing
@@ -384,6 +402,28 @@ def subprocess_echo_dc_nested_patch(monkeypatch, tmp_path):
     )
 
 
+@pytest.fixture
+def subprocess_future_echo_patch(monkeypatch, tmp_path):
+    """Create sitecustomize.py so spawned subprocesses have future_echo method."""
+    sc = tmp_path / "sitecustomize.py"
+    sc.write_text(
+        "\n".join(
+            [
+                "from concurrent.futures import Future",
+                "from typing import Any",
+                "",
+                "from vllm.v1.engine.core import EngineCore",
+                inspect.getsource(future_echo),
+                "EngineCore.future_echo = future_echo",
+            ]
+        )
+    )
+    monkeypatch.setenv(
+        "PYTHONPATH",
+        os.pathsep.join(filter(None, [str(tmp_path), os.getenv("PYTHONPATH")])),
+    )
+
+
 @create_new_process_for_each_test()
 @pytest.mark.parametrize("multiprocessing_mode", [True, False])
 def test_engine_core_client(
@@ -787,6 +827,48 @@ async def test_engine_core_client_util_method_nested_structures(
             client.shutdown()
 
 
+@pytest.mark.asyncio(loop_scope="function")
+async def test_engine_core_client_future_utility_async(
+    monkeypatch: pytest.MonkeyPatch,
+    subprocess_future_echo_patch,
+):
+    """Test that a utility returning a Future (completed by a per_step_hook
+    after N steps) completes when the future is done (engine uses add_done_callback).
+    """
+    with monkeypatch.context() as m:
+        m.setattr(EngineCore, "future_echo", future_echo, raising=False)
+
+        engine_args = EngineArgs(model=MODEL_NAME, enforce_eager=True)
+        vllm_config = engine_args.create_engine_config(
+            usage_context=UsageContext.UNKNOWN_CONTEXT
+        )
+        executor_class = Executor.get_class(vllm_config)
+
+        with set_default_torch_num_threads(1):
+            client = EngineCoreClient.make_client(
+                multiprocess_mode=True,
+                asyncio_mode=True,
+                vllm_config=vllm_config,
+                executor_class=executor_class,
+                log_stats=True,
+            )
+
+        try:
+            core_client: AsyncMPClient = client
+
+            # Completes after 2 engine steps (num_wait_loops=2)
+            result = await core_client.call_utility_async(
+                "future_echo", "future_result", 2
+            )
+            assert result == "future_result"
+
+            # None is a valid result (num_wait_loops=0 → completes on first step)
+            result = await core_client.call_utility_async("future_echo", None, 0)
+            assert result is None
+        finally:
+            client.shutdown()
+
+
 @pytest.mark.parametrize(
     "multiprocessing_mode,publisher_config",
     [(True, "tcp"), (False, "inproc")],

tests/v1/engine/test_fast_incdec_prefix_err.py

@@ -32,7 +32,6 @@ def test_fast_inc_detok_invalid_utf8_err_case():
         mm_features=None,
         sampling_params=params,
         pooling_params=None,
-        eos_token_id=None,
         arrival_time=0.0,
         lora_request=None,
         cache_salt=None,

tests/v1/engine/test_output_processor.py

@@ -66,7 +66,6 @@ def test_incremental_detokenization(
             external_req_id=f"request-{idx}",
             prompt_token_ids=prompt_tokens,
             mm_features=None,
-            eos_token_id=None,
             arrival_time=0,
             lora_request=None,
             cache_salt=None,
@@ -487,7 +486,6 @@ def test_logprobs_processor(
             external_req_id=request_id_list[idx],
             prompt_token_ids=prompt_tokens,
             mm_features=None,
-            eos_token_id=None,
             arrival_time=0,
             lora_request=None,
             cache_salt=None,
@@ -663,6 +661,19 @@ def test_stop_token(
     prompt_string = dummy_test_vectors.prompt_strings[0]
     prompt_tokens = dummy_test_vectors.prompt_tokens[0]
 
+    sampling_params = SamplingParams(
+        skip_special_tokens=False,
+        spaces_between_special_tokens=False,
+        output_kind=RequestOutputKind.DELTA,
+        stop=[],
+        stop_token_ids=stop_token_ids,
+        include_stop_str_in_output=include_stop_str_in_output,
+        logprobs=num_sample_logprobs,
+        prompt_logprobs=None,
+        ignore_eos=ignore_eos,
+    )
+    sampling_params.update_from_generation_config({}, eos_token_id)
+
     # Make request.
     request_id = "request-0"
     request = EngineCoreRequest(
@@ -670,22 +681,11 @@ def test_stop_token(
         external_req_id=request_id + "-ext",
         prompt_token_ids=prompt_tokens,
         mm_features=None,
-        eos_token_id=eos_token_id,
         arrival_time=0,
         lora_request=None,
         cache_salt=None,
         data_parallel_rank=None,
-        sampling_params=SamplingParams(
-            skip_special_tokens=False,
-            spaces_between_special_tokens=False,
-            output_kind=RequestOutputKind.DELTA,
-            stop=[],
-            stop_token_ids=stop_token_ids,
-            include_stop_str_in_output=include_stop_str_in_output,
-            logprobs=num_sample_logprobs,
-            prompt_logprobs=None,
-            ignore_eos=ignore_eos,
-        ),
+        sampling_params=sampling_params,
         pooling_params=None,
     )
 
@@ -693,9 +693,8 @@ def test_stop_token(
         tokens_list=[generation_tokens],
         generated_logprobs_raw=[generation_logprobs] if do_logprobs else None,
         prompt_logprobs_raw=None,
-        eos_token_id=eos_token_id,
-        stop_token_ids=stop_token_ids,
-        ignore_eos=ignore_eos,
+        eos_token_id=sampling_params.eos_token_id,
+        stop_token_ids=sampling_params.stop_token_ids,
         request_ids=[request.request_id],
     )
 
@@ -775,7 +774,6 @@ def test_stop_string(
             external_req_id=request_id_list[idx],
             prompt_token_ids=prompt_tokens,
             mm_features=None,
-            eos_token_id=None,
             arrival_time=0,
             lora_request=None,
             cache_salt=None,
@@ -907,7 +905,6 @@ def test_iteration_stats(dummy_test_vectors):
             external_req_id=f"request-{idx}-ext",
             prompt_token_ids=prompt_tokens,
             mm_features=None,
-            eos_token_id=None,
             arrival_time=0,
             lora_request=None,
             cache_salt=None,
@@ -994,7 +991,6 @@ def test_lora_request_tracking(log_stats: bool, dummy_test_vectors):
             external_req_id=f"request-{idx}",
             prompt_token_ids=prompt_tokens,
             mm_features=None,
-            eos_token_id=None,
             arrival_time=0,
             lora_request=lora_assignments[idx],
             cache_salt=None,
@@ -1315,7 +1311,6 @@ def test_abort_requests(runner: str, abort_by: str, dummy_test_vectors):
             external_req_id=f"external-{idx}",
             prompt_token_ids=prompt_tokens,
             mm_features=None,
-            eos_token_id=None,
             arrival_time=0,
             lora_request=None,
             cache_salt=None,

tests/v1/engine/test_parallel_sampling.py

@@ -76,7 +76,6 @@ def make_request(sampling_params: SamplingParams) -> EngineCoreRequest:
         mm_features=None,
         sampling_params=sampling_params,
         pooling_params=None,
-        eos_token_id=None,
         arrival_time=0.0,
         lora_request=None,
         cache_salt=None,

tests/v1/engine/utils.py

@@ -342,7 +342,6 @@ class MockEngineCore:
         prompt_logprobs_raw: list[LogprobsTensors] | None = None,
         eos_token_id: int | None = None,
         stop_token_ids: list[int] | None = None,
-        ignore_eos: bool = False,
         request_ids: list[str] | None = None,
     ) -> None:
         self.num_requests = len(tokens_list)
@@ -355,7 +354,6 @@ class MockEngineCore:
         self.request_finished = [False for _ in range(self.num_requests)]
         self.eos_token_id = eos_token_id
         self.stop_token_ids = stop_token_ids
-        self.ignore_eos = ignore_eos
         self.request_ids = (
             request_ids
             if request_ids is not None
@@ -400,7 +398,7 @@ class MockEngineCore:
                 if token_idx == len(token_ids) - 1:
                     output.finish_reason = FinishReason.LENGTH
                     self.request_finished[req_idx] = True
-                if not self.ignore_eos and new_token_id == self.eos_token_id:
+                if new_token_id == self.eos_token_id:
                     output.finish_reason = FinishReason.STOP
                     self.request_finished[req_idx] = True
                 if new_token_id in (self.stop_token_ids or ()):

tests/v1/kv_connector/unit/test_decode_bench_connector.py

@@ -93,12 +93,14 @@ class DecodeBenchTestRunner:
         """Create a new request with given token IDs."""
         self.req_id += 1
 
+        sampling_params = SamplingParams(max_tokens=100)
+        sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
+
         req = Request(
             request_id=str(self.req_id),
             prompt_token_ids=token_ids,
-            sampling_params=SamplingParams(max_tokens=100),
+            sampling_params=sampling_params,
             pooling_params=None,
-            eos_token_id=EOS_TOKEN_ID,
             block_hasher=self._block_hasher,
         )
 

tests/v1/kv_connector/unit/test_lmcache_integration.py

@@ -142,12 +142,14 @@ def test_request_interface():
     from vllm.sampling_params import SamplingParams
     from vllm.v1.request import Request
 
+    sampling_params = SamplingParams(max_tokens=10)
+    sampling_params.update_from_generation_config({}, eos_token_id=100)
+
     req = Request(
         request_id="test_request",
         prompt_token_ids=[1, 2, 3],
-        sampling_params=SamplingParams(max_tokens=10),
+        sampling_params=sampling_params,
         pooling_params=None,
-        eos_token_id=100,
         lora_request=None,
     )
     assumes(req, "mm_features", is_instance_of=(list, NoneType))

tests/v1/kv_connector/unit/test_offloading_connector.py

@@ -226,12 +226,14 @@ class RequestRunner:
     def new_request(self, token_ids: list[int]):
         self.req_id += 1
 
+        sampling_params = SamplingParams(max_tokens=1000)
+        sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
+
         req = Request(
             request_id=str(self.req_id),
             prompt_token_ids=token_ids,
-            sampling_params=SamplingParams(max_tokens=1000),
+            sampling_params=sampling_params,
             pooling_params=None,
-            eos_token_id=EOS_TOKEN_ID,
             block_hasher=self._block_hasher,
         )
 

tests/v1/kv_connector/unit/utils.py

@@ -212,6 +212,7 @@ def create_request(
 
     max_tokens = 1 if do_remote_decode else max_tokens
     sampling_params = SamplingParams(max_tokens=max_tokens)
+    sampling_params.update_from_generation_config({}, EOS_TOKEN_ID)
 
     common_prefix = [1] * common_prefix_len if common_prefix_len > 0 else []
     suffix = [i * request_id for i in range(num_tokens - common_prefix_len)]
@@ -223,7 +224,6 @@ def create_request(
         sampling_params=sampling_params,
         pooling_params=None,
         mm_features=None,
-        eos_token_id=EOS_TOKEN_ID,
         block_hasher=get_request_block_hasher(block_size, hash_fn),
     )
     req.kv_transfer_params = kv_transfer_params

tests/v1/streaming_input/test_async_llm_streaming.py

@@ -7,7 +7,7 @@ from unittest.mock import AsyncMock, MagicMock
 
 import pytest
 
-from vllm.inputs import StreamingInput
+from vllm.engine.protocol import StreamingInput
 from vllm.outputs import RequestOutput
 from vllm.sampling_params import RequestOutputKind, SamplingParams
 from vllm.v1.engine.async_llm import AsyncLLM

tests/v1/streaming_input/test_scheduler_streaming.py

@@ -43,7 +43,6 @@ class DummyRequest(Request):
                 stop_token_ids=[STOP_TOKEN], max_tokens=max_tokens
             ),
             pooling_params=None,
-            eos_token_id=None,
             mm_features=mm_features,
             resumable=resumable,
         )

tests/v1/structured_output/test_backend_guidance.py

@@ -83,6 +83,7 @@ def test_grammar_bitmask_with_specdec():
             ),
         )
         sampling_params.structured_outputs._backend = "guidance"
+        sampling_params.update_from_generation_config({}, tokenizer.eos_token_id)
 
         my_req_id = f"my_req_id_{i}"
         request = Request(
@@ -90,7 +91,6 @@ def test_grammar_bitmask_with_specdec():
             prompt_token_ids=prompt[:i],
             sampling_params=sampling_params,
             pooling_params=None,
-            eos_token_id=tokenizer.eos_token_id,
         )
 
         structured_output_manager.grammar_init(request)
@@ -147,13 +147,13 @@ def test_grammar_init_async_and_sync(async_grammar):
         ),
     )
     sampling_params.structured_outputs._backend = "guidance"
+    sampling_params.update_from_generation_config({}, tokenizer.eos_token_id)
 
     request = Request(
         "test_request",
         prompt_token_ids=prompt,
         sampling_params=sampling_params,
         pooling_params=None,
-        eos_token_id=tokenizer.eos_token_id,
     )
 
     structured_output_manager.grammar_init(request)

tests/v1/structured_output/test_gptoss_structural_tags.py

@@ -25,6 +25,7 @@ class TestGptOssReasoningParser:
         """Create a mock tokenizer for testing."""
         tokenizer = Mock()
         tokenizer.encode = Mock(return_value=[1, 2, 3, 4, 5])
+        tokenizer.vocab = {"<|end|>": 6}
         return tokenizer
 
     @pytest.fixture

vllm/benchmarks/datasets.py

@@ -380,7 +380,7 @@ def gen_prompt_decode_to_target_len(
     max_retry: int = 10,
     add_special_tokens: bool = False,
     rng: np.random.Generator | None = None,
-) -> tuple[str, list[int]]:
+) -> tuple[str, list[int], int]:
     """
     Ensure decoded-then-encoded prompt length matches the target token length.
 
@@ -392,7 +392,9 @@ def gen_prompt_decode_to_target_len(
     [6880, 6881] -> ['Ġcalls', 'here'] ->
     [1650, 939, 486] -> ['Ġcall', 'sh', 'ere']
 
-    Returns a tuple of the final prompt string and the adjusted token sequence.
+    Returns a tuple of the final prompt string, the adjusted token sequence,
+    and the token mismatch (final_len - target_token_len) if the retry budget
+    is exhausted.
     """
     remain_num_try = max_retry
     token_mismatch = 0
@@ -499,7 +501,7 @@ class RandomDataset(BenchmarkDataset):
         allowed_tokens = np.array(list(set(all_tokens) - set(prohibited_tokens)))
 
         # Generate prefix once
-        prefix_token_ids = self.get_prefix(allowed_tokens, prefix_len)
+        prefix_token_ids = self.get_prefix(tokenizer, allowed_tokens, prefix_len)
 
         requests = []
         token_mismatch_total = 0
@@ -554,19 +556,36 @@ class RandomDataset(BenchmarkDataset):
 
     def get_prefix(
         self,
+        tokenizer: TokenizerLike,
         allowed_tokens: np.ndarray,
         prefix_len: int,
     ) -> list[int]:
         """
         Get the prefix for the dataset.
         """
-        return (
-            allowed_tokens[
-                self._rng.integers(0, len(allowed_tokens), size=prefix_len)
-            ].tolist()
-            if prefix_len > 0
-            else []
+        if prefix_len <= 0:
+            return []
+
+        prefix_tokens = allowed_tokens[
+            self._rng.integers(0, len(allowed_tokens), size=prefix_len)
+        ].tolist()
+        _, adjusted_tokens, token_mismatch = gen_prompt_decode_to_target_len(
+            tokenizer=tokenizer,
+            token_sequence=prefix_tokens,
+            target_token_len=prefix_len,
+            add_special_tokens=False,
+            rng=self._rng,
         )
+        if token_mismatch != 0:
+            sign = "more" if token_mismatch > 0 else "fewer"
+            logger.warning(
+                "Prefix tokenization produced %d %s tokens than expected "
+                "after decoding and re-encoding. This is expected due to "
+                "the imperfect nature of the sampling procedure",
+                abs(token_mismatch),
+                sign,
+            )
+        return adjusted_tokens
 
     def get_sampling_params(
         self,
@@ -1128,7 +1147,7 @@ class RandomMultiModalDataset(RandomDataset):
             "Sampling from %d out of %d (vocab size)", len(allowed_tokens), vocab_size
         )
         # Generate prefix once
-        prefix_token_ids = self.get_prefix(allowed_tokens, prefix_len)
+        prefix_token_ids = self.get_prefix(tokenizer, allowed_tokens, prefix_len)
         # Add synthetic multimodal items to each request
         mm_requests = []
         token_mismatch_total = 0

vllm/benchmarks/mm_processor.py

@@ -18,7 +18,7 @@ import dataclasses
 import json
 import time
 from datetime import datetime
-from typing import Any
+from typing import TYPE_CHECKING, Any
 
 import numpy as np
 
@@ -28,9 +28,6 @@ from vllm.benchmarks.datasets import (
 )
 from vllm.benchmarks.throughput import get_requests
 from vllm.engine.arg_utils import EngineArgs
-from vllm.multimodal.processing.context import (
-    get_timing_stats_from_engine_client,
-)
 from vllm.utils.gc_utils import freeze_gc_heap
 from vllm.utils.import_utils import PlaceholderModule
 
@@ -39,16 +36,103 @@ try:
 except ImportError:
     pd = PlaceholderModule("pandas")
 
+if TYPE_CHECKING:  # Avoid having to mock during docs build
+    from vllm.v1.engine.llm_engine import LLMEngine
+else:
+    LLMEngine = object
+
+
+def get_timing_stats_from_engine(llm_engine: LLMEngine) -> dict[str, dict[str, float]]:
+    """
+    Get all multimodal timing stats from the LLM engine.
+
+    Collects both preprocessing stats (HF processor, hashing, cache lookup,
+    prompt update) and encoder forward pass timing, merged by request_id.
+
+    Args:
+        llm_engine: The LLM engine (has input_processor and workers).
+
+    Returns:
+        Dictionary mapping request_id to merged stats dict containing
+        both preprocessing and encoder timing metrics.
+
+    Example:
+        {
+            'request-123': {
+                'hf_processor_time': 0.45,
+                'hashing_time': 0.02,
+                'cache_lookup_time': 0.01,
+                'prompt_update_time': 0.03,
+                'preprocessor_total_time': 0.51,
+                'encoder_forward_time': 0.23,
+                'num_encoder_calls': 1
+            }
+        }
+    """
+    observability_config = llm_engine.vllm_config.observability_config
+    if not observability_config or not observability_config.enable_mm_processor_stats:
+        return {}
+
+    renderer = llm_engine.renderer
+    mm_processor = renderer.get_mm_processor()
+    preprocessing_stats = mm_processor.info.ctx.get_all_timing_stats()
+
+    encoder_stats = dict[str, dict[str, float]]()
+    for worker_stats in llm_engine.collective_rpc("get_encoder_timing_stats"):
+        if not worker_stats:
+            continue
+
+        for request_id, stats_dict in worker_stats.items():
+            if request_id not in encoder_stats:
+                encoder_stats[request_id] = dict(stats_dict)
+            else:
+                # Aggregate timing metrics across workers
+                current_time = encoder_stats[request_id].get(
+                    "encoder_forward_time", 0.0
+                )
+                new_time = stats_dict.get("encoder_forward_time", 0.0)
+                encoder_stats[request_id]["encoder_forward_time"] = max(
+                    current_time, new_time
+                )
+
+                current_calls = encoder_stats[request_id].get("num_encoder_calls", 0)
+                new_calls = stats_dict.get("num_encoder_calls", 0)
+                encoder_stats[request_id]["num_encoder_calls"] = max(
+                    current_calls, new_calls
+                )
+
+    merged_stats = dict[str, dict[str, float]]()
+
+    for request_id, prep_dict in preprocessing_stats.items():
+        merged_stats[request_id] = dict(prep_dict)
+
+    for request_id, enc_dict in encoder_stats.items():
+        if request_id in merged_stats:
+            merged_stats[request_id].update(enc_dict)
+            continue
+
+        # In V1 engine, the request_id in encoder_stats has a suffix
+        # appended to the original request_id (which is used in
+        # preprocessing_stats).
+        # We try to strip the suffix to find the matching request.
+        possible_original_id = request_id.rpartition("-")[0]
+        if possible_original_id and possible_original_id in merged_stats:
+            merged_stats[possible_original_id].update(enc_dict)
+        else:
+            merged_stats[request_id] = dict(enc_dict)
+
+    return merged_stats
+
 
 def collect_mm_processor_stats(
-    llm_engine: Any,
+    llm_engine: LLMEngine,
     num_warmup_reqs: int = 0,
 ) -> dict[str, list[float]]:
     """
     Collect multimodal processor timing stats.
     Returns a dictionary mapping stage names to lists of timing values (in seconds).
     """
-    all_stats = get_timing_stats_from_engine_client(llm_engine)
+    all_stats = get_timing_stats_from_engine(llm_engine)
 
     stat_keys = [
         "hf_processor_time",

vllm/config/cache.py

@@ -101,6 +101,17 @@ class CacheConfig:
     Note that this requires fast CPU-GPU interconnect, as part of the model is
     loaded from CPU memory to GPU memory on the fly in each model forward pass.
     """
+    cpu_offload_params: set[str] = Field(default_factory=set)
+    """ The set of parameter name segments to target for CPU offloading.
+    Unmatched parameters are not offloaded. If this set is empty, parameters
+    are offloaded non-selectively until the memory limit defined by
+    `cpu_offload_gb` is reached.
+    Examples:
+        - For parameter name "mlp.experts.w2_weight":
+            - "experts" or "experts.w2_weight" will match.
+            - "expert" or "w2" will NOT match (must be exact segments).
+    This allows distinguishing parameters like "w2_weight" and "w2_weight_scale".
+    """
     calculate_kv_scales: bool = False
     """This enables dynamic calculation of `k_scale` and `v_scale` when
     kv_cache_dtype is fp8. If `False`, the scales will be loaded from the model

vllm/config/model.py

@@ -357,6 +357,12 @@ class ModelConfig:
         from vllm.config.utils import get_hash_factors, hash_factors
 
         factors = get_hash_factors(self, ignored_factors)
+
+        # NOTE: For some models (e.g, Qwen3-VL), whether the MM code path is enabled
+        # affects the computation graph of the language model, therefore we add it
+        # here early.
+        if self.multimodal_config:
+            factors["language_model_only"] = self.multimodal_config.language_model_only
         return hash_factors(factors)
 
     def _update_nested(

vllm/config/multimodal.py

@@ -219,7 +219,6 @@ class MultiModalConfig:
         the final hidden states.
         """
         factors: list[Any] = [
-            self.language_model_only,
             self.mm_encoder_attn_backend.name
             if self.mm_encoder_attn_backend is not None
             else None,

vllm/config/vllm.py

@@ -648,11 +648,6 @@ class VllmConfig:
                     "`external_launcher` distributed executor backend, but you chose "
                     f"`{executor_backend}`."
                 )
-            if self.cache_config.mamba_cache_mode != "none":
-                raise ValueError(
-                    "Currently, async scheduling is not compatible with "
-                    "prefix caching for Mamba models."
-                )
         elif self.scheduler_config.async_scheduling is None:
             # Enable async scheduling unless there is an incompatible option.
             if (
@@ -685,13 +680,6 @@ class VllmConfig:
                     scope="local",
                 )
                 self.scheduler_config.async_scheduling = False
-            elif self.cache_config.mamba_cache_mode != "none":
-                logger.warning_once(
-                    "Async scheduling is not compatible with "
-                    "prefix caching for Mamba models and will be disabled.",
-                    scope="local",
-                )
-                self.scheduler_config.async_scheduling = False
             else:
                 self.scheduler_config.async_scheduling = True
 
@@ -728,13 +716,13 @@ class VllmConfig:
                 "precision for chunked prefill triton kernels."
             )
 
-        if (
-            self.optimization_level > OptimizationLevel.O0
-            and self.model_config is not None
-            and self.model_config.enforce_eager
-        ):
-            logger.warning("Enforce eager set, overriding optimization level to -O0")
-            self.optimization_level = OptimizationLevel.O0
+        if self.model_config is not None and self.model_config.enforce_eager:
+            logger.warning(
+                "Enforce eager set, disabling torch.compile and CUDAGraphs. "
+                "This is equivalent to setting -cc.mode=none -cc.cudagraph_mode=none"
+            )
+            self.compilation_config.mode = CompilationMode.NONE
+            self.compilation_config.cudagraph_mode = CUDAGraphMode.NONE
 
         if self.compilation_config.backend == "eager" or (
             self.compilation_config.mode is not None
@@ -1110,6 +1098,15 @@ class VllmConfig:
             self.scheduler_config.disable_hybrid_kv_cache_manager = False
 
         if self.cache_config.mamba_cache_mode == "align":
+            assert (
+                self.cache_config.block_size
+                <= self.scheduler_config.max_num_batched_tokens
+            ), (
+                "In Mamba cache align mode, block_size "
+                f"({self.cache_config.block_size}) must be <= "
+                "max_num_batched_tokens "
+                f"({self.scheduler_config.max_num_batched_tokens})."
+            )
             if self.scheduler_config.long_prefill_token_threshold > 0:
                 assert (
                     self.scheduler_config.long_prefill_token_threshold

vllm/distributed/kv_transfer/kv_connector/v1/__init__.py

@@ -6,7 +6,7 @@ from vllm.distributed.kv_transfer.kv_connector.v1.base import (
     SupportsHMA,
     supports_hma,
 )
-from vllm.distributed.kv_transfer.kv_connector.v1.decode_bench_connector import (  # noqa E:501
+from vllm.distributed.kv_transfer.kv_connector.v1.decode_bench_connector import (  # noqa: E501
     DecodeBenchConnector,
 )
 

vllm/distributed/kv_transfer/kv_connector/v1/example_connector.py

@@ -145,7 +145,6 @@ class ExampleConnector(KVConnectorBase_V1):
                     num_pages * page_size, -1
                 )
                 dst_kv_cache_layer[slot_mapping, ...] = src_kv_cache
-                dst_kv_cache_layer.reshape(dst_kv_cache_layer_shape)
             else:
                 num_pages = dst_kv_cache_layer_shape[1]
                 page_size = dst_kv_cache_layer_shape[2]
@@ -153,18 +152,11 @@ class ExampleConnector(KVConnectorBase_V1):
                     2, num_pages * page_size, -1
                 )
                 dst_kv_cache_layer[:, slot_mapping, ...] = src_kv_cache
-                dst_kv_cache_layer.reshape(dst_kv_cache_layer_shape)
 
         # Get the metadata
         metadata: KVConnectorMetadata = self._get_connector_metadata()
         assert isinstance(metadata, ExampleConnectorMetadata)
 
-        if metadata is None:
-            logger.warning(
-                "In connector.start_load_kv, but the connector metadata is None"
-            )
-            return
-
         attn_metadata = forward_context.attn_metadata
         if attn_metadata is None:
             logger.warning("In connector.start_load_kv, but the attn_metadata is None")

vllm/distributed/kv_transfer/kv_connector/v1/lmcache_mp_connector.py

@@ -36,7 +36,6 @@ except ImportError:
     )
 
 if TYPE_CHECKING:
-    from vllm.config import VllmConfig
     from vllm.distributed.kv_events import KVCacheEvent
     from vllm.distributed.kv_transfer.kv_connector.v1.metrics import (
         KVConnectorPromMetrics,

vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py

@@ -85,7 +85,7 @@ class OffloadingConnectorStats(KVConnectorStats):
         for transfer_type, ops_list in self.data.items():
             assert isinstance(ops_list, list)
             total_bytes = 0
-            total_time = 0
+            total_time = 0.0
             for op in ops_list:
                 assert isinstance(op, dict)
                 total_bytes += op["op_size"]

vllm/distributed/parallel_state.py

@@ -33,7 +33,7 @@ from contextlib import contextmanager, nullcontext
 from dataclasses import dataclass
 from datetime import timedelta
 from multiprocessing import shared_memory
-from typing import Any
+from typing import Any, Protocol
 from unittest.mock import patch
 
 import torch
@@ -64,6 +64,14 @@ class GraphCaptureContext:
 TensorMetadata = namedtuple("TensorMetadata", ["device", "dtype", "size"])
 
 
+class Handle(Protocol):
+    """Minimal async work handle used by P2P send/recv methods."""
+
+    def is_completed(self) -> bool: ...
+
+    def wait(self) -> None: ...
+
+
 def _split_tensor_dict(
     tensor_dict: dict[str, torch.Tensor | Any],
 ) -> tuple[list[tuple[str, Any]], list[torch.Tensor]]:
@@ -780,6 +788,20 @@ class GroupCoordinator:
                 async_handle.wait()
         return tensor_dict
 
+    def _should_use_all_gather(
+        self,
+        key: str,
+        numel: int,
+        all_gather_group: "GroupCoordinator | None",
+        all_gather_tensors: dict[str, bool] | None,
+    ) -> bool:
+        if all_gather_group is None:
+            return False
+        use_all_gather = numel % all_gather_group.world_size == 0
+        if all_gather_tensors is not None:
+            use_all_gather = all_gather_tensors.get(key, use_all_gather)
+        return use_all_gather
+
     def send_tensor_dict(
         self,
         tensor_dict: dict[str, torch.Tensor | Any],
@@ -808,6 +830,35 @@ class GroupCoordinator:
         # Bypass the function if we are using only 1 GPU.
         if not torch.distributed.is_initialized() or self.world_size == 1:
             return tensor_dict
+        handles = self.isend_tensor_dict(
+            tensor_dict,
+            dst=dst,
+            all_gather_group=all_gather_group,
+            all_gather_tensors=all_gather_tensors,
+        )
+        for handle in handles:
+            handle.wait()
+        return None
+
+    def isend_tensor_dict(
+        self,
+        tensor_dict: dict[str, torch.Tensor | Any],
+        dst: int | None = None,
+        all_gather_group: "GroupCoordinator | None" = None,
+        all_gather_tensors: dict[str, bool] | None = None,
+    ) -> list[Handle]:
+        if self.world_size <= 1:
+            return []
+
+        if self.use_cpu_custom_send_recv:
+            if self.device_communicator is None:
+                raise ValueError("No device communicator found")
+            # custom device communicator path is synchronous
+            self.device_communicator.send_tensor_dict(  # type: ignore
+                tensor_dict, dst
+            )
+            return []
+
         all_gather_size = 1 if all_gather_group is None else all_gather_group.world_size
         all_gather_rank = (
             0 if all_gather_group is None else all_gather_group.rank_in_group
@@ -820,53 +871,31 @@ class GroupCoordinator:
             dst = (self.rank_in_group + 1) % self.world_size
         assert dst < self.world_size, f"Invalid dst rank ({dst})"
 
-        if self.use_cpu_custom_send_recv:
-            if self.device_communicator is None:
-                raise ValueError("No device communicator found")
-            self.device_communicator.send_tensor_dict(  # type: ignore
-                tensor_dict, dst
-            )
-            return None
-
-        metadata_list: list[tuple[Any, Any]] = []
-        assert isinstance(tensor_dict, dict), (
-            f"Expecting a dictionary, got {type(tensor_dict)}"
-        )
         metadata_list, tensor_list = _split_tensor_dict(tensor_dict)
-        # `metadata_list` lives in CPU memory.
-        # `send_object_list` has serialization & deserialization,
-        # all happening on CPU. Therefore, we can use the CPU group.
         self.send_object(metadata_list, dst=dst)
 
         tensor_keys = [k for k, v in tensor_dict.items() if isinstance(v, torch.Tensor)]
         assert len(tensor_keys) == len(tensor_list)
 
+        handles: list[Handle] = []
         for key, tensor in zip(tensor_keys, tensor_list):
             if tensor.numel() == 0:
-                # Skip sending empty tensors.
                 continue
 
-            # send-allgather: send only a slice, then do allgather.
-            use_all_gather = (
-                all_gather_group is not None and tensor.numel() % all_gather_size == 0
-            )
-            use_all_gather = (
-                all_gather_tensors.get(key, use_all_gather)
-                if all_gather_tensors
-                else use_all_gather
-            )
-            if use_all_gather:
+            if self._should_use_all_gather(
+                key, tensor.numel(), all_gather_group, all_gather_tensors
+            ):
                 tensor = tensor.reshape(all_gather_size, -1)[all_gather_rank]
 
-            if tensor.is_cpu:
-                # use metadata_group for CPU tensors
-                torch.distributed.send(
-                    tensor, dst=self.ranks[dst], group=metadata_group
-                )
-            else:
-                # use group for GPU tensors
-                torch.distributed.send(tensor, dst=self.ranks[dst], group=group)
-        return None
+            comm_group = metadata_group if tensor.is_cpu else group
+            handle = torch.distributed.isend(
+                tensor, dst=self.ranks[dst], group=comm_group
+            )
+            if tensor.is_cuda:
+                tensor.record_stream(torch.cuda.current_stream(tensor.device))
+            handles.append(handle)
+
+        return handles
 
     def recv_tensor_dict(
         self,
@@ -895,6 +924,38 @@ class GroupCoordinator:
         # Bypass the function if we are using only 1 GPU.
         if not torch.distributed.is_initialized() or self.world_size == 1:
             return None
+        tensor_dict, handles, postprocess = self.irecv_tensor_dict(
+            src=src,
+            all_gather_group=all_gather_group,
+            all_gather_tensors=all_gather_tensors,
+        )
+        for handle in handles:
+            handle.wait()
+        for fn in postprocess:
+            fn()
+        return tensor_dict
+
+    def irecv_tensor_dict(
+        self,
+        src: int | None = None,
+        all_gather_group: "GroupCoordinator | None" = None,
+        all_gather_tensors: dict[str, bool] | None = None,
+    ) -> tuple[
+        dict[str, torch.Tensor | Any] | None,
+        list[Handle],
+        list[Callable[[], None]],
+    ]:
+        if not torch.distributed.is_initialized() or self.world_size == 1:
+            return None, [], []
+        if self.use_cpu_custom_send_recv:
+            if self.device_communicator is None:
+                raise ValueError("No device communicator found")
+            # custom device communicator path is synchronous
+            sync_tensor_dict = self.device_communicator.recv_tensor_dict(  # type: ignore
+                src
+            )
+            return sync_tensor_dict, [], []
+
         all_gather_size = 1 if all_gather_group is None else all_gather_group.world_size
         all_gather_rank = (
             0 if all_gather_group is None else all_gather_group.rank_in_group
@@ -907,57 +968,57 @@ class GroupCoordinator:
             src = (self.rank_in_group - 1) % self.world_size
         assert src < self.world_size, f"Invalid src rank ({src})"
 
-        if self.use_cpu_custom_send_recv:
-            if self.device_communicator is None:
-                raise ValueError("No device communicator found")
-            return self.device_communicator.recv_tensor_dict(  # type: ignore
-                src
-            )
-
         recv_metadata_list = self.recv_object(src=src)
         tensor_dict: dict[str, Any] = {}
+        handles: list[Handle] = []
+        postprocess: list[Callable[[], None]] = []
+
         for key, value in recv_metadata_list:
             if isinstance(value, TensorMetadata):
-                tensor = torch.empty(value.size, dtype=value.dtype, device=value.device)
-                if tensor.numel() == 0:
-                    # Skip broadcasting empty tensors.
-                    tensor_dict[key] = tensor
+                full_tensor = torch.empty(
+                    value.size, dtype=value.dtype, device=value.device
+                )
+                if full_tensor.numel() == 0:
+                    tensor_dict[key] = full_tensor
                     continue
 
-                # send-allgather: send only a slice, then do allgather.
-                use_all_gather = (
-                    all_gather_group is not None
-                    and tensor.numel() % all_gather_size == 0
-                )
-                use_all_gather = (
-                    all_gather_tensors.get(key, use_all_gather)
-                    if all_gather_tensors
-                    else use_all_gather
-                )
-
-                if use_all_gather:
-                    orig_shape = tensor.shape
-                    tensor = tensor.reshape(all_gather_size, -1)[all_gather_rank]
-
-                if tensor.is_cpu:
-                    # use metadata_group for CPU tensors
-                    torch.distributed.recv(
-                        tensor, src=self.ranks[src], group=metadata_group
+                if self._should_use_all_gather(
+                    key, full_tensor.numel(), all_gather_group, all_gather_tensors
+                ):
+                    orig_shape = full_tensor.shape
+                    slice_tensor = full_tensor.reshape(all_gather_size, -1)[
+                        all_gather_rank
+                    ]
+                    comm_group = metadata_group if slice_tensor.is_cpu else group
+                    handle = torch.distributed.irecv(
+                        slice_tensor, src=self.ranks[src], group=comm_group
                     )
+                    handles.append(handle)
+
+                    def _postprocess(
+                        key: str = key,
+                        slice_tensor: torch.Tensor = slice_tensor,
+                        orig_shape: tuple[int, ...] = tuple(orig_shape),
+                        all_gather_group=all_gather_group,
+                    ) -> None:
+                        assert all_gather_group is not None
+                        tensor_dict[key] = all_gather_group.all_gather(
+                            slice_tensor, dim=0
+                        ).reshape(orig_shape)
+
+                    postprocess.append(_postprocess)
+                    tensor_dict[key] = slice_tensor
                 else:
-                    # use group for GPU tensors
-                    torch.distributed.recv(tensor, src=self.ranks[src], group=group)
-                if use_all_gather:
-                    # do the allgather
-                    tensor = all_gather_group.all_gather(  # type: ignore
-                        tensor, dim=0
+                    comm_group = metadata_group if full_tensor.is_cpu else group
+                    handle = torch.distributed.irecv(
+                        full_tensor, src=self.ranks[src], group=comm_group
                     )
-                    tensor = tensor.reshape(orig_shape)
-
-                tensor_dict[key] = tensor
+                    handles.append(handle)
+                    tensor_dict[key] = full_tensor
             else:
                 tensor_dict[key] = value
-        return tensor_dict
+
+        return tensor_dict, handles, postprocess
 
     def barrier(self):
         """Barrier synchronization among the group.

vllm/distributed/utils.py

@@ -524,3 +524,43 @@ def stateless_destroy_torch_distributed_process_group(pg: ProcessGroup) -> None:
     """
     pg.shutdown()
     _unregister_process_group(pg.group_name)
+
+
+def get_worker_rank_suffix(global_rank: int | None = None) -> str:
+    """Generate a descriptive rank suffix for worker identification.
+
+    Returns a string like 'dp0_pp0_tp0_dcp0_ep0_rank0' including all
+    parallel dimensions: DP, PP, TP, DCP, EP.
+
+    Args:
+        global_rank: Optional global rank to append. If not provided,
+                     only parallel dimension ranks are included.
+
+    Returns:
+        A string suffix identifying the worker's position in the
+        distributed topology.
+    """
+    from vllm.distributed.parallel_state import (
+        get_dcp_group,
+        get_dp_group,
+        get_ep_group,
+        get_pp_group,
+        get_tp_group,
+    )
+
+    try:
+        dp_rank = get_dp_group().rank_in_group
+        pp_rank = get_pp_group().rank_in_group
+        tp_rank = get_tp_group().rank_in_group
+        dcp_rank = get_dcp_group().rank_in_group
+        ep_rank = get_ep_group().rank_in_group
+
+        suffix = f"dp{dp_rank}_pp{pp_rank}_tp{tp_rank}_dcp{dcp_rank}_ep{ep_rank}"
+        if global_rank is not None:
+            suffix = f"{suffix}_rank{global_rank}"
+        return suffix
+    except Exception:
+        # Fallback if parallel state not initialized
+        if global_rank is not None:
+            return f"rank{global_rank}"
+        return ""

vllm/engine/arg_utils.py

@@ -434,6 +434,7 @@ class EngineArgs:
     disable_cascade_attn: bool = ModelConfig.disable_cascade_attn
     swap_space: float = CacheConfig.swap_space
     cpu_offload_gb: float = CacheConfig.cpu_offload_gb
+    cpu_offload_params: set[str] = get_field(CacheConfig, "cpu_offload_params")
     gpu_memory_utilization: float = CacheConfig.gpu_memory_utilization
     kv_cache_memory_bytes: int | None = CacheConfig.kv_cache_memory_bytes
     max_num_batched_tokens: int | None = None
@@ -942,6 +943,9 @@ class EngineArgs:
             "--prefix-caching-hash-algo", **cache_kwargs["prefix_caching_hash_algo"]
         )
         cache_group.add_argument("--cpu-offload-gb", **cache_kwargs["cpu_offload_gb"])
+        cache_group.add_argument(
+            "--cpu-offload-params", **cache_kwargs["cpu_offload_params"]
+        )
         cache_group.add_argument(
             "--calculate-kv-scales", **cache_kwargs["calculate_kv_scales"]
         )
@@ -1453,6 +1457,7 @@ class EngineArgs:
             enable_prefix_caching=self.enable_prefix_caching,
             prefix_caching_hash_algo=self.prefix_caching_hash_algo,
             cpu_offload_gb=self.cpu_offload_gb,
+            cpu_offload_params=self.cpu_offload_params,
             calculate_kv_scales=self.calculate_kv_scales,
             kv_sharing_fast_prefill=self.kv_sharing_fast_prefill,
             mamba_cache_dtype=self.mamba_cache_dtype,

vllm/engine/protocol.py

@@ -3,6 +3,7 @@
 
 from abc import ABC, abstractmethod
 from collections.abc import AsyncGenerator, Iterable, Mapping
+from dataclasses import dataclass
 from typing import TYPE_CHECKING, Any
 
 from vllm.config import ModelConfig, VllmConfig
@@ -10,7 +11,7 @@ from vllm.distributed.weight_transfer.base import (
     WeightTransferInitRequest,
     WeightTransferUpdateRequest,
 )
-from vllm.inputs.data import PromptType, StreamingInput
+from vllm.inputs.data import PromptType
 from vllm.lora.request import LoRARequest
 from vllm.outputs import PoolingRequestOutput, RequestOutput
 from vllm.plugins.io_processors import IOProcessor
@@ -26,6 +27,18 @@ if TYPE_CHECKING:
     from vllm.v1.engine import PauseMode
 
 
+@dataclass
+class StreamingInput:
+    """Input data for a streaming generation request.
+
+    This is used with generate() to support multi-turn streaming sessions
+    where inputs are provided via an async generator.
+    """
+
+    prompt: PromptType
+    sampling_params: SamplingParams | None = None
+
+
 class EngineClient(ABC):
     """Protocol class for Clients to Engine"""
 

vllm/entrypoints/llm.py

@@ -72,7 +72,7 @@ from vllm.outputs import (
 )
 from vllm.platforms import current_platform
 from vllm.pooling_params import PoolingParams
-from vllm.renderers import ChatParams, TokenizeParams, merge_kwargs
+from vllm.renderers import ChatParams, merge_kwargs
 from vllm.renderers.inputs import DictPrompt, TokPrompt
 from vllm.renderers.inputs.preprocess import (
     conversation_to_seq,
@@ -384,7 +384,7 @@ class LLM:
         return parallel_config.world_size
 
     def reset_mm_cache(self) -> None:
-        self.input_processor.clear_mm_cache()
+        self.renderer.clear_mm_cache()
         self.llm_engine.reset_mm_cache()
 
     def get_default_sampling_params(self) -> SamplingParams:
@@ -458,6 +458,93 @@ class LLM:
 
         return self.engine_class.validate_outputs(outputs, RequestOutput)
 
+    def enqueue(
+        self,
+        prompts: PromptType | Sequence[PromptType],
+        sampling_params: SamplingParams | Sequence[SamplingParams] | None = None,
+        lora_request: list[LoRARequest] | LoRARequest | None = None,
+        priority: list[int] | None = None,
+        use_tqdm: bool | Callable[..., tqdm] = True,
+        tokenization_kwargs: dict[str, Any] | None = None,
+    ) -> list[str]:
+        """Enqueue prompts for generation without waiting for completion.
+
+        This method adds requests to the engine queue but does not start
+        processing them. Use wait_for_completion() to process the queued
+        requests and get results.
+
+        Args:
+            prompts: The prompts to the LLM. See generate() for details.
+            sampling_params: The sampling parameters for text generation.
+            lora_request: LoRA request to use for generation, if any.
+            priority: The priority of the requests, if any.
+            use_tqdm: If True, shows a tqdm progress bar while adding requests.
+            tokenization_kwargs: Overrides for `tokenizer.encode`.
+
+        Returns:
+            A list of request IDs for the enqueued requests.
+        """
+        model_config = self.model_config
+        runner_type = model_config.runner_type
+        if runner_type != "generate":
+            raise ValueError("LLM.enqueue() is only supported for generative models.")
+
+        if sampling_params is None:
+            sampling_params = self.get_default_sampling_params()
+
+        # Use the same preprocessing as _run_completion
+        seq_prompts = prompt_to_seq(prompts)
+        seq_params = self._params_to_seq(sampling_params, len(seq_prompts))
+
+        if any(param.truncate_prompt_tokens is not None for param in seq_params):
+            engine_prompts: Sequence[DictPrompt | TokPrompt] = [
+                engine_prompt
+                for prompt, param in zip(seq_prompts, seq_params)
+                for engine_prompt in self._preprocess_cmpl(
+                    [prompt],
+                    tokenization_kwargs=merge_kwargs(
+                        tokenization_kwargs,
+                        dict(truncate_prompt_tokens=param.truncate_prompt_tokens),
+                    ),
+                )
+            ]
+        else:
+            engine_prompts = self._preprocess_cmpl(
+                seq_prompts,
+                tokenization_kwargs=tokenization_kwargs,
+            )
+
+        request_ids = self._validate_and_add_requests(
+            prompts=engine_prompts,
+            params=seq_params,
+            use_tqdm=use_tqdm,
+            lora_request=self._get_modality_specific_lora_reqs(
+                engine_prompts, lora_request
+            ),
+            tokenization_kwargs=tokenization_kwargs,
+            priority=priority,
+        )
+
+        return request_ids
+
+    def wait_for_completion(
+        self,
+        use_tqdm: bool | Callable[..., tqdm] = True,
+    ) -> list[RequestOutput]:
+        """Wait for all enqueued requests to complete and return results.
+
+        This method processes all requests currently in the engine queue
+        and returns their outputs. Use after enqueue() to get results.
+
+        Args:
+            use_tqdm: If True, shows a tqdm progress bar.
+
+        Returns:
+            A list of RequestOutput objects for all completed requests.
+        """
+        outputs = self._run_engine(use_tqdm=use_tqdm)
+        return self.engine_class.validate_outputs(outputs, RequestOutput)
+
     def _get_modality_specific_lora_reqs(
         self,
         prompts: Sequence[DictPrompt | TokPrompt],
@@ -789,20 +876,7 @@ class LLM:
 
         return outputs
 
-    def _get_cmpl_tok_params(self, tokenization_kwargs: dict[str, Any] | None):
-        model_config = self.model_config
-        encoder_config = model_config.encoder_config or {}
-
-        return TokenizeParams(
-            max_total_tokens=model_config.max_model_len,
-            do_lower_case=encoder_config.get("do_lower_case", False),
-            # For Whisper, special tokens should be provided by the user based
-            # on the task and language of their request. Also needed to avoid
-            # appending an EOS token to the prompt which disrupts generation.
-            add_special_tokens=not model_config.is_encoder_decoder,
-        ).with_kwargs(tokenization_kwargs)
-
-    def _preprocess_completion(
+    def _preprocess_cmpl(
         self,
         prompts: Sequence[PromptType],
         tokenization_kwargs: dict[str, Any] | None = None,
@@ -814,7 +888,7 @@ class LLM:
         Refer to [LLM.generate][] for a complete description of the arguments.
 
         Returns:
-            A list of `TokensPrompts` objects containing the tokenized prompt
+            A list of `TokPrompt` objects containing the tokenized prompt
             after chat template interpolation, and the raw multi-modal inputs.
         """
         renderer = self.renderer
@@ -823,20 +897,12 @@ class LLM:
         parsed_prompts = [
             parse_model_prompt(model_config, prompt) for prompt in prompts
         ]
-        tok_params = self._get_cmpl_tok_params(tokenization_kwargs)
+        tok_params = renderer.default_cmpl_tok_params.with_kwargs(
+            **(tokenization_kwargs or {})
+        )
 
         return renderer.render_cmpl(parsed_prompts, tok_params)
 
-    def _get_chat_tok_params(self, tokenization_kwargs: dict[str, Any] | None):
-        model_config = self.model_config
-        encoder_config = model_config.encoder_config or {}
-
-        return TokenizeParams(
-            max_total_tokens=model_config.max_model_len,
-            do_lower_case=encoder_config.get("do_lower_case", False),
-            add_special_tokens=False,
-        ).with_kwargs(tokenization_kwargs)
-
     def _preprocess_chat(
         self,
         conversations: Sequence[list[ChatCompletionMessageParam]],
@@ -856,7 +922,7 @@ class LLM:
         Refer to [LLM.chat][] for a complete description of the arguments.
 
         Returns:
-            A list of `TokensPrompts` objects containing the tokenized prompt
+            A list of `TokPrompt` objects containing the tokenized prompt
             after chat template interpolation, and the raw multi-modal inputs.
         """
         renderer = self.renderer
@@ -874,7 +940,9 @@ class LLM:
                 ),
             ),
         )
-        tok_params = self._get_chat_tok_params(tokenization_kwargs)
+        tok_params = renderer.default_chat_tok_params.with_kwargs(
+            **(tokenization_kwargs or {})
+        )
 
         _, engine_prompts = renderer.render_chat(
             conversations,
@@ -1566,7 +1634,10 @@ class LLM:
             architecture=architecture,
         )
 
-        tok_params = self._get_cmpl_tok_params(tokenization_kwargs)
+        renderer = self.renderer
+        tok_params = renderer.default_cmpl_tok_params.with_kwargs(
+            **(tokenization_kwargs or {})
+        )
         encode_kwargs = tok_params.get_encode_kwargs()
 
         if model_config.is_cross_encoder:
@@ -1598,8 +1669,15 @@ class LLM:
                 tokenization_kwargs=encode_kwargs,
             )
 
-    def start_profile(self) -> None:
-        self.llm_engine.start_profile()
+    def start_profile(self, profile_prefix: str | None = None) -> None:
+        """Start profiling with optional custom trace prefix.
+
+        Args:
+            profile_prefix: Optional prefix for the trace file names. If provided,
+                           trace files will be named as "<prefix>_dp<X>_pp<Y>_tp<Z>".
+                           If not provided, default naming will be used.
+        """
+        self.llm_engine.start_profile(profile_prefix)
 
     def stop_profile(self) -> None:
         self.llm_engine.stop_profile()
@@ -1618,19 +1696,22 @@ class LLM:
         during the sleep period, before `wake_up` is called.
 
         Args:
-            level: The sleep level. Level 1 sleep will offload the model
-                weights and discard the kv cache. The content of kv cache
-                is forgotten. Level 1 sleep is good for sleeping and waking
-                up the engine to run the same model again. The model weights
-                are backed up in CPU memory. Please make sure there's enough
-                CPU memory to store the model weights. Level 2 sleep will
-                discard both the model weights and the kv cache. The content
-                of both the model weights and kv cache is forgotten. Level 2
-                sleep is good for sleeping and waking up the engine to run a
-                different model or update the model, where previous model
-                weights are not needed. It reduces CPU memory pressure.
+            level: The sleep level.
+                - Level 0: Pause scheduling but continue accepting requests.
+                           Requests are queued but not processed.
+                - Level 1: Offload model weights to CPU, discard KV cache.
+                           The content of kv cache is forgotten. Good for
+                           sleeping and waking up the engine to run the same
+                           model again. Please make sure there's enough CPU
+                           memory to store the model weights.
+                - Level 2: Discard all GPU memory (weights + KV cache).
+                           Good for sleeping and waking up the engine to run
+                           a different model or update the model, where
+                           previous model weights are not needed. It reduces
+                           CPU memory pressure.
         """
-        self.reset_prefix_cache()
+        if level > 0:
+            self.reset_prefix_cache()
         self.llm_engine.sleep(level=level)
 
     def wake_up(self, tags: list[str] | None = None):
@@ -1641,9 +1722,10 @@ class LLM:
         Args:
             tags: An optional list of tags to reallocate the engine memory
                 for specific memory allocations. Values must be in
-                `("weights", "kv_cache")`. If None, all memory is reallocated.
-                wake_up should be called with all tags (or None) before the
-                engine is used again.
+                `("weights", "kv_cache", "scheduling")`. If None, all memory
+                is reallocated. wake_up should be called with all tags
+                (or None) before the engine is used again.
+                Use tags=["scheduling"] to resume from level 0 sleep.
         """
         self.llm_engine.wake_up(tags)
 
@@ -1725,11 +1807,11 @@ class LLM:
         if any(param.truncate_prompt_tokens is not None for param in seq_params):
             # TODO: Remove this after deprecating `param.truncate_prompt_tokens`
             # Then, move the code from the `else` block to the top and let
-            # `self._preprocess_completion` handle prompt normalization
+            # `self._preprocess_cmpl` handle prompt normalization
             engine_prompts: Sequence[DictPrompt | TokPrompt] = [
                 engine_prompt
                 for prompt, param in zip(seq_prompts, seq_params)
-                for engine_prompt in self._preprocess_completion(
+                for engine_prompt in self._preprocess_cmpl(
                     [prompt],
                     tokenization_kwargs=merge_kwargs(
                         tokenization_kwargs,
@@ -1738,7 +1820,7 @@ class LLM:
                 )
             ]
         else:
-            engine_prompts = self._preprocess_completion(
+            engine_prompts = self._preprocess_cmpl(
                 seq_prompts,
                 tokenization_kwargs=tokenization_kwargs,
             )
@@ -1810,7 +1892,7 @@ class LLM:
         lora_request: Sequence[LoRARequest | None] | LoRARequest | None,
         tokenization_kwargs: dict[str, Any] | None = None,
         priority: list[int] | None = None,
-    ) -> None:
+    ) -> list[str]:
         num_requests = len(prompts)
         seq_params = self._params_to_seq(params, num_requests)
         seq_lora_requests = self._lora_request_to_seq(lora_request, num_requests)
@@ -1844,6 +1926,8 @@ class LLM:
                 self.llm_engine.abort_request(added_request_ids, internal=True)
             raise e
 
+        return added_request_ids
+
     def _add_request(
         self,
         prompt: PromptType | DictPrompt | TokPrompt,
@@ -1870,7 +1954,10 @@ class LLM:
                 dict(truncate_prompt_tokens=params.truncate_prompt_tokens),
             )
 
-        tok_params = self._get_cmpl_tok_params(tokenization_kwargs)
+        renderer = self.renderer
+        tok_params = renderer.default_cmpl_tok_params.with_kwargs(
+            **(tokenization_kwargs or {})
+        )
 
         tokenization_kwargs = tok_params.get_encode_kwargs()
         engine_request = self.input_processor.process_inputs(
@@ -1895,7 +1982,9 @@ class LLM:
         return engine_request.request_id
 
     def _run_engine(
-        self, *, use_tqdm: bool | Callable[..., tqdm] = True
+        self,
+        *,
+        use_tqdm: bool | Callable[..., tqdm] = True,
     ) -> list[RequestOutput | PoolingRequestOutput]:
         # Initialize tqdm.
         if use_tqdm:

vllm/entrypoints/openai/chat_completion/api_router.py

@@ -57,7 +57,7 @@ async def create_chat_completion(request: ChatCompletionRequest, raw_request: Re
     try:
         generator = await handler.create_chat_completion(request, raw_request)
     except Exception as e:
-        return handler.create_error_response(e)
+        generator = handler.create_error_response(e)
 
     if isinstance(generator, ErrorResponse):
         return JSONResponse(
@@ -96,7 +96,7 @@ async def render_chat_completion(request: ChatCompletionRequest, raw_request: Re
     try:
         result = await handler.render_chat_request(request)
     except Exception as e:
-        return handler.create_error_response(e)
+        result = handler.create_error_response(e)
 
     if isinstance(result, ErrorResponse):
         return JSONResponse(content=result.model_dump(), status_code=result.error.code)