[Core] Enable extensible KV cache for all attention backends and Mamba

Port of internal D110967544. The extensible KV cache flow (reserve KV virtual address space up front, capture CUDA graphs first, then size and commit the KV cache from post-capture free memory) previously required a block-major attention backend and rejected Mamba models. This enables it for every backend layout and for Mamba / linear attention: - ExtensibleTensor gains num_segments: the reservation is divided into equal segments that grow in lockstep, with committed bytes forming a prefix of each segment. Physical pages are mapped at allocation-granularity granules and deduped across overlapping ranges, so a granule straddling a segment boundary is mapped exactly once. resize_per_segment_(bytes, zero_new=True) zeroes only the newly committed logical range of each segment. - Each KV cache buffer keeps its layers' physical layout and is committed as one prefix per layout segment. The segment count is derived from the backend's get_kv_cache_shape / get_kv_cache_block_dim / stride order: K/V-split layouts (e.g. FlashAttention) get one prefix per half, block-major layouts (e.g. FlashInfer, MLA) a single prefix. Mamba state pages are block-major per layer, and hybrid-model attention caches are re-strided to block-major, so both use a single segment. - Removed the supports_extensible_kv_cache gate plumbing from EngineCore, Executor, Worker, WorkerBase and GPUModelRunner; a CUDA platform check remains in EngineCore. - enable_extensible_kv_cache is reported as unsupported by the V2 model runner so V2-default models fall back to the V1 runner (which implements the flow); also fixed initialize_kv_cache being called with the extensible kwarg on runners that do not accept it, which broke every default V2-runner boot on this branch. Tested on H100: - tests/utils_/test_extensible_tensor.py (5 passed, incl. new segmented lockstep-grow/zero, granule-dedup and invalid-usage tests) - tests/v1/worker/test_extensible_kv_cache.py (new, 6 passed: segment derivation, split grows both halves, block-major, legacy full commit, Mamba per-layer growth, hybrid attention+Mamba) - E2E Qwen3-0.6B greedy with VLLM_ATTENTION_BACKEND=FLASH_ATTN (a K/V-split backend the old gate rejected): extensible generations byte-identical to the legacy path; log shows reserve then "Extended KV cache to 34663 blocks". V2->V1 auto-fallback path verified as well.
add new files
2026-07-10 17:08:34 -07:00 · 2026-07-01 17:32:19 -07:00 · 2026-07-01 17:30:53 -07:00 · 2026-07-01 12:43:00 -07:00 · 2026-07-01 12:36:48 -07:00 · 2026-07-01 19:17:05 +00:00
323 changed files with 12334 additions and 6533 deletions
@@ -81,7 +81,7 @@ steps:
      'cd tests &&
      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
      set -o pipefail &&
-      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-3-43264-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype1-1-2049-64-128-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]"'
+      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-3-43264-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype1-1-2049-64-128-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]" --deselect="tests/lora/test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-32000-32-4-4]" --deselect="tests/lora/test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-32000-32-4-4]"'

 - label: LoRA Punica FP8/XPU Ops
  timeout_in_minutes: 45
@@ -103,6 +103,31 @@ steps:
      pytest -v -s v1/kv_offload &&
      pytest -v -s v1/kv_connector/unit/test_offloading_connector.py'

+- label: NixlConnector PD accuracy (2 GPUs)
+  timeout_in_minutes: 60
+  num_devices: 2
+  device: intel_gpu
+  agent_tags:
+    label: production
+    gpu: 2+
+    mem: 24+
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/v1/worker/kv_connector_model_runner_mixin.py
+    - tests/v1/kv_connector/nixl_integration/
+    - vllm/platforms/xpu.py
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'cd tests &&
+      bash v1/kv_connector/nixl_integration/run_xpu_disagg_accuracy_test.sh'
+
 - label: Regression
  key: regression
  timeout_in_minutes: 30
@@ -22,7 +22,7 @@ steps:
  commands:
    - >-
      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      'pip install av &&
      cd tests &&
      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2" &&
      pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model'
@@ -47,8 +47,7 @@ steps:
  commands:
    - >-
      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
-      cd tests &&
+      'cd tests &&
      pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model'

 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
@@ -71,8 +70,7 @@ steps:
  commands:
    - >-
      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
-      cd tests &&
+      'cd tests &&
      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma" &&
      pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model'

@@ -96,7 +94,7 @@ steps:
  commands:
    - >-
      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      'pip install av &&
      cd tests &&
      pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing'

@@ -121,7 +119,7 @@ steps:
  commands:
    - >-
      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install av matplotlib ftfy git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      'pip install av matplotlib ftfy &&
      pip install open-clip-torch --no-deps &&
      cd tests &&
      pytest -v -s models/multimodal/processing/test_tensor_schema.py
@@ -68,7 +68,6 @@ steps:
      - >-
        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
        'cd tests &&
-        bash v1/kv_connector/nixl_integration/run_xpu_disagg_accuracy_test.sh &&
        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 --ignore=v1/engine/test_output_processor.py &&
        pytest -v -s v1/sample --ignore=v1/sample/test_logprobs.py --ignore=v1/sample/test_logprobs_e2e.py -k "not test_topk_only and not test_topp_only and not test_topk_and_topp" &&
@@ -534,6 +534,20 @@ else
  echo "--- Single-node job"
  echo "Render devices: $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES"

+  ulimit_core_hard=$(ulimit -H -c)
+  if [[ "$ulimit_core_hard" == "unlimited" ]]; then
+    # docker run can't pass "unlimited" to --ulimit
+    ulimit_core_hard="-1"
+  fi
+   # Disable core dumps in the ROCm test container unless the ROCm debug agent is enabled
+  coredump_flags="--ulimit core=0:$ulimit_core_hard"
+  if [[ "$commands" == *"ROCm debug agent enabled"* ]]; then
+    # Works around https://github.com/rocm/rocm-systems/issues/6206
+    coredump_flags='-e HSA_COREDUMP_PATTERN="/tmp/gpucore.%p"'
+  else
+    echo "ROCm debug agent not enabled, coredumps are disabled in the test container."
+  fi
+
  docker run \
    --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
    $RDMA_FLAGS \
@@ -541,6 +555,7 @@ else
    --shm-size=16gb \
    --group-add "$render_gid" \
    --rm \
+    $coredump_flags \
    -e HF_TOKEN \
    -e "HF_HUB_DOWNLOAD_TIMEOUT=${HF_HUB_DOWNLOAD_TIMEOUT}" \
    -e "HF_HUB_ETAG_TIMEOUT=${HF_HUB_ETAG_TIMEOUT}" \
@@ -114,26 +114,6 @@ steps:

 #----------------------------------------------------------  mi250 · compile  ----------------------------------------------------------#

- label: PyTorch Compilation Unit Tests # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
-  agent_pool: mi250_1
-  torch_nightly: true
-  optional: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/compilation/
-  - vllm/model_executor/layers/
-  - vllm/v1/worker/
-  - vllm/v1/attention/
-  - vllm/v1/cudagraph_dispatcher.py
-  - vllm/config/compilation.py
-  - csrc/
-  - tests/compile
-  - vllm/platforms/rocm.py
-  commands:
-  - "find compile/ -maxdepth 1 -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
-
 - label: PyTorch Fullgraph Smoke Test # TBD
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -259,7 +239,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
  - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model

@@ -438,7 +417,7 @@ steps:
 #-----------------------------------------------------  mi300 · basic_correctness  -----------------------------------------------------#

 - label: Basic Correctness # TBD
-  timeout_in_minutes: 50
+  timeout_in_minutes: 95
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
  fast_check: true
@@ -456,7 +435,7 @@ steps:
  - pytest -v -s basic_correctness/test_cpu_offload.py

 - label: Distributed Model Tests (2 GPUs) # TBD
-  timeout_in_minutes: 65
+  timeout_in_minutes: 110
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_2
  num_gpus: 2
@@ -498,6 +477,26 @@ steps:

 #----------------------------------------------------------  mi300 · compile  ----------------------------------------------------------#

+- label: PyTorch Compilation Unit Tests # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  torch_nightly: true
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/compilation/
+  - vllm/model_executor/layers/
+  - vllm/v1/worker/
+  - vllm/v1/attention/
+  - vllm/v1/cudagraph_dispatcher.py
+  - vllm/config/compilation.py
+  - csrc/
+  - tests/compile
+  - vllm/platforms/rocm.py
+  commands:
+  - "find compile/ -maxdepth 1 -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
+
 - label: Fusion E2E Config Sweep (H100-MI300) # TBD
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -678,7 +677,7 @@ steps:
  - pytest -v -s distributed/test_eplb_spec_decode.py

 - label: Distributed Tests (2xH100-2xMI300) # TBD
-  timeout_in_minutes: 30
+  timeout_in_minutes: 75
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_2
  num_gpus: 2
@@ -1222,7 +1221,7 @@ steps:
 #---------------------------------------------------------  mi300 · examples  ----------------------------------------------------------#

 - label: Examples # TBD
-  timeout_in_minutes: 45
+  timeout_in_minutes: 90
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
  optional: true
@@ -1258,7 +1257,7 @@ steps:
 #----------------------------------------------------------  mi300 · kernels  ----------------------------------------------------------#

 - label: Kernels Attention Test %N # TBD
-  timeout_in_minutes: 55
+  timeout_in_minutes: 100
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
  optional: true
@@ -1292,7 +1291,7 @@ steps:
  - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py  kernels/test_concat_mla_q.py kernels/test_top_k_per_row.py

 - label: Kernels MoE Test %N # TBD
-  timeout_in_minutes: 50
+  timeout_in_minutes: 95
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
  optional: true
@@ -1374,8 +1373,6 @@ steps:
  - tests/lora
  - vllm/platforms/rocm.py
  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
  - pytest -v -s -x lora/test_chatglm3_tp.py
  - pytest -v -s -x lora/test_llama_tp.py
  - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
@@ -1439,7 +1436,7 @@ steps:
  - pytest -v -s models/test_initialization.py::test_can_initialize_small_subset

 - label: Basic Models Tests (Other) # TBD
-  timeout_in_minutes: 45
+  timeout_in_minutes: 90
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
  optional: true
@@ -1516,7 +1513,6 @@ steps:
  - tests/models/multimodal/generation
  - tests/models/multimodal/test_mapping.py
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
  - pytest -v -s models/multimodal/test_mapping.py

@@ -1530,7 +1526,6 @@ steps:
  - vllm/
  - tests/models/multimodal/generation
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'


@@ -1544,7 +1539,6 @@ steps:
  - vllm/
  - tests/models/multimodal/generation
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'

 - label: "Multi-Modal Models (Standard) 1: qwen2" # TBD
@@ -1558,7 +1552,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
  - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model

@@ -1574,7 +1567,6 @@ steps:
  - tests/models/multimodal/generation
  - tests/models/multimodal/test_mapping.py
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
  - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model

@@ -1589,7 +1581,6 @@ steps:
  - tests/models/multimodal/generation
  - tests/models/multimodal/test_mapping.py
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing
  - pytest -v -s models/multimodal/generation/test_memory_leak.py -m core_model
  - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model
@@ -1605,7 +1596,6 @@ steps:
  - tests/models/multimodal
  - tests/models/registry.py
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/processing/test_tensor_schema.py

 - label: Multi-Modal Processor (CPU) %N # TBD
@@ -1621,7 +1611,6 @@ steps:
  - tests/models/multimodal
  - tests/models/registry.py
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB

 #-----------------------------------------------------  mi300 · models / quantized  -----------------------------------------------------#
@@ -1903,7 +1892,7 @@ steps:
  - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"

 - label: Spec Decode Eagle # TBD
-  timeout_in_minutes: 45
+  timeout_in_minutes: 90
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
  optional: true
@@ -2119,7 +2108,7 @@ steps:
  - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py

 - label: Metrics, Tracing (2 GPUs) # TBD
-  timeout_in_minutes: 20
+  timeout_in_minutes: 65
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_2
  optional: true
@@ -2272,7 +2261,7 @@ steps:
 #------------------------------------------------------  mi300 · weight_loading  -------------------------------------------------------#

 - label: Weight Loading Multiple GPU # TBD
-  timeout_in_minutes: 30
+  timeout_in_minutes: 75
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_2
  num_gpus: 2
@@ -2284,7 +2273,7 @@ steps:
  - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt

 - label: Weight Loading Multiple GPU - Large Models # TBD
-  timeout_in_minutes: 30
+  timeout_in_minutes: 75
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_2
  num_gpus: 2
@@ -2469,7 +2458,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
  - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model

@@ -2831,7 +2819,7 @@ steps:
  - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py

 - label: Kernels Attention Test %N # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 100
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
  agent_pool: mi355_1
  parallelism: 2
@@ -2980,7 +2968,6 @@ steps:
  - tests/models/multimodal/generation
  - tests/models/multimodal/test_mapping.py
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
  - pytest -v -s models/multimodal/test_mapping.py

@@ -2994,7 +2981,6 @@ steps:
  - vllm/
  - tests/models/multimodal/generation
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'

 - label: Multi-Modal Models (Extended Pooling) # TBD
@@ -3020,7 +3006,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
  - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model

@@ -3035,7 +3020,6 @@ steps:
  - vllm/
  - tests/models/multimodal/generation
  commands:
-  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
  - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing
  - pytest -v -s models/multimodal/generation/test_memory_leak.py -m core_model
  - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model
@@ -3178,7 +3162,7 @@ steps:
 #------------------------------------------------------  mi355 · weight_loading  -------------------------------------------------------#

 - label: Weight Loading Multiple GPU # TBD
-  timeout_in_minutes: 30
+  timeout_in_minutes: 75
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
  agent_pool: mi355_2
  num_gpus: 2
@@ -3190,7 +3174,7 @@ steps:
  - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt

 - label: Weight Loading Multiple GPU - Large Models # TBD
-  timeout_in_minutes: 30
+  timeout_in_minutes: 75
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
  agent_pool: mi355_2
  working_dir: "/vllm-workspace/tests"
@@ -233,7 +233,7 @@ steps:
  num_devices: 2
  commands:
    - pytest -v -s tests/distributed/test_context_parallel.py
-    - pytest -v -s tests/distributed/test_nccl_symm_mem_allreduce.py
+    - pytest -v -s tests/distributed/test_nccl_symm_mem.py
    - pytest -v -s tests/v1/distributed/test_dbo.py
    - pytest -v -s tests/distributed/test_mnnvl_alltoall.py

@@ -54,8 +54,8 @@ steps:
    - export VLLM_USE_DEEP_GEMM=0  # We found Triton is faster than DeepGEMM for H100
    - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-hopper.txt --tp-size=4

- label: LM Eval Small Models (2xB200)
-  key: lm-eval-small-models-2xb200
+- label: LM Eval Small Models (1xB200)
+  key: lm-eval-small-models-1xb200
  timeout_in_minutes: 120
  device: b200-k8s
  optional: true
@@ -65,9 +65,10 @@ steps:
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt

- label: LM Eval Small Models (2xL4)
-  key: lm-eval-small-models-tp
-  timeout_in_minutes: 10
+- label: LM Eval Small Models Distributed (2xB200)
+  key: lm-eval-small-models-distributed-2xb200
+  timeout_in_minutes: 120
+  device: b200-k8s
  num_devices: 2
  optional: true
  source_file_dependencies:
@@ -10,7 +10,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
    - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
  mirror:
@@ -27,7 +26,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
    - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
  mirror:
@@ -44,7 +42,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
    - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
  mirror:
@@ -61,7 +58,6 @@ steps:
  - vllm/
  - tests/models/multimodal
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
    - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
    - pytest models/multimodal/generation/test_memory_leak.py -m core_model
@@ -83,7 +79,6 @@ steps:
  - tests/models/registry.py
  device: cpu-medium
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py

 - label: Multi-Modal Processor # 44min
@@ -95,7 +90,6 @@ steps:
  - tests/models/multimodal
  - tests/models/registry.py
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/processing/test_tensor_schema.py

 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
@@ -129,7 +123,6 @@ steps:
  - tests/models/multimodal/generation
  - tests/models/multimodal/test_mapping.py
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
    - pytest -v -s models/multimodal/test_mapping.py
  mirror:
@@ -146,7 +139,6 @@ steps:
  - vllm/
  - tests/models/multimodal/generation
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'

 - label: Multi-Modal Models (Extended Generation 3)
@@ -157,7 +149,6 @@ steps:
  - vllm/
  - tests/models/multimodal/generation
  commands:
-    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'

 - label: Multi-Modal Models (Extended Pooling)
@@ -132,10 +132,8 @@ def benchmark_function(
    reset_memory_stats()

    # Benchmark
-    start_events = [
-        torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)
-    ]
-    end_events = [torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)]
+    start_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
+    end_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]

    for i in range(benchmark_iters):
        logits_copy = logits.clone()
@@ -134,8 +134,8 @@ def benchmark_config(
    torch.accelerator.synchronize()

    # Benchmark
-    start = torch.cuda.Event(enable_timing=True)
-    end = torch.cuda.Event(enable_timing=True)
+    start = torch.Event(enable_timing=True)
+    end = torch.Event(enable_timing=True)
    start.record()
    for _ in range(num_iters):
        with override_config(config):
@@ -170,8 +170,8 @@ def benchmark_config(
                graph.replay()
            torch.accelerator.synchronize()

-            start = torch.cuda.Event(enable_timing=True)
-            end = torch.cuda.Event(enable_timing=True)
+            start = torch.Event(enable_timing=True)
+            end = torch.Event(enable_timing=True)
            latencies: list[float] = []
            for _ in range(num_iters):
                start.record()
@@ -427,6 +427,7 @@ if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
    set(VLLM_EXT_SRC
        "csrc/cpu/shm.cpp"
        "csrc/cpu/activation_lut_bf16.cpp"
+        "csrc/cpu/cpu_tanhf_neon.hpp"
        "csrc/cpu/cpu_fused_moe.cpp"
        ${VLLM_EXT_SRC})
 endif()
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG b3964b1d8b95d8e8447435668ab169a2700bab65
+          GIT_TAG 2c839c33742309ec41e620bf837495ec9926c56e
          GIT_PROGRESS TRUE
          # Don't share the vllm-flash-attn build between build types
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
@@ -126,6 +126,18 @@ void gelu_tanh_and_mul(torch::Tensor& out,    // [..., d]
      });
 }

+void gelu_tanh(torch::Tensor& out, torch::Tensor& input) {
+  int num_tokens = input.numel() / input.size(-1);
+  int d = input.size(-1);
+
+  VLLM_DISPATCH_FLOATING_TYPES(input.scalar_type(), "gelu_tanh_impl", [&] {
+    CPU_KERNEL_GUARD_IN(gelu_tanh_impl)
+    activation_kernel<scalar_t, gelu_tanh_act, false>(
+        num_tokens, d, input.data_ptr<scalar_t>(), out.data_ptr<scalar_t>());
+    CPU_KERNEL_GUARD_OUT(gelu_tanh_impl)
+  });
+}
+
 void gelu_new(torch::Tensor& out, torch::Tensor& input) {
  int num_tokens = input.numel() / input.size(-1);
  int d = input.size(-1);
@@ -0,0 +1,128 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+#ifndef CPU_TANHF_NEON_HPP
+#define CPU_TANHF_NEON_HPP
+
+#include <cstdint>
+#include <arm_neon.h>
+
+namespace vec_op {
+
+namespace {
+
+struct TanhfConstants {
+  float32x4_t special_bound;
+  float32x4_t two;
+  float32x4_t c0;
+  float32x4_t c2;
+  int32x4_t exponent_bias;
+  float c1;
+  float c3;
+  float two_over_ln2;
+  float c4;
+  float ln2_hi;
+  float ln2_lo;
+};
+
+const TanhfConstants kTanhfConstants = {
+    // 9.01, above which tanhf rounds to 1 (or -1 for  negative).
+    .special_bound = vdupq_n_f32(0x1.205966p+3f),
+    .two = vdupq_n_f32(0x1.0p+1f),
+    .c0 = vdupq_n_f32(0x1.fffffep-2f),
+    .c2 = vdupq_n_f32(0x1.555736p-5f),
+    .exponent_bias = vdupq_n_s32(0x3f800000),
+    .c1 = 0x1.5554aep-3f,
+    .c3 = 0x1.12287cp-7f,
+    .two_over_ln2 = 0x1.715476p+1f,
+    .c4 = 0x1.6b55a2p-10f,
+    .ln2_hi = 0x1.62e4p-1f,
+    .ln2_lo = 0x1.7f7d1cp-20f,
+};
+
+// Return the ptr but hide it's value from the compiler so accesses
+// through it can't be optimised based on contents.
+template <typename T>
+inline const T* ptr_barrier(const T* ptr) {
+  const T* opaque_ptr = ptr;
+  __asm__("" : "+r"(opaque_ptr));
+  return opaque_ptr;
+}
+
+// Check whether any lanes in the mask are set
+inline bool any_u32(uint32x4_t x) { return vmaxvq_u32(x) != 0; }
+
+// e^2x - 1 inline helper
+inline float32x4_t e2xm1f_inline(float32x4_t x, const TanhfConstants* d) {
+  float32x2_t ln2 = vld1_f32(&d->ln2_hi);
+  float32x4_t lane_consts = vld1q_f32(&d->c1);
+
+  // Reduce argument: f in [-ln2/2, ln2/2], i is exact.
+  float32x4_t j = vrndaq_f32(vmulq_laneq_f32(x, lane_consts, 2));
+  int32x4_t i = vcvtq_s32_f32(j);
+  float32x4_t f = vaddq_f32(x, x);
+  f = vfmsq_lane_f32(f, j, ln2, 0);
+  f = vfmsq_lane_f32(f, j, ln2, 1);
+
+  // Approximate expm1(f) with polynomial P, expm1(f) ~= f + f^2 * P(f)
+  float32x4_t f2 = vmulq_f32(f, f);
+  float32x4_t f4 = vmulq_f32(f2, f2);
+  float32x4_t p01 = vfmaq_laneq_f32(d->c0, f, lane_consts, 0);
+  float32x4_t p23 = vfmaq_laneq_f32(d->c2, f, lane_consts, 1);
+  float32x4_t poly = vfmaq_f32(p01, f2, p23);
+  poly = vfmaq_laneq_f32(poly, f4, lane_consts, 3);
+  poly = vfmaq_f32(f, f2, poly);
+
+  // scale = 2^i
+  int32x4_t u = vaddq_s32(vshlq_n_s32(i, 23), d->exponent_bias);
+  float32x4_t scale = vreinterpretq_f32_s32(u);
+  return vfmaq_f32(vsubq_f32(scale, vdupq_n_f32(1.0f)), poly, scale);
+}
+
+// Calculate the result tanh(x) = q / (q+2) and set special lanes to ±1
+inline float32x4_t special_case(float32x4_t x, float32x4_t q,
+                                uint32x4_t special) {
+  const TanhfConstants* d = ptr_barrier(&kTanhfConstants);
+
+  float32x4_t y = vdivq_f32(q, vaddq_f32(q, d->two));
+  uint32x4_t ix = vreinterpretq_u32_f32(x);
+  uint32x4_t one_bits = vreinterpretq_u32_s32(d->exponent_bias);
+  uint32x4_t sign_mask = vdupq_n_u32(0x80000000u);
+  uint32x4_t special_bits = vbslq_u32(sign_mask, ix, one_bits);
+  float32x4_t special_y = vreinterpretq_f32_u32(special_bits);
+  return vbslq_f32(special, special_y, y);
+}
+
+}  // namespace
+
+// Implementation of tanhf adapted from Arm Optimized Routines (tanhf
+// AdvSIMD)
+// https://github.com/ARM-software/optimized-routines/blob/master/math/aarch64/advsimd/tanhf.c
+//
+// Approximation for single-precision vector tanh(x), using a simplified
+// version of expm1f. The maximum error is 2.08 + 0.5 ULP:
+// _ZGVnN4v_tanhf (0x1.fa5eep-5) got 0x1.f9ba02p-5 want 0x1.f9ba08p-5.
+inline float32x4_t fast_tanhf_f32x4(float32x4_t x) {
+  const TanhfConstants* d = ptr_barrier(&kTanhfConstants);
+
+  // tanh(x) = (e^2x - 1) / (e^2x + 1)
+  // q = e^2x -1
+  float32x4_t q = e2xm1f_inline(x, d);
+
+  // Check for special cases
+  uint32x4_t special = vcagtq_f32(x, d->special_bound);
+
+  // Fall back to vectorised special case for any lanes which would cause
+  // expm1 to overflow
+  if (any_u32(special)) {
+    return special_case(x, q, special);
+  }
+
+  // Complete fast path if no special lanes
+  // tanh(x) = q / (q+2)
+  return vdivq_f32(q, vaddq_f32(q, d->two));
+}
+
+}  // namespace vec_op
+
+#endif  // CPU_TANHF_NEON_HPP
@@ -3,6 +3,8 @@

 #include <arm_neon.h>

+#include "cpu/cpu_tanhf_neon.hpp"
+
 #include <torch/all.h>
 #include <ATen/cpu/vec/functional.h>
 #include <ATen/cpu/vec/vec.h>
@@ -345,6 +347,10 @@ struct FP32Vec4 : public VectorizedRegWrapper<FP32Vec4, 1, float> {
  explicit FP32Vec4(float32x4_t data) : Base(VectorizedT(data)) {};

  explicit FP32Vec4(const FP32Vec4& data) : Base(data) {};
+
+  FORCE_INLINE FP32Vec4 tanh() const {
+    return FP32Vec4(fast_tanhf_f32x4(reg.val[0]));
+  }
 };

 struct FP32Vec8 : public VectorizedRegWrapper<FP32Vec8, 2, float> {
@@ -391,6 +397,13 @@ struct FP32Vec8 : public VectorizedRegWrapper<FP32Vec8, 2, float> {
    reg.val[1] = Vectorized<float>(data.val[1]);
  }

+  FORCE_INLINE FP32Vec8 tanh() const {
+    FP32Vec8 r(uninit);
+    r.reg.val[0] = Vectorized<float>(fast_tanhf_f32x4(reg.val[0]));
+    r.reg.val[1] = Vectorized<float>(fast_tanhf_f32x4(reg.val[1]));
+    return r;
+  }
+
  FORCE_INLINE float reduce_sum() const noexcept {
    float answer = 0;
    std::plus<VectorizedT> add;
@@ -497,6 +510,15 @@ struct FP32Vec16 : public VectorizedRegWrapper<FP32Vec16, 4, float> {
    reg.val[3] = Vectorized<float>(vcvt_f32_f16(vget_high_f16(v.reg.val[1])));
  };

+  FORCE_INLINE FP32Vec16 tanh() const {
+    FP32Vec16 r(uninit);
+    r.reg.val[0] = Vectorized<float>(fast_tanhf_f32x4(reg.val[0]));
+    r.reg.val[1] = Vectorized<float>(fast_tanhf_f32x4(reg.val[1]));
+    r.reg.val[2] = Vectorized<float>(fast_tanhf_f32x4(reg.val[2]));
+    r.reg.val[3] = Vectorized<float>(fast_tanhf_f32x4(reg.val[3]));
+    return r;
+  }
+
  static FORCE_INLINE void load_even_odd(const float* ptr, FP32Vec16& even,
                                         FP32Vec16& odd) noexcept {
    const float32x4x2_t x01 = vuzpq_f32(vld1q_f32(ptr), vld1q_f32(ptr + 4));
@@ -298,6 +298,10 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
  ops.def("gelu_tanh_and_mul(Tensor! out, Tensor input) -> ()");
  ops.impl("gelu_tanh_and_mul", torch::kCPU, &gelu_tanh_and_mul);

+  // GELU tanh implementation.
+  ops.def("gelu_tanh(Tensor! out, Tensor input) -> ()");
+  ops.impl("gelu_tanh", torch::kCPU, &gelu_tanh);
+
  // GELU implementation used in GPT-2.
  ops.def("gelu_new(Tensor! out, Tensor input) -> ()");
  ops.impl("gelu_new", torch::kCPU, &gelu_new);
@@ -804,35 +804,6 @@ void minimax_reduce_rms_op(MiniMaxReduceRMSParams const& params) {
 }  // namespace tensorrt_llm
 }  // namespace vllm

-torch::stable::Tensor minimax_allreduce_rms(
-    torch::stable::Tensor const& input,
-    torch::stable::Tensor const& norm_weight, torch::stable::Tensor workspace,
-    int64_t const rank, int64_t const nranks, double const eps) {
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      input.get_device_index());
-  auto allreduce_params = vllm::tensorrt_llm::MiniMaxReduceRMSParams();
-
-  allreduce_params.nranks = static_cast<int>(nranks);
-  allreduce_params.rank = static_cast<int>(rank);
-  allreduce_params.dtype = input.scalar_type();
-  allreduce_params.size_q = static_cast<int>(input.numel());
-  allreduce_params.hidden_dim = static_cast<int>(input.size(-1));
-  allreduce_params.stride_q = allreduce_params.hidden_dim;
-  allreduce_params.workspace =
-      reinterpret_cast<void**>(workspace.mutable_data_ptr());
-  allreduce_params.allreduce_in = const_cast<void*>(input.const_data_ptr());
-  allreduce_params.rms_gamma = const_cast<void*>(norm_weight.const_data_ptr());
-  allreduce_params.rms_eps = static_cast<float>(eps);
-  allreduce_params.stream = get_current_cuda_stream(input.get_device_index());
-
-  torch::stable::Tensor rms_norm_out = torch::stable::empty_like(input);
-  allreduce_params.rms_norm_out = rms_norm_out.mutable_data_ptr();
-
-  vllm::tensorrt_llm::minimax_reduce_rms_op(allreduce_params);
-
-  return rms_norm_out;
-}
-
 std::tuple<torch::stable::Tensor, torch::stable::Tensor>
 minimax_allreduce_rms_qk(torch::stable::Tensor qkv,
                         torch::stable::Tensor const& norm_weight_q,
@@ -288,10 +288,6 @@ void fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert(
    int64_t cache_block_size);

 #ifndef USE_ROCM
-torch::stable::Tensor minimax_allreduce_rms(
-    torch::stable::Tensor const& input,
-    torch::stable::Tensor const& norm_weight, torch::stable::Tensor workspace,
-    int64_t const rank, int64_t const nranks, double const eps);
 std::tuple<torch::stable::Tensor, torch::stable::Tensor>
 minimax_allreduce_rms_qk(torch::stable::Tensor qkv,
                         torch::stable::Tensor const& norm_weight_q,
@@ -449,10 +449,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
      "int cache_block_size) -> ()");

 #ifndef USE_ROCM
-  ops.def(
-      "minimax_allreduce_rms("
-      "Tensor input, Tensor norm_weight, Tensor workspace, "
-      "int rank, int nranks, float eps) -> Tensor");
  ops.def(
      "minimax_allreduce_rms_qk("
      "Tensor qkv, Tensor norm_weight_q, Tensor norm_weight_k, "
@@ -705,7 +701,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
      "fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert",
      TORCH_BOX(&fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert));
 #ifndef USE_ROCM
-  ops.impl("minimax_allreduce_rms", TORCH_BOX(&minimax_allreduce_rms));
  ops.impl("minimax_allreduce_rms_qk", TORCH_BOX(&minimax_allreduce_rms_qk));
 #endif
  ops.impl("fused_minimax_m3_qknorm_rope_kv_insert",
@@ -35,6 +35,8 @@ void gelu_and_mul(torch::Tensor& out, torch::Tensor& input);

 void gelu_tanh_and_mul(torch::Tensor& out, torch::Tensor& input);

+void gelu_tanh(torch::Tensor& out, torch::Tensor& input);
+
 void gelu_new(torch::Tensor& out, torch::Tensor& input);

 void gelu_fast(torch::Tensor& out, torch::Tensor& input);
@@ -222,7 +222,7 @@ MLA decode backends are selected using the standard
 | ------- | ------ | --------- | ----------- | ---------- | ---- | ---------- | ------ | --------- | --- | --------------- | ------------ |
 | `CUTLASS_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 128 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 10.x |
 | `FLASHINFER_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 10.x |
-| `FLASHINFER_MLA_SPARSE` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | 10.x |
+| `FLASHINFER_MLA_SPARSE` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 10.x |
 | `FLASHINFER_MLA_SPARSE_SM120` | bf16 | `auto`, `fp8`, `fp8_e4m3`, `fp8_ds_mla` | 64, 256 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | 12.x |
 | `FLASHMLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 64 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 9.x-10.x |
 | `FLASHMLA_SPARSE` | bf16 | `auto`, `bfloat16`, `fp8_ds_mla` | 64 | 576 | ❌ | ❌ | ✅ | ❌ | ❌ | Decoder | 9.x-10.x |
@@ -405,7 +405,6 @@ th {
 | `Glm4MoeLiteForCausalLM` | GLM-4.7-Flash | `zai-org/GLM-4.7-Flash`, etc. | ✅︎ | ✅︎ |
 | `GlmMoeDsaForCausalLM` | GLM-5, GLM-5.1, GLM-5.2 | `zai-org/GLM-5`, etc. | ✅︎ | ✅︎ |
 | `GPT2LMHeadModel` | GPT-2 | `openai-community/gpt2`, `openai-community/gpt2-xl`, etc. | | ✅︎ |
-| `GPTBigCodeForCausalLM` | StarCoder, SantaCoder, WizardCoder | `bigcode/starcoder`, `bigcode/gpt_bigcode-santacoder`, `WizardLM/WizardCoder-15B-V1.0`, etc. | ✅︎ | ✅︎ |
 | `GPTJForCausalLM` | GPT-J | `EleutherAI/gpt-j-6b`, `nomic-ai/gpt4all-j`, etc. | | ✅︎ |
 | `GPTNeoXForCausalLM` | GPT-NeoX, Pythia, OpenAssistant, Dolly V2, StableLM | `EleutherAI/gpt-neox-20b`, `EleutherAI/pythia-12b`, `OpenAssistant/oasst-sft-4-pythia-12b-epoch-3.5`, `databricks/dolly-v2-12b`, `stabilityai/stablelm-tuned-alpha-7b`, etc. | | ✅︎ |
 | `GptOssForCausalLM` | GPT-OSS | `openai/gpt-oss-120b`, `openai/gpt-oss-20b` | ✅︎ | ✅︎ |
@@ -477,7 +476,6 @@ th {
 | `SolarForCausalLM` | Solar Pro | `upstage/solar-pro-preview-instruct`, etc. | ✅︎ | ✅︎ |
 | `StableLmForCausalLM` | StableLM | `stabilityai/stablelm-3b-4e1t`, `stabilityai/stablelm-base-alpha-7b-v2`, etc. | | |
 | `StableLMEpochForCausalLM` | StableLM Epoch | `stabilityai/stablelm-zephyr-3b`, etc. | | ✅︎ |
-| `Starcoder2ForCausalLM` | Starcoder2 | `bigcode/starcoder2-3b`, `bigcode/starcoder2-7b`, `bigcode/starcoder2-15b`, etc. | | ✅︎ |
 | `Step1ForCausalLM` | Step-Audio | `stepfun-ai/Step-Audio-EditX`, etc. | ✅︎ | ✅︎ |
 | `Step3p5ForCausalLM` | Step-3.5-flash | `stepfun-ai/Step-3.5-Flash`, etc. | | ✅︎ |
 | `TeleChatForCausalLM` | TeleChat | `chuhac/TeleChat2-35B`, etc. | ✅︎ | ✅︎ |
@@ -490,7 +488,9 @@ Some models are supported only via the [Transformers modeling backend](#transfor

 | Architecture | Models | Example HF Models | [LoRA](../features/lora.md) | [PP](../serving/parallelism_scaling.md) |
 | ------------ | ------ | ----------------- | -------------------- | ------------------------- |
+| `GPTBigCodeForCausalLM` | StarCoder, SantaCoder, WizardCoder | `bigcode/starcoder`, `bigcode/gpt_bigcode-santacoder`, `WizardLM/WizardCoder-15B-V1.0`, etc. | ✅︎ | |
 | `SmolLM3ForCausalLM` | SmolLM3 | `HuggingFaceTB/SmolLM3-3B` | ✅︎ | ✅︎ |
+| `Starcoder2ForCausalLM` | Starcoder2 | `bigcode/starcoder2-3b`, `bigcode/starcoder2-7b`, `bigcode/starcoder2-15b`, etc. | ✅︎ | ✅︎ |

 !!! note
    Currently, the ROCm version of vLLM supports Mistral and Mixtral only for context lengths up to 4096.
@@ -532,7 +532,6 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | ------------ | ------ | ------ | ----------------- | -------------------- | ------------------------- |
 | `AriaForConditionalGeneration` | Aria | T + I<sup>+</sup> | `rhymes-ai/Aria` | | |
 | `AudioFlamingo3ForConditionalGeneration` | AudioFlamingo3 | T + A | `nvidia/audio-flamingo-3-hf`, `nvidia/music-flamingo-hf` | ✅︎ | ✅︎ |
-| `AyaVisionForConditionalGeneration` | Aya Vision | T + I<sup>+</sup> | `CohereLabs/aya-vision-8b`, `CohereLabs/aya-vision-32b`, etc. | | ✅︎ |
 | `BagelForConditionalGeneration` | BAGEL | T + I<sup>+</sup> | `ByteDance-Seed/BAGEL-7B-MoT` | ✅︎ | ✅︎ |
 | `BeeForConditionalGeneration` | Bee-8B | T + I<sup>E+</sup> | `Open-Bee/Bee-8B-RL`, `Open-Bee/Bee-8B-SFT` | | ✅︎ |
 | `Blip2ForConditionalGeneration` | BLIP-2 | T + I<sup>E</sup> | `Salesforce/blip2-opt-2.7b`, `Salesforce/blip2-opt-6.7b`, etc. | ✅︎ | ✅︎ |
@@ -579,7 +578,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `Lfm2VlForConditionalGeneration` | LFM2-VL | T + I<sup>+</sup> | `LiquidAI/LFM2-VL-450M`, `LiquidAI/LFM2-VL-3B`, `LiquidAI/LFM2-VL-8B-A1B`, etc. | ✅︎ | ✅︎ |
 | `Llama4ForConditionalGeneration` | Llama 4 | T + I<sup>+</sup> | `meta-llama/Llama-4-Scout-17B-16E-Instruct`, `meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8`, `meta-llama/Llama-4-Maverick-17B-128E-Instruct`, etc. | ✅︎ | ✅︎ |
 | `Llama_Nemotron_Nano_VL` | Llama Nemotron Nano VL | T + I<sup>E+</sup> | `nvidia/Llama-3.1-Nemotron-Nano-VL-8B-V1` | ✅︎ | ✅︎ |
-| `LlavaForConditionalGeneration` | LLaVA-1.5, Pixtral (HF Transformers) | T + I<sup>E+</sup> | `llava-hf/llava-1.5-7b-hf`, `TIGER-Lab/Mantis-8B-siglip-llama3` (see note), `mistral-community/pixtral-12b`, etc. | ✅︎ | ✅︎ |
+| `LlavaForConditionalGeneration` | LLaVA-1.5, Pixtral (HF Transformers) | T + I<sup>E+</sup> | `llava-hf/llava-1.5-7b-hf`, `mistral-community/pixtral-12b`, etc. | ✅︎ | ✅︎ |
 | `LlavaNextForConditionalGeneration` | LLaVA-NeXT, Granite Vision | T + I<sup>E+</sup> | `llava-hf/llava-v1.6-mistral-7b-hf`, `llava-hf/llava-v1.6-vicuna-7b-hf`, `ibm-granite/granite-vision-3.3-2b`, etc. | | ✅︎ |
 | `LlavaNextVideoForConditionalGeneration` | LLaVA-NeXT-Video | T + V | `llava-hf/LLaVA-NeXT-Video-7B-hf`, etc. | | ✅︎ |
 | `LlavaOnevisionForConditionalGeneration` | LLaVA-Onevision | T + I<sup>+</sup> + V<sup>+</sup> | `llava-hf/llava-onevision-qwen2-7b-ov-hf`, `llava-hf/llava-onevision-qwen2-0.5b-ov-hf`, etc. | | ✅︎ |
@@ -594,7 +593,6 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `Molmo2ForConditionalGeneration` | Molmo2 | T + I<sup>+</sup> / V | `allenai/Molmo2-4B`, `allenai/Molmo2-8B`, `allenai/Molmo2-O-7B`, `allenai/MolmoWeb-4B`<sup>^</sup>, `allenai/MolmoWeb-8B`<sup>^</sup> | ✅︎ | ✅︎ |
 | `MossAudioModel` | MOSS-Audio | T + A<sup>+</sup> | `OpenMOSS-Team/MOSS-Audio-4B-Instruct`, `OpenMOSS-Team/MOSS-Audio-4B-Thinking`, `OpenMOSS-Team/MOSS-Audio-8B-Instruct`, `OpenMOSS-Team/MOSS-Audio-8B-Thinking` | ✅︎ | ✅︎ |
 | `Moondream3ForCausalLM` | Moondream3 | T + I | `moondream/moondream3-preview` | | ✅︎ |
-| `MusicFlamingoForConditionalGeneration` | MusicFlamingo | T + A | `nvidia/music-flamingo-2601-hf`, `nvidia/music-flamingo-think-2601-hf` | ✅︎ | ✅︎ |
 | `NVLM_D_Model` | NVLM-D 1.0 | T + I<sup>+</sup> | `nvidia/NVLM-D-72B`, etc. | | ✅︎ |
 | `OpenCUAForConditionalGeneration` | OpenCUA-7B | T + I<sup>E+</sup> | `xlangai/OpenCUA-7B` | ✅︎ | ✅︎ |
 | `OpenPanguVLForConditionalGeneration` | openpangu-VL | T + I<sup>E+</sup> + V<sup>E+</sup> | `FreedomIntelligence/openPangu-VL-7B` | ✅︎ | ✅︎ |
@@ -678,9 +676,6 @@ Some models are supported only via the [Transformers modeling backend](#transfor
    coordinate decoding and are not exposed by this vLLM implementation.
    See [Moondream3 prompt recipes](../features/multimodal_inputs.md#moondream3-prompt-recipes).

-!!! note
-    To use `TIGER-Lab/Mantis-8B-siglip-llama3`, you have to pass `--hf_overrides '{"architectures": ["MantisForConditionalGeneration"]}'` when running vLLM.
-
 !!! note
    The official `openbmb/MiniCPM-V-2` doesn't work yet, so we need to use a fork (`HwwwH/MiniCPM-V-2`) for now.
    For more details, please see: <https://github.com/vllm-project/vllm/pull/4087#issuecomment-2250397630>
@@ -58,7 +58,7 @@ class Fp8PerTensorOnlineLinearMethod(LinearMethodBase):

 ### High Level Weight Transfer API

-The layerwise reloading system is integrated with the post-training weight transfer system. To use layerwise reloading in conjunction to the weight transfer system, follow the examples found [here](../../examples/rl/). Layerwise reloading is controlled by the `WeightTransferUpdateInfo.is_checkpoint_format` flag and is set to `True` by default.
+The layerwise reloading system is integrated with the post-training weight transfer system. To use layerwise reloading in conjunction to the weight transfer system, follow the examples found [here](../../examples/rl/). Checkpoint-format weight transfer engines (e.g. the NCCL and IPC backends) run layerwise reloading automatically inside their `start_weight_update`/`finish_weight_update` lifecycle.

 ### Mid Level `reload_weights` API

@@ -17,6 +17,7 @@ The weight transfer system follows a **four-phase protocol** with a pluggable ba
 | ------- | --------- | -------- |
 | [NCCL](nccl.md) | NCCL broadcast | Separate GPUs for training and inference |
 | [IPC](ipc.md) | CUDA IPC handles | Colocated training and inference on same GPU |
+| [sparse_nccl](nccl.md#sparse-nccl) | NCCL broadcast | Sparse flat-index weight patches (TP=1/PP=1) |

 ## Configuration

@@ -41,7 +42,7 @@ vllm serve my-model \
    --weight-transfer-config '{"backend": "nccl"}'
 ```

-The `backend` field accepts `"nccl"` (default) or `"ipc"`.
+The `backend` field accepts `"nccl"` (default), `"ipc"`, or `"sparse_nccl"`.

 ## API Endpoints

@@ -69,7 +70,7 @@ Both backends provide static methods that the trainer calls to send weights. The
 EngineClass.trainer_init(init_info)

 # 2. Start weight update on inference side
-llm.start_weight_update(is_checkpoint_format=True)
+llm.start_weight_update()

 # 3. Send weights to inference workers
 EngineClass.trainer_send_weights(
@@ -11,15 +11,23 @@ The `WeightTransferEngine` is a generic abstract class parameterized by two data

 ### Abstract Methods

-Subclasses must implement these four methods:
+Subclasses must implement these 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 |
+| `start_weight_update()` | Inference | Prepare for an update (e.g. begin layerwise reload); no-op for in-place engines |
+| `finish_weight_update()` | Inference | Finalize the update (e.g. finalize layerwise reload); no-op for in-place engines |
+| `receive_weights(update_info)` | Inference | Receive weights and load them into `self.model` |
 | `shutdown()` | Inference | Clean up resources |
 | `trainer_send_weights(iterator, trainer_args)` | Trainer | Static method to send weights from the trainer process |

+The base class provides two methods:
+
+1. `__init__` : Engines receive `config` (`WeightTransferConfig`),  `vllm_config` (`VllmConfig`), `device` (`torch.device`) and  `model` (`nn.Module`)  
+2. `update_weights(update_info_dict)`:  Thin wrapper for `receive_weights`: parses
+the dict into user-specified data type, calls `receive_weights`, and synchronizes the device. Subclasses implement `receive_weights`.
+
 ### 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.
@@ -81,7 +89,7 @@ class MyUpdateInfo(WeightTransferUpdateInfo):
 ### 2. Implement the Engine

 ```python
-from collections.abc import Callable, Iterator
+from collections.abc import Iterator
 from typing import Any
 import torch

@@ -93,18 +101,25 @@ class MyWeightTransferEngine(WeightTransferEngine[MyInitInfo, MyUpdateInfo]):
        # 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
+    def start_weight_update(self) -> None:
+        # Checkpoint-format engines: run initialize_layerwise_reload(self.model).
+        # In-place engines: no-op
+        ...
+
+    def finish_weight_update(self) -> None:
+        # Checkpoint-format engines: run finalize_layerwise_reload(...).
+        # In-place engines: no-op
+        ...
+
+    def receive_weights(self, update_info: MyUpdateInfo) -> None:
+        weights = []
        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)])
+            weights.append((name, weight))
+        self.model.load_weights(weights)

    def shutdown(self) -> None:
        # Clean up resources
@@ -121,9 +136,6 @@ class MyWeightTransferEngine(WeightTransferEngine[MyInitInfo, MyUpdateInfo]):
            ...
 ```

-!!! 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
@@ -147,7 +159,7 @@ Once registered, users can select your backend via `WeightTransferConfig(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.
+The factory uses a registry pattern with lazy loading. Built-in engines (`nccl`, `ipc`, and `sparse_nccl`) 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
@@ -155,7 +167,8 @@ 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,
+    vllm_config=vllm_config,
+    device=device,
    model=model,
 )
 ```
@@ -55,7 +55,7 @@ trainer_args = IPCTrainerSendWeightsArgs(
    llm_handle=llm_actor_handle,
 )
 # start
-ray.get(llm_actor_handle.start_weight_update.remote(is_checkpoint_format=True))
+ray.get(llm_actor_handle.start_weight_update.remote())
 # send weights
 IPCWeightTransferEngine.trainer_send_weights(
    iterator=model.named_parameters(),
@@ -80,7 +80,7 @@ trainer_args = IPCTrainerSendWeightsArgs(
 # start
 base_url = "http://localhost:8000"
 url = f"{base_url}/start_weight_update"
-response = requests.post(url, json={"is_checkpoint_format": True}, timeout=60)
+response = requests.post(url, json={}, timeout=60)
 response.raise_for_status()
 # send weights
 IPCWeightTransferEngine.trainer_send_weights(
@@ -11,7 +11,7 @@ The NCCL weight transfer engine uses [NCCL](https://developer.nvidia.com/nccl) b
 ## 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.
+2. The trainer broadcasts weights to all workers simultaneously. Each worker receives and loads the weights.
 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
@@ -93,7 +93,7 @@ remaining three steps are:
 from vllm.distributed.weight_transfer.base import WeightTransferUpdateRequest

 # 1. Start the weight update
-llm.start_weight_update(is_checkpoint_format=True)
+llm.start_weight_update()

 # 2. Receive weights (can be called multiple times for chunked transfers)
 llm.update_weights(
@@ -116,19 +116,22 @@ must match the order in which the trainer iterates over its parameters.

 `start_weight_update` must be called before `update_weights`, and
 `finish_weight_update` must be called after all weight chunks have been
-transferred. The `is_checkpoint_format` flag controls whether layerwise reload
-processing is applied (`True` for checkpoint-format weights, `False` for
-pre-processed kernel-format weights).
+transferred. The NCCL engine receives checkpoint-format weights and applies
+layerwise reload processing automatically inside `start_weight_update` /
+`finish_weight_update`.

-Sparse NCCL patches still use `update_kind="sparse_flat"` inside
-`update_info`, but they should be wrapped in
-`start_weight_update(is_checkpoint_format=False)` because sparse patches apply
-directly to runtime/kernel-format parameters. The current sparse MVP requires
-`TP=1` and `PP=1`.
+## Sparse NCCL
+
+Sparse, flat-index weight patches use a separate backend,
+`WeightTransferConfig(backend="sparse_nccl")`, implemented by
+`SparseNCCLWeightTransferEngine`. It shares only NCCL process-group
+initialization with the dense engine; patches are applied directly in place to
+existing parameters (no layerwise reload). The current sparse MVP requires
+`TP=1` and `PP=1`. See the example below.

 ## Examples

 - [RLHF with NCCL weight syncing (offline, Ray)](../../../examples/rl/rlhf_nccl.py) - Trainer on one GPU, 2x tensor-parallel vLLM engine on two others, with packed NCCL weight broadcast
- [RLHF with sparse NCCL weight syncing (offline, Ray)](../../../examples/rl/rlhf_sparse_nccl.py) - Dense-vs-sparse equivalence demo with a real model on a 2-GPU trainer/inference setup; sparse patches use `start_weight_update(is_checkpoint_format=False)` and currently require `TP=1` and `PP=1`
+- [RLHF with sparse NCCL weight syncing (offline, Ray)](../../../examples/rl/rlhf_sparse_nccl.py) - Dense-vs-sparse equivalence demo with a real model on a 2-GPU trainer/inference setup; sparse patches use `backend="sparse_nccl"` and currently require `TP=1` and `PP=1`
 - [RLHF with async weight syncing (offline, Ray)](../../../examples/rl/rlhf_async_new_apis.py) - 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.py) - Weight transfer with a running vLLM HTTP server using HTTP control plane and NCCL data plane
@@ -326,6 +326,27 @@ vLLM supports dynamically loading and unloading LoRA adapters at runtime via the

 **Warning:** Dynamic LoRA loading is not a secure operation and should not be enabled in deployments exposed to untrusted clients. If you must enable dynamic LoRA loading, restrict access to the `/v1/load_lora_adapter` and `/v1/unload_lora_adapter` endpoints to trusted administrators only, using a reverse proxy or network-level access controls. Do not expose these endpoints to end users. For details on configuring LoRA adapters, see the [LoRA Adapters documentation](../features/lora.md).

+## gRPC Interface
+
+vLLM provides an optional gRPC Generate service on a separate TCP port, enabled via the `--grpc-port` flag. When not specified, no gRPC server is started. The gRPC listener binds to the same host address as the HTTP server.
+
+**Warning:** The gRPC interface is **insecure by default** — it does not implement authentication, authorization, or encryption. It should be considered a private, internal interface intended for use only between co-located services within a trusted network. Do not expose the gRPC port to the public internet or untrusted clients. If you enable the gRPC interface, protect it via network-level access controls such as firewall rules, network segmentation, or deployment on an isolated private network.
+
+### Security Implications
+
+An attacker who can reach the gRPC port can:
+
+1. **Run arbitrary inference** via the `Generate` and `GenerateStream` RPCs without any credentials
+2. **Consume GPU and compute resources** by submitting unbounded generation requests
+3. **Cause Denial of Service** by exploiting bugs in the gRPC interface that can crash vLLM.
+
+### Recommendations
+
+- Only enable `--grpc-port` when you have a specific need for gRPC-based inference
+- Ensure the gRPC port is only accessible from trusted hosts or services
+- Use firewall rules to block external access to the gRPC port
+- Consider deploying the gRPC interface on a dedicated internal network interface
+
 ## Cache Directory Security

 vLLM assumes that its cache directories are **private and trusted**. Cache contents are loaded without cryptographic integrity verification, including formats that support arbitrary code execution. If an untrusted user or process can write to vLLM's cache directories, they may be able to crash vLLM or cause it to execute arbitrary code.
@@ -327,6 +327,9 @@ async def handle_request(api: str, request: Request):
        session, decode_response = await decode_request_task
        stream_generator = stream_decode_response(session, decode_response, request_id)
        response = await make_response(stream_generator)
+        response.headers["Content-Type"] = decode_response.headers.get(
+            "Content-Type", "application/json"
+        )
        return response
    except Exception as e:
        logger.exception("An error occurred while handling the request: %s", e)
@@ -91,44 +91,6 @@ def run_cohere_asr(question: str, audio_count: int) -> ModelRequestData:
    )


-# MusicFlamingo
-def run_musicflamingo(question: str, audio_count: int) -> ModelRequestData:
-    model_name = "nvidia/music-flamingo-2601-hf"
-    engine_args = EngineArgs(
-        model=model_name,
-        max_model_len=4096,
-        max_num_seqs=2,
-        limit_mm_per_prompt={"audio": audio_count},
-        enforce_eager=True,
-    )
-
-    # MusicFlamingo prompt placeholders use <sound>; vLLM's MusicFlamingo
-    # multimodal processor expands each one into <|sound_bos|> + audio tokens +
-    # <|sound_eos|> based on extracted audio feature lengths.
-    audio_placeholder = "<sound>" * audio_count
-    system_prompt = (
-        "You are Music Flamingo, a multimodal assistant for language and music. "
-        "On each turn you receive an audio clip which contains music and optional "
-        "text, you will receive at least one or both; use your world knowledge and "
-        "reasoning to help the user with any task. Interpret the entirety of the "
-        "content any input music--regardlenss of whether the user calls it audio, "
-        "music, or sound."
-    )
-
-    prompt = (
-        "<|im_start|>system\n"
-        f"{system_prompt}<|im_end|>\n"
-        "<|im_start|>user\n"
-        f"{audio_placeholder}{question}<|im_end|>\n"
-        "<|im_start|>assistant\n"
-    )
-
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompt=prompt,
-    )
-
-
 # Gemma3N
 def run_gemma3n(question: str, audio_count: int) -> ModelRequestData:
    model_name = "google/gemma-3n-E2B-it"
@@ -565,7 +527,6 @@ model_example_map = {
    "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,
@@ -74,39 +74,6 @@ def load_aria(question: str, image_urls: list[str]) -> ModelRequestData:
    )


-def load_aya_vision(question: str, image_urls: list[str]) -> ModelRequestData:
-    model_name = "CohereLabs/aya-vision-8b"
-
-    engine_args = EngineArgs(
-        model=model_name,
-        max_num_seqs=2,
-        limit_mm_per_prompt={"image": len(image_urls)},
-    )
-
-    placeholders = [{"type": "image", "image": url} for url in image_urls]
-    messages = [
-        {
-            "role": "user",
-            "content": [
-                *placeholders,
-                {"type": "text", "text": question},
-            ],
-        }
-    ]
-
-    processor = AutoProcessor.from_pretrained(model_name)
-
-    prompt = processor.apply_chat_template(
-        messages, tokenize=False, add_generation_prompt=True
-    )
-
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompt=prompt,
-        image_data=[fetch_image(url) for url in image_urls],
-    )
-
-
 def load_bee(question: str, image_urls: list[str]) -> ModelRequestData:
    model_name = "Open-Bee/Bee-8B-RL"

@@ -1420,7 +1387,6 @@ def load_molmo2(question: str, image_urls: list[str]) -> ModelRequestData:

 model_example_map = {
    "aria": load_aria,
-    "aya_vision": load_aya_vision,
    "bee": load_bee,
    "command_a_vision": load_command_a_vision,
    "deepseek_vl_v2": load_deepseek_vl2,
@@ -68,28 +68,6 @@ def run_aria(questions: list[str], modality: str) -> ModelRequestData:
    )


-# Aya Vision
-def run_aya_vision(questions: list[str], modality: str) -> ModelRequestData:
-    assert modality == "image"
-    model_name = "CohereLabs/aya-vision-8b"
-
-    engine_args = EngineArgs(
-        model=model_name,
-        max_model_len=2048,
-        max_num_seqs=2,
-        mm_processor_kwargs={"crop_to_patches": True},
-        limit_mm_per_prompt={modality: 1},
-    )
-    prompts = [
-        f"<|START_OF_TURN_TOKEN|><|USER_TOKEN|><image>{question}<|END_OF_TURN_TOKEN|><|START_OF_TURN_TOKEN|><|CHATBOT_TOKEN|>"
-        for question in questions
-    ]
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompts=prompts,
-    )
-
-
 # Bee-8B
 def run_bee(questions: list[str], modality: str) -> ModelRequestData:
    assert modality == "image"
@@ -1377,28 +1355,6 @@ def run_llava_onevision(questions: list[str], modality: str) -> ModelRequestData
    )


-# Mantis
-def run_mantis(questions: list[str], modality: str) -> ModelRequestData:
-    assert modality == "image"
-
-    llama3_template = "<|start_header_id|>user<|end_header_id|>\n\n{}<|eot_id|><|start_header_id|>assistant<|end_header_id|>\n\n"  # noqa: E501
-    prompts = [llama3_template.format(f"{question}\n<image>") for question in questions]
-
-    engine_args = EngineArgs(
-        model="TIGER-Lab/Mantis-8B-siglip-llama3",
-        max_model_len=4096,
-        hf_overrides={"architectures": ["MantisForConditionalGeneration"]},
-        limit_mm_per_prompt={modality: 1},
-    )
-    stop_token_ids = [128009]
-
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompts=prompts,
-        stop_token_ids=stop_token_ids,
-    )
-
-
 # MiniCPM-V
 def run_minicpmv_base(questions: list[str], modality: str, model_name):
    assert modality in ["image", "video", "image+video"]
@@ -2349,7 +2305,6 @@ def run_step_vl(questions: list[str], modality: str) -> ModelRequestData:

 model_example_map = {
    "aria": run_aria,
-    "aya_vision": run_aya_vision,
    "bagel": run_bagel,
    "cheers": run_cheers,
    "bee": run_bee,
@@ -2390,7 +2345,6 @@ model_example_map = {
    "llava-next": run_llava_next,
    "llava-next-video": run_llava_next_video,
    "llava-onevision": run_llava_onevision,
-    "mantis": run_mantis,
    "minicpmo": run_minicpmo,
    "minicpmv": run_minicpmv,
    "mistral3": run_mistral3,
@@ -306,7 +306,7 @@ gen_futures = [

 ray.get(llm.pause_after_n_tokens.remote())

-ray.get(llm.start_weight_update.remote(is_checkpoint_format=True))
+ray.get(llm.start_weight_update.remote())

 inference_handle = llm.update_weights.remote(
    WeightTransferUpdateRequest(
@@ -80,14 +80,10 @@ def init_weight_transfer_engine(base_url: str) -> None:
    response.raise_for_status()


-def start_weight_update(
-    base_url: str,
-    is_checkpoint_format: bool = True,
-) -> None:
+def start_weight_update(base_url: str) -> None:
    """Start a weight update via HTTP endpoint."""
    url = f"{base_url}/start_weight_update"
-    payload = {"is_checkpoint_format": is_checkpoint_format}
-    response = requests.post(url, json=payload, timeout=60)
+    response = requests.post(url, json={}, timeout=60)
    response.raise_for_status()


@@ -170,7 +166,7 @@ def main():
    pause_generation(BASE_URL)

    # Start weight update, broadcast via IPC, then finish
-    start_weight_update(BASE_URL, is_checkpoint_format=False)
+    start_weight_update(BASE_URL)

    print("Broadcasting weights via CUDA IPC (HTTP)...")
    trainer_args = IPCTrainerSendWeightsArgs(send_mode="http", url=BASE_URL)
@@ -83,14 +83,10 @@ def init_weight_transfer_engine(
    response.raise_for_status()


-def start_weight_update(
-    base_url: str,
-    is_checkpoint_format: bool = True,
-) -> None:
+def start_weight_update(base_url: str) -> None:
    """Start a weight update via HTTP endpoint."""
    url = f"{base_url}/start_weight_update"
-    payload = {"is_checkpoint_format": is_checkpoint_format}
-    response = requests.post(url, json=payload, timeout=60)
+    response = requests.post(url, json={}, timeout=60)
    response.raise_for_status()


@@ -223,7 +219,7 @@ def main():
        shapes.append(list(p.shape))

    # Start weight update
-    start_weight_update(BASE_URL, is_checkpoint_format=True)
+    start_weight_update(BASE_URL)

    # Start the update_weights call in a separate thread since it will block
    # waiting for NCCL broadcasts
@@ -139,7 +139,7 @@ ray.get(llm.sleep.remote(level=0))

 ray.get(train_model.init_weight_transfer.remote())
 # Start weight update, sync weights, then finish
-ray.get(llm.start_weight_update.remote(is_checkpoint_format=True))
+ray.get(llm.start_weight_update.remote())
 ray.get(train_model.broadcast_weights.remote(llm))
 ray.get(llm.finish_weight_update.remote())

@@ -277,15 +277,8 @@ class DataParallelInferenceEngine:
            ]
        )

-    def start_weight_update(self, is_checkpoint_format: bool = True):
-        ray.get(
-            [
-                actor.start_weight_update.remote(
-                    is_checkpoint_format=is_checkpoint_format
-                )
-                for actor in self.llm_actors
-            ]
-        )
+    def start_weight_update(self):
+        ray.get([actor.start_weight_update.remote() for actor in self.llm_actors])

    def finish_weight_update(self):
        ray.get([actor.finish_weight_update.remote() for actor in self.llm_actors])
@@ -392,7 +385,7 @@ def main():
    ray.get(inference_engine.wake_up.remote(tags=["weights"]))

    print("[sync] Starting weight update...")
-    ray.get(inference_engine.start_weight_update.remote(is_checkpoint_format=True))
+    ray.get(inference_engine.start_weight_update.remote())

    print("[sync] Packed IPC transfer FSDP → vLLM...")
    ray.get(
@@ -202,7 +202,7 @@ ray.get([train_handle, inference_handle])
 names, dtype_names, shapes = ray.get(train_model.get_weight_metadata.remote())

 # Start weight update
-ray.get(llm.start_weight_update.remote(is_checkpoint_format=True))
+ray.get(llm.start_weight_update.remote())

 # Issue update_weights call with NCCL-specific update info
 # packed=True enables efficient batched tensor broadcasting
@@ -299,7 +299,7 @@ async def main():
    print(f"[sync] Got metadata for {len(names)} parameters.")

    print("[sync] Starting weight update...")
-    await engine.start_weight_update(is_checkpoint_format=True)
+    await engine.start_weight_update()

    print("[sync] Broadcasting weights from FSDP → vLLM...")
    broadcast_handles = [
@@ -44,11 +44,14 @@ from transformers import AutoModelForCausalLM, AutoTokenizer

 from vllm import LLM, SamplingParams
 from vllm.config import WeightTransferConfig
-from vllm.distributed.weight_transfer.base import SparseWeightPatch
 from vllm.distributed.weight_transfer.nccl_engine import (
    NCCLTrainerSendWeightsArgs,
    NCCLWeightTransferEngine,
 )
+from vllm.distributed.weight_transfer.sparse_nccl_engine import (
+    SparseNCCLWeightTransferEngine,
+    SparseWeightPatch,
+)
 from vllm.utils.network_utils import get_ip, get_open_port

 MODEL_NAME = "Qwen/Qwen2.5-0.5B-Instruct"
@@ -244,7 +247,6 @@ class TrainModel:
            dtype_names=[str(self.patched_param.dtype).split(".")[-1]],
            shapes=[list(self.patched_param.shape)],
            num_updates_list=[flat_indices.numel()],
-            update_kind="sparse_flat",
        )
        return update_info, selected_token_ids, patch_digest, sparse_payload_bytes

@@ -271,7 +273,7 @@ class TrainModel:
            raise RuntimeError("Sparse patch has not been prepared")

        start = time.perf_counter()
-        NCCLWeightTransferEngine.trainer_send_sparse_weights(
+        SparseNCCLWeightTransferEngine.trainer_send_weights(
            iter(self.pending_sparse_patches),
            NCCLTrainerSendWeightsArgs(group=self.model_update_group),
        )
@@ -282,6 +284,7 @@ class TrainModel:

 def launch_llm(
    scheduling_inference: PlacementGroupSchedulingStrategy,
+    backend: str = "nccl",
 ):
    return ray.remote(
        num_cpus=0,
@@ -293,7 +296,7 @@ def launch_llm(
        tensor_parallel_size=1,
        distributed_executor_backend="ray",
        gpu_memory_utilization=0.7,
-        weight_transfer_config=WeightTransferConfig(backend="nccl"),
+        weight_transfer_config=WeightTransferConfig(backend=backend),
    )


@@ -332,7 +335,7 @@ def run_dense_phase(
    scheduling_inference: PlacementGroupSchedulingStrategy,
 ) -> dict[str, object]:
    ray.get(train_model.reset_model.remote())
-    llm = launch_llm(scheduling_inference)
+    llm = launch_llm(scheduling_inference, backend="nccl")
    try:
        dense_before = collect_vllm_generations(llm)

@@ -351,7 +354,7 @@ def run_dense_phase(
        )
        trainer_init = train_model.init_weight_transfer_group.remote(world_size)
        ray.get([trainer_init, inference_init])
-        ray.get(llm.start_weight_update.remote(is_checkpoint_format=True))
+        ray.get(llm.start_weight_update.remote())

        dense_update_info, dense_payload_bytes = ray.get(
            train_model.get_dense_update_info.remote()
@@ -391,7 +394,7 @@ def run_sparse_phase(
    scheduling_inference: PlacementGroupSchedulingStrategy,
 ) -> dict[str, object]:
    ray.get(train_model.reset_model.remote())
-    llm = launch_llm(scheduling_inference)
+    llm = launch_llm(scheduling_inference, backend="sparse_nccl")
    try:
        sparse_before = collect_vllm_generations(llm)

@@ -410,7 +413,7 @@ def run_sparse_phase(
        )
        trainer_init = train_model.init_weight_transfer_group.remote(world_size)
        ray.get([trainer_init, inference_init])
-        ray.get(llm.start_weight_update.remote(is_checkpoint_format=False))
+        ray.get(llm.start_weight_update.remote())

        sparse_update_info, selected_token_ids, patch_digest, sparse_payload_bytes = (
            ray.get(train_model.prepare_sparse_patch.remote(PROMPTS))
@@ -50,7 +50,8 @@ mod request;
 mod stream;

 use vllm_engine_core_client::EngineCoreClient;
-use vllm_engine_core_client::protocol::{ModelDtype, ReasoningParserKwargs};
+use vllm_engine_core_client::protocol::dtype::ModelDtype;
+use vllm_engine_core_client::protocol::request::ReasoningParserKwargs;
 use vllm_llm::Llm;
 use vllm_text::{Prompt, TextLlm, TextRequest};

@@ -22,7 +22,7 @@ use llm_multimodal::{
    TrackedMedia,
 };
 use tracing::warn;
-use vllm_engine_core_client::protocol::ModelDtype;
+use vllm_engine_core_client::protocol::dtype::ModelDtype;
 use vllm_engine_core_client::protocol::multimodal::{
    MmBatchedField, MmFeatureSpec, MmFeatures, MmField, MmFieldElem, MmFlatField, MmKwargsItem,
    MmSharedField, MmSlice, PlaceholderRange, SliceSpec,
@@ -2,7 +2,7 @@ use std::collections::HashMap;

 use half::{bf16, f16};
 use llm_multimodal::{ModelSpecificValue, PreprocessedImages};
-use vllm_engine_core_client::protocol::ModelDtype;
+use vllm_engine_core_client::protocol::dtype::ModelDtype;
 use vllm_engine_core_client::protocol::multimodal::MmKwargValue as ProtocolKwargValue;
 use vllm_engine_core_client::protocol::tensor::{ShapeExt as _, WireTensor};

@@ -1,7 +1,9 @@
 //! Applies xgrammar structural-tag constraints for strict tool calling.

 use thiserror_ext::AsReport;
-use vllm_engine_core_client::protocol::{StructuredOutputBackend, StructuredOutputsParams};
+use vllm_engine_core_client::protocol::structured_outputs::{
+    StructuredOutputBackend, StructuredOutputsParams,
+};
 use vllm_parser::tool::StructuralTagModel;
 use xgrammar_structural_tag::{
    FunctionDefinition, FunctionToolParam, ToolChoice as StructuralTagToolChoice, ToolParam,
@@ -76,7 +78,9 @@ fn structural_tag_tool_choice(request: &ChatRequest) -> Option<StructuralTagTool
 #[cfg(test)]
 mod tests {
    use serde_json::{Value, json};
-    use vllm_engine_core_client::protocol::{StructuredOutputBackend, StructuredOutputsParams};
+    use vllm_engine_core_client::protocol::structured_outputs::{
+        StructuredOutputBackend, StructuredOutputsParams,
+    };
    use vllm_parser::tool::{Qwen3CoderToolParser, Tool, ToolParser};

    use super::*;
@@ -15,9 +15,10 @@ use vllm_chat::{
 use vllm_engine_core_client::protocol::logprobs::{
    Logprobs, MaybeWireLogprobs, PositionLogprobs, TokenLogprob,
 };
-use vllm_engine_core_client::protocol::{
-    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs, EngineCoreRequest, StopReason,
+use vllm_engine_core_client::protocol::output::{
+    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs, StopReason,
 };
+use vllm_engine_core_client::protocol::request::EngineCoreRequest;
 use vllm_engine_core_client::test_utils::{IpcNamespace, spawn_mock_engine_task};
 use vllm_engine_core_client::{EngineCoreClient, EngineCoreClientConfig};
 use vllm_llm::Llm;
@@ -5,9 +5,9 @@ use clap::Parser;
 use futures::StreamExt as _;
 use tokio::time::timeout;
 use tracing_subscriber::EnvFilter;
-use vllm_engine_core_client::protocol::{
-    EngineCoreFinishReason, EngineCoreRequest, EngineCoreSamplingParams,
-};
+use vllm_engine_core_client::protocol::output::EngineCoreFinishReason;
+use vllm_engine_core_client::protocol::request::EngineCoreRequest;
+use vllm_engine_core_client::protocol::sampling::EngineCoreSamplingParams;
 use vllm_engine_core_client::{
    EngineCoreClient, EngineCoreClientConfig, EngineCoreStreamOutput, TransportMode,
 };
@@ -11,10 +11,11 @@ use tracing::{debug, info, trace};
 use crate::client::imp::{ClientInner, run_abort_loop, run_output_dispatcher_loop};
 use crate::coordinator::CoordinatorHandle;
 use crate::error::{Error, Result};
+use crate::protocol::dtype::ModelDtype;
 use crate::protocol::handshake::EngineCoreReadyResponse;
 use crate::protocol::lora::LoraRequest;
+use crate::protocol::request::{EngineCoreRequest, EngineCoreRequestType};
 use crate::protocol::utility::{EngineCoreUtilityRequest, PauseMode};
-use crate::protocol::{EngineCoreRequest, EngineCoreRequestType, ModelDtype};
 use crate::runtime::{BackgroundShutdownRuntime, build_zmq_runtime};
 use crate::transport::{self, ConnectedEngine};

@@ -16,12 +16,11 @@ use crate::client::stream::EngineCoreStreamOutput;
 use crate::client::{AbortCause, AbortRequest};
 use crate::error::{client_closed, dispatcher_closed, unexpected_dispatcher_output};
 use crate::metrics::{LoraInfoExporter, record_scheduler_stats};
+use crate::protocol::encode_msgpack;
+use crate::protocol::output::{ClassifiedEngineCoreOutputs, EngineCoreOutput, EngineCoreOutputs};
+use crate::protocol::request::EngineCoreRequestType;
 use crate::protocol::stats::SchedulerStats;
 use crate::protocol::utility::UtilityOutput;
-use crate::protocol::{
-    ClassifiedEngineCoreOutputs, EngineCoreOutput, EngineCoreOutputs, EngineCoreRequestType,
-    encode_msgpack,
-};
 use crate::transport::{ConnectedEngine, EngineId};
 use crate::{Error, Result, transport};

@@ -7,9 +7,9 @@ use tracing::trace;
 use crate::EngineId;
 use crate::client::stream::EngineCoreStreamOutput;
 use crate::error::{Error, Result};
+use crate::protocol::output::{EngineCoreEventType, EngineCoreFinishReason, EngineCoreOutput};
 use crate::protocol::stats::SchedulerStats;
 use crate::protocol::utility::UtilityOutput;
-use crate::protocol::{EngineCoreEventType, EngineCoreFinishReason, EngineCoreOutput};
 use crate::transport::ConnectedEngine;

 pub type OutputSender = mpsc::UnboundedSender<Result<EngineCoreStreamOutput>>;
@@ -452,7 +452,7 @@ mod tests {
        EngineLoadSnapshot, EngineRoutingState, RequestRegistry, UtilityRegistry,
    };
    use crate::mock_engine::default_ready_response;
-    use crate::protocol::{
+    use crate::protocol::output::{
        EngineCoreEvent, EngineCoreEventType, EngineCoreFinishReason, EngineCoreOutput,
    };
    use crate::transport::ConnectedEngine;
@@ -10,7 +10,7 @@ use tracing::{debug, error, warn};

 use crate::client::AbortRequest;
 use crate::client::state::OutputReceiver;
-use crate::protocol::{EngineCoreFinishReason, EngineCoreOutput};
+use crate::protocol::output::{EngineCoreFinishReason, EngineCoreOutput};
 use crate::{AbortCause, Error, Result};

 #[derive(Debug, Clone, Copy, PartialEq, Eq)]
@@ -10,10 +10,9 @@ use zeromq::{XPubSocket, ZmqMessage};
 use crate::client::imp::ClientInner;
 use crate::coordinator::handle::{CoordinatorCommand, CoordinatorState};
 use crate::error::{Error, Result, bail_unexpected_coordinator_output};
-use crate::protocol::{
-    ClassifiedEngineCoreOutputs, DpControlMessage, EngineCoreOutputs, EngineCoreRequestType,
-    encode_msgpack,
-};
+use crate::protocol::encode_msgpack;
+use crate::protocol::output::{ClassifiedEngineCoreOutputs, DpControlMessage, EngineCoreOutputs};
+use crate::protocol::request::EngineCoreRequestType;

 /// Coordinator-to-engine `START_DP_WAVE` control payload encoded on the
 /// engine-facing coordinator socket.
@@ -8,8 +8,9 @@ use zeromq::{DealerSocket, PushSocket, SocketOptions, SubSocket, ZmqMessage};

 use crate::EngineId;
 use crate::error::{Error, Result, bail_unexpected_handshake_message};
+use crate::protocol::dtype::ModelDtype;
 use crate::protocol::handshake::{EngineCoreReadyResponse, HandshakeInitMessage, ReadyMessage};
-use crate::protocol::{ModelDtype, decode_msgpack, encode_msgpack};
+use crate::protocol::{decode_msgpack, encode_msgpack};

 /// Default model length advertised by reusable mock engine helpers.
 pub const DEFAULT_MOCK_MAX_MODEL_LEN: u64 = 1024 * 1024;
@@ -2,7 +2,8 @@ use std::collections::BTreeMap;

 use serde::{Deserialize, Serialize};

-use crate::protocol::{ModelDtype, OpaqueValue};
+use crate::protocol::OpaqueValue;
+use crate::protocol::dtype::ModelDtype;

 /// Decoded engine startup-handshake payload sent on the handshake socket.
 ///
@@ -9,8 +9,7 @@ use enum_as_inner::EnumAsInner;
 use serde::{Deserialize, Deserializer, Serialize};

 use self::wire::*;
-use super::{EngineCoreOutput, EngineCoreOutputs, decode_msgpack};
-use crate::error::{Error, Result, bail_ext_value_decode, ext_value_decode};
+use crate::error::{Error, Result, bail_ext_value_decode};
 use crate::protocol::tensor::{WireArrayData, WireNdArray};

 /// One token candidate and its logprob metadata for a single sequence position.
@@ -160,7 +159,7 @@ impl Serialize for MaybeWireLogprobs {
 impl MaybeWireLogprobs {
    /// Resolve the wire representation into decoded logprobs by looking up aux
    /// frames and decoding raw views as needed.
-    fn resolve<Frame>(self, frames: &[Frame], field_prefix: &str) -> Result<Self>
+    pub(super) fn resolve<Frame>(self, frames: &[Frame], field_prefix: &str) -> Result<Self>
    where
        Frame: AsRef<[u8]>,
    {
@@ -171,37 +170,6 @@ impl MaybeWireLogprobs {
    }
 }

-impl EngineCoreOutputs {
-    /// Resolve all wire-format fields in-place by looking up aux frames and
-    /// decoding raw-view payloads as needed.
-    fn resolve_in_place<Frame>(&mut self, frames: &[Frame]) -> Result<()>
-    where
-        Frame: AsRef<[u8]>,
-    {
-        for output in &mut self.outputs {
-            output.resolve_in_place(frames)?;
-        }
-        Ok(())
-    }
-}
-
-impl EngineCoreOutput {
-    /// Resolve all wire-format fields in-place by looking up aux frames and
-    /// decoding raw-view payloads as needed.
-    fn resolve_in_place<Frame>(&mut self, frames: &[Frame]) -> Result<()>
-    where
-        Frame: AsRef<[u8]>,
-    {
-        self.new_logprobs = (self.new_logprobs.take())
-            .map(|value| value.resolve(frames, "new_logprobs"))
-            .transpose()?;
-        self.new_prompt_logprobs_tensors = (self.new_prompt_logprobs_tensors.take())
-            .map(|value| value.resolve(frames, "new_prompt_logprobs_tensors"))
-            .transpose()?;
-        Ok(())
-    }
-}
-
 impl WireLogprobs {
    /// Convert semantic per-position logprobs into the Python wire tuple shape.
    ///
@@ -315,16 +283,3 @@ impl WireLogprobs {
        Ok(Logprobs { positions })
    }
 }
-
-/// Decode one ordinary or multipart engine-core output message into the strong
-/// typed public protocol shape.
-pub fn decode_engine_core_outputs<Frame>(frames: &[Frame]) -> Result<EngineCoreOutputs>
-where
-    Frame: AsRef<[u8]>,
-{
-    let first_frame = frames.first().ok_or_else(|| ext_value_decode!("missing output frame"))?;
-
-    let mut outputs: EngineCoreOutputs = decode_msgpack(first_frame.as_ref())?;
-    outputs.resolve_in_place(frames)?;
-    Ok(outputs)
-}
@@ -3,8 +3,8 @@ use std::collections::BTreeSet;
 use bytes::Bytes;
 use rmpv::Value;

-use super::{Logprobs, PositionLogprobs, TokenLogprob, decode_engine_core_outputs};
-use crate::protocol::EngineCoreFinishReason;
+use super::{Logprobs, PositionLogprobs, TokenLogprob};
+use crate::protocol::output::{EngineCoreFinishReason, decode_engine_core_outputs};

 fn encode_value(value: &Value) -> Vec<u8> {
    let mut out = Vec::new();
@@ -1,28 +1,11 @@
 use std::any::type_name;
-use std::collections::{BTreeMap, BTreeSet, HashMap};
 use std::io::Cursor;

-use bytes::Bytes;
 use rmpv::Value;
 use serde::{Deserialize, Serialize};
-use serde_default::DefaultFromSerde;
-use serde_repr::{Deserialize_repr, Serialize_repr};
-use serde_tuple::{Deserialize_tuple, Serialize_tuple};
 use thiserror_ext::AsReport;

 use crate::error::{Error, Result};
-use crate::protocol::logprobs::MaybeWireLogprobs;
-use crate::protocol::multimodal::MmFeatures;
-use crate::protocol::stats::{PrefillStats, SchedulerStats};
-use crate::protocol::utility::UtilityOutput;
-
-// TODO: This module currently mixes reusable frontend-facing semantic types
-// (for example `FinishReason`, `StopReason`, `RequestOutputKind`, and future
-// cleaned-up frontend sampling types) with engine-core-specific wire DTOs and
-// handshake/control messages. While the Rust frontend is still evolving
-// quickly, keep them co-located here for iteration speed. Once the higher-level
-// API boundary stabilizes, move the truly reusable semantic types into a
-// lower-level common crate and keep the engine transport/wire messages here.

 /// Dynamic msgpack value used for schema positions that are preserved but not
 /// yet strongly typed in the early-stage Rust client.
@@ -36,499 +19,18 @@ fn is_false(v: &bool) -> bool {
    !v
 }

-fn default_top_p() -> f32 {
-    1.0
-}
-
-fn default_repetition_penalty() -> f32 {
-    1.0
-}
-
-fn default_temperature() -> f32 {
-    1.0
-}
-
-fn default_max_tokens() -> u32 {
-    16
-}
-
-mod classified_outputs;
 pub mod dtype;
 pub mod handshake;
 pub mod logprobs;
 pub mod lora;
 pub mod multimodal;
+pub mod output;
+pub mod request;
+pub mod sampling;
 pub mod stats;
+pub mod structured_outputs;
 pub mod tensor;
 pub mod utility;
-pub use classified_outputs::{
-    ClassifiedEngineCoreOutputs, DpControlMessage, RequestBatchOutputs, UtilityCallOutput,
-};
-pub use dtype::ModelDtype;
-pub use logprobs::decode_engine_core_outputs;
-
-/// Request types are encoded as single-byte protocol constants so they can be
-/// sent over the ZMQ socket without an extra encoding step.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L217-L228>
-#[derive(Debug, Clone, Copy, PartialEq, Eq)]
-#[repr(u8)]
-pub enum EngineCoreRequestType {
-    Add = 0,
-    Abort = 1,
-    StartDpWave = 2,
-    Utility = 3,
-}
-
-impl EngineCoreRequestType {
-    /// Decode the single-byte request type frame used on the engine input
-    /// socket. Returns `None` for unrecognized values.
-    pub fn from_frame(frame: &[u8]) -> Option<Self> {
-        let [value] = frame else {
-            return None;
-        };
-
-        match value {
-            0 => Some(Self::Add),
-            1 => Some(Self::Abort),
-            2 => Some(Self::StartDpWave),
-            3 => Some(Self::Utility),
-            _ => None,
-        }
-    }
-
-    /// Encode the request type as the single-byte frame used on the engine
-    /// input socket.
-    pub fn to_frame(self) -> Bytes {
-        Bytes::from_static(match self {
-            Self::Add => b"\x00",
-            Self::Abort => b"\x01",
-            Self::StartDpWave => b"\x02",
-            Self::Utility => b"\x03",
-        })
-    }
-}
-
-/// Reason a request finished: stop, length, abort, error, or repetition.
-///
-/// This mirrors the Python enum and uses integer encoding for compact wire
-/// representation.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L41-L63>
-#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize_repr, Deserialize_repr)]
-#[repr(u8)]
-pub enum EngineCoreFinishReason {
-    /// A stop string was emitted.
-    Stop = 0,
-    /// `max_tokens` or `max_model_len` was reached.
-    Length = 1,
-    /// The request was aborted by the client.
-    Abort = 2,
-    /// A retryable request-level internal error occurred.
-    Error = 3,
-    /// A repetitive token pattern was detected.
-    Repetition = 4,
-}
-
-/// Event types emitted by engine-core for one request.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L113-L118>
-#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize_repr, Deserialize_repr)]
-#[repr(u8)]
-pub enum EngineCoreEventType {
-    Queued = 1,
-    Scheduled = 2,
-    Preempted = 3,
-}
-
-/// A timestamped engine-core event associated with one request.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L121-L130>
-#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
-pub struct EngineCoreEvent {
-    pub r#type: EngineCoreEventType,
-    pub timestamp: f64,
-}
-
-/// Controls how intermediate outputs are returned to the frontend.
-///
-/// `Cumulative = 0` is intentionally not supported in Rust frontend.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/sampling_params.py#L146-L152>
-#[derive(Debug, Clone, Copy, PartialEq, Eq, Default, Serialize_repr, Deserialize_repr)]
-#[repr(u8)]
-pub enum RequestOutputKind {
-    /// Return only token deltas in each update.
-    #[default]
-    Delta = 1,
-    /// Suppress intermediate updates and return only the final output.
-    FinalOnly = 2,
-}
-
-/// Structured-output backend selected for EngineCore grammar compilation.
-///
-/// Python vLLM stores this in `StructuredOutputsParams._backend` after request
-/// validation. The Rust frontend currently always lowers structured-output
-/// requests to guidance, while ignoring any user-supplied `_backend` value.
-#[derive(Debug, Clone, Copy, Default, PartialEq, Eq, Serialize, Deserialize)]
-#[serde(rename_all = "kebab-case")]
-pub enum StructuredOutputBackend {
-    Xgrammar,
-    #[default]
-    Guidance,
-    Outlines,
-    LmFormatEnforcer,
-}
-
-/// The stop reason associated with a finished output.
-///
-/// Python models this as the union-typed `stop_reason: int | str | None`
-/// field on `EngineCoreOutput`; the Rust client narrows it into a tagged enum.
-///
-/// Original Python field:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L155>
-#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
-#[serde(untagged)]
-pub enum StopReason {
-    TokenId(u32),
-    Text(String),
-}
-
-/// Parameters for configuring structured outputs (guided decoding).
-///
-/// Exactly one constraint field (`json`, `regex`, `choice`, `grammar`,
-/// `json_object`, or `structural_tag`) should be set. The engine-core
-/// backend selects the appropriate grammar compiler based on which field
-/// is present.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/sampling_params.py#L36-L107>
-#[serde_with::skip_serializing_none]
-#[derive(Debug, Clone, Default, PartialEq, Serialize, Deserialize)]
-#[serde(default)]
-pub struct StructuredOutputsParams {
-    /// JSON schema (as a dict/object or JSON string) constraining the output.
-    pub json: Option<serde_json::Value>,
-    /// Regular expression the output must match.
-    pub regex: Option<String>,
-    /// List of allowed output strings (the model must produce one of these).
-    pub choice: Option<Vec<String>>,
-    /// Context-free grammar (in EBNF-like notation) the output must conform to.
-    pub grammar: Option<String>,
-    /// When `true`, output must be valid JSON (free-form, no schema).
-    pub json_object: Option<bool>,
-    /// Disable any additional whitespace in guided JSON output.
-    #[serde(skip_serializing_if = "crate::protocol::is_false")]
-    pub disable_any_whitespace: bool,
-    /// Disable `additionalProperties` in JSON schema output.
-    #[serde(skip_serializing_if = "crate::protocol::is_false")]
-    pub disable_additional_properties: bool,
-    /// Custom whitespace pattern for guided JSON output.
-    pub whitespace_pattern: Option<String>,
-    /// Structural tag configuration (JSON-encoded string).
-    pub structural_tag: Option<String>,
-    /// Structured-output backend, mirroring Python's internal `_backend`.
-    ///
-    /// User-supplied values are ignored during deserialization. This matches
-    /// Python's request boundary, where `_backend` is set by validation rather
-    /// than accepted as a request-level backend selector.
-    #[serde(
-        default,
-        rename = "_backend",
-        deserialize_with = "serde_with::rust::deserialize_ignore_any"
-    )]
-    pub backend: StructuredOutputBackend,
-}
-
-/// Engine-core-facing sampling parameters for text generation.
-///
-/// This is the normalized southbound subset used by the Rust frontend when it
-/// talks to Python engine-core over the wire. User-facing request semantics
-/// such as `stop` strings, `n`, `ignore_eos`, and output aggregation mode are
-/// intentionally handled by higher layers before values reach this DTO.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/sampling_params.py#L155-L291>
-// Python's SamplingParams is `omit_defaults=True`, so msgpack drops
-// default-valued keys; default the whole struct. Per-field fns cover the
-// non-zero defaults.
-#[serde_with::skip_serializing_none]
-#[derive(Debug, Clone, PartialEq, Serialize, Deserialize, DefaultFromSerde)]
-#[serde(default)]
-pub struct EngineCoreSamplingParams {
-    /// Controls randomness. Lower values are more deterministic; zero means
-    /// greedy sampling.
-    #[serde(default = "default_temperature")]
-    pub temperature: f32,
-    /// Cumulative probability threshold for nucleus sampling.
-    #[serde(default = "default_top_p")]
-    pub top_p: f32,
-    /// Maximum number of top tokens to consider. `0` means all tokens.
-    pub top_k: u32,
-    /// Random seed used by the sampler when present.
-    pub seed: Option<i64>,
-    /// Maximum number of tokens to generate per output sequence.
-    #[serde(default = "default_max_tokens")]
-    pub max_tokens: u32,
-    /// Minimum number of tokens to generate before EOS or stop-token handling.
-    pub min_tokens: u32,
-    /// Maximum number of reasoning ("thinking") tokens to emit before the
-    /// reasoning section is force-closed. `None` means unlimited; the
-    /// user-facing `-1` sentinel is normalized to `None` by the frontend before
-    /// reaching this DTO, so only non-negative values are sent. Enforced
-    /// engine-side (and only when a reasoning parser is configured).
-    pub thinking_token_budget: Option<u64>,
-    /// Number of log probabilities to return per generated token.
-    ///
-    /// `None` disables sample logprobs. `-1` requests the full vocabulary.
-    pub logprobs: Option<i32>,
-    /// Number of log probabilities to return per prompt token.
-    ///
-    /// `None` disables prompt logprobs. `-1` requests the full vocabulary.
-    pub prompt_logprobs: Option<i32>,
-    /// Minimum probability threshold for token sampling.
-    pub min_p: f32,
-    /// Frequency penalty applied by the sampler.
-    pub frequency_penalty: f32,
-    /// Presence penalty applied by the sampler.
-    pub presence_penalty: f32,
-    /// Repetition penalty applied by the sampler.
-    #[serde(default = "default_repetition_penalty")]
-    pub repetition_penalty: f32,
-    /// Token IDs that stop generation.
-    pub stop_token_ids: Vec<u32>,
-    /// Primary EOS token ID used by engine-core's dedicated EOS stop path.
-    ///
-    /// This mirrors Python's internal `_eos_token_id` field and is derived by
-    /// the frontend from tokenizer/model metadata rather than supplied directly
-    /// by end users.
-    #[serde(rename = "_eos_token_id")]
-    pub eos_token_id: Option<u32>,
-    /// Complete stop-token set used by engine-core for `min_tokens` masking.
-    ///
-    /// This mirrors Python's internal `_all_stop_token_ids` field and should
-    /// contain explicit `stop_token_ids` plus any frontend-derived EOS token
-    /// IDs.
-    #[serde(rename = "_all_stop_token_ids")]
-    pub all_stop_token_ids: BTreeSet<u32>,
-    /// Logit biases to apply during sampling.
-    /// Keys are token IDs
-    pub logit_bias: Option<HashMap<u32, f32>>,
-    /// Restrict output to these token IDs only.
-    pub allowed_token_ids: Option<Vec<u32>>,
-    /// Tokenized bad words to avoid during generation.
-    #[serde(rename = "_bad_words_token_ids")]
-    pub bad_words_token_ids: Option<Vec<Vec<u32>>>,
-    /// Parameters for configuring structured outputs (guided decoding).
-    pub structured_outputs: Option<StructuredOutputsParams>,
-    /// Specific token IDs for which log probabilities should be returned at
-    /// each position.
-    ///
-    /// When set, the engine returns logprobs for exactly these tokens in
-    /// addition to the sampled/scored token. Mutually exclusive with the
-    /// `logprobs` count field in practice.
-    pub logprob_token_ids: Option<Vec<u32>>,
-    /// If `Some(true)`, the request will not attempt to read from the prefix
-    /// cache; newly computed blocks may still populate the cache. `None`
-    /// defers to engine-core defaults.
-    pub skip_reading_prefix_cache: Option<bool>,
-    /// Additional request parameters for custom extensions (from `vllm_xargs`).
-    pub extra_args: Option<HashMap<String, serde_json::Value>>,
-}
-
-impl EngineCoreSamplingParams {
-    /// Constructs a default sampling params for testing purposes only.
-    pub fn for_test() -> Self {
-        Self {
-            temperature: 1.0,
-            top_p: 1.0,
-            top_k: 0,
-            seed: None,
-            max_tokens: 65536,
-            min_tokens: 0,
-            thinking_token_budget: None,
-            logprobs: None,
-            prompt_logprobs: None,
-            min_p: 0.0,
-            frequency_penalty: 0.0,
-            presence_penalty: 0.0,
-            repetition_penalty: 1.0,
-            stop_token_ids: Vec::new(),
-            eos_token_id: None,
-            all_stop_token_ids: BTreeSet::new(),
-            logit_bias: None,
-            allowed_token_ids: None,
-            bad_words_token_ids: None,
-            structured_outputs: None,
-            logprob_token_ids: None,
-            skip_reading_prefix_cache: None,
-            extra_args: None,
-        }
-    }
-}
-
-/// Extra kwargs consumed by engine-side reasoning parsers.
-///
-/// Original Python construction point:
-/// <https://github.com/vllm-project/vllm/blob/cec2ec11760f9f3beabd4c90451936078bf91533/vllm/entrypoints/openai/chat_completion/serving.py#L367-L369>
-#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
-pub struct ReasoningParserKwargs {
-    /// Effective kwargs visible to the chat template for this request.
-    pub chat_template_kwargs: HashMap<String, serde_json::Value>,
-}
-
-/// Engine-core add-request payload sent from frontend to engine.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/3f5bd482f5c1a5dbdffbbf68d624e20bb7032013/vllm/v1/engine/__init__.py#L80-L129>
-#[derive(Debug, Clone, PartialEq, Serialize_tuple, Deserialize_tuple, DefaultFromSerde)]
-pub struct EngineCoreRequest {
-    pub request_id: String,
-    pub prompt_token_ids: Option<Vec<u32>>,
-    /// Multimodal features attached to the request.
-    pub mm_features: Option<MmFeatures>,
-    pub sampling_params: Option<EngineCoreSamplingParams>,
-    /// Pooling parameters are preserved in the schema but not yet strongly
-    /// typed.
-    pub pooling_params: Option<OpaqueValue>,
-    pub arrival_time: f64,
-    #[serde(default)]
-    pub lora_request: Option<lora::LoraRequest>,
-    #[serde(default)]
-    pub cache_salt: Option<String>,
-    #[serde(default)]
-    pub data_parallel_rank: Option<u32>,
-    /// Unsupported in the first-stage Rust client because Python uses a custom
-    /// tensor/aux-frame encoding path for this field.
-    #[serde(default)]
-    pub prompt_embeds: Option<OpaqueValue>,
-    /// Per-position mask for mixed-mode inputs (e.g. chat completion with
-    /// `prompt_embeds` content parts). `Some(true)` means real token id;
-    /// `Some(false)` means the position uses a pre-computed entry from
-    /// `prompt_embeds`. `None` for pure-tokens and pure-embeds requests.
-    #[serde(default)]
-    pub prompt_is_token_ids: Option<Vec<bool>>,
-    /// Index of the client, used to ensure outputs are sent back to the same
-    /// client when scaling out the frontend.
-    #[serde(default)]
-    pub client_index: u32,
-    /// In DP mode, indicates which wave this request is expected to belong to.
-    #[serde(default)]
-    pub current_wave: u32,
-    #[serde(default)]
-    pub priority: i32,
-    #[serde(default)]
-    pub trace_headers: Option<BTreeMap<String, String>>,
-    #[serde(default)]
-    pub resumable: bool,
-    /// Original user-provided request ID, used for output reporting and aborts.
-    #[serde(default)]
-    pub external_req_id: Option<String>,
-    #[serde(default)]
-    pub reasoning_ended: Option<bool>,
-    /// Reasoning-parser kwargs forwarded from the frontend to the
-    /// structured-output backend.
-    #[serde(default)]
-    pub reasoning_parser_kwargs: Option<ReasoningParserKwargs>,
-    /// If `true`, the request should be added to the scheduler's waiting queue
-    /// and immediately aborted, so connector-side cleanup runs via the
-    /// standard `request_finished` hook.
-    #[serde(default)]
-    pub abort_immediately: bool,
-}
-
-impl EngineCoreRequest {
-    /// Validate fields intentionally not supported in the first-stage client.
-    pub fn validate(&self) -> Result<()> {
-        if self.prompt_embeds.is_some() {
-            return Err(Error::UnsupportedField {
-                context: "EngineCoreRequest",
-                field: "prompt_embeds",
-            });
-        }
-        Ok(())
-    }
-}
-
-/// Engine-core output for a single request.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/d3af8c18317c0dc008d42e4367fbb9045cfb7bf6/vllm/v1/engine/__init__.py#L154-L184>
-#[derive(Debug, Clone, PartialEq, Serialize_tuple, Deserialize_tuple, DefaultFromSerde)]
-pub struct EngineCoreOutput {
-    pub request_id: String,
-    pub new_token_ids: Vec<u32>,
-    /// Decoded sample logprobs for the newly generated positions in this
-    /// output.
-    #[serde(default)]
-    pub new_logprobs: Option<MaybeWireLogprobs>,
-    /// Decoded prompt logprobs for the scored prompt positions emitted in this
-    /// output.
-    #[serde(default)]
-    pub new_prompt_logprobs_tensors: Option<MaybeWireLogprobs>,
-    #[serde(default)]
-    pub pooling_output: Option<OpaqueValue>,
-    #[serde(default)]
-    pub finish_reason: Option<EngineCoreFinishReason>,
-    #[serde(default)]
-    pub stop_reason: Option<StopReason>,
-    #[serde(default)]
-    pub events: Option<Vec<EngineCoreEvent>>,
-    #[serde(default)]
-    pub kv_transfer_params: Option<serde_json::Value>,
-    #[serde(default)]
-    pub trace_headers: Option<OpaqueValue>,
-    /// Breakdown of the scheduled prefill computation, set on the first output
-    /// of a newly scheduled prefill and elided for subsequent decode outputs.
-    #[serde(default)]
-    pub prefill_stats: Option<PrefillStats>,
-    #[serde(default)]
-    pub routed_experts: Option<OpaqueValue>,
-    /// Number of NaNs seen in logits. Values above zero indicate corruption.
-    #[serde(default)]
-    pub num_nans_in_logits: u32,
-}
-
-impl EngineCoreOutput {
-    /// Returns whether this output is terminal for the request.
-    pub fn finished(&self) -> bool {
-        self.finish_reason.is_some()
-    }
-}
-
-/// Batch of engine-core outputs returned to a frontend client.
-///
-/// Original Python definition:
-/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L186-L214>
-#[derive(Debug, Clone, PartialEq, Serialize_tuple, Deserialize_tuple, DefaultFromSerde)]
-pub struct EngineCoreOutputs {
-    #[serde(default)]
-    pub engine_index: u32,
-    /// Outputs grouped for this client in the current engine tick.
-    #[serde(default)]
-    pub outputs: Vec<EngineCoreOutput>,
-    #[serde(default)]
-    pub scheduler_stats: Option<Box<SchedulerStats>>,
-    #[serde(default)]
-    pub timestamp: f64,
-    #[serde(default)]
-    pub utility_output: Option<UtilityOutput>,
-    #[serde(default)]
-    pub finished_requests: Option<BTreeSet<String>>,
-    /// In DP mode, signals that the current wave finished and engines are
-    /// paused.
-    #[serde(default)]
-    pub wave_complete: Option<u32>,
-    /// In DP mode, signals that a request arrived for an old wave and the next
-    /// wave needs to start in other engines.
-    #[serde(default)]
-    pub start_wave: Option<u32>,
-}

 /// Encode a Rust value into msgpack using the protocol crate's serde model.
 pub fn encode_msgpack<T>(value: &T) -> Result<Vec<u8>>
@@ -564,81 +66,17 @@ where
    })
 }

+/// Decode a msgpack payload into a dynamic value for diagnostics and tests.
 pub fn decode_value(bytes: &[u8]) -> Result<Value> {
    Ok(rmpv::decode::read_value(&mut Cursor::new(bytes))?)
 }

 #[cfg(test)]
 mod tests {
-    use std::collections::BTreeSet;
+    use std::collections::BTreeMap;

    use super::*;

-    #[test]
-    fn engine_core_request_serializes_as_full_array() {
-        let request = EngineCoreRequest {
-            request_id: "req-1".to_string(),
-            prompt_token_ids: Some(vec![1, 2, 3]),
-            sampling_params: Some(EngineCoreSamplingParams {
-                max_tokens: 8,
-                ..EngineCoreSamplingParams::for_test()
-            }),
-            arrival_time: 1234.5,
-            client_index: 7,
-            ..EngineCoreRequest::default()
-        };
-
-        let encoded = encode_msgpack(&request).unwrap();
-        let value = decode_value(&encoded).unwrap();
-        let array = match value {
-            Value::Array(array) => array,
-            other => panic!("expected array, got {other:?}"),
-        };
-
-        assert_eq!(array.len(), 20);
-        assert_eq!(array[0], Value::from("req-1"));
-        assert_eq!(array[2], Value::Nil);
-        assert_eq!(array[4], Value::Nil);
-        assert_eq!(array[10], Value::Nil);
-        assert_eq!(array[11], Value::from(7));
-    }
-
-    #[test]
-    fn engine_core_outputs_roundtrip_finished_fields() {
-        let outputs = EngineCoreOutputs {
-            outputs: vec![EngineCoreOutput {
-                request_id: "req-1".to_string(),
-                new_token_ids: vec![42],
-                new_logprobs: None,
-                new_prompt_logprobs_tensors: None,
-                pooling_output: None,
-                finish_reason: Some(EngineCoreFinishReason::Length),
-                stop_reason: Some(StopReason::Text("stop".to_string())),
-                events: None,
-                kv_transfer_params: None,
-                trace_headers: None,
-                prefill_stats: None,
-                routed_experts: None,
-                num_nans_in_logits: 0,
-            }],
-            finished_requests: Some(BTreeSet::from(["req-1".to_string()])),
-            ..Default::default()
-        };
-
-        let encoded = encode_msgpack(&outputs).unwrap();
-        let decoded: EngineCoreOutputs = decode_msgpack(&encoded).unwrap();
-
-        assert_eq!(decoded.outputs.len(), 1);
-        assert_eq!(
-            decoded.outputs[0].finish_reason,
-            Some(EngineCoreFinishReason::Length)
-        );
-        assert_eq!(
-            decoded.finished_requests,
-            Some(BTreeSet::from(["req-1".to_string()]))
-        );
-    }
-
    #[test]
    fn decode_msgpack_includes_type_name_and_value_fallback() {
        let error = decode_msgpack::<u64>(
@@ -648,72 +86,4 @@ mod tests {

        expect_test::expect![[r#"messagepack decode failed for u64: wrong msgpack marker FixMap(1); value fallback: {"status": "READY"}"#]].assert_eq(&error.to_report_string());
    }
-
-    #[test]
-    fn structured_outputs_backend_ignores_deserialized_value() {
-        let params: StructuredOutputsParams = serde_json::from_value(serde_json::json!({
-            "json_object": true,
-            "_backend": "xgrammar",
-        }))
-        .unwrap();
-
-        assert_eq!(params.backend, StructuredOutputBackend::Guidance);
-
-        let value = serde_json::to_value(params).unwrap();
-        assert_eq!(value["_backend"], "guidance");
-    }
-
-    /// A real `sampling_params` is a sparse `omit_defaults` map; absent fields
-    /// must fall back to defaults. `python_compat` can't catch this since Rust
-    /// encodes full maps (see `engine_core_request_serializes_as_full_array`).
-    #[test]
-    fn decodes_sampling_params_with_omitted_defaults() {
-        let sampling_params = Value::Map(vec![
-            (
-                Value::from("stop_token_ids"),
-                Value::Array(vec![Value::from(151643u32)]),
-            ),
-            (Value::from("skip_reading_prefix_cache"), Value::from(false)),
-        ]);
-        let request = Value::Array(vec![
-            Value::from("req-omit-defaults"),
-            Value::Array(vec![
-                Value::from(1u32),
-                Value::from(2u32),
-                Value::from(3u32),
-            ]),
-            Value::Nil,
-            sampling_params,
-            Value::Nil,
-            Value::from(1.0f64),
-        ]);
-
-        let mut bytes = Vec::new();
-        rmpv::encode::write_value(&mut bytes, &request).unwrap();
-
-        let decoded: EngineCoreRequest = decode_msgpack(&bytes)
-            .expect("a real omit_defaults request must decode (regression: missing field)");
-
-        assert_eq!(decoded.request_id, "req-omit-defaults");
-        let sampling = decoded.sampling_params.expect("sampling params present");
-
-        assert_eq!(sampling.stop_token_ids, vec![151643]);
-        assert_eq!(sampling.skip_reading_prefix_cache, Some(false));
-
-        // Omitted fields -> Python defaults.
-        assert_eq!(sampling.temperature, 1.0);
-        assert_eq!(sampling.top_p, 1.0);
-        assert_eq!(sampling.top_k, 0);
-        assert_eq!(sampling.seed, None);
-        assert_eq!(sampling.max_tokens, 16);
-        assert_eq!(sampling.min_tokens, 0);
-        assert_eq!(sampling.min_p, 0.0);
-        assert_eq!(sampling.frequency_penalty, 0.0);
-        assert_eq!(sampling.presence_penalty, 0.0);
-        assert_eq!(sampling.repetition_penalty, 1.0);
-        assert_eq!(sampling.logprobs, None);
-        assert_eq!(sampling.prompt_logprobs, None);
-        assert_eq!(sampling.eos_token_id, None);
-        assert!(sampling.all_stop_token_ids.is_empty());
-    }
 }
@@ -1,10 +1,164 @@
 use std::collections::BTreeSet;

 use enum_as_inner::EnumAsInner;
+use serde::{Deserialize, Serialize};
+use serde_default::DefaultFromSerde;
+use serde_repr::{Deserialize_repr, Serialize_repr};
+use serde_tuple::{Deserialize_tuple, Serialize_tuple};

 use super::utility::UtilityOutput;
-use super::{EngineCoreOutput, EngineCoreOutputs};
-use crate::protocol::stats::SchedulerStats;
+use crate::error::{Error, Result, ext_value_decode};
+use crate::protocol::logprobs::MaybeWireLogprobs;
+use crate::protocol::stats::{PrefillStats, SchedulerStats};
+use crate::protocol::{OpaqueValue, decode_msgpack};
+
+/// The stop reason associated with a finished output.
+///
+/// Python models this as the union-typed `stop_reason: int | str | None`
+/// field on `EngineCoreOutput`; the Rust client narrows it into a tagged enum.
+///
+/// Original Python field:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L155>
+#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
+#[serde(untagged)]
+pub enum StopReason {
+    TokenId(u32),
+    Text(String),
+}
+
+/// Reason a request finished: stop, length, abort, error, or repetition.
+///
+/// This mirrors the Python enum and uses integer encoding for compact wire
+/// representation.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L41-L63>
+#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize_repr, Deserialize_repr)]
+#[repr(u8)]
+pub enum EngineCoreFinishReason {
+    /// A stop string was emitted.
+    Stop = 0,
+    /// `max_tokens` or `max_model_len` was reached.
+    Length = 1,
+    /// The request was aborted by the client.
+    Abort = 2,
+    /// A retryable request-level internal error occurred.
+    Error = 3,
+    /// A repetitive token pattern was detected.
+    Repetition = 4,
+}
+
+/// Event types emitted by engine-core for one request.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L113-L118>
+#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize_repr, Deserialize_repr)]
+#[repr(u8)]
+pub enum EngineCoreEventType {
+    Queued = 1,
+    Scheduled = 2,
+    Preempted = 3,
+}
+
+/// A timestamped engine-core event associated with one request.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L121-L130>
+#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
+pub struct EngineCoreEvent {
+    pub r#type: EngineCoreEventType,
+    pub timestamp: f64,
+}
+
+/// Engine-core output for a single request.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/d3af8c18317c0dc008d42e4367fbb9045cfb7bf6/vllm/v1/engine/__init__.py#L154-L184>
+#[derive(Debug, Clone, PartialEq, Serialize_tuple, Deserialize_tuple, DefaultFromSerde)]
+pub struct EngineCoreOutput {
+    pub request_id: String,
+    pub new_token_ids: Vec<u32>,
+    /// Decoded sample logprobs for the newly generated positions in this
+    /// output.
+    #[serde(default)]
+    pub new_logprobs: Option<MaybeWireLogprobs>,
+    /// Decoded prompt logprobs for the scored prompt positions emitted in this
+    /// output.
+    #[serde(default)]
+    pub new_prompt_logprobs_tensors: Option<MaybeWireLogprobs>,
+    #[serde(default)]
+    pub pooling_output: Option<OpaqueValue>,
+    #[serde(default)]
+    pub finish_reason: Option<EngineCoreFinishReason>,
+    #[serde(default)]
+    pub stop_reason: Option<StopReason>,
+    #[serde(default)]
+    pub events: Option<Vec<EngineCoreEvent>>,
+    #[serde(default)]
+    pub kv_transfer_params: Option<serde_json::Value>,
+    #[serde(default)]
+    pub trace_headers: Option<OpaqueValue>,
+    /// Breakdown of the scheduled prefill computation, set on the first output
+    /// of a newly scheduled prefill and elided for subsequent decode outputs.
+    #[serde(default)]
+    pub prefill_stats: Option<PrefillStats>,
+    #[serde(default)]
+    pub routed_experts: Option<OpaqueValue>,
+    /// Number of NaNs seen in logits. Values above zero indicate corruption.
+    #[serde(default)]
+    pub num_nans_in_logits: u32,
+}
+
+impl EngineCoreOutput {
+    /// Returns whether this output is terminal for the request.
+    pub fn finished(&self) -> bool {
+        self.finish_reason.is_some()
+    }
+
+    /// Resolve all wire-format fields in-place by looking up aux frames and
+    /// decoding raw-view payloads as needed.
+    fn resolve_in_place<Frame>(&mut self, frames: &[Frame]) -> Result<()>
+    where
+        Frame: AsRef<[u8]>,
+    {
+        self.new_logprobs = (self.new_logprobs.take())
+            .map(|value| value.resolve(frames, "new_logprobs"))
+            .transpose()?;
+        self.new_prompt_logprobs_tensors = (self.new_prompt_logprobs_tensors.take())
+            .map(|value| value.resolve(frames, "new_prompt_logprobs_tensors"))
+            .transpose()?;
+        Ok(())
+    }
+}
+
+/// Batch of engine-core outputs returned to a frontend client.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L186-L214>
+#[derive(Debug, Clone, PartialEq, Serialize_tuple, Deserialize_tuple, DefaultFromSerde)]
+pub struct EngineCoreOutputs {
+    #[serde(default)]
+    pub engine_index: u32,
+    /// Outputs grouped for this client in the current engine tick.
+    #[serde(default)]
+    pub outputs: Vec<EngineCoreOutput>,
+    #[serde(default)]
+    pub scheduler_stats: Option<Box<SchedulerStats>>,
+    #[serde(default)]
+    pub timestamp: f64,
+    #[serde(default)]
+    pub utility_output: Option<UtilityOutput>,
+    #[serde(default)]
+    pub finished_requests: Option<BTreeSet<String>>,
+    /// In DP mode, signals that the current wave finished and engines are
+    /// paused.
+    #[serde(default)]
+    pub wave_complete: Option<u32>,
+    /// In DP mode, signals that a request arrived for an old wave and the next
+    /// wave needs to start in other engines.
+    #[serde(default)]
+    pub start_wave: Option<u32>,
+}

 /// Data-parallel control notifications multiplexed through `EngineCoreOutputs`.
 #[derive(Debug, Clone, PartialEq, Eq)]
@@ -49,6 +203,18 @@ pub enum ClassifiedEngineCoreOutputs {
 }

 impl EngineCoreOutputs {
+    /// Resolve all wire-format fields in-place by looking up aux frames and
+    /// decoding raw-view payloads as needed.
+    fn resolve_in_place<Frame>(&mut self, frames: &[Frame]) -> Result<()>
+    where
+        Frame: AsRef<[u8]>,
+    {
+        for output in &mut self.outputs {
+            output.resolve_in_place(frames)?;
+        }
+        Ok(())
+    }
+
    /// Classify the raw wire message into a more semantic Rust enum.
    pub fn classify(self) -> ClassifiedEngineCoreOutputs {
        let has_request_payload = !self.outputs.is_empty()
@@ -92,12 +258,62 @@ impl EngineCoreOutputs {
    }
 }

+/// Decode one ordinary or multipart engine-core output message into the strong
+/// typed public protocol shape.
+pub fn decode_engine_core_outputs<Frame>(frames: &[Frame]) -> Result<EngineCoreOutputs>
+where
+    Frame: AsRef<[u8]>,
+{
+    let first_frame = frames.first().ok_or_else(|| ext_value_decode!("missing output frame"))?;
+
+    let mut outputs: EngineCoreOutputs = decode_msgpack(first_frame.as_ref())?;
+    outputs.resolve_in_place(frames)?;
+    Ok(outputs)
+}
+
 #[cfg(test)]
 mod tests {
    use std::collections::BTreeSet;

    use super::*;
-    use crate::protocol::EngineCoreOutput;
+    use crate::protocol::output::EngineCoreOutput;
+    use crate::protocol::{decode_msgpack, encode_msgpack};
+
+    #[test]
+    fn engine_core_outputs_roundtrip_finished_fields() {
+        let outputs = EngineCoreOutputs {
+            outputs: vec![EngineCoreOutput {
+                request_id: "req-1".to_string(),
+                new_token_ids: vec![42],
+                new_logprobs: None,
+                new_prompt_logprobs_tensors: None,
+                pooling_output: None,
+                finish_reason: Some(EngineCoreFinishReason::Length),
+                stop_reason: Some(StopReason::Text("stop".to_string())),
+                events: None,
+                kv_transfer_params: None,
+                trace_headers: None,
+                prefill_stats: None,
+                routed_experts: None,
+                num_nans_in_logits: 0,
+            }],
+            finished_requests: Some(BTreeSet::from(["req-1".to_string()])),
+            ..Default::default()
+        };
+
+        let encoded = encode_msgpack(&outputs).unwrap();
+        let decoded: EngineCoreOutputs = decode_msgpack(&encoded).unwrap();
+
+        assert_eq!(decoded.outputs.len(), 1);
+        assert_eq!(
+            decoded.outputs[0].finish_reason,
+            Some(EngineCoreFinishReason::Length)
+        );
+        assert_eq!(
+            decoded.finished_requests,
+            Some(BTreeSet::from(["req-1".to_string()]))
+        );
+    }

    #[test]
    fn engine_core_outputs_classify_request_batch() {
@@ -0,0 +1,175 @@
+use std::collections::{BTreeMap, HashMap};
+
+use bytes::Bytes;
+use serde::{Deserialize, Serialize};
+use serde_default::DefaultFromSerde;
+use serde_tuple::{Deserialize_tuple, Serialize_tuple};
+
+use crate::protocol::multimodal::MmFeatures;
+use crate::protocol::sampling::EngineCoreSamplingParams;
+use crate::protocol::{OpaqueValue, lora};
+use crate::{Error, Result};
+
+/// Request types are encoded as single-byte protocol constants so they can be
+/// sent over the ZMQ socket without an extra encoding step.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L217-L228>
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+#[repr(u8)]
+pub enum EngineCoreRequestType {
+    Add = 0,
+    Abort = 1,
+    StartDpWave = 2,
+    Utility = 3,
+}
+
+impl EngineCoreRequestType {
+    /// Decode the single-byte request type frame used on the engine input
+    /// socket. Returns `None` for unrecognized values.
+    pub fn from_frame(frame: &[u8]) -> Option<Self> {
+        let [value] = frame else {
+            return None;
+        };
+
+        match value {
+            0 => Some(Self::Add),
+            1 => Some(Self::Abort),
+            2 => Some(Self::StartDpWave),
+            3 => Some(Self::Utility),
+            _ => None,
+        }
+    }
+
+    /// Encode the request type as the single-byte frame used on the engine
+    /// input socket.
+    pub fn to_frame(self) -> Bytes {
+        Bytes::from_static(match self {
+            Self::Add => b"\x00",
+            Self::Abort => b"\x01",
+            Self::StartDpWave => b"\x02",
+            Self::Utility => b"\x03",
+        })
+    }
+}
+
+/// Extra kwargs consumed by engine-side reasoning parsers.
+///
+/// Original Python construction point:
+/// <https://github.com/vllm-project/vllm/blob/cec2ec11760f9f3beabd4c90451936078bf91533/vllm/entrypoints/openai/chat_completion/serving.py#L367-L369>
+#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
+pub struct ReasoningParserKwargs {
+    /// Effective kwargs visible to the chat template for this request.
+    pub chat_template_kwargs: HashMap<String, serde_json::Value>,
+}
+
+/// Engine-core add-request payload sent from frontend to engine.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/3f5bd482f5c1a5dbdffbbf68d624e20bb7032013/vllm/v1/engine/__init__.py#L80-L129>
+#[derive(Debug, Clone, PartialEq, Serialize_tuple, Deserialize_tuple, DefaultFromSerde)]
+pub struct EngineCoreRequest {
+    pub request_id: String,
+    pub prompt_token_ids: Option<Vec<u32>>,
+    /// Multimodal features attached to the request.
+    pub mm_features: Option<MmFeatures>,
+    pub sampling_params: Option<EngineCoreSamplingParams>,
+    /// Pooling parameters are preserved in the schema but not yet strongly
+    /// typed.
+    pub pooling_params: Option<OpaqueValue>,
+    pub arrival_time: f64,
+    #[serde(default)]
+    pub lora_request: Option<lora::LoraRequest>,
+    #[serde(default)]
+    pub cache_salt: Option<String>,
+    #[serde(default)]
+    pub data_parallel_rank: Option<u32>,
+    /// Unsupported in the first-stage Rust client because Python uses a custom
+    /// tensor/aux-frame encoding path for this field.
+    #[serde(default)]
+    pub prompt_embeds: Option<OpaqueValue>,
+    /// Per-position mask for mixed-mode inputs (e.g. chat completion with
+    /// `prompt_embeds` content parts). `Some(true)` means real token id;
+    /// `Some(false)` means the position uses a pre-computed entry from
+    /// `prompt_embeds`. `None` for pure-tokens and pure-embeds requests.
+    #[serde(default)]
+    pub prompt_is_token_ids: Option<Vec<bool>>,
+    /// Index of the client, used to ensure outputs are sent back to the same
+    /// client when scaling out the frontend.
+    #[serde(default)]
+    pub client_index: u32,
+    /// In DP mode, indicates which wave this request is expected to belong to.
+    #[serde(default)]
+    pub current_wave: u32,
+    #[serde(default)]
+    pub priority: i32,
+    #[serde(default)]
+    pub trace_headers: Option<BTreeMap<String, String>>,
+    #[serde(default)]
+    pub resumable: bool,
+    /// Original user-provided request ID, used for output reporting and aborts.
+    #[serde(default)]
+    pub external_req_id: Option<String>,
+    #[serde(default)]
+    pub reasoning_ended: Option<bool>,
+    /// Reasoning-parser kwargs forwarded from the frontend to the
+    /// structured-output backend.
+    #[serde(default)]
+    pub reasoning_parser_kwargs: Option<ReasoningParserKwargs>,
+    /// If `true`, the request should be added to the scheduler's waiting queue
+    /// and immediately aborted, so connector-side cleanup runs via the
+    /// standard `request_finished` hook.
+    #[serde(default)]
+    pub abort_immediately: bool,
+}
+
+impl EngineCoreRequest {
+    /// Validate fields intentionally not supported in the first-stage client.
+    pub fn validate(&self) -> Result<()> {
+        if self.prompt_embeds.is_some() {
+            return Err(Error::UnsupportedField {
+                context: "EngineCoreRequest",
+                field: "prompt_embeds",
+            });
+        }
+        Ok(())
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use rmpv::Value;
+
+    use super::*;
+    use crate::protocol::sampling::EngineCoreSamplingParams;
+    use crate::protocol::{decode_value, encode_msgpack};
+
+    #[test]
+    fn engine_core_request_serializes_as_full_array() {
+        let request = EngineCoreRequest {
+            request_id: "req-1".to_string(),
+            prompt_token_ids: Some(vec![1, 2, 3]),
+            sampling_params: Some(EngineCoreSamplingParams {
+                max_tokens: 8,
+                ..EngineCoreSamplingParams::for_test()
+            }),
+            arrival_time: 1234.5,
+            client_index: 7,
+            ..EngineCoreRequest::default()
+        };
+
+        let encoded = encode_msgpack(&request).unwrap();
+        let value = decode_value(&encoded).unwrap();
+        let array = match value {
+            Value::Array(array) => array,
+            other => panic!("expected array, got {other:?}"),
+        };
+
+        assert_eq!(array.len(), 20);
+        assert_eq!(array[0], Value::from("req-1"));
+        assert_eq!(array[2], Value::Nil);
+        assert_eq!(array[4], Value::Nil);
+        assert_eq!(array[10], Value::Nil);
+        assert_eq!(array[11], Value::from(7));
+    }
+}
@@ -0,0 +1,211 @@
+use std::collections::{BTreeSet, HashMap};
+
+use serde::{Deserialize, Serialize};
+use serde_default::DefaultFromSerde;
+
+use crate::protocol::structured_outputs::StructuredOutputsParams;
+
+fn default_top_p() -> f32 {
+    1.0
+}
+
+fn default_repetition_penalty() -> f32 {
+    1.0
+}
+
+fn default_temperature() -> f32 {
+    1.0
+}
+
+fn default_max_tokens() -> u32 {
+    16
+}
+
+/// Engine-core-facing sampling parameters for text generation.
+///
+/// This is the normalized southbound subset used by the Rust frontend when it
+/// talks to Python engine-core over the wire. User-facing request semantics
+/// such as `stop` strings, `n`, `ignore_eos`, and output aggregation mode are
+/// intentionally handled by higher layers before values reach this DTO.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/sampling_params.py#L155-L291>
+// Python's SamplingParams is `omit_defaults=True`, so msgpack drops
+// default-valued keys; default the whole struct. Per-field fns cover the
+// non-zero defaults.
+#[serde_with::skip_serializing_none]
+#[derive(Debug, Clone, PartialEq, Serialize, Deserialize, DefaultFromSerde)]
+#[serde(default)]
+pub struct EngineCoreSamplingParams {
+    /// Controls randomness. Lower values are more deterministic; zero means
+    /// greedy sampling.
+    #[serde(default = "default_temperature")]
+    pub temperature: f32,
+    /// Cumulative probability threshold for nucleus sampling.
+    #[serde(default = "default_top_p")]
+    pub top_p: f32,
+    /// Maximum number of top tokens to consider. `0` means all tokens.
+    pub top_k: u32,
+    /// Random seed used by the sampler when present.
+    pub seed: Option<i64>,
+    /// Maximum number of tokens to generate per output sequence.
+    #[serde(default = "default_max_tokens")]
+    pub max_tokens: u32,
+    /// Minimum number of tokens to generate before EOS or stop-token handling.
+    pub min_tokens: u32,
+    /// Maximum number of reasoning ("thinking") tokens to emit before the
+    /// reasoning section is force-closed. `None` means unlimited; the
+    /// user-facing `-1` sentinel is normalized to `None` by the frontend before
+    /// reaching this DTO, so only non-negative values are sent. Enforced
+    /// engine-side (and only when a reasoning parser is configured).
+    pub thinking_token_budget: Option<u64>,
+    /// Number of log probabilities to return per generated token.
+    ///
+    /// `None` disables sample logprobs. `-1` requests the full vocabulary.
+    pub logprobs: Option<i32>,
+    /// Number of log probabilities to return per prompt token.
+    ///
+    /// `None` disables prompt logprobs. `-1` requests the full vocabulary.
+    pub prompt_logprobs: Option<i32>,
+    /// Minimum probability threshold for token sampling.
+    pub min_p: f32,
+    /// Frequency penalty applied by the sampler.
+    pub frequency_penalty: f32,
+    /// Presence penalty applied by the sampler.
+    pub presence_penalty: f32,
+    /// Repetition penalty applied by the sampler.
+    #[serde(default = "default_repetition_penalty")]
+    pub repetition_penalty: f32,
+    /// Token IDs that stop generation.
+    pub stop_token_ids: Vec<u32>,
+    /// Primary EOS token ID used by engine-core's dedicated EOS stop path.
+    ///
+    /// This mirrors Python's internal `_eos_token_id` field and is derived by
+    /// the frontend from tokenizer/model metadata rather than supplied directly
+    /// by end users.
+    #[serde(rename = "_eos_token_id")]
+    pub eos_token_id: Option<u32>,
+    /// Complete stop-token set used by engine-core for `min_tokens` masking.
+    ///
+    /// This mirrors Python's internal `_all_stop_token_ids` field and should
+    /// contain explicit `stop_token_ids` plus any frontend-derived EOS token
+    /// IDs.
+    #[serde(rename = "_all_stop_token_ids")]
+    pub all_stop_token_ids: BTreeSet<u32>,
+    /// Logit biases to apply during sampling.
+    /// Keys are token IDs
+    pub logit_bias: Option<HashMap<u32, f32>>,
+    /// Restrict output to these token IDs only.
+    pub allowed_token_ids: Option<Vec<u32>>,
+    /// Tokenized bad words to avoid during generation.
+    #[serde(rename = "_bad_words_token_ids")]
+    pub bad_words_token_ids: Option<Vec<Vec<u32>>>,
+    /// Parameters for configuring structured outputs (guided decoding).
+    pub structured_outputs: Option<StructuredOutputsParams>,
+    /// Specific token IDs for which log probabilities should be returned at
+    /// each position.
+    ///
+    /// When set, the engine returns logprobs for exactly these tokens in
+    /// addition to the sampled/scored token. Mutually exclusive with the
+    /// `logprobs` count field in practice.
+    pub logprob_token_ids: Option<Vec<u32>>,
+    /// If `Some(true)`, the request will not attempt to read from the prefix
+    /// cache; newly computed blocks may still populate the cache. `None`
+    /// defers to engine-core defaults.
+    pub skip_reading_prefix_cache: Option<bool>,
+    /// Additional request parameters for custom extensions (from `vllm_xargs`).
+    pub extra_args: Option<HashMap<String, serde_json::Value>>,
+}
+
+impl EngineCoreSamplingParams {
+    /// Constructs a default sampling params for testing purposes only.
+    pub fn for_test() -> Self {
+        Self {
+            temperature: 1.0,
+            top_p: 1.0,
+            top_k: 0,
+            seed: None,
+            max_tokens: 65536,
+            min_tokens: 0,
+            thinking_token_budget: None,
+            logprobs: None,
+            prompt_logprobs: None,
+            min_p: 0.0,
+            frequency_penalty: 0.0,
+            presence_penalty: 0.0,
+            repetition_penalty: 1.0,
+            stop_token_ids: Vec::new(),
+            eos_token_id: None,
+            all_stop_token_ids: BTreeSet::new(),
+            logit_bias: None,
+            allowed_token_ids: None,
+            bad_words_token_ids: None,
+            structured_outputs: None,
+            logprob_token_ids: None,
+            skip_reading_prefix_cache: None,
+            extra_args: None,
+        }
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use rmpv::Value;
+
+    use crate::protocol::decode_msgpack;
+    use crate::protocol::request::EngineCoreRequest;
+
+    /// A real `sampling_params` is a sparse `omit_defaults` map; absent fields
+    /// must fall back to defaults. `python_compat` can't catch this since Rust
+    /// encodes full maps (see `engine_core_request_serializes_as_full_array`).
+    #[test]
+    fn decodes_sampling_params_with_omitted_defaults() {
+        let sampling_params = Value::Map(vec![
+            (
+                Value::from("stop_token_ids"),
+                Value::Array(vec![Value::from(151643u32)]),
+            ),
+            (Value::from("skip_reading_prefix_cache"), Value::from(false)),
+        ]);
+        let request = Value::Array(vec![
+            Value::from("req-omit-defaults"),
+            Value::Array(vec![
+                Value::from(1u32),
+                Value::from(2u32),
+                Value::from(3u32),
+            ]),
+            Value::Nil,
+            sampling_params,
+            Value::Nil,
+            Value::from(1.0f64),
+        ]);
+
+        let mut bytes = Vec::new();
+        rmpv::encode::write_value(&mut bytes, &request).unwrap();
+
+        let decoded: EngineCoreRequest = decode_msgpack(&bytes)
+            .expect("a real omit_defaults request must decode (regression: missing field)");
+
+        assert_eq!(decoded.request_id, "req-omit-defaults");
+        let sampling = decoded.sampling_params.expect("sampling params present");
+
+        assert_eq!(sampling.stop_token_ids, vec![151643]);
+        assert_eq!(sampling.skip_reading_prefix_cache, Some(false));
+
+        // Omitted fields -> Python defaults.
+        assert_eq!(sampling.temperature, 1.0);
+        assert_eq!(sampling.top_p, 1.0);
+        assert_eq!(sampling.top_k, 0);
+        assert_eq!(sampling.seed, None);
+        assert_eq!(sampling.max_tokens, 16);
+        assert_eq!(sampling.min_tokens, 0);
+        assert_eq!(sampling.min_p, 0.0);
+        assert_eq!(sampling.frequency_penalty, 0.0);
+        assert_eq!(sampling.presence_penalty, 0.0);
+        assert_eq!(sampling.repetition_penalty, 1.0);
+        assert_eq!(sampling.logprobs, None);
+        assert_eq!(sampling.prompt_logprobs, None);
+        assert_eq!(sampling.eos_token_id, None);
+        assert!(sampling.all_stop_token_ids.is_empty());
+    }
+}
@@ -0,0 +1,81 @@
+use serde::{Deserialize, Serialize};
+
+/// Structured-output backend selected for EngineCore grammar compilation.
+///
+/// Python vLLM stores this in `StructuredOutputsParams._backend` after request
+/// validation. The Rust frontend currently always lowers structured-output
+/// requests to guidance, while ignoring any user-supplied `_backend` value.
+#[derive(Debug, Clone, Copy, Default, PartialEq, Eq, Serialize, Deserialize)]
+#[serde(rename_all = "kebab-case")]
+pub enum StructuredOutputBackend {
+    Xgrammar,
+    #[default]
+    Guidance,
+    Outlines,
+    LmFormatEnforcer,
+}
+
+/// Parameters for configuring structured outputs (guided decoding).
+///
+/// Exactly one constraint field (`json`, `regex`, `choice`, `grammar`,
+/// `json_object`, or `structural_tag`) should be set. The engine-core
+/// backend selects the appropriate grammar compiler based on which field
+/// is present.
+///
+/// Original Python definition:
+/// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/sampling_params.py#L36-L107>
+#[serde_with::skip_serializing_none]
+#[derive(Debug, Clone, Default, PartialEq, Serialize, Deserialize)]
+#[serde(default)]
+pub struct StructuredOutputsParams {
+    /// JSON schema (as a dict/object or JSON string) constraining the output.
+    pub json: Option<serde_json::Value>,
+    /// Regular expression the output must match.
+    pub regex: Option<String>,
+    /// List of allowed output strings (the model must produce one of these).
+    pub choice: Option<Vec<String>>,
+    /// Context-free grammar (in EBNF-like notation) the output must conform to.
+    pub grammar: Option<String>,
+    /// When `true`, output must be valid JSON (free-form, no schema).
+    pub json_object: Option<bool>,
+    /// Disable any additional whitespace in guided JSON output.
+    #[serde(skip_serializing_if = "crate::protocol::is_false")]
+    pub disable_any_whitespace: bool,
+    /// Disable `additionalProperties` in JSON schema output.
+    #[serde(skip_serializing_if = "crate::protocol::is_false")]
+    pub disable_additional_properties: bool,
+    /// Custom whitespace pattern for guided JSON output.
+    pub whitespace_pattern: Option<String>,
+    /// Structural tag configuration (JSON-encoded string).
+    pub structural_tag: Option<String>,
+    /// Structured-output backend, mirroring Python's internal `_backend`.
+    ///
+    /// User-supplied values are ignored during deserialization. This matches
+    /// Python's request boundary, where `_backend` is set by validation rather
+    /// than accepted as a request-level backend selector.
+    #[serde(
+        default,
+        rename = "_backend",
+        deserialize_with = "serde_with::rust::deserialize_ignore_any"
+    )]
+    pub backend: StructuredOutputBackend,
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn structured_outputs_backend_ignores_deserialized_value() {
+        let params: StructuredOutputsParams = serde_json::from_value(serde_json::json!({
+            "json_object": true,
+            "_backend": "xgrammar",
+        }))
+        .unwrap();
+
+        assert_eq!(params.backend, StructuredOutputBackend::Guidance);
+
+        let value = serde_json::to_value(params).unwrap();
+        assert_eq!(value["_backend"], "guidance");
+    }
+}
@@ -22,13 +22,14 @@ use crate::protocol::multimodal::{
    MmFeatureSpec, MmField, MmFieldElem, MmFlatField, MmKwargValue, MmSlice, PlaceholderRange,
    SliceSpec,
 };
+use crate::protocol::output::{
+    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs, decode_engine_core_outputs,
+};
+use crate::protocol::request::{EngineCoreRequest, EngineCoreRequestType};
+use crate::protocol::sampling::EngineCoreSamplingParams;
 use crate::protocol::stats::SchedulerStats;
 use crate::protocol::tensor::WireTensor;
 use crate::protocol::utility::{UtilityOutput, UtilityResultEnvelope};
-use crate::protocol::{
-    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs, EngineCoreRequest,
-    EngineCoreRequestType, EngineCoreSamplingParams, decode_engine_core_outputs,
-};
 use crate::test_utils::{
    IpcNamespace, setup_bootstrapped_mock_engine, setup_mock_engine_sockets,
    setup_mock_engine_with_init, spawn_mock_engine_task,
@@ -18,9 +18,8 @@ use crate::error::{Error, Result, bail_unexpected_handshake_message};
 use crate::protocol::handshake::{
    EngineCoreReadyResponse, HandshakeAddresses, HandshakeInitMessage, ReadyMessage,
 };
-use crate::protocol::{
-    EngineCoreOutputs, decode_engine_core_outputs, decode_msgpack, encode_msgpack,
-};
+use crate::protocol::output::{EngineCoreOutputs, decode_engine_core_outputs};
+use crate::protocol::{decode_msgpack, encode_msgpack};

 /// Dedicated single-frame sentinel emitted by Python `EngineCoreProc` when the
 /// engine dies.
@@ -5,7 +5,7 @@ use clap::Parser;
 use futures::StreamExt as _;
 use tokio::time::timeout;
 use tracing_subscriber::EnvFilter;
-use vllm_engine_core_client::protocol::EngineCoreSamplingParams;
+use vllm_engine_core_client::protocol::sampling::EngineCoreSamplingParams;
 use vllm_engine_core_client::{EngineCoreClient, EngineCoreClientConfig, TransportMode};
 use vllm_llm::{FinishReason, GenerateOutputStream, GenerateRequest, Llm};

@@ -8,7 +8,7 @@ use futures::stream::FusedStream;
 use futures::{Stream, StreamExt as _, pin_mut};
 use serde::{Deserialize, Serialize};
 use vllm_engine_core_client::protocol::logprobs::Logprobs;
-use vllm_engine_core_client::protocol::{EngineCoreFinishReason, StopReason};
+use vllm_engine_core_client::protocol::output::{EngineCoreFinishReason, StopReason};
 use vllm_engine_core_client::{AbortCause, EngineCoreOutputStream};

 use crate::error::Result;
@@ -4,9 +4,8 @@ use std::time::{SystemTime, UNIX_EPOCH};
 use uuid::Uuid;
 use vllm_engine_core_client::protocol::lora::LoraRequest;
 use vllm_engine_core_client::protocol::multimodal::MmFeatures;
-use vllm_engine_core_client::protocol::{
-    EngineCoreRequest, EngineCoreSamplingParams, ReasoningParserKwargs,
-};
+use vllm_engine_core_client::protocol::request::{EngineCoreRequest, ReasoningParserKwargs};
+use vllm_engine_core_client::protocol::sampling::EngineCoreSamplingParams;

 use crate::error::{Error, Result};

@@ -134,7 +133,8 @@ fn current_unix_timestamp_secs() -> f64 {
 mod tests {
    use std::collections::BTreeMap;

-    use vllm_engine_core_client::protocol::{EngineCoreSamplingParams, ReasoningParserKwargs};
+    use vllm_engine_core_client::protocol::request::ReasoningParserKwargs;
+    use vllm_engine_core_client::protocol::sampling::EngineCoreSamplingParams;

    use super::GenerateRequest;
    use crate::error::Error;
@@ -1,7 +1,9 @@
 use std::time::{SystemTime, UNIX_EPOCH};

+use vllm_engine_core_client::protocol::output::{
+    EngineCoreEvent, EngineCoreEventType, EngineCoreOutput,
+};
 use vllm_engine_core_client::protocol::stats::PrefillStats;
-use vllm_engine_core_client::protocol::{EngineCoreEvent, EngineCoreEventType, EngineCoreOutput};
 use vllm_metrics::{
    EngineLabels, FinishedReasonLabels, METRICS, PromptTokenSourceLabels, RequestMetrics,
 };
@@ -328,8 +330,8 @@ pub(crate) fn current_unix_timestamp_secs() -> f64 {

 #[cfg(test)]
 mod tests {
+    use vllm_engine_core_client::protocol::output::{EngineCoreEvent, EngineCoreEventType};
    use vllm_engine_core_client::protocol::stats::PrefillStats;
-    use vllm_engine_core_client::protocol::{EngineCoreEvent, EngineCoreEventType};

    use super::{RequestMetricsTracker, diff_or_zero};

@@ -341,7 +343,7 @@ mod tests {
            2,
            10.0,
            100.2,
-            &vllm_engine_core_client::protocol::EngineCoreOutput {
+            &vllm_engine_core_client::protocol::output::EngineCoreOutput {
                request_id: "req-1".to_string(),
                new_token_ids: vec![1],
                finish_reason: None,
@@ -369,7 +371,7 @@ mod tests {
            2,
            11.5,
            100.4,
-            &vllm_engine_core_client::protocol::EngineCoreOutput {
+            &vllm_engine_core_client::protocol::output::EngineCoreOutput {
                request_id: "req-1".to_string(),
                new_token_ids: vec![2, 3],
                finish_reason: None,
@@ -9,11 +9,13 @@ use uuid::Uuid;
 use vllm_engine_core_client::protocol::logprobs::{
    Logprobs, MaybeWireLogprobs, PositionLogprobs, TokenLogprob,
 };
-use vllm_engine_core_client::protocol::stats::PrefillStats;
-use vllm_engine_core_client::protocol::{
+use vllm_engine_core_client::protocol::output::{
    EngineCoreEvent, EngineCoreEventType, EngineCoreFinishReason, EngineCoreOutput,
-    EngineCoreOutputs, EngineCoreRequest, EngineCoreSamplingParams,
+    EngineCoreOutputs,
 };
+use vllm_engine_core_client::protocol::request::EngineCoreRequest;
+use vllm_engine_core_client::protocol::sampling::EngineCoreSamplingParams;
+use vllm_engine_core_client::protocol::stats::PrefillStats;
 use vllm_engine_core_client::test_utils::{IpcNamespace, spawn_mock_engine_task};
 use vllm_engine_core_client::{EngineCoreClient, EngineCoreClientConfig};
 use vllm_llm::{
@@ -11,12 +11,13 @@ use tokio::sync::mpsc;
 use tokio::task::yield_now;
 use tokio_util::sync::CancellationToken;
 use tracing::{debug, info, warn};
+use vllm_engine_core_client::protocol::output::{
+    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs,
+};
+use vllm_engine_core_client::protocol::request::EngineCoreRequest;
 use vllm_engine_core_client::protocol::utility::{
    EngineCoreUtilityRequest, UtilityOutput, UtilityResultEnvelope,
 };
-use vllm_engine_core_client::protocol::{
-    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs, EngineCoreRequest,
-};

 use super::Opt;

@@ -4,10 +4,9 @@ use tokio::sync::mpsc;
 use tokio_util::sync::CancellationToken;
 use tracing::warn;
 use vllm_engine_core_client::mock_engine::MockEngineDataSockets;
+use vllm_engine_core_client::protocol::request::{EngineCoreRequest, EngineCoreRequestType};
 use vllm_engine_core_client::protocol::utility::EngineCoreUtilityRequest;
-use vllm_engine_core_client::protocol::{
-    EngineCoreRequest, EngineCoreRequestType, decode_msgpack, encode_msgpack,
-};
+use vllm_engine_core_client::protocol::{decode_msgpack, encode_msgpack};
 use zeromq::{DealerSocket, PushSocket, SocketRecv as _, SocketSend as _, ZmqMessage};

 use crate::engine::{EngineInput, EngineOutput};
@@ -5,9 +5,9 @@ use anyhow::Result;
 use futures::StreamExt as _;
 use tokio::time::timeout;
 use tokio_util::sync::CancellationToken;
-use vllm_engine_core_client::protocol::{
-    EngineCoreFinishReason, EngineCoreRequest, EngineCoreSamplingParams,
-};
+use vllm_engine_core_client::protocol::output::EngineCoreFinishReason;
+use vllm_engine_core_client::protocol::request::EngineCoreRequest;
+use vllm_engine_core_client::protocol::sampling::EngineCoreSamplingParams;
 use vllm_engine_core_client::test_utils::IpcNamespace;
 use vllm_engine_core_client::{EngineCoreClient, EngineCoreClientConfig, TransportMode};

@@ -5,14 +5,13 @@

 mod adapter;

+pub(super) use adapter::UnifiedToolParserAdapter;
 use futures::FutureExt as _;
 use openai_protocol::common::{Function as OpenAiFunction, Tool as OpenAiTool};
 use tool_parser::traits::ToolParser as ExternalToolParser;
 use vllm_parser::tool::test_utils::collect_stream;
 use vllm_parser::tool::{Tool, ToolParser};

-pub(super) use adapter::UnifiedToolParserAdapter;
-
 pub(super) fn openai_tools(tools: &[Tool]) -> Vec<OpenAiTool> {
    tools
        .iter()
@@ -231,9 +231,10 @@ fn _rust_tool_parser(m: &Bound<'_, PyModule>) -> PyResult<()> {

 #[cfg(test)]
 mod tests {
-    use super::*;
    use serde_json::json;

+    use super::*;
+
    fn with_python<R>(f: impl for<'py> FnOnce(Python<'py>) -> R) -> R {
        Python::initialize();
        Python::attach(f)
@@ -49,10 +49,8 @@ mod tests {
    use std::sync::Arc;

    use super::SeedOssReasoningParser;
-    use crate::reasoning::{
-        ReasoningParser,
-        tests::{SEED_THINK_END_ID, SEED_THINK_START_ID, fake_tokenizer},
-    };
+    use crate::reasoning::ReasoningParser;
+    use crate::reasoning::tests::{SEED_THINK_END_ID, SEED_THINK_START_ID, fake_tokenizer};

    #[test]
    fn without_prompt_markers_expects_start_token() {
@@ -127,10 +127,8 @@ mod tests {
    use std::sync::Arc;

    use super::Step3p5ReasoningParser;
-    use crate::reasoning::{
-        ReasoningParser,
-        tests::{THINK_START_ID, fake_tokenizer},
-    };
+    use crate::reasoning::ReasoningParser;
+    use crate::reasoning::tests::{THINK_START_ID, fake_tokenizer};

    #[test]
    fn picks_up_prompt_start_boundary() {
@@ -238,6 +238,9 @@ mod tests {

        let error = parser.parse_chunk(&input).unwrap_err();

-        expect!["tool parser parsing failed: "].assert_eq(&error.to_report_string());
+        expect![[
+            r#"tool parser parsing failed: near "tool<｜tool▁sep｜>get_weather\n```json\n{}": "#
+        ]]
+        .assert_eq(&error.to_report_string());
    }
 }
@@ -242,6 +242,7 @@ mod tests {

        let error = parser.parse_chunk(&input).unwrap_err();

-        expect!["tool parser parsing failed: "].assert_eq(&error.to_report_string());
+        expect![[r#"tool parser parsing failed: near "<｜tool▁sep｜>{}": "#]]
+            .assert_eq(&error.to_report_string());
    }
 }
@@ -540,8 +540,10 @@ mod tests {
            .parse_chunk(r#"<tool_call>{"name":"f","arguments":42}</tool_call>"#)
            .unwrap_err();

-        expect!["tool parser parsing failed: invalid Granite4 arguments"]
-            .assert_eq(&error.to_report_string());
+        expect![[
+            r#"tool parser parsing failed: near "42}</tool_call>": invalid Granite4 arguments"#
+        ]]
+        .assert_eq(&error.to_report_string());
    }

    #[test]
@@ -335,7 +335,7 @@ mod tests {
        let error = parser.parse_chunk(&input).unwrap_err();

        expect![[r#"
-            tool parser parsing failed: invalid InternLM2
+            tool parser parsing failed: near "{\"name\":\"get_weather\",\"params\":{\"location\":\"Tokyo\"}}<|action_end|>": invalid InternLM2
            expected `parameters`, `arguments`"#]]
        .assert_eq(&error.to_report_string());
    }
@@ -336,7 +336,7 @@ mod tests {
            .unwrap_err();

        expect![[r#"
-            tool parser parsing failed: invalid Llama JSON
+            tool parser parsing failed: near "{\"name\":\"get_weather\",\"arguments\":{\"location\":\"Tokyo\"}}": invalid Llama JSON
            expected `parameters`"#]]
        .assert_eq(&error.to_report_string());
    }
@@ -474,7 +474,7 @@ mod tests {
        let error = parser.parse_chunk(r#"{"parameters":{},"name":"get_weather"}"#).unwrap_err();

        expect![[r#"
-            tool parser parsing failed: invalid Llama JSON
+            tool parser parsing failed: near "{\"parameters\":{},\"name\":\"get_weather\"}": invalid Llama JSON
            expected `name`"#]]
        .assert_eq(&error.to_report_string());
    }
@@ -489,7 +489,7 @@ mod tests {
            ))
            .unwrap_err();

-        expect!["tool parser parsing failed: invalid Llama JSON"]
+        expect![[r#"tool parser parsing failed: near " trailing": invalid Llama JSON"#]]
            .assert_eq(&error.to_report_string());
    }
 }
@@ -240,7 +240,7 @@ mod tests {
            .unwrap_err();

        expect![[r#"
-            tool parser parsing failed: invalid Mistral
+            tool parser parsing failed: near "{\"arguments\":{},\"name\":\"get_weather\"}]": invalid Mistral
            expected `name`"#]]
        .assert_eq(&error.to_report_string());
    }
@@ -280,7 +280,7 @@ mod tests {
            .unwrap_err();

        expect![[r#"
-            tool parser parsing failed: invalid Qwen XML
+            tool parser parsing failed: near "{\"arguments\":{},\"name\":\"get_weather\"}\n</tool_call>": invalid Qwen XML
            expected `name`"#]]
        .assert_eq(&error.to_report_string());
    }
@@ -594,6 +594,9 @@ mod tests {

        let error = parser.parse_chunk(&input).unwrap_err();

-        expect!["tool parser parsing failed: "].assert_eq(&error.to_report_string());
+        expect![[
+            r#"tool parser parsing failed: near "get_weather<|tool_call_argument_begin|>{}": "#
+        ]]
+        .assert_eq(&error.to_report_string());
    }
 }
@@ -593,6 +593,7 @@ mod tests {
        let mut parser = MinimaxM2ToolParser::new(&test_tools());
        let error = parser.parse_chunk("<minimax:tool_call><bad></minimax:tool_call>").unwrap_err();

-        expect!["tool parser parsing failed: "].assert_eq(&error.to_report_string());
+        expect![[r#"tool parser parsing failed: near "<bad></minimax:tool_call>": "#]]
+            .assert_eq(&error.to_report_string());
    }
 }
@@ -878,7 +878,10 @@ mod tests {
            ))
            .unwrap_err();

-        expect!["tool parser parsing failed: "].assert_eq(&error.to_report_string());
+        expect![[
+            r#"tool parser parsing failed: near "]<]minimax[>[<bad>]<]minimax[>[</tool_call>": "#
+        ]]
+        .assert_eq(&error.to_report_string());
    }

    #[test]
@@ -14,8 +14,6 @@ mod parameters;
 mod qwen_coder;
 #[cfg(any(test, feature = "test-util"))]
 pub mod test_utils;
-use crate::utils;
-
 use std::collections::{BTreeMap, btree_map};

 pub use deepseek_dsml::{DeepSeekV4ToolParser, DeepSeekV32ToolParser};
@@ -35,6 +33,8 @@ use serde::{Deserialize, Serialize};
 use serde_json::Value;
 pub use xgrammar_structural_tag::Model as StructuralTagModel;

+use crate::utils;
+
 /// One function-style tool made available to the model.
 #[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
 pub struct Tool {
@@ -685,7 +685,8 @@ mod tests {
        let mut parser = Qwen3CoderToolParser::new(&test_tools());
        let error = parser.parse_chunk("<tool_call>\n<bad>\n</tool_call>").unwrap_err();

-        expect!["tool parser parsing failed: "].assert_eq(&error.to_report_string());
+        expect![[r#"tool parser parsing failed: near "\n<bad>\n</tool_call>": "#]]
+            .assert_eq(&error.to_report_string());
    }

    #[test]
@@ -697,7 +698,7 @@ mod tests {
            )
            .unwrap_err();

-        expect!["tool parser parsing failed: "].assert_eq(&error.to_report_string());
+        expect![[r#"tool parser parsing failed: near "\n<function=get_weather>\n<parameter=location>SF</function>\n</tool_call>": "#]].assert_eq(&error.to_report_string());
    }

    #[test]
@@ -2,11 +2,10 @@

 use vllm_tokenizer::DynTokenizer;

+use super::{Result, UnifiedParser, UnifiedParserError, UnifiedParserOutput};
 use crate::reasoning::ReasoningParser;
 use crate::tool::{StructuralTagModel, Tool, ToolParser, ToolParserOutput};

-use super::{Result, UnifiedParser, UnifiedParserError, UnifiedParserOutput};
-
 /// Unified parser that composes existing reasoning and tool parsers.
 pub struct CombinedParser {
    reasoning: Option<Box<dyn ReasoningParser>>,
@@ -1,4 +1,5 @@
 use serde_json::{Map, Number, Value};
+use vllm_tokenizer::DynTokenizer;
 use winnow::ascii::multispace0 as ws0;
 use winnow::combinator::{alt, delimited, eof, opt, separated, seq, terminated};
 use winnow::error::{ContextError, ErrMode, ModalResult};
@@ -6,8 +7,6 @@ use winnow::prelude::*;
 use winnow::stream::{Partial, Stream};
 use winnow::token::{literal, take_till, take_until};

-use vllm_tokenizer::DynTokenizer;
-
 use super::{Result, UnifiedParser, UnifiedParserError, UnifiedParserOutput};
 use crate::reasoning::last_reasoning_boundary;
 use crate::tool::{Tool, ToolCallDelta};
@@ -3,13 +3,12 @@
 mod combined;
 mod gemma4;

+pub use combined::CombinedParser;
+pub use gemma4::Gemma4UnifiedParser;
 use thiserror::Error;
 use thiserror_ext::Macro;
 use vllm_tokenizer::DynTokenizer;

-pub use combined::CombinedParser;
-pub use gemma4::Gemma4UnifiedParser;
-
 use crate::reasoning::ReasoningError;
 use crate::tool::{
    StructuralTagModel, Tool, ToolCallDelta, ToolParserError, ToolParserEvent, ToolParserOutput,
@@ -395,9 +395,9 @@ pub fn parse_buffered_event<E>(
        Ok(event) => event,
        Err(ErrMode::Incomplete(_)) => return Ok(None),
        Err(ErrMode::Backtrack(e) | ErrMode::Cut(e)) => {
-            // TODO: enrich context for error reporting
+            let snippet = buffer.char_indices().nth(80).map_or(buffer, |(i, _)| &buffer[..i]);
            return Err(ToolParserError::ParsingFailed {
-                message: e.to_string(),
+                message: format!("near {snippet:?}: {e}"),
            });
        }
    };
@@ -423,8 +423,9 @@ mod tests {
    use winnow::stream::{Offset, Partial, Stream};

    use super::{
-        JsonObjectScanState, JsonStringScanState, MarkerScanState, json_str, partial_prefix_len,
-        safe_text_len, safe_text_len_mul, take_json_object, take_json_string, take_until_marker,
+        JsonObjectScanState, JsonStringScanState, MarkerScanState, json_str, parse_buffered_event,
+        partial_prefix_len, safe_text_len, safe_text_len_mul, take_json_object, take_json_string,
+        take_until_marker,
    };

    #[test]
@@ -832,4 +833,27 @@ mod tests {

        assert!(matches!(error, ErrMode::Incomplete(_)));
    }
+
+    #[test]
+    fn parse_buffered_event_error_includes_input_snippet() {
+        let result = parse_buffered_event(" {\"x\":1}", |input| {
+            take_json_object(input, &mut JsonObjectScanState::default())
+        });
+        let err = result.unwrap_err().to_string();
+        assert!(err.contains("near \""), "error must include snippet");
+    }
+
+    #[test]
+    fn parse_buffered_event_error_truncates_long_input() {
+        let long_input = format!(" {}", "x".repeat(100));
+        let result = parse_buffered_event(&long_input, |input| {
+            take_json_object(input, &mut JsonObjectScanState::default())
+        });
+        let err = result.unwrap_err().to_string();
+        assert!(err.contains("near \""), "error must include snippet");
+        assert!(
+            !err.contains(&long_input),
+            "snippet must be truncated for long input"
+        );
+    }
 }
@@ -1,8 +1,7 @@
 use axum::Json;
 use axum::http::StatusCode;
 use axum::response::{IntoResponse, Response};
-use thiserror_ext::AsReport as _;
-use thiserror_ext::{Construct, Macro};
+use thiserror_ext::{AsReport as _, Construct, Macro};

 use crate::routes::openai::utils::types::{ErrorDetail, ErrorResponse};

@@ -3,7 +3,8 @@

 use tonic::Status;
 use uuid::Uuid;
-use vllm_engine_core_client::protocol::{StopReason, StructuredOutputsParams};
+use vllm_engine_core_client::protocol::output::StopReason;
+use vllm_engine_core_client::protocol::structured_outputs::StructuredOutputsParams;
 use vllm_text::{
    DecodedLogprobs, DecodedPromptLogprobs, FinishReason, Finished, Prompt, SamplingParams,
    TextDecodeOptions, TextRequest,
@@ -502,7 +503,7 @@ impl ResponseOpts {

 #[cfg(test)]
 mod tests {
-    use vllm_engine_core_client::protocol::StopReason;
+    use vllm_engine_core_client::protocol::output::StopReason;
    use vllm_text::{FinishReason, Finished, Prompt};

    use super::pb::finish_info::{FinishReason as PbFinishReason, StopReason as PbStopReason};
@@ -18,9 +18,10 @@ use vllm_chat::{
    ChatBackend, ChatLlm, ChatRenderer, ChatRequest, ChatTextBackend, DefaultChatOutputProcessor,
    DynChatOutputProcessor, DynChatRenderer, NewChatOutputProcessorOptions, RenderedPrompt,
 };
-use vllm_engine_core_client::protocol::{
-    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs, EngineCoreRequest,
+use vllm_engine_core_client::protocol::output::{
+    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs,
 };
+use vllm_engine_core_client::protocol::request::EngineCoreRequest;
 use vllm_engine_core_client::test_utils::{IpcNamespace, spawn_mock_engine_task};
 use vllm_engine_core_client::{EngineCoreClient, EngineCoreClientConfig, EngineId};
 use vllm_llm::Llm;
@@ -152,8 +152,8 @@ impl axum::serve::Listener for Listener {

 /// Allow the unified listener to be adaptable to `tls_listener`.
 impl AsyncAccept for Listener {
-    type Connection = ListenerIo;
    type Address = ListenerAddr;
+    type Connection = ListenerIo;
    type Error = std::io::Error;

    fn poll_accept(
@@ -57,9 +57,9 @@ where
    S::Error: Send + 'static,
    B: Send + 'static,
 {
-    type Response = S::Response;
    type Error = S::Error;
    type Future = BoxFuture<'static, Result<Self::Response, Self::Error>>;
+    type Response = S::Response;

    fn poll_ready(&mut self, cx: &mut Context<'_>) -> Poll<Result<(), Self::Error>> {
        self.inner.poll_ready(cx)
@@ -18,9 +18,10 @@ use vllm_chat::{
    ChatBackend, ChatLlm, ChatRenderer, ChatRequest, ChatTextBackend, DefaultChatOutputProcessor,
    DynChatOutputProcessor, DynChatRenderer, NewChatOutputProcessorOptions, RenderedPrompt,
 };
-use vllm_engine_core_client::protocol::{
-    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs, EngineCoreRequest,
+use vllm_engine_core_client::protocol::output::{
+    EngineCoreFinishReason, EngineCoreOutput, EngineCoreOutputs,
 };
+use vllm_engine_core_client::protocol::request::EngineCoreRequest;
 use vllm_engine_core_client::test_utils::{IpcNamespace, spawn_mock_engine_task};
 use vllm_engine_core_client::{EngineCoreClient, EngineCoreClientConfig, EngineId};
 use vllm_llm::Llm;
@@ -21,7 +21,7 @@ use vllm_chat::{
    AssistantBlockKind, AssistantMessageExt as _, ChatEvent, ChatEventStream, ChatEventStreamTrait,
    CollectedAssistantMessage, FinishReason,
 };
-use vllm_engine_core_client::protocol::StopReason;
+use vllm_engine_core_client::protocol::output::StopReason;

 use self::convert::{ResponseOptions, prepare_chat_request};
 use crate::config::ApiServerOptions;
@@ -825,7 +825,7 @@ mod tests {
    use vllm_chat::{
        AssistantBlockKind, AssistantContentBlock, AssistantToolCall, ChatEvent, FinishReason,
    };
-    use vllm_engine_core_client::protocol::StopReason;
+    use vllm_engine_core_client::protocol::output::StopReason;
    use vllm_text::{DecodedLogprobs, DecodedPositionLogprobs, DecodedTokenLogprob};

    use super::{
@@ -199,6 +199,7 @@ mod tests {
    use super::prepare_completion_request;
    use crate::lora::LoraModelResolution;
    use crate::routes::openai::completions::types::CompletionRequest;
+    use crate::routes::openai::utils::types::Normalizable;
    use crate::utils::{ResolvedRequestContext, resolve_request_context};

    fn request_context(headers: &HeaderMap, request_id: Option<&str>) -> ResolvedRequestContext {
@@ -249,6 +250,28 @@ mod tests {
        assert!(request.ignore_eos);
    }

+    #[test]
+    fn normalize_coerces_null_max_tokens_to_default() {
+        // An absent `max_tokens` already gets the serde default.
+        let absent: CompletionRequest =
+            serde_json::from_value(base_request_json()).expect("parse request");
+        assert_eq!(absent.max_tokens, Some(16));
+
+        // An explicit `null` deserializes to `None`, bypassing the default;
+        // `normalize` must coerce it back to match Python vLLM.
+        let mut request: CompletionRequest = serde_json::from_value(json!({
+            "model": "Qwen/Qwen1.5-0.5B-Chat",
+            "prompt": "hello",
+            "stream": true,
+            "max_tokens": null
+        }))
+        .expect("parse request");
+        assert_eq!(request.max_tokens, None);
+
+        request.normalize();
+        assert_eq!(request.max_tokens, Some(16));
+    }
+
    #[test]
    fn prepare_completion_request_maps_sampling_fields() {
        let request: CompletionRequest = serde_json::from_value(json!({
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Zhuohan Li	0116e1cedc	[Core] Enable extensible KV cache for all attention backends and Mamba Port of internal D110967544. The extensible KV cache flow (reserve KV virtual address space up front, capture CUDA graphs first, then size and commit the KV cache from post-capture free memory) previously required a block-major attention backend and rejected Mamba models. This enables it for every backend layout and for Mamba / linear attention: - ExtensibleTensor gains num_segments: the reservation is divided into equal segments that grow in lockstep, with committed bytes forming a prefix of each segment. Physical pages are mapped at allocation-granularity granules and deduped across overlapping ranges, so a granule straddling a segment boundary is mapped exactly once. resize_per_segment_(bytes, zero_new=True) zeroes only the newly committed logical range of each segment. - Each KV cache buffer keeps its layers' physical layout and is committed as one prefix per layout segment. The segment count is derived from the backend's get_kv_cache_shape / get_kv_cache_block_dim / stride order: K/V-split layouts (e.g. FlashAttention) get one prefix per half, block-major layouts (e.g. FlashInfer, MLA) a single prefix. Mamba state pages are block-major per layer, and hybrid-model attention caches are re-strided to block-major, so both use a single segment. - Removed the supports_extensible_kv_cache gate plumbing from EngineCore, Executor, Worker, WorkerBase and GPUModelRunner; a CUDA platform check remains in EngineCore. - enable_extensible_kv_cache is reported as unsupported by the V2 model runner so V2-default models fall back to the V1 runner (which implements the flow); also fixed initialize_kv_cache being called with the extensible kwarg on runners that do not accept it, which broke every default V2-runner boot on this branch. Tested on H100: - tests/utils_/test_extensible_tensor.py (5 passed, incl. new segmented lockstep-grow/zero, granule-dedup and invalid-usage tests) - tests/v1/worker/test_extensible_kv_cache.py (new, 6 passed: segment derivation, split grows both halves, block-major, legacy full commit, Mamba per-layer growth, hybrid attention+Mamba) - E2E Qwen3-0.6B greedy with VLLM_ATTENTION_BACKEND=FLASH_ATTN (a K/V-split backend the old gate rejected): extensible generations byte-identical to the legacy path; log shows reserve then "Extended KV cache to 34663 blocks". V2->V1 auto-fallback path verified as well.	2026-07-10 17:08:34 -07:00
Zhuohan Li	80f66afd81	add new files	2026-07-01 17:32:19 -07:00
Zhuohan Li	1fa4c3adb7	[Core] Demo implementation of extensible kv cache memory	2026-07-01 17:30:53 -07:00
Nick HillandGitHub	4787f2dd1b	[Bugfix] Don't read KV cache past `seq_len` in triton paged attn kernels (#47305 )	2026-07-01 12:43:00 -07:00
Nick HillandGitHub	8cfeb84dba	[ModelRunner V2] Warmup cross-attn properly in encoder-decoder case (#47308 )	2026-07-01 12:36:48 -07:00
Chaitanya Sri Krishna LollaandGitHub	5fd442187c	[ROCm][P/D] MoRIIO toy proxy: support JSON Content-Type for OpenAI clients. (#46482 ) Signed-off-by: lcskrishna <lollachaitanya@gmail.com>	2026-07-01 19:17:05 +00:00
Yiliu Dong GitHub Nick Hill	00eb7cefa3	[Bugfix] Prevent padding placeholders from reaching embeddings (#47029 ) Signed-off-by: qianlihuang <91178480+qianlihuang@users.noreply.github.com> Signed-off-by: Yiliu Dong <91178480+qianlihuang@users.noreply.github.com> Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Nick Hill <nickhill123@gmail.com>	2026-07-01 09:26:03 -07:00
Michał GanczarenkoandGitHub	c8bdcc0116	[Bench][BugFix] Fix empty decoder prompt for Cohere ASR in throughput benchmark (#47135 ) Signed-off-by: Michal Ganczarenko <michal.ganczarenko@intel.com>	2026-07-01 15:42:27 +00:00
Benjamin Chislett GitHub Giancarlo Delfin mgoin	f5a8d73377	[Spec Decode] DSpark (#46995 ) Signed-off-by: Benjamin Chislett <bchislett@nvidia.com> Signed-off-by: Giancarlo Delfin <gdelfin@inferact.ai> Signed-off-by: mgoin <mgoin64@gmail.com> Co-authored-by: Giancarlo Delfin <gdelfin@inferact.ai> Co-authored-by: mgoin <mgoin64@gmail.com>	2026-07-01 08:30:24 -07:00
Michał Ganczarenko GitHub Kunshang Ji	63fcce4de1	[Bugfix] Fix GraniteMoeShared weight loading broken by #41184 (#47031 ) Signed-off-by: <Michal Ganczarenko> <michal.ganczarenko@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-07-01 22:39:12 +08:00
Bugen ZhaoandGitHub	c638f9216a	[Rust Frontend] Split engine core DTOs into separate modules (#47265 ) Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-01 15:28:21 +01:00
Chaojun ZhangandGitHub	13c49f9845	[xpu][lora]: Align LoRA implementation with Punica GPU: fix _apply_expand rank mismatch, add_inputs hardcode, and MoE EP (#45368 ) Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>	2026-07-01 22:14:04 +08:00
Nick HillandGitHub	f1cf6b0086	[CI] Fix segfault in tracing test (#47299 ) Signed-off-by: Nick Hill <nickhill123@gmail.com>	2026-07-01 14:00:37 +00:00
Harry MellorandGitHub	a78c15616f	Migrate GPTBigCode and Starcoder2 to the Transformers modeling backend (#30966 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-01 13:41:36 +00:00
Juan Pérez de Algaba GitHub Russell Bryant Russell Bryant	5c4db60f01	docs(security): document gRPC interface as insecure for private use only (#45903 ) Signed-off-by: jperezde <jperezde@redhat.com> Signed-off-by: Russell Bryant <russell.bryant@gmail.com> Co-authored-by: Russell Bryant <russell.bryant@gmail.com> Co-authored-by: Russell Bryant <rbryant@redhat.com>	2026-07-01 12:39:57 +00:00
Fangzhou Ai GitHub Claude	4e5ca89cfe	[ROCm][MiniMax-M3] Cross-layer lightning-indexer top-k sharing (#47269 ) Signed-off-by: Fangzhou Ai <fangzhouai@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-01 10:50:09 +00:00
Harry MellorandGitHub	a22e0dfc69	[Model] Remove AyaVision, MusicFlamingo (#47263 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-01 10:39:33 +00:00
stevenkuangandGitHub	cc56379e28	[Model] Support Hy3 token suffix and JSON Schema array types (#47192 ) Signed-off-by: stevenkuang-tencent <stevenkuang@tencent.com>	2026-07-01 10:16:07 +00:00
Aleksei Ivashov GitHub NickLucche	024b06b0dc	[Bugfix] Expose usage field in GenerateResponse for disaggregated serving (#42748 ) Signed-off-by: AIvashov <ivashov.aleksey@proton.me> Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai> Co-authored-by: NickLucche <nicolo.lucchesi@mistral.ai>	2026-07-01 10:00:19 +00:00
Harry MellorandGitHub	e7d0fcbc09	[CI] Fix various failures on `main` (#47197 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-01 10:35:34 +01:00
akii96andGitHub	aa8bb5562e	[ROCm][Perf][Bugfix] DSv4 indexer: use platform FP8 dtype (fnuz) for Q-quant on gfx942 (#46730 ) Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com>	2026-07-01 17:33:55 +08:00
Andy LoandGitHub	fa4bec9056	[Bugfix] Fix pooled Whisper sliding-window KV sizing (#47071 ) Signed-off-by: Andy Lo <andy@mistral.ai>	2026-07-01 11:33:19 +02:00
Jyothirmai Kottu GitHub Claude	dee5da1dec	[Test] Run SageMaker handler-override tests in-process via TestClient (#47250 ) Signed-off-by: Jyothirmai Kottu <jkottu@amazon.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-01 09:14:00 +00:00
Fangzhou Ai GitHub Cursor Claude	ed41aa270a	[ROCm][DSV4] Use aiter mHC pre/post as the default ROCm path (#43950 ) Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com> Signed-off-by: Fangzhou-Ai <fangzhouai@gmail.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-01 16:27:42 +08:00
Aaron Hao GitHub Claude Sonnet 4.6	77a9c5ae28	Weight sync refactor + move sparse nccl engine (#44353 ) Signed-off-by: hao-aaron <ahao@anyscale.com> Signed-off-by: haoaaron <ahao@anyscale.com> Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>	2026-07-01 01:25:19 -07:00
Yejing Lai GitHub Kunshang Ji	f651a8a9a4	[XPU][UT]Enable ut qk_norm_rope_fusion (#42486 ) Signed-off-by: Lai, Yejing <yejing.lai@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-07-01 07:38:03 +00:00
Jee Jee LiandGitHub	8f82be5705	[CI/Build] Fix LoRA testing (#47242 ) Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>	2026-07-01 15:36:13 +08:00
Nils MattesonandGitHub	a461070d1c	[Core] Make sleep-mode backend capability flags communicator-agnostic (#47243 )	2026-07-01 07:17:44 +00:00
Tiezhen WANG GitHub Claude	4470ae84de	Remove mantis (#46806 ) Signed-off-by: Xianbao QIAN <xianbao.qian@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-01 07:13:58 +00:00
ChaunceyandGitHub	697c34b97b	[Bugfix] Fix beam search candidate indexing when logprobs count varies (#47126 ) Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>	2026-07-01 07:07:06 +00:00
Blas Rodriguez IrizarandGitHub	5b431b905c	[Rust Frontend] Coerce completion `max_tokens: null` to default (#47166 ) Signed-off-by: Blas Rodriguez Irizar <rodrigblas@gmail.com>	2026-07-01 06:41:33 +00:00
almayne GitHubLi, Jiang <jiang1.li@intel.com>	89e99202f2	[CPU][Perf]Added tanh AOR for faster gelu activations. (#44639 ) Signed-off-by: Anna Mayne <anna.mayne@arm.com> Signed-off-by: almayne <anna.mayne@arm.com> Co-authored-by: Li, Jiang <jiang1.li@intel.com>	2026-06-30 23:24:40 -07:00
Micah WilliamsonandGitHub	b446792306	[ROCm][Bugfix] Fix Triton "out of resource: shared memory" Error In One-Shot LoRA MoE (#47209 ) Signed-off-by: Micah Williamson <micah.williamson@amd.com>	2026-06-30 23:24:36 -07:00
Micah WilliamsonandGitHub	c3b1f9e827	[ROCm][CI] Enable LoRA TP Distributed Test Group In AMD CI (#47193 ) Signed-off-by: Micah Williamson <micah.williamson@amd.com>	2026-06-30 23:24:32 -07:00
Jonathan MamouandGitHub	df802a87b7	[CPU] Remove speculative decoding stream overrides from CPUModelRunner (#47162 ) Signed-off-by: jmamou <jonathan.mamou@intel.com>	2026-07-01 06:12:49 +00:00
Nils MattesonandGitHub	93d8f834dd	[Core] Pluggable sleep-mode backend abstraction (RFC #34303 ) (#44074 )	2026-06-30 22:00:53 -07:00
Maria GuevaraandGitHub	aeb35b90f0	[Rust Frontend] Add error context in tool parser failures (#46512 ) Signed-off-by: Maria Guevara <kawaiiplush14@gmail.com>	2026-07-01 12:48:55 +08:00
Gabriel WuandGitHub	9a08a5118e	fix: skip cooperative top-K on SM120 (#47164 ) Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>	2026-06-30 21:32:54 -07:00
Jiangyun Zhu GitHub Jingyi Yang Lucas Wilkinson Claude Opus 4.8	c5200d3565	[Attention][DSA] support dcp for FLASHINFER_MLA_SPARSE (#46076 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Signed-off-by: Jingyi Yang <girasoleyang@gmail.com> Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Signed-off-by: GirasoleY <girasoleyang@gmail.com> Co-authored-by: Jingyi Yang <girasoleyang@gmail.com> Co-authored-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>	2026-07-01 00:32:20 -04:00
MattandGitHub	3c1396bab6	[Hardware][AMD][CI] Toggle test coredumps on ROCm debug agent (#47222 ) Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>	2026-06-30 23:30:10 -05:00
Benjamin ChislettandGitHub	9969466a59	[Spec Decode] Support SWA + DFlash for MiMo (#46104 )	2026-06-30 20:34:47 -07:00
achyuthan.sandGitHub	3406e8f83d	[Bugfix][Frontend][gpt-oss] Return raw output when Harmony parser ends non-terminal (#47062 ) Signed-off-by: Achyuthan Sivasankar <achyuthan.sivasankar@gmail.com>	2026-07-01 01:46:01 +00:00
Woosuk Kwon GitHub Claude Opus 4.8	a264e41975	[Distributed] Default FlashInfer allreduce to mnnvl on single node (#47219 ) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>	2026-06-30 18:35:56 -07:00
Woosuk KwonandGitHub	f098ee70c7	[GLM5] Support FlashMLA FP8 KV cache (Hopper & Blackwell) (#47090 ) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>	2026-06-30 18:13:21 -07:00
hcl GitHub Flora Feng	9294dd27eb	fix(reasoning): guard rfind in ernie45 streaming </response> branch (#46255 ) Signed-off-by: Chenglun Hu <chenglunhu@gmail.com> Co-authored-by: Flora Feng <4florafeng@gmail.com>	2026-07-01 01:01:14 +00:00
yzong-rhandGitHub	b1190d03cc	[Refactor][GPT-OSS] Harmony Responses API Refactor to use HarmonyParser (#47185 ) Signed-off-by: Yifan Zong <yzong@redhat.com>	2026-06-30 19:23:20 -04:00
Albert Cheng GitHub Vadim Gimpelson	92c7fac640	[Perf] Restore zero-init of swizzled NVFP4 scale buffer to recover Blackwell decode throughput (#45739 ) Signed-off-by: Albert Cheng <albertching0112@gmail.com> Co-authored-by: Vadim Gimpelson <156319763+vadiklyutiy@users.noreply.github.com>	2026-06-30 22:56:56 +00:00
Ting SUNandGitHub	ac521f6237	[Bugfix][Structured Outputs] Reject degenerate `structured_outputs` that crash EngineCore (#45346 ) Signed-off-by: Ting Sun <suntcrick@gmail.com>	2026-06-30 22:41:33 +00:00
tarjan1 GitHub Ben Browning	28242824e0	[Bugfix][Frontend] Normalize constrained Harmony recipients (#45657 ) Signed-off-by: shaojunjie <626650687@qq.com> Co-authored-by: Ben Browning <bbrownin@redhat.com>	2026-06-30 17:33:10 -04:00
VectorPeakandGitHub	68294739d1	[Bugfix] Align OpenCV video metadata timeline (#47099 ) Signed-off-by: VectorPeak <73048950+VectorPeak@users.noreply.github.com>	2026-06-30 20:43:42 +00:00
Joe Rowell GitHub Robert Shaw	c8d2f3cb14	[Bugfix] compressed-tensors: allow int8 grouped WNA16 MoE on Marlin (#47154 ) Signed-off-by: Joe Rowell <joerowell4@gmail.com> Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>	2026-06-30 12:50:46 -07:00
MattandGitHub	345b28ff2f	[Hardware][AMD][CI] Bump timeouts of various test groups on AMD CI (#47195 ) Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>	2026-06-30 14:30:53 -05:00
Roberto L. Castro GitHub Lucas Wilkinson	248d1fbb71	[Feat][1/N] CuTeDSL warmup infrastructure, FA4 MLA (#46182 ) Signed-off-by: LopezCastroRoberto <rocastro@redhat.com> Signed-off-by: Roberto L. Castro <38211239+LopezCastroRoberto@users.noreply.github.com> Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com>	2026-06-30 12:17:34 -07:00
Joe Rowell GitHub Robert Shaw	11b26c5528	[Bugfix][Tool Parser] PoolsideV1: fix logprobs AttributeError on Responses API (#47138 ) Signed-off-by: Joe Rowell <joerowell4@gmail.com> Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>	2026-06-30 19:14:09 +00:00
Roberto L. CastroandGitHub	20434c472e	[Feat] Improve Triton JIT diagnostics (#46621 ) Signed-off-by: LopezCastroRoberto <rocastro@redhat.com>	2026-06-30 18:50:15 +00:00
Andreas KaratzasandGitHub	c8f9c156a5	[ROCm][V1][MLA] Clone prefill backend state per metadata builder (#46993 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-06-30 11:43:54 -07:00
Woosuk Kwon GitHub snordmann	953bba488d	[PERF] Extend NCCL symmetric memory to AllGather and ReduceScatter (#46703 ) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai> Co-authored-by: snordmann <snordmann@nvidia.com>	2026-06-30 11:38:18 -07:00
Wentao YeandGitHub	3a9784b82c	[Feature] DP supervisor using rust frontend (#47076 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-30 14:34:05 -04:00
Giancarlo DelfinandGitHub	3cecee40f3	[Model Runner V2][Spec Decode] Fix stale values in idx_mapping from CG num reqs padding (#47066 )	2026-06-30 11:25:32 -07:00
Nikita Shapovalov GitHub Zhanda Zhu Shang Wang Michael Goin	a7732537f4	[Bugfix] Restore part of bugfix #42650 after accidental deletion in #43241 (#47039 ) Signed-off-by: zhanda <zhandazhu@gmail.com> Signed-off-by: Nikita Shapovalov <nikita@poolside.ai> Co-authored-by: Zhanda Zhu <49645678+zhandaz@users.noreply.github.com> Co-authored-by: Shang Wang <shangw@nvidia.com> Co-authored-by: Michael Goin <mgoin64@gmail.com>	2026-06-30 11:07:59 -07:00
Rishi Puri GitHub Anshika Ojha Claudegemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>Stefano Castagnetta	727971f1c1	Add Medusa speculative decoding e2e test (#41396 ) Signed-off-by: Anshika Ojha <anshikao@nvidia.com> Signed-off-by: Rishi Puri <riship@nvidia.com> Signed-off-by: Rishi Puri <puririshi98@berkeley.edu> Signed-off-by: Stefano Castagnetta <scastagnetta@nvidia.com> Co-authored-by: Anshika Ojha <215760622+ojhaanshika@users.noreply.github.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> Co-authored-by: Stefano Castagnetta <scastagnetta@nvidia.com>	2026-06-30 18:02:22 +00:00
Ben Browning GitHub Claude Opus 4.6	25671cb520	[Parser][Bugfix] Ensure tool call or other special tokens don't leak in non-streaming tool parsing (#46875 ) Signed-off-by: Ben Browning <bbrownin@redhat.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>	2026-06-30 13:46:53 -04:00
Lucas Wilkinson GitHub OpenAI Codexmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	27d5f78b63	[CI] Move distributed small LM eval to B200 (#47048 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-06-30 13:34:25 -04:00
liuzhenweiandGitHub	7a341fa109	[XPU] Support ZE_AFFINITY_MASK passthrough in xpu_disagg_acc_test (#47105 ) Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com>	2026-06-30 17:06:12 +00:00
Charlie FuandGitHub	f41e8ddc97	[ROCm][CI] Move PyTorch Compilation Unit Tests to MI300(gfx942) (#47065 ) Signed-off-by: charlifu <charlifu@amd.com>	2026-06-30 11:32:58 -05:00
fangyuchu GitHub Tyler Michael Smithmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	245888ff77	[Feature] Detect all2all peer fault with fault tolerance backend and prevent corrupted output (#43637 ) Signed-off-by: fangyuchu <fangyuchu@qq.com> Co-authored-by: Tyler Michael Smith <tlrmchlsmth@gmail.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-06-30 09:00:25 -07:00
Kunshang Ji GitHub Harry Mellor	e840f0d3f5	[Platform] Replace `torch.cuda.Event` with `torch.Event` (#47140 ) Signed-off-by: Kunshang Ji <kunshang.ji@intel.com> Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-06-30 08:39:59 -07:00
Nick Hill GitHub ziminghuang	fcaa84efa7	[BugFix] Gate MRV2 mixed sparse-MLA warmup on `max_num_seqs` > 1 (#47050 ) Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: ziminghuang <ziminghuang@inferact.ai>	2026-06-30 16:31:27 +01:00
Wentao YeandGitHub	9e84ec8648	[Refactor] Remove dead minimax allreduce rms kernel (#46842 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-30 08:29:21 -07:00