minor refactor

Signed-off-by: Bugen Zhao <i@bugenzhao.com>

.buildkite/scripts/ci-bake-rocm.sh

@@ -17,7 +17,7 @@ DEFAULT_REPO_SLUG="vllm-project/vllm"
 DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
 DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tools/install_protoc.sh rust-toolchain.toml tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
 DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
-DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_rixl build_rocshmem build_deepep mori_base ci_base"
+DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_nixl build_rocshmem build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="1"
 IMAGE_EXISTED_BEFORE_BUILD=0
 
@@ -1159,8 +1159,8 @@ ci_base_metadata_pairs() {
     metadata_pair "vllm.rocm.nic_backend" "$(resolve_dockerfile_arg_value "${dockerfile}" "NIC_BACKEND")"
     metadata_pair "vllm.rocm.ainic_version" "$(resolve_dockerfile_arg_value "${dockerfile}" "AINIC_VERSION")"
     metadata_pair "vllm.rocm.ubuntu_codename" "$(resolve_dockerfile_arg_value "${dockerfile}" "UBUNTU_CODENAME")"
-    metadata_pair "vllm.rocm.rixl_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_REPO")"
-    metadata_pair "vllm.rocm.rixl_commit" "${RIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_BRANCH")}"
+    metadata_pair "vllm.rocm.nixl_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "NIXL_REPO")"
+    metadata_pair "vllm.rocm.nixl_commit" "${NIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "NIXL_BRANCH")}"
     metadata_pair "vllm.rocm.ucx_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_REPO")"
     metadata_pair "vllm.rocm.ucx_commit" "${UCX_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_BRANCH")}"
     metadata_pair "vllm.rocm.rocshmem_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "ROCSHMEM_REPO")"
@@ -1169,7 +1169,7 @@ ci_base_metadata_pairs() {
     metadata_pair "vllm.rocm.deepep_commit" "${DEEPEP_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_BRANCH")}"
     metadata_pair "vllm.rocm.deepep_nic" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_NIC")"
     metadata_pair "vllm.rocm.deepep_rocm_arch" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_ROCM_ARCH")"
-    metadata_pair "vllm.rocm.rixl_cache_key" "${RIXL_CACHE_KEY:-}"
+    metadata_pair "vllm.rocm.nixl_cache_key" "${NIXL_CACHE_KEY:-}"
     metadata_pair "vllm.rocm.rocshmem_cache_key" "${ROCSHMEM_CACHE_KEY:-}"
     metadata_pair "vllm.rocm.deepep_cache_key" "${DEEPEP_CACHE_KEY:-}"
 
@@ -1686,7 +1686,7 @@ extract_dependency_pins() {
         return 0
     fi
 
-    for var in RIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
+    for var in NIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
         if [[ -n "${!var:-}" ]]; then
             echo "Using provided ${var}: ${!var}"
             continue
@@ -1706,30 +1706,30 @@ extract_dependency_pins() {
 compute_dependency_cache_keys() {
     local bake_dir=""
     local dockerfile_rocm=""
-    local rixl_branch=""
+    local nixl_branch=""
     local ucx_branch=""
     local rocshmem_branch=""
     local deepep_branch=""
-    local rixl_material=""
+    local nixl_material=""
     local rocshmem_material=""
     local deepep_material=""
 
     bake_dir=$(dirname "${VLLM_BAKE_FILE}")
     dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
-    rixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "RIXL_BRANCH")
+    nixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "NIXL_BRANCH")
     ucx_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "UCX_BRANCH")
     rocshmem_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "ROCSHMEM_BRANCH")
     deepep_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "DEEPEP_BRANCH")
 
-    if [[ -n "${rixl_branch}" && -n "${ucx_branch}" ]]; then
-        rixl_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_rixl")
-        RIXL_CACHE_KEY=$(
+    if [[ -n "${nixl_branch}" && -n "${ucx_branch}" ]]; then
+        nixl_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_nixl")
+        NIXL_CACHE_KEY=$(
             compose_dependency_cache_key \
-                "${rixl_branch}-ucx-${ucx_branch}" \
-                "${rixl_material}"
+                "${nixl_branch}-ucx-${ucx_branch}" \
+                "${nixl_material}"
         )
-        export RIXL_CACHE_KEY
-        echo "RIXL dependency cache key: ${RIXL_CACHE_KEY}"
+        export NIXL_CACHE_KEY
+        echo "NIXL dependency cache key: ${NIXL_CACHE_KEY}"
     fi
 
     if [[ -n "${rocshmem_branch}" ]]; then
@@ -1780,11 +1780,11 @@ dependency_cache_ref_for_target() {
     local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
 
     case "${target}" in
-        rixl-rocm-ci)
-            if [[ -n "${RIXL_CACHE_KEY:-}" ]]; then
-                printf '%s\n' "${cache_repo}:rixl-rocm-${RIXL_CACHE_KEY}"
-            elif [[ -n "${RIXL_BRANCH:-}" ]]; then
-                printf '%s\n' "${cache_repo}:rixl-rocm-${RIXL_BRANCH}-ucx-${UCX_BRANCH:-}"
+        nixl-rocm-ci)
+            if [[ -n "${NIXL_CACHE_KEY:-}" ]]; then
+                printf '%s\n' "${cache_repo}:nixl-rocm-${NIXL_CACHE_KEY}"
+            elif [[ -n "${NIXL_BRANCH:-}" ]]; then
+                printf '%s\n' "${cache_repo}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH:-}"
             fi
             ;;
         rocshmem-rocm-ci)
@@ -1815,7 +1815,7 @@ add_dependency_cache_target() {
 
 resolve_ci_base_dependency_targets() {
     local mode="${ROCM_DEP_CACHE_EXPORT_MODE:-missing}"
-    local rixl_ref=""
+    local nixl_ref=""
     local rocshmem_ref=""
     local deepep_ref=""
 
@@ -1824,7 +1824,7 @@ resolve_ci_base_dependency_targets() {
     case "${mode}" in
         always)
             echo "ROCM_DEP_CACHE_EXPORT_MODE=always; exporting all dependency caches serially"
-            for target in rixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
+            for target in nixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
                 if [[ -n "$(dependency_cache_ref_for_target "${target}")" ]]; then
                     add_dependency_cache_target "${target}"
                 fi
@@ -1844,13 +1844,13 @@ resolve_ci_base_dependency_targets() {
             ;;
     esac
 
-    if [[ "${mode}" != "always" && -n "${RIXL_CACHE_KEY:-}" ]]; then
-        rixl_ref=$(dependency_cache_ref_for_target "rixl-rocm-ci")
-        if dependency_cache_ref_exists "${rixl_ref}"; then
-            echo "RIXL dependency cache exists: ${rixl_ref}"
+    if [[ "${mode}" != "always" && -n "${NIXL_CACHE_KEY:-}" ]]; then
+        nixl_ref=$(dependency_cache_ref_for_target "nixl-rocm-ci")
+        if dependency_cache_ref_exists "${nixl_ref}"; then
+            echo "NIXL dependency cache exists: ${nixl_ref}"
         else
-            echo "RIXL dependency cache missing; will seed: ${rixl_ref}"
-            add_dependency_cache_target "rixl-rocm-ci"
+            echo "NIXL dependency cache missing; will seed: ${nixl_ref}"
+            add_dependency_cache_target "nixl-rocm-ci"
         fi
     fi
 

.buildkite/test_areas/basic_correctness.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
   key: basic-correctness
-  timeout_in_minutes: 45
+  timeout_in_minutes: 68
   device: h200_18gb
   source_file_dependencies:
   - vllm/

.buildkite/test_areas/benchmarks.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
   key: benchmarks-cli-test
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
   - vllm/

.buildkite/test_areas/cuda.yaml

@@ -26,7 +26,10 @@ steps:
   - vllm/v1/cudagraph_dispatcher.py
   - vllm/config/compilation.py
   - vllm/compilation
+  - vllm/v1/worker/encoder_cudagraph.py
+  - vllm/v1/worker/encoder_cudagraph_defs.py
   commands:
     - pytest -v -s v1/cudagraph/test_cudagraph_dispatch.py
     - pytest -v -s v1/cudagraph/test_cudagraph_mode.py
     - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
+    - pytest -v -s v1/cudagraph/test_encoder_cudagraph.py

.buildkite/test_areas/engine.yaml

@@ -51,7 +51,7 @@ steps:
 
 - label: e2e Scheduling (1 GPU)
   key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 35
+  timeout_in_minutes: 53
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/

.buildkite/test_areas/entrypoints.yaml

@@ -39,7 +39,7 @@ steps:
 - label: Entrypoints Integration (API Server)
   key: entrypoints-integration-api-server
   device: h200_35gb
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -59,7 +59,7 @@ steps:
 - label: Entrypoints Integration (API Server OpenAI - Part 1)
   device: h200_35gb
   key: entrypoints-integration-api-server-openai-part-1
-  timeout_in_minutes: 45
+  timeout_in_minutes: 68
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -78,7 +78,7 @@ steps:
 - label: Entrypoints Integration (API Server OpenAI - Part 2)
   device: h200_35gb
   key: entrypoints-integration-api-server-openai-part-2
-  timeout_in_minutes: 45
+  timeout_in_minutes: 83
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -156,7 +156,7 @@ steps:
 - label: Entrypoints Integration (Pooling)
   device: h200_35gb
   key: entrypoints-integration-pooling
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/

.buildkite/test_areas/misc.yaml

@@ -31,7 +31,7 @@ steps:
 
 - label: V1 Sample + Logits
   key: v1-sample-logits
-  timeout_in_minutes: 45
+  timeout_in_minutes: 83
   device: h200_18gb
   source_file_dependencies:
     - vllm/config/
@@ -90,6 +90,7 @@ steps:
     - tests/v1/kv_offload
     - tests/v1/simple_kv_offload
     - tests/v1/worker
+    - tests/v1/streaming_input
     - tests/v1/kv_connector/unit
     - tests/v1/ec_connector/unit
     - tests/v1/metrics
@@ -103,6 +104,7 @@ steps:
     - pytest -v -s v1/kv_offload
     - pytest -v -s v1/simple_kv_offload
     - pytest -v -s v1/worker
+    - pytest -v -s v1/streaming_input
     - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/ec_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/metrics
@@ -143,6 +145,7 @@ steps:
     - pytest -v -s -m 'cpu_test' v1/core
     - pytest -v -s v1/structured_output
     - pytest -v -s v1/test_serial_utils.py
+    - pytest -v -s v1/cudagraph/test_cudagraph_manager.py
     - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'cpu_test' v1/metrics
 

.buildkite/test_areas/model_executor.yaml

@@ -5,7 +5,7 @@ steps:
 - label: Model Executor
   device: h200_35gb
   key: model-executor
-  timeout_in_minutes: 45
+  timeout_in_minutes: 60
   source_file_dependencies:
   - vllm/engine/arg_utils.py
   - vllm/config/model.py

.buildkite/test_areas/models_basic.yaml

@@ -46,6 +46,19 @@ steps:
       depends_on:
       - image-build-amd
 
+- label: Inkling Unit Tests (B200)
+  key: inkling-unit-tests-b200
+  timeout_in_minutes: 40
+  device: b200-k8s
+  source_file_dependencies:
+  - vllm/models/inkling/
+  - vllm/cute_utils/
+  - cmake/external_projects/tml_fa4.cmake
+  - tests/models/inkling/
+  commands:
+    # FA4 kernel tests require SM100; the suite skips them elsewhere.
+    - pytest -v -s models/inkling
+
 - label: Basic Models Test (Other CPU) # 5min
   key: basic-models-test-other-cpu
   depends_on:

.buildkite/test_areas/models_language.yaml

@@ -137,7 +137,7 @@ steps:
 
 - label: Language Models Test (MTEB)
   key: language-models-test-mteb
-  timeout_in_minutes: 45
+  timeout_in_minutes: 68
   device: h200_18gb
   optional: true
   source_file_dependencies:

.buildkite/test_areas/models_multimodal.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
   key: multi-modal-models-standard-1-qwen2
-  timeout_in_minutes: 45
+  timeout_in_minutes: 68
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -20,7 +20,7 @@ steps:
 
 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
   key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -54,7 +54,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
   device: h200_35gb
   key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
@@ -85,7 +85,7 @@ steps:
 
 - label: Multi-Modal Processor # 44min
   key: multi-modal-processor
-  timeout_in_minutes: 65
+  timeout_in_minutes: 98
   device: h200_18gb
   source_file_dependencies:
   - vllm/

.buildkite/test_areas/pytorch.yaml

@@ -5,7 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
   device: h200_35gb
   key: pytorch-compilation-unit-tests
-  timeout_in_minutes: 90
+  timeout_in_minutes: 110
   source_file_dependencies:
     - vllm/__init__.py
     - vllm/_aiter_ops.py

.buildkite/test_areas/spec_decode.yaml

@@ -170,3 +170,19 @@ steps:
     - tests/v1/e2e/spec_decode/
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "qwen3_5-hybrid"
+
+- label: Spec Decode DeepSeek MTP Parallel Load (B200)
+  key: spec-decode-deepseek-mtp-parallel-load-b200
+  timeout_in_minutes: 30
+  device: b200-k8s
+  optional: true
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/v1/spec_decode/llm_base_proposer.py
+    - vllm/v1/spec_decode/eagle.py
+    - vllm/v1/worker/gpu/spec_decode/eagle/
+    - vllm/model_executor/models/deepseek_mtp.py
+    - vllm/model_executor/models/deepseek_v2.py
+    - tests/v1/e2e/spec_decode/test_mtp_parallel_load.py
+  commands:
+    - pytest -v -s v1/e2e/spec_decode/test_mtp_parallel_load.py

.github/mergify.yml

@@ -181,6 +181,18 @@ pull_request_rules:
       add:
         - performance
 
+- name: label-quantization
+  description: Automatically apply quantization label
+  conditions:
+    - label != stale
+    - or:
+      - files~=^vllm/model_executor/layers/quantization/
+      - title~=(?i)quant
+  actions:
+    label:
+      add:
+        - quantization
+
 - name: label-qwen
   description: Automatically apply qwen label
   conditions:

.github/workflows/issue_autolabel.yml

@@ -130,6 +130,47 @@ jobs:
                   },
                 ],
               },
+              quantization: {
+                keywords: [
+                  {
+                    term: "quantization",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "quantized",
+                    searchIn: "both"
+                  },
+                ],
+              },
+              "intel-gpu": {
+                // Keyword search - matches whole words only (with word boundaries)
+                keywords: [
+                  {
+                    term: "B50",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "B60",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "B70",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "intel gpu",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "Arc GPU",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "BMG",
+                    searchIn: "both"
+                  },
+                ],
+              },
               // Add more label configurations here as needed
               // example: {
               //   keywords: [...],
@@ -491,4 +532,4 @@ jobs:
               issue_number: context.issue.number,
               body: message,
             });
-            core.notice(`Requested missing ROCm info from @${author}: ${missing.map(m => m.name).join(', ')}`);
+            core.notice(`Requested missing ROCm info from @${author}: ${missing.map(m => m.name).join(', ')}`);

cmake/external_projects/fmha_sm100.cmake

@@ -17,7 +17,7 @@ else()
   FetchContent_Declare(
     fmha_sm100
     GIT_REPOSITORY https://github.com/vllm-project/MSA.git
-    GIT_TAG 2e63ec37a0fc29bc20f39cd1a52e0f5affc33a73
+    GIT_TAG 890aaa1a37a598ad17ccff0827fea21540d381fa
     GIT_PROGRESS TRUE
     CONFIGURE_COMMAND ""
     BUILD_COMMAND ""

cmake/external_projects/vllm_flash_attn.cmake

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

csrc/libtorch_stable/cache_kernels.cu

@@ -1025,6 +1025,9 @@ __global__ void gather_and_maybe_dequant_cache(
     batch_offset += offset;
     int32_t block_table_id = batch_offset / block_size;
     int32_t slot_id = batch_offset % block_size;
+    // seq_starts may push the block index past the end of the batch's block
+    // table row.
+    if (block_table_id >= block_table_stride) continue;
     int32_t block_table_offset = batch_id * block_table_stride + block_table_id;
     int32_t block_id = block_table[block_table_offset];
     int64_t cache_offset =

csrc/libtorch_stable/moe/moe_ops.h

@@ -9,14 +9,16 @@ void topk_softmax(torch::stable::Tensor& topk_weights,
                   torch::stable::Tensor& topk_indices,
                   torch::stable::Tensor& token_expert_indices,
                   torch::stable::Tensor& gating_output, bool renormalize,
-                  std::optional<torch::stable::Tensor> bias);
+                  std::optional<torch::stable::Tensor> bias,
+                  std::optional<torch::stable::Tensor> is_padding);
 
 void topk_sigmoid(torch::stable::Tensor& topk_weights,
                   torch::stable::Tensor& topk_indices,
                   torch::stable::Tensor& token_expert_indices,
                   torch::stable::Tensor& gating_output, bool renormalize,
                   std::optional<torch::stable::Tensor> bias,
-                  double routed_scaling_factor);
+                  double routed_scaling_factor,
+                  std::optional<torch::stable::Tensor> is_padding);
 
 void topk_softplus_sqrt(
     torch::stable::Tensor& topk_weights, torch::stable::Tensor& topk_indices,
@@ -25,7 +27,8 @@ void topk_softplus_sqrt(
     double routed_scaling_factor,
     const std::optional<torch::stable::Tensor>& correction_bias,
     const std::optional<torch::stable::Tensor>& input_ids,
-    const std::optional<torch::stable::Tensor>& tid2eid);
+    const std::optional<torch::stable::Tensor>& tid2eid,
+    const std::optional<torch::stable::Tensor>& is_padding);
 
 void moe_sum(torch::stable::Tensor& input, torch::stable::Tensor& output,
              std::optional<torch::stable::Tensor> topk_ids,

csrc/libtorch_stable/moe/topk_softmax_kernels.cu

@@ -174,7 +174,8 @@ __launch_bounds__(TPB) __global__ void moeTopK(
     const int end_expert,
     const bool renormalize,
     const float* bias,
-    const double routed_scaling_factor)
+    const double routed_scaling_factor,
+    const bool* is_padding)
 {
 
     using cub_kvp = cub::KeyValuePair<int, float>;
@@ -228,12 +229,14 @@ __launch_bounds__(TPB) __global__ void moeTopK(
             const int expert = result_kvp.key;
             const bool node_uses_expert = expert >= start_expert && expert < end_expert;
             const bool should_process_row = row_is_active && node_uses_expert;
+            const bool is_pad_row = is_padding != nullptr && is_padding[block_row];
 
             const int idx = k * block_row + k_idx;
             // Return the unbiased scores for output weights
             output[idx] = inputs_after_softmax[thread_read_offset + expert];
-            indices[idx] = should_process_row ? (expert - start_expert) : num_experts;
-            assert(indices[idx] >= 0);
+            indices[idx] = is_pad_row ? static_cast<IndType>(-1)
+                                       : (should_process_row ? (expert - start_expert) : num_experts);
+            assert(is_pad_row || indices[idx] >= 0);
             source_rows[idx] = k_idx * num_rows + block_row;
             if (renormalize) {
                 selected_sum += inputs_after_softmax[thread_read_offset + expert];
@@ -277,7 +280,7 @@ template <int VPT, int NUM_EXPERTS, int WARPS_PER_CTA, int BYTES_PER_LDG, int WA
 __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
     void topkGating(const InputType* input, const bool* finished, float* output, const int num_rows, IndType* indices,
         int* source_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
-        const float* bias, const double routed_scaling_factor)
+        const float* bias, const double routed_scaling_factor, const bool* is_padding)
 {
     static_assert(std::is_same_v<InputType, float> || std::is_same_v<InputType, __nv_bfloat16> ||
                       std::is_same_v<InputType, __half>,
@@ -545,12 +548,14 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
             // Add a guard to ignore experts not included by this node
             const bool node_uses_expert = expert >= start_expert && expert < end_expert;
             const bool should_process_row = row_is_active && node_uses_expert;
+            const bool is_pad_row = is_padding != nullptr && is_padding[thread_row];
 
             // The lead thread from each sub-group will write out the final results to global memory. (This will be a
             // single) thread per row of the input/output matrices.
             const int idx = k * thread_row + k_idx;
             output[idx] = max_val;
-            indices[idx] = should_process_row ? (expert - start_expert) : NUM_EXPERTS;
+            indices[idx] = is_pad_row ? static_cast<IndType>(-1)
+                                       : (should_process_row ? (expert - start_expert) : NUM_EXPERTS);
             source_rows[idx] = k_idx * num_rows + thread_row;
             if (renormalize) {
                 selected_sum += max_val;
@@ -605,7 +610,7 @@ struct TopkConstants
 template <int EXPERTS, int WARPS_PER_TB, int WARP_SIZE_PARAM, int MAX_BYTES_PER_LDG, typename IndType, typename InputType, ScoringFunc SF>
 void topkGatingLauncherHelper(const InputType* input, const bool* finished, float* output, IndType* indices,
     int* source_row, const int num_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
-    const float* bias, const double routed_scaling_factor, cudaStream_t stream)
+    const float* bias, const double routed_scaling_factor, cudaStream_t stream, const bool* is_padding)
 {
     static constexpr int BYTES_PER_LDG = MIN(MAX_BYTES_PER_LDG, sizeof(InputType) * EXPERTS);
     using Constants = detail::TopkConstants<EXPERTS, BYTES_PER_LDG, WARP_SIZE_PARAM, InputType>;
@@ -616,7 +621,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
 
     dim3 block_dim(WARP_SIZE_PARAM, WARPS_PER_TB);
     topkGating<VPT, EXPERTS, WARPS_PER_TB, BYTES_PER_LDG, WARP_SIZE_PARAM, IndType, InputType, SF><<<num_blocks, block_dim, 0, stream>>>(
-        input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias, routed_scaling_factor);
+        input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias, routed_scaling_factor, is_padding);
 }
 
 #ifndef USE_ROCM
@@ -627,7 +632,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
                              IndType, InputType, SF>(                         \
         gating_output, nullptr, topk_weights, topk_indices,                   \
         token_expert_indices, num_tokens, topk, 0, num_experts, renormalize,  \
-        bias, routed_scaling_factor, stream);
+        bias, routed_scaling_factor, stream, is_padding);
 #else
   #define LAUNCH_TOPK(NUM_EXPERTS, WARPS_PER_TB, MAX_BYTES)                    \
     if (WARP_SIZE == 64) {                                                     \
@@ -635,13 +640,13 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
                                IndType, InputType, SF>(                        \
           gating_output, nullptr, topk_weights, topk_indices,                  \
           token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
-          bias, routed_scaling_factor, stream);                                                       \
+          bias, routed_scaling_factor, stream, is_padding);                                            \
     } else if (WARP_SIZE == 32) {                                              \
       topkGatingLauncherHelper<NUM_EXPERTS, WARPS_PER_TB, 32, MAX_BYTES,       \
                                IndType, InputType, SF>(                        \
           gating_output, nullptr, topk_weights, topk_indices,                  \
           token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
-          bias, routed_scaling_factor, stream);                                                       \
+          bias, routed_scaling_factor, stream, is_padding);                                            \
     } else {                                                                   \
       assert(false &&                                                          \
              "Unsupported warp size. Only 32 and 64 are supported for ROCm");  \
@@ -661,7 +666,8 @@ void topkGatingKernelLauncher(
     const bool renormalize,
     const float* bias,
     const double routed_scaling_factor,
-    cudaStream_t stream) {
+    cudaStream_t stream,
+    const bool* is_padding) {
     static constexpr int WARPS_PER_TB = 4;
     static constexpr int BYTES_PER_LDG_POWER_OF_2 = 16;
 #ifndef USE_ROCM
@@ -736,7 +742,7 @@ void topkGatingKernelLauncher(
             }
             moeTopK<TPB><<<num_tokens, TPB, 0, stream>>>(
                 workspace, nullptr, topk_weights, topk_indices, token_expert_indices,
-                num_experts, topk, 0, num_experts, renormalize, bias, routed_scaling_factor);
+                num_experts, topk, 0, num_experts, renormalize, bias, routed_scaling_factor, is_padding);
         }
     }
 }
@@ -755,7 +761,8 @@ void dispatch_topk_launch(
     int num_tokens, int num_experts, int topk, bool renormalize,
     std::optional<torch::stable::Tensor> bias,
     double routed_scaling_factor,
-    cudaStream_t stream)
+    cudaStream_t stream,
+    std::optional<torch::stable::Tensor> is_padding)
  {
     const float* bias_ptr = nullptr;
     if (bias.has_value()) {
@@ -769,6 +776,18 @@ void dispatch_topk_launch(
       bias_ptr = bias_tensor.const_data_ptr<float>();
     }
 
+    const bool* is_padding_ptr = nullptr;
+    if (is_padding.has_value()) {
+      const torch::stable::Tensor& is_padding_tensor = is_padding.value();
+      STD_TORCH_CHECK(is_padding_tensor.scalar_type() == torch::headeronly::ScalarType::Bool,
+                      "is_padding tensor must be bool");
+      STD_TORCH_CHECK(is_padding_tensor.dim() == 1, "is_padding tensor must be 1D");
+      STD_TORCH_CHECK(is_padding_tensor.size(0) == num_tokens,
+                      "is_padding size mismatch, expected: ", num_tokens);
+      STD_TORCH_CHECK(is_padding_tensor.is_contiguous(), "is_padding tensor must be contiguous");
+      is_padding_ptr = is_padding_tensor.const_data_ptr<bool>();
+    }
+
     if (topk_indices.scalar_type() == torch::headeronly::ScalarType::Int) {
         vllm::moe::topkGatingKernelLauncher<int, ComputeType, SF>(
             reinterpret_cast<const ComputeType*>(gating_output.const_data_ptr()),
@@ -777,7 +796,7 @@ void dispatch_topk_launch(
             token_expert_indices.mutable_data_ptr<int>(),
             softmax_workspace.mutable_data_ptr<float>(),
             num_tokens, num_experts, topk, renormalize,
-            bias_ptr, routed_scaling_factor, stream);
+            bias_ptr, routed_scaling_factor, stream, is_padding_ptr);
     } else if (topk_indices.scalar_type() == torch::headeronly::ScalarType::UInt32) {
         vllm::moe::topkGatingKernelLauncher<uint32_t, ComputeType, SF>(
             reinterpret_cast<const ComputeType*>(gating_output.const_data_ptr()),
@@ -786,7 +805,7 @@ void dispatch_topk_launch(
             token_expert_indices.mutable_data_ptr<int>(),
             softmax_workspace.mutable_data_ptr<float>(),
             num_tokens, num_experts, topk, renormalize,
-            bias_ptr, routed_scaling_factor, stream);
+            bias_ptr, routed_scaling_factor, stream, is_padding_ptr);
     } else {
         STD_TORCH_CHECK(topk_indices.scalar_type() == torch::headeronly::ScalarType::Long);
         vllm::moe::topkGatingKernelLauncher<int64_t, ComputeType, SF>(
@@ -796,7 +815,7 @@ void dispatch_topk_launch(
             token_expert_indices.mutable_data_ptr<int>(),
             softmax_workspace.mutable_data_ptr<float>(),
             num_tokens, num_experts, topk, renormalize,
-            bias_ptr, routed_scaling_factor, stream);
+            bias_ptr, routed_scaling_factor, stream, is_padding_ptr);
     }
 }
 
@@ -806,7 +825,8 @@ void topk_softmax(
     torch::stable::Tensor& token_expert_indices,        // [num_tokens, topk]
     torch::stable::Tensor& gating_output,               // [num_tokens, num_experts]
     bool renormalize,
-    std::optional<torch::stable::Tensor> bias)
+    std::optional<torch::stable::Tensor> bias,
+    std::optional<torch::stable::Tensor> is_padding)
 {
     const int num_experts = gating_output.size(-1);
     const auto num_tokens = gating_output.numel() / num_experts;
@@ -825,15 +845,15 @@ void topk_softmax(
     if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
         dispatch_topk_launch<float, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
             token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, 1.0, stream);
+            bias, 1.0, stream, is_padding);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
         dispatch_topk_launch<__half, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
             token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, 1.0, stream);
+            bias, 1.0, stream, is_padding);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
         dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
             token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, 1.0, stream);
+            bias, 1.0, stream, is_padding);
     } else {
         STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
     }
@@ -846,7 +866,8 @@ void topk_sigmoid(
     torch::stable::Tensor& gating_output,               // [num_tokens, num_experts]
     bool renormalize,
     std::optional<torch::stable::Tensor> bias,
-    double routed_scaling_factor)
+    double routed_scaling_factor,
+    std::optional<torch::stable::Tensor> is_padding)
 {
     const int num_experts = gating_output.size(-1);
     const auto num_tokens = gating_output.numel() / num_experts;
@@ -865,15 +886,15 @@ void topk_sigmoid(
     if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
         dispatch_topk_launch<float, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
             token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, routed_scaling_factor, stream);
+            bias, routed_scaling_factor, stream, is_padding);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
         dispatch_topk_launch<__half, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
             token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, routed_scaling_factor, stream);
+            bias, routed_scaling_factor, stream, is_padding);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
         dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
             token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, routed_scaling_factor, stream);
+            bias, routed_scaling_factor, stream, is_padding);
     } else {
         STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
     }

csrc/libtorch_stable/moe/topk_softplus_sqrt_kernels.cu

@@ -80,22 +80,27 @@ __launch_bounds__(128) __global__
                                   OutIndType* indices, int num_rows,
                                   int num_experts, float routed_scaling_factor,
                                   const HashIndType* input_ids,
-                                  const HashIndType* tid2eid) {
+                                  const HashIndType* tid2eid,
+                                  const bool* is_padding) {
   const int warp = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
   const int lane = threadIdx.x % 32;
   if (warp >= num_rows) return;
   const int64_t token_id = load_index_as_int64(input_ids, warp);
+  const bool is_pad_row = is_padding != nullptr && is_padding[warp];
 
   #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
   cudaGridDependencySynchronize();
   #endif
   int expert = 0;
   float weight = 0.f;
-  if (lane < 6) {
+  if (lane < 6 && !is_pad_row) {
     // only load and calculate for 6 experts
     expert = static_cast<int>(tid2eid[token_id * 6 + lane]);
     const float x = input[warp * num_experts + expert];
     weight = sqrtf(fmaxf(x, 0.f) + __logf(1.f + __expf(-fabsf(x))));
+    if (isnan(weight)) {
+      weight = 0.f;
+    }
   }
   float weight_sum = weight;
   #pragma unroll
@@ -111,7 +116,8 @@ __launch_bounds__(128) __global__
     const int offset = warp * 6 + lane;
     output[offset] =
         weight * routed_scaling_factor / (weight_sum > 0.f ? weight_sum : 1.f);
-    indices[offset] = static_cast<OutIndType>(expert);
+    indices[offset] = !is_pad_row ? static_cast<OutIndType>(expert)
+                                  : static_cast<OutIndType>(-1);
   }
 }
 
@@ -120,7 +126,8 @@ void launchDsv4HashTopk(const float* input, float* output, OutIndType* indices,
                         int num_rows, int num_experts,
                         double routed_scaling_factor,
                         const HashIndType* input_ids,
-                        const HashIndType* tid2eid, cudaStream_t stream) {
+                        const HashIndType* tid2eid, cudaStream_t stream,
+                        const bool* is_padding) {
   if (num_rows == 0) return;
   auto* kernel = &dsv4HashTopkSoftplusSqrt<OutIndType, HashIndType>;
   cudaLaunchConfig_t config = {};
@@ -134,7 +141,7 @@ void launchDsv4HashTopk(const float* input, float* output, OutIndType* indices,
   config.numAttrs = 1;
   const float scale = static_cast<float>(routed_scaling_factor);
   cudaLaunchKernelEx(&config, kernel, input, output, indices, num_rows,
-                     num_experts, scale, input_ids, tid2eid);
+                     num_experts, scale, input_ids, tid2eid, is_padding);
 }
 #endif
 
@@ -166,7 +173,8 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
         const int num_rows, IndType* indices, int* source_rows, const int k,
         const int start_expert, const int end_expert, const bool renormalize,
         double routed_scaling_factor, const float* correction_bias,
-        const HashIndType* input_ids, const HashIndType* tid2eid) {
+        const HashIndType* input_ids, const HashIndType* tid2eid,
+        const bool* is_padding) {
   static_assert(std::is_same_v<InputType, float> ||
                     std::is_same_v<InputType, __nv_bfloat16> ||
                     std::is_same_v<InputType, __half>,
@@ -231,6 +239,7 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
     return;
   }
   const bool row_is_active = finished ? !finished[thread_row] : true;
+  const bool is_pad_row = is_padding != nullptr && is_padding[thread_row];
 
   // We finally start setting up the read pointers for each thread. First, each
   // thread jumps to the start of the row it will read.
@@ -249,9 +258,12 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
   cudaGridDependencySynchronize();
 #endif
 
-  // NOTE(zhuhaoran): dispatch different input types loading, BF16/FP16 convert
-  // to float
-  if constexpr (std::is_same_v<InputType, float>) {
+  if (is_pad_row) {
+#pragma unroll
+    for (int ii = 0; ii < VPT; ++ii) {
+      row_chunk[ii] = 0.f;
+    }
+  } else if constexpr (std::is_same_v<InputType, float>) {
     using VecType = AlignedArray<float, ELTS_PER_LDG>;
     VecType* row_chunk_vec_ptr = reinterpret_cast<VecType*>(&row_chunk);
     const VecType* vec_thread_read_ptr =
@@ -315,12 +327,22 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
   if constexpr (USE_HASH) {
     const int64_t token_id = load_index_as_int64(input_ids, thread_row);
     const int64_t token_expert_offset = token_id * static_cast<int64_t>(k);
+    if (!is_pad_row) {
 #pragma unroll
-    for (int ii = 0; ii < VPT; ++ii) {
-      float val = row_chunk[ii];
-      float val_b = val * beta;
-      val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
-      row_chunk[ii] = sqrtf(val);
+      for (int ii = 0; ii < VPT; ++ii) {
+        float val = row_chunk[ii];
+        float val_b = val * beta;
+        val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
+        val = sqrtf(val);
+
+        // Dummy/padding tokens can result in NaN values, so
+        // clamp them to 0.0. Note: this clamp could likely be removed if
+        // 'is_padding' is made mandatory
+        if (isnan(val)) {
+          val = 0.f;
+        }
+        row_chunk[ii] = val;
+      }
     }
     float selected_sum = 0.f;
 #pragma unroll
@@ -335,7 +357,8 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
                                group_id * THREADS_PER_ROW * ELTS_PER_LDG +
                                local_id;
         if (expert == expert_idx) {
-          indices[idx] = static_cast<IndType>(expert);
+          indices[idx] = !is_pad_row ? static_cast<IndType>(expert)
+                                     : static_cast<IndType>(-1);
           selected_sum += row_chunk[ii];
           break;
         }
@@ -379,23 +402,31 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
 #endif
     return;
   } else {
+    if (!is_pad_row) {
 #pragma unroll
-    for (int ii = 0; ii < VPT; ++ii) {
-      float val = row_chunk[ii];
-      float val_b = val * beta;
-      // Compute softplus: log(1 + exp(val)) with numerical stability
-      // When val > threshold, softplus(x) ≈ x to avoid exp overflow
-      val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
-      val = sqrtf(val);
-      if (correction_bias) {
-        const int group_id = ii / ELTS_PER_LDG;
-        const int local_id = ii % ELTS_PER_LDG;
-        const int expert_idx = first_elt_read_by_thread +
-                               group_id * THREADS_PER_ROW * ELTS_PER_LDG +
-                               local_id;
-        val = val + correction_bias[expert_idx];
+      for (int ii = 0; ii < VPT; ++ii) {
+        float val = row_chunk[ii];
+        float val_b = val * beta;
+        // Compute softplus: log(1 + exp(val)) with numerical stability
+        // When val > threshold, softplus(x) ≈ x to avoid exp overflow
+        val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
+        val = sqrtf(val);
+        // Dummy/padding tokens can result in NaN values, so
+        // clamp them to 0.0. Note: this clamp could likely be removed if
+        // 'is_padding' is made mandatory
+        if (isnan(val)) {
+          val = 0.f;
+        }
+        if (correction_bias) {
+          const int group_id = ii / ELTS_PER_LDG;
+          const int local_id = ii % ELTS_PER_LDG;
+          const int expert_idx = first_elt_read_by_thread +
+                                 group_id * THREADS_PER_ROW * ELTS_PER_LDG +
+                                 local_id;
+          val = val + correction_bias[expert_idx];
+        }
+        row_chunk[ii] = val;
       }
-      row_chunk[ii] = val;
     }
 
     // Original TopK path: find top-k experts by score
@@ -450,18 +481,19 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
         // Add a guard to ignore experts not included by this node
         const bool node_uses_expert =
             expert >= start_expert && expert < end_expert;
-        const bool should_process_row = row_is_active && node_uses_expert;
+        const bool should_process_row =
+            row_is_active && node_uses_expert && !is_pad_row;
 
         // The lead thread from each sub-group will write out the final results
         // to global memory. (This will be a single) thread per row of the
         // input/output matrices.
         const int idx = k * thread_row + k_idx;
-        if (correction_bias != nullptr) {
+        if (correction_bias != nullptr && should_process_row) {
           max_val -= correction_bias[expert];
         }
         output[idx] = max_val;
         indices[idx] =
-            should_process_row ? (expert - start_expert) : NUM_EXPERTS;
+            !is_pad_row ? expert - start_expert : static_cast<IndType>(-1);
         source_rows[idx] = k_idx * num_rows + thread_row;
         if (renormalize) {
           selected_sum += max_val;
@@ -544,7 +576,7 @@ void topkGatingSoftplusSqrtLauncherHelper(
     const int start_expert, const int end_expert, const bool renormalize,
     double routed_scaling_factor, const float* correction_bias,
     const bool use_hash, const HashIndType* input_ids,
-    const HashIndType* tid2eid, cudaStream_t stream) {
+    const HashIndType* tid2eid, cudaStream_t stream, const bool* is_padding) {
   static constexpr int BYTES_PER_LDG =
       MIN(MAX_BYTES_PER_LDG, sizeof(InputType) * EXPERTS);
   using Constants =
@@ -573,12 +605,12 @@ void topkGatingSoftplusSqrtLauncherHelper(
     cudaLaunchKernelEx(&config, kernel, input, finished, output, num_rows,
                        indices, source_row, k, start_expert, end_expert,
                        renormalize, routed_scaling_factor, correction_bias,
-                       input_ids, tid2eid);
+                       input_ids, tid2eid, is_padding);
 #else
     kernel<<<num_blocks, block_dim, 0, stream>>>(
         input, finished, output, num_rows, indices, source_row, k, start_expert,
         end_expert, renormalize, routed_scaling_factor, correction_bias,
-        input_ids, tid2eid);
+        input_ids, tid2eid, is_padding);
 #endif
   })
 }
@@ -592,7 +624,7 @@ void topkGatingSoftplusSqrtLauncherHelper(
         gating_output, nullptr, topk_weights, topk_indices,                    \
         token_expert_indices, num_tokens, topk, 0, num_experts, renormalize,   \
         routed_scaling_factor, correction_bias, use_hash, input_ids, tid2eid,  \
-        stream);
+        stream, is_padding);
 #else
   #define LAUNCH_SOFTPLUS_SQRT(NUM_EXPERTS, WARPS_PER_TB, MAX_BYTES)           \
     if (WARP_SIZE == 64) {                                                     \
@@ -601,14 +633,14 @@ void topkGatingSoftplusSqrtLauncherHelper(
           gating_output, nullptr, topk_weights, topk_indices,                  \
           token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
           routed_scaling_factor, correction_bias, use_hash, input_ids,         \
-          tid2eid, stream);                                                    \
+          tid2eid, stream, is_padding);                                        \
     } else if (WARP_SIZE == 32) {                                              \
       topkGatingSoftplusSqrtLauncherHelper<NUM_EXPERTS, WARPS_PER_TB, 32,      \
                                            MAX_BYTES>(                         \
           gating_output, nullptr, topk_weights, topk_indices,                  \
           token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
           routed_scaling_factor, correction_bias, use_hash, input_ids,         \
-          tid2eid, stream);                                                    \
+          tid2eid, stream, is_padding);                                        \
     } else {                                                                   \
       assert(false &&                                                          \
              "Unsupported warp size. Only 32 and 64 are supported for ROCm");  \
@@ -622,14 +654,14 @@ void topkGatingSoftplusSqrtKernelLauncher(
     const int topk, const bool renormalize, double routed_scaling_factor,
     const float* correction_bias, const bool use_hash,
     const HashIndType* input_ids, const HashIndType* tid2eid,
-    cudaStream_t stream) {
+    cudaStream_t stream, const bool* is_padding) {
 #ifndef USE_ROCM
   if constexpr (std::is_same_v<InputType, float>) {
     if (use_hash && topk == 6 && renormalize &&
         (num_experts == 256 || num_experts == 384)) {
       launchDsv4HashTopk<IndType, HashIndType>(
           gating_output, topk_weights, topk_indices, num_tokens, num_experts,
-          routed_scaling_factor, input_ids, tid2eid, stream);
+          routed_scaling_factor, input_ids, tid2eid, stream, is_padding);
       return;
     }
   }
@@ -728,7 +760,8 @@ void dispatch_topk_softplus_sqrt_launch(
     int num_experts, int topk, bool renormalize, double routed_scaling_factor,
     const std::optional<torch::stable::Tensor>& correction_bias,
     const std::optional<torch::stable::Tensor>& input_ids,
-    const std::optional<torch::stable::Tensor>& tid2eid, cudaStream_t stream) {
+    const std::optional<torch::stable::Tensor>& tid2eid, cudaStream_t stream,
+    const std::optional<torch::stable::Tensor>& is_padding) {
   const float* bias_ptr = nullptr;
   if (correction_bias.has_value()) {
     bias_ptr = correction_bias.value().const_data_ptr<float>();
@@ -737,6 +770,22 @@ void dispatch_topk_softplus_sqrt_launch(
   auto launch = [&](auto* topk_indices_ptr) {
     using OutIndType =
         typename std::remove_pointer<decltype(topk_indices_ptr)>::type;
+
+    const bool* is_padding_ptr = nullptr;
+    if (is_padding.has_value()) {
+      const torch::stable::Tensor& is_padding_tensor = is_padding.value();
+      STD_TORCH_CHECK(is_padding_tensor.scalar_type() ==
+                          torch::headeronly::ScalarType::Bool,
+                      "is_padding tensor must be bool");
+      STD_TORCH_CHECK(is_padding_tensor.dim() == 1,
+                      "is_padding tensor must be 1D");
+      STD_TORCH_CHECK(is_padding_tensor.size(0) == num_tokens,
+                      "is_padding size mismatch, expected: ", num_tokens);
+      STD_TORCH_CHECK(is_padding_tensor.is_contiguous(),
+                      "is_padding tensor must be contiguous");
+      is_padding_ptr = is_padding_tensor.const_data_ptr<bool>();
+    }
+
     if (tid2eid.has_value()) {
       STD_TORCH_CHECK(input_ids.has_value(),
                       "input_ids is required for hash MoE");
@@ -751,7 +800,7 @@ void dispatch_topk_softplus_sqrt_launch(
             topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
             num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
             bias_ptr, true, input_ids.value().const_data_ptr<int64_t>(),
-            tid2eid.value().const_data_ptr<int64_t>(), stream);
+            tid2eid.value().const_data_ptr<int64_t>(), stream, is_padding_ptr);
       } else {
         STD_TORCH_CHECK(tid2eid.value().scalar_type() ==
                         torch::headeronly::ScalarType::Int);
@@ -761,7 +810,7 @@ void dispatch_topk_softplus_sqrt_launch(
             topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
             num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
             bias_ptr, true, input_ids.value().const_data_ptr<int>(),
-            tid2eid.value().const_data_ptr<int>(), stream);
+            tid2eid.value().const_data_ptr<int>(), stream, is_padding_ptr);
       }
     } else {
       vllm::moe::topkGatingSoftplusSqrtKernelLauncher<OutIndType, ComputeType>(
@@ -769,7 +818,7 @@ void dispatch_topk_softplus_sqrt_launch(
           topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
           num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
           bias_ptr, false, static_cast<const OutIndType*>(nullptr),
-          static_cast<const OutIndType*>(nullptr), stream);
+          static_cast<const OutIndType*>(nullptr), stream, is_padding_ptr);
     }
   };
 
@@ -793,7 +842,8 @@ void topk_softplus_sqrt(
     bool renormalize, double routed_scaling_factor,
     const std::optional<torch::stable::Tensor>& correction_bias,
     const std::optional<torch::stable::Tensor>& input_ids,
-    const std::optional<torch::stable::Tensor>& tid2eid) {
+    const std::optional<torch::stable::Tensor>& tid2eid,
+    const std::optional<torch::stable::Tensor>& is_padding) {
   const int num_experts = gating_output.size(-1);
   const auto num_tokens = gating_output.numel() / num_experts;
   const int topk = topk_weights.size(-1);
@@ -806,21 +856,22 @@ void topk_softplus_sqrt(
     dispatch_topk_softplus_sqrt_launch<float>(
         gating_output.const_data_ptr<float>(), topk_weights, topk_indices,
         token_expert_indices, num_tokens, num_experts, topk, renormalize,
-        routed_scaling_factor, correction_bias, input_ids, tid2eid, stream);
+        routed_scaling_factor, correction_bias, input_ids, tid2eid, stream,
+        is_padding);
   } else if (gating_output.scalar_type() ==
              torch::headeronly::ScalarType::Half) {
     dispatch_topk_softplus_sqrt_launch<__half>(
         reinterpret_cast<const __half*>(gating_output.const_data_ptr()),
         topk_weights, topk_indices, token_expert_indices, num_tokens,
         num_experts, topk, renormalize, routed_scaling_factor, correction_bias,
-        input_ids, tid2eid, stream);
+        input_ids, tid2eid, stream, is_padding);
   } else if (gating_output.scalar_type() ==
              torch::headeronly::ScalarType::BFloat16) {
     dispatch_topk_softplus_sqrt_launch<__nv_bfloat16>(
         reinterpret_cast<const __nv_bfloat16*>(gating_output.const_data_ptr()),
         topk_weights, topk_indices, token_expert_indices, num_tokens,
         num_experts, topk, renormalize, routed_scaling_factor, correction_bias,
-        input_ids, tid2eid, stream);
+        input_ids, tid2eid, stream, is_padding);
   } else {
     STD_TORCH_CHECK(false, "Unsupported gating_output data type: ",
                     gating_output.scalar_type());

csrc/libtorch_stable/moe/torch_bindings.cpp

@@ -8,19 +8,19 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_moe_C, m) {
   m.def(
       "topk_softmax(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
       "token_expert_indices, Tensor gating_output, bool renormalize, Tensor? "
-      "bias) -> ()");
+      "bias, Tensor? is_padding) -> ()");
 
   // Apply topk sigmoid to the gating outputs.
   m.def(
       "topk_sigmoid(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
       "token_expert_indices, Tensor gating_output, bool renormalize, "
-      "Tensor? bias, float routed_scaling_factor) -> ()");
+      "Tensor? bias, float routed_scaling_factor, Tensor? is_padding) -> ()");
 
   m.def(
       "topk_softplus_sqrt(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
       "token_expert_indices, Tensor gating_output, bool renormalize, float "
       "routed_scaling_factor, Tensor? "
-      "bias, Tensor? input_ids, Tensor? tid2eid) -> ()");
+      "bias, Tensor? input_ids, Tensor? tid2eid, Tensor? is_padding) -> ()");
 
   // Calculate the result of moe by summing up the partial results
   // from all selected experts. topk_ids/expert_map are optional and, when

docker/Dockerfile

@@ -793,7 +793,7 @@ RUN --mount=type=cache,target=/opt/uv/cache \
 # Install FlashInfer JIT cache (requires CUDA-version-specific index URL)
 # https://docs.flashinfer.ai/installation.html
 # From versions.json: .flashinfer.version
-ARG FLASHINFER_VERSION=0.6.14
+ARG FLASHINFER_VERSION=0.6.15.post1
 RUN --mount=type=cache,target=/opt/uv/cache \
     uv pip install --system flashinfer-jit-cache==${FLASHINFER_VERSION} \
         --index-url https://flashinfer.ai/whl/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')

docker/Dockerfile.rocm

@@ -339,18 +339,17 @@ COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust /rust
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust-toolchain.toml /rust-toolchain.toml
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1
 
-# RIXL/UCX build stages
-FROM base AS build_rixl
-ARG RIXL_BRANCH="39be1de8"
-ARG RIXL_REPO="https://github.com/ROCm/RIXL.git"
-ARG UCX_BRANCH="bfb51733"
+# NIXL/UCX build stages
+FROM base AS build_nixl
+ARG NIXL_BRANCH="231d56753047c989062a5cb2ac703a1ad761c7d2"
+ARG NIXL_REPO="https://github.com/ai-dynamo/nixl.git"
+ARG UCX_BRANCH="96e58a16039f6d7d213bc967b8069238742c5194"
 ARG UCX_REPO="https://github.com/openucx/ucx.git"
 ENV ROCM_PATH=/opt/rocm
 ENV UCX_HOME=/usr/local/ucx
-ENV RIXL_HOME=/usr/local/rixl
-ENV RIXL_BENCH_HOME=/usr/local/rixl_bench
+ENV NIXL_HOME=/usr/local/nixl
 
-# RIXL build system dependences and RDMA support
+# NIXL build system dependencies and RDMA support
 RUN apt-get -y update && apt-get -y install autoconf libtool pkg-config \
     libgrpc-dev \
     libgrpc++-dev \
@@ -368,7 +367,8 @@ RUN apt-get -y update && apt-get -y install autoconf libtool pkg-config \
     && rm -rf /var/lib/apt/lists/*
 
 RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install --system meson auditwheel patchelf tomlkit
+    uv pip install --system meson meson-python pybind11 pyyaml types-PyYAML \
+        auditwheel build patchelf pytest tomlkit "setuptools>=80.9.0"
 
 RUN --mount=type=cache,target=/root/.cache/ccache \
     cd /usr/local/src && \
@@ -396,30 +396,50 @@ ENV PATH=/usr/local/ucx/bin:$PATH
 ENV LD_LIBRARY_PATH=${UCX_HOME}/lib:${LD_LIBRARY_PATH}
 
 RUN --mount=type=cache,target=/root/.cache/ccache \
-    git clone ${RIXL_REPO} /opt/rixl && \
-    cd /opt/rixl && \
-    git checkout ${RIXL_BRANCH} && \
+    git clone ${NIXL_REPO} /opt/nixl && \
+    cd /opt/nixl && \
+    git checkout ${NIXL_BRANCH} && \
     CC="ccache gcc" CXX="ccache g++" \
-    meson setup build --prefix=${RIXL_HOME} \
+    meson setup build --prefix=${NIXL_HOME} \
                      -Ducx_path=${UCX_HOME} \
-                     -Drocm_path=${ROCM_PATH} && \
+                     -Dwheel_variant=rocm \
+                     -Dbuild_tests=false \
+                     -Dbuild_examples=false && \
     cd build && \
     ninja -j$(nproc) && \
-    ninja install
+    ninja install && \
+    echo "${NIXL_HOME}/lib/$(uname -m)-linux-gnu" \
+        > /etc/ld.so.conf.d/nixl.conf && \
+    echo "${NIXL_HOME}/lib/$(uname -m)-linux-gnu/plugins" \
+        >> /etc/ld.so.conf.d/nixl.conf && \
+    ldconfig
 
-# Generate RIXL wheel
+# Generate the ROCm NIXL wheel. Upstream's generic wheel helper detects CUDA,
+# so configure the ROCm wheel variant directly through Meson.
 # Exclude libcore and libpull from auditwheel: transitive dependencies
 # that are not shipped in the wheel and vary across base images.
-RUN cd /opt/rixl && \
-    sed -i "s/--exclude 'libamdhip64\*'/--exclude 'libamdhip64*' --exclude 'libcore*' --exclude 'libpull*'/" \
-        contrib/build-wheel.sh && \
-    mkdir -p /app/install && \
-    _ucx_install_dir=${UCX_HOME} \
-    ./contrib/build-wheel.sh \
-        --output-dir /app/install \
-        --rocm-dir ${ROCM_PATH} \
+RUN cd /opt/nixl && \
+    ./contrib/tomlutil.py --wheel-name nixl-rocm pyproject.toml && \
+    CC="ccache gcc" CXX="ccache g++" \
+    uv build --wheel --no-build-isolation --out-dir /tmp/nixl_wheels \
+        --python ${PYTHON_VERSION} \
+        -Csetup-args=-Ducx_path=${UCX_HOME} \
+        -Csetup-args=-Dwheel_variant=rocm \
+        -Csetup-args=-Dbuild_tests=false \
+        -Csetup-args=-Dbuild_examples=false && \
+    mkdir -p /tmp/nixl_wheels/repaired /app/install && \
+    auditwheel repair \
+        --exclude 'libamdhip64*' \
+        --exclude 'libcore*' \
+        --exclude 'libpull*' \
+        /tmp/nixl_wheels/nixl_rocm*.whl \
+        --plat manylinux_2_34_$(uname -m) \
+        --wheel-dir /tmp/nixl_wheels/repaired && \
+    ./contrib/wheel_add_ucx_plugins.py \
         --ucx-plugins-dir ${UCX_HOME}/lib/ucx \
-        --nixl-plugins-dir ${RIXL_HOME}/lib/x86_64-linux-gnu/plugins
+        --nixl-plugins-dir ${NIXL_HOME}/lib/$(uname -m)-linux-gnu/plugins \
+        /tmp/nixl_wheels/repaired/*.whl && \
+    cp /tmp/nixl_wheels/repaired/*.whl /app/install
 
 # ROCShmem build stage - split from DeepEP so changing DEEPEP_BRANCH does not
 # invalidate the slow ROCShmem build.
@@ -660,10 +680,10 @@ RUN if [ "${DEEPEP_NIC}" = "cx7" ] || [ "${DEEPEP_NIC}" = "io" ]; then \
     ninja && ninja install && ldconfig && rm -rf /tmp/rdma-core; \
 fi
 
-# Install RIXL + DeepEP wheels.
-RUN --mount=type=bind,from=build_rixl,src=/app/install,target=/rixl_install \
+# Install NIXL + DeepEP wheels.
+RUN --mount=type=bind,from=build_nixl,src=/app/install,target=/nixl_install \
     --mount=type=bind,from=build_deepep,src=/app/deep_install,target=/deep_install \
-    uv pip install --system /rixl_install/*.whl /deep_install/*.whl
+    uv pip install --system /nixl_install/*.whl /deep_install/*.whl
 
 # Copy ROCShmem runtime libraries.
 COPY --from=build_rocshmem /opt/rocshmem /opt/rocshmem
@@ -724,6 +744,7 @@ ENV MIOPEN_DEBUG_CONV_GEMM=0
 # Use legacy IPC mode for HSA to avoid GPU memory pinning issues with UCX rocm_ipc.
 # See: https://github.com/ROCm/rocm-libraries/issues/6266
 ENV HSA_ENABLE_IPC_MODE_LEGACY=1
+ENV UCX_RMA_PPLN_ENABLE=y
 
 # ROCm profiler limits workaround.
 RUN echo "ROCTRACER_MAX_EVENTS=10000000" > ${COMMON_WORKDIR}/libkineto.conf
@@ -796,9 +817,9 @@ RUN --mount=type=bind,from=export_vllm,src=/,target=/install \
     && pip uninstall -y vllm \
     && uv pip install --system *.whl
 
-# Install RIXL wheel
-RUN --mount=type=bind,from=build_rixl,src=/app/install,target=/rixl_install \
-    uv pip install --system /rixl_install/*.whl
+# Install NIXL ROCm wheel
+RUN --mount=type=bind,from=build_nixl,src=/app/install,target=/nixl_install \
+    uv pip install --system /nixl_install/*.whl
 
 ARG COMMON_WORKDIR
 ARG BASE_IMAGE
@@ -813,6 +834,7 @@ COPY --from=export_vllm /docker ${COMMON_WORKDIR}/vllm/docker
 # Use legacy IPC mode for HSA to avoid GPU memory pinning issues with UCX rocm_ipc
 # See: https://github.com/ROCm/rocm-libraries/issues/6266
 ENV HSA_ENABLE_IPC_MODE_LEGACY=1
+ENV UCX_RMA_PPLN_ENABLE=y
 
 ENV TOKENIZERS_PARALLELISM=false
 

docker/Dockerfile.rocm_base

@@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.9.0"
 ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git"
 ARG FA_BRANCH="0e60e394"
 ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git"
-ARG AITER_BRANCH="v0.1.16.post3"
+ARG AITER_BRANCH="v0.1.16.post5"
 ARG AITER_REPO="https://github.com/ROCm/aiter.git"
 ARG MORI_BRANCH="v1.1.0"
 ARG MORI_REPO="https://github.com/ROCm/mori.git"

docker/Dockerfile.s390x

@@ -86,6 +86,29 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     mkdir -p /tmp/hf-xet/dist && \
     cp dist/*.whl /tmp/hf-xet/dist/
 
+# Build LLVM 20 from source for llvmlite (system repos ship LLVM 21 which
+# llvmlite v0.47 does not support; only SystemZ target is needed).
+FROM base AS llvm20-build
+ARG LLVM_VERSION=20.1.8
+WORKDIR /tmp
+RUN microdnf install -y ninja-build gcc gcc-c++ python3 xz && \
+    curl -LO https://github.com/llvm/llvm-project/releases/download/llvmorg-${LLVM_VERSION}/llvm-project-${LLVM_VERSION}.src.tar.xz && \
+    tar -xf llvm-project-${LLVM_VERSION}.src.tar.xz && \
+    cmake -G Ninja -S llvm-project-${LLVM_VERSION}.src/llvm -B build \
+        -DCMAKE_BUILD_TYPE=Release \
+        -DCMAKE_INSTALL_PREFIX=/opt/llvm20 \
+        -DLLVM_TARGETS_TO_BUILD="SystemZ" \
+        -DLLVM_ENABLE_RTTI=ON \
+        -DLLVM_BUILD_TOOLS=OFF \
+        -DLLVM_BUILD_UTILS=ON \
+        -DLLVM_BUILD_EXAMPLES=OFF \
+        -DLLVM_BUILD_TESTS=OFF \
+        -DLLVM_INCLUDE_TESTS=OFF \
+        -DLLVM_INCLUDE_EXAMPLES=OFF \
+        -DLLVM_INCLUDE_BENCHMARKS=OFF && \
+    ninja -C build install && \
+    rm -rf build llvm-project-${LLVM_VERSION}.src*
+
 # Build numba
 FROM python-install AS numba-builder
 
@@ -96,11 +119,13 @@ WORKDIR /tmp
 
 # Clone all required dependencies
 RUN --mount=type=cache,target=/root/.cache/uv \
+    --mount=type=bind,from=llvm20-build,source=/opt/llvm20,target=/opt/llvm20 \
     microdnf install ninja-build gcc gcc-c++ -y && \
     git clone --recursive https://github.com/numba/llvmlite.git -b v0.47.0 && \
     git clone --recursive https://github.com/numba/numba.git -b ${NUMBA_VERSION} && \
     cd llvmlite && \
     uv pip install 'cmake<4' 'setuptools<70' numpy && \
+    CMAKE_PREFIX_PATH=/opt/llvm20 LLVM_CONFIG=/opt/llvm20/bin/llvm-config \
     python setup.py bdist_wheel && \
     cd ../numba && \
     if ! grep '#include "dynamic_annotations.h"' numba/_dispatcher.cpp; then \
@@ -158,7 +183,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
      NUMBA_WHL_FILE=$(ls /tmp/numba-wheels/*.whl) && \
      OPENCV_WHL_FILE=$(ls /tmp/opencv-wheels/*.whl) && \
      uv pip install -v \
-        $ARROW_WHL_FILE \
         $VISION_WHL_FILE \
         $HF_XET_WHL_FILE \
         $LLVM_WHL_FILE \

docker/ci-rocm.hcl

@@ -59,7 +59,7 @@ variable "PYTORCH_ROCM_ARCH" {
 }
 
 # Pre-built CI base image (Tier 1). Per-PR builds pull this instead of
-# rebuilding RIXL/DeepEP/torchcodec from scratch. The ci_base stage in
+# rebuilding NIXL/DeepEP/torchcodec from scratch. The ci_base stage in
 # Dockerfile.rocm inherits from base, so CI_BASE_IMAGE only affects the test
 # stage and is irrelevant when building --target ci_base itself.
 variable "CI_BASE_IMAGE" {
@@ -75,7 +75,7 @@ variable "CI_MAX_JOBS" {
 # Upstream dependency commit pins -- extracted from Dockerfile.rocm by
 # ci-bake-rocm.sh at build time. Empty defaults are safe: the cache
 # functions produce no entries when the variable is empty.
-variable "RIXL_BRANCH" {
+variable "NIXL_BRANCH" {
   default = ""
 }
 
@@ -91,7 +91,7 @@ variable "DEEPEP_BRANCH" {
   default = ""
 }
 
-variable "RIXL_CACHE_KEY" {
+variable "NIXL_CACHE_KEY" {
   default = ""
 }
 
@@ -236,7 +236,7 @@ function "get_cache_to_rocm_rust" {
   ])
 }
 
-# Cache functions for upstream dependency stages (RIXL/UCX, ROCShmem, DeepEP).
+# Cache functions for upstream dependency stages (NIXL/UCX, ROCShmem, DeepEP).
 # These stages are pinned to specific upstream commit hashes, so cache keys use
 # those hashes rather than the Buildkite commit. This means the cache persists
 # across all vLLM commits as long as the upstream dependency pins don't change.
@@ -244,16 +244,16 @@ function "get_cache_to_rocm_rust" {
 function "get_cache_from_rocm_deps" {
   params = []
   result = compact([
-    RIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_CACHE_KEY}" : (RIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_BRANCH}-ucx-${UCX_BRANCH}" : ""),
+    NIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_CACHE_KEY}" : (NIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH}" : ""),
     ROCSHMEM_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocshmem-rocm-${ROCSHMEM_CACHE_KEY}" : (ROCSHMEM_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocshmem-rocm-${ROCSHMEM_BRANCH}" : ""),
     DEEPEP_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:deepep-rocm-${DEEPEP_CACHE_KEY}" : (DEEPEP_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:deepep-rocm-${DEEPEP_BRANCH}-rocshmem-${ROCSHMEM_BRANCH}" : ""),
   ])
 }
 
-function "get_cache_to_rocm_rixl" {
+function "get_cache_to_rocm_nixl" {
   params = []
   result = compact([
-    RIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_CACHE_KEY},mode=min" : (RIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_BRANCH}-ucx-${UCX_BRANCH},mode=min" : ""),
+    NIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_CACHE_KEY},mode=min" : (NIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH},mode=min" : ""),
   ])
 }
 
@@ -372,11 +372,11 @@ variable "CI_BASE_IMAGE_TAG_STABLE" {
 # in the registry cache keyed by its upstream commit hash. When ci_base rebuilds
 # (e.g., requirements change), these stages are cache hits if their upstream
 # pins haven't changed -- saving ~35min of compilation.
-target "rixl-rocm-ci" {
+target "nixl-rocm-ci" {
   inherits   = ["_common-rocm", "_ci-rocm"]
-  target     = "build_rixl"
+  target     = "build_nixl"
   cache-from = get_cache_from_rocm_deps()
-  cache-to   = get_cache_to_rocm_rixl()
+  cache-to   = get_cache_to_rocm_nixl()
   output     = ["type=cacheonly"]
 }
 
@@ -396,7 +396,7 @@ target "deepep-rocm-ci" {
   output     = ["type=cacheonly"]
 }
 
-# Builds only the ci_base stage (RIXL, DeepEP, torchcodec, etc.)
+# Builds only the ci_base stage (NIXL, DeepEP, torchcodec, etc.)
 # Invoked by the ensure-ci-base step when the content hash of ci_base-affecting
 # files drifts from the remote image label. Per-PR builds then pull the result
 # as CI_BASE_IMAGE instead of rebuilding those slow layers on every commit.
@@ -412,7 +412,7 @@ target "ci-base-rocm-ci" {
       CI_BASE_IMAGE_TAG_CONTENT_EXTRA != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_CONTENT_EXTRA}" : "",
       CI_BASE_IMAGE_TAG_STABLE != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_STABLE}" : "",
     ]),
-    # Import upstream dependency caches so RIXL/ROCShmem/DeepEP stages
+    # Import upstream dependency caches so NIXL/ROCShmem/DeepEP stages
     # are cache hits even when ci_base itself needs rebuilding.
     get_cache_from_rocm_deps(),
   )
@@ -424,5 +424,5 @@ target "ci-base-rocm-ci" {
 # Group for ci_base builds -- exports dependency stage caches alongside the
 # ci_base image so future rebuilds can reuse them independently.
 group "ci-base-rocm-ci-with-deps" {
-  targets = ["rixl-rocm-ci", "rocshmem-rocm-ci", "deepep-rocm-ci", "ci-base-rocm-ci"]
+  targets = ["nixl-rocm-ci", "rocshmem-rocm-ci", "deepep-rocm-ci", "ci-base-rocm-ci"]
 }

docker/docker-bake-rocm.hcl

@@ -53,7 +53,7 @@ variable "CI_BASE_IMAGE" {
 # Upstream dependency commit pins. Plain local bake builds use the Dockerfile
 # ARG defaults. ci-bake-rocm.sh resolves those defaults (plus any env
 # overrides) and writes a small HCL override before invoking CI targets.
-variable "RIXL_BRANCH" {
+variable "NIXL_BRANCH" {
   default = ""
 }
 
@@ -106,7 +106,7 @@ target "test-rocm" {
   output   = ["type=docker"]
 }
 
-# CI base image target - builds only the ci_base stage (RIXL, DeepEP,
+# CI base image target - builds only the ci_base stage (NIXL, DeepEP,
 # torchcodec, requirements, etc.). Used by the weekly scheduled build and
 # the auto-rebuild trigger when requirements change in a PR.
 target "ci-base-rocm" {

docker/versions.json

@@ -68,7 +68,7 @@
       "default": "true"
     },
     "FLASHINFER_VERSION": {
-      "default": "0.6.14"
+      "default": "0.6.15.post1"
     },
     "GDRCOPY_CUDA_VERSION": {
       "default": "12.8"

docs/features/nixl_connector_usage.md

@@ -13,11 +13,7 @@ Install the NIXL library: `uv pip install nixl`, as a quick start on Nvidia plat
 - Refer to [NIXL official repository](https://github.com/ai-dynamo/nixl) for more installation instructions
 - The specified required NIXL version can be found in [requirements/kv_connectors.txt](../../requirements/kv_connectors.txt) and other relevant config files
 
-For ROCm platform, the [ROCm docker file](../../docker/Dockerfile.rocm) includes RIXL and ucx already.
-
-- Refer to [RIXL official repository](https://github.com/rocm/rixl) for more information
-- The supportive libraries for RIXL can be found in [requirements/kv_connectors_rocm.txt](../../requirements/kv_connectors_rocm.txt)
-- In the future we may remove RIXL from docker image file and users will be able to install from pre-compiled binary packages
+For ROCm, the [ROCm Dockerfile](../../docker/Dockerfile.rocm) builds NIXL and UCX with ROCm support from source.
 
 For non-cuda platform, please install nixl with ucx build from source, instructed as below.
 

docs/getting_started/installation/cpu.md

@@ -315,7 +315,7 @@ vLLM CPU supports data parallel (DP), tensor parallel (TP) and pipeline parallel
 - vLLM CPU supports quantizations:
     - AWQ (x86 only)
     - GPTQ (x86 only)
-    - compressed-tensor INT8 W8A8 (x86, s390x)
+    - compressed-tensor INT8 W8A8 (x86 only)
 
 ### Why do I see `get_mempolicy: Operation not permitted` when running in Docker?
 

docs/getting_started/installation/cpu.s390x.inc.md

@@ -11,7 +11,7 @@ Currently, the CPU implementation for s390x architecture supports FP32, BF16 and
 - OS: `Linux`
 - SDK: `gcc/g++ >= 14.0.0` or later with Command Line Tools
 - Instruction Set Architecture (ISA): VXE support is required. Works with Z14 and above.
-- Build install python packages: `torchvision`, `llvmlite`, `numba`, `pyarrow (for testing)`, `opencv-headless`
+- Build from source python packages (no pre-built s390x wheels): `torchvision`, `llvmlite`, `numba`, `opencv-python-headless`, `hf-xet`
 
 --8<-- [end:requirements]
 --8<-- [start:set-up-using-python]
@@ -28,13 +28,24 @@ Install the following packages from the package manager before building the vLLM
 
 ```bash
 dnf install -y \
-    which procps findutils tar vim git gcc-toolset-14 gcc-toolset-14-binutils gcc-toolset-14-libatomic-devel zlib-devel \
+    which procps findutils tar vim git patch xz ninja-build \
+    gcc-toolset-14 gcc-toolset-14-binutils gcc-toolset-14-libatomic-devel zlib-devel \
     libjpeg-turbo-devel libtiff-devel libpng-devel libwebp-devel freetype-devel harfbuzz-devel \
     openssl-devel openblas openblas-devel autoconf automake libtool cmake numpy libsndfile \
     clang llvm-devel llvm-static clang-devel
 ```
 
-Install rust>=1.80 which is needed for `outlines-core` and `uvloop` python packages installation.
+Build and install `numactl` from source:
+
+```bash
+curl -LO https://github.com/numactl/numactl/archive/refs/tags/v2.0.19.tar.gz
+tar -xvzf v2.0.19.tar.gz
+cd numactl-2.0.19
+./autogen.sh && ./configure && make && make install
+cd ..
+```
+
+Install rust>=1.80 which is needed for `outlines-core`, `uvloop`, and `hf-xet` python packages installation.
 
 ```bash
 curl https://sh.rustup.rs -sSf | sh -s -- -y && \
@@ -44,26 +55,79 @@ curl https://sh.rustup.rs -sSf | sh -s -- -y && \
 Execute the following commands to build and install vLLM from source.
 
 !!! tip
-    Please build the following dependencies, `torchvision`, `llvmlite`, `numba`, `llguidance`, `pyarrow`, `opencv-headless` from source before building vLLM.
+    Pre-built wheels are not available for s390x for the following packages. Build them from source before building vLLM: `torchvision`, `llvmlite`, `numba`, `opencv-python-headless`, `hf-xet`.
+    See `docker/Dockerfile.s390x` for exact versions and build commands used in each multi-stage build.
+
+!!! note "LLVM 20 required for llvmlite"
+    `llvmlite v0.47` requires LLVM 20, but UBI 9.6 repos ship LLVM 21 which is
+    not compatible. You must build LLVM 20 from source before building `llvmlite`:
+
+    ```bash
+    curl -LO https://github.com/llvm/llvm-project/releases/download/llvmorg-20.1.8/llvm-project-20.1.8.src.tar.xz
+    tar -xf llvm-project-20.1.8.src.tar.xz
+    cmake -G Ninja -S llvm-project-20.1.8.src/llvm -B llvm-build \
+        -DCMAKE_BUILD_TYPE=Release \
+        -DCMAKE_INSTALL_PREFIX=/opt/llvm20 \
+        -DLLVM_TARGETS_TO_BUILD="SystemZ" \
+        -DLLVM_ENABLE_RTTI=ON \
+        -DLLVM_BUILD_TOOLS=OFF \
+        -DLLVM_BUILD_UTILS=ON \
+        -DLLVM_BUILD_EXAMPLES=OFF \
+        -DLLVM_BUILD_TESTS=OFF \
+        -DLLVM_INCLUDE_TESTS=OFF \
+        -DLLVM_INCLUDE_EXAMPLES=OFF \
+        -DLLVM_INCLUDE_BENCHMARKS=OFF
+    ninja -C llvm-build install
+    ```
+
+    Then build `llvmlite` pointing to LLVM 20:
+
+    ```bash
+    CMAKE_PREFIX_PATH=/opt/llvm20 LLVM_CONFIG=/opt/llvm20/bin/llvm-config \
+        python setup.py bdist_wheel
+    ```
 
 ```bash
-    uv pip install -v \
-        -r requirements/build/cpu.txt \
-        -r requirements/cpu.txt \
-        --torch-backend cpu \
-        --index-strategy unsafe-best-match && \
-    VLLM_TARGET_DEVICE=cpu python setup.py bdist_wheel && \
-        uv pip install dist/*.whl
+uv pip install -v \
+    /path/to/torchvision.whl \
+    /path/to/llvmlite.whl \
+    /path/to/numba.whl \
+    /path/to/opencv_python_headless.whl \
+    /path/to/hf_xet.whl \
+    -r requirements/build/cpu.txt \
+    -r requirements/cpu.txt \
+    --torch-backend cpu \
+    --index-strategy unsafe-best-match && \
+VLLM_TARGET_DEVICE=cpu VLLM_CPU_MOE_PREPACK=0 python setup.py bdist_wheel && \
+    uv pip install dist/*.whl
 ```
 
 ??? console "pip"
     ```bash
-        pip install -v \
-            --extra-index-url https://download.pytorch.org/whl/cpu \
-            -r requirements/build/cpu.txt \
-            -r requirements/cpu.txt \
-        VLLM_TARGET_DEVICE=cpu python setup.py bdist_wheel && \
-            pip install dist/*.whl
+    pip install -v \
+        --extra-index-url https://download.pytorch.org/whl/cpu \
+        /path/to/torchvision.whl \
+        /path/to/llvmlite.whl \
+        /path/to/numba.whl \
+        /path/to/opencv_python_headless.whl \
+        /path/to/hf_xet.whl \
+        -r requirements/build/cpu.txt \
+        -r requirements/cpu.txt && \
+    VLLM_TARGET_DEVICE=cpu VLLM_CPU_MOE_PREPACK=0 python setup.py bdist_wheel && \
+        pip install dist/*.whl
+    ```
+
+!!! warning "Protobuf workaround for s390x"
+    The C++ protobuf extension crashes on s390x. After installation, set the
+    following environment variable and remove the C++ extensions:
+
+    ```bash
+    export PROTOCOL_BUFFERS_PYTHON_IMPLEMENTATION=python
+
+    # Remove C++ protobuf extensions that crash on s390x
+    SITE_PKGS=$(python -c "import site; print(site.getsitepackages()[0])")
+    rm -rf "$SITE_PKGS/google/_upb/"*.so \
+           "$SITE_PKGS/google/protobuf/pyext/"*.so 2>/dev/null || true
     ```
 
 --8<-- [end:build-wheel-from-source]
@@ -80,19 +144,20 @@ docker build -f docker/Dockerfile.s390x \
 
 # Launch OpenAI server
 docker run --rm \
-    --privileged true \
+    --security-opt seccomp=unconfined \
+    --cap-add SYS_NICE \
     --shm-size 4g \
     -p 8000:8000 \
     -e VLLM_CPU_KVCACHE_SPACE=<KV cache space> \
     -e VLLM_CPU_OMP_THREADS_BIND=<CPU cores for inference> \
     vllm-cpu-env \
     --model meta-llama/Llama-3.2-1B-Instruct \
-    --dtype float \
+    --dtype bfloat16 \
     other vLLM OpenAI server arguments
 ```
 
 !!! tip
-    An alternative of `--privileged true` is `--cap-add SYS_NICE --security-opt seccomp=unconfined`.
+    Alternatively, `--privileged=true` also works but is broader and not generally recommended.
 
 --8<-- [end:build-image-from-source]
 --8<-- [start:extra-information]

docs/mkdocs/javascript/reo.js

@@ -0,0 +1,3 @@
+// Reo.Dev documentation tracking
+// https://docs.reo.dev/integrations/input-sources/developer-insights/documentation
+!function(){var e,t,n;e="d5c4337961ef0ac",t=function(){Reo.init({clientID:"d5c4337961ef0ac", enableThirdPartyTracking: true})},(n=document.createElement("script")).src="https://static.reo.dev/"+e+"/reo.js",n.defer=!0,n.onload=t,document.head.appendChild(n)}();

mkdocs.yaml

@@ -160,3 +160,4 @@ extra_javascript:
   - https://unpkg.com/mathjax@3.2.2/es5/tex-mml-chtml.js
   - mkdocs/javascript/edit_and_feedback.js
   - mkdocs/javascript/slack_and_forum.js
+  - mkdocs/javascript/reo.js

requirements/cuda.txt

@@ -14,8 +14,8 @@ PyNvVideoCodec==2.0.4
 # flashinfer-cubin is not on PyPI since 0.6.14; setup.py excludes it from
 # install_requires so the published wheel does not carry an unresolvable pin
 --extra-index-url https://flashinfer.ai/whl/
-flashinfer-python==0.6.14
-flashinfer-cubin==0.6.14
+flashinfer-python==0.6.15.post1
+flashinfer-cubin==0.6.15.post1
 apache-tvm-ffi==0.1.10
 tilelang==0.1.9
 nvidia-cudnn-frontend>=1.19.1
@@ -26,7 +26,7 @@ fastsafetensors >= 0.3.2
 
 # QuACK and Cutlass DSL for FA4 (cute-DSL implementation)
 nvidia-cutlass-dsl[cu13]==4.6.0
-quack-kernels>=0.4.0  # Required for tml-fa4
+quack-kernels>=0.6.1  # Required for CUTLASS DSL 4.6 by MSA
 
 # Tokenspeed_MLA for faster mla with spec decode
 tokenspeed-mla==0.1.8; platform_system == "Linux"

requirements/tpu.txt

@@ -12,4 +12,4 @@ ray[data]
 setuptools==78.1.0
 setuptools-rust>=1.9.0
 nixl==0.3.0
-tpu-inference==0.24.0
+tpu-inference==0.25.0

rust/Cargo.lock

@@ -4252,18 +4252,6 @@ dependencies = [
  "syn 2.0.117",
 ]
 
-[[package]]
-name = "subenum"
-version = "1.1.3"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "ec3d08fe7078c57309d5c3d938e50eba95ba1d33b9c3a101a8465fc6861a5416"
-dependencies = [
- "heck",
- "proc-macro2",
- "quote",
- "syn 2.0.117",
-]
-
 [[package]]
 name = "subtle"
 version = "2.6.1"
@@ -5516,15 +5504,11 @@ dependencies = [
  "asynk-strim-attr",
  "bytes",
  "clap",
- "easy-ext",
  "expect-test",
  "futures",
  "half",
- "indexmap 2.13.0",
  "itertools 0.14.0",
  "llm-multimodal",
- "minijinja",
- "minijinja-contrib",
  "ndarray 0.17.2",
  "oss-harmony",
  "paste",
@@ -5535,17 +5519,16 @@ dependencies = [
  "serde_json",
  "serde_with",
  "serial_test",
- "strum",
- "subenum",
  "tempfile",
  "thiserror 2.0.18",
  "thiserror-ext",
- "time",
  "tokio",
  "tracing",
  "tracing-subscriber",
  "trait-set",
  "uuid",
+ "vllm-chat-renderer",
+ "vllm-chat-types",
  "vllm-engine-core-client",
  "vllm-llm",
  "vllm-parser",
@@ -5555,6 +5538,42 @@ dependencies = [
  "zeromq",
 ]
 
+[[package]]
+name = "vllm-chat-renderer"
+version = "0.1.0"
+dependencies = [
+ "anyhow",
+ "enum-as-inner",
+ "expect-test",
+ "indexmap 2.13.0",
+ "itertools 0.14.0",
+ "minijinja",
+ "minijinja-contrib",
+ "oss-harmony",
+ "serde",
+ "serde-json-fmt",
+ "serde_json",
+ "serde_with",
+ "strum",
+ "tempfile",
+ "thiserror 2.0.18",
+ "thiserror-ext",
+ "time",
+ "tracing",
+ "vllm-chat-types",
+ "vllm-tokenizer",
+]
+
+[[package]]
+name = "vllm-chat-types"
+version = "0.1.0"
+dependencies = [
+ "easy-ext",
+ "serde",
+ "serde_json",
+ "serde_with",
+]
+
 [[package]]
 name = "vllm-cmd"
 version = "0.1.0"
@@ -5695,11 +5714,11 @@ dependencies = [
  "expect-test",
  "futures",
  "openai-protocol",
- "serde",
  "serde_json",
  "thiserror 2.0.18",
  "thiserror-ext",
  "tool-parser",
+ "vllm-chat-types",
  "vllm-tokenizer",
  "winnow",
  "xgrammar-structural-tag",
@@ -5806,6 +5825,7 @@ dependencies = [
  "rustc-hash 1.1.0",
  "serde",
  "serde_json",
+ "serde_with",
  "tekken-rs",
  "tempfile",
  "thiserror 2.0.18",

rust/Cargo.toml

@@ -2,6 +2,8 @@
 members = [
     "src/bench",
     "src/chat",
+    "src/chat-renderer",
+    "src/chat-types",
     "src/cmd",
     "src/engine-core-client",
     "src/llm",
@@ -135,6 +137,8 @@ uuid = { version = "1.22.0", features = ["v4"] }
 validator = { version = "0.20.0", features = ["derive"] }
 vllm-bench = { path = "src/bench" }
 vllm-chat = { path = "src/chat" }
+vllm-chat-renderer = { path = "src/chat-renderer" }
+vllm-chat-types = { path = "src/chat-types" }
 vllm-engine-core-client = { path = "src/engine-core-client" }
 vllm-llm = { path = "src/llm" }
 vllm-managed-engine = { path = "src/managed-engine" }

rust/proto/vllm_grpc.proto

@@ -14,6 +14,10 @@ service Generate {
   rpc GenerateStream (GenerateRequest) returns (stream GenerateResponse) {}
 }
 
+service Control {
+  rpc Abort (AbortRequest) returns (AbortResponse) {}
+}
+
 // ======================================================================================
 // Generate Request
 // ======================================================================================
@@ -201,3 +205,12 @@ message TokenIds {
   repeated uint32 ids = 1;
 }
 
+// ======================================================================================
+// Control
+// ======================================================================================
+
+message AbortRequest {
+  repeated string request_ids = 1;
+}
+
+message AbortResponse {}

rust/src/chat-renderer/Cargo.toml

@@ -0,0 +1,32 @@
+[package]
+name = "vllm-chat-renderer"
+version.workspace = true
+edition.workspace = true
+license.workspace = true
+
+[dependencies]
+anyhow.workspace = true
+enum-as-inner.workspace = true
+indexmap.workspace = true
+itertools.workspace = true
+minijinja.workspace = true
+minijinja-contrib.workspace = true
+openai-harmony.workspace = true
+serde.workspace = true
+serde-json-fmt.workspace = true
+serde_json.workspace = true
+serde_with.workspace = true
+strum.workspace = true
+thiserror.workspace = true
+thiserror-ext.workspace = true
+time.workspace = true
+tracing.workspace = true
+vllm-chat-types.workspace = true
+vllm-tokenizer.workspace = true
+
+[dev-dependencies]
+expect-test.workspace = true
+tempfile.workspace = true
+
+[lints]
+workspace = true

rust/src/chat-renderer/src/deepseek_v32/encoding.rs

@@ -11,8 +11,10 @@ use serde_json::Value;
 use serde_json_fmt::JsonFormat;
 
 use crate::error::{Error, Result};
-use crate::request::{ChatContent, ChatMessage, ChatRequest, ChatRole, ChatTool};
-use crate::{AssistantContentBlock, AssistantMessageExt, AssistantToolCall};
+use crate::{
+    AssistantContentBlock, AssistantMessageExt, AssistantToolCall, ChatContent, ChatMessage,
+    ChatRole, RenderRequest, Tool,
+};
 
 const BOS_TOKEN: &str = "<｜begin▁of▁sentence｜>";
 const EOS_TOKEN: &str = "<｜end▁of▁sentence｜>";
@@ -39,7 +41,7 @@ struct RenderedToolSchema<'a> {
 }
 
 /// Render one chat request into the final prompt string.
-pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
+pub(super) fn render_request(request: &RenderRequest<'_>) -> Result<String> {
     let thinking_mode = match request.enable_thinking()?.unwrap_or(false) {
         true => ThinkingMode::Thinking,
         false => ThinkingMode::Chat,
@@ -49,20 +51,19 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
         Some(ChatRole::User | ChatRole::Developer)
     );
     let render_offset = isize::from(request.tool_parsing_enabled());
-    let last_user_render_index =
-        find_last_user_render_index(request.messages.as_slice(), render_offset);
-    let last_user_actual_index = find_last_user_actual_index(request.messages.as_slice());
+    let last_user_render_index = find_last_user_render_index(request.messages, render_offset);
+    let last_user_actual_index = find_last_user_actual_index(request.messages);
     let continue_final_message = request.chat_options.continue_final_message();
     let mut prompt = String::from(BOS_TOKEN);
 
     if request.tool_parsing_enabled() {
-        render_system_message(&mut prompt, None, &request.tools)?;
+        render_system_message(&mut prompt, None, request.tools)?;
     }
 
     for (message_index, message) in request.messages.iter().enumerate() {
         render_message(
             &mut prompt,
-            request.messages.as_slice(),
+            request.messages,
             message_index,
             message,
             render_offset,
@@ -168,7 +169,7 @@ fn find_last_user_actual_index(messages: &[ChatMessage]) -> usize {
 fn render_system_message(
     out: &mut String,
     content: Option<&ChatContent>,
-    tools: &[ChatTool],
+    tools: &[Tool],
 ) -> Result<()> {
     if let Some(content) = content {
         write_chat_content(out, content)?;
@@ -185,7 +186,7 @@ fn render_system_message(
 fn render_developer_message(
     out: &mut String,
     content: &ChatContent,
-    tools: &[ChatTool],
+    tools: &[Tool],
     opens_thinking: bool,
 ) -> Result<()> {
     if content.is_empty() {
@@ -452,7 +453,7 @@ fn encode_arguments_to_dsml(out: &mut String, tool_call: &AssistantToolCall) ->
 }
 
 /// Render the full tool preamble shown to the model.
-fn render_tools(out: &mut String, tools: &[ChatTool]) -> Result<()> {
+fn render_tools(out: &mut String, tools: &[Tool]) -> Result<()> {
     out.push_str(
         r#"## Tools
 
@@ -500,7 +501,7 @@ Here are the functions available in JSONSchema format:
 
 /// Serialize one typed tool schema into the JSON shape embedded inside
 /// `<functions>`.
-fn render_tool_schema(out: &mut String, tool: &ChatTool) -> Result<()> {
+fn render_tool_schema(out: &mut String, tool: &Tool) -> Result<()> {
     out.push_str(&json_dumps(&RenderedToolSchema {
         name: &tool.name,
         description: tool.description.as_deref(),
@@ -517,7 +518,7 @@ fn write_chat_content(out: &mut String, content: &ChatContent) -> Result<()> {
         ChatContent::Text(text) => out.push_str(text),
         ChatContent::Parts(parts) => {
             for part in parts {
-                out.push_str(part.as_text()?);
+                out.push_str(part.as_text().map_err(Error::UnsupportedMultimodalContent)?);
             }
         }
     }

rust/src/chat-renderer/src/deepseek_v32/mod.rs

@@ -3,11 +3,10 @@
 
 mod encoding;
 
-use vllm_text::Prompt;
-
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{
+    ChatRenderer, RenderRequest, RenderedPrompt, RenderedPromptContent, request_template_kwargs,
+};
 use crate::Result;
-use crate::request::ChatRequest;
 
 /// Dedicated DeepSeek V3.2 renderer.
 #[derive(Debug, Clone, Copy, Default)]
@@ -21,12 +20,12 @@ impl DeepSeekV32ChatRenderer {
 }
 
 impl ChatRenderer for DeepSeekV32ChatRenderer {
-    fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
+    fn render(&self, request: RenderRequest<'_>) -> Result<RenderedPrompt> {
         request.validate()?;
 
         Ok(RenderedPrompt {
-            prompt: Prompt::Text(encoding::render_request(request)?),
-            effective_template_kwargs: request_template_kwargs(request),
+            content: RenderedPromptContent::Text(encoding::render_request(&request)?),
+            effective_template_kwargs: request_template_kwargs(&request),
         })
     }
 }

rust/src/chat-renderer/src/deepseek_v32/tests.rs

@@ -9,36 +9,35 @@ use thiserror_ext::AsReport;
 
 use super::DeepSeekV32ChatRenderer;
 use crate::error::Error;
-use crate::event::{AssistantContentBlock, AssistantToolCall};
-use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
-use crate::request::{
-    ChatContentPart, ChatMessage, ChatRequest, ChatTool, ChatToolChoice, GenerationPromptMode,
+use crate::test_utils::{FixtureRequestOptions, fixture_chat_request};
+use crate::{AssistantContentBlock, AssistantToolCall};
+use crate::{
+    ChatContentPart, ChatMessage, ChatToolChoice, GenerationPromptMode, TestRenderRequest, Tool,
 };
 use crate::{ChatRenderer, ChatRole};
 
-fn render_request(request: &ChatRequest) -> String {
+fn render_request(request: &TestRenderRequest) -> String {
     DeepSeekV32ChatRenderer::new()
-        .render(request)
+        .render(request.as_request())
         .unwrap()
-        .prompt
+        .content
         .into_text()
         .expect("deepseek renderer should return text prompt")
 }
 
-fn render_result(request: &ChatRequest) -> Result<String, Error> {
-    DeepSeekV32ChatRenderer::new().render(request).map(|rendered| {
+fn render_result(request: &TestRenderRequest) -> Result<String, Error> {
+    DeepSeekV32ChatRenderer::new().render(request.as_request()).map(|rendered| {
         rendered
-            .prompt
+            .content
             .into_text()
             .expect("deepseek renderer should return text prompt")
     })
 }
 
-fn thinking_request(messages: Vec<ChatMessage>) -> ChatRequest {
-    let mut request = ChatRequest {
-        request_id: "deepseek-v32-small-test".to_string(),
+fn thinking_request(messages: Vec<ChatMessage>) -> TestRenderRequest {
+    let mut request = TestRenderRequest {
         messages,
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     if matches!(
         request.messages.last().map(ChatMessage::role),
@@ -53,7 +52,7 @@ fn thinking_request(messages: Vec<ChatMessage>) -> ChatRequest {
     request
 }
 
-fn fixture_request(input_name: &str) -> ChatRequest {
+fn fixture_request(input_name: &str) -> TestRenderRequest {
     fixture_chat_request(&fixture_path(input_name), deepseek_fixture_options())
 }
 
@@ -66,7 +65,7 @@ fn deepseek_fixture_options() -> FixtureRequestOptions {
 
 fn fixture_path(name: &str) -> PathBuf {
     PathBuf::from(env!("CARGO_MANIFEST_DIR"))
-        .join("src/renderer/deepseek_v32")
+        .join("src/deepseek_v32")
         .join("fixtures")
         .join(name)
 }
@@ -103,13 +102,12 @@ fn renders_official_search_fixture_with_date() {
 
 #[test]
 fn request_level_tools_are_lowered_as_synthetic_leading_system_message() {
-    let mut request = ChatRequest {
-        request_id: "deepseek-v32-tools".to_string(),
+    let mut request = TestRenderRequest {
         messages: vec![
             ChatMessage::system("System prompt."),
             ChatMessage::text(ChatRole::User, "Hello"),
         ],
-        tools: vec![ChatTool {
+        tools: vec![Tool {
             name: "lookup".to_string(),
             description: Some("Look things up".to_string()),
             parameters: json!({
@@ -124,7 +122,7 @@ fn request_level_tools_are_lowered_as_synthetic_leading_system_message() {
             strict: None,
         }],
         tool_choice: ChatToolChoice::Auto,
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     request
         .chat_options
@@ -272,12 +270,12 @@ fn assistant_after_last_user_requires_reasoning_or_tool_calls() {
 
 #[test]
 fn continue_final_assistant_omits_final_eos() {
-    let mut request = ChatRequest {
+    let mut request = TestRenderRequest {
         messages: vec![
             ChatMessage::user("write"),
             ChatMessage::assistant_text("partial answer"),
         ],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     request.chat_options.generation_prompt_mode = GenerationPromptMode::ContinueFinalAssistant;
 
@@ -289,14 +287,14 @@ fn continue_final_assistant_omits_final_eos() {
 
 #[test]
 fn render_rejects_multimodal_input() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![ChatMessage::user(vec![ChatContentPart::image_url(
             "data:image/png;base64,test",
         )])],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
-    let error = DeepSeekV32ChatRenderer::new().render(&request).unwrap_err();
+    let error = DeepSeekV32ChatRenderer::new().render(request.as_request()).unwrap_err();
 
     assert!(matches!(
         error,

rust/src/chat-renderer/src/deepseek_v4/encoding.rs

@@ -14,8 +14,10 @@ use serde_json::Value;
 use serde_json_fmt::JsonFormat;
 
 use crate::error::{Error, Result};
-use crate::request::{ChatContent, ChatMessage, ChatRequest, ChatTool, ReasoningEffort};
-use crate::{AssistantContentBlock, AssistantMessageExt, AssistantToolCall};
+use crate::{
+    AssistantContentBlock, AssistantMessageExt, AssistantToolCall, ChatContent, ChatMessage,
+    ReasoningEffort, RenderRequest, Tool,
+};
 
 const BOS_TOKEN: &str = "<｜begin▁of▁sentence｜>";
 const EOS_TOKEN: &str = "<｜end▁of▁sentence｜>";
@@ -46,14 +48,14 @@ struct RenderedToolSchema<'a> {
 }
 
 /// Render one chat request into the final prompt string.
-pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
+pub(super) fn render_request(request: &RenderRequest<'_>) -> Result<String> {
     let (thinking_mode, max_reasoning_effort) = resolve_thinking_options(request)?;
     let request_tools = request_tools(request);
     let synthetic_tool_system = needs_synthetic_tool_system(request, request_tools);
     let drop_thinking = request.parse_template_bool("drop_thinking")?.unwrap_or(true)
         && !rendered_tools_present(request, request_tools);
     let last_user_render_index =
-        find_last_user_render_index(request.messages.as_slice(), synthetic_tool_system);
+        find_last_user_render_index(request.messages, synthetic_tool_system);
     let mut out = String::from(BOS_TOKEN);
     if thinking_mode == ThinkingMode::Thinking && max_reasoning_effort {
         out.push_str(REASONING_EFFORT_MAX);
@@ -68,7 +70,7 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
     }
 
     for (message_index, message) in request.messages.iter().enumerate() {
-        if is_following_tool_response(request.messages.as_slice(), message_index) {
+        if is_following_tool_response(request.messages, message_index) {
             continue;
         }
 
@@ -101,12 +103,12 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
                 render_assistant_message(&mut out, emit_thinking_block, append_eos, content)?;
             }
             ChatMessage::ToolResponse { .. } => {
-                render_tool_response_block(&mut out, request.messages.as_slice(), message_index)?;
+                render_tool_response_block(&mut out, request.messages, message_index)?;
             }
         }
 
         if is_user_like_entry(message)
-            && next_rendered_entry_is_assistant_or_end(request.messages.as_slice(), message_index)
+            && next_rendered_entry_is_assistant_or_end(request.messages, message_index)
         {
             write_assistant_transition(
                 &mut out,
@@ -124,7 +126,7 @@ pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
 /// wrapper, the Rust renderer only consumes the typed top-level
 /// `reasoning_effort`; the generic template-kwargs map is left for HF
 /// templates.
-fn resolve_thinking_options(request: &ChatRequest) -> Result<(ThinkingMode, bool)> {
+fn resolve_thinking_options(request: &RenderRequest<'_>) -> Result<(ThinkingMode, bool)> {
     let mut thinking_mode = match request.enable_thinking()?.unwrap_or(false) {
         true => ThinkingMode::Thinking,
         false => ThinkingMode::Chat,
@@ -141,16 +143,16 @@ fn resolve_thinking_options(request: &ChatRequest) -> Result<(ThinkingMode, bool
 }
 
 /// Return request-level tools only when native tool parsing is enabled.
-fn request_tools(request: &ChatRequest) -> &[ChatTool] {
+fn request_tools<'a>(request: &RenderRequest<'a>) -> &'a [Tool] {
     if request.tool_parsing_enabled() {
-        request.tools.as_slice()
+        request.tools
     } else {
         &[]
     }
 }
 
 /// Return whether request tools need a synthetic leading system entry.
-fn needs_synthetic_tool_system(request: &ChatRequest, request_tools: &[ChatTool]) -> bool {
+fn needs_synthetic_tool_system(request: &RenderRequest<'_>, request_tools: &[Tool]) -> bool {
     !request_tools.is_empty()
         && !request
             .messages
@@ -159,7 +161,7 @@ fn needs_synthetic_tool_system(request: &ChatRequest, request_tools: &[ChatTool]
 }
 
 /// Return whether any rendered message carries tool schemas.
-fn rendered_tools_present(request: &ChatRequest, request_tools: &[ChatTool]) -> bool {
+fn rendered_tools_present(request: &RenderRequest<'_>, request_tools: &[Tool]) -> bool {
     !request_tools.is_empty()
         || request.messages.iter().any(|message| {
             matches!(
@@ -228,7 +230,7 @@ fn next_rendered_entry_is_assistant_or_end(messages: &[ChatMessage], message_ind
 }
 
 /// Render the tool preamble shown to the model, V4 flavor.
-fn render_tools(out: &mut String, tools: &[ChatTool]) -> Result<()> {
+fn render_tools(out: &mut String, tools: &[Tool]) -> Result<()> {
     out.push_str(
         r#"## Tools
 
@@ -269,7 +271,7 @@ Otherwise, output directly after </think> with tool calls or final response.
 }
 
 /// Serialize one typed tool schema into the JSON shape embedded in the prompt.
-fn render_tool_schema(out: &mut String, tool: &ChatTool) -> Result<()> {
+fn render_tool_schema(out: &mut String, tool: &Tool) -> Result<()> {
     out.push_str(&json_dumps(&RenderedToolSchema {
         name: &tool.name,
         description: tool.description.as_deref(),
@@ -283,7 +285,7 @@ fn render_tool_schema(out: &mut String, tool: &ChatTool) -> Result<()> {
 fn render_system_message(
     out: &mut String,
     content: Option<&ChatContent>,
-    tools: &[ChatTool],
+    tools: &[Tool],
 ) -> Result<()> {
     if let Some(content) = content {
         write_chat_content(out, content)?;
@@ -296,11 +298,7 @@ fn render_system_message(
 }
 
 /// Developer messages are rendered as user-like turns with optional tools.
-fn render_developer_message(
-    out: &mut String,
-    content: &ChatContent,
-    tools: &[ChatTool],
-) -> Result<()> {
+fn render_developer_message(out: &mut String, content: &ChatContent, tools: &[Tool]) -> Result<()> {
     if content.is_empty() {
         return Err(Error::ChatTemplate(
             "invalid DeepSeek V4 developer message: empty content".to_string(),
@@ -519,7 +517,7 @@ fn write_chat_content(out: &mut String, content: &ChatContent) -> Result<()> {
         ChatContent::Text(text) => out.push_str(text),
         ChatContent::Parts(parts) => {
             for part in parts {
-                out.push_str(part.as_text()?);
+                out.push_str(part.as_text().map_err(Error::UnsupportedMultimodalContent)?);
             }
         }
     }

rust/src/chat-renderer/src/deepseek_v4/mod.rs

@@ -3,29 +3,29 @@
 
 mod encoding;
 
-use vllm_text::Prompt;
-
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{
+    ChatRenderer, RenderRequest, RenderedPrompt, RenderedPromptContent, request_template_kwargs,
+};
 use crate::Result;
-use crate::request::ChatRequest;
 
 /// Dedicated DeepSeek V4 renderer.
 #[derive(Debug, Clone, Copy, Default)]
 pub struct DeepSeekV4ChatRenderer;
 
 impl DeepSeekV4ChatRenderer {
+    /// Create the dedicated DeepSeek V4 renderer.
     pub fn new() -> Self {
         Self
     }
 }
 
 impl ChatRenderer for DeepSeekV4ChatRenderer {
-    fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
+    fn render(&self, request: RenderRequest<'_>) -> Result<RenderedPrompt> {
         request.validate()?;
 
         Ok(RenderedPrompt {
-            prompt: Prompt::Text(encoding::render_request(request)?),
-            effective_template_kwargs: request_template_kwargs(request),
+            content: RenderedPromptContent::Text(encoding::render_request(&request)?),
+            effective_template_kwargs: request_template_kwargs(&request),
         })
     }
 }

rust/src/chat-renderer/src/deepseek_v4/tests.rs

@@ -8,20 +8,20 @@ use serde_json::Value;
 
 use super::DeepSeekV4ChatRenderer;
 use crate::ChatRenderer;
-use crate::event::{AssistantContentBlock, AssistantToolCall};
-use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
-use crate::request::{ChatMessage, ChatRequest, GenerationPromptMode, ReasoningEffort};
+use crate::test_utils::{FixtureRequestOptions, fixture_chat_request};
+use crate::{AssistantContentBlock, AssistantToolCall};
+use crate::{ChatMessage, GenerationPromptMode, ReasoningEffort, TestRenderRequest};
 
-fn render_request(request: &ChatRequest) -> String {
+fn render_request(request: &TestRenderRequest) -> String {
     DeepSeekV4ChatRenderer::new()
-        .render(request)
+        .render(request.as_request())
         .unwrap()
-        .prompt
+        .content
         .into_text()
         .expect("deepseek v4 renderer should return text prompt")
 }
 
-fn fixture_request(input_name: &str) -> ChatRequest {
+fn fixture_request(input_name: &str) -> TestRenderRequest {
     fixture_chat_request(&fixture_path(input_name), deepseek_fixture_options())
 }
 
@@ -34,7 +34,7 @@ fn deepseek_fixture_options() -> FixtureRequestOptions {
 
 fn fixture_path(name: &str) -> PathBuf {
     PathBuf::from(env!("CARGO_MANIFEST_DIR"))
-        .join("src/renderer/deepseek_v4")
+        .join("src/deepseek_v4")
         .join("fixtures")
         .join(name)
 }
@@ -63,9 +63,9 @@ fn renders_v4_fixture_2_multi_turn_drop_thinking() {
 
 #[test]
 fn reasoning_effort_max_adds_prefix_when_thinking_is_enabled() {
-    let mut request = ChatRequest {
+    let mut request = TestRenderRequest {
         messages: vec![ChatMessage::user("solve it")],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     request
         .chat_options
@@ -86,9 +86,9 @@ fn reasoning_effort_max_adds_prefix_when_thinking_is_enabled() {
 
 #[test]
 fn reasoning_effort_none_disables_thinking() {
-    let mut request = ChatRequest {
+    let mut request = TestRenderRequest {
         messages: vec![ChatMessage::user("answer directly")],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     request
         .chat_options
@@ -104,9 +104,9 @@ fn reasoning_effort_none_disables_thinking() {
 
 #[test]
 fn reasoning_effort_template_kwarg_is_ignored() {
-    let mut request = ChatRequest {
+    let mut request = TestRenderRequest {
         messages: vec![ChatMessage::user("solve it")],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     request
         .chat_options
@@ -124,7 +124,7 @@ fn reasoning_effort_template_kwarg_is_ignored() {
 
 #[test]
 fn tool_results_are_sorted_by_previous_assistant_tool_call_order() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![
             ChatMessage::assistant_blocks(vec![
                 AssistantContentBlock::ToolCall(AssistantToolCall {
@@ -141,7 +141,7 @@ fn tool_results_are_sorted_by_previous_assistant_tool_call_order() {
             ChatMessage::tool_response("first result", "first"),
             ChatMessage::tool_response("second result", "second"),
         ],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
     let rendered = render_request(&request);
@@ -164,7 +164,7 @@ fn tool_results_are_sorted_by_previous_assistant_tool_call_order() {
 
 #[test]
 fn drop_thinking_false_keeps_prior_assistant_reasoning() {
-    let mut request = ChatRequest {
+    let mut request = TestRenderRequest {
         messages: vec![
             ChatMessage::assistant_blocks(vec![
                 AssistantContentBlock::Reasoning {
@@ -176,7 +176,7 @@ fn drop_thinking_false_keeps_prior_assistant_reasoning() {
             ]),
             ChatMessage::user("next"),
         ],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     request
         .chat_options
@@ -197,16 +197,16 @@ fn drop_thinking_false_keeps_prior_assistant_reasoning() {
 
 #[test]
 fn continue_final_assistant_omits_final_eos() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![
             ChatMessage::user("write"),
             ChatMessage::assistant_text("partial answer"),
         ],
-        chat_options: crate::request::ChatOptions {
+        chat_options: crate::ChatOptions {
             generation_prompt_mode: GenerationPromptMode::ContinueFinalAssistant,
             ..Default::default()
         },
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
     let rendered = render_request(&request);

rust/src/chat-renderer/src/error.rs

@@ -0,0 +1,54 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use thiserror::Error;
+
+type BoxedError = Box<dyn std::error::Error + Send + Sync>;
+
+/// Error returned while constructing or applying a chat renderer.
+#[derive(Debug, Error)]
+pub enum Error {
+    /// Rendering requires at least one chat message.
+    #[error("chat request must contain at least one message")]
+    EmptyMessages,
+    /// Continuation mode requires the final message to be an assistant turn.
+    #[error("cannot continue the final message when the last message is not from the assistant")]
+    ContinueFinalAssistantWithoutFinalAssistant,
+    /// The selected renderer requires a chat template and none was configured.
+    #[error("chat template is required but none was configured")]
+    MissingChatTemplate,
+    /// A chat template could not be compiled or applied.
+    #[error("chat template error: {0}")]
+    ChatTemplate(String),
+    /// The selected renderer cannot represent the given multimodal part.
+    #[error("unsupported multimodal content: {0}")]
+    UnsupportedMultimodalContent(&'static str),
+    /// The process-wide GPT-OSS Harmony encoding could not be initialized.
+    #[error("failed to initialize the Harmony encoding")]
+    HarmonyEncoding {
+        /// Underlying Harmony initialization failure.
+        #[source]
+        error: BoxedError,
+    },
+    /// Tokenizer construction or encoding failed.
+    #[error(transparent)]
+    Tokenizer(#[from] vllm_tokenizer::TokenizerError),
+}
+
+impl Error {
+    /// Whether this error should be reported as invalid request input when
+    /// raised while rendering.
+    pub fn is_request_validation_error(&self) -> bool {
+        matches!(
+            self,
+            Self::EmptyMessages
+                | Self::ContinueFinalAssistantWithoutFinalAssistant
+                | Self::MissingChatTemplate
+                | Self::ChatTemplate(_)
+                | Self::UnsupportedMultimodalContent(_)
+        )
+    }
+}
+
+/// Result returned by chat renderer operations.
+pub type Result<T> = std::result::Result<T, Error>;

rust/src/chat-renderer/src/harmony/encoding.rs

@@ -12,13 +12,13 @@ use thiserror_ext::AsReport as _;
 use crate::error::{Error, Result};
 
 /// Lazily load the shared GPT-OSS Harmony encoding once per process.
-pub(crate) fn harmony_encoding() -> Result<&'static HarmonyEncoding> {
+pub fn harmony_encoding() -> Result<&'static HarmonyEncoding> {
     static ENCODING: LazyLock<anyhow::Result<HarmonyEncoding>> = LazyLock::new(|| {
         load_harmony_encoding(HarmonyEncodingName::HarmonyGptOss)
             .context("failed to load harmony encoding for gpt-oss")
     });
 
-    ENCODING.as_ref().map_err(|error| Error::HarmonyOutputParsing {
+    ENCODING.as_ref().map_err(|error| Error::HarmonyEncoding {
         error: error.to_report_string().into(),
     })
 }

rust/src/chat-renderer/src/harmony/mod.rs

@@ -3,7 +3,7 @@
 
 //! Native Harmony chat renderer for `gpt_oss`.
 
-pub(crate) mod encoding;
+mod encoding;
 
 use openai_harmony::HarmonyEncoding;
 use openai_harmony::chat::{
@@ -12,14 +12,16 @@ use openai_harmony::chat::{
 };
 use thiserror_ext::AsReport as _;
 use time::macros::format_description;
-use vllm_text::Prompt;
 
-use self::encoding::harmony_encoding;
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+pub use self::encoding::harmony_encoding;
+use super::{
+    ChatRenderer, RenderRequest, RenderedPrompt, RenderedPromptContent, request_template_kwargs,
+};
 use crate::error::{Error, Result};
-use crate::event::AssistantContentBlock;
-use crate::request::{ChatContent, ChatMessage, ChatRequest, ChatTool, GenerationPromptMode};
-use crate::{AssistantMessageExt as _, ReasoningEffort};
+use crate::{
+    AssistantContentBlock, AssistantMessageExt as _, ChatContent, ChatMessage,
+    GenerationPromptMode, ReasoningEffort, Tool,
+};
 
 const SYSTEM_START_DATE_ENV: &str = "VLLM_SYSTEM_START_DATE";
 const HARMONY_SYSTEM_INSTRUCTIONS_ENV: &str = "VLLM_GPT_OSS_HARMONY_SYSTEM_INSTRUCTIONS";
@@ -73,7 +75,7 @@ impl HarmonyChatRenderer {
     ///
     /// Harmony owns both prompt formatting and tokenization, so the Rust
     /// frontend bypasses the generic HF tokenizer path for GPT-OSS input.
-    fn render_token_ids(&self, request: &ChatRequest) -> Result<Vec<u32>> {
+    fn render_token_ids(&self, request: &RenderRequest<'_>) -> Result<Vec<u32>> {
         if request.has_multimodal() {
             return Err(Error::UnsupportedMultimodalContent("image_url"));
         }
@@ -111,12 +113,12 @@ impl HarmonyChatRenderer {
 }
 
 impl ChatRenderer for HarmonyChatRenderer {
-    /// Render a chat request as [`Prompt::TokenIds`] with template kwargs echoed
-    /// for downstream accounting/debugging.
-    fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
+    /// Render a chat request as token IDs with template kwargs echoed for
+    /// downstream accounting/debugging.
+    fn render(&self, request: RenderRequest<'_>) -> Result<RenderedPrompt> {
         Ok(RenderedPrompt {
-            prompt: Prompt::TokenIds(self.render_token_ids(request)?),
-            effective_template_kwargs: request_template_kwargs(request),
+            content: RenderedPromptContent::TokenIds(self.render_token_ids(&request)?),
+            effective_template_kwargs: request_template_kwargs(&request),
         })
     }
 }
@@ -126,10 +128,10 @@ impl ChatRenderer for HarmonyChatRenderer {
 /// This adds the Harmony system/developer preamble, peels at most one leading
 /// system/developer instruction message, and then lowers the remaining chat
 /// history message-by-message.
-fn to_harmony_messages(request: &ChatRequest, options: &Options) -> Result<Vec<Message>> {
+fn to_harmony_messages(request: &RenderRequest<'_>, options: &Options) -> Result<Vec<Message>> {
     let (instructions, leading_developer_tools, remaining_messages) =
-        peel_leading_instructions(&request.messages)?;
-    let tool_call_names = tool_call_names(&request.messages);
+        peel_leading_instructions(request.messages)?;
+    let tool_call_names = tool_call_names(request.messages);
     let mut messages =
         build_harmony_preamble(request, instructions, leading_developer_tools, options)?;
 
@@ -147,7 +149,7 @@ fn to_harmony_messages(request: &ChatRequest, options: &Options) -> Result<Vec<M
 #[allow(clippy::type_complexity)]
 fn peel_leading_instructions(
     messages: &[ChatMessage],
-) -> Result<(Option<String>, Option<&[ChatTool]>, &[ChatMessage])> {
+) -> Result<(Option<String>, Option<&[Tool]>, &[ChatMessage])> {
     let Some(first) = messages.first() else {
         return Ok((None, None, messages));
     };
@@ -172,9 +174,9 @@ fn peel_leading_instructions(
 /// a developer message depending on `use_system_instructions`; request-level and
 /// leading developer tools are attached to the developer message.
 fn build_harmony_preamble(
-    request: &ChatRequest,
+    request: &RenderRequest<'_>,
     instructions: Option<String>,
-    leading_developer_tools: Option<&[ChatTool]>,
+    leading_developer_tools: Option<&[Tool]>,
     options: &Options,
 ) -> Result<Vec<Message>> {
     let mut messages = vec![Message::from_role_and_content(
@@ -211,12 +213,12 @@ fn build_harmony_preamble(
 
 /// Collect request-level and leading developer function tools for the preamble.
 fn preamble_tool_descriptions(
-    request: &ChatRequest,
-    leading_developer_tools: Option<&[ChatTool]>,
+    request: &RenderRequest<'_>,
+    leading_developer_tools: Option<&[Tool]>,
 ) -> Vec<ToolDescription> {
     let mut tools = Vec::new();
     if request.tool_parsing_enabled() {
-        tools.extend(to_tool_descriptions(&request.tools));
+        tools.extend(to_tool_descriptions(request.tools));
     }
     if let Some(leading_developer_tools) = leading_developer_tools {
         tools.extend(to_tool_descriptions(leading_developer_tools));
@@ -311,7 +313,7 @@ fn to_harmony_message(
 fn system_or_developer_message(
     role: &str,
     instructions: String,
-    tools: Option<&[ChatTool]>,
+    tools: Option<&[Tool]>,
     options: &Options,
 ) -> Result<Message> {
     if role == "system" && options.use_system_instructions {
@@ -325,7 +327,7 @@ fn system_or_developer_message(
 }
 
 /// Build a Harmony developer message with optional instructions and function tools.
-fn developer_message(instructions: Option<String>, tools: Option<&[ChatTool]>) -> Message {
+fn developer_message(instructions: Option<String>, tools: Option<&[Tool]>) -> Message {
     let mut content = DeveloperContent::new();
     if let Some(instructions) = instructions.filter(|text| !text.is_empty()) {
         content = content.with_instructions(instructions);
@@ -427,11 +429,11 @@ fn auto_drop_analysis_messages(messages: Vec<Message>) -> Vec<Message> {
 
 /// Flatten vLLM text content and reject unsupported multimodal parts.
 fn flatten_text(content: &ChatContent) -> Result<String> {
-    content.try_flatten_to_text()
+    content.try_flatten_to_text().map_err(Error::UnsupportedMultimodalContent)
 }
 
 /// Convert vLLM function tool definitions to Harmony tool descriptions.
-fn to_tool_descriptions(tools: &[ChatTool]) -> Vec<ToolDescription> {
+fn to_tool_descriptions(tools: &[Tool]) -> Vec<ToolDescription> {
     tools
         .iter()
         .map(|tool| {

rust/src/chat-renderer/src/harmony/tests.rs

@@ -10,15 +10,15 @@ use super::HarmonyChatRenderer;
 use super::encoding::harmony_encoding;
 use crate::ChatRenderer;
 use crate::error::Error;
-use crate::event::{AssistantContentBlock, AssistantToolCall};
-use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
-use crate::request::{
-    ChatContentPart, ChatMessage, ChatRequest, GenerationPromptMode, ReasoningEffort,
+use crate::test_utils::{FixtureRequestOptions, fixture_chat_request};
+use crate::{AssistantContentBlock, AssistantToolCall};
+use crate::{
+    ChatContentPart, ChatMessage, GenerationPromptMode, ReasoningEffort, TestRenderRequest,
 };
 
 const PINNED_DATE: &str = "2025-06-28";
 
-fn fixture_request(input_name: &str) -> ChatRequest {
+fn fixture_request(input_name: &str) -> TestRenderRequest {
     fixture_chat_request(
         &fixture_path(input_name),
         FixtureRequestOptions {
@@ -30,7 +30,7 @@ fn fixture_request(input_name: &str) -> ChatRequest {
 
 fn fixture_path(name: &str) -> PathBuf {
     PathBuf::from(env!("CARGO_MANIFEST_DIR"))
-        .join("src/renderer/harmony")
+        .join("src/harmony")
         .join("fixtures")
         .join(name)
 }
@@ -39,24 +39,24 @@ fn test_renderer(use_system_instructions: bool) -> HarmonyChatRenderer {
     HarmonyChatRenderer::with_options(PINNED_DATE, use_system_instructions).unwrap()
 }
 
-fn render_token_ids(request: &ChatRequest) -> Vec<u32> {
+fn render_token_ids(request: &TestRenderRequest) -> Vec<u32> {
     render_token_ids_with(&test_renderer(false), request)
 }
 
-fn render_token_ids_with(renderer: &HarmonyChatRenderer, request: &ChatRequest) -> Vec<u32> {
+fn render_token_ids_with(renderer: &HarmonyChatRenderer, request: &TestRenderRequest) -> Vec<u32> {
     renderer
-        .render(request)
+        .render(request.as_request())
         .unwrap()
-        .prompt
+        .content
         .into_token_ids()
         .expect("Harmony renderer returns token IDs")
 }
 
-fn render_prompt_text(request: &ChatRequest) -> String {
+fn render_prompt_text(request: &TestRenderRequest) -> String {
     render_prompt_text_with(&test_renderer(false), request)
 }
 
-fn render_prompt_text_with(renderer: &HarmonyChatRenderer, request: &ChatRequest) -> String {
+fn render_prompt_text_with(renderer: &HarmonyChatRenderer, request: &TestRenderRequest) -> String {
     let token_ids = render_token_ids_with(renderer, request);
     harmony_encoding().unwrap().tokenizer().decode_utf8(&token_ids).unwrap()
 }
@@ -137,10 +137,10 @@ fn drops_stale_analysis_fixture() {
 
 #[test]
 fn rejects_invalid_reasoning_effort() {
-    let mut request = ChatRequest::for_test();
+    let mut request = TestRenderRequest::for_test();
     request.chat_options.reasoning_effort = Some(ReasoningEffort::None);
 
-    let error = test_renderer(false).render(&request).unwrap_err();
+    let error = test_renderer(false).render(request.as_request()).unwrap_err();
 
     expect![[r#"chat template error: reasoning_effort="none" is not supported by Harmony. Supported values are: low, medium, high."#]]
         .assert_eq(&error.to_report_string());
@@ -148,7 +148,7 @@ fn rejects_invalid_reasoning_effort() {
 
 #[test]
 fn rejects_unknown_tool_response_id() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![
             ChatMessage::assistant_blocks(vec![AssistantContentBlock::ToolCall(
                 AssistantToolCall {
@@ -159,10 +159,10 @@ fn rejects_unknown_tool_response_id() {
             )]),
             ChatMessage::tool_response("{}", "call-unknown"),
         ],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
-    let error = test_renderer(false).render(&request).unwrap_err();
+    let error = test_renderer(false).render(request.as_request()).unwrap_err();
 
     expect![
         "chat template error: invalid Harmony tool message: unknown tool_call_id `call-unknown`"
@@ -172,14 +172,14 @@ fn rejects_unknown_tool_response_id() {
 
 #[test]
 fn rejects_multimodal_input() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![ChatMessage::user(vec![ChatContentPart::image_url(
             "data:image/png;base64,test",
         )])],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
-    let error = test_renderer(false).render(&request).unwrap_err();
+    let error = test_renderer(false).render(request.as_request()).unwrap_err();
 
     assert!(matches!(
         error,
@@ -189,16 +189,16 @@ fn rejects_multimodal_input() {
 
 #[test]
 fn rejects_continue_final_assistant() {
-    let mut request = ChatRequest {
+    let mut request = TestRenderRequest {
         messages: vec![
             ChatMessage::user("write"),
             ChatMessage::assistant_text("partial"),
         ],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
     request.chat_options.generation_prompt_mode = GenerationPromptMode::ContinueFinalAssistant;
 
-    let error = test_renderer(false).render(&request).unwrap_err();
+    let error = test_renderer(false).render(request.as_request()).unwrap_err();
 
     expect!["chat template error: Harmony renderer does not support continue_final_message"]
         .assert_eq(&error.to_report_string());

rust/src/chat-renderer/src/hf/format.rs

@@ -33,8 +33,11 @@ pub enum ChatTemplateContentFormatOption {
 }
 
 impl ChatTemplateContentFormatOption {
+    /// CLI/config literal for automatic content-format detection.
     pub const AUTO_LITERAL: &str = "auto";
+    /// CLI/config literal for OpenAI-compatible structured content.
     pub const OPENAI_LITERAL: &str = "openai";
+    /// CLI/config literal for flattened string content.
     pub const STRING_LITERAL: &str = "string";
 }
 

rust/src/chat-renderer/src/hf/mod.rs

@@ -1,27 +1,31 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+//! Hugging Face Jinja chat-template renderer.
+
 use std::collections::HashMap;
+use std::path::Path;
 
 use serde::Serialize;
 use serde_json::Value as JsonValue;
 use thiserror_ext::AsReport as _;
 use tracing::{info, trace, warn};
-use vllm_text::Prompt;
-use vllm_text::backend::hf::{
-    HfSpecialTokens, HfTokenizerConfig, ResolvedModelFiles, load_tokenizer_config,
-};
+use vllm_tokenizer::{HfSpecialTokens, HfTokenizerConfig, load_tokenizer_config};
 
 use self::format::{
     ChatTemplateContentFormat, ChatTemplateContentFormatOption as ContentFormatOption,
 };
-use self::template::{CompiledChatTemplate, TemplateContext};
+use self::template::{
+    CompiledChatTemplate, TemplateContext, load_chat_template, resolve_chat_template,
+};
 use self::value::{TemplateValue, to_template_value};
-use super::{ChatRenderer, RenderedPrompt, effective_template_kwargs};
+use super::{
+    ChatRenderer, RenderRequest, RenderedPrompt, RenderedPromptContent, effective_template_kwargs,
+};
 use crate::error::Result;
-use crate::request::{ChatContent, ChatContentPart, ChatMessage, ChatRequest};
 use crate::{
-    AssistantContentBlock, AssistantMessageExt, ChatTool, Error, LoadModelBackendsOptions,
+    AssistantContentBlock, AssistantMessageExt, ChatContent, ChatContentPart, ChatMessage, Error,
+    Tool,
 };
 
 mod error;
@@ -30,8 +34,6 @@ mod template;
 mod tojson;
 mod value;
 
-pub use template::{load_chat_template, resolve_chat_template};
-
 pub use self::format::ChatTemplateContentFormatOption;
 
 /// Template-visible placeholder tokens per supported modality.
@@ -40,11 +42,36 @@ pub use self::format::ChatTemplateContentFormatOption;
 /// content parts of that modality are rejected during rendering.
 #[derive(Debug, Clone, Default)]
 pub struct MultimodalRenderInfo {
+    /// Template-visible image placeholder token.
     pub image_token: Option<String>,
+    /// Template-visible video placeholder token.
     pub video_token: Option<String>,
+    /// Template-visible audio placeholder token.
     pub audio_token: Option<String>,
 }
 
+/// Resolved local files consumed while constructing an HF renderer.
+#[derive(Debug, Clone, Copy, Default)]
+pub struct HfRendererFiles<'a> {
+    /// Resolved `tokenizer_config.json`, when available.
+    pub tokenizer_config: Option<&'a Path>,
+    /// Resolved dedicated chat-template file, when available.
+    pub chat_template: Option<&'a Path>,
+}
+
+/// Model-level options used to construct an HF renderer.
+#[derive(Debug, Clone, Default)]
+pub struct HfRendererConfig {
+    /// Optional inline or file-backed chat-template override.
+    pub chat_template: Option<String>,
+    /// Default template kwargs merged before request-level kwargs.
+    pub default_template_kwargs: HashMap<String, JsonValue>,
+    /// How `message.content` should be represented in the template.
+    pub content_format: ContentFormatOption,
+    /// Template-visible multimodal placeholder tokens.
+    pub multimodal: Option<MultimodalRenderInfo>,
+}
+
 /// Hugging Face chat-template renderer backed by the local Jinja chat-template
 /// state.
 pub struct HfChatRenderer {
@@ -76,27 +103,25 @@ impl HfChatRenderer {
         })
     }
 
+    /// Attach named special tokens exposed to chat templates.
     pub fn with_special_tokens(mut self, special_tokens: Option<HfSpecialTokens>) -> Self {
         self.special_tokens = special_tokens;
         self
     }
 
+    /// Attach template-visible multimodal placeholder tokens.
     pub fn with_multimodal(mut self, multimodal: Option<MultimodalRenderInfo>) -> Self {
         self.multimodal = multimodal;
         self
     }
 
-    /// Create a renderer from the given model files and loading options.
-    pub fn load(
-        files: &ResolvedModelFiles,
-        options: LoadModelBackendsOptions,
-        multimodal: Option<MultimodalRenderInfo>,
-    ) -> Result<Self> {
+    /// Create a renderer from resolved local model files and renderer options.
+    pub fn load(files: HfRendererFiles<'_>, options: HfRendererConfig) -> Result<Self> {
         let HfTokenizerConfig {
             special_tokens,
             chat_template,
             ..
-        } = load_tokenizer_config(files.tokenizer_config_path.as_deref())?;
+        } = load_tokenizer_config(files.tokenizer_config)?;
         let mut template = chat_template;
         let special_tokens = (!special_tokens.is_empty()).then_some(special_tokens);
 
@@ -106,7 +131,7 @@ impl HfChatRenderer {
                     .map_err(|error| Error::ChatTemplate(error.to_report_string()))?,
             );
             info!("using configured chat template override");
-        } else if let Some(chat_template_path) = files.chat_template_path.as_deref() {
+        } else if let Some(chat_template_path) = files.chat_template {
             // If independent chat template file(s) exist and contain non-empty content,
             // they take priority over template entries in the tokenizer config
             let file_template = load_chat_template(chat_template_path)
@@ -128,11 +153,11 @@ impl HfChatRenderer {
 
         Ok(Self::new(
             template,
-            options.default_chat_template_kwargs,
-            options.chat_template_content_format,
+            options.default_template_kwargs,
+            options.content_format,
         )?
         .with_special_tokens(special_tokens)
-        .with_multimodal(multimodal))
+        .with_multimodal(options.multimodal))
     }
 
     /// Apply the chat template to one chat request, rendering the prompt string
@@ -141,7 +166,7 @@ impl HfChatRenderer {
     /// If the request carries a per-request `chat_template` override, a
     /// temporary template is compiled from that string and used instead of
     /// the model's default.
-    fn apply_chat_template(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
+    fn apply_chat_template(&self, request: &RenderRequest<'_>) -> Result<RenderedPrompt> {
         let override_template = request
             .chat_options
             .chat_template
@@ -162,10 +187,10 @@ impl HfChatRenderer {
     fn apply_chat_template_inner(
         &self,
         effective_template: &CompiledChatTemplate,
-        request: &ChatRequest,
+        request: &RenderRequest<'_>,
     ) -> Result<RenderedPrompt> {
         let mut messages = to_template_messages(
-            &request.messages,
+            request.messages,
             effective_template.content_format(),
             self.multimodal.as_ref(),
         )?;
@@ -181,7 +206,7 @@ impl HfChatRenderer {
             None
         };
 
-        let tools = request.tool_parsing_enabled().then(|| to_template_tools(&request.tools));
+        let tools = request.tool_parsing_enabled().then(|| to_template_tools(request.tools));
         trace!(
             message_count = messages.len(),
             content_format = ?effective_template.content_format(),
@@ -198,7 +223,7 @@ impl HfChatRenderer {
                 add_generation_prompt: request.chat_options.add_generation_prompt(),
                 continue_final_message: request.chat_options.continue_final_message(),
                 tools: tools.as_deref(),
-                documents: request.documents.as_deref(),
+                documents: request.documents,
                 template_kwargs: Some(&effective_template_kwargs),
                 special_tokens: self.special_tokens.as_ref(),
             })
@@ -217,15 +242,15 @@ impl HfChatRenderer {
         );
 
         Ok(RenderedPrompt {
-            prompt: Prompt::Text(prompt),
+            content: RenderedPromptContent::Text(prompt),
             effective_template_kwargs,
         })
     }
 }
 
 impl ChatRenderer for HfChatRenderer {
-    fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
-        self.apply_chat_template(request)
+    fn render(&self, request: RenderRequest<'_>) -> Result<RenderedPrompt> {
+        self.apply_chat_template(&request)
     }
 }
 
@@ -556,7 +581,7 @@ fn truncate_prompt_at_continue_final_message_tag(
     Ok(rendered)
 }
 
-fn to_template_tools(tools: &[ChatTool]) -> Vec<TemplateTool> {
+fn to_template_tools(tools: &[Tool]) -> Vec<TemplateTool> {
     tools
         .iter()
         .map(|tool| TemplateTool {
@@ -577,42 +602,39 @@ mod tests {
 
     use expect_test::expect;
     use serde_json::Value;
-    use vllm_text::Prompt;
-    use vllm_text::backend::hf::{HfSpecialTokens, NamedSpecialToken};
+    use vllm_tokenizer::{HfSpecialTokens, NamedSpecialToken};
 
     use super::{ChatTemplateContentFormatOption, HfChatRenderer, MultimodalRenderInfo};
-    use crate::request::{
-        ChatContentPart, ChatMessage, ChatRequest, ChatRole, ChatTool, ChatToolChoice,
-        GenerationPromptMode, ReasoningEffort,
+    use crate::{
+        AssistantContentBlock, ChatContentPart, ChatMessage, ChatRenderer, ChatRole,
+        ChatToolChoice, Error, GenerationPromptMode, ReasoningEffort, RenderedPromptContent,
+        Result, TestRenderRequest, Tool,
     };
-    use crate::{AssistantContentBlock, ChatRenderer, Error, Result};
 
-    const QWEN3_0_6B_TEMPLATE: &str = include_str!("../../../tests/templates/qwen3.jinja");
-    const QWEN3_5_0_8B_TEMPLATE: &str = include_str!("../../../tests/templates/qwen35.jinja");
+    const QWEN3_0_6B_TEMPLATE: &str = include_str!("../../../chat/tests/templates/qwen3.jinja");
+    const QWEN3_5_0_8B_TEMPLATE: &str = include_str!("../../../chat/tests/templates/qwen35.jinja");
 
-    fn sample_request(messages: Vec<ChatMessage>) -> ChatRequest {
-        ChatRequest {
+    fn sample_request(messages: Vec<ChatMessage>) -> TestRenderRequest {
+        TestRenderRequest {
             messages,
-            request_id: "render-test".to_string(),
-            ..ChatRequest::for_test()
+            ..TestRenderRequest::for_test()
         }
     }
 
-    fn render(template: Option<&str>, request: &ChatRequest) -> Result<String> {
-        HfChatRenderer::new(
+    fn render(template: Option<&str>, request: &TestRenderRequest) -> Result<String> {
+        let rendered = HfChatRenderer::new(
             template.map(str::to_owned),
             HashMap::new(),
             ChatTemplateContentFormatOption::Auto,
         )?
-        .render(request)?
-        .prompt
-        .into_text()
-        .map_err(|_| unreachable!("HF renderer should return text prompt"))
+        .render(request.as_request())?
+        .content;
+        Ok(rendered.into_text().expect("HF renderer should return text prompt"))
     }
 
     fn render_mm(
         template: &str,
-        request: &ChatRequest,
+        request: &TestRenderRequest,
         content_format: ChatTemplateContentFormatOption,
     ) -> Result<crate::RenderedPrompt> {
         HfChatRenderer::new(Some(template.to_string()), HashMap::new(), content_format)?
@@ -621,10 +643,10 @@ mod tests {
                 video_token: Some("<video>".to_string()),
                 audio_token: Some("<audio>".to_string()),
             }))
-            .render(request)
+            .render(request.as_request())
     }
 
-    fn image_request() -> ChatRequest {
+    fn image_request() -> TestRenderRequest {
         sample_request(vec![ChatMessage::user(vec![
             ChatContentPart::text("a"),
             ChatContentPart::image_url("data:image/png;base64,test"),
@@ -632,7 +654,7 @@ mod tests {
         ])])
     }
 
-    fn video_request() -> ChatRequest {
+    fn video_request() -> TestRenderRequest {
         sample_request(vec![ChatMessage::user(vec![
             ChatContentPart::text("a"),
             ChatContentPart::video_url("https://example.com/demo.mp4"),
@@ -640,7 +662,7 @@ mod tests {
         ])])
     }
 
-    fn audio_request() -> ChatRequest {
+    fn audio_request() -> TestRenderRequest {
         sample_request(vec![ChatMessage::user(vec![
             ChatContentPart::text("a"),
             ChatContentPart::input_audio("dGVzdA==", Some("wav".to_string())),
@@ -658,7 +680,10 @@ mod tests {
         )
         .unwrap();
 
-        assert_eq!(rendered.prompt, Prompt::Text("a<image>b".to_string()));
+        assert_eq!(
+            rendered.content,
+            RenderedPromptContent::Text("a<image>b".to_string())
+        );
     }
 
     #[test]
@@ -670,7 +695,10 @@ mod tests {
         )
         .unwrap();
 
-        assert_eq!(rendered.prompt, Prompt::Text("a<video>b".to_string()));
+        assert_eq!(
+            rendered.content,
+            RenderedPromptContent::Text("a<video>b".to_string())
+        );
     }
 
     #[test]
@@ -683,8 +711,8 @@ mod tests {
         .unwrap();
 
         assert_eq!(
-            rendered.prompt,
-            Prompt::Text("a<audio><audio>b".to_string())
+            rendered.content,
+            RenderedPromptContent::Text("a<audio><audio>b".to_string())
         );
     }
 
@@ -697,7 +725,10 @@ mod tests {
         )
         .unwrap();
 
-        assert_eq!(rendered.prompt, Prompt::Text("a<|image_pad|>b".to_string()));
+        assert_eq!(
+            rendered.content,
+            RenderedPromptContent::Text("a<|image_pad|>b".to_string())
+        );
     }
 
     #[test]
@@ -709,7 +740,10 @@ mod tests {
         )
         .unwrap();
 
-        assert_eq!(rendered.prompt, Prompt::Text("a<|video_pad|>b".to_string()));
+        assert_eq!(
+            rendered.content,
+            RenderedPromptContent::Text("a<|video_pad|>b".to_string())
+        );
     }
 
     #[test]
@@ -722,8 +756,8 @@ mod tests {
         .unwrap();
 
         assert_eq!(
-            rendered.prompt,
-            Prompt::Text("a<|audio_pad|><|audio_pad|>b".to_string())
+            rendered.content,
+            RenderedPromptContent::Text("a<|audio_pad|><|audio_pad|>b".to_string())
         );
     }
 
@@ -740,7 +774,7 @@ mod tests {
             video_token: None,
             audio_token: None,
         }))
-        .render(&video_request())
+        .render(video_request().as_request())
         .unwrap_err();
 
         assert!(matches!(
@@ -842,12 +876,12 @@ mod tests {
             ChatTemplateContentFormatOption::OpenAi,
         )
         .unwrap()
-        .prompt;
+        .content;
 
         // Anything rendered after the continued text (here the image
         // placeholder and the end marker) is truncated away, matching
         // transformers.
-        assert_eq!(rendered, Prompt::Text("Sure,".to_string()));
+        assert_eq!(rendered, RenderedPromptContent::Text("Sure,".to_string()));
     }
 
     #[test]
@@ -904,7 +938,7 @@ mod tests {
     fn chat_template_exposes_developer_tools() {
         let request = sample_request(vec![ChatMessage::developer(
             "policy",
-            Some(vec![ChatTool {
+            Some(vec![Tool {
                 name: "get_weather".to_string(),
                 description: Some("Get weather".to_string()),
                 parameters: serde_json::json!({
@@ -1005,10 +1039,13 @@ mod tests {
         )
         .unwrap()
         .with_special_tokens(Some(special_tokens))
-        .apply_chat_template(&request)
+        .apply_chat_template(&request.as_request())
         .unwrap();
 
-        assert_eq!(rendered.prompt, Prompt::Text("<bos>|true".to_string()));
+        assert_eq!(
+            rendered.content,
+            RenderedPromptContent::Text("<bos>|true".to_string())
+        );
     }
 
     #[test]
@@ -1046,11 +1083,14 @@ mod tests {
             ChatTemplateContentFormatOption::String,
         )
         .unwrap()
-        .render(&request)
+        .render(request.as_request())
         .unwrap()
-        .prompt;
+        .content;
 
-        assert_eq!(rendered, Prompt::Text("hello world".to_string()));
+        assert_eq!(
+            rendered,
+            RenderedPromptContent::Text("hello world".to_string())
+        );
     }
 
     #[test]
@@ -1066,11 +1106,14 @@ mod tests {
             ChatTemplateContentFormatOption::OpenAi,
         )
         .unwrap()
-        .render(&request)
+        .render(request.as_request())
         .unwrap()
-        .prompt;
+        .content;
 
-        assert_eq!(rendered, Prompt::Text("hello world".to_string()));
+        assert_eq!(
+            rendered,
+            RenderedPromptContent::Text("hello world".to_string())
+        );
     }
 
     #[test]
@@ -1091,9 +1134,9 @@ mod tests {
         )
         .unwrap();
 
-        let rendered = renderer.render(&request).unwrap().prompt;
+        let rendered = renderer.render(request.as_request()).unwrap().content;
 
-        assert_eq!(rendered, Prompt::Text("true|x".to_string()));
+        assert_eq!(rendered, RenderedPromptContent::Text("true|x".to_string()));
     }
 
     #[test]
@@ -1115,9 +1158,12 @@ mod tests {
         )
         .unwrap();
 
-        let rendered = renderer.render(&request).unwrap();
+        let rendered = renderer.render(request.as_request()).unwrap();
 
-        assert_eq!(rendered.prompt, Prompt::Text("max".to_string()));
+        assert_eq!(
+            rendered.content,
+            RenderedPromptContent::Text("max".to_string())
+        );
         assert_eq!(
             rendered.effective_template_kwargs.get("reasoning_effort"),
             Some(&Value::String("max".to_string()))
@@ -1144,9 +1190,12 @@ mod tests {
         )
         .unwrap();
 
-        let rendered = renderer.render(&request).unwrap();
+        let rendered = renderer.render(request.as_request()).unwrap();
 
-        assert_eq!(rendered.prompt, Prompt::Text("none|true".to_string()));
+        assert_eq!(
+            rendered.content,
+            RenderedPromptContent::Text("none|true".to_string())
+        );
         assert_eq!(
             rendered.effective_template_kwargs.get("reasoning_effort"),
             Some(&Value::String("none".to_string()))
@@ -1222,7 +1271,7 @@ mod tests {
     #[test]
     fn chat_template_exposes_tools_to_templates_when_auto_enabled() {
         let mut request = sample_request(vec![ChatMessage::text(ChatRole::User, "hello")]);
-        request.tools = vec![ChatTool {
+        request.tools = vec![Tool {
             name: "get_weather".to_string(),
             description: Some("Get weather".to_string()),
             parameters: serde_json::json!({
@@ -1326,7 +1375,7 @@ mod tests {
                 "<|im_start|>user\na<|vision_start|><|image_pad|><|vision_end|>b<|im_end|>\n",
             )
         "#]]
-        .assert_debug_eq(&rendered.prompt);
+        .assert_debug_eq(&rendered.content);
     }
 
     #[test]

rust/src/chat-renderer/src/hf/template.rs

@@ -5,7 +5,7 @@
 //!
 //! This module is inlined from SMG's tokenizer crate with local adaptations:
 //! - thinking-related detection/state is removed
-//! - special tokens are wired to `vllm_text::backends::hf::HfSpecialTokens`
+//! - special tokens are wired to `vllm_tokenizer::HfSpecialTokens`
 
 use std::collections::HashMap;
 use std::fs;
@@ -14,14 +14,14 @@ use std::path::Path;
 use minijinja::Environment;
 use serde::{Deserialize, Serialize};
 use serde_json::{self};
-use vllm_text::backend::hf::HfSpecialTokens;
+use vllm_tokenizer::HfSpecialTokens;
 
 use super::error::TemplateError;
 use super::format::{
     ChatTemplateContentFormat, ChatTemplateContentFormatOption, detect_chat_template_content_format,
 };
 use super::tojson::hf_tojson_filter;
-use crate::renderer::hf::{TemplateMessage, TemplateTool};
+use crate::hf::{TemplateMessage, TemplateTool};
 
 type Result<T> = std::result::Result<T, TemplateError>;
 
@@ -55,7 +55,7 @@ pub(super) struct TemplateContext<'a> {
 }
 
 /// Load chat template from a file (`.jinja` or `.json` containing Jinja).
-pub fn load_chat_template(template_path: &Path) -> Result<Option<String>> {
+pub(crate) fn load_chat_template(template_path: &Path) -> Result<Option<String>> {
     let content = fs::read_to_string(template_path).map_err(TemplateError::ReadTemplateFile)?;
 
     if template_path.extension().is_some_and(|ext| ext == "json") {
@@ -82,7 +82,7 @@ pub fn load_chat_template(template_path: &Path) -> Result<Option<String>> {
 }
 
 /// Resolve a configured chat template value into a template string.
-pub fn resolve_chat_template(chat_template: &str) -> Result<String> {
+pub(crate) fn resolve_chat_template(chat_template: &str) -> Result<String> {
     let path = Path::new(chat_template);
     if path.exists() {
         return load_chat_template(path).map(|template| template.unwrap_or_default());
@@ -136,7 +136,7 @@ mod tests {
     use std::fs;
 
     use tempfile::TempDir;
-    use vllm_text::backend::hf::{HfSpecialTokens, NamedSpecialToken};
+    use vllm_tokenizer::{HfSpecialTokens, NamedSpecialToken};
 
     use super::*;
 

rust/src/chat-renderer/src/inkling/mod.rs

@@ -5,13 +5,15 @@ use std::collections::HashMap;
 
 use serde_json::{Map, Value, json};
 use thiserror_ext::AsReport as _;
-use vllm_text::Prompt;
-use vllm_text::tokenizer::{DynTokenizer, Tokenizer};
+use vllm_tokenizer::{DynTokenizer, Tokenizer};
 
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{
+    ChatRenderer, RenderRequest, RenderedPrompt, RenderedPromptContent, request_template_kwargs,
+};
 use crate::error::{Error, Result};
-use crate::request::{ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatTool};
-use crate::{AssistantContentBlock, AssistantToolCall};
+use crate::{
+    AssistantContentBlock, AssistantToolCall, ChatContent, ChatContentPart, ChatMessage, Tool,
+};
 
 const MESSAGE_USER: &str = "<|message_user|>";
 const MESSAGE_MODEL: &str = "<|message_model|>";
@@ -170,7 +172,7 @@ impl InklingChatRenderer {
         )
     }
 
-    fn write_tool_declarations(&self, out: &mut Vec<u32>, tools: &[&ChatTool]) -> Result<()> {
+    fn write_tool_declarations(&self, out: &mut Vec<u32>, tools: &[&Tool]) -> Result<()> {
         if tools.is_empty() {
             return Ok(());
         }
@@ -277,14 +279,14 @@ impl InklingChatRenderer {
         tool_call_id: &str,
         tool_call_id_to_name: &HashMap<String, String>,
     ) -> Result<()> {
-        let text = content.try_flatten_to_text()?;
+        let text = content.try_flatten_to_text().map_err(Error::UnsupportedMultimodalContent)?;
         let tool_name = tool_call_id_to_name.get(tool_call_id).map(String::as_str).unwrap_or("");
         self.write_text_block(out, self.special.message_tool, Some(tool_name), &text)
     }
 }
 
 impl ChatRenderer for InklingChatRenderer {
-    fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
+    fn render(&self, request: RenderRequest<'_>) -> Result<RenderedPrompt> {
         request.validate()?;
         if request.chat_options.continue_final_message() {
             return Err(Error::ChatTemplate(
@@ -294,13 +296,13 @@ impl ChatRenderer for InklingChatRenderer {
 
         let mut out = Vec::new();
         let mut tool_call_id_to_name = HashMap::new();
-        let effective_template_kwargs = request_template_kwargs(request);
-        let tools = rendered_tools(request);
+        let effective_template_kwargs = request_template_kwargs(&request);
+        let tools = rendered_tools(&request);
         self.write_tool_declarations(&mut out, &tools)?;
         let mut reasoning_effort =
             resolve_reasoning_effort(effective_template_kwargs.get("reasoning_effort"));
 
-        for message in &request.messages {
+        for message in request.messages {
             if !matches!(
                 message,
                 ChatMessage::System { .. } | ChatMessage::Developer { .. }
@@ -345,7 +347,7 @@ impl ChatRenderer for InklingChatRenderer {
         }
 
         Ok(RenderedPrompt {
-            prompt: Prompt::TokenIds(out),
+            content: RenderedPromptContent::TokenIds(out),
             effective_template_kwargs,
         })
     }
@@ -378,14 +380,14 @@ fn resolve_special_token(tokenizer: &dyn Tokenizer, token: &str) -> Result<u32>
     })
 }
 
-fn rendered_tools(request: &ChatRequest) -> Vec<&ChatTool> {
+fn rendered_tools<'a>(request: &RenderRequest<'a>) -> Vec<&'a Tool> {
     if !request.tool_parsing_enabled() {
         return Vec::new();
     }
 
     let mut tools = Vec::with_capacity(request.tools.len());
     tools.extend(request.tools.iter());
-    for message in &request.messages {
+    for message in request.messages {
         if let ChatMessage::Developer {
             tools: Some(local_tools),
             ..

rust/src/chat-renderer/src/inkling/tests.rs

@@ -7,27 +7,23 @@ use std::sync::Arc;
 use expect_test::{ExpectFile, expect_file};
 use serde_json::json;
 use thiserror_ext::AsReport;
-use vllm_text::tokenizer::Tokenizer;
+use vllm_tokenizer::Tokenizer;
 
-use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
+use crate::test_utils::{FixtureRequestOptions, fixture_chat_request};
 
 use super::{
     AUDIO_END, CONTENT_AUDIO_INPUT, CONTENT_IMAGE, CONTENT_INVOKE_TOOL_JSON, CONTENT_TEXT,
     CONTENT_THINKING, CONTENT_XML, END_MESSAGE, InklingChatRenderer, MESSAGE_MODEL, MESSAGE_SYSTEM,
     MESSAGE_TOOL, MESSAGE_USER,
 };
-use crate::event::{AssistantContentBlock, AssistantToolCall};
-use crate::request::{ChatMessage, ChatRequest, GenerationPromptMode, ReasoningEffort};
+use crate::{AssistantContentBlock, AssistantToolCall};
+use crate::{ChatMessage, GenerationPromptMode, ReasoningEffort, TestRenderRequest};
 use crate::{ChatRenderer, Error};
 
 struct FixtureTokenizer;
 
 impl Tokenizer for FixtureTokenizer {
-    fn encode(
-        &self,
-        text: &str,
-        _add_special_tokens: bool,
-    ) -> vllm_text::tokenizer::Result<Vec<u32>> {
+    fn encode(&self, text: &str, _add_special_tokens: bool) -> vllm_tokenizer::Result<Vec<u32>> {
         Ok(text.bytes().map(u32::from).collect())
     }
 
@@ -35,7 +31,7 @@ impl Tokenizer for FixtureTokenizer {
         &self,
         token_ids: &[u32],
         _skip_special_tokens: bool,
-    ) -> vllm_text::tokenizer::Result<String> {
+    ) -> vllm_tokenizer::Result<String> {
         Ok(token_ids
             .iter()
             .map(|token_id| char::from_u32(*token_id).unwrap_or('\u{FFFD}'))
@@ -86,16 +82,16 @@ fn renderer() -> InklingChatRenderer {
     InklingChatRenderer::new(Arc::new(FixtureTokenizer)).unwrap()
 }
 
-fn render_token_ids(request: &ChatRequest) -> Vec<u32> {
+fn render_token_ids(request: &TestRenderRequest) -> Vec<u32> {
     renderer()
-        .render(request)
+        .render(request.as_request())
         .unwrap()
-        .prompt
+        .content
         .into_token_ids()
         .expect("Inkling renderer returns token IDs")
 }
 
-fn fixture_request(name: &str) -> ChatRequest {
+fn fixture_request(name: &str) -> TestRenderRequest {
     fixture_chat_request(&fixture_path(name), inkling_fixture_options())
 }
 
@@ -108,7 +104,7 @@ fn inkling_fixture_options() -> FixtureRequestOptions {
 
 fn fixture_path(name: &str) -> PathBuf {
     PathBuf::from(env!("CARGO_MANIFEST_DIR"))
-        .join("src/renderer/inkling")
+        .join("src/inkling")
         .join("fixtures")
         .join(name)
 }
@@ -229,14 +225,14 @@ fn renders_named_reasoning_effort_after_tool_declarations() {
 
 #[test]
 fn emits_one_reasoning_effort_for_multi_turn_conversation() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![
             ChatMessage::system("rules"),
             ChatMessage::user("user1"),
             ChatMessage::assistant_text("assistant1"),
             ChatMessage::user("user2"),
         ],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
     let rendered = render_symbolic_tokens(&render_token_ids(&request));
@@ -253,7 +249,7 @@ fn emits_one_reasoning_effort_for_multi_turn_conversation() {
 #[test]
 fn canonicalizes_numeric_zero_reasoning_effort() {
     for value in [0.0, -0.0] {
-        let mut request = ChatRequest::for_test();
+        let mut request = TestRenderRequest::for_test();
         request
             .chat_options
             .template_kwargs
@@ -270,7 +266,7 @@ fn canonicalizes_numeric_zero_reasoning_effort() {
 
 #[test]
 fn defaults_reasoning_effort_to_high() {
-    let request = ChatRequest::for_test();
+    let request = TestRenderRequest::for_test();
 
     assert!(
         render_symbolic_tokens(&render_token_ids(&request)).starts_with(
@@ -282,7 +278,7 @@ fn defaults_reasoning_effort_to_high() {
 
 #[test]
 fn renders_numeric_reasoning_effort_template_kwarg() {
-    let mut request = ChatRequest::for_test();
+    let mut request = TestRenderRequest::for_test();
     request
         .chat_options
         .template_kwargs
@@ -299,7 +295,7 @@ fn renders_numeric_reasoning_effort_template_kwarg() {
 #[test]
 fn ignores_unsupported_reasoning_effort_values() {
     for value in [json!(true), json!("invalid"), json!(null)] {
-        let mut request = ChatRequest::for_test();
+        let mut request = TestRenderRequest::for_test();
         request
             .chat_options
             .template_kwargs
@@ -315,38 +311,38 @@ fn ignores_unsupported_reasoning_effort_values() {
 #[test]
 fn rejects_out_of_range_reasoning_effort() {
     for value in [0.990_000_1, 1.0, 1.5, -0.1] {
-        let mut request = ChatRequest::for_test();
+        let mut request = TestRenderRequest::for_test();
         request
             .chat_options
             .template_kwargs
             .insert("reasoning_effort".to_string(), json!(value));
 
-        let error = renderer().render(&request).unwrap_err();
+        let error = renderer().render(request.as_request()).unwrap_err();
         assert!(error.as_report().to_string().contains("must be in [0.0, 0.99]"));
     }
 }
 
 #[test]
 fn rejects_continue_final_message() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![ChatMessage::assistant_text("partial")],
         chat_options: crate::ChatOptions {
             generation_prompt_mode: GenerationPromptMode::ContinueFinalAssistant,
             ..Default::default()
         },
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
-    let error = renderer().render(&request).unwrap_err();
+    let error = renderer().render(request.as_request()).unwrap_err();
     assert!(matches!(error, Error::ChatTemplate(_)));
     assert!(error.as_report().to_string().contains("continue_final_message"));
 }
 
 #[test]
 fn renders_developer_messages_as_system() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![ChatMessage::developer("rules", None)],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
     assert_eq!(
@@ -359,7 +355,7 @@ fn renders_developer_messages_as_system() {
 
 #[test]
 fn rejects_non_object_tool_call_arguments() {
-    let request = ChatRequest {
+    let request = TestRenderRequest {
         messages: vec![ChatMessage::assistant_blocks(vec![
             AssistantContentBlock::ToolCall(AssistantToolCall {
                 id: "call_1".to_string(),
@@ -367,9 +363,9 @@ fn rejects_non_object_tool_call_arguments() {
                 arguments: "[]".to_string(),
             }),
         ])],
-        ..ChatRequest::for_test()
+        ..TestRenderRequest::for_test()
     };
 
-    let error = renderer().render(&request).unwrap_err();
+    let error = renderer().render(request.as_request()).unwrap_err();
     assert!(error.as_report().to_string().contains("JSON object"));
 }

rust/src/chat-renderer/src/lib.rs

@@ -1,45 +1,70 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+//! Engine-independent chat prompt renderers.
+//!
+//! The crate accepts borrowed chat-domain inputs and produces text or token-ID
+//! prompt artifacts. Serving requests, sampling policy, multimodal
+//! preprocessing, parser state, and engine transport remain in consumer
+//! crates.
+
 use std::collections::HashMap;
 use std::sync::Arc;
 
+use enum_as_inner::EnumAsInner;
+use serde::{Deserialize, Serialize};
 use serde_json::{Value, json};
-use vllm_text::Prompt;
 
-use crate::error::Result;
-use crate::request::{ChatRequest, ReasoningEffort};
-
-pub mod deepseek_v32;
-pub mod deepseek_v4;
+mod deepseek_v32;
+mod deepseek_v4;
+mod error;
 pub mod harmony;
 pub mod hf;
 mod inkling;
+mod request;
 mod selection;
 #[cfg(test)]
 mod test_utils;
+#[cfg(test)]
+pub(crate) use test_utils::TestRenderRequest;
 
 pub use deepseek_v4::DeepSeekV4ChatRenderer;
 pub use deepseek_v32::DeepSeekV32ChatRenderer;
+pub use error::{Error, Result};
 pub use harmony::HarmonyChatRenderer;
 pub use inkling::InklingChatRenderer;
+pub use request::RenderRequest;
 pub use selection::RendererSelection;
+pub use vllm_chat_types::{
+    AssistantBlockKind, AssistantContentBlock, AssistantMessage, AssistantMessageExt,
+    AssistantToolCall, ChatContent, ChatContentPart, ChatMessage, ChatOptions, ChatRole,
+    ChatToolChoice, GenerationPromptMode, ImageDetail, ReasoningEffort, Tool,
+};
 
-/// Rendered chat prompt submitted to the text backend.
-#[derive(Debug, Clone, PartialEq)]
+/// Engine-independent prompt content produced by a chat renderer.
+#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize, EnumAsInner)]
+#[serde(untagged)]
+pub enum RenderedPromptContent {
+    /// Untokenized prompt text.
+    Text(String),
+    /// Prompt token IDs produced by a format-specific renderer.
+    TokenIds(Vec<u32>),
+}
+
+/// Rendered chat prompt plus effective template metadata.
+#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
 pub struct RenderedPrompt {
-    /// The rendered prompt, either as text or already tokenized.
-    pub prompt: Prompt,
+    /// Engine-independent prompt content.
+    pub content: RenderedPromptContent,
     /// Effective chat-template kwargs visible to the renderer after applying
     /// server defaults, request overrides, and typed reasoning controls.
     pub effective_template_kwargs: HashMap<String, Value>,
 }
 
-/// Minimal chat-prompt renderer used by `vllm-chat`.
+/// Synchronous chat-prompt renderer.
 pub trait ChatRenderer: Send + Sync {
-    /// Render one chat request into the text prompt submitted to the text
-    /// backend.
-    fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt>;
+    /// Render one borrowed chat request into text or token IDs.
+    fn render(&self, request: RenderRequest<'_>) -> Result<RenderedPrompt>;
 }
 
 /// Shared trait-object form of [`ChatRenderer`].
@@ -49,7 +74,7 @@ pub type DynChatRenderer = Arc<dyn ChatRenderer>;
 /// using the provided defaults as the base.
 pub(crate) fn effective_template_kwargs(
     default_template_kwargs: &HashMap<String, Value>,
-    request: &ChatRequest,
+    request: &RenderRequest<'_>,
 ) -> HashMap<String, Value> {
     let mut kwargs = default_template_kwargs.clone();
     kwargs.extend(request.chat_options.template_kwargs.clone());
@@ -71,6 +96,6 @@ pub(crate) fn effective_template_kwargs(
 }
 
 /// Extract the effective chat-template kwargs visible to the renderer from the request.
-pub(crate) fn request_template_kwargs(request: &ChatRequest) -> HashMap<String, Value> {
+pub(crate) fn request_template_kwargs(request: &RenderRequest<'_>) -> HashMap<String, Value> {
     effective_template_kwargs(&HashMap::new(), request)
 }

rust/src/chat-renderer/src/request.rs

@@ -0,0 +1,172 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use serde_json::Value;
+use vllm_chat_types::{
+    ChatMessage, ChatOptions, ChatRole, ChatToolChoice, GenerationPromptMode, Tool,
+};
+
+use crate::{Error, Result};
+
+/// Borrowed input consumed by one chat renderer.
+///
+/// This view contains only chat-domain values that affect prompt construction.
+/// Serving, sampling, decoding, scheduling, and engine metadata stay with the
+/// caller.
+#[derive(Debug, Clone, Copy)]
+pub struct RenderRequest<'a> {
+    /// Ordered chat history to render.
+    pub messages: &'a [ChatMessage],
+    /// Chat-template and generation-prompt controls.
+    pub chat_options: &'a ChatOptions,
+    /// Request-level tools available to the model.
+    pub tools: &'a [Tool],
+    /// Tool-choice behavior used to decide whether tools are exposed.
+    pub tool_choice: &'a ChatToolChoice,
+    /// Optional retrieval documents exposed to HF chat templates.
+    pub documents: Option<&'a [Value]>,
+}
+
+impl RenderRequest<'_> {
+    /// Validate renderer-owned request invariants.
+    pub fn validate(&self) -> Result<()> {
+        if self.messages.is_empty() {
+            return Err(Error::EmptyMessages);
+        }
+        match (
+            self.chat_options.generation_prompt_mode,
+            self.messages.last().map(ChatMessage::role),
+        ) {
+            (GenerationPromptMode::ContinueFinalAssistant, Some(ChatRole::Assistant)) => {}
+            (GenerationPromptMode::ContinueFinalAssistant, _) => {
+                return Err(Error::ContinueFinalAssistantWithoutFinalAssistant);
+            }
+            (GenerationPromptMode::NoGenerationPrompt, _)
+            | (GenerationPromptMode::StartNewAssistant, _) => {}
+        }
+        Ok(())
+    }
+
+    /// Return whether any message contains multimodal content.
+    pub fn has_multimodal(&self) -> bool {
+        self.messages.iter().any(ChatMessage::has_multimodal)
+    }
+
+    /// Return whether request-level tools should be exposed to the renderer.
+    pub fn tool_parsing_enabled(&self) -> bool {
+        !matches!(self.tool_choice, ChatToolChoice::None) && !self.tools.is_empty()
+    }
+
+    /// Return the request-level thinking toggle when explicitly requested.
+    ///
+    /// The `thinking` and `enable_thinking` kwargs must be booleans when
+    /// present and must carry the same value when both are set.
+    pub fn enable_thinking(&self) -> Result<Option<bool>> {
+        let thinking = self.parse_template_bool("thinking")?;
+        let enable_thinking = self.parse_template_bool("enable_thinking")?;
+
+        match (thinking, enable_thinking) {
+            (None, None) => Ok(None),
+            (Some(thinking), Some(enable_thinking)) if thinking != enable_thinking => {
+                Err(Error::ChatTemplate(
+                    "template kwargs `thinking` and `enable_thinking` must match when both are set"
+                        .to_string(),
+                ))
+            }
+            (Some(thinking), _) => Ok(Some(thinking)),
+            (None, Some(enable_thinking)) => Ok(Some(enable_thinking)),
+        }
+    }
+
+    /// Parse one optional boolean chat-template kwarg.
+    pub fn parse_template_bool(&self, key: &str) -> Result<Option<bool>> {
+        match self.chat_options.template_kwargs.get(key) {
+            None => Ok(None),
+            Some(Value::Bool(value)) => Ok(Some(*value)),
+            Some(other) => Err(Error::ChatTemplate(format!(
+                "template kwarg `{key}` must be a boolean, got {other}"
+            ))),
+        }
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use serde_json::json;
+
+    use super::RenderRequest;
+    use crate::{ChatMessage, ChatToolChoice, Error, GenerationPromptMode, TestRenderRequest};
+
+    #[test]
+    fn rejects_empty_message_history() {
+        let mut request = TestRenderRequest::for_test();
+        request.messages.clear();
+
+        assert!(matches!(
+            request.as_request().validate(),
+            Err(Error::EmptyMessages)
+        ));
+    }
+
+    #[test]
+    fn continuation_requires_a_final_assistant_message() {
+        let mut request = TestRenderRequest::for_test();
+        request.chat_options.generation_prompt_mode = GenerationPromptMode::ContinueFinalAssistant;
+
+        assert!(matches!(
+            request.as_request().validate(),
+            Err(Error::ContinueFinalAssistantWithoutFinalAssistant)
+        ));
+
+        request.messages.push(ChatMessage::assistant_text("partial"));
+        request.as_request().validate().unwrap();
+    }
+
+    #[test]
+    fn thinking_aliases_must_match() {
+        let mut request = TestRenderRequest::for_test();
+        request
+            .chat_options
+            .template_kwargs
+            .insert("thinking".to_string(), json!(false));
+        request
+            .chat_options
+            .template_kwargs
+            .insert("enable_thinking".to_string(), json!(true));
+
+        assert!(matches!(
+            request.as_request().enable_thinking(),
+            Err(Error::ChatTemplate(message))
+                if message.contains("`thinking` and `enable_thinking` must match")
+        ));
+    }
+
+    #[test]
+    fn tool_exposure_requires_tools_and_an_enabled_choice() {
+        let mut request = TestRenderRequest::for_test();
+        assert!(!request.as_request().tool_parsing_enabled());
+
+        request.tools.push(crate::Tool {
+            name: "lookup".to_string(),
+            description: None,
+            parameters: json!({"type": "object"}),
+            strict: None,
+        });
+        assert!(!request.as_request().tool_parsing_enabled());
+
+        request.tool_choice = ChatToolChoice::Auto;
+        assert!(request.as_request().tool_parsing_enabled());
+    }
+
+    #[test]
+    fn borrowed_request_carries_optional_documents() {
+        let mut request = TestRenderRequest::for_test();
+        request.documents = Some(vec![json!({"title": "doc"})]);
+
+        let borrowed: RenderRequest<'_> = request.as_request();
+        assert_eq!(
+            borrowed.documents.unwrap(),
+            request.documents.as_deref().unwrap()
+        );
+    }
+}

rust/src/chat-renderer/src/selection.rs

@@ -29,13 +29,21 @@ pub enum RendererSelection {
 }
 
 impl RendererSelection {
+    /// CLI/config literal for automatic renderer selection.
     pub const AUTO_LITERAL: &str = "auto";
+    /// CLI/config literal and model type for DeepSeek V3.2.
     pub const DEEPSEEK_V32_LITERAL: &str = "deepseek_v32";
+    /// CLI/config literal and model type for DeepSeek V4.
     pub const DEEPSEEK_V4_LITERAL: &str = "deepseek_v4";
+    /// Hugging Face model type routed to the Harmony renderer.
     pub const GPT_OSS_MODEL_TYPE: &str = "gpt_oss";
+    /// CLI/config literal for the Harmony renderer.
     pub const HARMONY_LITERAL: &str = "harmony";
+    /// CLI/config literal for the Hugging Face renderer.
     pub const HF_LITERAL: &str = "hf";
+    /// CLI/config literal for the Inkling renderer.
     pub const INKLING_LITERAL: &str = "inkling";
+    /// Hugging Face model type routed to the Inkling renderer.
     pub const INKLING_MODEL_TYPE: &str = "inkling_mm_model";
 
     /// Resolve the renderer selection using the given model type string, if

rust/src/chat-renderer/src/test_utils.rs

@@ -7,13 +7,45 @@ use std::path::Path;
 use serde::Deserialize;
 use serde_json::Value;
 
-use crate::event::{AssistantContentBlock, AssistantToolCall};
-use crate::request::{
-    ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatTool, ChatToolChoice,
-    GenerationPromptMode, ReasoningEffort,
+use crate::{
+    AssistantContentBlock, AssistantToolCall, ChatContent, ChatContentPart, ChatMessage,
+    ChatOptions, ChatRole, ChatToolChoice, GenerationPromptMode, ReasoningEffort, RenderRequest,
+    Tool,
 };
 
-/// Options for constructing a [`ChatRequest`] from a fixture file.
+/// Owned test input used to construct borrowed [`RenderRequest`] values.
+#[derive(Debug, Clone)]
+pub(crate) struct TestRenderRequest {
+    pub messages: Vec<ChatMessage>,
+    pub chat_options: ChatOptions,
+    pub tools: Vec<Tool>,
+    pub tool_choice: ChatToolChoice,
+    pub documents: Option<Vec<Value>>,
+}
+
+impl TestRenderRequest {
+    pub(crate) fn for_test() -> Self {
+        Self {
+            messages: vec![ChatMessage::text(ChatRole::User, "test")],
+            chat_options: ChatOptions::default(),
+            tools: Vec::new(),
+            tool_choice: ChatToolChoice::None,
+            documents: None,
+        }
+    }
+
+    pub(crate) fn as_request(&self) -> RenderRequest<'_> {
+        RenderRequest {
+            messages: &self.messages,
+            chat_options: &self.chat_options,
+            tools: &self.tools,
+            tool_choice: &self.tool_choice,
+            documents: self.documents.as_deref(),
+        }
+    }
+}
+
+/// Options for constructing a [`TestRenderRequest`] from a fixture file.
 #[derive(Debug, Clone, Copy)]
 pub(crate) struct FixtureRequestOptions {
     /// Whether to set the template kwarg `[enable_]thinking=true`.
@@ -23,9 +55,11 @@ pub(crate) struct FixtureRequestOptions {
     pub no_generation_prompt_when_last_assistant: bool,
 }
 
-/// Read a fixture file from the given path and convert it into a [`ChatRequest`]
-/// using the provided options.
-pub(crate) fn fixture_chat_request(path: &Path, options: FixtureRequestOptions) -> ChatRequest {
+/// Read a fixture file and convert it into an owned renderer test request.
+pub(crate) fn fixture_chat_request(
+    path: &Path,
+    options: FixtureRequestOptions,
+) -> TestRenderRequest {
     let fixture = fs::read_to_string(path).unwrap();
     let fixture: FixtureFile = serde_json::from_str(&fixture).unwrap();
     fixture.into_request().into_chat_request(options)
@@ -135,9 +169,8 @@ struct FixtureToolCallFunction {
 }
 
 impl FixtureRequest {
-    fn into_chat_request(self, options: FixtureRequestOptions) -> ChatRequest {
-        let mut request = ChatRequest {
-            request_id: "renderer-fixture".to_string(),
+    fn into_chat_request(self, options: FixtureRequestOptions) -> TestRenderRequest {
+        let mut request = TestRenderRequest {
             messages: self
                 .messages
                 .into_iter()
@@ -150,7 +183,7 @@ impl FixtureRequest {
             } else {
                 ChatToolChoice::Auto
             },
-            ..ChatRequest::for_test()
+            ..TestRenderRequest::for_test()
         };
 
         if options.no_generation_prompt_when_last_assistant
@@ -243,10 +276,10 @@ fn to_chat_content(content: FixtureContent) -> ChatContent {
     }
 }
 
-fn to_chat_tools(tools: &[FixtureTool]) -> Vec<ChatTool> {
+fn to_chat_tools(tools: &[FixtureTool]) -> Vec<Tool> {
     tools
         .iter()
-        .map(|tool| ChatTool {
+        .map(|tool| Tool {
             name: tool.function.name.clone(),
             description: tool.function.description.clone(),
             parameters: tool.function.parameters.clone(),

rust/src/chat-renderer/tests/templates/vllm_examples/README.md

@@ -0,0 +1,6 @@
+# vLLM Example Templates
+
+These fixtures are copied from `vllm/examples/`.
+
+They are used by `src/chat-renderer/src/hf/format.rs` tests to keep
+chat-template content-format detection aligned with Python vLLM behavior.