[Metrics] Add Prometheus counter for CUDA graph iteration mode

Add `vllm:cudagraph_iterations` counter with `runtime_mode` label
(NONE/PIECEWISE/FULL) so CUDA graph usage percentage can be computed
in PromQL. Also always populate CUDAGraphStat regardless of the
`--cudagraph-metrics` flag, since the stat is trivially cheap and that
flag should only gate the verbose text log table.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>

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

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

.buildkite/test-amd.yaml

@@ -1573,7 +1573,7 @@ steps:
   - tests/compile/fullgraph/test_basic_correctness.py
   - examples/offline_inference/rlhf.py
   - examples/offline_inference/rlhf_colocate.py
-  - examples/offline_inference/new_weight_syncing/
+  - examples/rl/
   - tests/examples/offline_inference/data_parallel.py
   - tests/v1/distributed
   - tests/v1/engine/test_engine_core_client.py
@@ -1615,7 +1615,7 @@ steps:
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 RAY_DEDUP_LOGS=0 python3 rlhf_colocate.py
   - popd
   # NEW rlhf examples
-  - pushd ../examples/offline_inference/new_weight_syncing
+  - pushd ../examples/rl
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_nccl.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_ipc.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_async_new_apis.py
@@ -2660,7 +2660,7 @@ steps:
   - tests/v1/entrypoints/openai/test_multi_api_servers.py
   - tests/v1/shutdown
   - tests/v1/worker/test_worker_memory_snapshot.py
-  - examples/offline_inference/new_weight_syncing/
+  - examples/rl/
   commands:
   # Work around HIP bug tracked here: https://github.com/ROCm/hip/issues/3876
   # TODO: Remove when the bug is fixed in a future ROCm release
@@ -3325,7 +3325,7 @@ steps:
   - tests/compile/fullgraph/test_basic_correctness.py
   - examples/offline_inference/rlhf.py
   - examples/offline_inference/rlhf_colocate.py
-  - examples/offline_inference/new_weight_syncing/
+  - examples/rl/
   - tests/examples/offline_inference/data_parallel.py
   - tests/v1/distributed
   - tests/v1/engine/test_engine_core_client.py
@@ -3367,7 +3367,7 @@ steps:
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 RAY_DEDUP_LOGS=0 python3 rlhf_colocate.py
   - popd
   # NEW rlhf examples
-  - pushd ../examples/offline_inference/new_weight_syncing
+  - pushd ../examples/rl
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_nccl.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_ipc.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_async_new_apis.py

.buildkite/test_areas/distributed.yaml

@@ -82,7 +82,7 @@ steps:
 
 - label: Distributed Torchrun + Examples (4 GPUs)
   timeout_in_minutes: 30
-  working_dir: "/vllm-workspace/tests"
+  working_dir: "/vllm-workspace"
   num_devices: 4
   source_file_dependencies:
   - vllm/distributed/
@@ -90,33 +90,28 @@ steps:
   - tests/distributed/test_torchrun_example_moe.py
   - examples/offline_inference/rlhf.py
   - examples/offline_inference/rlhf_colocate.py
-  - examples/offline_inference/new_weight_syncing/
+  - examples/rl/
   - tests/examples/offline_inference/data_parallel.py
   commands:
   # https://github.com/NVIDIA/nccl/issues/1838
   - export NCCL_CUMEM_HOST_ENABLE=0
   # test with torchrun tp=2 and external_dp=2
-  - torchrun --nproc-per-node=4 distributed/test_torchrun_example.py
+  - torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example.py
   # test with torchrun tp=2 and pp=2
-  - PP_SIZE=2 torchrun --nproc-per-node=4 distributed/test_torchrun_example.py
+  - PP_SIZE=2 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example.py
   # test with torchrun tp=4 and dp=1
-  - TP_SIZE=4 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
+  - TP_SIZE=4 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
   # test with torchrun tp=2, pp=2 and dp=1
-  - PP_SIZE=2 TP_SIZE=2 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
+  - PP_SIZE=2 TP_SIZE=2 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
   # test with torchrun tp=1 and dp=4 with ep
-  - DP_SIZE=4 ENABLE_EP=1 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
+  - DP_SIZE=4 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
   # test with torchrun tp=2 and dp=2 with ep
-  - TP_SIZE=2 DP_SIZE=2 ENABLE_EP=1 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
+  - TP_SIZE=2 DP_SIZE=2 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
   # test with internal dp
-  - python3 ../examples/offline_inference/data_parallel.py --enforce-eager
-  # OLD rlhf examples
-  - cd ../examples/offline_inference
-  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf.py
-  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 RAY_DEDUP_LOGS=0 python3 rlhf_colocate.py
-  # NEW rlhf examples
-  - cd new_weight_syncing
-  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_nccl.py
-  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 rlhf_ipc.py
+  - python3 examples/offline_inference/data_parallel.py --enforce-eager
+  # rlhf examples
+  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_nccl.py
+  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_ipc.py
 
 - label: Distributed DP Tests (4 GPUs)
   timeout_in_minutes: 30

.buildkite/test_areas/expert_parallelism.yaml

@@ -24,8 +24,7 @@ steps:
 
 - label: Elastic EP Scaling Test
   timeout_in_minutes: 20
-  device: b200
-  optional: true
+  device: h100
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:

CMakeLists.txt

@@ -999,6 +999,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   list(APPEND VLLM_MOE_EXT_SRC
     "csrc/moe/moe_wna16.cu"
     "csrc/moe/grouped_topk_kernels.cu"
+    "csrc/moe/gpt_oss_router_gemm.cu"
     "csrc/moe/router_gemm.cu")
 endif()
 

benchmarks/kernels/benchmark_moe.py

@@ -750,17 +750,20 @@ def get_weight_block_size_safety(config, default_value=None):
 
 
 def get_model_params(config):
-    if config.architectures[0] == "DbrxForCausalLM":
+    architectures = getattr(config, "architectures", None) or [type(config).__name__]
+    architecture = architectures[0]
+
+    if architecture == "DbrxForCausalLM":
         E = config.ffn_config.moe_num_experts
         topk = config.ffn_config.moe_top_k
         intermediate_size = config.ffn_config.ffn_hidden_size
         hidden_size = config.hidden_size
-    elif config.architectures[0] == "JambaForCausalLM":
+    elif architecture == "JambaForCausalLM":
         E = config.num_experts
         topk = config.num_experts_per_tok
         intermediate_size = config.intermediate_size
         hidden_size = config.hidden_size
-    elif config.architectures[0] in (
+    elif architecture in (
         "DeepseekV2ForCausalLM",
         "DeepseekV3ForCausalLM",
         "DeepseekV32ForCausalLM",
@@ -774,7 +777,7 @@ def get_model_params(config):
         topk = config.num_experts_per_tok
         intermediate_size = config.moe_intermediate_size
         hidden_size = config.hidden_size
-    elif config.architectures[0] in (
+    elif architecture in (
         "Qwen2MoeForCausalLM",
         "Qwen3MoeForCausalLM",
         "Qwen3NextForCausalLM",
@@ -783,23 +786,27 @@ def get_model_params(config):
         topk = config.num_experts_per_tok
         intermediate_size = config.moe_intermediate_size
         hidden_size = config.hidden_size
-    elif config.architectures[0] == "Qwen3VLMoeForConditionalGeneration":
+    elif architecture in (
+        "Qwen3VLMoeForConditionalGeneration",
+        "Qwen3_5MoeForConditionalGeneration",
+        "Qwen3_5MoeTextConfig",
+    ):
         text_config = config.get_text_config()
         E = text_config.num_experts
         topk = text_config.num_experts_per_tok
         intermediate_size = text_config.moe_intermediate_size
         hidden_size = text_config.hidden_size
-    elif config.architectures[0] == "HunYuanMoEV1ForCausalLM":
+    elif architecture == "HunYuanMoEV1ForCausalLM":
         E = config.num_experts
         topk = config.moe_topk[0]
         intermediate_size = config.moe_intermediate_size[0]
         hidden_size = config.hidden_size
-    elif config.architectures[0] == "Qwen3OmniMoeForConditionalGeneration":
+    elif architecture == "Qwen3OmniMoeForConditionalGeneration":
         E = config.thinker_config.text_config.num_experts
         topk = config.thinker_config.text_config.num_experts_per_tok
         intermediate_size = config.thinker_config.text_config.moe_intermediate_size
         hidden_size = config.thinker_config.text_config.hidden_size
-    elif config.architectures[0] == "PixtralForConditionalGeneration":
+    elif architecture == "PixtralForConditionalGeneration":
         # Pixtral can contain different LLM architectures,
         # recurse to get their parameters
         return get_model_params(config.get_text_config())
@@ -814,6 +821,23 @@ def get_model_params(config):
     return E, topk, intermediate_size, hidden_size
 
 
+def resolve_dtype(config) -> torch.dtype:
+    if current_platform.is_rocm():
+        return torch.float16
+
+    dtype = getattr(config, "dtype", None)
+    if dtype is not None:
+        return dtype
+
+    if hasattr(config, "get_text_config"):
+        text_config = config.get_text_config()
+        dtype = getattr(text_config, "dtype", None)
+        if dtype is not None:
+            return dtype
+
+    return torch.bfloat16
+
+
 def get_quantization_group_size(config) -> int | None:
     """Extract the quantization group size from the HF model config.
 
@@ -861,7 +885,7 @@ def main(args: argparse.Namespace):
     else:
         ensure_divisibility(intermediate_size, args.tp_size, "intermediate_size")
         shard_intermediate_size = 2 * intermediate_size // args.tp_size
-    dtype = torch.float16 if current_platform.is_rocm() else config.dtype
+    dtype = resolve_dtype(config)
     use_fp8_w8a8 = args.dtype == "fp8_w8a8"
     use_int8_w8a16 = args.dtype == "int8_w8a16"
     use_int4_w4a16 = args.dtype == "int4_w4a16"

benchmarks/kernels/benchmark_router_gemm.py

@@ -0,0 +1,134 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import torch
+import torch.nn.functional as F
+
+from vllm import _custom_ops as ops
+from vllm.platforms import current_platform
+from vllm.transformers_utils.config import get_config
+from vllm.triton_utils import triton
+from vllm.utils.argparse_utils import FlexibleArgumentParser
+
+# Dimensions supported by the DSV3 specialized kernel
+DSV3_SUPPORTED_NUM_EXPERTS = [256, 384]
+DSV3_SUPPORTED_HIDDEN_SIZES = [7168]
+
+# Dimensions supported by the gpt-oss specialized kernel
+GPT_OSS_SUPPORTED_NUM_EXPERTS = [32, 128]
+GPT_OSS_SUPPORTED_HIDDEN_SIZES = [2880]
+
+
+def get_batch_size_range(max_batch_size):
+    return [2**x for x in range(14) if 2**x <= max_batch_size]
+
+
+def get_model_params(config):
+    if config.architectures[0] in (
+        "DeepseekV2ForCausalLM",
+        "DeepseekV3ForCausalLM",
+        "DeepseekV32ForCausalLM",
+    ):
+        num_experts = config.n_routed_experts
+        hidden_size = config.hidden_size
+    elif config.architectures[0] in ("GptOssForCausalLM",):
+        num_experts = config.num_local_experts
+        hidden_size = config.hidden_size
+    else:
+        raise ValueError(f"Unsupported architecture: {config.architectures}")
+    return num_experts, hidden_size
+
+
+def get_benchmark(model, max_batch_size, trust_remote_code):
+    @triton.testing.perf_report(
+        triton.testing.Benchmark(
+            x_names=["batch_size"],
+            x_vals=get_batch_size_range(max_batch_size),
+            x_log=False,
+            line_arg="provider",
+            line_vals=[
+                "torch",
+                "vllm",
+            ],
+            line_names=["PyTorch", "vLLM"],
+            styles=([("blue", "-"), ("red", "-")]),
+            ylabel="TFLOPs",
+            plot_name=f"{model} router gemm throughput",
+            args={},
+        )
+    )
+    def benchmark(batch_size, provider):
+        config = get_config(model=model, trust_remote_code=trust_remote_code)
+        num_experts, hidden_size = get_model_params(config)
+
+        mat_a = torch.randn(
+            (batch_size, hidden_size), dtype=torch.bfloat16, device="cuda"
+        ).contiguous()
+        mat_b = torch.randn(
+            (num_experts, hidden_size), dtype=torch.bfloat16, device="cuda"
+        ).contiguous()
+        bias = torch.randn(
+            num_experts, dtype=torch.bfloat16, device="cuda"
+        ).contiguous()
+
+        is_hopper_or_blackwell = current_platform.is_device_capability(
+            90
+        ) or current_platform.is_device_capability_family(100)
+        allow_dsv3_router_gemm = (
+            is_hopper_or_blackwell
+            and num_experts in DSV3_SUPPORTED_NUM_EXPERTS
+            and hidden_size in DSV3_SUPPORTED_HIDDEN_SIZES
+        )
+        allow_gpt_oss_router_gemm = (
+            is_hopper_or_blackwell
+            and num_experts in GPT_OSS_SUPPORTED_NUM_EXPERTS
+            and hidden_size in GPT_OSS_SUPPORTED_HIDDEN_SIZES
+        )
+
+        has_bias = False
+        if allow_gpt_oss_router_gemm:
+            has_bias = True
+
+        quantiles = [0.5, 0.2, 0.8]
+
+        if provider == "torch":
+
+            def runner():
+                if has_bias:
+                    F.linear(mat_a, mat_b, bias)
+                else:
+                    F.linear(mat_a, mat_b)
+        elif provider == "vllm":
+
+            def runner():
+                if allow_dsv3_router_gemm:
+                    ops.dsv3_router_gemm(mat_a, mat_b, torch.bfloat16)
+                elif allow_gpt_oss_router_gemm:
+                    ops.gpt_oss_router_gemm(mat_a, mat_b, bias)
+                else:
+                    raise ValueError("Unsupported router gemm")
+
+        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+            runner, quantiles=quantiles
+        )
+
+        def tflops(t_ms):
+            flops = 2 * batch_size * hidden_size * num_experts
+            return flops / (t_ms * 1e-3) / 1e12
+
+        return tflops(ms), tflops(max_ms), tflops(min_ms)
+
+    return benchmark
+
+
+if __name__ == "__main__":
+    parser = FlexibleArgumentParser()
+    parser.add_argument("--model", type=str, default="openai/gpt-oss-20b")
+    parser.add_argument("--max-batch-size", default=16, type=int)
+    parser.add_argument("--trust-remote-code", action="store_true")
+    args = parser.parse_args()
+
+    # Get the benchmark function
+    benchmark = get_benchmark(args.model, args.max_batch_size, args.trust_remote_code)
+    # Run performance benchmark
+    benchmark.run(print_data=True)

csrc/moe/gpt_oss_router_gemm.cu

@@ -0,0 +1,144 @@
+/*
+ * Adapted from
+ * https://github.com/NVIDIA/TensorRT-LLM/blob/v1.3.0rc7/cpp/tensorrt_llm/kernels/tinygemm2/tinygemm2_cuda.cu
+ * Copyright (c) 2025, The vLLM team.
+ * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * All rights reserved. SPDX-License-Identifier: Apache-2.0
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAStream.h>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <torch/all.h>
+#include "gpt_oss_router_gemm.cuh"
+
+void launch_gpt_oss_router_gemm(__nv_bfloat16* gA, __nv_bfloat16* gB,
+                                __nv_bfloat16* gC, __nv_bfloat16* bias,
+                                int batch_size, int output_features,
+                                int input_features, cudaStream_t stream) {
+  static int const WARP_TILE_M = 16;
+  static int const TILE_M = WARP_TILE_M;
+  static int const TILE_N = 8;
+  static int const TILE_K = 64;
+  static int const STAGES = 16;
+  static int const STAGE_UNROLL = 4;
+  static bool const PROFILE = false;
+
+  CUtensorMap weight_map{};
+  CUtensorMap activation_map{};
+
+  constexpr uint32_t rank = 2;
+  uint64_t size[rank] = {(uint64_t)input_features, (uint64_t)output_features};
+  uint64_t stride[rank - 1] = {input_features * sizeof(__nv_bfloat16)};
+  uint32_t box_size[rank] = {TILE_K, TILE_M};
+  uint32_t elem_stride[rank] = {1, 1};
+
+  CUresult res = cuTensorMapEncodeTiled(
+      &weight_map, CUtensorMapDataType::CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, rank,
+      gB, size, stride, box_size, elem_stride,
+      CUtensorMapInterleave::CU_TENSOR_MAP_INTERLEAVE_NONE,
+      CUtensorMapSwizzle::CU_TENSOR_MAP_SWIZZLE_128B,
+      CUtensorMapL2promotion::CU_TENSOR_MAP_L2_PROMOTION_NONE,
+      CUtensorMapFloatOOBfill::CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);
+  TORCH_CHECK(res == CUDA_SUCCESS,
+              "cuTensorMapEncodeTiled failed for weight_map, error code=",
+              static_cast<int>(res));
+
+  size[1] = batch_size;
+  box_size[1] = TILE_N;
+
+  res = cuTensorMapEncodeTiled(
+      &activation_map, CUtensorMapDataType::CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,
+      rank, gA, size, stride, box_size, elem_stride,
+      CUtensorMapInterleave::CU_TENSOR_MAP_INTERLEAVE_NONE,
+      CUtensorMapSwizzle::CU_TENSOR_MAP_SWIZZLE_128B,
+      CUtensorMapL2promotion::CU_TENSOR_MAP_L2_PROMOTION_NONE,
+      CUtensorMapFloatOOBfill::CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);
+  TORCH_CHECK(res == CUDA_SUCCESS,
+              "cuTensorMapEncodeTiled failed for activation_map, error code=",
+              static_cast<int>(res));
+
+  int smem_size = STAGES * STAGE_UNROLL *
+                  (TILE_M * TILE_K * sizeof(__nv_bfloat16) +
+                   TILE_N * TILE_K * sizeof(__nv_bfloat16));
+
+  gpuErrChk(cudaFuncSetAttribute(
+      gpt_oss_router_gemm_kernel<WARP_TILE_M, TILE_M, TILE_N, TILE_K, STAGES,
+                                 STAGE_UNROLL, PROFILE>,
+      cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size));
+
+  int tiles_m = (output_features + TILE_M - 1) / TILE_M;
+  int tiles_n = (batch_size + TILE_N - 1) / TILE_N;
+
+  dim3 grid(tiles_m, tiles_n);
+  dim3 block(384);
+
+  cudaLaunchConfig_t config;
+  cudaLaunchAttribute attrs[1];
+  config.gridDim = grid;
+  config.blockDim = block;
+  config.dynamicSmemBytes = smem_size;
+  config.stream = stream;
+  config.attrs = attrs;
+  attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
+  attrs[0].val.programmaticStreamSerializationAllowed = 1;
+  config.numAttrs = 1;
+
+  cudaLaunchKernelEx(
+      &config,
+      &gpt_oss_router_gemm_kernel<WARP_TILE_M, TILE_M, TILE_N, TILE_K, STAGES,
+                                  STAGE_UNROLL, PROFILE>,
+      gC, gA, gB, bias, output_features, batch_size, input_features, weight_map,
+      activation_map, nullptr);
+}
+
+void gpt_oss_router_gemm_cuda_forward(torch::Tensor& output,
+                                      torch::Tensor input, torch::Tensor weight,
+                                      torch::Tensor bias) {
+  auto const batch_size = input.size(0);
+  auto const input_dim = input.size(1);
+  auto const output_dim = weight.size(0);
+
+  auto stream = at::cuda::getCurrentCUDAStream();
+
+  if (input.scalar_type() == at::ScalarType::BFloat16) {
+    launch_gpt_oss_router_gemm((__nv_bfloat16*)input.data_ptr(),
+                               (__nv_bfloat16*)weight.data_ptr(),
+                               (__nv_bfloat16*)output.mutable_data_ptr(),
+                               (__nv_bfloat16*)bias.data_ptr(), batch_size,
+                               output_dim, input_dim, stream);
+  } else {
+    throw std::invalid_argument("Unsupported dtype, only supports bfloat16");
+  }
+}
+
+void gpt_oss_router_gemm(torch::Tensor& output, torch::Tensor input,
+                         torch::Tensor weight, torch::Tensor bias) {
+  TORCH_CHECK(input.dim() == 2, "input must be 2D");
+  TORCH_CHECK(weight.dim() == 2, "weight must be 2D");
+  TORCH_CHECK(bias.dim() == 1, "bias must be 1D");
+  TORCH_CHECK(input.sizes()[1] == weight.sizes()[1],
+              "input.size(1) must match weight.size(1)");
+  TORCH_CHECK(weight.sizes()[0] == bias.sizes()[0],
+              "weight.size(0) must match bias.size(0)");
+  TORCH_CHECK(input.scalar_type() == at::ScalarType::BFloat16,
+              "input tensor must be bfloat16");
+  TORCH_CHECK(weight.scalar_type() == at::ScalarType::BFloat16,
+              "weight tensor must be bfloat16");
+  TORCH_CHECK(bias.scalar_type() == at::ScalarType::BFloat16,
+              "bias tensor must be bfloat16");
+  gpt_oss_router_gemm_cuda_forward(output, input, weight, bias);
+}

csrc/moe/gpt_oss_router_gemm.cuh

@@ -0,0 +1,447 @@
+/*
+ * Adapted from
+ * https://github.com/NVIDIA/TensorRT-LLM/blob/v1.3.0rc7/cpp/tensorrt_llm/kernels/tinygemm2/tinygemm2_kernel.cuh
+ * Copyright (c) 2025, The vLLM team.
+ * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * All rights reserved. SPDX-License-Identifier: Apache-2.0
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "cuda_bf16.h"
+#include <stdint.h>
+#include <stdio.h>
+#include <vector>
+
+#include "cuda_pipeline.h"
+#include <cuda.h>
+#include <cuda/barrier>
+#include <cuda/std/utility>
+#include <cuda_runtime.h>
+
+using barrier = cuda::barrier<cuda::thread_scope_block>;
+namespace cde = cuda::device::experimental;
+namespace ptx = cuda::ptx;
+
+#define gpuErrChk(ans)                    \
+  {                                       \
+    gpuAssert((ans), __FILE__, __LINE__); \
+  }
+
+inline void gpuAssert(cudaError_t code, char const* file, int line,
+                      bool abort = true) {
+  if (code != cudaSuccess) {
+    fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file,
+            line);
+    if (abort) {
+      throw std::runtime_error(cudaGetErrorString(code));
+    }
+  }
+}
+
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+__device__ uint64_t gclock64() {
+  unsigned long long int rv;
+  asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(rv));
+  return rv;
+}
+
+__device__ void ldmatrix(__nv_bfloat16 rv[2], uint32_t smem_ptr) {
+  int dst;
+  asm volatile("ldmatrix.sync.aligned.x1.m8n8.shared.b16 {%0}, [%1];\n"
+               : "=r"(dst)
+               : "r"(smem_ptr));
+  int* rvi = reinterpret_cast<int*>(&rv[0]);
+  rvi[0] = dst;
+}
+
+__device__ void ldmatrix2(__nv_bfloat16 rv[4], uint32_t smem_ptr) {
+  int x, y;
+  asm volatile("ldmatrix.sync.aligned.x2.m8n8.shared.b16 {%0, %1}, [%2];\n"
+               : "=r"(x), "=r"(y)
+               : "r"(smem_ptr));
+
+  int* rvi = reinterpret_cast<int*>(&rv[0]);
+  rvi[0] = x;
+  rvi[1] = y;
+}
+
+__device__ void ldmatrix4(__nv_bfloat16 rv[8], uint32_t smem_ptr) {
+  int x, y, z, w;
+  asm volatile(
+      "ldmatrix.sync.aligned.x4.m8n8.shared.b16 {%0, %1, %2, %3}, [%4];"
+      : "=r"(x), "=r"(y), "=r"(z), "=r"(w)
+      : "r"(smem_ptr));
+  int* rvi = reinterpret_cast<int*>(&rv[0]);
+  rvi[0] = x;
+  rvi[1] = y;
+  rvi[2] = z;
+  rvi[3] = w;
+}
+
+__device__ void HMMA_1688(float d[4], __nv_bfloat16 a[4], __nv_bfloat16 b[2],
+                          float c[4]) {
+  uint32_t const* A = reinterpret_cast<uint32_t const*>(&a[0]);
+  uint32_t const* B = reinterpret_cast<uint32_t const*>(&b[0]);
+  float const* C = reinterpret_cast<float const*>(&c[0]);
+  float* D = reinterpret_cast<float*>(&d[0]);
+
+  asm volatile(
+      "mma.sync.aligned.m16n8k8.row.col.f32.bf16.bf16.f32 "
+      "{%0,%1,%2,%3}, {%4,%5}, {%6}, {%7,%8,%9,%10};\n"
+      : "=f"(D[0]), "=f"(D[1]), "=f"(D[2]), "=f"(D[3])
+      : "r"(A[0]), "r"(A[1]), "r"(B[0]), "f"(C[0]), "f"(C[1]), "f"(C[2]),
+        "f"(C[3]));
+}
+
+__device__ void HMMA_16816(float d[4], __nv_bfloat16 a[8], __nv_bfloat16 b[4],
+                           float c[4]) {
+  uint32_t const* A = reinterpret_cast<uint32_t const*>(&a[0]);
+  uint32_t const* B = reinterpret_cast<uint32_t const*>(&b[0]);
+  float const* C = reinterpret_cast<float const*>(&c[0]);
+  float* D = reinterpret_cast<float*>(&d[0]);
+
+  asm volatile(
+      "mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 "
+      "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n"
+      : "=f"(D[0]), "=f"(D[1]), "=f"(D[2]), "=f"(D[3])
+      : "r"(A[0]), "r"(A[1]), "r"(A[2]), "r"(A[3]), "r"(B[0]), "r"(B[1]),
+        "f"(C[0]), "f"(C[1]), "f"(C[2]), "f"(C[3]));
+}
+
+__device__ void bar_wait(uint32_t bar_ptr, int phase) {
+  asm volatile(
+      "{\n"
+      ".reg .pred                P1;\n"
+      "LAB_WAIT:\n"
+      "mbarrier.try_wait.parity.shared::cta.b64 P1, [%0], %1;\n"
+      "@P1                       bra.uni DONE;\n"
+      "bra.uni                   LAB_WAIT;\n"
+      "DONE:\n"
+      "}\n" ::"r"(bar_ptr),
+      "r"(phase));
+}
+
+__device__ bool bar_try_wait(uint32_t bar_ptr, int phase) {
+  uint32_t success;
+  #ifdef INTERNAL
+  asm volatile(".pragma \"set knob DontInsertYield\";\n" : : : "memory");
+  #endif
+  asm volatile(
+      "{\n\t"
+      ".reg .pred P1; \n\t"
+      "mbarrier.try_wait.parity.shared::cta.b64 P1, [%1], %2; \n\t"
+      "selp.b32 %0, 1, 0, P1; \n\t"
+      "}"
+      : "=r"(success)
+      : "r"(bar_ptr), "r"(phase));
+  return success;
+}
+
+__device__ uint32_t elect_one_sync() {
+  uint32_t pred = 0;
+  uint32_t laneid = 0;
+  asm volatile(
+      "{\n"
+      ".reg .b32 %%rx;\n"
+      ".reg .pred %%px;\n"
+      "     elect.sync %%rx|%%px, %2;\n"
+      "@%%px mov.s32 %1, 1;\n"
+      "     mov.s32 %0, %%rx;\n"
+      "}\n"
+      : "+r"(laneid), "+r"(pred)
+      : "r"(0xFFFFFFFF));
+  return pred;
+}
+#endif
+
+struct Profile {
+  uint64_t start;
+  uint64_t weight_load_start;
+  uint64_t act_load_start;
+  uint64_t compute_start;
+  uint64_t complete;
+};
+
+template <int WARP_TILE_M, int TILE_M, int TILE_N, int TILE_K, int STAGES,
+          int STAGE_UNROLL, bool PROFILE>
+__global__ __launch_bounds__(384, 1) void gpt_oss_router_gemm_kernel(
+    __nv_bfloat16* output, __nv_bfloat16* weights, __nv_bfloat16* activations,
+    __nv_bfloat16* bias, int M, int N, int K,
+    const __grid_constant__ CUtensorMap weight_map,
+    const __grid_constant__ CUtensorMap activation_map,
+    Profile* profile = nullptr) {
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+
+  if (PROFILE && threadIdx.x == 0 && blockIdx.y == 0)
+    profile[blockIdx.x].start = gclock64();
+
+  extern __shared__ __align__(128) char smem[];
+
+  __nv_bfloat16* sh_weights = (__nv_bfloat16*)&smem[0];
+  __nv_bfloat16* sh_activations =
+      (__nv_bfloat16*)&smem[STAGES * STAGE_UNROLL * TILE_M * TILE_K *
+                            sizeof(__nv_bfloat16)];
+
+  #pragma nv_diag_suppress static_var_with_dynamic_init
+  __shared__ barrier bar_wt_ready[STAGES];
+  __shared__ barrier bar_act_ready[STAGES];
+  __shared__ barrier bar_data_consumed[STAGES];
+
+  __shared__ float4 reduction_buffer[128];
+
+  __shared__ nv_bfloat16 sh_bias[TILE_M];
+
+  if (threadIdx.x == 0) {
+    for (int i = 0; i < STAGES; i++) {
+      init(&bar_wt_ready[i], 1);
+      init(&bar_act_ready[i], 1);
+      init(&bar_data_consumed[i], 32);
+    }
+    ptx::fence_proxy_async(ptx::space_shared);
+    asm volatile("prefetch.tensormap [%0];"
+                 :
+                 : "l"(reinterpret_cast<uint64_t>(&weight_map))
+                 : "memory");
+    asm volatile("prefetch.tensormap [%0];"
+                 :
+                 : "l"(reinterpret_cast<uint64_t>(&activation_map))
+                 : "memory");
+  }
+  __syncthreads();
+
+  int warp_id = threadIdx.x / 32;
+  int lane_id = threadIdx.x % 32;
+
+  int phase = 0;
+
+  int mib = blockIdx.x * TILE_M;
+  int ni = blockIdx.y * TILE_N;
+
+  float accum[4];
+  for (int i = 0; i < 4; i++) accum[i] = 0.f;
+
+  int const K_LOOPS_DMA =
+      (K + 4 * TILE_K * STAGE_UNROLL - 1) / (4 * (TILE_K * STAGE_UNROLL));
+  int const K_LOOPS_COMPUTE = K_LOOPS_DMA;
+
+  // Data loading thread
+  if (warp_id >= 4 && elect_one_sync()) {
+    int stage = warp_id % 4;
+
+    bool weight_warp = warp_id < 8;
+    if (!weight_warp) {
+      cudaGridDependencySynchronize();
+      cudaTriggerProgrammaticLaunchCompletion();
+    }
+
+    for (int ki = 0; ki < K_LOOPS_DMA; ki++) {
+      int k = (ki * 4 + (warp_id % 4)) * TILE_K * STAGE_UNROLL;
+
+      uint64_t desc_ptr_wt = reinterpret_cast<uint64_t>(&weight_map);
+      uint64_t desc_ptr_act = reinterpret_cast<uint64_t>(&activation_map);
+
+      uint32_t bar_ptr_wt = __cvta_generic_to_shared(&bar_wt_ready[stage]);
+      uint32_t bar_ptr_act = __cvta_generic_to_shared(&bar_act_ready[stage]);
+      int bytes_wt = TILE_M * TILE_K * sizeof(__nv_bfloat16);
+      int bytes_act = TILE_N * TILE_K * sizeof(__nv_bfloat16);
+
+      bar_wait(__cvta_generic_to_shared(&bar_data_consumed[stage]), phase ^ 1);
+
+      if (weight_warp)
+        asm volatile("mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;"
+                     :
+                     : "r"(bar_ptr_wt), "r"(STAGE_UNROLL * bytes_wt));
+      if (!weight_warp)
+        asm volatile("mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;"
+                     :
+                     : "r"(bar_ptr_act), "r"(STAGE_UNROLL * bytes_act));
+
+      if (PROFILE && blockIdx.y == 0 && ki == 0 && weight_warp)
+        profile[blockIdx.x].weight_load_start = gclock64();
+      if (PROFILE && blockIdx.y == 0 && ki == 0 && !weight_warp)
+        profile[blockIdx.x].act_load_start = gclock64();
+
+      for (int i = 0; i < STAGE_UNROLL; i++) {
+        uint32_t smem_ptr_wt = __cvta_generic_to_shared(
+            &sh_weights[(stage * STAGE_UNROLL + i) * TILE_M * TILE_K]);
+        uint32_t crd0 = k + i * TILE_K;
+        uint32_t crd1 = mib;
+        if (weight_warp)
+          asm volatile(
+              "cp.async.bulk.tensor.2d.shared::cta.global.mbarrier::complete_"
+              "tx::bytes [%0], [%1, {%3,%4}], "
+              "[%2];"
+              :
+              : "r"(smem_ptr_wt), "l"(desc_ptr_wt), "r"(bar_ptr_wt), "r"(crd0),
+                "r"(crd1)
+              : "memory");
+
+        uint32_t smem_ptr_act = __cvta_generic_to_shared(
+            &sh_activations[(stage * STAGE_UNROLL + i) * TILE_N * TILE_K]);
+        crd0 = k + i * TILE_K;
+        crd1 = ni;
+        if (!weight_warp)
+          asm volatile(
+              "cp.async.bulk.tensor.2d.shared::cta.global.mbarrier::complete_"
+              "tx::bytes [%0], [%1, {%3,%4}], "
+              "[%2];"
+              :
+              : "r"(smem_ptr_act), "l"(desc_ptr_act), "r"(bar_ptr_act),
+                "r"(crd0), "r"(crd1)
+              : "memory");
+      }
+
+      stage += 4;
+      if (stage >= STAGES) {
+        stage = warp_id % 4;
+        phase ^= 1;
+      }
+    }
+    // Wait for pending loads to be consumed before exiting, to avoid race
+    for (int i = 0; i < (STAGES / 4) - 1; i++) {
+      bar_wait(__cvta_generic_to_shared(&bar_data_consumed[stage]), phase ^ 1);
+      stage += 4;
+      if (stage >= STAGES) {
+        stage = warp_id % 4;
+        phase ^= 1;
+      }
+    }
+  }
+  // Compute threads
+  else if (warp_id < 4) {
+    // Sneak the bias load into the compute warps since they're just waiting for
+    // stuff anyway
+    if (threadIdx.x < TILE_M) sh_bias[threadIdx.x] = bias[mib + threadIdx.x];
+
+    int stage = warp_id;
+
+    int phase = 0;
+    int lane_id_div8 = lane_id / 8;
+    int lane_id_mod8 = lane_id % 8;
+
+    int lane_row_offset_wt = (lane_id_div8 % 2) ? 8 : 0;
+    int lane_col_offset_wt = (lane_id_div8 / 2) ? 1 : 0;
+
+    int row_wt = lane_id_mod8 + lane_row_offset_wt;
+    int row_act = lane_id_mod8;
+
+    int row_offset_wt = (reinterpret_cast<uintptr_t>(sh_weights) / 128) % 8;
+    int row_offset_act = row_offset_wt;
+
+    uint32_t bar_ptr_wt = __cvta_generic_to_shared(&bar_wt_ready[stage]);
+    uint32_t bar_ptr_act = __cvta_generic_to_shared(&bar_act_ready[stage]);
+
+    bool weight_ready = bar_try_wait(bar_ptr_wt, phase);
+    bool act_ready = bar_try_wait(bar_ptr_act, phase);
+
+  #pragma unroll 2
+    for (int ki = 0; ki < K_LOOPS_COMPUTE; ki++) {
+      int next_stage = stage + 4;
+      int next_phase = phase;
+      if (next_stage >= STAGES) {
+        next_stage = warp_id;
+        next_phase ^= 1;
+      }
+
+      while (!weight_ready || !act_ready) {
+        weight_ready = bar_try_wait(bar_ptr_wt, phase);
+        act_ready = bar_try_wait(bar_ptr_act, phase);
+      }
+
+      if (PROFILE && blockIdx.y == 0 && threadIdx.x == 0 && ki == 0)
+        profile[blockIdx.x].compute_start = gclock64();
+
+      if (ki + 1 < K_LOOPS_COMPUTE) {
+        weight_ready = bar_try_wait(
+            __cvta_generic_to_shared(&bar_wt_ready[next_stage]), next_phase);
+        act_ready = bar_try_wait(
+            __cvta_generic_to_shared(&bar_act_ready[next_stage]), next_phase);
+      }
+
+  #pragma unroll
+      for (int su = 0; su < STAGE_UNROLL; su++) {
+        __nv_bfloat16* ptr_weights =
+            &sh_weights[(stage * STAGE_UNROLL + su) * TILE_M * TILE_K];
+        __nv_bfloat16* ptr_act =
+            &sh_activations[(stage * STAGE_UNROLL + su) * TILE_N * TILE_K];
+
+  #pragma unroll
+        for (int kii = 0; kii < TILE_K / 16; kii++) {
+          __nv_bfloat16 a[8];
+          __nv_bfloat16 b[4];
+
+          int col = 2 * kii + lane_col_offset_wt;
+          int col_sw = ((row_wt + row_offset_wt) % 8) ^ col;
+
+          ldmatrix4(a, __cvta_generic_to_shared(
+                           &ptr_weights[row_wt * TILE_K + col_sw * 8]));
+
+          col = 2 * kii + lane_id_div8;
+          col_sw = ((row_act + row_offset_act) % 8) ^ col;
+
+          ldmatrix2(b, __cvta_generic_to_shared(
+                           &ptr_act[row_act * TILE_K + 8 * col_sw]));
+
+          HMMA_16816(accum, a, b, accum);
+        }
+      }
+
+      uint32_t bar_c = __cvta_generic_to_shared(&bar_data_consumed[stage]);
+      asm volatile("mbarrier.arrive.shared::cta.b64 _, [%0];" : : "r"(bar_c));
+
+      stage = next_stage;
+      phase = next_phase;
+    }
+
+    float4 accum4;
+    accum4.x = accum[0];
+    accum4.y = accum[1];
+    accum4.z = accum[2];
+    accum4.w = accum[3];
+    reduction_buffer[threadIdx.x] = accum4;
+
+    __syncthreads();
+
+    if (warp_id == 0) {
+      int mi = mib + warp_id * WARP_TILE_M;
+      int tm = mi + lane_id / 4;
+      int tn = ni + 2 * (lane_id % 4);
+
+      float4 accum1 = reduction_buffer[32 + threadIdx.x];
+      float4 accum2 = reduction_buffer[64 + threadIdx.x];
+      float4 accum3 = reduction_buffer[96 + threadIdx.x];
+
+      accum[0] = accum[0] + accum1.x + accum2.x + accum3.x;
+      accum[1] = accum[1] + accum1.y + accum2.y + accum3.y;
+      accum[2] = accum[2] + accum1.z + accum2.z + accum3.z;
+      accum[3] = accum[3] + accum1.w + accum2.w + accum3.w;
+
+      float bias_lo = __bfloat162float(sh_bias[tm - mib]);
+      float bias_hi = __bfloat162float(sh_bias[tm + 8 - mib]);
+
+      if (tn < N && tm < M)
+        output[tn * M + tm] = __float2bfloat16(accum[0] + bias_lo);
+      if (tn + 1 < N && tm < M)
+        output[(tn + 1) * M + tm] = __float2bfloat16(accum[1] + bias_lo);
+      if (tn < N && tm + 8 < M)
+        output[tn * M + tm + 8] = __float2bfloat16(accum[2] + bias_hi);
+      if (tn + 1 < N && tm + 8 < M)
+        output[(tn + 1) * M + tm + 8] = __float2bfloat16(accum[3] + bias_hi);
+
+      if (PROFILE && blockIdx.y == 0 && threadIdx.x == 0)
+        profile[blockIdx.x].complete = gclock64();
+    }
+  }
+#endif  // end if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+}

csrc/moe/moe_ops.h

@@ -70,4 +70,8 @@ torch::Tensor router_gemm_bf16_fp32(torch::Tensor const& input,
 // Supports num_tokens in [1, 16], num_experts in {256, 384}, hidden_dim = 7168
 void dsv3_router_gemm(torch::Tensor& output, const torch::Tensor& mat_a,
                       const torch::Tensor& mat_b);
+
+// gpt-oss optimized router GEMM kernel for SM90+
+void gpt_oss_router_gemm(torch::Tensor& output, torch::Tensor input,
+                         torch::Tensor weight, torch::Tensor bias);
 #endif

csrc/moe/torch_bindings.cpp

@@ -132,6 +132,12 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, m) {
   // DeepSeek V3 optimized router GEMM for SM90+
   m.def("dsv3_router_gemm(Tensor! output, Tensor mat_a, Tensor mat_b) -> ()");
   // conditionally compiled so impl registration is in source file
+
+  // gpt-oss optimized router GEMM kernel for SM90+
+  m.def(
+      "gpt_oss_router_gemm(Tensor! output, Tensor input, Tensor weights, "
+      "Tensor bias) -> ()");
+  m.impl("gpt_oss_router_gemm", torch::kCUDA, &gpt_oss_router_gemm);
 #endif
 }
 

csrc/ops.h

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

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

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

csrc/quantization/w8a8/cutlass/scaled_mm_entry.cu

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

csrc/torch_bindings.cpp

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

docker/Dockerfile.xpu

@@ -76,19 +76,22 @@ ENV UV_LINK_MODE="copy"
 RUN --mount=type=cache,target=/root/.cache/uv \
     --mount=type=bind,src=requirements/common.txt,target=/workspace/vllm/requirements/common.txt \
     --mount=type=bind,src=requirements/xpu.txt,target=/workspace/vllm/requirements/xpu.txt \
+    --mount=type=bind,src=requirements/xpu-test.in,target=/workspace/vllm/requirements/xpu-test.in \
     uv pip install --upgrade pip && \
-    uv pip install -r requirements/xpu.txt
-
- # used for suffix method speculative decoding
- # build deps for proto + nanobind-based extensions to set up the build environment
-RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install grpcio-tools protobuf nanobind
- # arctic-inference is built from source which needs torch-xpu properly installed first
-RUN --mount=type=cache,target=/root/.cache/uv \
+    uv pip install -r requirements/xpu.txt && \
+    uv pip compile /workspace/vllm/requirements/xpu-test.in \
+        -o /workspace/vllm/requirements/xpu-test.txt \
+        -c /workspace/vllm/requirements/xpu.txt \
+        --index-strategy unsafe-best-match \
+        --extra-index-url ${PIP_EXTRA_INDEX_URL} \
+        --python-version ${PYTHON_VERSION} && \
+    uv pip install grpcio-tools protobuf nanobind && \
     source /opt/intel/oneapi/setvars.sh --force && \
     source /opt/intel/oneapi/ccl/2021.15/env/vars.sh --force && \
-    export CMAKE_PREFIX_PATH="$(python -c 'import site; print(site.getsitepackages()[0])'):${CMAKE_PREFIX_PATH}" && \
-    uv pip install --no-build-isolation arctic-inference==0.1.1
+    export CMAKE_PREFIX_PATH="$(python3 -c 'import site; print(site.getsitepackages()[0])'):${CMAKE_PREFIX_PATH}" && \
+    uv pip install --no-build-isolation -r /workspace/vllm/requirements/xpu-test.txt
+
+
 
 ENV LD_LIBRARY_PATH="$LD_LIBRARY_PATH:/usr/local/lib/"
 

docs/mkdocs/hooks/generate_examples.py

@@ -23,15 +23,18 @@ def title(text: str) -> str:
     # Custom substitutions
     subs = {
         "io": "IO",
-        "api": "API",
+        "rl": "RL",
+        "api(s?)": r"API\1",
         "cli": "CLI",
         "cpu": "CPU",
+        "ipc": "IPC",
         "llm": "LLM",
         "mae": "MAE",
         "ner": "NER",
         "tpu": "TPU",
         "gguf": "GGUF",
         "lora": "LoRA",
+        "nccl": "NCCL",
         "rlhf": "RLHF",
         "vllm": "vLLM",
         "openai": "OpenAI",
@@ -196,6 +199,11 @@ class Example:
 
 
 def on_startup(command: Literal["build", "gh-deploy", "serve"], dirty: bool):
+    # Monkey-patch dirname_to_title in awesome-nav so that sub-directory names are
+    # title-cased (e.g. "Offline Inference" instead of "Offline inference").
+    import mkdocs_awesome_nav.nav.directory as _nav_dir
+
+    _nav_dir.dirname_to_title = title
     logger.info("Generating example documentation")
     logger.debug("Root directory: %s", ROOT_DIR.resolve())
     logger.debug("Example directory: %s", EXAMPLE_DIR.resolve())

docs/models/pooling_models.md

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

docs/models/supported_models.md

@@ -707,7 +707,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `GraniteSpeechForConditionalGeneration` | Granite Speech | T + A | `ibm-granite/granite-speech-3.3-8b` | ✅︎ | ✅︎ |
 | `HCXVisionForCausalLM` | HyperCLOVAX-SEED-Vision-Instruct-3B | T + I<sup>+</sup> + V<sup>+</sup> | `naver-hyperclovax/HyperCLOVAX-SEED-Vision-Instruct-3B` | | |
 | `HCXVisionV2ForCausalLM` | HyperCLOVAX-SEED-Think-32B | T + I<sup>+</sup> + V<sup>+</sup> | `naver-hyperclovax/HyperCLOVAX-SEED-Think-32B` | | |
-| `H2OVLChatModel` | H2OVL | T + I<sup>E+</sup> | `h2oai/h2ovl-mississippi-800m`, `h2oai/h2ovl-mississippi-2b`, etc. | | ✅︎ |
+| `H2OVLChatModel` | H2OVL | T + I<sup>E+</sup> | `h2oai/h2ovl-mississippi-800m`, `h2oai/h2ovl-mississippi-2b`, etc. | ✅︎ | ✅︎ |
 | `HunYuanVLForConditionalGeneration` | HunyuanOCR | T + I<sup>E+</sup> | `tencent/HunyuanOCR`, etc. | ✅︎ | ✅︎ |
 | `Idefics3ForConditionalGeneration` | Idefics3 | T + I | `HuggingFaceM4/Idefics3-8B-Llama3`, etc. | ✅︎ | |
 | `IsaacForConditionalGeneration` | Isaac | T + I<sup>+</sup> | `PerceptronAI/Isaac-0.1` | ✅︎ | ✅︎ |
@@ -834,6 +834,7 @@ The following table lists those that are tested in vLLM.
 | `CLIPModel` | CLIP | T / I | `openai/clip-vit-base-patch32`, `openai/clip-vit-large-patch14`, etc. | | |
 | `ColModernVBertForRetrieval` | ColModernVBERT | T / I | `ModernVBERT/colmodernvbert-merged` | | |
 | `ColPaliForRetrieval` | ColPali | T / I | `vidore/colpali-v1.3-hf` | | |
+| `ColQwen3_5` | ColQwen3.5 | T + I + V | `athrael-soju/colqwen3.5-4.5B-v3` | | |
 | `LlamaNemotronVLModel` | Llama Nemotron Embedding + SigLIP | T + I | `nvidia/llama-nemotron-embed-vl-1b-v2` | | |
 | `LlavaNextForConditionalGeneration`<sup>C</sup> | LLaVA-NeXT-based | T / I | `royokong/e5-v` | | ✅︎ |
 | `Phi3VForCausalLM`<sup>C</sup> | Phi-3-Vision-based | T + I | `TIGER-Lab/VLM2Vec-Full` | | ✅︎ |

docs/training/async_rl.md

@@ -0,0 +1,63 @@
+# Async Reinforcement Learning
+
+## Overview
+
+In a standard RL training loop, generation and training happen sequentially: the policy generates rollouts, then training runs on those rollouts, and the cycle repeats. During generation the training accelerators sit idle, and vice versa.
+
+The **one-off pipelining** approach separates the generation and training phases into two parallel coroutines, allowing the model to generate new samples while simultaneously training on previously generated data. This can lead to better GPU utilization and greater training throughput.
+
+However, this overlap introduces a complication: weights must be updated in the inference engine mid-flight, while requests may still be in progress.
+
+## The Pause and Resume API
+
+To safely update weights while the inference engine is running, vLLM provides `pause_generation` and `resume_generation` methods. These let the trainer coordinate a clean window for weight synchronization without losing in-flight work.
+
+### pause_generation
+
+```python
+await engine.pause_generation(mode="keep", clear_cache=True)
+```
+
+The `mode` parameter controls how in-flight requests are handled:
+
+| Mode | Behavior |
+| ---- | -------- |
+| `"abort"` | Abort all in-flight requests immediately and return partial results (default) |
+| `"wait"` | Wait for all in-flight requests to finish before pausing |
+| `"keep"` | Freeze requests in the queue; they resume when `resume_generation` is called |
+
+The `clear_cache` parameter controls whether to clear the KV cache and prefix cache after pausing.
+
+### resume_generation
+
+```python
+await engine.resume_generation()
+```
+
+Resumes the scheduler after a pause. Any requests frozen with `mode="keep"` will continue generating.
+
+### HTTP Endpoints
+
+When using the vLLM HTTP server, the same functionality is available via:
+
+- `POST /pause?mode=keep` - Pause generation
+- `POST /resume` - Resume generation
+
+!!! note "Data Parallelism"
+    When using data parallelism with vLLM's **internal load balancer** (i.e. `data_parallel_backend="ray"`), pause and resume are handled automatically across all DP ranks -- a single call is sufficient. When using an **external load balancer** (i.e. multiple independent vLLM instances behind a proxy), you must send pause and resume requests to **every** engine instance individually before and after the weight update.
+
+## Typical Async RL Flow
+
+A typical async RL loop with weight syncing looks like this:
+
+1. Start generating rollouts from the current policy
+2. Once trainer has new weights to update to, pause generation with `mode="keep"`
+3. Sync the updated weights from the trainer to the inference engine (see [Weight Transfer](weight_transfer/README.md))
+4. Resume generation -- in-flight requests continue with the new weights
+5. Repeat
+
+The key insight is that requests paused with `mode="keep"` will produce tokens from the **old** weights before the pause and tokens from the **new** weights after resume. The `clear_cache` parameter controls whether the KV cache is invalidated during the pause. When `clear_cache=True`, previously cached key-value entries are discarded, so all tokens generated after resume will be computed entirely with the new weights. When `clear_cache=False`, existing KV cache entries are retained, meaning some tokens in context may still reflect the old weights (stale KV cache).
+
+## Example
+
+The [async RLHF example](../examples/rl/rlhf_async_new_apis.md) demonstrates this pattern with `vllm.AsyncLLMEngine`, NCCL weight transfer, and mid-flight pause/resume with validation.

docs/training/rlhf.md

@@ -16,11 +16,9 @@ The following open-source RL libraries use vLLM for fast rollouts (sorted alphab
 - [Unsloth](https://github.com/unslothai/unsloth)
 - [verl](https://github.com/volcengine/verl)
 
-See the following basic examples to get started if you don't want to use an existing library:
+For weight synchronization between training and inference, see the [Weight Transfer](weight_transfer/README.md) documentation, which covers the pluggable backend system with [NCCL](weight_transfer/nccl.md) (multi-GPU) and [IPC](weight_transfer/ipc.md) (same-GPU) engines.
 
-- [Training and inference processes are located on separate GPUs (inspired by OpenRLHF)](../examples/offline_inference/rlhf.md)
-- [Training and inference processes are colocated on the same GPUs using Ray](../examples/offline_inference/rlhf_colocate.md)
-- [Utilities for performing RLHF with vLLM](../examples/offline_inference/rlhf_utils.md)
+For pipelining generation and training to improve GPU utilization and throughput, see the [Async Reinforcement Learning](async_rl.md) guide, which covers the pause/resume API for safely updating weights mid-flight.
 
 See the following notebooks showing how to use vLLM for GRPO:
 

docs/training/weight_transfer/README.md

@@ -0,0 +1,78 @@
+# Weight Transfer
+
+vLLM provides a pluggable weight transfer system for synchronizing model weights from a training process to the inference engine during reinforcement learning (RL) workflows. This is essential for RLHF, GRPO, and other online RL methods where the policy model is iteratively updated during training and the updated weights must be reflected in the inference engine for rollout generation.
+
+## Architecture
+
+The weight transfer system follows a **two-phase protocol** with a pluggable backend design:
+
+1. **Initialization** (`init_weight_transfer_engine`): Establishes the communication channel between the trainer and inference workers. Called once before the training loop begins.
+2. **Weight Update** (`update_weights`): Transfers updated weights from the trainer to the inference engine. Called after each training step (or batch of steps).
+
+## Available Backends
+
+| Backend | Transport | Use Case |
+| ------- | --------- | -------- |
+| [NCCL](nccl.md) | NCCL broadcast | Separate GPUs for training and inference |
+| [IPC](ipc.md) | CUDA IPC handles | Colocated training and inference on same GPU |
+
+## Configuration
+
+Specify the weight transfer backend through `WeightTransferConfig`. The backend determines which engine handles the weight synchronization.
+
+### Programmatic (Offline Inference)
+
+```python
+from vllm import LLM
+from vllm.config import WeightTransferConfig
+
+llm = LLM(
+    model="my-model",
+    weight_transfer_config=WeightTransferConfig(backend="nccl"),  # or "ipc"
+)
+```
+
+### CLI (Online Serving)
+
+```bash
+vllm serve my-model \
+    --weight-transfer-config '{"backend": "nccl"}'
+```
+
+The `backend` field accepts `"nccl"` (default) or `"ipc"`.
+
+## API Endpoints
+
+When running vLLM as an HTTP server, the following endpoints are available for weight transfer:
+
+| Endpoint | Method | Description |
+| -------- | ------ | ----------- |
+| `/init_weight_transfer_engine` | POST | Initialize the weight transfer engine with backend-specific info |
+| `/update_weights` | POST | Trigger a weight update with backend-specific metadata |
+| `/pause` | POST | Pause generation before weight sync to handle inflight requests |
+| `/resume` | POST | Resume generation after weight sync |
+| `/get_world_size` | GET | Get the number of inference workers (useful for NCCL world size calculation) |
+
+!!! note
+    The HTTP weight transfer endpoints require `VLLM_SERVER_DEV_MODE=1` to be set.
+
+## Trainer-Side API
+
+Both backends provide static methods that the trainer calls to send weights. The general pattern is:
+
+```python
+# 1. Initialize the transfer engine (backend-specific)
+EngineClass.trainer_init(init_info)
+
+# 2. Send weights to inference workers
+EngineClass.trainer_send_weights(
+    iterator=model.named_parameters(),
+    trainer_args=backend_specific_args,
+)
+```
+
+See the [NCCL](nccl.md) and [IPC](ipc.md) pages for backend-specific trainer APIs and full examples.
+
+## Extending the System
+
+The weight transfer system is designed to be extensible. You can implement custom backends by subclassing `WeightTransferEngine` and registering them with the factory. See the [Base Class](base.md) page for details.

docs/training/weight_transfer/base.md

@@ -0,0 +1,162 @@
+# Base Class and Custom Engines
+
+The weight transfer system is built on an abstract base class that defines the contract between vLLM's worker infrastructure and the transport backend. You can implement custom backends by subclassing `WeightTransferEngine` and registering them with the `WeightTransferEngineFactory`.
+
+## WeightTransferEngine
+
+The `WeightTransferEngine` is a generic abstract class parameterized by two dataclass types:
+
+- **`TInitInfo`** (extends `WeightTransferInitInfo`): Backend-specific initialization parameters.
+- **`TUpdateInfo`** (extends `WeightTransferUpdateInfo`): Backend-specific weight update metadata.
+
+### Abstract Methods
+
+Subclasses must implement these four methods:
+
+| Method | Side | Description |
+| ------ | ---- | ----------- |
+| `init_transfer_engine(init_info)` | Inference | Initialize the communication channel on each inference worker |
+| `receive_weights(update_info, load_weights)` | Inference | Receive weights and call `load_weights` incrementally |
+| `shutdown()` | Inference | Clean up resources |
+| `trainer_send_weights(iterator, trainer_args)` | Trainer | Static method to send weights from the trainer process |
+
+### Request Classes
+
+The API-level request classes provide backend-agnostic serialization using plain dictionaries. The engine's `parse_init_info` and `parse_update_info` methods convert these dictionaries into typed dataclasses.
+
+```python
+from vllm.distributed.weight_transfer.base import (
+    WeightTransferInitRequest,
+    WeightTransferUpdateRequest,
+)
+
+# Init request (dict is converted to backend-specific TInitInfo)
+init_request = WeightTransferInitRequest(
+    init_info={"master_address": "10.0.0.1", "master_port": 29500, ...}
+)
+
+# Update request (dict is converted to backend-specific TUpdateInfo)
+update_request = WeightTransferUpdateRequest(
+    update_info={"names": [...], "dtype_names": [...], "shapes": [...]}
+)
+```
+
+### WeightTransferUpdateInfo
+
+The base `WeightTransferUpdateInfo` includes an `is_checkpoint_format` flag:
+
+```python
+@dataclass
+class WeightTransferUpdateInfo(ABC):
+    is_checkpoint_format: bool = True
+```
+
+When `is_checkpoint_format=True` (the default), vLLM applies layerwise weight processing (repacking, renaming, etc.) on the received weights before loading them. Set to `False` if the trainer has already converted weights to the kernel format expected by the model.
+
+## Implementing a Custom Engine
+
+To create a custom weight transfer backend:
+
+### 1. Define Info Dataclasses
+
+```python
+from dataclasses import dataclass
+from vllm.distributed.weight_transfer.base import (
+    WeightTransferEngine,
+    WeightTransferInitInfo,
+    WeightTransferUpdateInfo,
+)
+
+@dataclass
+class MyInitInfo(WeightTransferInitInfo):
+    endpoint: str
+    token: str
+
+@dataclass
+class MyUpdateInfo(WeightTransferUpdateInfo):
+    names: list[str]
+    dtype_names: list[str]
+    shapes: list[list[int]]
+    # Add custom fields as needed
+```
+
+### 2. Implement the Engine
+
+```python
+from collections.abc import Callable, Iterator
+from typing import Any
+import torch
+
+class MyWeightTransferEngine(WeightTransferEngine[MyInitInfo, MyUpdateInfo]):
+    init_info_cls = MyInitInfo
+    update_info_cls = MyUpdateInfo
+
+    def init_transfer_engine(self, init_info: MyInitInfo) -> None:
+        # Set up connection to trainer using init_info.endpoint, etc.
+        ...
+
+    def receive_weights(
+        self,
+        update_info: MyUpdateInfo,
+        load_weights: Callable[[list[tuple[str, torch.Tensor]]], None],
+    ) -> None:
+        # Receive each weight and call load_weights incrementally
+        for name, dtype_name, shape in zip(
+            update_info.names, update_info.dtype_names, update_info.shapes
+        ):
+            dtype = getattr(torch, dtype_name)
+            weight = self._fetch_weight(name, shape, dtype)
+            load_weights([(name, weight)])
+
+    def shutdown(self) -> None:
+        # Clean up resources
+        ...
+
+    @staticmethod
+    def trainer_send_weights(
+        iterator: Iterator[tuple[str, torch.Tensor]],
+        trainer_args: dict[str, Any],
+    ) -> None:
+        # Send weights from the trainer process
+        for name, tensor in iterator:
+            # Send tensor via custom transport
+            ...
+```
+
+!!! important
+    The `load_weights` callable passed to `receive_weights` should be called **incrementally** (one or a few weights at a time) rather than accumulating all weights first. This avoids GPU out-of-memory errors with large models.
+
+### 3. Register with the Factory
+
+```python
+from vllm.distributed.weight_transfer.factory import WeightTransferEngineFactory
+
+# Option 1: Lazy loading (recommended for built-in engines)
+WeightTransferEngineFactory.register_engine(
+    "my_backend",
+    "my_package.my_module",
+    "MyWeightTransferEngine",
+)
+
+# Option 2: Direct class registration
+WeightTransferEngineFactory.register_engine(
+    "my_backend",
+    MyWeightTransferEngine,
+)
+```
+
+Once registered, users can select your backend via `WeightTransferConfig(backend="my_backend")`.
+
+## WeightTransferEngineFactory
+
+The factory uses a registry pattern with lazy loading. Built-in engines (`nccl` and `ipc`) are registered at import time but their modules are only loaded when the backend is actually requested. This avoids importing heavy dependencies (like NCCL communicators) when they aren't needed.
+
+```python
+from vllm.distributed.weight_transfer.factory import WeightTransferEngineFactory
+
+# Create an engine from config
+engine = WeightTransferEngineFactory.create_engine(
+    config=weight_transfer_config,
+    parallel_config=parallel_config,
+)
+```

docs/training/weight_transfer/ipc.md

@@ -0,0 +1,73 @@
+# IPC Engine
+
+The IPC weight transfer engine uses **CUDA IPC** (Inter-Process Communication) handles to share GPU memory directly between the trainer and inference workers on the **same node and same GPU**. This avoids any data copying, making it a efficient option when colocating training and inference.
+
+## When to Use IPC
+
+- Training and inference on the **same GPU** (colocated)
+- You want to minimize memory overhead by sharing tensors in-place
+
+## How It Works
+
+1. The trainer creates CUDA tensors for each weight and generates IPC handles using `torch.multiprocessing.reductions.reduce_tensor`.
+2. IPC handles are sent to the inference engine via **Ray.remote()** or **HTTP POST**.
+3. The inference worker reconstructs the tensors from the handles, reading directly from the trainer's GPU memory.
+
+!!! warning
+    IPC handles involve sending serialized Python objects. When using HTTP transport, you must set `VLLM_ALLOW_INSECURE_SERIALIZATION=1` on both the server and client. This is because IPC handles are pickled and base64-encoded for HTTP transmission.
+
+## Initialization
+
+The IPC backend requires no initialization on either side. The `init_transfer_engine` call is a no-op for IPC.
+
+## Sending Weights
+
+IPC supports two transport modes for delivering the handles:
+
+### Ray Mode
+
+Used when vLLM is running as a Ray actor:
+
+```python
+from vllm.distributed.weight_transfer.ipc_engine import (
+    IPCTrainerSendWeightsArgs,
+    IPCWeightTransferEngine,
+)
+
+trainer_args = IPCTrainerSendWeightsArgs(
+    mode="ray",
+    llm_handle=llm_actor_handle,
+)
+
+IPCWeightTransferEngine.trainer_send_weights(
+    iterator=model.named_parameters(),
+    trainer_args=trainer_args,
+)
+```
+
+In Ray mode, the engine calls `llm_handle.update_weights.remote(...)` directly, passing the IPC handles via Ray's serialization.
+
+### HTTP Mode
+
+Used when vLLM is running as an HTTP server:
+
+```python
+trainer_args = IPCTrainerSendWeightsArgs(
+    mode="http",
+    url="http://localhost:8000",
+)
+
+IPCWeightTransferEngine.trainer_send_weights(
+    iterator=model.named_parameters(),
+    trainer_args=trainer_args,
+)
+```
+
+In HTTP mode, IPC handles are pickled, base64-encoded, and sent as JSON to the `/update_weights` endpoint.
+
+See [`IPCTrainerSendWeightsArgs`](https://github.com/vllm-project/vllm/blob/main/vllm/distributed/weight_transfer/ipc_engine.py) for the full list of configurable fields.
+
+## Examples
+
+- [RLHF with IPC weight syncing (offline, Ray)](../../examples/rl/rlhf_ipc.md) - Colocated training and inference on a single GPU using Ray placement groups and CUDA IPC handles
+- [RLHF with IPC weight syncing (online serving, HTTP)](../../examples/rl/rlhf_http_ipc.md) - Weight transfer with a vLLM HTTP server where both server and trainer share the same GPU

docs/training/weight_transfer/nccl.md

@@ -0,0 +1,110 @@
+# NCCL Engine
+
+The NCCL weight transfer engine uses [NCCL](https://developer.nvidia.com/nccl) broadcast operations to transfer weights from the trainer to inference workers. It supports **multi-node** and **multi-GPU** setups where the trainer and inference engine run on separate GPUs.
+
+## When to Use NCCL
+
+- Training and inference on **separate GPUs** (possibly across nodes)
+- **Tensor-parallel** inference with multiple workers that all need the updated weights
+- You need high-bandwidth, low-latency weight transfer over NVLink or InfiniBand
+
+## How It Works
+
+1. The trainer and all inference workers join a shared NCCL process group using `StatelessProcessGroup` (vLLM's torch.distributed-independent group abstraction).
+2. The trainer broadcasts weights to all workers simultaneously. Each worker receives and loads weights incrementally.
+3. Optionally, **packed tensor broadcasting** batches multiple small tensors into larger buffers with double/triple buffering and CUDA stream overlap for higher throughput. This implementation is based on [NeMo-RL's packed tensor](https://github.com/NVIDIA-NeMo/RL/blob/main/nemo_rl/utils/packed_tensor.py).
+
+## Initialization
+
+NCCL requires explicit process group setup. The trainer and inference workers must agree on a master address, port, and world size.
+
+### Inference Side
+
+```python
+from vllm.distributed.weight_transfer.base import WeightTransferInitRequest
+
+# rank_offset accounts for the trainer occupying rank 0
+llm.init_weight_transfer_engine(
+    WeightTransferInitRequest(
+        init_info=dict(
+            master_address=master_address,
+            master_port=master_port,
+            rank_offset=1,
+            world_size=world_size,  # trainer + all inference workers
+        )
+    )
+)
+```
+
+### Trainer Side
+
+```python
+from vllm.distributed.weight_transfer.nccl_engine import (
+    NCCLWeightTransferEngine,
+)
+
+group = NCCLWeightTransferEngine.trainer_init(
+    dict(
+        master_address=master_address,
+        master_port=master_port,
+        world_size=world_size,
+    )
+)
+```
+
+!!! note
+    `trainer_init` always assigns the trainer to rank 0. Inference workers start at `rank_offset` (typically 1).
+
+## Sending Weights
+
+```python
+from vllm.distributed.weight_transfer.nccl_engine import (
+    NCCLTrainerSendWeightsArgs,
+    NCCLWeightTransferEngine,
+)
+
+trainer_args = NCCLTrainerSendWeightsArgs(
+    group=group,
+    packed=True,  # use packed broadcasting for efficiency
+)
+
+NCCLWeightTransferEngine.trainer_send_weights(
+    iterator=model.named_parameters(),
+    trainer_args=trainer_args,
+)
+```
+
+See [`NCCLTrainerSendWeightsArgs`](https://github.com/vllm-project/vllm/blob/main/vllm/distributed/weight_transfer/nccl_engine.py) for the full list of configurable fields.
+
+### Packed Tensor Broadcasting
+
+When `packed=True`, multiple weight tensors are packed into large contiguous buffers before broadcasting. This reduces the number of NCCL operations and uses double/triple buffering with dedicated CUDA streams for overlap between packing, broadcasting, and unpacking.
+
+Both the trainer (`NCCLTrainerSendWeightsArgs`) and inference side (`NCCLWeightTransferUpdateInfo`) must use matching `packed_buffer_size_bytes` and `packed_num_buffers` values.
+
+## Receiving Weights (Inference Side)
+
+The inference side triggers weight reception by calling `update_weights`:
+
+```python
+from vllm.distributed.weight_transfer.base import WeightTransferUpdateRequest
+
+llm.update_weights(
+    WeightTransferUpdateRequest(
+        update_info=dict(
+            names=names,
+            dtype_names=dtype_names,
+            shapes=shapes,
+            packed=True,
+        )
+    )
+)
+```
+
+The `names`, `dtype_names`, and `shapes` lists describe each parameter. These must match the order in which the trainer iterates over its parameters.
+
+## Examples
+
+- [RLHF with NCCL weight syncing (offline, Ray)](../../examples/rl/rlhf_nccl.md) - Trainer on one GPU, 2x tensor-parallel vLLM engine on two others, with packed NCCL weight broadcast
+- [RLHF with async weight syncing (offline, Ray)](../../examples/rl/rlhf_async_new_apis.md) - Async generation with mid-flight pause, weight sync, resume, and validation against a fresh model
+- [RLHF with NCCL weight syncing (online serving, HTTP)](../../examples/rl/rlhf_http_nccl.md) - Weight transfer with a running vLLM HTTP server using HTTP control plane and NCCL data plane

examples/offline_inference/audio_language.py

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

examples/offline_inference/rlhf.py

@@ -1,147 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Demonstrates reinforcement learning from human feedback (RLHF) using vLLM and Ray.
-
-The script separates training and inference workloads onto distinct GPUs
-so that Ray can manage process placement and inter-process communication.
-A Hugging Face Transformer model occupies GPU 0 for training, whereas a
-tensor-parallel vLLM inference engine occupies GPU 1–2.
-
-The example performs the following steps:
-
-* Load the training model on GPU 0.
-* Split the inference model across GPUs 1–2 using vLLM's tensor parallelism
-  and Ray placement groups.
-* Generate text from a list of prompts using the inference engine.
-* Update the weights of the training model and broadcast the updated weights
-  to the inference engine by using a Ray collective RPC group. Note that
-  for demonstration purposes we simply zero out the weights.
-
-For a production-ready implementation that supports multiple training and
-inference replicas, see the OpenRLHF framework:
-https://github.com/OpenRLHF/OpenRLHF
-
-This example assumes a single-node cluster with three GPUs, but Ray
-supports multi-node clusters. vLLM expects the GPUs are only used for vLLM
-workloads. Residual GPU activity interferes with vLLM memory profiling and
-causes unexpected behavior.
-"""
-
-import os
-
-import ray
-import torch
-from ray.util.placement_group import placement_group
-from ray.util.scheduling_strategies import PlacementGroupSchedulingStrategy
-from rlhf_utils import stateless_init_process_group
-from transformers import AutoModelForCausalLM
-
-from vllm import LLM, SamplingParams
-from vllm.utils.network_utils import get_ip, get_open_port
-
-
-class MyLLM(LLM):
-    """Configure the vLLM worker for Ray placement group execution."""
-
-    def __init__(self, *args, **kwargs):
-        # Remove the top-level CUDA_VISIBLE_DEVICES variable set by Ray
-        # so that vLLM can manage its own device placement within the worker.
-        os.environ.pop("CUDA_VISIBLE_DEVICES", None)
-        super().__init__(*args, **kwargs)
-
-
-# Load the OPT-125M model onto GPU 0 for the training workload.
-train_model = AutoModelForCausalLM.from_pretrained("facebook/opt-125m")
-train_model.to("cuda:0")
-
-# Initialize Ray and set the visible devices. The vLLM engine will
-# be placed on GPUs 1 and 2.
-os.environ["CUDA_VISIBLE_DEVICES"] = "1,2"
-ray.init()
-
-# Create a placement group that reserves GPU 1–2 for the vLLM inference engine.
-# Learn more about Ray placement groups:
-# https://docs.ray.io/en/latest/ray-core/scheduling/placement-group.html
-pg_inference = placement_group([{"GPU": 1, "CPU": 0}] * 2)
-ray.get(pg_inference.ready())
-scheduling_inference = PlacementGroupSchedulingStrategy(
-    placement_group=pg_inference,
-    placement_group_capture_child_tasks=True,
-    placement_group_bundle_index=0,
-)
-
-# Launch the vLLM inference engine. The `enforce_eager` flag reduces
-# start-up latency.
-llm = ray.remote(
-    num_cpus=0,
-    num_gpus=0,
-    scheduling_strategy=scheduling_inference,
-)(MyLLM).remote(
-    model="facebook/opt-125m",
-    enforce_eager=True,
-    worker_extension_cls="rlhf_utils.WorkerExtension",
-    tensor_parallel_size=2,
-    distributed_executor_backend="ray",
-)
-
-# Generate text from the prompts.
-prompts = [
-    "Hello, my name is",
-    "The president of the United States is",
-    "The capital of France is",
-    "The future of AI is",
-]
-
-sampling_params = SamplingParams(temperature=0)
-
-outputs = ray.get(llm.generate.remote(prompts, sampling_params))
-
-print("-" * 50)
-for output in outputs:
-    prompt = output.prompt
-    generated_text = output.outputs[0].text
-    print(f"Prompt: {prompt!r}\nGenerated text: {generated_text!r}")
-    print("-" * 50)
-
-# Set up the communication channel between the training process and the
-# inference engine.
-master_address = get_ip()
-master_port = get_open_port()
-
-handle = llm.collective_rpc.remote(
-    "init_weight_update_group", args=(master_address, master_port, 1, 3)
-)
-
-model_update_group = stateless_init_process_group(
-    master_address, master_port, 0, 3, torch.device("cuda:0")
-)
-ray.get(handle)
-
-# Simulate a training step by zeroing out all model weights.
-# In a real RLHF training loop the weights would be updated using the gradient
-# from an RL objective such as PPO on a reward model.
-for name, p in train_model.named_parameters():
-    p.data.zero_()
-
-# Synchronize the updated weights to the inference engine.
-for name, p in train_model.named_parameters():
-    dtype_name = str(p.dtype).split(".")[-1]
-    handle = llm.collective_rpc.remote(
-        "update_weight", args=(name, dtype_name, p.shape)
-    )
-    model_update_group.broadcast(p, src=0, stream=torch.cuda.current_stream())
-    ray.get(handle)
-
-# Verify that the inference weights have been updated.
-assert all(ray.get(llm.collective_rpc.remote("check_weights_changed")))
-
-# Generate text with the updated model. The output is expected to be nonsense
-# because the weights are zero.
-outputs_updated = ray.get(llm.generate.remote(prompts, sampling_params))
-print("-" * 50)
-for output in outputs_updated:
-    prompt = output.prompt
-    generated_text = output.outputs[0].text
-    print(f"Prompt: {prompt!r}\nGenerated text: {generated_text!r}")
-    print("-" * 50)

examples/offline_inference/rlhf_colocate.py

@@ -1,256 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Demonstrates how to co-locate a vLLM inference worker and training
-actors on the same set of GPUs for reinforcement learning from human feedback
-(RLHF) workloads.
-
-Ray serves as the distributed execution framework in this example. Ray
-placement groups allocate both training actors and vLLM workers to the
-same GPU bundles, enabling fast, in-GPU communication between the two
-components.
-
-The script shows how to do the following:
-
-* Configure environment variables (`VLLM_RAY_PER_WORKER_GPUS` and
-  `VLLM_RAY_BUNDLE_INDICES`) so that vLLM workers land on the desired
-  devices.
-* Exchange tensors between processes by means of CUDA inter-process
-  communication (IPC). CUDA IPC sidesteps NCCL limitations that occur
-  when multiple processes share a single GPU.
-
-Note that this example assumes a single-node cluster with four GPUs, but Ray
-supports multi-node clusters. vLLM expects exclusive use of the GPUs during
-its initialization for memory profiling. Residual GPU activity interferes
-with vLLM memory profiling and causes unexpected behavior.
-
-Learn more about Ray placement groups:
-https://docs.ray.io/en/latest/placement-groups.html
-"""
-
-import gc
-import os
-import sys
-
-import ray
-import torch
-import zmq
-from ray.util.placement_group import placement_group
-from ray.util.scheduling_strategies import PlacementGroupSchedulingStrategy
-from torch.multiprocessing.reductions import reduce_tensor
-
-from vllm import LLM
-
-if torch.version.hip is not None:
-    print("Skipping test for ROCm. Ray is unsupported on vLLM ROCm.")
-    sys.exit(0)
-
-
-class MyLLM(LLM):
-    """Configure the vLLM worker for Ray placement group execution.
-
-    The constructor sets environment variables that allow multiple vLLM
-    workers to share a single physical GPU and that encode the bundle
-    indices assigned by the placement group.
-
-    Args:
-        *args: Positional arguments forwarded to `vllm.LLM`.
-        bundle_indices (list[int]): Placement-group bundle indices
-            assigned to this worker.
-        **kwargs: Keyword arguments forwarded to `vllm.LLM`.
-    """
-
-    def __init__(self, *args, bundle_indices: list[int], **kwargs):
-        # Prevent Ray from manipulating the top-level CUDA_VISIBLE_DEVICES variable
-        # so that vLLM can its own device placement inside the worker.
-        os.environ.pop("CUDA_VISIBLE_DEVICES", None)
-        # Each worker uses 0.4 GPU so that two instances fit on the same GPUs.
-        os.environ["VLLM_RAY_PER_WORKER_GPUS"] = "0.4"
-        os.environ["VLLM_RAY_BUNDLE_INDICES"] = ",".join(map(str, bundle_indices))
-        print(f"creating LLM with bundle_indices={bundle_indices}")
-        super().__init__(*args, **kwargs)
-
-
-class RayTrainingActor:
-    """Training actor that hosts a Facebook OPT-125M model from Hugging Face.
-
-    The model is loaded onto the first GPU assigned to this actor, and expose
-    the CUDA IPC handles so that colocated vLLM workers can map tensors
-    directly.
-    """
-
-    def __init__(self):
-        # Ray sets CUDA_VISIBLE_DEVICES to the GPUs assigned to this actor.
-        from transformers import AutoModelForCausalLM
-
-        self.model = AutoModelForCausalLM.from_pretrained("facebook/opt-125m")
-        self.model.to("cuda:0")
-        # Zero out all the parameters.
-        for name, p in self.model.named_parameters():
-            p.data.zero_()
-        torch.accelerator.synchronize()
-        # The argument for `get_device_uuid` is the index of the GPU in the
-        # list of visible devices.
-        from vllm.platforms import current_platform
-
-        self.device_uuid = current_platform.get_device_uuid(0)
-        self.zmq_context = zmq.Context()
-        self.zmq_address_counter = 0
-        self.zmq_handle = None
-
-    def report_device_id(self) -> str:
-        return self.device_uuid
-
-    def get_zmq_handles(self) -> dict[str, str]:
-        suffix = f"{self.device_uuid}-{self.zmq_address_counter}"
-        self.zmq_handle = f"ipc:///tmp/rl-colocate-zmq-{suffix}.sock"
-        self.zmq_address_counter += 1
-        return {self.device_uuid: self.zmq_handle}
-
-    def update_weights(self):
-        # align size to avoid misaligned address
-        align_size = 256
-
-        def get_size(p: torch.Tensor) -> int:
-            return (p.nbytes + align_size - 1) // align_size * align_size
-
-        named_parameters: dict[str, torch.nn.Parameter] = dict(
-            self.model.named_parameters()
-        )
-        max_tensor_size = max(get_size(p) for p in named_parameters.values())
-        # use max_tensor_size * 2 as buffer size
-        buffer = torch.empty(max_tensor_size * 2, dtype=torch.uint8, device="cuda:0")
-        s = self.zmq_context.socket(zmq.REQ)
-        s.bind(self.zmq_handle)
-        handle = reduce_tensor(buffer)
-
-        offset = 0
-        buckets: list[tuple[list[dict], list[torch.Tensor]]] = []
-        named_tensors: list[dict] = []
-        real_tensors: list[torch.Tensor] = []
-        for name, p in named_parameters.items():
-            size = get_size(p)
-            if offset + size > buffer.numel():
-                buckets.append((named_tensors, real_tensors))
-                named_tensors, real_tensors = [], []
-                offset = 0
-            # assume tensors are contiguous
-            named_tensors.append(
-                {"name": name, "dtype": p.dtype, "shape": p.shape, "offset": offset}
-            )
-            real_tensors.append(p)
-            offset += size
-        if named_tensors:
-            buckets.append((named_tensors, real_tensors))
-        s.send_pyobj(handle)
-        s.recv()
-        for named_tensors, real_tensors in buckets:
-            offset = 0
-            for p in real_tensors:
-                buffer[offset : offset + p.nbytes].data.copy_(
-                    p.data.view(-1).view(dtype=torch.uint8), non_blocking=True
-                )
-                offset += get_size(p)
-            torch.accelerator.synchronize()
-            s.send_pyobj(named_tensors)
-            s.recv()
-        s.send_pyobj(None)
-        s.recv()
-        s.close()
-        del buffer
-        gc.collect()
-        torch.accelerator.empty_cache()
-
-
-# Ray manages four GPUs.
-
-os.environ["CUDA_VISIBLE_DEVICES"] = "0,1,2,3"
-ray.init()
-
-# Co-locate vLLM instances and training actors on the same set of GPUs:
-#   * GPU 0 and 1: training actor 0, training actor 1, and vLLM instance 0
-#     (tensor parallelism = 2).
-#   * GPU 2 and 3: training actor 2, training actor 3, and vLLM instance 1
-#     (tensor parallelism = 2).
-
-pg = placement_group([{"GPU": 1, "CPU": 0}] * 4)
-ray.get(pg.ready())
-print(f"placement group has bundles {pg.bundle_specs=}")
-
-training_actors = []
-training_actor_device_ids = []
-inference_engines = []
-inference_engine_device_ids = []
-
-for bundle_index in [0, 1, 2, 3]:
-    training_actor = ray.remote(
-        num_cpus=0,
-        num_gpus=0.4,
-        scheduling_strategy=PlacementGroupSchedulingStrategy(
-            placement_group=pg,
-            placement_group_capture_child_tasks=True,
-            placement_group_bundle_index=bundle_index,
-        ),
-    )(RayTrainingActor).remote()
-    training_actors.append(training_actor)
-
-for bundle_index, training_actor in enumerate(training_actors):
-    device_id = ray.get(training_actor.report_device_id.remote())
-    print(f"training actor {bundle_index} is on {device_id}")
-    training_actor_device_ids.append(device_id)
-
-for i, bundle_indices in enumerate([[0, 1], [2, 3]]):
-    # Use the following syntax instead of the @ray.remote decorator so that
-    # the placement group is customized for each bundle.
-    llm = ray.remote(
-        num_cpus=0,
-        num_gpus=0,
-        scheduling_strategy=PlacementGroupSchedulingStrategy(
-            placement_group=pg,
-            placement_group_capture_child_tasks=True,
-        ),
-    )(MyLLM).remote(
-        model="facebook/opt-125m",
-        enforce_eager=True,
-        worker_extension_cls="rlhf_utils.ColocateWorkerExtension",
-        tensor_parallel_size=2,
-        distributed_executor_backend="ray",
-        gpu_memory_utilization=0.4,
-        bundle_indices=bundle_indices,
-    )
-    inference_engines.append(llm)
-    # Do not call any method on the inference engine at this point; the call
-    # blocks until the vLLM instance finishes initialization.
-
-for i, llm in enumerate(inference_engines):
-    inference_engine_device_ids.append(
-        ray.get(llm.collective_rpc.remote("report_device_id", args=tuple()))
-    )
-    print(f"inference engine {i} is on {inference_engine_device_ids[-1]}")
-
-# Verify placement: the first two training actors share the same GPUs as
-# the first inference engine.
-assert training_actor_device_ids[:2] == inference_engine_device_ids[0]
-# Verify placement: the last two training actors share the same GPUs as
-# the second inference engine.
-assert training_actor_device_ids[2:] == inference_engine_device_ids[1]
-
-print("Gather all the ZMQ handles from the training actors.")
-zmq_handles = {}
-for actor in training_actors:
-    zmq_handles.update(ray.get(actor.get_zmq_handles.remote()))
-
-print(f"ZMQ handles: {zmq_handles}")
-
-print("Update the weights of the inference engines.")
-ray.get(
-    [actor.update_weights.remote() for actor in training_actors]
-    + [
-        llm.collective_rpc.remote("update_weights_from_ipc", args=(zmq_handles,))
-        for llm in inference_engines
-    ]
-)
-
-print("Check if the weights are updated.")
-for llm in inference_engines:
-    assert ray.get(llm.collective_rpc.remote("check_weights_changed", args=tuple()))

examples/offline_inference/rlhf_online_quant.py

@@ -1,162 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Demonstrates reinforcement learning from human feedback (RLHF) using vLLM and Ray.
-
-The script separates training and inference workloads onto distinct GPUs
-so that Ray can manage process placement and inter-process communication.
-A Hugging Face Transformer model occupies GPU 0 for training, whereas a
-tensor-parallel vLLM inference engine occupies GPU 1–2.
-
-The example performs the following steps:
-
-* Load the training model on GPU 0.
-* Split the inference model across GPUs 1–2 using vLLM's tensor parallelism
-  and Ray placement groups.
-* Generate text from a list of prompts using the inference engine.
-* Update the weights of the training model and broadcast the updated weights
-  to the inference engine by using a Ray collective RPC group. Note that
-  for demonstration purposes we simply zero out the weights.
-
-For a production-ready implementation that supports multiple training and
-inference replicas, see the OpenRLHF framework:
-https://github.com/OpenRLHF/OpenRLHF
-
-This example assumes a single-node cluster with three GPUs, but Ray
-supports multi-node clusters. vLLM expects the GPUs are only used for vLLM
-workloads. Residual GPU activity interferes with vLLM memory profiling and
-causes unexpected behavior.
-"""
-
-import json
-import os
-
-import ray
-import torch
-from ray.util.placement_group import placement_group
-from ray.util.scheduling_strategies import PlacementGroupSchedulingStrategy
-from rlhf_utils import stateless_init_process_group
-from torchao.core.config import config_to_dict
-from torchao.quantization import (
-    Float8DynamicActivationFloat8WeightConfig,
-    PerRow,
-)
-from transformers import AutoModelForCausalLM
-
-from vllm import LLM, SamplingParams
-from vllm.utils.network_utils import get_ip, get_open_port
-
-
-class MyLLM(LLM):
-    """Configure the vLLM worker for Ray placement group execution."""
-
-    def __init__(self, *args, **kwargs):
-        # Remove the top-level CUDA_VISIBLE_DEVICES variable set by Ray
-        # so that vLLM can manage its own device placement within the worker.
-        os.environ.pop("CUDA_VISIBLE_DEVICES", None)
-        super().__init__(*args, **kwargs)
-
-
-# Load the OPT-125M model onto GPU 0 for the training workload.
-train_model = AutoModelForCausalLM.from_pretrained("facebook/opt-125m")
-train_model.to("cuda:0")
-
-# Initialize Ray and set the visible devices. The vLLM engine will
-# be placed on GPUs 1 and 2.
-os.environ["CUDA_VISIBLE_DEVICES"] = "1,2"
-ray.init()
-
-# Create a placement group that reserves GPU 1–2 for the vLLM inference engine.
-# Learn more about Ray placement groups:
-# https://docs.ray.io/en/latest/ray-core/scheduling/placement-group.html
-pg_inference = placement_group([{"GPU": 1, "CPU": 0}] * 2)
-ray.get(pg_inference.ready())
-scheduling_inference = PlacementGroupSchedulingStrategy(
-    placement_group=pg_inference,
-    placement_group_capture_child_tasks=True,
-    placement_group_bundle_index=0,
-)
-
-# Launch the vLLM inference engine. The `enforce_eager` flag reduces
-# start-up latency.
-
-# generate torchao quantization config for RL rollout
-# see https://github.com/vllm-project/vllm/pull/23014 for instructions to
-# use serialized config files instead of passing around json string
-config = Float8DynamicActivationFloat8WeightConfig(granularity=PerRow())
-
-json_str = json.dumps(config_to_dict(config))
-
-llm = ray.remote(
-    num_cpus=0,
-    num_gpus=0,
-    scheduling_strategy=scheduling_inference,
-)(MyLLM).remote(
-    model="facebook/opt-125m",
-    hf_overrides={"quantization_config_dict_json": json_str},
-    enforce_eager=True,
-    worker_extension_cls="rlhf_utils.WorkerExtension",
-    tensor_parallel_size=2,
-    distributed_executor_backend="ray",
-)
-
-# Generate text from the prompts.
-prompts = [
-    "Hello, my name is",
-    "The president of the United States is",
-    "The capital of France is",
-    "The future of AI is",
-]
-
-sampling_params = SamplingParams(temperature=0)
-
-outputs = ray.get(llm.generate.remote(prompts, sampling_params))
-
-print("-" * 50)
-for output in outputs:
-    prompt = output.prompt
-    generated_text = output.outputs[0].text
-    print(f"Prompt: {prompt!r}\nGenerated text: {generated_text!r}")
-    print("-" * 50)
-
-# Set up the communication channel between the training process and the
-# inference engine.
-master_address = get_ip()
-master_port = get_open_port()
-
-handle = llm.collective_rpc.remote(
-    "init_weight_update_group", args=(master_address, master_port, 1, 3)
-)
-
-model_update_group = stateless_init_process_group(
-    master_address, master_port, 0, 3, torch.device("cuda:0")
-)
-ray.get(handle)
-
-# Simulate a training step by zeroing out all model weights.
-# In a real RLHF training loop the weights would be updated using the gradient
-# from an RL objective such as PPO on a reward model.
-for name, p in train_model.named_parameters():
-    p.data.zero_()
-
-# Synchronize the updated weights to the inference engine.
-for name, p in train_model.named_parameters():
-    dtype_name = str(p.dtype).split(".")[-1]
-    handle = llm.collective_rpc.remote(
-        "update_weight", args=(name, dtype_name, p.shape)
-    )
-    model_update_group.broadcast(p, src=0, stream=torch.cuda.current_stream())
-    ray.get(handle)
-
-# Verify that the inference weights have been updated.
-assert all(ray.get(llm.collective_rpc.remote("check_weights_changed")))
-
-# Generate text with the updated model. The output is expected to be nonsense
-# because the weights are zero.
-outputs_updated = ray.get(llm.generate.remote(prompts, sampling_params))
-print("-" * 50)
-for output in outputs_updated:
-    prompt = output.prompt
-    generated_text = output.outputs[0].text
-    print(f"Prompt: {prompt!r}\nGenerated text: {generated_text!r}")
-    print("-" * 50)

examples/offline_inference/rlhf_utils.py

@@ -1,168 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import gc
-from collections.abc import Callable
-from typing import TypedDict
-
-import torch
-import zmq
-
-
-def stateless_init_process_group(master_address, master_port, rank, world_size, device):
-    """
-    vLLM provides `StatelessProcessGroup` to create a process group
-    without considering the global process group in torch.distributed.
-    It is recommended to create `StatelessProcessGroup`, and then initialize
-    the data-plane communication (NCCL) between external (train processes)
-    and vLLM workers.
-    """
-    from vllm.distributed.device_communicators.pynccl import PyNcclCommunicator
-    from vllm.distributed.utils import StatelessProcessGroup
-
-    pg = StatelessProcessGroup.create(
-        host=master_address, port=master_port, rank=rank, world_size=world_size
-    )
-    pynccl = PyNcclCommunicator(pg, device=device)
-    return pynccl
-
-
-class WorkerExtension:
-    """
-    The class for vLLM's worker to inherit from.
-    By defining an extension class, the code can work no matter what is
-    the underlying worker class.
-
-    NOTE: we define this class in a separate module, and the main module
-    should pass the full qualified name as `worker_extension_cls` argument.
-    """
-
-    def init_weight_update_group(
-        self, master_address, master_port, rank_offset, world_size
-    ):
-        from vllm.distributed.parallel_state import get_world_group
-
-        rank = get_world_group().rank + rank_offset
-        self.model_update_group = stateless_init_process_group(
-            master_address,
-            master_port,
-            rank,
-            world_size,
-            self.device,
-        )
-
-    def update_weight(self, name, dtype_name, shape):
-        dtype = getattr(torch, dtype_name)
-        weight = torch.empty(shape, dtype=dtype, device="cuda")
-        self.model_update_group.broadcast(
-            weight, src=0, stream=torch.cuda.current_stream()
-        )
-
-        self.model_runner.model.load_weights(weights=[(name, weight)])
-
-        del weight
-
-    def check_weights_changed(self):
-        """
-        Check if the weights are updated to 0.
-        """
-        weights_updated = True
-        for name, p in self.model_runner.model.named_parameters():
-            weights_updated = weights_updated and torch.allclose(p, torch.zeros_like(p))
-        return weights_updated
-
-
-def rebuild_ipc(
-    handle: tuple[Callable, tuple], device_id: int | None = None
-) -> torch.Tensor:
-    func, args = handle
-    list_args = list(args)
-    if device_id is not None:
-        # the key is to change device id to the current device id
-        # in case two processes have different CUDA_VISIBLE_DEVICES
-        list_args[6] = device_id
-    buffer = func(*list_args)
-    return buffer
-
-
-class FlattenedTensorMetadata(TypedDict):
-    name: str
-    shape: torch.Size
-    dtype: torch.dtype
-    # specify the start offset of this tensor in shared ipc_buffer tensor
-    offset: int
-
-
-class ColocateWorkerExtension:
-    """
-    The class for vLLM's worker to inherit from, in the colocate setting.
-    By defining an extension class, the code can work no matter what is
-    the underlying worker class.
-
-    NOTE: we define this class in a separate module, and the main module
-    should pass the full qualified name as `worker_extension_cls` argument.
-    """
-
-    def update_weights_from_ipc(self, zmq_handles: dict[str, str]):
-        from vllm.model_executor.model_loader.utils import process_weights_after_loading
-
-        assert self.device is not None
-        if not hasattr(self, "_zmq_ctx") or self._zmq_ctx is None:
-            self._zmq_ctx = zmq.Context()
-        socket = self._zmq_ctx.socket(zmq.REP)
-        socket.connect(zmq_handles[self.report_device_id()])
-        buffer: torch.Tensor | None = None
-        while True:
-            payload: tuple[Callable, tuple] | list[FlattenedTensorMetadata] | None = (
-                socket.recv_pyobj()
-            )
-            if payload is None:
-                # means the update is done
-                process_weights_after_loading(
-                    self.model_runner.model, self.model_config, self.device
-                )
-                torch.accelerator.synchronize()
-                socket.send(b"")
-                break
-            if isinstance(payload, tuple):
-                # an ipc handle that vLLM can use `func, args = handle`
-                # and `func(*args)` to rebuild GPU tensor.
-                buffer = rebuild_ipc(payload, self.device.index)
-                assert buffer.dtype == torch.uint8
-                socket.send(b"")
-                continue
-            assert isinstance(payload, list)
-            assert buffer is not None
-            weights = []
-            for item in payload:
-                shape = item["shape"]
-                if isinstance(shape, (list, tuple)):
-                    shape = torch.Size(shape)
-                assert isinstance(shape, torch.Size)
-                dtype, offset = item["dtype"], item["offset"]
-                size = dtype.itemsize * shape.numel()
-                tensor = buffer[offset : offset + size].view(dtype=dtype).view(shape)
-                weights.append((item["name"], tensor))
-            self.model_runner.model.load_weights(weights=weights)
-            del weights
-            torch.accelerator.synchronize()
-            socket.send(b"")
-
-        socket.close()
-        del buffer
-        gc.collect()
-        torch.accelerator.empty_cache()
-
-    def report_device_id(self) -> str:
-        from vllm.platforms import current_platform
-
-        self.device_uuid = current_platform.get_device_uuid(self.device.index)
-        return self.device_uuid
-
-    def check_weights_changed(self):
-        """
-        Check if the weights are updated to 0.
-        """
-        weights_updated = True
-        for name, p in self.model_runner.model.named_parameters():
-            weights_updated = weights_updated and torch.allclose(p, torch.zeros_like(p))
-        return weights_updated

examples/pooling/score/colqwen3_5_rerank_online.py

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

examples/rl/rlhf_nccl_fsdp_ep.py

@@ -0,0 +1,339 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+RLHF with FSDP2 training (4 GPUs) and vLLM expert-parallel inference (4 GPUs).
+
+8-GPU layout:
+  Training  — 4 GPUs, PyTorch FSDP2 (fully_shard)
+  Inference — 4 GPUs, vLLM AsyncLLMEngine with expert parallelism +
+              data parallelism (TP=1, DP=4, enable_expert_parallel
+              → EP_SIZE = TP×DP = 4)
+
+FSDP workers are Ray actors that form a single FSDP2 process group.
+Rank 0 gathers full parameters via DTensor.full_tensor() and broadcasts
+them to the vLLM inference engine through the NCCL weight-transfer API.
+
+The inference engine uses AsyncLLMEngine which automatically spawns
+DP worker processes (no manual placement group needed).  Weight sync
+uses pause_generation / resume_generation.
+
+Steps:
+  1. Launch 4 FSDP training workers.
+  2. Launch AsyncLLMEngine with EP+DP (dummy weights).
+  3. Generate from prompts → gibberish (random weights).
+  4. Pause generation, transfer weights from FSDP, resume.
+  5. Generate from prompts → sensible output (synced weights).
+
+Assumes a single-node cluster with 8 GPUs.
+"""
+
+import asyncio
+import os
+import uuid
+from dataclasses import asdict
+
+import ray
+import torch
+import torch.distributed as dist
+from huggingface_hub import snapshot_download
+from torch.distributed.fsdp import fully_shard
+from transformers import AutoModelForCausalLM
+
+import vllm
+from vllm import SamplingParams
+from vllm.config import WeightTransferConfig
+from vllm.distributed.weight_transfer.base import (
+    WeightTransferInitRequest,
+    WeightTransferUpdateRequest,
+)
+from vllm.distributed.weight_transfer.nccl_engine import (
+    NCCLTrainerSendWeightsArgs,
+    NCCLWeightTransferEngine,
+    NCCLWeightTransferInitInfo,
+    NCCLWeightTransferUpdateInfo,
+)
+from vllm.utils.network_utils import get_ip, get_open_port
+from vllm.v1.executor import Executor
+
+MODEL_NAME = "Qwen/Qwen3-30B-A3B"
+
+FSDP_WORLD_SIZE = 4
+INFERENCE_TP_SIZE = 1
+INFERENCE_DP_SIZE = 4
+
+
+@ray.remote(num_gpus=1)
+class FSDPTrainWorker:
+    """
+    One FSDP2 training worker per GPU.  Four of these form the FSDP group.
+    Rank 0 additionally handles weight transfer to the vLLM engine.
+    """
+
+    def __init__(
+        self,
+        model_name: str,
+        rank: int,
+        fsdp_world_size: int,
+        fsdp_master_addr: str,
+        fsdp_master_port: int,
+    ):
+        self.rank = rank
+
+        os.environ["MASTER_ADDR"] = fsdp_master_addr
+        os.environ["MASTER_PORT"] = str(fsdp_master_port)
+
+        dist.init_process_group(backend="nccl", rank=rank, world_size=fsdp_world_size)
+        torch.accelerator.set_device_index(0)
+
+        model = AutoModelForCausalLM.from_pretrained(
+            model_name, torch_dtype=torch.bfloat16
+        )
+
+        self.weight_names = [n for n, _ in model.named_parameters()]
+        self.weight_dtype_names = [
+            str(p.dtype).split(".")[-1] for _, p in model.named_parameters()
+        ]
+        self.weight_shapes = [list(p.shape) for _, p in model.named_parameters()]
+
+        for layer in model.model.layers:
+            fully_shard(layer)
+        fully_shard(model)
+
+        self.model = model
+
+        self.transfer_port = None
+        self.transfer_master_address = None
+        self.model_update_group = None
+
+    def get_rank(self):
+        return self.rank
+
+    # ---- weight-transfer setup (rank 0 only) ----
+
+    def setup_transfer_endpoint(self):
+        """Create the NCCL rendezvous endpoint for weight transfer."""
+        assert self.rank == 0
+        self.transfer_port = get_open_port()
+        self.transfer_master_address = get_ip()
+        return self.transfer_master_address, self.transfer_port
+
+    def init_weight_transfer_group(self, transfer_world_size: int):
+        """Join the weight-transfer NCCL group as rank 0 (the source)."""
+        assert self.rank == 0
+        self.model_update_group = NCCLWeightTransferEngine.trainer_init(
+            dict(
+                master_address=self.transfer_master_address,
+                master_port=self.transfer_port,
+                world_size=transfer_world_size,
+            ),
+        )
+
+    def get_weight_metadata(self):
+        """Return weight names, dtypes, and shapes captured before FSDP wrapping."""
+        return self.weight_names, self.weight_dtype_names, self.weight_shapes
+
+    # ---- collective ops (ALL FSDP ranks must call concurrently) ----
+
+    def gather_and_broadcast_weights(self, packed: bool = True):
+        """
+        All-gather full parameters and broadcast them to vLLM.
+        Only rank 0 performs the actual NCCL broadcast; others just
+        participate in the FSDP all-gather.
+
+        full_tensor() is a collective — all FSDP ranks must call it
+        for each parameter in the same order.  Rank 0 additionally
+        feeds each gathered tensor to the weight-transfer engine.
+        """
+        if self.rank == 0:
+
+            def _full_param_iter():
+                for name, param in self.model.named_parameters():
+                    yield name, param.full_tensor()
+
+            trainer_args = NCCLTrainerSendWeightsArgs(
+                group=self.model_update_group,
+                packed=packed,
+            )
+            NCCLWeightTransferEngine.trainer_send_weights(
+                iterator=_full_param_iter(),
+                trainer_args=trainer_args,
+            )
+        else:
+            for _, param in self.model.named_parameters():
+                param.full_tensor()
+
+
+def create_async_engine(**kwargs):
+    """Create an AsyncLLMEngine directly (no subclass needed)."""
+    engine_args = vllm.AsyncEngineArgs(**kwargs)
+    vllm_config = engine_args.create_engine_config()
+    executor_class = Executor.get_class(vllm_config)
+    return vllm.AsyncLLMEngine(
+        vllm_config=vllm_config,
+        executor_class=executor_class,
+        log_requests=engine_args.enable_log_requests,
+        log_stats=not engine_args.disable_log_stats,
+    )
+
+
+async def generate_batch(engine, prompts, sampling_params):
+    """Generate completions for a batch of prompts."""
+
+    async def gen_one(prompt):
+        output = None
+        async for request_output in engine.generate(
+            {"prompt": prompt},
+            sampling_params,
+            request_id=str(uuid.uuid4()),
+        ):
+            output = request_output
+        return output
+
+    return await asyncio.gather(*[gen_one(p) for p in prompts])
+
+
+async def main():
+    ray.init()
+
+    # Download model weights to local/shared disk once.
+    local_model_path = snapshot_download(MODEL_NAME)
+    print(f"[init] Model downloaded to {local_model_path}")
+
+    # FSDP rendezvous address (single-node)
+    fsdp_master_addr = get_ip()
+    fsdp_master_port = get_open_port()
+
+    # Launch 4 FSDP training workers.
+    # Ray allocates 1 GPU per worker; AsyncLLMEngine's internal DP
+    # placement groups will land on the remaining 4 GPUs.
+    fsdp_workers = [
+        FSDPTrainWorker.remote(
+            local_model_path,
+            rank,
+            FSDP_WORLD_SIZE,
+            fsdp_master_addr,
+            fsdp_master_port,
+        )
+        for rank in range(FSDP_WORLD_SIZE)
+    ]
+    ray.get([w.get_rank.remote() for w in fsdp_workers])
+    print(f"[init] {FSDP_WORLD_SIZE} FSDP training workers ready.")
+
+    # Launch vLLM with expert parallelism + data parallelism.
+    # AsyncLLMEngine with data_parallel_backend="ray" creates its own
+    # placement groups internally — no manual placement group needed.
+    print("[engine] Creating AsyncLLMEngine...")
+    engine = create_async_engine(
+        model=local_model_path,
+        enforce_eager=True,
+        tensor_parallel_size=INFERENCE_TP_SIZE,
+        data_parallel_size=INFERENCE_DP_SIZE,
+        enable_expert_parallel=True,
+        distributed_executor_backend="ray",
+        data_parallel_backend="ray",
+        weight_transfer_config=WeightTransferConfig(backend="nccl"),
+        load_format="dummy",
+        gpu_memory_utilization=0.7,
+    )
+    print("[engine] AsyncLLMEngine created.")
+
+    prompts = [
+        "Hello, my name is",
+        "The president of the United States is",
+        "The capital of France is",
+        "The future of AI is",
+    ]
+    sampling_params = SamplingParams(temperature=0)
+
+    # Generate with dummy weights — expect gibberish.
+    print("[generate] Starting generation with dummy weights...")
+    outputs = await generate_batch(engine, prompts, sampling_params)
+    print("[generate] Generation complete.")
+
+    print("-" * 60)
+    print("BEFORE weight sync (dummy weights):")
+    print("-" * 60)
+    for output in outputs:
+        print(f"Prompt: {output.prompt!r}")
+        print(f"Generated: {output.outputs[0].text!r}")
+        print("-" * 60)
+
+    # --- Weight-transfer setup ---
+    print("[transfer] Setting up weight-transfer endpoint...")
+    transfer_addr, transfer_port = ray.get(
+        fsdp_workers[0].setup_transfer_endpoint.remote()
+    )
+    print(f"[transfer] Endpoint ready at {transfer_addr}:{transfer_port}")
+
+    transfer_world_size = INFERENCE_TP_SIZE * INFERENCE_DP_SIZE + 1
+    print(
+        f"[transfer] World size: {transfer_world_size} "
+        f"(1 trainer + {INFERENCE_TP_SIZE * INFERENCE_DP_SIZE} vLLM workers)"
+    )
+
+    print("[transfer] Initializing NCCL groups...")
+    train_handle = fsdp_workers[0].init_weight_transfer_group.remote(
+        transfer_world_size
+    )
+    await engine.init_weight_transfer_engine(
+        WeightTransferInitRequest(
+            init_info=asdict(
+                NCCLWeightTransferInitInfo(
+                    master_address=transfer_addr,
+                    master_port=transfer_port,
+                    rank_offset=1,
+                    world_size=transfer_world_size,
+                )
+            )
+        )
+    )
+    ray.get(train_handle)
+    print("[transfer] NCCL groups initialized.")
+
+    # --- Pause, transfer weights, resume ---
+    print("[sync] Pausing generation...")
+    await engine.pause_generation(mode="abort")
+    print("[sync] Generation paused.")
+
+    names, dtype_names, shapes = ray.get(fsdp_workers[0].get_weight_metadata.remote())
+    print(f"[sync] Got metadata for {len(names)} parameters.")
+
+    print("[sync] Broadcasting weights from FSDP → vLLM...")
+    broadcast_handles = [
+        w.gather_and_broadcast_weights.remote(packed=True) for w in fsdp_workers
+    ]
+    await engine.update_weights(
+        WeightTransferUpdateRequest(
+            update_info=asdict(
+                NCCLWeightTransferUpdateInfo(
+                    names=names,
+                    dtype_names=dtype_names,
+                    shapes=shapes,
+                    packed=True,
+                )
+            )
+        )
+    )
+    ray.get(broadcast_handles)
+    print("[sync] Weight broadcast complete.")
+
+    print("[sync] Resuming generation...")
+    await engine.resume_generation()
+    print("[sync] Generation resumed.")
+
+    # Generate with synced weights — expect sensible output.
+    print("[generate] Starting generation with synced weights...")
+    outputs_updated = await generate_batch(engine, prompts, sampling_params)
+    print("[generate] Generation complete.")
+
+    print("-" * 60)
+    print("AFTER weight sync (real weights):")
+    print("-" * 60)
+    for output in outputs_updated:
+        print(f"Prompt: {output.prompt!r}")
+        print(f"Generated: {output.outputs[0].text!r}")
+        print("-" * 60)
+
+
+if __name__ == "__main__":
+    asyncio.run(main())

requirements/common.txt

@@ -12,7 +12,7 @@ tokenizers >= 0.21.1  # Required for fast incremental detokenization.
 protobuf >= 5.29.6, !=6.30.*, !=6.31.*, !=6.32.*, !=6.33.0.*, !=6.33.1.*, !=6.33.2.*, !=6.33.3.*, !=6.33.4.* # Required by LlamaTokenizer, gRPC. CVE-2026-0994
 fastapi[standard] >= 0.115.0 # Required by FastAPI's form models in the OpenAI API server's audio transcriptions endpoint.
 aiohttp >= 3.13.3
-openai >= 1.99.1, < 2.25.0  # For Responses API with reasoning content
+openai >= 2.0.0  # For Responses API with reasoning content
 pydantic >= 2.12.0
 prometheus_client >= 0.18.0
 pillow  # Required for image processing

requirements/xpu-test.in

@@ -0,0 +1,35 @@
+# --- Test Infrastructure ---
+tblib
+pytest-timeout
+pytest-cov
+pytest-forked
+pytest-rerunfailures
+pytest-shard
+
+# --- Core Tools & Bindings ---
+absl-py
+arctic-inference
+
+# --- Audio Processing ---
+librosa
+audioread
+soxr
+pooch
+soundfile
+
+# --- Tool Parsing & Evaluation ---
+blobfile
+rapidfuzz
+gpt-oss
+schemathesis
+jiwer
+bm25s
+pystemmer
+mteb[bm25s]
+num2words
+pqdm
+
+# --- Vision & Multimodal ---
+timm
+albumentations
+mistral-common[image,audio]

requirements/xpu-test.txt

@@ -0,0 +1,42 @@
+# XPU Test Dependencies
+# NOTE: Base image already has common.txt + xpu.txt installed,
+#       and vllm-openai stage has pytest, pytest-asyncio, lm-eval[api].
+#       This file only adds incremental test-specific packages.
+
+# Additional test infrastructure (pytest/pytest-asyncio already in base)
+# This file was autogenerated by uv via the following command:
+#    uv pip compile /workspace/vllm/requirements/xpu-test.in -o /workspace/vllm/requirements/xpu-test.txt -c /workspace/vllm/requirements/xpu.txt --index-strategy unsafe-best-match --extra-index-url ${PIP_EXTRA_INDEX_URL} --python-version ${PYTHON_VERSION} 
+tblib==3.1.0
+pytest-timeout==2.3.1
+pytest-cov==6.3.0
+pytest-forked==1.6.0
+pytest-rerunfailures==14.0
+pytest-shard==0.1.2
+
+arctic-inference==0.1.1
+
+# Required for audio processing tests
+librosa==0.10.2.post1
+audioread==3.0.1
+soxr==0.5.0.post1
+pooch==1.8.2
+soundfile==0.13.1
+
+# Required for Mistral's streaming tool parser
+blobfile==3.0.0
+rapidfuzz==3.12.1
+
+# Required for Mistral's streaming tool parser and some evaluation scripts
+gpt-oss==0.0.8
+schemathesis==3.39.15
+jiwer==4.0.0
+bm25s==0.2.13
+pystemmer==3.0.0
+mteb[bm25s]>=2, <3
+num2words==0.5.14
+pqdm==0.2.0
+
+# Required for some evaluation scripts
+timm==1.0.17
+albumentations==1.4.6
+mistral-common[image,audio]==1.9.1

tests/compile/fusions_e2e/conftest.py

@@ -82,6 +82,10 @@ def run_e2e_fusion_test(monkeypatch, caplog_mp_spawn):
                 f"attention backend '{attn_backend.backend.name}'"
             )
 
+        # TODO: remove this after finishing migration from envs to model kwargs
+        if model_name == "openai/gpt-oss-20b":
+            monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_MXFP4_MXFP8", "1")
+
         # Disable, compile cache to make sure custom passes run.
         # Otherwise, we can't verify fusion happened through the logs.
         monkeypatch.setenv("VLLM_DISABLE_COMPILE_CACHE", "1")

tests/compile/fusions_e2e/models.py

@@ -162,3 +162,12 @@ deepseek_v3_fp8 = ModelFusionInfo(
         # async_tp=n_layers * 2,
     ),
 )
+
+gpt_oss_20b = ModelFusionInfo(
+    model_name="openai/gpt-oss-20b",
+    matches=lambda n_layers: Matches(
+        ar_rms_fusion=n_layers * 2 + 1,
+        sequence_parallel=n_layers * 2 + 1,
+        async_tp=n_layers * 2,
+    ),
+)

tests/compile/fusions_e2e/test_tp2_ar_rms.py

@@ -20,6 +20,7 @@ from .models import (
     FLASHINFER_MLA_ATTN,
     TRITON_ATTN,
     deepseek_v3_fp8,
+    gpt_oss_20b,
     llama3_8b,
     llama3_8b_fp4,
     llama3_8b_fp8,
@@ -158,7 +159,7 @@ def test_tp2_ar_rms_fp4_fusions(
 @multi_gpu_test(num_gpus=2)
 @pytest.mark.parametrize(
     "model_name, matches_fn, model_kwargs, hf_overrides",
-    [llama3_8b, qwen3_a3b],
+    [llama3_8b, qwen3_a3b, gpt_oss_20b],
 )
 @pytest.mark.parametrize("attn_backend", [TRITON_ATTN])
 @pytest.mark.parametrize("n_layers", [4])

tests/compile/passes/test_rope_kvcache_fusion.py

@@ -295,7 +295,7 @@ def test_rope_kvcache_fusion(
             }
             q_unfused, k_unfused, v_unfused, dummy = model(qkv_unfused, pos_unfused)
             attn_layer = forward_context.no_compile_layers[model.layer_name]
-            kv_cache_unfused = attn_layer.kv_cache[forward_context.virtual_engine]
+            kv_cache_unfused = attn_layer.kv_cache[0]
         del dummy
 
         torch._dynamo.mark_dynamic(qkv, 0)
@@ -309,7 +309,7 @@ def test_rope_kvcache_fusion(
             }
             q_fused, k_fused, v_fused, dummy = model_fused(qkv, pos)
             attn_layer = forward_context.no_compile_layers[model.layer_name]
-            kv_cache_fused = attn_layer.kv_cache[forward_context.virtual_engine]
+            kv_cache_fused = attn_layer.kv_cache[0]
         del dummy
 
         assert fusion_pass.matched_count == 1

tests/distributed/test_eplb_algo.py

@@ -5,6 +5,7 @@ import numpy as np
 import pytest
 import torch
 
+from vllm.distributed.eplb.eplb_state import compute_logical_maps
 from vllm.distributed.eplb.policy.default import DefaultEplbPolicy
 
 
@@ -24,9 +25,10 @@ def test_basic_rebalance():
     num_nodes = 2
     num_gpus = 8
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    log2phy, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Verify output shapes
     assert phy2log.shape == (
@@ -78,9 +80,10 @@ def test_single_gpu_case():
     num_nodes = 1
     num_gpus = 1
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    log2phy, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Verify shapes
     assert phy2log.shape == (1, 4)
@@ -100,9 +103,10 @@ def test_equal_weights():
     num_nodes = 2
     num_gpus = 4
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    _, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Verify shapes
     assert phy2log.shape == (1, 8)
@@ -123,9 +127,10 @@ def test_extreme_weight_imbalance():
     num_nodes = 2
     num_gpus = 4
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    _, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Verify shapes
     assert phy2log.shape == (1, 12)
@@ -151,9 +156,10 @@ def test_multiple_layers():
     num_nodes = 2
     num_gpus = 4
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    _, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Verify shapes
     assert phy2log.shape == (3, 8)
@@ -176,7 +182,8 @@ def test_parameter_validation():
     # Test non-divisible case - this should handle normally without throwing
     # errors because the function will fall back to global load balancing
     # strategy
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(weight, 8, 3, 2, 4)
+    phy2log = DefaultEplbPolicy.rebalance_experts(weight, 8, 3, 2, 4)
+    _, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
     assert phy2log.shape == (1, 8)
     assert logcnt.shape == (1, 4)
 
@@ -198,9 +205,10 @@ def test_small_scale_hierarchical():
     num_nodes = 2  # 2 nodes
     num_gpus = 4  # 4 GPUs
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    _, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Verify basic constraints
     assert phy2log.shape == (1, 12)
@@ -225,9 +233,10 @@ def test_global_load_balance_fallback():
     num_nodes = 2
     num_gpus = 4
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    _, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Should work normally, just using global load balancing strategy
     assert phy2log.shape == (1, 8)
@@ -247,9 +256,10 @@ def test_device_compatibility(device):
     num_nodes = 1
     num_gpus = 2
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
+    _, logcnt = compute_logical_maps(phy2log, weight.shape[-1])
 
     # Function will convert to CPU internally, but should handle different
     # device inputs normally
@@ -264,9 +274,8 @@ def test_additional_cases():
     weight1 = torch.tensor(
         [[50, 100, 75, 120, 90, 60, 80, 110, 40, 70, 95, 85, 65, 55, 45, 35]]
     )
-    phy2log1, log2phy1, logcnt1 = DefaultEplbPolicy.rebalance_experts(
-        weight1, 24, 8, 4, 8
-    )
+    phy2log1 = DefaultEplbPolicy.rebalance_experts(weight1, 24, 8, 4, 8)
+    _, logcnt1 = compute_logical_maps(phy2log1, weight1.shape[-1])
 
     assert phy2log1.shape == (1, 24)
     assert logcnt1.shape == (1, 16)
@@ -279,9 +288,8 @@ def test_additional_cases():
             [12, 25, 50, 100, 150, 200],  # Increasing weights
         ]
     )
-    phy2log2, log2phy2, logcnt2 = DefaultEplbPolicy.rebalance_experts(
-        weight2, 10, 3, 1, 2
-    )
+    phy2log2 = DefaultEplbPolicy.rebalance_experts(weight2, 10, 3, 1, 2)
+    _, logcnt2 = compute_logical_maps(phy2log2, weight2.shape[-1])
 
     assert phy2log2.shape == (2, 10)
     assert logcnt2.shape == (2, 6)
@@ -292,6 +300,42 @@ def test_additional_cases():
         assert logcnt2[layer, max_weight_idx] >= 2
 
 
+def test_compute_logical_maps_with_negative_indices():
+    """
+    Test that compute_logical_maps correctly handles physical slots containing
+    -1 (unused slots).
+    """
+    # 2 layers, 6 physical slots, 4 logical experts.
+    # Slots 2 and 5 are unused (-1).
+    phy2log = torch.tensor(
+        [
+            [0, 1, -1, 2, 3, -1],
+            [3, -1, 2, 1, 0, -1],
+        ]
+    )
+    num_layers = 2
+    num_logical_experts = 4
+
+    log2phy, logcnt = compute_logical_maps(phy2log, num_logical_experts)
+
+    assert logcnt.shape == (num_layers, num_logical_experts)
+    assert log2phy.shape == (num_layers, num_logical_experts, 1)
+
+    expected_logcnt = torch.ones(num_layers, num_logical_experts, dtype=phy2log.dtype)
+    assert torch.all(logcnt == expected_logcnt), (
+        f"Expected that all replica counts == 1, got {logcnt}"
+    )
+
+    assert torch.all(log2phy >= 0), (
+        "log2phy should only contain valid physical indices, not -1"
+    )
+
+    assert log2phy[0, 0, 0] == 0
+    assert log2phy[0, 1, 0] == 1
+    assert log2phy[0, 2, 0] == 3
+    assert log2phy[0, 3, 0] == 4
+
+
 if __name__ == "__main__":
     weight = torch.tensor(
         [
@@ -305,7 +349,7 @@ if __name__ == "__main__":
     num_nodes = 2
     num_gpus = 8
 
-    phy2log, log2phy, logcnt = DefaultEplbPolicy.rebalance_experts(
+    phy2log = DefaultEplbPolicy.rebalance_experts(
         weight, num_replicas, num_groups, num_nodes, num_gpus
     )
     print(phy2log)
@@ -434,9 +478,10 @@ def test_preserve_intragpu_slots(
     """Experts that stay on a GPU keep their old slots; incoming not lost."""
     phy_replicas_idx = _make_phy_replicas_idx_from_phy2log(new_phy2log)
 
-    post_phy2log, post_phy_replicas_idx = DefaultEplbPolicy.preserve_intragpu_slots(
-        new_phy2log, phy_replicas_idx, num_ranks, old_phy2log
+    post_phy2log = DefaultEplbPolicy.preserve_intragpu_slots(
+        new_phy2log, num_ranks, old_phy2log
     )
+    post_phy_replicas_idx = _make_phy_replicas_idx_from_phy2log(post_phy2log)
 
     # Shapes preserved
     assert post_phy2log.shape == new_phy2log.shape

tests/entrypoints/openai/chat_completion/test_audio_in_video.py

@@ -9,7 +9,7 @@ import pytest
 import pytest_asyncio
 
 from tests.conftest import VideoTestAssets
-from tests.utils import RemoteOpenAIServer
+from tests.utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer
 
 MODEL_NAME = "Qwen/Qwen2.5-Omni-3B"
 
@@ -22,6 +22,7 @@ def server():
         "--enforce-eager",
         "--limit-mm-per-prompt",
         json.dumps({"audio": 3, "video": 3}),
+        *ROCM_EXTRA_ARGS,
     ]
 
     with RemoteOpenAIServer(

tests/entrypoints/openai/correctness/test_transcription_api_correctness.py

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

tests/entrypoints/openai/responses/conftest.py

@@ -370,7 +370,7 @@ def log_response_diagnostics(
 def default_server_args():
     return [
         "--max-model-len",
-        "8192",
+        "18192",
         "--enforce-eager",  # For faster startup.
         "--enable-auto-tool-choice",
         "--structured-outputs-config.backend",

tests/entrypoints/openai/responses/test_function_call.py

@@ -118,7 +118,6 @@ async def test_function_tool_use(
         tool_choice=tool_choice,
         temperature=0.0,
     )
-
     assert len(response.output) >= 1
     tool_call = None
     reasoning = None
@@ -127,11 +126,15 @@ async def test_function_tool_use(
             tool_call = out
         if out.type == "reasoning":
             reasoning = out
-    assert tool_call is not None
-    assert tool_call.type == "function_call"
-    assert json.loads(tool_call.arguments) is not None
-    assert reasoning is not None
-    assert reasoning.type == "reasoning"
+    if response.incomplete_details is None:
+        assert tool_call is not None
+        assert tool_call.type == "function_call"
+        assert json.loads(tool_call.arguments) is not None
+        assert reasoning is not None
+        assert reasoning.type == "reasoning"
+    else:
+        print(response.model_dump_json(indent=2))
+        assert response.incomplete_details.reason == "max_output_tokens"
 
 
 @pytest.mark.asyncio

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

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

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

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

tests/kernels/attention/test_trtllm_kvfp8_dequant.py

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

tests/kernels/helion/helpers.py

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

tests/kernels/helion/test_autotune.py

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

tests/kernels/helion/test_pattern_matching.py

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

tests/kernels/helion/test_register.py

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

tests/kernels/moe/test_router_gemm.py

@@ -0,0 +1,37 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for optimized router GEMM kernel
+
+Run `pytest tests/kernels/moe/test_router_gemm.py`.
+"""
+
+import pytest
+import torch
+
+import vllm._custom_ops as ops
+from vllm.platforms import current_platform
+from vllm.utils.torch_utils import set_random_seed
+
+
+@pytest.mark.skipif(
+    not (
+        current_platform.is_cuda()
+        and (
+            current_platform.is_device_capability(90)
+            or current_platform.is_device_capability_family(100)
+        )
+    ),
+    reason="This test only runs on Hopper or Blackwell GPUs.",
+)
+@pytest.mark.parametrize("batch_size", [1, 2, 4, 8])
+@pytest.mark.parametrize("input_dim", [360, 720, 1440, 2880])
+@pytest.mark.parametrize("output_dim", [32, 64, 128])
+def test_gpt_oss_router_gemm(batch_size, input_dim, output_dim):
+    set_random_seed(0)
+    x = torch.randn(batch_size, input_dim, device="cuda", dtype=torch.bfloat16)
+    weight = torch.randn(output_dim, input_dim, device="cuda", dtype=torch.bfloat16)
+    bias = torch.randn(output_dim, device="cuda", dtype=torch.bfloat16)
+
+    output = ops.gpt_oss_router_gemm(x, weight, bias)
+    output_ref = torch.nn.functional.linear(x, weight, bias)
+    torch.testing.assert_close(output, output_ref, atol=1e-2, rtol=1e-2)

tests/models/multimodal/generation/test_common.py

@@ -777,6 +777,7 @@ VLM_TEST_SETTINGS = {
         max_model_len=8192,
         max_num_seqs=2,
         auto_cls=AutoModelForCausalLM,
+        patch_hf_runner=model_utils.paddleocr_vl_patch_hf_runner,
         image_size_factors=[(0.25,)],
         marks=[
             pytest.mark.skipif(

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

@@ -489,13 +489,14 @@ def h2ovl_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
             self.image_size = self.vision_config.image_size
 
         def __call__(self, text: str, images: Image | list[Image], **kwargs):
-            from vllm.model_executor.models.h2ovl import (
-                IMG_CONTEXT,
-                IMG_END,
-                IMG_START,
+            from vllm.transformers_utils.processors.h2ovl import (
                 image_to_pixel_values_h2ovl,
             )
 
+            IMG_START = "<img>"
+            IMG_END = "</img>"
+            IMG_CONTEXT = "<IMG_CONTEXT>"
+
             images = [images] if isinstance(images, Image) else images
             pixel_values = [
                 image_to_pixel_values_h2ovl(
@@ -751,16 +752,17 @@ def skyworkr1v_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
             self.image_size = self.vision_config.image_size
 
         def __call__(self, text: str, images: Image | list[Image], **kwargs):
-            from vllm.model_executor.models.skyworkr1v import (
-                IMG_CONTEXT,
-                IMG_END,
-                IMG_START,
-                image_to_pixel_values_skyworkr1v,
+            from vllm.transformers_utils.processors.internvl import (
+                image_to_pixel_values_internvl,
             )
 
+            IMG_START = "<img>"
+            IMG_END = "</img>"
+            IMG_CONTEXT = "<IMG_CONTEXT>"
+
             images = [images] if isinstance(images, Image) else images
             pixel_values = [
-                image_to_pixel_values_skyworkr1v(
+                image_to_pixel_values_internvl(
                     image,
                     input_size=self.image_size,
                     min_num=self.min_num,
@@ -815,14 +817,15 @@ def internvl_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
             videos: npt.NDArray | list[npt.NDArray] = None,
             **kwargs,
         ):
-            from vllm.model_executor.models.internvl import (
-                IMG_CONTEXT,
-                IMG_END,
-                IMG_START,
+            from vllm.transformers_utils.processors.internvl import (
                 image_to_pixel_values_internvl,
                 video_to_pixel_values_internvl,
             )
 
+            IMG_START = "<img>"
+            IMG_END = "</img>"
+            IMG_CONTEXT = "<IMG_CONTEXT>"
+
             images = [images] if isinstance(images, Image) else images
             videos = [videos] if isinstance(videos, np.ndarray) else videos
             if images is not None:
@@ -1149,6 +1152,31 @@ def ovis2_5_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
     return hf_model
 
 
+def paddleocr_vl_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
+    """Patches the HfRunner to fix create_causal_mask API mismatch.
+
+    The PaddleOCR-VL HF model passes `inputs_embeds` to create_causal_mask,
+    but transformers renamed this parameter to `input_embeds`.
+    """
+    import sys
+
+    model_module = sys.modules.get(type(hf_model.model.model).__module__)
+    if model_module is None:
+        return hf_model
+
+    original_create_causal_mask = getattr(model_module, "create_causal_mask", None)
+    if original_create_causal_mask is None:
+        return hf_model
+
+    def patched_create_causal_mask(*args, **kwargs):
+        if "inputs_embeds" in kwargs:
+            kwargs["input_embeds"] = kwargs.pop("inputs_embeds")
+        return original_create_causal_mask(*args, **kwargs)
+
+    model_module.create_causal_mask = patched_create_causal_mask  # type: ignore[attr-defined]
+    return hf_model
+
+
 def qwen2_5_omni_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
     """Patches and returns an instance of the HfRunner for Qwen2.5-Omni."""
     thinker = hf_model.model.thinker

tests/models/multimodal/pooling/test_colqwen3_5.py

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

tests/models/registry.py

@@ -639,6 +639,11 @@ _LATE_INTERACTION_EXAMPLE_MODELS = {
     "OpsColQwen3Model": _HfExamplesInfo(
         "OpenSearch-AI/Ops-Colqwen3-4B", trust_remote_code=True
     ),
+    "ColQwen3_5": _HfExamplesInfo(
+        "athrael-soju/colqwen3.5-4.5B-v3",
+        trust_remote_code=True,
+        max_model_len=4096,
+    ),
     "Qwen3VLNemotronEmbedModel": _HfExamplesInfo(
         "nvidia/nemotron-colembed-vl-4b-v2",
     ),
@@ -774,7 +779,8 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         "rednote-hilab/dots.ocr", trust_remote_code=True
     ),
     "Eagle2_5_VLForConditionalGeneration": _HfExamplesInfo(
-        "nvidia/Eagle2.5-8B", trust_remote_code=True, is_available_online=False
+        "nvidia/Eagle2.5-8B",
+        trust_remote_code=True,
     ),
     "Emu3ForConditionalGeneration": _HfExamplesInfo("BAAI/Emu3-Chat-hf"),
     "Ernie4_5_VLMoeForConditionalGeneration": _HfExamplesInfo(
@@ -1116,6 +1122,11 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         tokenizer_mode="mistral",
     ),
     # [Encoder-decoder]
+    "CohereASRForConditionalGeneration": _HfExamplesInfo(
+        "/host/engines/vllm/audio/2b-release",
+        trust_remote_code=True,
+        is_available_online=False,  # TODO (ekagra): revert after asr release
+    ),
     "NemotronParseForConditionalGeneration": _HfExamplesInfo(
         "nvidia/NVIDIA-Nemotron-Parse-v1.1", trust_remote_code=True
     ),

tests/multimodal/media/test_video.py

@@ -1,9 +1,11 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import io
 from pathlib import Path
 
 import numpy as np
 import numpy.typing as npt
+import pybase64
 import pytest
 from PIL import Image
 
@@ -235,3 +237,53 @@ def test_video_media_io_backend_env_var_fallback(monkeypatch: pytest.MonkeyPatch
         frames_missing, metadata_missing = videoio_missing.load_bytes(b"test")
         np.testing.assert_array_equal(frames_missing, FAKE_OUTPUT_2)
         assert metadata_missing["video_backend"] == "test_video_backend_override_2"
+
+
+def test_load_base64_jpeg_returns_metadata():
+    """Regression test: load_base64 with video/jpeg must return metadata.
+
+    Previously, base64 JPEG frame sequences returned an empty dict for
+    metadata, which broke downstream consumers that rely on fields like
+    total_num_frames and fps. See PR #37301.
+    """
+
+    num_test_frames = 3
+    frame_width, frame_height = 8, 8
+
+    # Build a few tiny JPEG frames and base64-encode them
+    b64_frames = []
+    for i in range(num_test_frames):
+        img = Image.new("RGB", (frame_width, frame_height), color=(i * 80, 0, 0))
+        buf = io.BytesIO()
+        img.save(buf, format="JPEG")
+        b64_frames.append(pybase64.b64encode(buf.getvalue()).decode("ascii"))
+
+    data = ",".join(b64_frames)
+
+    imageio = ImageMediaIO()
+    videoio = VideoMediaIO(imageio, num_frames=num_test_frames)
+    frames, metadata = videoio.load_base64("video/jpeg", data)
+
+    # Frames array shape: (num_frames, H, W, 3)
+    assert frames.shape[0] == num_test_frames
+
+    # All required metadata keys must be present
+    required_keys = {
+        "total_num_frames",
+        "fps",
+        "duration",
+        "video_backend",
+        "frames_indices",
+        "do_sample_frames",
+    }
+    assert required_keys.issubset(metadata.keys()), (
+        f"Missing metadata keys: {required_keys - metadata.keys()}"
+    )
+
+    assert metadata["total_num_frames"] == num_test_frames
+    assert metadata["video_backend"] == "jpeg_sequence"
+    assert metadata["frames_indices"] == list(range(num_test_frames))
+    assert metadata["do_sample_frames"] is False
+    # Default fps=1 → duration == num_frames
+    assert metadata["fps"] == 1.0
+    assert metadata["duration"] == float(num_test_frames)

tests/tool_parsers/test_glm47_moe_tool_parser.py

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

tests/tool_parsers/test_glm4_moe_tool_parser.py

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

tests/v1/attention/test_mla_backends.py

@@ -266,22 +266,6 @@ def create_and_prepopulate_kv_cache(
     return kv_cache
 
 
-class MockAttentionLayer:
-    """A mock attention layer for testing."""
-
-    def __init__(self, device: torch.device):
-        self._q_scale = torch.tensor(1.0, device=device)
-        self._k_scale = torch.tensor(1.0, device=device)
-        self._v_scale = torch.tensor(1.0, device=device)
-        self._prob_scale = torch.tensor(1.0, device=device)
-        self._q_scale_float = 1.0
-        self._k_scale_float = 1.0
-        self._v_scale_float = 1.0
-
-    def forward(self, *_args, **_kwargs):
-        raise NotImplementedError
-
-
 class MockSparseMLAAttentionLayer:
     """A mock sparse MLA attention layer for testing.
 
@@ -304,6 +288,8 @@ class MockSparseMLAAttentionLayer:
         device: torch.device,
         W_UK: torch.Tensor,
         W_UV: torch.Tensor,
+        q_scale: float,
+        k_scale: float,
     ):
         self.impl = impl
         self.num_heads = num_heads
@@ -319,13 +305,13 @@ class MockSparseMLAAttentionLayer:
         self.W_UV = W_UV.transpose(0, 1)
 
         # Scale attributes needed by attention backends
-        self._q_scale = torch.tensor(1.0, device=device)
-        self._k_scale = torch.tensor(1.0, device=device)
-        self._v_scale = torch.tensor(1.0, device=device)
+        self._q_scale = torch.tensor(q_scale, device=device)
+        self._k_scale = torch.tensor(k_scale, device=device)
+        self._v_scale = torch.tensor(float("nan"), device=device)
         self._prob_scale = torch.tensor(1.0, device=device)
-        self._q_scale_float = 1.0
-        self._k_scale_float = 1.0
-        self._v_scale_float = 1.0
+        self._q_scale_float = q_scale
+        self._k_scale_float = k_scale
+        self._v_scale_float = float("nan")
 
         self._decode_concat_quant_fp8_op = _DecodeConcatQuantFP8(
             static=True,
@@ -420,6 +406,8 @@ class MockMLAAttentionLayer(AttentionLayerBase):
         kv_lora_rank: int,
         device: torch.device,
         kv_b_proj,
+        q_scale: float,
+        k_scale: float,
     ):
         self.impl = impl
         self.num_heads = num_heads
@@ -443,13 +431,13 @@ class MockMLAAttentionLayer(AttentionLayerBase):
         self.W_UK_T = W_UK.permute(1, 2, 0)
 
         # Scale attributes needed by attention backends
-        self._q_scale = torch.tensor(1.0, device=device)
-        self._k_scale = torch.tensor(1.0, device=device)
-        self._v_scale = torch.tensor(1.0, device=device)
+        self._q_scale = torch.tensor(q_scale, device=device)
+        self._k_scale = torch.tensor(k_scale, device=device)
+        self._v_scale = torch.tensor(float("nan"), device=device)
         self._prob_scale = torch.tensor(1.0, device=device)
-        self._q_scale_float = 1.0
-        self._k_scale_float = 1.0
-        self._v_scale_float = 1.0
+        self._q_scale_float = q_scale
+        self._k_scale_float = k_scale
+        self._v_scale_float = float("nan")
 
         self._decode_concat_quant_fp8_op = _DecodeConcatQuantFP8(
             static=True,
@@ -568,6 +556,8 @@ def run_attention_backend(
     qk_rope_head_dim: int,
     v_head_dim: int,
     mock_kv_b_proj,
+    q_scale: float,
+    k_scale: float,
     kv_cache_dtype: str = "auto",
 ) -> torch.Tensor:
     """Run attention computation using the specified backend's AttentionImpl."""
@@ -625,6 +615,8 @@ def run_attention_backend(
             kv_lora_rank=kv_lora_rank,
             device=device,
             kv_b_proj=mock_kv_b_proj,
+            q_scale=q_scale,
+            k_scale=k_scale,
         )
 
         # Populate static_forward_context with mock attention layers
@@ -674,6 +666,7 @@ def run_attention_backend(
 @pytest.mark.parametrize("model", ["deepseek-ai/DeepSeek-R1"])
 @pytest.mark.parametrize("tensor_parallel_size", [1, 4, 8, 16])
 @pytest.mark.parametrize("kv_cache_dtype", ["auto", "fp8", "fp8_e4m3"])
+@pytest.mark.parametrize(("q_scale", "k_scale"), [(1.0, 1.0), (2.0, 3.0)])
 def test_backend_correctness(
     default_vllm_config,
     dist_init,
@@ -681,6 +674,8 @@ def test_backend_correctness(
     model: str,
     tensor_parallel_size: int,
     kv_cache_dtype: str,
+    q_scale: float,
+    k_scale: float,
 ):
     """
     Test that all backends produce similar outputs to a reference implementation
@@ -709,6 +704,11 @@ def test_backend_correctness(
         for b in BACKENDS_TO_TEST
         if kv_cache_dtype in b.get_class().supported_kv_cache_dtypes
     ]
+    if (
+        q_scale != 1.0 or k_scale != 1.0
+    ) and AttentionBackendEnum.CUTLASS_MLA in backends_to_test:
+        # CUTLASS_MLA does not support non-1 Q/K scales
+        backends_to_test.remove(AttentionBackendEnum.CUTLASS_MLA)
     if not backends_to_test:
         pytest.skip(f"No backends support kv_cache_dtype={kv_cache_dtype}")
 
@@ -1029,6 +1029,7 @@ def test_backend_correctness(
             common_attn_metadata=common_attn_metadata,
             randomize_blocks=True,
             kv_cache_dtype=kv_cache_dtype,
+            scale=k_scale,
         )
         kv_cache_per_block_size[block_size] = kv_cache
 
@@ -1072,6 +1073,8 @@ def test_backend_correctness(
             qk_rope_head_dim,
             v_head_dim,
             mock_kv_b_proj,
+            q_scale=q_scale,
+            k_scale=k_scale,
             kv_cache_dtype=kv_cache_dtype,
         )
 

tests/v1/attention/test_sparse_mla_backends.py

@@ -178,6 +178,7 @@ def _quantize_dequantize_fp8_ds_mla(
 @pytest.mark.parametrize("kv_cache_dtype", ["auto", "fp8", "fp8_ds_mla"])
 @pytest.mark.parametrize("tensor_parallel_size", [1, 2, 4])
 @pytest.mark.parametrize("block_size", [32, 64])
+@pytest.mark.parametrize(("q_scale", "k_scale"), [(1.0, 1.0), (2.0, 3.0)])
 def test_sparse_backend_decode_correctness(
     default_vllm_config,
     dist_init,
@@ -187,6 +188,8 @@ def test_sparse_backend_decode_correctness(
     tensor_parallel_size,
     block_size,
     workspace_init,
+    q_scale: float,
+    k_scale: float,
 ):
     if kv_cache_dtype not in backend_cls.supported_kv_cache_dtypes:
         pytest.skip(f"{backend_cls.get_name()} does not support {kv_cache_dtype}")
@@ -332,7 +335,7 @@ def test_sparse_backend_decode_correctness(
     kv_c_contexts, k_pe_contexts = [], []
     reference_outputs = []
 
-    kv_cache_scale = torch.tensor(1.0, dtype=torch.float32, device=device)
+    kv_cache_scale = torch.tensor(k_scale, dtype=torch.float32, device=device)
     global_token_idx = 0
 
     for i in range(batch_spec.batch_size):
@@ -490,6 +493,8 @@ def test_sparse_backend_decode_correctness(
             device=device,
             W_UK=W_UK,
             W_UV=W_UV,
+            q_scale=q_scale,
+            k_scale=k_scale,
         )
 
     out_buffer = torch.empty(
@@ -513,7 +518,9 @@ def test_sparse_backend_decode_correctness(
     # FP8 quantization introduces some error, but should be within reasonable bounds
     # BF16 (auto) should be very accurate, FP8 allows slightly more tolerance
     if kv_cache_dtype.startswith("fp8"):
-        torch.testing.assert_close(backend_output, sdpa_reference, rtol=0.05, atol=0.05)
+        torch.testing.assert_close(
+            backend_output, sdpa_reference, rtol=0.065, atol=0.05
+        )
     else:
         torch.testing.assert_close(backend_output, sdpa_reference, rtol=0.01, atol=0.01)
 

tests/v1/attention/test_trtllm_attention_integration.py

@@ -43,12 +43,12 @@ class MockAttentionLayer:
     """Minimal mock of an attention layer for testing."""
 
     def __init__(self, device: torch.device):
-        self._q_scale = torch.tensor(1.0, device=device)
-        self._k_scale = torch.tensor(1.0, device=device)
-        self._v_scale = torch.tensor(1.0, device=device)
-        self._q_scale_float = 1.0
-        self._k_scale_float = 1.0
-        self._v_scale_float = 1.0
+        self._q_scale = torch.tensor(2.0, device=device)
+        self._k_scale = torch.tensor(3.0, device=device)
+        self._v_scale = torch.tensor(4.0, device=device)
+        self._q_scale_float = 2.0
+        self._k_scale_float = 3.0
+        self._v_scale_float = 4.0
         self._o_scale_float = None
 
 

tests/v1/core/utils.py

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

tests/v1/executor/test_executor.py

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

tests/v1/kv_connector/unit/test_decode_bench_connector.py

@@ -86,7 +86,7 @@ class DecodeBenchTestRunner:
         self._block_hasher = get_request_block_hasher(block_size, sha256)
 
         self._dummy_ctx: ForwardContext = ForwardContext(
-            no_compile_layers={}, attn_metadata={}, virtual_engine=0, slot_mapping={}
+            no_compile_layers={}, attn_metadata={}, slot_mapping={}
         )
 
     def new_request(self, token_ids: list[int]) -> Request:

tests/v1/kv_connector/unit/test_lmcache_integration.py

@@ -211,7 +211,6 @@ def test_forward_context_interface():
     from vllm.forward_context import ForwardContext
 
     assumes(ForwardContext, "no_compile_layers", is_instance_of=dict)
-    assumes(ForwardContext, "virtual_engine")
     assumes(ForwardContext, "attn_metadata")
 
 

tests/v1/kv_connector/unit/test_multi_connector.py

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

tests/v1/kv_connector/unit/test_nixl_connector.py

@@ -599,7 +599,6 @@ class TestNixlHandshake:
             dummy_ctx = ForwardContext(
                 no_compile_layers={},
                 attn_metadata={},
-                virtual_engine=0,
                 slot_mapping={},
             )
             _before_load = time.perf_counter()
@@ -672,7 +671,6 @@ class TestNixlHandshake:
             dummy_ctx = ForwardContext(
                 no_compile_layers={},
                 attn_metadata={},
-                virtual_engine=0,
                 slot_mapping={},
             )
             _before_load = time.perf_counter()
@@ -694,16 +692,18 @@ class TestNixlHandshake:
     )
     @pytest.mark.parametrize("local_tp_size", [1, 2])
     def test_prefill_tp_size_greater_than_decode_tp_size(
-        self, local_tp_size: int, default_vllm_config, dist_init
+        self, local_tp_size: int, default_vllm_config, dist_init, monkeypatch
     ):
         """
         Verify remote TP > local TP handshake succeeds with different
         remote configurations.
         """
+        monkeypatch.setattr(
+            "vllm.distributed.kv_transfer.kv_connector.v1.nixl_connector.get_tensor_model_parallel_world_size",
+            lambda: local_tp_size,
+        )
 
         vllm_config = create_vllm_config()
-        local_tp_size = 1
-        vllm_config.parallel_config.tensor_parallel_size = local_tp_size
 
         connector = NixlConnector(
             vllm_config, KVConnectorRole.WORKER, make_kv_cache_config(block_size=16)
@@ -738,10 +738,10 @@ class TestNixlHandshake:
         remote_agents = worker._nixl_handshake(
             host="localhost",
             port=1234,
-            remote_tp_size=2,
+            remote_tp_size=4,
             expected_engine_id=worker.REMOTE_ENGINE_ID,
         )
-        check_handshake(2)
+        check_handshake(4)
 
         # NOTE flexibility: a second remote with higher number of ranks is
         # discovered. This is not a scenario we actively support right now, but
@@ -759,9 +759,8 @@ class TestNixlHandshake:
         "vllm.distributed.kv_transfer.kv_connector.v1.nixl_connector.NixlWrapper",
         FakeNixlWrapper,
     )
-    @pytest.mark.parametrize("local_tp_size", [1, 2])
     def test_prefill_tp_size_greater_than_decode_tp_size_mla(
-        self, local_tp_size: int, default_vllm_config, dist_init
+        self, default_vllm_config, dist_init
     ):
         """
         Verify remote TP > local TP handshake succeeds with different
@@ -907,7 +906,6 @@ class TestNixlHandshake:
             dummy_ctx = ForwardContext(
                 no_compile_layers={},
                 attn_metadata={},
-                virtual_engine=0,
                 slot_mapping={},
             )
             _before_load = time.perf_counter()
@@ -1078,7 +1076,6 @@ def test_kv_connector_stats(default_vllm_config, dist_init):
     dummy_ctx = ForwardContext(
         no_compile_layers={},
         attn_metadata={},
-        virtual_engine=0,
         slot_mapping={},
     )
     connector.start_load_kv(dummy_ctx)
@@ -1369,7 +1366,13 @@ def test_abort_timeout_on_prefiller(monkeypatch, distributed_executor_backend):
                     "NIXL_TELEMETRY_ENABLE": "1",
                 },
             }
-            ray.init(runtime_env=runtime_env)
+            # On XPU/ROCm, vLLM expects Ray's device key to be "GPU".
+            # Explicitly reserving GPU resources here prevents false negatives
+            # when Ray cannot auto-detect accelerator resources in test envs.
+            ray_init_kwargs: dict[str, Any] = {"runtime_env": runtime_env}
+            if not current_platform.is_cuda():
+                ray_init_kwargs["num_gpus"] = 1
+            ray.init(**ray_init_kwargs)
             try:
                 run_test_and_cleanup()
             finally:
@@ -1883,7 +1886,6 @@ def test_aborted_request_removed_from_worker_in_batch(default_vllm_config, dist_
     dummy_ctx = ForwardContext(
         no_compile_layers={},
         attn_metadata={},
-        virtual_engine=0,
         slot_mapping={},
     )
     connector.start_load_kv(dummy_ctx)
@@ -2052,7 +2054,6 @@ def test_transfer_failure_logging(
     dummy_ctx = ForwardContext(
         no_compile_layers={},
         attn_metadata={},
-        virtual_engine=0,
         slot_mapping={},
     )
 
@@ -2155,7 +2156,6 @@ def test_handshake_failure_returns_finished(default_vllm_config, dist_init):
     dummy_ctx = ForwardContext(
         no_compile_layers={},
         attn_metadata={},
-        virtual_engine=0,
         slot_mapping={},
     )
     connector.start_load_kv(dummy_ctx)
@@ -2208,7 +2208,6 @@ def test_transfer_setup_failure_returns_finished(default_vllm_config, dist_init)
     dummy_ctx = ForwardContext(
         no_compile_layers={},
         attn_metadata={},
-        virtual_engine=0,
         slot_mapping={},
     )
     connector.start_load_kv(dummy_ctx)

tests/v1/kv_connector/unit/test_offloading_connector.py

@@ -13,10 +13,14 @@ from vllm import SamplingParams
 from vllm.config import KVTransferConfig, VllmConfig
 from vllm.distributed.kv_events import BlockRemoved, BlockStored
 from vllm.distributed.kv_transfer.kv_connector.v1 import KVConnectorRole
+from vllm.distributed.kv_transfer.kv_connector.v1.offloading.common import (
+    OffloadingConnectorMetadata,
+)
+from vllm.distributed.kv_transfer.kv_connector.v1.offloading.metrics import (
+    OffloadingConnectorStats,
+)
 from vllm.distributed.kv_transfer.kv_connector.v1.offloading_connector import (
     OffloadingConnector,
-    OffloadingConnectorMetadata,
-    OffloadingConnectorStats,
 )
 from vllm.forward_context import ForwardContext
 from vllm.utils.hashing import sha256
@@ -257,7 +261,6 @@ class RequestRunner:
         self._dummy_ctx: ForwardContext = ForwardContext(
             no_compile_layers={},
             attn_metadata={},
-            virtual_engine=0,
             slot_mapping={},
         )
 
@@ -363,10 +366,7 @@ class RequestRunner:
             assert kv_connector_metadata is not None
             assert isinstance(kv_connector_metadata, OffloadingConnectorMetadata)
 
-            if scheduler_output.preempted_req_ids:
-                self.worker_connector.handle_preemptions(
-                    scheduler_output.preempted_req_ids
-                )
+            self.worker_connector.handle_preemptions(kv_connector_metadata)
 
             self.worker_connector.bind_connector_metadata(kv_connector_metadata)
             self.worker_connector.start_load_kv(self._dummy_ctx)

tests/v1/kv_connector/unit/test_scheduler_kv_connector_override.py

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

tests/v1/kv_offload/test_cpu_gpu.py

@@ -135,19 +135,19 @@ def test_transfer(
     # set transfer direction
     if gpu_to_cpu:
         handler = handlers.gpu_to_cpu_handler
-        src_spec_class = GPULoadStoreSpec
-        dst_spec_class = CPULoadStoreSpec
         src_blocks = gpu_blocks
         dst_blocks = cpu_blocks
+        src_spec = GPULoadStoreSpec(src_blocks, group_sizes=(len(src_blocks),))
+        dst_spec = CPULoadStoreSpec(dst_blocks)
         src_blocks_in_kernel_block_size = gpu_blocks_in_kernel_block_size
         dst_blocks_in_kernel_block_size = cpu_blocks_in_kernel_block_size
         dst_size_in_kernel_blocks = num_cpu_blocks * kernel_blocks_per_cpu_block
     else:
         handler = handlers.cpu_to_gpu_handler
-        src_spec_class = CPULoadStoreSpec
-        dst_spec_class = GPULoadStoreSpec
         src_blocks = cpu_blocks
         dst_blocks = gpu_blocks
+        src_spec = CPULoadStoreSpec(src_blocks)
+        dst_spec = GPULoadStoreSpec(dst_blocks, group_sizes=(len(dst_blocks),))
         src_blocks_in_kernel_block_size = cpu_blocks_in_kernel_block_size
         dst_blocks_in_kernel_block_size = gpu_blocks_in_kernel_block_size
         dst_size_in_kernel_blocks = num_gpu_blocks * kernel_blocks_per_gpu_block
@@ -159,10 +159,6 @@ def test_transfer(
     ):
         dst_to_src[dst_block] = src_block
 
-    # build transfer specs
-    src_spec = src_spec_class(src_blocks)
-    dst_spec = dst_spec_class(dst_blocks)
-
     # clone src and dst tensors before transfer
     orig_src_caches = [x.clone() for x in handler.src_tensors]
     orig_dst_caches = [x.clone() for x in handler.dst_tensors]

tests/v1/kv_offload/test_cpu_offloading.py

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

tests/v1/metrics/test_perf_metrics.py

@@ -7,6 +7,7 @@ Tests for the analytic estimators in metrics/flops.py.
 import types
 from types import SimpleNamespace
 
+import pytest
 from transformers.models.deepseek_v3.configuration_deepseek_v3 import DeepseekV3Config
 from transformers.models.llama4.configuration_llama4 import (
     Llama4Config,
@@ -21,10 +22,12 @@ from vllm.transformers_utils.model_arch_config_convertor import (
     ModelArchConfigConvertorBase,
 )
 from vllm.v1.metrics.perf import (
+    _QUANT_WEIGHT_BYTE_SIZE,
     AttentionMetrics,
     BaseConfigParser,
     ExecutionContext,
     FfnMetrics,
+    InvalidComponent,
     ModelMetrics,
     ParsedArgs,
     UnembedMetrics,
@@ -905,3 +908,116 @@ def test_attention_per_gpu_heads_not_evenly_divisible():
     assert per_gpu_flops > 0
     assert global_flops > 0
     assert global_flops > per_gpu_flops
+
+
+# INT4 / FP4 quantization methods (weight_byte_size == 0.5)
+_INT4_FP4_METHODS = [m for m, s in _QUANT_WEIGHT_BYTE_SIZE.items() if s == 0.5]
+
+
+@pytest.mark.parametrize("quant_method", _INT4_FP4_METHODS)
+def test_quantization_config_parser_int4_methods(quant_method):
+    """Test quantization parsers with INT4/FP4 methods (0.5 bytes)."""
+
+    class MockQuantConfig:
+        def get_name(self):
+            return quant_method
+
+    hf_config = Qwen3Config(
+        hidden_size=2048,
+        num_attention_heads=16,
+        intermediate_size=8192,
+        num_hidden_layers=1,
+    )
+    vllm_config = create_mock_vllm_config(hf_config, quant_config=MockQuantConfig())
+
+    attn_result = AttentionMetrics.get_parser().parse(vllm_config)
+    assert attn_result.weight_byte_size == 0.5, (
+        f"Expected 0.5 for {quant_method}, got {attn_result.weight_byte_size}"
+    )
+
+    ffn_result = FfnMetrics.get_parser().parse(vllm_config)
+    assert ffn_result.weight_byte_size == 0.5, (
+        f"Expected 0.5 for {quant_method}, got {ffn_result.weight_byte_size}"
+    )
+
+
+# FP8 / INT8 quantization methods (weight_byte_size == 1)
+_FP8_INT8_METHODS = [m for m, s in _QUANT_WEIGHT_BYTE_SIZE.items() if s == 1]
+
+
+@pytest.mark.parametrize("quant_method", _FP8_INT8_METHODS)
+def test_quantization_config_parser_fp8_methods(quant_method):
+    """Test quantization parsers with FP8/INT8 methods (1 byte)."""
+
+    class MockQuantConfig:
+        def get_name(self):
+            return quant_method
+
+    hf_config = Qwen3Config(
+        hidden_size=2048,
+        num_attention_heads=16,
+        intermediate_size=8192,
+        num_hidden_layers=1,
+    )
+    vllm_config = create_mock_vllm_config(hf_config, quant_config=MockQuantConfig())
+
+    attn_result = AttentionMetrics.get_parser().parse(vllm_config)
+    assert attn_result.weight_byte_size == 1, (
+        f"Expected 1 for {quant_method}, got {attn_result.weight_byte_size}"
+    )
+
+    ffn_result = FfnMetrics.get_parser().parse(vllm_config)
+    assert ffn_result.weight_byte_size == 1, (
+        f"Expected 1 for {quant_method}, got {ffn_result.weight_byte_size}"
+    )
+
+
+def test_quantization_config_parser_unknown_method():
+    """Test that an unrecognized quant method raises InvalidComponent."""
+
+    class MockQuantConfig:
+        def get_name(self):
+            return "unknown_quant_method"
+
+    hf_config = Qwen3Config(
+        hidden_size=2048,
+        num_attention_heads=16,
+        intermediate_size=8192,
+        num_hidden_layers=1,
+    )
+    vllm_config = create_mock_vllm_config(hf_config, quant_config=MockQuantConfig())
+
+    with pytest.raises(InvalidComponent):
+        AttentionMetrics.get_parser().parse(vllm_config)
+
+    with pytest.raises(InvalidComponent):
+        FfnMetrics.get_parser().parse(vllm_config)
+
+
+def test_quantized_model_metrics_aggregation():
+    """Test that ModelMetrics works end-to-end with a quantized model config."""
+
+    class MockQuantConfig:
+        def get_name(self):
+            return "gptq"
+
+    hf_config = Qwen3Config(
+        hidden_size=2048,
+        num_attention_heads=16,
+        num_hidden_layers=12,
+        vocab_size=32000,
+        intermediate_size=8192,
+    )
+    vllm_config = create_mock_vllm_config(hf_config, quant_config=MockQuantConfig())
+
+    model_metrics = ModelMetrics(vllm_config)
+    ctx = ExecutionContext.from_single_request(
+        num_tokens=100, context_len=512, is_prefill=True
+    )
+
+    # Should not crash and should produce valid metrics
+    total_flops = model_metrics.get_num_flops(ctx)
+    breakdown = model_metrics.get_num_flops_breakdown(ctx)
+
+    assert total_flops > 0
+    assert total_flops == sum(breakdown.values())

vllm/_custom_ops.py

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

vllm/benchmarks/datasets.py

@@ -183,6 +183,68 @@ class BenchmarkDataset(ABC):
         )
         return lora_request
 
+    def get_round_robin_lora_request(
+        self,
+        index: int,
+        max_loras: int | None = None,
+        lora_path: str | None = None,
+    ) -> LoRARequest | None:
+        """
+        Optionally select a LoRA request using deterministic round-robin.
+
+        This method cycles through LoRA IDs in order based on the request
+        index, providing reproducible LoRA assignment.
+
+        Args:
+            index (int): The request index used for round-robin selection.
+            max_loras (Optional[int]): The maximum number of LoRAs available.
+                If `None`, LoRA is not used.
+            lora_path (Optional[str]): Path to the LoRA parameters on disk.
+                If `None`, LoRA is not used.
+
+        Returns:
+            A new [`LoRARequest`][vllm.lora.request.LoRARequest]
+            (or `None` if not applicable).
+        """
+        if max_loras is None or lora_path is None:
+            return None
+
+        # Deterministic round-robin: cycle through [1, max_loras]
+        lora_id = index % max_loras + 1
+        lora_request = LoRARequest(
+            lora_name=str(lora_id),
+            lora_int_id=lora_id,
+            lora_path=lora_path_on_disk(lora_path),
+        )
+        return lora_request
+
+    def get_lora_request(
+        self,
+        index: int,
+        max_loras: int | None = None,
+        lora_path: str | None = None,
+        lora_assignment: str = "random",
+    ) -> LoRARequest | None:
+        """
+        Select a LoRA request using the specified assignment strategy.
+
+        Args:
+            index (int): The request index (used for round-robin).
+            max_loras (Optional[int]): The maximum number of LoRAs available.
+            lora_path (Optional[str]): Path to the LoRA parameters on disk.
+            lora_assignment (str): Strategy for LoRA selection.
+                'random' (default) or 'round-robin'.
+
+        Returns:
+            A new [`LoRARequest`][vllm.lora.request.LoRARequest]
+            (or `None` if not applicable).
+        """
+        if lora_assignment == "round-robin":
+            return self.get_round_robin_lora_request(
+                index=index, max_loras=max_loras, lora_path=lora_path
+            )
+        return self.get_random_lora_request(max_loras=max_loras, lora_path=lora_path)
+
     @abstractmethod
     def sample(
         self,
@@ -478,6 +540,9 @@ class RandomDataset(BenchmarkDataset):
         input_len: int = DEFAULT_INPUT_LEN,
         output_len: int = DEFAULT_OUTPUT_LEN,
         batchsize: int = 1,
+        max_loras: int | None = None,
+        lora_path: str | None = None,
+        lora_assignment: str = "random",
         **kwargs,
     ) -> list[SampleRequest]:
         # validate total input tokens (prefix + sampled) is at least 1.
@@ -522,11 +587,18 @@ class RandomDataset(BenchmarkDataset):
                 allowed_tokens=allowed_tokens,
             )
             token_mismatch_total += token_mismatch
+            lora_req = self.get_lora_request(
+                index=i,
+                max_loras=max_loras,
+                lora_path=lora_path,
+                lora_assignment=lora_assignment,
+            )
             requests.append(
                 SampleRequest(
                     prompt=prompt,
                     prompt_len=total_input_len,
                     expected_output_len=int(output_lens[i]),
+                    lora_request=lora_req,
                     request_id=request_id_prefix + str(i),
                 )
             )
@@ -1263,6 +1335,7 @@ class ShareGPTDataset(BenchmarkDataset):
         enable_multimodal_chat: bool = False,
         request_id_prefix: str = "",
         no_oversample: bool = False,
+        lora_assignment: str = "random",
         **kwargs,
     ) -> list:
         samples: list = []
@@ -1275,8 +1348,11 @@ class ShareGPTDataset(BenchmarkDataset):
                 entry["conversations"][1]["value"],
             )
 
-            lora_request = self.get_random_lora_request(
-                max_loras=max_loras, lora_path=lora_path
+            lora_request = self.get_lora_request(
+                index=ind,
+                max_loras=max_loras,
+                lora_path=lora_path,
+                lora_assignment=lora_assignment,
             )
             prompt_ids = tokenizer(prompt).input_ids
             completion_ids = tokenizer(completion).input_ids
@@ -2413,6 +2489,7 @@ class BurstGPTDataset(BenchmarkDataset):
         lora_path: str | None = None,
         request_id_prefix: str = "",
         no_oversample: bool = False,
+        lora_assignment: str = "random",
         **kwargs,
     ) -> list[SampleRequest]:
         samples = []
@@ -2420,8 +2497,11 @@ class BurstGPTDataset(BenchmarkDataset):
         for i in range(num_requests):
             input_len = int(data[i][2])
             output_len = int(data[i][3])
-            lora_req = self.get_random_lora_request(
-                max_loras=max_loras, lora_path=lora_path
+            lora_req = self.get_lora_request(
+                index=i,
+                max_loras=max_loras,
+                lora_path=lora_path,
+                lora_assignment=lora_assignment,
             )
             vocab_size = tokenizer.vocab_size
             # Generate a synthetic prompt: a list of token IDs computed as (i +
@@ -3157,7 +3237,7 @@ class ASRDataset(HuggingFaceDataset):
         **kwargs,
     ) -> list:
         output_len = output_len if output_len is not None else self.DEFAULT_OUTPUT_LEN
-        if "openai" in tokenizer.name_or_path:
+        if "openai" in getattr(tokenizer, "name_or_path", ""):
             prompt = "<|startoftranscript|><|en|><|transcribe|><|notimestamps|>"
         else:
             prompt = ""

vllm/benchmarks/serve.py

@@ -624,6 +624,7 @@ async def benchmark(
     lora_modules: Iterable[str] | None,
     extra_headers: dict | None,
     extra_body: dict | None,
+    lora_assignment: Literal["random", "round-robin"] = "random",
     ramp_up_strategy: Literal["linear", "exponential"] | None = None,
     ramp_up_start_rps: int | None = None,
     ramp_up_end_rps: int | None = None,
@@ -731,10 +732,20 @@ async def benchmark(
     print("Starting main benchmark run...")
 
     if lora_modules:
-        # For each input request, choose a LoRA module at random.
-        lora_modules = iter(
-            [random.choice(lora_modules) for _ in range(len(input_requests))]
-        )
+        lora_modules_list = list(lora_modules)
+        if lora_assignment == "round-robin":
+            # Deterministic round-robin assignment across requests.
+            lora_modules = iter(
+                [
+                    lora_modules_list[i % len(lora_modules_list)]
+                    for i in range(len(input_requests))
+                ]
+            )
+        else:
+            # For each input request, choose a LoRA module at random.
+            lora_modules = iter(
+                [random.choice(lora_modules_list) for _ in range(len(input_requests))]
+            )
 
     if profile:
         print("Starting profiler...")
@@ -1523,7 +1534,18 @@ def add_cli_args(parser: argparse.ArgumentParser):
         default=None,
         help="A subset of LoRA module names passed in when "
         "launching the server. For each request, the "
-        "script chooses a LoRA module at random.",
+        "script chooses a LoRA module at random by default. "
+        "Use --lora-assignment to control selection strategy.",
+    )
+
+    parser.add_argument(
+        "--lora-assignment",
+        type=str,
+        default="random",
+        choices=["random", "round-robin"],
+        help="Strategy for assigning LoRA modules to requests. "
+        "'random' (default) selects a LoRA at random for each request. "
+        "'round-robin' cycles through LoRA modules deterministically.",
     )
 
     parser.add_argument(
@@ -1788,6 +1810,7 @@ async def main_async(args: argparse.Namespace) -> dict[str, Any]:
         goodput_config_dict=goodput_config_dict,
         max_concurrency=args.max_concurrency,
         lora_modules=args.lora_modules,
+        lora_assignment=args.lora_assignment,
         extra_headers=headers,
         extra_body=extra_body,
         ramp_up_strategy=args.ramp_up_strategy,

vllm/benchmarks/throughput.py

@@ -350,6 +350,7 @@ def get_requests(args, tokenizer):
         "tokenizer": tokenizer,
         "lora_path": args.lora_path,
         "max_loras": args.max_loras,
+        "lora_assignment": getattr(args, "lora_assignment", "random"),
         "num_requests": args.num_prompts,
     }
 
@@ -778,6 +779,15 @@ def add_cli_args(parser: argparse.ArgumentParser):
         help="Path to the lora adapters to use. This can be an absolute path, "
         "a relative path, or a Hugging Face model identifier.",
     )
+    parser.add_argument(
+        "--lora-assignment",
+        type=str,
+        default="random",
+        choices=["random", "round-robin"],
+        help="Strategy for assigning LoRA adapters to requests. "
+        "'random' (default) selects a LoRA at random for each request. "
+        "'round-robin' cycles through LoRAs deterministically.",
+    )
     parser.add_argument(
         "--prefix-len",
         type=int,

vllm/compilation/backends.py

@@ -371,13 +371,15 @@ class CompilerManager:
                 logger.info_once(
                     "Cache the graph of compile range %s for later use",
                     str(compile_range),
+                    scope="local",
                 )
-            logger.debug(
+            logger.debug_once(
                 "Store the %s-th graph for compile range%s from %s via handle %s",
                 graph_index,
                 str(compile_range),
                 self.compiler.name,
                 handle,
+                scope="local",
             )
 
         # after compiling the last graph, record the end time

vllm/config/parallel.py

@@ -2,6 +2,7 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import os
+import socket
 from collections.abc import Callable
 from typing import TYPE_CHECKING, Any, Literal, overload
 
@@ -138,6 +139,13 @@ class ParallelConfig:
     """Whether the deployed model is MoE (if known)."""
     enable_expert_parallel: bool = False
     """Use expert parallelism instead of tensor parallelism for MoE layers."""
+    enable_ep_weight_filter: bool = False
+    """Skip non-local expert weights during model loading when expert
+    parallelism is active.  Each rank only reads its own expert shard from
+    disk, which can drastically reduce storage I/O for MoE models with
+    per-expert weight tensors (e.g. DeepSeek, Mixtral, Kimi-K2.5).  Has no
+    effect on 3D fused-expert checkpoints (e.g. GPT-OSS) or non-MoE
+    models."""
     enable_eplb: bool = False
     """Enable expert parallelism load balancing for MoE layers."""
     eplb_config: EPLBConfig = Field(default_factory=EPLBConfig)
@@ -259,33 +267,9 @@ class ParallelConfig:
     Set to be private as it's not intended to be configured by users.
     """
 
-    _stateless_dp_group_port_list: list[list[int]] = Field(default_factory=list)
-    """List of open ports for stateless DP groups when enable_elastic_ep is True.
-    Set to be private as it's not intended to be configured by users.
-    It is a list of list[int], with each inner list contains a set of 3 ports
-    to be used for setting up the stateless CPU/device/TCPStore groups
-    in StatelessGroupCoordinator. The number of inner lists is equal to
-    the number of DP groups, 
-    i.e., len(self._stateless_dp_group_port_list) == world_size_across_dp // dp_size,
-    and len(self._stateless_dp_group_port_list[i]) == 3 for all i.
-    """
-
-    _stateless_ep_group_port_list: list[list[int]] = Field(default_factory=list)
-    """List of open ports for stateless EP groups when enable_elastic_ep is True.
-    Set to be private as it's not intended to be configured by users.
-    len(self._stateless_ep_group_port_list) == world_size_across_dp // ep_size,
-    """
-
-    _stateless_eplb_group_port_list: list[list[int]] = Field(default_factory=list)
-    """List of open ports for stateless EPLB groups when enable_elastic_ep is True.
-    Same topology as EP but separate NCCL communicator to avoid deadlocks.
-    """
-
-    _stateless_world_group_port_list: list[list[int]] = Field(default_factory=list)
-    """List of open ports for stateless world group when enable_elastic_ep is True.
-    Set to be private as it's not intended to be configured by users.
-    len(self._stateless_world_group_port_list) == 1,
-    """
+    _coord_store_port: int = 0
+    """Port of the coordination TCPStore. Can be set by the API server; workers
+    connect as clients to exchange self-picked group ports at runtime."""
 
     decode_context_parallel_size: int = 1
     """Number of decode context parallel groups, because the world size does
@@ -458,65 +442,32 @@ class ParallelConfig:
 
         return answer
 
-    def allocate_elastic_ep_ports(self) -> None:
-        """Allocate all ports for elastic EP (stateless groups + DP master).
+    def _pick_stateless_dp_port(self) -> tuple[int, socket.socket | None]:
+        """Return ``(port, listen_socket)`` for DP group init.
 
-        Must be called AFTER ray.init() so that ports claimed by Ray's
-        idle worker pool are already in use and won't be returned by
-        get_open_ports_list().
+        With a coord store, rank 0 binds a socket and publishes the port;
+        others read it.  Without one, pops a pre-allocated port and
+        returns ``listen_socket=None``.
         """
-        if not self.enable_elastic_ep:
-            return
-        if self._stateless_world_group_port_list:
-            return
+        if not self._coord_store_port:
+            return self.get_next_dp_init_port(), None
 
-        num_world_groups = 1
-        dp_size = self.data_parallel_size
-        ep_size = self.data_parallel_size * self.world_size_across_dp
-        num_dp_groups = max(1, self.world_size_across_dp // dp_size)
-        num_ep_groups = max(1, self.world_size_across_dp // ep_size)
-        num_eplb_groups = num_ep_groups
-        total_stateless_ports = (
-            num_world_groups + num_dp_groups + num_ep_groups + num_eplb_groups
-        ) * 3
-        num_dp_master_ports = 5
+        from vllm.distributed.utils import get_cached_tcp_store_client
 
-        all_ports = get_open_ports_list(total_stateless_ports + num_dp_master_ports)
+        store = get_cached_tcp_store_client(
+            self.data_parallel_master_ip, self._coord_store_port
+        )
 
-        self._data_parallel_master_port_list = all_ports[-num_dp_master_ports:]
-        self.data_parallel_master_port = self._data_parallel_master_port_list.pop()
-        all_ports = all_ports[:-num_dp_master_ports]
-
-        self._stateless_world_group_port_list = [
-            all_ports[i : i + 3] for i in range(0, num_world_groups * 3, 3)
-        ]
-        start_idx = num_world_groups * 3
-        self._stateless_dp_group_port_list = [
-            all_ports[i : i + 3]
-            for i in range(start_idx, start_idx + num_dp_groups * 3, 3)
-        ]
-        start_idx += num_dp_groups * 3
-        self._stateless_ep_group_port_list = [
-            all_ports[i : i + 3]
-            for i in range(start_idx, start_idx + num_ep_groups * 3, 3)
-        ]
-        start_idx += num_ep_groups * 3
-        self._stateless_eplb_group_port_list = [
-            all_ports[i : i + 3]
-            for i in range(start_idx, start_idx + num_eplb_groups * 3, 3)
-        ]
-
-    def get_next_stateless_world_group_port(self) -> list[int]:
-        return self._stateless_world_group_port_list.pop()
-
-    def get_next_stateless_dp_group_port(self) -> list[int]:
-        return self._stateless_dp_group_port_list.pop()
-
-    def get_next_stateless_ep_group_port(self) -> list[int]:
-        return self._stateless_ep_group_port_list.pop()
-
-    def get_next_stateless_eplb_group_port(self) -> list[int]:
-        return self._stateless_eplb_group_port_list.pop()
+        key = "dp_master_port"
+        if self.data_parallel_rank == 0:
+            s = socket.socket(socket.AF_INET, socket.SOCK_STREAM)
+            s.bind((self.data_parallel_master_ip, 0))
+            s.listen()
+            port = s.getsockname()[1]
+            store.set(key, str(port).encode())
+            return port, s
+        else:
+            return int(store.get(key).decode()), None
 
     @overload
     def stateless_init_dp_group(
@@ -546,14 +497,16 @@ class ParallelConfig:
         last_exc: Exception | None = None
         for _ in range(max_retries):
             try:
+                port, listen_socket = self._pick_stateless_dp_port()
                 # use gloo since the engine process might not have cuda device
                 return stateless_init_torch_distributed_process_group(
                     self.data_parallel_master_ip,
-                    self.get_next_dp_init_port(),
+                    port,
                     self.data_parallel_rank,
                     self.data_parallel_size,
                     backend="gloo",
                     return_store=return_store,
+                    listen_socket=listen_socket,
                 )
             except DistNetworkError as e:
                 # We only want to retry when the root cause is EADDRINUSE.

vllm/config/scheduler.py

@@ -228,9 +228,10 @@ class SchedulerConfig:
         self.encoder_cache_size = self.max_num_batched_tokens
 
         if self.enable_chunked_prefill:
-            logger.info(
+            logger.info_once(
                 "Chunked prefill is enabled with max_num_batched_tokens=%d.",
                 self.max_num_batched_tokens,
+                scope="local",
             )
 
         if self.max_num_partial_prefills > 1:

vllm/config/vllm.py

@@ -682,12 +682,11 @@ class VllmConfig:
                 self.model_config, self.load_config
             )
 
+        from vllm.v1.executor.abstract import Executor
+
         executor_backend = self.parallel_config.distributed_executor_backend
-        executor_supports_async_sched = executor_backend in (
-            "mp",
-            "uni",
-            "external_launcher",
-        )
+        executor_class = Executor.get_class(self)
+        executor_supports_async_sched = executor_class.supports_async_scheduling()
 
         if self.scheduler_config.async_scheduling:
             # Async scheduling explicitly enabled, hard fail any incompatibilities.
@@ -711,9 +710,7 @@ class VllmConfig:
                     )
             if not executor_supports_async_sched:
                 raise ValueError(
-                    "Currently, async scheduling only supports `mp`, `uni`, or "
-                    "`external_launcher` distributed executor backend, but you chose "
-                    f"`{executor_backend}`."
+                    f"`{executor_backend}` does not support async scheduling yet."
                 )
         elif self.scheduler_config.async_scheduling is None:
             # Enable async scheduling unless there is an incompatible option.
@@ -742,8 +739,7 @@ class VllmConfig:
             elif not executor_supports_async_sched:
                 logger.warning_once(
                     "Async scheduling will be disabled because it is not supported "
-                    "with the `%s` distributed executor backend (only `mp`, `uni`, and "
-                    "`external_launcher` are supported).",
+                    "with the `%s` distributed executor backend. ",
                     executor_backend,
                     scope="local",
                 )

vllm/distributed/elastic_ep/elastic_execute.py

@@ -162,10 +162,8 @@ class ElasticEPScalingExecutor:
                 new_dp_size=new_dp_size,
                 new_world_size_across_dp=new_world_size_across_dp,
                 master_ip=reconfig_request.new_data_parallel_master_ip,
-                world_group_ports=reconfig_request.new_stateless_world_group_port_list,
-                dp_group_ports=reconfig_request.new_stateless_dp_group_port_list,
-                ep_group_ports=reconfig_request.new_stateless_ep_group_port_list,
-                eplb_group_ports=reconfig_request.new_stateless_eplb_group_port_list,
+                coord_store_port=reconfig_request.coord_store_port,
+                enable_eplb=updated_config.parallel_config.enable_eplb,
             )
         self.worker.model_runner.eep_eplb_suppressed = True
         standby_ep_group = get_standby_ep_group()

vllm/distributed/elastic_ep/elastic_state.py

@@ -563,15 +563,4 @@ class ElasticEPScalingState:
         parallel_config._data_parallel_master_port_list = (
             reconfig_request.new_data_parallel_master_port_list
         )
-        parallel_config._stateless_world_group_port_list = (
-            reconfig_request.new_stateless_world_group_port_list
-        )
-        parallel_config._stateless_dp_group_port_list = (
-            reconfig_request.new_stateless_dp_group_port_list
-        )
-        parallel_config._stateless_ep_group_port_list = (
-            reconfig_request.new_stateless_ep_group_port_list
-        )
-        parallel_config._stateless_eplb_group_port_list = (
-            reconfig_request.new_stateless_eplb_group_port_list
-        )
+        parallel_config._coord_store_port = reconfig_request.coord_store_port

vllm/distributed/elastic_ep/standby_state.py

@@ -38,10 +38,8 @@ def create_standby_groups(
     new_dp_size: int,
     new_world_size_across_dp: int,
     master_ip: str,
-    world_group_ports: list[list[int]],
-    dp_group_ports: list[list[int]],
-    ep_group_ports: list[list[int]],
-    eplb_group_ports: list[list[int]] | None = None,
+    coord_store_port: int,
+    enable_eplb: bool = True,
     backend: str | None = None,
 ) -> None:
     global \
@@ -51,19 +49,23 @@ def create_standby_groups(
         _STANDBY_EP, \
         _STANDBY_EPLB
 
+    from vllm.distributed.utils import get_cached_tcp_store_client
+
     assert new_world_size_across_dp == torch.distributed.get_world_size() * new_dp_size
     world_group = get_world_group()
     assert isinstance(world_group, StatelessGroupCoordinator)
     backend = backend or world_group.backend
 
+    coord_store = get_cached_tcp_store_client(master_ip, coord_store_port)
+
     standby_world_ranks = [list(range(new_world_size_across_dp))]
     _STANDBY_WORLD = _init_stateless_group(
         standby_world_ranks,
         "world",
-        world_group_ports,
         master_ip,
         backend,
         use_device_communicator=False,
+        coord_store=coord_store,
     )
     _STANDBY_WORLD_NODE_COUNT = _node_count(_STANDBY_WORLD.tcp_store_group)
 
@@ -76,7 +78,7 @@ def create_standby_groups(
     standby_dp_ranks = all_ranks.transpose(1, 3).reshape(-1, new_dp_size).unbind(0)
     standby_dp_ranks = [x.tolist() for x in standby_dp_ranks]
     _STANDBY_DP = _init_stateless_group(
-        standby_dp_ranks, "dp", dp_group_ports, master_ip, backend
+        standby_dp_ranks, "dp", master_ip, backend, coord_store=coord_store
     )
 
     standby_ep_ranks = (
@@ -84,12 +86,16 @@ def create_standby_groups(
     )
     standby_ep_ranks = [x.tolist() for x in standby_ep_ranks]
     _STANDBY_EP = _init_stateless_group(
-        standby_ep_ranks, "ep", ep_group_ports, master_ip, backend
+        standby_ep_ranks, "ep", master_ip, backend, coord_store=coord_store
     )
 
-    if eplb_group_ports is not None:
+    if enable_eplb:
         _STANDBY_EPLB = _init_stateless_group(
-            standby_ep_ranks, "eplb", eplb_group_ports, master_ip, backend
+            standby_ep_ranks,
+            "eplb",
+            master_ip,
+            backend,
+            coord_store=coord_store,
         )
 
 

vllm/distributed/eplb/async_worker.py

@@ -73,11 +73,7 @@ def run_rebalance_experts(
     # Move the global expert load window to CPU for computation.
     global_expert_load_window = eplb_stats.global_expert_load_window.cpu()
     # Compute new expert mappings for the model
-    (
-        new_physical_to_logical_map,
-        new_logical_to_physical_map,
-        new_logical_replica_count,
-    ) = eplb_state.policy.rebalance_experts(
+    new_physical_to_logical_map = eplb_state.policy.rebalance_experts(
         global_expert_load_window,
         eplb_stats.num_replicas,
         eplb_stats.num_groups,
@@ -89,16 +85,6 @@ def run_rebalance_experts(
 
     model_state.new_physical_to_logical_map = new_physical_to_logical_map
 
-    max_slots = model_state.logical_to_physical_map.shape[-1]
-    padded_logical = torch.nn.functional.pad(
-        new_logical_to_physical_map,
-        (0, max(0, max_slots - new_logical_to_physical_map.shape[-1])),
-        value=-1,
-    ).to(model_state.logical_to_physical_map.device)
-    new_replica = new_logical_replica_count.to(model_state.logical_replica_count.device)
-    model_state.new_logical_to_physical_map = padded_logical
-    model_state.new_logical_replica_count = new_replica
-
 
 async def transfer_run_periodically(
     state: "EplbState",

vllm/distributed/eplb/eplb_state.py

@@ -235,16 +235,6 @@ class EplbModelState:
     intermediate variable between `move_to_buffer` and `move_to_workspace`.
     the size is same as physical_to_logical_map
     """
-    new_logical_to_physical_map: torch.Tensor | None = None
-    """
-    intermediate variable between `move_to_buffer` and `move_to_workspace`.
-    the size is same as logical_to_physical_map
-    """
-    new_logical_replica_count: torch.Tensor | None = None
-    """
-    intermediate variable between `move_to_buffer` and `move_to_workspace`.
-    the size is same as logical_replica_count
-    """
 
 
 class EplbState:
@@ -508,8 +498,6 @@ class EplbState:
             ),
             cuda_device_index=self.cuda_device_index,
             new_physical_to_logical_map=None,
-            new_logical_to_physical_map=None,
-            new_logical_replica_count=None,
         )
         self.model_states[model_config.compute_hash()] = model_state
         self.num_valid_physical_experts = model.num_physical_experts
@@ -738,17 +726,20 @@ class EplbState:
         ):
             if not self.is_async or is_profile:
                 # Get new expert mappings for the model
-                (
-                    new_physical_to_logical_map,
-                    new_logical_to_physical_map,
-                    new_logical_replica_count,
-                ) = self.policy.rebalance_experts(
-                    global_expert_load_window,
+                new_physical_to_logical_map = self.policy.rebalance_experts(
+                    global_expert_load_window.cpu(),
                     num_replicas,
                     num_groups,
                     num_nodes,
                     num_gpus,
-                    eplb_model_state.physical_to_logical_map,
+                    eplb_model_state.physical_to_logical_map.cpu(),
+                )
+
+                num_logical_experts = global_expert_load_window.shape[-1]
+                (new_logical_to_physical_map, new_logical_replica_count) = (
+                    compute_logical_maps(
+                        new_physical_to_logical_map, num_logical_experts
+                    )
                 )
 
                 # Update expert weights
@@ -847,11 +838,7 @@ class EplbState:
     def _update_layer_mapping_from_new(
         self, model_state: EplbModelState, layer: int
     ) -> None:
-        if (
-            model_state.new_physical_to_logical_map is None
-            or model_state.new_logical_to_physical_map is None
-            or model_state.new_logical_replica_count is None
-        ):
+        if model_state.new_physical_to_logical_map is None:
             return
 
         target_device = model_state.physical_to_logical_map.device
@@ -865,19 +852,23 @@ class EplbState:
                 new_physical[layer].to(target_device, non_blocking=True)
             )
 
+        num_logical_experts = model_state.logical_to_physical_map.shape[1]
+        new_logical, new_replica_count = compute_logical_maps(
+            new_physical[layer], num_logical_experts
+        )
+
         logical_device = model_state.logical_to_physical_map.device
-        new_logical = model_state.new_logical_to_physical_map[layer].to(logical_device)
         max_slots = model_state.logical_to_physical_map.shape[-1]
         slot_delta = max_slots - new_logical.shape[-1]
         if slot_delta > 0:
             new_logical = torch.nn.functional.pad(
                 new_logical, (0, slot_delta), value=-1
             )
-        model_state.logical_to_physical_map[layer].copy_(new_logical)
+        model_state.logical_to_physical_map[layer].copy_(new_logical.to(logical_device))
 
         replica_device = model_state.logical_replica_count.device
         model_state.logical_replica_count[layer].copy_(
-            model_state.new_logical_replica_count[layer].to(replica_device)
+            new_replica_count.to(replica_device)
         )
 
     def _all_ranks_buffer_ready(self, model_state: EplbModelState) -> bool:
@@ -966,7 +957,7 @@ class EplbState:
                 transferred_layer,
             )
             if model_state.layer_to_transfer >= model_state.model.num_moe_layers:
-                self.post_eplb(model_state, is_profile)
+                self.post_eplb(model_state)
                 model_state.rebalanced = False
                 model_state.layer_to_transfer = 0
                 model_state.pending_global_ready_check = False
@@ -987,14 +978,9 @@ class EplbState:
                     str(e),
                 )
 
-    def post_eplb(self, model_state: EplbModelState, is_profile: bool = False) -> None:
+    def post_eplb(self, model_state: EplbModelState) -> None:
         assert model_state.new_physical_to_logical_map is not None
-        assert model_state.new_logical_to_physical_map is not None
-        assert model_state.new_logical_replica_count is not None
-
         model_state.new_physical_to_logical_map = None
-        model_state.new_logical_to_physical_map = None
-        model_state.new_logical_replica_count = None
 
     def _allreduce_list(self, tensor_list: list[torch.Tensor]) -> list[torch.Tensor]:
         """
@@ -1052,39 +1038,28 @@ class EplbState:
             model_config=model_config,
         )
         eplb_state.num_valid_physical_experts = num_valid_physical_experts
-        num_moe_layers = expanded_physical_to_logical.shape[0]
-        num_physical_experts = expanded_physical_to_logical.shape[1]
         eplb_model_state = eplb_state.model_states[model_config.compute_hash()]
         eplb_model_state.physical_to_logical_map.copy_(expanded_physical_to_logical)
 
-        logical_to_physical_map = torch.full(
-            (
-                num_moe_layers,
-                model.num_logical_experts,
-                eplb_model_state.logical_to_physical_map.shape[2],
-            ),
-            -1,
-            dtype=torch.int64,
+        (logical_to_physical_map_cpu, logical_replica_count_cpu) = compute_logical_maps(
+            expanded_physical_to_logical.cpu(), model.num_logical_experts
         )
-        logical_replica_count = torch.zeros(
-            (num_moe_layers, model.num_logical_experts),
-            dtype=torch.int64,
-        )
-        expanded_physical_to_logical_numpy = expanded_physical_to_logical.cpu().numpy()
-        for layer_idx in range(num_moe_layers):
-            for phys_idx in range(num_physical_experts):
-                logical_idx = expanded_physical_to_logical_numpy[layer_idx, phys_idx]
-                if logical_idx >= 0:
-                    replica_idx = logical_replica_count[layer_idx, logical_idx]
-                    logical_to_physical_map[layer_idx, logical_idx, replica_idx] = (
-                        phys_idx
-                    )
-                    logical_replica_count[layer_idx, logical_idx] += 1
 
-        logical_to_physical_map = logical_to_physical_map.to(device)
-        logical_replica_count = logical_replica_count.to(device)
+        max_num_replicas = eplb_model_state.logical_to_physical_map.shape[-1]
+        num_replicas = logical_to_physical_map_cpu.shape[-1]
+        logical_to_physical_map = torch.nn.functional.pad(
+            logical_to_physical_map_cpu,
+            (
+                0,
+                max_num_replicas - num_replicas,
+            ),
+            value=-1,
+        ).to(device)
+        logical_replica_count = logical_replica_count_cpu.to(device)
+
         eplb_model_state.logical_to_physical_map.copy_(logical_to_physical_map)
         eplb_model_state.logical_replica_count.copy_(logical_replica_count)
+
         return eplb_state
 
 
@@ -1132,3 +1107,82 @@ def _node_count_with_rank_mapping(
                 node_assignment[other_rank] = next_node_id
 
     return next_node_id
+
+
+def compute_logical_maps(
+    physical_to_logical_map: torch.Tensor,
+    num_logical_experts: int,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """
+    Derive logical_to_physical_map and logical_replica_count from
+    physical_to_logical_map.
+
+    Args:
+        physical_to_logical_map: [num_layers, num_physical_experts], logical
+            expert index for each physical expert slot
+        num_logical_experts: total number of logical experts
+
+    Returns:
+        logical_to_physical_map: [num_layers, num_logical_experts, max_replicas],
+            physical slots per logical expert; -1 where unused
+        logical_replica_count: [num_layers, num_logical_experts], number of
+            physical replicas per logical expert
+    """
+    device = physical_to_logical_map.device
+    assert physical_to_logical_map.device.type == "cpu"
+
+    dtype = physical_to_logical_map.dtype
+
+    # If computing maps for a single layer, unsqueeze a single element layer dimension
+    per_layer = physical_to_logical_map.dim() == 1
+    physical_to_logical_map_view = physical_to_logical_map
+    if per_layer:
+        physical_to_logical_map_view = physical_to_logical_map.unsqueeze(0)
+    assert len(physical_to_logical_map_view.shape) == 2
+    num_layers, num_physical = physical_to_logical_map_view.shape
+
+    valid_mask = physical_to_logical_map_view >= 0
+    logical_replica_count = torch.zeros(
+        num_layers,
+        num_logical_experts,
+        dtype=dtype,
+        device=device,
+    )
+    logical_replica_count.scatter_add_(
+        1,
+        physical_to_logical_map_view.clamp(min=0),
+        valid_mask.to(dtype),
+    )
+
+    max_replicas = int(logical_replica_count.max().item())
+    logical_to_physical_map_out = torch.full(
+        (num_layers, num_logical_experts, max_replicas),
+        -1,
+        dtype=dtype,
+        device=device,
+    )
+
+    running_count = torch.zeros_like(logical_replica_count)
+    layer_indices = torch.arange(num_layers, device=device)
+    for phys_idx in range(num_physical):
+        # Logical expert at physical slot phys_idx for each layer
+        logical_expert_ids = physical_to_logical_map_view[:, phys_idx]  # [num_layers]
+
+        # Scale up will set the logical expert ids to -1 for all new physical experts.
+        # Only consider "valid" experts when setting up the logical_to_physical map.
+        valid_expert_mask = logical_expert_ids >= 0
+        if not valid_expert_mask.any():
+            continue
+        valid_layers = layer_indices[valid_expert_mask]
+        valid_experts = logical_expert_ids[valid_expert_mask]
+
+        # Use the current running count as the replica index, then increment it.
+        replica_idx = running_count[valid_layers, valid_experts]
+        logical_to_physical_map_out[valid_layers, valid_experts, replica_idx] = phys_idx
+        running_count[valid_layers, valid_experts] += 1
+
+    # If computing maps for a single layer, squeeze out the extra layer dimension
+    # before returning
+    if per_layer:
+        return logical_to_physical_map_out.squeeze(0), logical_replica_count.squeeze(0)
+    return logical_to_physical_map_out, logical_replica_count

vllm/distributed/eplb/policy/abstract.py

@@ -17,7 +17,7 @@ class AbstractEplbPolicy(ABC):
         num_nodes: int,
         num_ranks: int,
         old_global_expert_indices: torch.Tensor | None = None,
-    ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    ) -> torch.Tensor:
         """
         Entry point for expert-parallelism load balancer.
 
@@ -35,9 +35,5 @@ class AbstractEplbPolicy(ABC):
         Returns:
             physical_to_logical_map: [layers, num_replicas], the expert
                 index of each replica
-            logical_to_physical_map: [layers, num_logical_experts, X],
-                the replica indices for each expert
-            expert_count: [layers, num_logical_experts], number of
-                physical replicas for each logical expert
         """
         raise NotImplementedError

vllm/distributed/eplb/policy/default.py

@@ -75,7 +75,7 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
     @classmethod
     def replicate_experts(
         cls, weight: np.ndarray, num_phy: int
-    ) -> tuple[np.ndarray, np.ndarray, np.ndarray]:
+    ) -> tuple[np.ndarray, np.ndarray]:
         """
         Replicate `num_log` experts to `num_phy` replicas, such that the maximum
         load of all replicas is minimized.
@@ -86,22 +86,19 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
 
         Returns:
             phy2log: [X, num_phy], logical expert id of each physical expert
-            replica_idx: [X, num_phy], the index of the replica for each logical expert
             logcnt: [X, num_log], number of replicas for each logical expert
         """
         n, num_log = weight.shape
         num_redundant = num_phy - num_log
         assert num_redundant >= 0
         phy2log = np.tile(np.arange(num_phy, dtype=np.int64), (n, 1))
-        replica_idx = np.zeros((n, num_phy), dtype=np.int64)
         logcnt = np.ones((n, num_log), dtype=np.int64)
         arangen = np.arange(n, dtype=np.int64)
         for i in range(num_log, num_phy):
             redundant_indices = np.argmax(weight / logcnt, axis=-1)
             phy2log[:, i] = redundant_indices
-            replica_idx[:, i] = logcnt[arangen, redundant_indices]
             logcnt[arangen, redundant_indices] += 1
-        return phy2log, replica_idx, logcnt
+        return phy2log, logcnt
 
     @classmethod
     def rebalance_experts_hierarchical(
@@ -111,7 +108,7 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         num_groups: int,
         num_nodes: int,
         num_gpus: int,
-    ) -> tuple[np.ndarray, np.ndarray, np.ndarray]:
+    ) -> np.ndarray:
         """
         Parameters:
             weight: [num_moe_layers, num_logical_experts]
@@ -124,10 +121,6 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         Returns:
             phy2log: [layers, num_replicas], the expert
                 index of each replica
-            pphy_replicas_idx: [layers, num_logical_experts, X],
-                the replica indices for each expert
-            logcnt: [layers, num_logical_experts], number of
-                physical replicas for each logical expert
         """
         num_layers, num_logical_experts = weight.shape
         assert num_logical_experts % num_groups == 0
@@ -167,7 +160,7 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         tokens_per_mlog = np.take_along_axis(weight, mlog2log, axis=1).reshape(
             -1, num_logical_experts // num_nodes
         )
-        phy2mlog, replicas_idx, mlogcnt = cls.replicate_experts(
+        phy2mlog, mlogcnt = cls.replicate_experts(
             tokens_per_mlog, num_physical_experts // num_nodes
         )
 
@@ -193,22 +186,15 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         ).reshape(num_layers, -1)
         # Map node-local logical indices back to global logical expert ids.
         pphy2log = np.take_along_axis(mlog2log, pphy2mlog, axis=1)
-        # Reorder replica ranks to the post-packing physical ordering.
-        pphy_replicas_idx = np.take_along_axis(replicas_idx, pphy2phy, axis=1).reshape(
-            num_layers, -1
-        )
-        # Convert replica counts back to the original logical ordering.
-        logcnt = np.take_along_axis(mlogcnt.reshape(num_layers, -1), log2mlog, axis=1)
-        return pphy2log, pphy_replicas_idx, logcnt
+        return pphy2log
 
     @classmethod
     def preserve_intragpu_slots(
         cls,
         phy2log: np.ndarray,
-        phy_replicas_idx: np.ndarray,
         num_ranks: int,
         old_phy2log: np.ndarray,
-    ) -> tuple[np.ndarray, np.ndarray]:
+    ) -> np.ndarray:
         """
         Reorder the new mapping per GPU so that experts that remain on the same GPU
         keep their previous slot positions when possible. Incoming experts to that GPU
@@ -218,14 +204,13 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         """
         num_phy_experts = phy2log.shape[1]
         if num_ranks <= 0 or num_phy_experts % num_ranks != 0:
-            return phy2log, phy_replicas_idx
+            return phy2log
 
         # Move to CPU and convert to NumPy for processing
         slots_per_gpu = num_phy_experts // num_ranks
         num_layers = phy2log.shape[0]
 
         post_phy2log = phy2log.copy()
-        post_phy_replicas_idx = phy_replicas_idx.copy()
 
         for gpu_idx in range(num_ranks):
             start = gpu_idx * slots_per_gpu
@@ -233,7 +218,6 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
             # Experts across all layers for this GPU
             old_local = old_phy2log[:, start:end]  # [layers, slots]
             new_local = phy2log[:, start:end]  # [layers, slots]
-            new_ridx = phy_replicas_idx[:, start:end]  # [layers, slots]
 
             used_new_indices = np.zeros((num_layers, slots_per_gpu), dtype=bool)
             preserved_positions = np.zeros((num_layers, slots_per_gpu), dtype=bool)
@@ -253,9 +237,6 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
                     post_phy2log[layer_indices, start + slot_idx] = new_local[
                         layer_indices, matched_new_positions
                     ]
-                    post_phy_replicas_idx[layer_indices, start + slot_idx] = new_ridx[
-                        layer_indices, matched_new_positions
-                    ]
                     used_new_indices[layer_indices, matched_new_positions] = True
                     preserved_positions[layer_indices, slot_idx] = True
 
@@ -287,11 +268,8 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
                     post_phy2log[layer_idx, start + dst_pos] = new_local[
                         layer_idx, src_pos
                     ]
-                    post_phy_replicas_idx[layer_idx, start + dst_pos] = new_ridx[
-                        layer_idx, src_pos
-                    ]
 
-        return post_phy2log, post_phy_replicas_idx
+        return post_phy2log
 
     @classmethod
     def rebalance_experts(
@@ -302,7 +280,7 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         num_nodes: int,
         num_ranks: int,
         old_global_expert_indices: torch.Tensor | None = None,
-    ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    ) -> torch.Tensor:
         """
         Entry point for expert-parallelism load balancer.
 
@@ -321,13 +299,7 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         Returns:
             phy2log: [layers, num_replicas], the expert
                 index of each replica
-            log2phy: [layers, num_logical_experts, X],
-                the replica indices for each expert
-            logcnt: [layers, num_logical_experts], number of
-                physical replicas for each logical expert
         """
-        device = weight.device
-        num_layers, num_logical_experts = weight.shape
         weight_np = weight.float().cpu().numpy()
         old_phy2log_np = (
             old_global_expert_indices.cpu().numpy()
@@ -337,17 +309,13 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
 
         if num_groups % num_nodes == 0:
             # use hierarchical load-balance policy
-            phy2log_np, phy_replicas_idx_np, logcnt_np = (
-                cls.rebalance_experts_hierarchical(
-                    weight_np, num_replicas, num_groups, num_nodes, num_ranks
-                )
+            phy2log_np = cls.rebalance_experts_hierarchical(
+                weight_np, num_replicas, num_groups, num_nodes, num_ranks
             )
         else:
             # use global load-balance policy
-            phy2log_np, phy_replicas_idx_np, logcnt_np = (
-                cls.rebalance_experts_hierarchical(
-                    weight_np, num_replicas, 1, 1, num_ranks
-                )
+            phy2log_np = cls.rebalance_experts_hierarchical(
+                weight_np, num_replicas, 1, 1, num_ranks
             )
 
         # Optional postprocessing to preserve slots for experts moving
@@ -355,22 +323,10 @@ class DefaultEplbPolicy(AbstractEplbPolicy):
         # Only apply when the number of GPUs and slots per GPU remain unchanged.
         # Helps to avoid unnecessary weight copying when experts move
         # within the same GPU.
-        if old_global_expert_indices is not None:
-            phy2log_np, phy_replicas_idx_np = cls.preserve_intragpu_slots(
-                phy2log_np, phy_replicas_idx_np, num_ranks, old_phy2log_np
+        if old_phy2log_np is not None:
+            phy2log_np = cls.preserve_intragpu_slots(
+                phy2log_np, num_ranks, old_phy2log_np
             )
-        num_redundant_experts = num_replicas - num_logical_experts
-        maxlogcnt = num_redundant_experts + 1
-        log2phy_np = np.full(
-            (num_layers, num_logical_experts, maxlogcnt), -1, dtype=np.int64
-        )
-        layer_indices = np.arange(num_layers)[:, None]
-        replica_indices = np.tile(
-            np.arange(num_replicas, dtype=np.int64), (num_layers, 1)
-        )
-        log2phy_np[layer_indices, phy2log_np, phy_replicas_idx_np] = replica_indices
 
-        phy2log = torch.from_numpy(phy2log_np).to(device)
-        log2phy = torch.from_numpy(log2phy_np).to(device)
-        logcnt = torch.from_numpy(logcnt_np).to(device)
-        return phy2log, log2phy, logcnt
+        phy2log = torch.from_numpy(phy2log_np)
+        return phy2log

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

@@ -25,8 +25,8 @@ The class provides the following primitives:
 
     Worker-side: runs in each worker, loads/saves KV cache to/from
     the Connector based on the metadata.
-        handle_preemptions() - called if there are preempted requests,
-            before their blocks are overwritten
+        handle_preemptions() - called for handling preempted requests
+            or request evicted blocks before they are overwritten
 
         start_load_kv() - starts loading all KVs (maybe async)
         wait_for_layer_load() - blocks until layer i load is done
@@ -288,9 +288,9 @@ class KVConnectorBase_V1(ABC):
         """
         return
 
-    def handle_preemptions(self, preempted_req_ids: set[str]):
+    def handle_preemptions(self, kv_connector_metadata: KVConnectorMetadata):
         """
-        Handle preempted requests BEFORE their blocks are overwritten.
+        Handle preempted requests or evicted blocks BEFORE they are overwritten.
         Needed for connectors which use async saves (e.g., OffloadingConnector)
         """
         return

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

@@ -185,7 +185,7 @@ class ExampleConnector(KVConnectorBase_V1):
                 if kv_cache_attr is None:
                     continue
 
-                kv_cache_layer = kv_cache_attr[forward_context.virtual_engine]
+                kv_cache_layer = kv_cache_attr[0]
 
                 filename = self._generate_filename_debug(
                     layer_name, request.token_ids, request.mm_hashes

vllm/distributed/kv_transfer/kv_connector/v1/lmcache_integration/vllm_v1_adapter.py

@@ -778,9 +778,7 @@ class LMCacheConnectorV1Impl:
                 continue
 
             if layer_name not in self.kv_caches:
-                self.kv_caches[layer_name] = attn_layer.kv_cache[
-                    forward_context.virtual_engine
-                ]
+                self.kv_caches[layer_name] = attn_layer.kv_cache[0]
 
     ####################
     # Worker side APIs

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

@@ -315,10 +315,11 @@ class MultiConnector(KVConnectorBase_V1):
         for c in self._connectors:
             c.set_host_xfer_buffer_ops(copy_operation)
 
-    def handle_preemptions(self, preempted_req_ids: set[str]):
+    def handle_preemptions(self, kv_connector_metadata: KVConnectorMetadata):
         """Handle preempted requests for all sub-connectors."""
-        for c in self._connectors:
-            c.handle_preemptions(preempted_req_ids)
+        assert isinstance(kv_connector_metadata, MultiKVConnectorMetadata)
+        for c, cm in zip(self._connectors, kv_connector_metadata.metadata):
+            c.handle_preemptions(cm)
 
     def get_finished_count(self) -> int | None:
         # TODO(https://github.com/vllm-project/vllm/issues/33400)

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

@@ -1318,12 +1318,12 @@ class NixlConnectorWorker:
                         f"Expected {expected_engine_id},"
                         f"received {metadata.engine_id}."
                     )
-                setup_agent_time = time.perf_counter()
 
                 # Register Remote agent.
                 remote_agent_name = self.add_remote_agent(
                     metadata, remote_rank, remote_tp_size
                 )
+                setup_agent_time = time.perf_counter()
                 logger.debug(
                     "NIXL handshake: add agent took: %s",
                     setup_agent_time - got_metadata_time,

vllm/distributed/kv_transfer/kv_connector/v1/offloading/common.py

@@ -0,0 +1,15 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+from dataclasses import dataclass
+
+from vllm.distributed.kv_transfer.kv_connector.v1.base import KVConnectorMetadata
+from vllm.v1.kv_offload.worker.worker import TransferSpec
+
+ReqId = str
+
+
+@dataclass
+class OffloadingConnectorMetadata(KVConnectorMetadata):
+    reqs_to_load: dict[ReqId, TransferSpec]
+    reqs_to_store: dict[ReqId, TransferSpec]
+    reqs_to_flush: set[str] | None = None