Merge remote-tracking branch 'origin/main' into sync/lesj/qwen4exp-bugfixes

Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>

.buildkite/scripts/crcr-report.sh

@@ -40,8 +40,7 @@ fi
 TOKEN_SECRET_KEY="${CRCR_BUILDKITE_TOKEN_SECRET_KEY:-CRCR_BUILDKITE_API_TOKEN}"
 BK_TOKEN="${BUILDKITE_API_TOKEN:-}"
 if [[ -z "${BK_TOKEN}" ]]; then
-    # Not in the job environment, so read it from a Buildkite secret. The agent
-    # redacts values fetched this way from the log.
+    # Not in the job environment, so read it from a Buildkite secret.
     #
     # Report why a lookup failed. Swallowing stderr made a missing secret, a
     # denied policy and an unusable agent indistinguishable, all surfacing as the
@@ -50,14 +49,9 @@ if [[ -z "${BK_TOKEN}" ]]; then
         echo "buildkite-agent is not on PATH; cannot read secret '${TOKEN_SECRET_KEY}'"
     else
         secret_err="$(mktemp)"
-        # Deliberately not passing --skip-redaction. On the deployed agent
-        # (v3.73.1) the flag cannot help: secret_get.go creates the Job API
-        # client unconditionally and only checks SkipRedaction afterwards, so
-        # under the docker plugin -- which does not expose the Job API socket to
-        # the container -- it fails before the flag is read. That ordering was
-        # only fixed in v3.107.0. Skipping redaction would also stop the token
-        # being registered with the log redactor, for no gain here.
-        if BK_TOKEN="$(buildkite-agent secret get "${TOKEN_SECRET_KEY}" 2>"${secret_err}")"; then
+        # Requires agent >= 3.107.0: the Docker plugin does not mount the Job API
+        # socket needed for redaction. Capture the value without logging it.
+        if BK_TOKEN="$(buildkite-agent secret get --skip-redaction "${TOKEN_SECRET_KEY}" 2>"${secret_err}")"; then
             if [[ -z "${BK_TOKEN}" ]]; then
                 echo "secret '${TOKEN_SECRET_KEY}' resolved but is empty"
             fi
@@ -77,7 +71,8 @@ if [[ -z "${BK_TOKEN}" ]]; then
 fi
 
 AUDIENCE="pytorch-cross-repo-ci-relay"
-OIDC_TOKEN="$(buildkite-agent oidc request-token --audience "${AUDIENCE}" 2>/dev/null)"
+# OIDC redaction also requires the unavailable Job API socket.
+OIDC_TOKEN="$(buildkite-agent oidc request-token --skip-redaction --audience "${AUDIENCE}" 2>/dev/null)"
 if [[ -z "${OIDC_TOKEN}" ]]; then
     echo "could not mint a Buildkite OIDC token -- skipping report"
     exit 0

.buildkite/test-amd.yaml

@@ -1704,8 +1704,9 @@ steps:
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
-  - pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --ignore=kernels/moe/test_modular_kernel_combinations.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
   - pytest -v -s kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_MOE=1 pytest -v -s kernels/moe/test_modular_kernel_combinations.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
 
 - label: ":amd: (MI300) Quantization Kernels" # TBD
   timeout_in_minutes: 55
@@ -2175,19 +2176,20 @@ steps:
   - pytest -v -s models/multimodal/generation/test_memory_leak.py -m core_model
   - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model
 
-- label: ":amd: (MI300) Multimodal Processor" # 1h 42m
+- label: ":amd: (MI300) Multimodal Processor Shard %N" # TBD
   timeout_in_minutes: 110
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   dind: false
   agent_pool: mi300_1
   optional: true
+  parallelism: 4
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   - tests/models/registry.py
   commands:
-  - pytest -v -s models/multimodal/processing/test_tensor_schema.py
+  - pytest -v -s models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
 - label: ":amd: (MI300) Multimodal Models (Extended Pooling)" # TBD
   timeout_in_minutes: 75

docker/Dockerfile.rock

@@ -5,7 +5,7 @@ ARG VLLM_BRANCH="main"
 ARG COMMON_WORKDIR=/app
 ARG BASE_IMAGE=rocm/vllm-dev:base
 ARG CI_BASE_IMAGE=rocm/vllm-dev:ci_base
-ARG ROCM_TRITON_KERNELS_COMMIT=0f380657dbf3ee86eb57558ff71df24f03b5d4e7
+ARG ROCM_TRITON_KERNELS_COMMIT=669b31acc1dd1b3fd93286afd5db67f65d9f7557
 ARG INSTALL_LMCACHE=false
 # NIC backend for MoRI RDMA support.
 # By default (all), drivers and userspace libraries for all supported NIC types
@@ -322,8 +322,9 @@ ENV SETUPTOOLS_SCM_PRETEND_VERSION="0.0.0+rocm.csrc.build"
 
 RUN --mount=type=bind,source=pyproject.toml,target=${COMMON_WORKDIR}/vllm/pyproject.toml \
     --mount=type=cache,id=vllm-rocm-ccache,target=/root/.cache/ccache \
-    grep -Fq "set(TRITON_KERNELS_TAG \"${ROCM_TRITON_KERNELS_COMMIT}\")" \
+    { grep -Fq "set(TRITON_KERNELS_TAG \"${ROCM_TRITON_KERNELS_COMMIT}\")" \
         cmake/external_projects/triton_kernels.cmake \
+        || echo "WARNING: Dockerfile.rock ROCM_TRITON_KERNELS_COMMIT (${ROCM_TRITON_KERNELS_COMMIT}) does not match cmake TRITON_KERNELS_TAG; the Docker build uses the Dockerfile pin, cmake governs only non-Docker builds"; } \
     && export CCACHE_BASEDIR="$PWD" \
     && echo "=== ccache stats before ROCm native build ===" \
     && (ccache --show-stats || true) \
@@ -669,8 +670,9 @@ RUN if [ "$GIT_REPO_CHECK" != "0" ]; then \
 
 # Install the packaging helpers before modifying requirements/rocm.txt.
 RUN --mount=type=cache,target=/root/.cache/uv \
-    grep -Fq "set(TRITON_KERNELS_TAG \"${ROCM_TRITON_KERNELS_COMMIT}\")" \
+    { grep -Fq "set(TRITON_KERNELS_TAG \"${ROCM_TRITON_KERNELS_COMMIT}\")" \
         vllm/cmake/external_projects/triton_kernels.cmake \
+        || echo "WARNING: Dockerfile.rock ROCM_TRITON_KERNELS_COMMIT (${ROCM_TRITON_KERNELS_COMMIT}) does not match cmake TRITON_KERNELS_TAG; the Docker build uses the Dockerfile pin, cmake governs only non-Docker builds"; } \
     && cd vllm \
     && uv pip install --system setuptools_scm regex
 

rust/src/chat/src/backend/hf.rs

@@ -70,12 +70,26 @@ impl HfChatBackend {
                 options,
                 multimodal_render_info,
             )?),
-            RendererSelection::DeepSeekV32 => Arc::new(DeepSeekV32ChatRenderer::new()),
-            RendererSelection::DeepSeekV4 => Arc::new(DeepSeekV4ChatRenderer::new()),
-            RendererSelection::DeepSeekV41 => Arc::new(DeepSeekV41ChatRenderer::new()),
-            RendererSelection::Harmony => Arc::new(HarmonyChatRenderer::new()?),
-            RendererSelection::Inkling => Arc::new(InklingChatRenderer::new(tokenizer.clone())?),
-            RendererSelection::KimiK3 => Arc::new(KimiK3ChatRenderer::new(tokenizer.clone())),
+            RendererSelection::DeepSeekV32 => Arc::new(DeepSeekV32ChatRenderer::new(
+                options.default_chat_template_kwargs,
+            )),
+            RendererSelection::DeepSeekV4 => Arc::new(DeepSeekV4ChatRenderer::new(
+                options.default_chat_template_kwargs,
+            )),
+            RendererSelection::DeepSeekV41 => Arc::new(DeepSeekV41ChatRenderer::new(
+                options.default_chat_template_kwargs,
+            )),
+            RendererSelection::Harmony => Arc::new(HarmonyChatRenderer::new(
+                options.default_chat_template_kwargs,
+            )?),
+            RendererSelection::Inkling => Arc::new(InklingChatRenderer::new(
+                tokenizer.clone(),
+                options.default_chat_template_kwargs,
+            )?),
+            RendererSelection::KimiK3 => Arc::new(KimiK3ChatRenderer::new(
+                tokenizer.clone(),
+                options.default_chat_template_kwargs,
+            )),
         };
 
         info!(
@@ -280,6 +294,39 @@ mod tests {
         assert_eq!(prompt, "hello");
     }
 
+    #[test]
+    fn native_renderer_inherits_deployment_reasoning_below_request_controls() {
+        let backend = HfChatBackend::from_resolved_model_files(
+            resolved_files(r#"{"model_type":"deepseek_v4"}"#, "{}"),
+            "test-model".to_string(),
+            LoadModelBackendsOptions {
+                default_chat_template_kwargs: [
+                    ("thinking".to_string(), serde_json::json!(false)),
+                    ("reasoning_effort".to_string(), serde_json::json!("low")),
+                ]
+                .into(),
+                ..Default::default()
+            },
+            test_tokenizer(),
+        )
+        .unwrap();
+        let mut request = request_with_user_text("hello");
+        let rendered = backend.chat_renderer().render(&request).unwrap();
+        assert!(rendered.prompt.into_text().unwrap().ends_with("</think>"));
+        assert_eq!(rendered.effective_template_kwargs["enable_thinking"], false);
+
+        request.chat_options.reasoning_effort = Some(crate::EffortValue::from("max"));
+        let rendered = backend.chat_renderer().render(&request).unwrap();
+        let prompt = rendered.prompt.into_text().unwrap();
+        assert!(prompt.contains("Reasoning Effort: Beyond maximum"));
+        assert!(prompt.ends_with("<think>"));
+        assert_eq!(
+            rendered.effective_template_kwargs["reasoning_effort"],
+            "max"
+        );
+        assert_eq!(rendered.effective_template_kwargs["thinking"], true);
+    }
+
     #[test]
     fn auto_uses_harmony_renderer_and_output_processor_for_gpt_oss_model_type() {
         let backend = backend_for_selection(

rust/src/chat/src/backend/mod.rs

@@ -78,8 +78,8 @@ pub struct LoadModelBackendsOptions {
     /// Optional server-default chat template override, provided either as an
     /// inline template or as a path to a template file.
     pub chat_template: Option<String>,
-    /// Optional server-default keyword arguments merged into every
-    /// chat-template render before request-level `chat_template_kwargs`.
+    /// Server-default keyword arguments. HF merges these before request kwargs;
+    /// native renderers inherit supported reasoning controls below request controls.
     pub default_chat_template_kwargs: HashMap<String, Value>,
     /// Maximum number of input items allowed per prompt for each modality.
     /// Unspecified modalities are unlimited.

rust/src/chat/src/error.rs

@@ -19,6 +19,8 @@ pub enum Error {
     ChatTemplate(String),
     #[error("{0}")]
     InvalidReasoningEffort(String),
+    #[error("{message}")]
+    InvalidReasoningControl { message: String },
     #[error("multimodal input is not supported by this chat renderer")]
     UnsupportedMultimodalRenderer,
     #[error("unsupported multimodal content: {0}")]
@@ -101,6 +103,7 @@ impl Error {
         match self {
             Self::PromptTooLong { .. }
             | Self::InvalidReasoningEffort(_)
+            | Self::InvalidReasoningControl { .. }
             | Self::DuplicateToolName { .. }
             | Self::ToolChoiceRequiresTools
             | Self::ToolChoiceFunctionNotFound { .. } => true,

rust/src/chat/src/lib.rs

@@ -31,6 +31,7 @@ pub use parser::reasoning::{
 };
 pub use parser::tool::{ToolParser, ToolParserError, ToolParserFactory};
 pub use parser::{ParserSelection, validate_parser_overrides};
+pub use reasoning::EffortValue;
 pub use renderer::hf::ChatTemplateContentFormatOption;
 pub use renderer::{
     ChatRenderer, DeepSeekV4ChatRenderer, DeepSeekV32ChatRenderer, DeepSeekV41ChatRenderer,
@@ -39,7 +40,7 @@ pub use renderer::{
 };
 pub use request::{
     ChatContent, ChatContentPart, ChatMessage, ChatOptions, ChatRequest, ChatRole, ChatTool,
-    ChatToolChoice, GenerationPromptMode, ReasoningEffort, ResolvedToolContext, SamplingParams,
+    ChatToolChoice, GenerationPromptMode, ResolvedToolContext, SamplingParams,
 };
 pub use stream::{ChatEventStream, ChatEventStreamTrait, CollectedAssistantMessage};
 pub use vllm_engine_core_client::protocol::multimodal::MmFeatures;
@@ -52,6 +53,7 @@ mod event;
 pub mod multimodal;
 mod output;
 mod parser;
+mod reasoning;
 mod renderer;
 mod request;
 mod stream;

rust/src/chat/src/reasoning.rs

@@ -0,0 +1,397 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+//! Shared reasoning-control resolution for native renderers.
+//!
+//! Each input source is parsed independently into [`ReasoningControl`]. Sources
+//! then fall back in priority order: request, deployment defaults, model default.
+//! Within a source, an explicit toggle decides the mode and effort supplies its
+//! intensity. Keeping sources separate lets request effort override a deployment
+//! toggle, and lets an enabled request inherit effort from enabled defaults.
+//!
+//! This module owns input types, precedence, and the standard `"none"` sentinel.
+//! Model adapters own supported effort names, numeric ranges, and prompt encoding.
+//! After model lowering, [`ReasoningControl::template_kwargs`] projects the prompt's
+//! decision for downstream consumers. HF prepares raw template inputs separately.
+
+use std::collections::HashMap;
+use std::fmt;
+
+use serde::{Deserialize, Serialize};
+use serde_json::{Number, Value, json};
+
+use crate::error::{Result, invalid_reasoning_control};
+use crate::request::ChatRequest;
+
+/// A model-specific reasoning effort, before or after renderer lowering.
+///
+/// Strings preserve model-specific names; JSON numbers preserve integer and
+/// floating-point representations for model-local range validation. Missing/null
+/// effort is represented by `Option::None` in [`crate::ChatOptions`].
+/// Supported names and ranges are validated by each renderer or HF template.
+#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
+#[serde(untagged)]
+pub enum EffortValue {
+    String(String),
+    Number(Number),
+}
+
+impl TryFrom<Value> for EffortValue {
+    type Error = crate::error::Error;
+
+    fn try_from(value: Value) -> Result<Self> {
+        match value {
+            Value::String(value) => Ok(Self::String(value)),
+            Value::Number(value) => Ok(Self::Number(value)),
+            _ => Err(invalid_reasoning_control!(
+                "reasoning effort must be a string or number, got {value}"
+            )),
+        }
+    }
+}
+
+impl TryFrom<f64> for EffortValue {
+    type Error = crate::error::Error;
+
+    /// Construct a numeric effort, rejecting NaN and infinities.
+    fn try_from(value: f64) -> Result<Self> {
+        Number::from_f64(value).map(Self::Number).ok_or_else(|| {
+            invalid_reasoning_control!("reasoning effort must be finite, got {value}")
+        })
+    }
+}
+
+impl From<&str> for EffortValue {
+    fn from(value: &str) -> Self {
+        Self::String(value.to_owned())
+    }
+}
+
+impl fmt::Display for EffortValue {
+    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
+        match self {
+            Self::String(value) => json!(value).fmt(f),
+            Self::Number(value) => value.fmt(f),
+        }
+    }
+}
+
+impl EffortValue {
+    pub fn as_str(&self) -> Option<&str> {
+        match self {
+            Self::String(value) => Some(value),
+            Self::Number(_) => None,
+        }
+    }
+
+    pub fn as_f64(&self) -> Option<f64> {
+        match self {
+            Self::Number(value) => value.as_f64(),
+            Self::String(_) => None,
+        }
+    }
+}
+
+/// Reasoning intent for one source, or the result of falling back across sources.
+#[derive(Debug, Clone, PartialEq)]
+pub(crate) enum ReasoningControl {
+    /// This source leaves the mode and effort to a lower-priority source.
+    Default,
+    /// Reasoning is explicitly disabled; lower-priority defaults preserve this mode.
+    Disabled,
+    /// Reasoning is enabled. Missing effort inherits from an enabled default.
+    Enabled { effort: Option<EffortValue> },
+}
+
+impl ReasoningControl {
+    /// Parse request controls; the typed effort field takes precedence over kwargs effort.
+    pub(crate) fn from_request(request: &ChatRequest) -> Result<Self> {
+        Self::from_kwargs(
+            &request.chat_options.template_kwargs,
+            request.chat_options.reasoning_effort.as_ref(),
+        )
+    }
+
+    /// Parse a deployment source, whose controls all come from template kwargs.
+    pub(crate) fn from_template_kwargs(kwargs: &HashMap<String, Value>) -> Result<Self> {
+        Self::from_kwargs(kwargs, None)
+    }
+
+    /// Select controls within one source, then derive its mode and optional effort.
+    fn from_kwargs(
+        kwargs: &HashMap<String, Value>,
+        typed_effort: Option<&EffortValue>,
+    ) -> Result<Self> {
+        // Presence sets alias priority: `thinking` wins, including when its value
+        // is invalid. Validate the selected spelling as a boolean.
+        let toggle = ["thinking", "enable_thinking"]
+            .into_iter()
+            .find_map(|key| kwargs.get(key).map(|value| (key, value)))
+            .map(|(key, value)| {
+                value.as_bool().ok_or_else(|| {
+                    invalid_reasoning_control!(
+                        "template kwarg `{key}` must be a boolean, got {value}"
+                    )
+                })
+            })
+            .transpose()?;
+
+        // The typed request field wins over kwargs. Parse kwargs only when used;
+        // disabled mode discards effort, and JSON null means omitted effort.
+        let effort = match (toggle, typed_effort) {
+            (Some(false), _) => None,
+            (_, Some(effort)) => Some(effort.clone()),
+            (_, None) => kwargs
+                .get("reasoning_effort")
+                .filter(|effort| !effort.is_null())
+                .cloned()
+                .map(EffortValue::try_from)
+                .transpose()?,
+        };
+
+        match (toggle, effort) {
+            // 1. Explicit false disables reasoning and discards any effort,
+            //    including values that would fail validation in an enabled mode.
+            (Some(false), _) => Ok(Self::Disabled),
+            // 2. `none` disables reasoning when the toggle is absent. Explicit true
+            //    enables reasoning and leaves effort to enabled deployment/model defaults.
+            (toggle, Some(EffortValue::String(effort))) if effort == "none" => {
+                Ok(if toggle == Some(true) {
+                    Self::Enabled { effort: None }
+                } else {
+                    Self::Disabled
+                })
+            }
+            // 3. Absent toggle and absent/null effort leave the entire decision
+            //    to the next source in the fallback chain.
+            (None, None) => Ok(Self::Default),
+            // 4. Explicit true or a concrete effort enables reasoning;
+            //    missing effort inherits later.
+            (_, effort) => Ok(Self::Enabled { effort }),
+        }
+    }
+
+    /// Parse request and deployment independently, then apply request precedence.
+    ///
+    /// Each source validates its selected input types. Model adapters subsequently
+    /// apply their model default and validate/map the resulting effective effort.
+    pub(crate) fn resolve(
+        request: &ChatRequest,
+        defaults: &HashMap<String, Value>,
+    ) -> Result<Self> {
+        Ok(Self::from_request(request)?.fallback(Self::from_template_kwargs(defaults)?))
+    }
+
+    /// Fill this higher-priority source's missing decision or effort from `defaults`.
+    ///
+    /// `Default` inherits the whole lower-priority state. `Enabled { effort: None }`
+    /// inherits effort only from enabled defaults and retains its enabled mode.
+    /// Explicit disabling and concrete effort are final at this priority level.
+    pub(crate) fn fallback(self, defaults: Self) -> Self {
+        match (self, defaults) {
+            (Self::Default, defaults) => defaults,
+            (Self::Enabled { effort: None }, Self::Enabled { effort }) => Self::Enabled { effort },
+            (request, _) => request,
+        }
+    }
+
+    /// Construct an enabled model default or a model-lowered concrete effort.
+    pub(crate) fn enabled(effort: impl Into<EffortValue>) -> Self {
+        Self::Enabled {
+            effort: Some(effort.into()),
+        }
+    }
+
+    /// Report the resolved mode; callers apply model defaults before rendering.
+    pub(crate) fn is_enabled(&self) -> bool {
+        matches!(self, Self::Enabled { .. })
+    }
+
+    /// Return the active effort, which may still be missing before model fallback.
+    pub(crate) fn effort(&self) -> Option<&EffortValue> {
+        match self {
+            Self::Enabled { effort } => effort.as_ref(),
+            Self::Default | Self::Disabled => None,
+        }
+    }
+
+    /// Project the same model-lowered decision used to render the native prompt.
+    ///
+    /// Preserve request extras and write both toggle aliases plus the effective
+    /// effort. Disabled mode emits `"none"`; enabled mode with no effective effort
+    /// clears the original effort key. `Default` preserves request kwargs.
+    pub(crate) fn template_kwargs(&self, request: &ChatRequest) -> HashMap<String, Value> {
+        let mut kwargs = request.chat_options.template_kwargs.clone();
+        if matches!(self, Self::Default) {
+            return kwargs;
+        }
+        kwargs.insert("thinking".to_string(), json!(self.is_enabled()));
+        kwargs.insert("enable_thinking".to_string(), json!(self.is_enabled()));
+        match self {
+            Self::Disabled => {
+                kwargs.insert("reasoning_effort".to_string(), json!("none"));
+            }
+            Self::Enabled {
+                effort: Some(effort),
+            } => {
+                kwargs.insert("reasoning_effort".to_string(), json!(effort));
+            }
+            Self::Enabled { effort: None } => {
+                kwargs.remove("reasoning_effort");
+            }
+            Self::Default => unreachable!(),
+        }
+        kwargs
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn kwargs(value: Value) -> HashMap<String, Value> {
+        serde_json::from_value(value).unwrap()
+    }
+
+    #[test]
+    fn numeric_effort_rejects_non_finite_values() {
+        for value in [f64::NAN, f64::INFINITY, f64::NEG_INFINITY] {
+            assert!(EffortValue::try_from(value).unwrap_err().is_request_validation_error());
+        }
+    }
+
+    #[test]
+    fn explicit_toggle_decides_mode_and_none_inherits_enabled_effort() {
+        use ReasoningControl::{Default, Disabled, Enabled};
+
+        for (source, expected) in [
+            (json!({}), Default),
+            (json!({"reasoning_effort": "none"}), Disabled),
+            (
+                json!({"reasoning_effort": "high"}),
+                ReasoningControl::enabled("high"),
+            ),
+            (json!({"enable_thinking": true}), Enabled { effort: None }),
+            (
+                json!({"enable_thinking": true, "reasoning_effort": "none"}),
+                Enabled { effort: None },
+            ),
+            (
+                json!({"enable_thinking": true, "reasoning_effort": "low"}),
+                ReasoningControl::enabled("low"),
+            ),
+            (
+                json!({"enable_thinking": false, "reasoning_effort": "high"}),
+                Disabled,
+            ),
+            (
+                json!({"thinking": true, "enable_thinking": false}),
+                Enabled { effort: None },
+            ),
+            (
+                json!({"thinking": false, "enable_thinking": true}),
+                Disabled,
+            ),
+            (
+                json!({"reasoning_effort": 37}),
+                ReasoningControl::enabled(EffortValue::Number(37.into())),
+            ),
+            (json!({"reasoning_effort": null}), Default),
+        ] {
+            let source = kwargs(source);
+            assert_eq!(
+                ReasoningControl::from_template_kwargs(&source).unwrap(),
+                expected,
+                "{source:?}"
+            );
+        }
+    }
+
+    #[test]
+    fn request_and_deployment_states_fall_back_before_model_defaults() {
+        for (request, defaults, expected) in [
+            (
+                json!({}),
+                json!({"enable_thinking": false}),
+                ReasoningControl::Disabled,
+            ),
+            (
+                json!({"reasoning_effort": "high"}),
+                json!({"thinking": false}),
+                ReasoningControl::enabled("high"),
+            ),
+            (
+                json!({"enable_thinking": false}),
+                json!({"reasoning_effort": "high"}),
+                ReasoningControl::Disabled,
+            ),
+            (
+                json!({"thinking": true, "reasoning_effort": "none"}),
+                json!({"reasoning_effort": "low"}),
+                ReasoningControl::enabled("low"),
+            ),
+            (
+                json!({"thinking": true}),
+                json!({"thinking": false, "reasoning_effort": "low"}),
+                ReasoningControl::enabled("max"),
+            ),
+            (json!({}), json!({}), ReasoningControl::enabled("max")),
+        ] {
+            let request = kwargs(request);
+            let defaults = kwargs(defaults);
+            let actual = ReasoningControl::from_template_kwargs(&request)
+                .unwrap()
+                .fallback(ReasoningControl::from_template_kwargs(&defaults).unwrap())
+                .fallback(ReasoningControl::enabled("max"));
+            assert_eq!(
+                actual, expected,
+                "request={request:?}, defaults={defaults:?}"
+            );
+        }
+    }
+
+    #[test]
+    fn selected_toggle_requires_a_boolean() {
+        for value in [
+            json!("true"),
+            json!(1),
+            json!(null),
+            json!({"type": "enabled"}),
+        ] {
+            assert!(
+                ReasoningControl::from_template_kwargs(&kwargs(json!({
+                    "thinking": value, "enable_thinking": true
+                })))
+                .unwrap_err()
+                .is_request_validation_error()
+            );
+        }
+    }
+
+    #[test]
+    fn invalid_effort_types_are_rejected_only_when_the_source_uses_them() {
+        for value in [json!(true), json!(false), json!([]), json!({})] {
+            for toggle in [None, Some(true), Some(false)] {
+                let mut source = kwargs(json!({"reasoning_effort": value}));
+                if let Some(toggle) = toggle {
+                    source.insert("thinking".to_owned(), json!(toggle));
+                }
+                let control = ReasoningControl::from_template_kwargs(&source);
+                if toggle == Some(false) {
+                    assert_eq!(control.unwrap(), ReasoningControl::Disabled);
+                } else {
+                    assert!(control.unwrap_err().is_request_validation_error());
+                }
+            }
+
+            // A typed request effort shadows malformed kwargs effort.
+            let mut request = ChatRequest::for_test();
+            request.chat_options.reasoning_effort = Some(EffortValue::from("high"));
+            request.chat_options.template_kwargs = kwargs(json!({"reasoning_effort": value}));
+            assert_eq!(
+                ReasoningControl::from_request(&request).unwrap(),
+                ReasoningControl::enabled("high")
+            );
+        }
+    }
+}

rust/src/chat/src/renderer/deepseek.rs

@@ -17,10 +17,10 @@ use serde_json_fmt::JsonFormat;
 use llm_multimodal::DEEPSEEK_V41_IMAGE_PLACEHOLDER;
 
 use super::MediaPartSource;
+use crate::EffortValue;
 use crate::error::{Error, Result};
-use crate::request::{
-    ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatTool, ReasoningEffort,
-};
+use crate::reasoning::ReasoningControl;
+use crate::request::{ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatTool};
 use crate::{AssistantContentBlock, AssistantMessageExt, AssistantToolCall};
 
 const BOS_TOKEN: &str = "<｜begin▁of▁sentence｜>";
@@ -87,21 +87,26 @@ struct RenderedToolSchema<'a> {
 }
 
 /// Render one chat request into the final prompt string.
-pub(super) fn render_request(request: &ChatRequest, dialect: DsDialect) -> Result<String> {
-    render_request_with_media_order(request, dialect).map(|(prompt, _)| prompt)
+pub(super) fn render_request(
+    request: &ChatRequest,
+    dialect: DsDialect,
+    reasoning: &ReasoningControl,
+) -> Result<String> {
+    render_request_with_media_order(request, dialect, reasoning).map(|(prompt, _)| prompt)
 }
 
 /// Render one chat request and record media in placeholder order.
 pub(super) fn render_request_with_media_order(
     request: &ChatRequest,
     dialect: DsDialect,
+    reasoning: &ReasoningControl,
 ) -> Result<(String, Vec<MediaPartSource>)> {
-    let (thinking_mode, reasoning_effort_prompt) = match dialect {
-        DsDialect::V4 => {
-            resolve_thinking_options(request).map(|(mode, prompt)| (mode, Cow::Borrowed(prompt)))?
-        }
-        DsDialect::V41 => resolve_v41_thinking_options(request)?,
+    let thinking_mode = if reasoning.is_enabled() {
+        ThinkingMode::Thinking
+    } else {
+        ThinkingMode::Chat
     };
+    let reasoning_effort_prompt = reasoning_prompt(reasoning, dialect);
     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)
@@ -254,76 +259,63 @@ pub(super) fn render_request_with_media_order(
     Ok((out, media_order))
 }
 
-/// Resolve DeepSeek V4's thinking controls. Unlike the Python tokenizer
-/// 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, &'static str)> {
-    let mut thinking_mode = match request.enable_thinking()?.unwrap_or(true) {
-        true => ThinkingMode::Thinking,
-        false => ThinkingMode::Chat,
-    };
-    let mut reasoning_effort_prompt = REASONING_EFFORT_HIGH;
-
-    match request.chat_options.reasoning_effort {
-        Some(ReasoningEffort::None) => thinking_mode = ThinkingMode::Chat,
-        Some(ReasoningEffort::Max) => {
-            reasoning_effort_prompt = REASONING_EFFORT_MAX;
-        }
-        Some(ReasoningEffort::XHigh | ReasoningEffort::High) => {
-            reasoning_effort_prompt = REASONING_EFFORT_HIGH;
-        }
-        Some(ReasoningEffort::Minimal | ReasoningEffort::Medium | ReasoningEffort::Low) => {
-            reasoning_effort_prompt = "";
-        }
-        None => {}
+/// Map the shared reasoning decision to the selected DeepSeek effort dialect.
+pub(super) fn resolve_reasoning(
+    control: ReasoningControl,
+    dialect: DsDialect,
+) -> Result<ReasoningControl> {
+    let control = control.fallback(ReasoningControl::enabled("high"));
+    if !control.is_enabled() {
+        return Ok(control);
     }
-
-    Ok((thinking_mode, reasoning_effort_prompt))
+    let effort = control.effort().expect("enabled DeepSeek reasoning has a default effort");
+    let effort = match dialect {
+        DsDialect::V4 => {
+            let effort = match effort.as_str() {
+                Some("minimal" | "medium" | "low") => "low",
+                Some("high" | "xhigh") => "high",
+                Some("max") => "max",
+                _ => {
+                    return Err(Error::InvalidReasoningEffort(format!(
+                        "DeepSeek V4 reasoning_effort must be minimal, low, medium, high, xhigh, or max, got {effort}"
+                    )));
+                }
+            };
+            EffortValue::from(effort)
+        }
+        DsDialect::V41 => {
+            let budget = match effort {
+                EffortValue::String(effort) => match effort.as_str() {
+                    "low" => 25,
+                    "high" => 50,
+                    "xhigh" => 75,
+                    "max" => 100,
+                    _ => return Err(invalid_v41_reasoning_effort()),
+                },
+                EffortValue::Number(budget) => budget
+                    .as_u64()
+                    .filter(|budget| (1..=100).contains(budget))
+                    .ok_or_else(invalid_v41_reasoning_effort)?,
+            };
+            EffortValue::Number(budget.into())
+        }
+    };
+    Ok(ReasoningControl::enabled(effort))
 }
 
-/// Resolve V4.1's numeric reasoning effort using the top-level value before template kwargs.
-fn resolve_v41_thinking_options(
-    request: &ChatRequest,
-) -> Result<(ThinkingMode, Cow<'static, str>)> {
-    let mut thinking = request.enable_thinking()?.unwrap_or(true);
-    let budget = match request.chat_options.reasoning_effort {
-        Some(ReasoningEffort::None) => {
-            thinking = false;
-            50
-        }
-        Some(ReasoningEffort::Low) => 25,
-        Some(ReasoningEffort::High) => 50,
-        Some(ReasoningEffort::XHigh) => 75,
-        Some(ReasoningEffort::Max) => 100,
-        Some(ReasoningEffort::Minimal | ReasoningEffort::Medium) => {
-            return Err(invalid_v41_reasoning_effort());
-        }
-        None => match request.chat_options.template_kwargs.get("reasoning_effort") {
-            None | Some(Value::Null) => 50,
-            Some(Value::String(effort)) => match effort.as_str() {
-                "low" => 25,
-                "high" => 50,
-                "xhigh" => 75,
-                "max" => 100,
-                _ => return Err(invalid_v41_reasoning_effort()),
-            },
-            Some(Value::Number(budget)) => budget
-                .as_u64()
-                .filter(|budget| (1..=100).contains(budget))
-                .ok_or_else(invalid_v41_reasoning_effort)?,
-            Some(_) => return Err(invalid_v41_reasoning_effort()),
-        },
+fn reasoning_prompt(control: &ReasoningControl, dialect: DsDialect) -> Cow<'static, str> {
+    let Some(effort) = control.effort() else {
+        return Cow::Borrowed("");
     };
-    if thinking {
-        Ok((
-            ThinkingMode::Thinking,
-            Cow::Owned(format!(
-                "Reasoning Effort: {budget} (range 1-100, the higher the value, the more thorough the reasoning)\n\n"
-            )),
-        ))
-    } else {
-        Ok((ThinkingMode::Chat, Cow::Borrowed("")))
+    match dialect {
+        DsDialect::V4 => Cow::Borrowed(match effort.as_str() {
+            Some("high") => REASONING_EFFORT_HIGH,
+            Some("max") => REASONING_EFFORT_MAX,
+            _ => "",
+        }),
+        DsDialect::V41 => Cow::Owned(format!(
+            "Reasoning Effort: {effort} (range 1-100, the higher the value, the more thorough the reasoning)\n\n"
+        )),
     }
 }
 

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

@@ -14,6 +14,7 @@ use serde_json::Value;
 use serde_json_fmt::JsonFormat;
 
 use crate::error::{Error, Result};
+use crate::reasoning::ReasoningControl;
 use crate::request::{ChatContent, ChatMessage, ChatRequest, ChatRole, ChatTool};
 use crate::{AssistantContentBlock, AssistantMessageExt, AssistantToolCall};
 
@@ -42,8 +43,11 @@ struct RenderedToolSchema<'a> {
 }
 
 /// Render one chat request into the final prompt string.
-pub(super) fn render_request(request: &ChatRequest) -> Result<String> {
-    let thinking_mode = match request.enable_thinking()?.unwrap_or(false) {
+pub(super) fn render_request(
+    request: &ChatRequest,
+    reasoning: &ReasoningControl,
+) -> Result<String> {
+    let thinking_mode = match reasoning.is_enabled() {
         true => ThinkingMode::Thinking,
         false => ThinkingMode::Chat,
     };

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

@@ -3,31 +3,41 @@
 
 mod encoding;
 
+use std::collections::HashMap;
+
+use serde_json::Value;
 use vllm_text::Prompt;
 
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{ChatRenderer, RenderedPrompt};
 use crate::Result;
+use crate::reasoning::ReasoningControl;
 use crate::request::ChatRequest;
 
 /// Dedicated DeepSeek V3.2 renderer.
-#[derive(Debug, Clone, Copy, Default)]
-pub struct DeepSeekV32ChatRenderer;
+#[derive(Debug, Clone, Default)]
+pub struct DeepSeekV32ChatRenderer {
+    default_template_kwargs: HashMap<String, Value>,
+}
 
 impl DeepSeekV32ChatRenderer {
     /// Create the dedicated DeepSeek V3.2 renderer.
-    pub fn new() -> Self {
-        Self
+    pub fn new(default_template_kwargs: HashMap<String, Value>) -> Self {
+        Self {
+            default_template_kwargs,
+        }
     }
 }
 
 impl ChatRenderer for DeepSeekV32ChatRenderer {
     fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
         request.validate()?;
+        let reasoning = ReasoningControl::resolve(request, &self.default_template_kwargs)?
+            .fallback(ReasoningControl::Disabled);
 
         Ok(RenderedPrompt {
-            prompt: Prompt::Text(encoding::render_request(request)?),
+            prompt: Prompt::Text(encoding::render_request(request, &reasoning)?),
             media_order: None,
-            effective_template_kwargs: request_template_kwargs(request),
+            effective_template_kwargs: reasoning.template_kwargs(request),
         })
     }
 }

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

@@ -8,6 +8,7 @@ use serde_json::{Value, json};
 use thiserror_ext::AsReport;
 
 use super::DeepSeekV32ChatRenderer;
+use crate::EffortValue;
 use crate::error::Error;
 use crate::event::{AssistantContentBlock, AssistantToolCall};
 use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
@@ -18,7 +19,7 @@ use crate::request::{
 use crate::{ChatRenderer, ChatRole};
 
 fn render_request(request: &ChatRequest) -> String {
-    DeepSeekV32ChatRenderer::new()
+    DeepSeekV32ChatRenderer::default()
         .render(request)
         .unwrap()
         .prompt
@@ -26,8 +27,30 @@ fn render_request(request: &ChatRequest) -> String {
         .expect("deepseek renderer should return text prompt")
 }
 
+#[test]
+fn effort_enables_binary_thinking_and_none_disables_it() {
+    let mut request = ChatRequest::for_test();
+    for (effort, enabled) in [
+        (EffortValue::from("high"), true),
+        (EffortValue::from("custom"), true),
+        (EffortValue::Number(42.into()), true),
+        (EffortValue::from("none"), false),
+    ] {
+        request.chat_options.reasoning_effort = Some(effort);
+        let rendered = DeepSeekV32ChatRenderer::default().render(&request).unwrap();
+        assert_eq!(
+            rendered.prompt.into_text().unwrap().ends_with("<think>"),
+            enabled
+        );
+        assert_eq!(
+            rendered.effective_template_kwargs["enable_thinking"],
+            enabled
+        );
+    }
+}
+
 fn render_result(request: &ChatRequest) -> Result<String, Error> {
-    DeepSeekV32ChatRenderer::new().render(request).map(|rendered| {
+    DeepSeekV32ChatRenderer::default().render(request).map(|rendered| {
         rendered
             .prompt
             .into_text()
@@ -345,7 +368,7 @@ fn render_rejects_multimodal_input() {
         ..ChatRequest::for_test()
     };
 
-    let error = DeepSeekV32ChatRenderer::new().render(&request).unwrap_err();
+    let error = DeepSeekV32ChatRenderer::default().render(&request).unwrap_err();
 
     assert!(matches!(
         error,

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

@@ -1,31 +1,47 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+use std::collections::HashMap;
+
+use serde_json::Value;
 use vllm_text::Prompt;
 
 use super::deepseek::{self, DsDialect};
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{ChatRenderer, RenderedPrompt};
 use crate::Result;
+use crate::reasoning::ReasoningControl;
 use crate::request::ChatRequest;
 
 /// Dedicated DeepSeek V4 renderer.
-#[derive(Debug, Clone, Copy, Default)]
-pub struct DeepSeekV4ChatRenderer;
+#[derive(Debug, Clone, Default)]
+pub struct DeepSeekV4ChatRenderer {
+    default_template_kwargs: HashMap<String, Value>,
+}
 
 impl DeepSeekV4ChatRenderer {
-    pub fn new() -> Self {
-        Self
+    pub fn new(default_template_kwargs: HashMap<String, Value>) -> Self {
+        Self {
+            default_template_kwargs,
+        }
     }
 }
 
 impl ChatRenderer for DeepSeekV4ChatRenderer {
     fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
         request.validate()?;
+        let reasoning = deepseek::resolve_reasoning(
+            ReasoningControl::resolve(request, &self.default_template_kwargs)?,
+            DsDialect::V4,
+        )?;
 
         Ok(RenderedPrompt {
-            prompt: Prompt::Text(deepseek::render_request(request, DsDialect::V4)?),
+            prompt: Prompt::Text(deepseek::render_request(
+                request,
+                DsDialect::V4,
+                &reasoning,
+            )?),
             media_order: None,
-            effective_template_kwargs: request_template_kwargs(request),
+            effective_template_kwargs: reasoning.template_kwargs(request),
         })
     }
 }

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

@@ -8,13 +8,14 @@ use serde_json::Value;
 
 use super::DeepSeekV4ChatRenderer;
 use crate::ChatRenderer;
+use crate::EffortValue;
 use crate::event::{AssistantContentBlock, AssistantToolCall};
 use crate::renderer::deepseek_v41::DeepSeekV41ChatRenderer;
 use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
-use crate::request::{ChatMessage, ChatRequest, ChatTool, GenerationPromptMode, ReasoningEffort};
+use crate::request::{ChatMessage, ChatRequest, ChatTool, GenerationPromptMode};
 
 fn render_request(request: &ChatRequest) -> String {
-    DeepSeekV4ChatRenderer::new()
+    DeepSeekV4ChatRenderer::default()
         .render(request)
         .unwrap()
         .prompt
@@ -31,13 +32,13 @@ fn thinking_request(messages: Vec<ChatMessage>) -> ChatRequest {
         .chat_options
         .template_kwargs
         .insert("thinking".to_string(), Value::Bool(true));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::Low);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("low"));
     request
 }
 
 fn fixture_request(input_name: &str) -> ChatRequest {
     let mut request = fixture_chat_request(&fixture_path(input_name), deepseek_fixture_options());
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::Low);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("low"));
     request
 }
 
@@ -203,7 +204,7 @@ fn reasoning_effort_max_adds_prefix_when_thinking_is_enabled() {
         .chat_options
         .template_kwargs
         .insert("thinking".to_string(), Value::Bool(true));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::Max);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("max"));
 
     let rendered = render_request(&request);
 
@@ -226,7 +227,7 @@ fn reasoning_effort_high_adds_0731_high_prefix() {
         .chat_options
         .template_kwargs
         .insert("thinking".to_string(), Value::Bool(true));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::High);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("high"));
 
     let rendered = render_request(&request);
 
@@ -258,7 +259,7 @@ fn reasoning_effort_xhigh_maps_to_high() {
         messages: vec![ChatMessage::user("solve it")],
         ..ChatRequest::for_test()
     };
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::XHigh);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("xhigh"));
 
     let rendered = render_request(&request);
 
@@ -268,7 +269,7 @@ fn reasoning_effort_xhigh_maps_to_high() {
 
 #[test]
 fn minimal_and_medium_reasoning_effort_map_to_low() {
-    for effort in [ReasoningEffort::Minimal, ReasoningEffort::Medium] {
+    for effort in [EffortValue::from("minimal"), EffortValue::from("medium")] {
         let mut request = ChatRequest {
             messages: vec![ChatMessage::user("solve it")],
             ..ChatRequest::for_test()
@@ -292,7 +293,7 @@ fn reasoning_effort_low_keeps_the_default_prompt() {
         .chat_options
         .template_kwargs
         .insert("thinking".to_string(), Value::Bool(true));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::Low);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("low"));
 
     let rendered = render_request(&request);
 
@@ -305,11 +306,7 @@ fn reasoning_effort_none_disables_thinking() {
         messages: vec![ChatMessage::user("answer directly")],
         ..ChatRequest::for_test()
     };
-    request
-        .chat_options
-        .template_kwargs
-        .insert("thinking".to_string(), Value::Bool(true));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::None);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("none"));
 
     let rendered = render_request(&request);
 
@@ -317,6 +314,49 @@ fn reasoning_effort_none_disables_thinking() {
         .assert_eq(&rendered);
 }
 
+#[test]
+fn normalized_reasoning_drives_both_prompt_and_compatibility_kwargs() {
+    for (kwargs, enabled, effort, prefix) in [
+        (
+            serde_json::json!({"enable_thinking": true, "reasoning_effort": "none"}),
+            true,
+            "high",
+            "Absolute maximum",
+        ),
+        (
+            serde_json::json!({"thinking": false, "enable_thinking": true, "reasoning_effort": "max"}),
+            false,
+            "none",
+            "",
+        ),
+        (
+            serde_json::json!({"reasoning_effort": "xhigh"}),
+            true,
+            "high",
+            "Absolute maximum",
+        ),
+    ] {
+        let mut request = ChatRequest::for_test();
+        request.chat_options.template_kwargs = serde_json::from_value(kwargs).unwrap();
+        let rendered = DeepSeekV4ChatRenderer::default().render(&request).unwrap();
+        let prompt = rendered.prompt.into_text().unwrap();
+        assert_eq!(prompt.ends_with("<think>"), enabled);
+        assert_eq!(prompt.contains("Reasoning Effort:"), enabled);
+        if enabled {
+            assert!(prompt.contains(prefix));
+        }
+        assert_eq!(rendered.effective_template_kwargs["thinking"], enabled);
+        assert_eq!(
+            rendered.effective_template_kwargs["enable_thinking"],
+            enabled
+        );
+        assert_eq!(
+            rendered.effective_template_kwargs["reasoning_effort"],
+            effort
+        );
+    }
+}
+
 #[test]
 fn system_to_assistant_transitions_in_chat_and_thinking_modes() {
     let rendered = [false, true].map(|enable_thinking| {
@@ -338,7 +378,7 @@ fn system_to_assistant_transitions_in_chat_and_thinking_modes() {
             .chat_options
             .template_kwargs
             .insert("enable_thinking".to_string(), Value::Bool(enable_thinking));
-        request.chat_options.reasoning_effort = Some(ReasoningEffort::Low);
+        request.chat_options.reasoning_effort = Some(EffortValue::from("low"));
         render_request(&request)
     });
 
@@ -366,7 +406,7 @@ fn trailing_system_transitions_in_chat_and_thinking_modes() {
             .chat_options
             .template_kwargs
             .insert("enable_thinking".to_string(), Value::Bool(enable_thinking));
-        request.chat_options.reasoning_effort = Some(ReasoningEffort::Low);
+        request.chat_options.reasoning_effort = Some(EffortValue::from("low"));
         render_request(&request)
     });
 
@@ -380,7 +420,7 @@ fn trailing_system_transitions_in_chat_and_thinking_modes() {
 }
 
 #[test]
-fn reasoning_effort_template_kwarg_is_ignored() {
+fn reasoning_effort_template_kwarg_selects_effort() {
     let mut request = ChatRequest {
         messages: vec![ChatMessage::user("solve it")],
         ..ChatRequest::for_test()
@@ -396,15 +436,15 @@ fn reasoning_effort_template_kwarg_is_ignored() {
 
     let rendered = render_request(&request);
 
-    assert!(rendered.starts_with("<｜begin▁of▁sentence｜>Reasoning Effort: Absolute maximum"));
+    assert!(rendered.starts_with("<｜begin▁of▁sentence｜>Reasoning Effort: Beyond maximum"));
     assert!(rendered.ends_with("<｜Assistant｜><think>"));
 }
 
 #[test]
 fn tool_call_arguments_preserve_non_objects_like_recipe() {
     let renderers: [(&dyn ChatRenderer, &str); 2] = [
-        (&DeepSeekV4ChatRenderer::new(), "parameter"),
-        (&DeepSeekV41ChatRenderer::new(), " parameter"),
+        (&DeepSeekV4ChatRenderer::default(), "parameter"),
+        (&DeepSeekV41ChatRenderer::default(), " parameter"),
     ];
     for arguments in [
         "not json",
@@ -600,7 +640,7 @@ fn drop_thinking_false_keeps_prior_assistant_reasoning() {
         .chat_options
         .template_kwargs
         .insert("drop_thinking".to_string(), Value::Bool(false));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::Low);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("low"));
 
     let rendered = render_request(&request);
 
@@ -693,7 +733,7 @@ fn trailing_system_omits_generation_prompt_when_disabled() {
         .chat_options
         .template_kwargs
         .insert("enable_thinking".to_string(), Value::Bool(false));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::Low);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("low"));
 
     let rendered = render_request(&request);
 

rust/src/chat/src/renderer/deepseek_v41/mod.rs

@@ -1,33 +1,45 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+use std::collections::HashMap;
+
+use serde_json::Value;
 use vllm_text::Prompt;
 
 use super::deepseek::{self, DsDialect};
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{ChatRenderer, RenderedPrompt};
 use crate::Result;
+use crate::reasoning::ReasoningControl;
 use crate::request::ChatRequest;
 
 /// Dedicated DeepSeek V4.1 renderer.
-#[derive(Debug, Clone, Copy, Default)]
-pub struct DeepSeekV41ChatRenderer;
+#[derive(Debug, Clone, Default)]
+pub struct DeepSeekV41ChatRenderer {
+    default_template_kwargs: HashMap<String, Value>,
+}
 
 impl DeepSeekV41ChatRenderer {
-    pub fn new() -> Self {
-        Self
+    pub fn new(default_template_kwargs: HashMap<String, Value>) -> Self {
+        Self {
+            default_template_kwargs,
+        }
     }
 }
 
 impl ChatRenderer for DeepSeekV41ChatRenderer {
     fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
         request.validate()?;
+        let reasoning = deepseek::resolve_reasoning(
+            ReasoningControl::resolve(request, &self.default_template_kwargs)?,
+            DsDialect::V41,
+        )?;
         let (prompt, media_order) =
-            deepseek::render_request_with_media_order(request, DsDialect::V41)?;
+            deepseek::render_request_with_media_order(request, DsDialect::V41, &reasoning)?;
 
         Ok(RenderedPrompt {
             prompt: Prompt::Text(prompt),
             media_order: Some(media_order),
-            effective_template_kwargs: request_template_kwargs(request),
+            effective_template_kwargs: reasoning.template_kwargs(request),
         })
     }
 }

rust/src/chat/src/renderer/deepseek_v41/tests.rs

@@ -8,14 +8,15 @@ use serde_json::{Value, json};
 
 use super::DeepSeekV41ChatRenderer;
 use crate::ChatRenderer;
+use crate::EffortValue;
 use crate::event::{AssistantContentBlock, AssistantToolCall};
 use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
 use crate::request::{
-    ChatContent, ChatContentPart, ChatMessage, ChatRequest, GenerationPromptMode, ReasoningEffort,
+    ChatContent, ChatContentPart, ChatMessage, ChatRequest, GenerationPromptMode,
 };
 
 fn render(request: &ChatRequest) -> String {
-    DeepSeekV41ChatRenderer::new()
+    DeepSeekV41ChatRenderer::default()
         .render(request)
         .unwrap()
         .prompt
@@ -64,7 +65,7 @@ fn inlines_image_placeholder_at_image_part_positions() {
     // The placeholder is unconditional (matching the Python encoding's
     // `IMAGE_PLACEHOLDER`); requests to a multimodal-less backend fail later,
     // at media preparation.
-    let renderer = DeepSeekV41ChatRenderer::new();
+    let renderer = DeepSeekV41ChatRenderer::default();
     let request = ChatRequest {
         messages: vec![ChatMessage::User {
             content: ChatContent::Parts(vec![
@@ -87,7 +88,7 @@ fn inlines_image_placeholder_at_image_part_positions() {
 
 #[test]
 fn media_order_follows_reordered_tool_results() {
-    let renderer = DeepSeekV41ChatRenderer::new();
+    let renderer = DeepSeekV41ChatRenderer::default();
     let request = ChatRequest {
         messages: vec![
             ChatMessage::assistant_blocks(vec![
@@ -136,7 +137,7 @@ fn media_order_follows_reordered_tool_results() {
 
 #[test]
 fn media_order_excludes_dropped_historical_developer_content() {
-    let renderer = DeepSeekV41ChatRenderer::new();
+    let renderer = DeepSeekV41ChatRenderer::default();
     let request = ChatRequest {
         messages: vec![
             ChatMessage::developer(vec![image_part("historical-image")], None),
@@ -165,10 +166,11 @@ fn image_part(uuid: &str) -> ChatContentPart {
 fn maps_reasoning_effort_to_reference_numeric_budget() {
     for (effort, budget) in [
         (None, 50),
-        (Some(ReasoningEffort::Low), 25),
-        (Some(ReasoningEffort::High), 50),
-        (Some(ReasoningEffort::XHigh), 75),
-        (Some(ReasoningEffort::Max), 100),
+        (Some(EffortValue::from("low")), 25),
+        (Some(EffortValue::from("high")), 50),
+        (Some(EffortValue::from("xhigh")), 75),
+        (Some(EffortValue::from("max")), 100),
+        (Some(EffortValue::Number(37.into())), 37),
     ] {
         let mut request = request();
         request.chat_options.reasoning_effort = effort;
@@ -193,11 +195,43 @@ fn accepts_numeric_template_effort_with_top_level_precedence() {
         request.chat_options.template_kwargs.insert("reasoning_effort".into(), effort);
         assert!(render(&request).contains(&format!("Reasoning Effort: {budget} (range 1-100")));
 
-        request.chat_options.reasoning_effort = Some(ReasoningEffort::XHigh);
+        request.chat_options.reasoning_effort = Some(EffortValue::from("xhigh"));
         assert!(render(&request).contains("Reasoning Effort: 75 (range 1-100"));
     }
 }
 
+#[test]
+fn enabling_with_none_inherits_deployment_effort_and_reports_numeric_control() {
+    let renderer =
+        DeepSeekV41ChatRenderer::new([("reasoning_effort".to_string(), json!(37))].into());
+    let mut request = request();
+    request.chat_options.reasoning_effort = Some(EffortValue::from("none"));
+    request
+        .chat_options
+        .template_kwargs
+        .insert("enable_thinking".into(), json!(true));
+    let rendered = renderer.render(&request).unwrap();
+    let prompt = rendered.prompt.into_text().unwrap();
+    assert!(prompt.contains("Reasoning Effort: 37 (range 1-100"));
+    assert!(prompt.ends_with("<think>"));
+    assert_eq!(rendered.effective_template_kwargs["reasoning_effort"], 37);
+    assert_eq!(rendered.effective_template_kwargs["enable_thinking"], true);
+
+    request.chat_options.template_kwargs.insert("thinking".into(), json!(false));
+    request
+        .chat_options
+        .template_kwargs
+        .insert("reasoning_effort".into(), json!(101));
+    request.chat_options.reasoning_effort = None;
+    let rendered = renderer.render(&request).unwrap();
+    assert!(rendered.prompt.into_text().unwrap().ends_with("</think>"));
+    assert_eq!(
+        rendered.effective_template_kwargs["reasoning_effort"],
+        "none"
+    );
+    assert_eq!(rendered.effective_template_kwargs["enable_thinking"], false);
+}
+
 #[test]
 fn rejects_undefined_effort_names_and_invalid_numeric_budgets() {
     for effort in [
@@ -211,13 +245,13 @@ fn rejects_undefined_effort_names_and_invalid_numeric_budgets() {
     ] {
         let mut request = request();
         request.chat_options.template_kwargs.insert("reasoning_effort".into(), effort);
-        let error = DeepSeekV41ChatRenderer::new().render(&request).unwrap_err();
+        let error = DeepSeekV41ChatRenderer::default().render(&request).unwrap_err();
         assert!(error.is_request_validation_error());
     }
-    for effort in [ReasoningEffort::Minimal, ReasoningEffort::Medium] {
+    for effort in [EffortValue::from("minimal"), EffortValue::from("medium")] {
         let mut request = request();
         request.chat_options.reasoning_effort = Some(effort);
-        let error = DeepSeekV41ChatRenderer::new().render(&request).unwrap_err();
+        let error = DeepSeekV41ChatRenderer::default().render(&request).unwrap_err();
         assert!(error.is_request_validation_error());
     }
 }
@@ -229,8 +263,8 @@ fn chat_mode_and_none_effort_use_closed_thinking_prefix() {
     let expected = expect!["<｜begin▁of▁sentence｜><｜User｜>question<｜Assistant｜></think>"];
     expected.assert_eq(&render(&request));
 
-    request.chat_options.template_kwargs.insert("thinking".into(), json!(true));
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::None);
+    request.chat_options.template_kwargs.remove("thinking");
+    request.chat_options.reasoning_effort = Some(EffortValue::from("none"));
     expected.assert_eq(&render(&request));
 }
 

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

@@ -5,21 +5,26 @@
 
 pub(crate) mod encoding;
 
+use std::collections::HashMap;
+
 use openai_harmony::HarmonyEncoding;
 use openai_harmony::chat::{
     Author, Conversation, DeveloperContent, Message, ReasoningEffort as HarmonyReasoningEffort,
     Role, SystemContent, ToolDescription,
 };
+use serde_json::Value;
 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};
+use super::{ChatRenderer, RenderedPrompt};
+use crate::AssistantMessageExt as _;
+use crate::EffortValue;
 use crate::error::{Error, Result};
 use crate::event::AssistantContentBlock;
+use crate::reasoning::ReasoningControl;
 use crate::request::{ChatContent, ChatMessage, ChatRequest, ChatTool, GenerationPromptMode};
-use crate::{AssistantMessageExt as _, ReasoningEffort};
 
 const SYSTEM_START_DATE_ENV: &str = "VLLM_SYSTEM_START_DATE";
 const HARMONY_SYSTEM_INSTRUCTIONS_ENV: &str = "VLLM_GPT_OSS_HARMONY_SYSTEM_INSTRUCTIONS";
@@ -28,6 +33,7 @@ const HARMONY_SYSTEM_INSTRUCTIONS_ENV: &str = "VLLM_GPT_OSS_HARMONY_SYSTEM_INSTR
 pub struct HarmonyChatRenderer {
     encoding: &'static HarmonyEncoding,
     options: Options,
+    default_template_kwargs: HashMap<String, Value>,
 }
 
 struct Options {
@@ -44,8 +50,9 @@ impl HarmonyChatRenderer {
     ///   unset, the renderer uses the current local date with a UTC fallback.
     /// - `VLLM_GPT_OSS_HARMONY_SYSTEM_INSTRUCTIONS` moves leading instructions
     ///   into the system model identity when set to a non-zero integer.
-    pub fn new() -> Result<Self> {
+    pub fn new(default_template_kwargs: HashMap<String, Value>) -> Result<Self> {
         Self::with_options(
+            default_template_kwargs,
             env_system_start_date(),
             env_use_harmony_system_instructions(),
         )
@@ -57,11 +64,13 @@ impl HarmonyChatRenderer {
     /// Production code should call [`Self::new`] so the renderer observes the
     /// same environment contract as the Python Harmony path.
     pub fn with_options(
+        default_template_kwargs: HashMap<String, Value>,
         system_start_date: impl Into<String>,
         use_system_instructions: bool,
     ) -> Result<Self> {
         Ok(Self {
             encoding: harmony_encoding()?,
+            default_template_kwargs,
             options: Options {
                 system_start_date: system_start_date.into(),
                 use_system_instructions,
@@ -73,7 +82,11 @@ 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: &ChatRequest,
+        effort: HarmonyReasoningEffort,
+    ) -> Result<Vec<u32>> {
         if request.has_multimodal() {
             return Err(Error::UnsupportedMultimodalContent("image_url"));
         }
@@ -86,7 +99,8 @@ impl HarmonyChatRenderer {
             ));
         }
 
-        let messages = auto_drop_analysis_messages(to_harmony_messages(request, &self.options)?);
+        let messages =
+            auto_drop_analysis_messages(to_harmony_messages(request, &self.options, effort)?);
         let conversation = Conversation::from_messages(messages);
         // Pass `None` so oss-harmony does not apply its narrower built-in
         // analysis-drop policy after the Rust-side Python-parity cleanup above.
@@ -114,10 +128,13 @@ 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> {
+        let reasoning = ReasoningControl::resolve(request, &self.default_template_kwargs)?
+            .fallback(ReasoningControl::enabled("medium"));
+        let effort = to_harmony_reasoning_effort(reasoning.effort())?;
         Ok(RenderedPrompt {
-            prompt: Prompt::TokenIds(self.render_token_ids(request)?),
+            prompt: Prompt::TokenIds(self.render_token_ids(request, effort)?),
             media_order: None,
-            effective_template_kwargs: request_template_kwargs(request),
+            effective_template_kwargs: reasoning.template_kwargs(request),
         })
     }
 }
@@ -127,12 +144,21 @@ 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: &ChatRequest,
+    options: &Options,
+    effort: HarmonyReasoningEffort,
+) -> Result<Vec<Message>> {
     let (instructions, leading_developer_tools, remaining_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)?;
+    let mut messages = build_harmony_preamble(
+        request,
+        instructions,
+        leading_developer_tools,
+        options,
+        effort,
+    )?;
 
     for message in remaining_messages {
         messages.extend(to_harmony_message(message, &tool_call_names, options)?);
@@ -177,12 +203,13 @@ fn build_harmony_preamble(
     instructions: Option<String>,
     leading_developer_tools: Option<&[ChatTool]>,
     options: &Options,
+    effort: HarmonyReasoningEffort,
 ) -> Result<Vec<Message>> {
     let mut messages = vec![Message::from_role_and_content(
         Role::System,
         system_content(
             instructions.as_deref().filter(|_| options.use_system_instructions),
-            request.chat_options.reasoning_effort,
+            Some(effort),
             &options.system_start_date,
         )?,
     )];
@@ -227,18 +254,18 @@ fn preamble_tool_descriptions(
 
 /// Construct the Harmony system content for the request preamble.
 ///
-/// Harmony defaults the reasoning effort to `medium` when none is provided, so
-/// this only sets an explicit effort after validating vLLM's request value.
+/// The preamble receives normalized effort; later system messages use Harmony's
+/// own default when their effort is omitted.
 fn system_content(
     instructions: Option<&str>,
-    reasoning_effort: Option<ReasoningEffort>,
+    reasoning_effort: Option<HarmonyReasoningEffort>,
     system_start_date: &str,
 ) -> Result<SystemContent> {
     let mut content =
         SystemContent::new().with_conversation_start_date(system_start_date.to_string());
 
     if let Some(reasoning_effort) = reasoning_effort {
-        content = content.with_reasoning_effort(to_harmony_reasoning_effort(reasoning_effort)?);
+        content = content.with_reasoning_effort(reasoning_effort);
     }
 
     if let Some(instructions) = instructions.filter(|text| !text.is_empty()) {
@@ -445,20 +472,17 @@ fn to_tool_descriptions(tools: &[ChatTool]) -> Vec<ToolDescription> {
         .collect()
 }
 
-/// Map supported OpenAI reasoning-effort values onto Harmony's enum.
+/// Map supported reasoning-effort values onto Harmony's enum.
 fn to_harmony_reasoning_effort(
-    reasoning_effort: ReasoningEffort,
+    reasoning_effort: Option<&EffortValue>,
 ) -> Result<HarmonyReasoningEffort> {
-    match reasoning_effort {
-        ReasoningEffort::Low => Ok(HarmonyReasoningEffort::Low),
-        ReasoningEffort::Medium => Ok(HarmonyReasoningEffort::Medium),
-        ReasoningEffort::High => Ok(HarmonyReasoningEffort::High),
-        ReasoningEffort::None
-        | ReasoningEffort::Minimal
-        | ReasoningEffort::XHigh
-        | ReasoningEffort::Max => Err(Error::ChatTemplate(format!(
-            "reasoning_effort={:?} is not supported by Harmony. Supported values are: low, medium, high.",
-            reasoning_effort.as_str()
+    match reasoning_effort.and_then(EffortValue::as_str) {
+        Some("low") => Ok(HarmonyReasoningEffort::Low),
+        Some("medium") => Ok(HarmonyReasoningEffort::Medium),
+        Some("high") => Ok(HarmonyReasoningEffort::High),
+        _ => Err(Error::InvalidReasoningEffort(format!(
+            "reasoning_effort={} is not supported by Harmony. Supported values are: low, medium, high.",
+            reasoning_effort.unwrap_or(&EffortValue::from("none"))
         ))),
     }
 }

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

@@ -9,12 +9,11 @@ use thiserror_ext::AsReport as _;
 use super::HarmonyChatRenderer;
 use super::encoding::harmony_encoding;
 use crate::ChatRenderer;
+use crate::EffortValue;
 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::request::{ChatContentPart, ChatMessage, ChatRequest, GenerationPromptMode};
 
 const PINNED_DATE: &str = "2025-06-28";
 
@@ -22,7 +21,7 @@ fn fixture_request(input_name: &str) -> ChatRequest {
     fixture_chat_request(
         &fixture_path(input_name),
         FixtureRequestOptions {
-            enable_thinking: Some(false),
+            enable_thinking: None,
             no_generation_prompt_when_last_assistant: false,
         },
     )
@@ -36,7 +35,8 @@ fn fixture_path(name: &str) -> PathBuf {
 }
 
 fn test_renderer(use_system_instructions: bool) -> HarmonyChatRenderer {
-    HarmonyChatRenderer::with_options(PINNED_DATE, use_system_instructions).unwrap()
+    HarmonyChatRenderer::with_options(Default::default(), PINNED_DATE, use_system_instructions)
+        .unwrap()
 }
 
 fn render_token_ids(request: &ChatRequest) -> Vec<u32> {
@@ -138,14 +138,48 @@ fn drops_stale_analysis_fixture() {
 #[test]
 fn rejects_invalid_reasoning_effort() {
     let mut request = ChatRequest::for_test();
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::None);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("none"));
 
     let error = test_renderer(false).render(&request).unwrap_err();
 
-    expect![[r#"chat template error: reasoning_effort="none" is not supported by Harmony. Supported values are: low, medium, high."#]]
+    expect![[r#"reasoning_effort="none" is not supported by Harmony. Supported values are: low, medium, high."#]]
         .assert_eq(&error.to_report_string());
 }
 
+#[test]
+fn normalized_reasoning_respects_harmony_capability_and_deployment_effort() {
+    let renderer = HarmonyChatRenderer::with_options(
+        [("reasoning_effort".to_string(), serde_json::json!("low"))].into(),
+        PINNED_DATE,
+        false,
+    )
+    .unwrap();
+    let mut request = ChatRequest::for_test();
+    request.chat_options.reasoning_effort = Some(EffortValue::from("none"));
+    request
+        .chat_options
+        .template_kwargs
+        .insert("enable_thinking".into(), serde_json::json!(true));
+    let rendered = renderer.render(&request).unwrap();
+    let prompt = harmony_encoding()
+        .unwrap()
+        .tokenizer()
+        .decode_utf8(rendered.prompt.into_token_ids().unwrap())
+        .unwrap();
+    assert!(prompt.contains("Reasoning: low"));
+    assert_eq!(
+        rendered.effective_template_kwargs["reasoning_effort"],
+        "low"
+    );
+    assert_eq!(rendered.effective_template_kwargs["enable_thinking"], true);
+
+    request
+        .chat_options
+        .template_kwargs
+        .insert("thinking".into(), serde_json::json!(false));
+    assert!(renderer.render(&request).is_err());
+}
+
 #[test]
 fn rejects_unknown_tool_response_id() {
     let request = ChatRequest {

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

@@ -16,7 +16,7 @@ use self::format::{
     ChatTemplateContentFormat, ChatTemplateContentFormatOption as ContentFormatOption,
 };
 use self::template::{CompiledChatTemplate, TemplateContext};
-use super::{ChatRenderer, RenderedPrompt, effective_template_kwargs};
+use super::{ChatRenderer, RenderedPrompt};
 use crate::error::Result;
 use crate::request::{ChatContent, ChatContentPart, ChatMessage, ChatRequest};
 use crate::{
@@ -33,6 +33,31 @@ pub use template::{load_chat_template, resolve_chat_template};
 
 pub use self::format::ChatTemplateContentFormatOption;
 
+/// Extract the effective chat-template kwargs visible to the renderer from the request,
+/// using the provided defaults as the base.
+fn effective_template_kwargs(
+    default_template_kwargs: &HashMap<String, JsonValue>,
+    request: &ChatRequest,
+) -> HashMap<String, JsonValue> {
+    let mut kwargs = default_template_kwargs.clone();
+    kwargs.extend(request.chat_options.template_kwargs.clone());
+
+    if let Some(reasoning_effort) = &request.chat_options.reasoning_effort {
+        kwargs.insert(
+            "reasoning_effort".to_string(),
+            serde_json::json!(reasoning_effort),
+        );
+        if !request.chat_options.template_kwargs.contains_key("enable_thinking") {
+            kwargs.insert(
+                "enable_thinking".to_string(),
+                serde_json::json!(reasoning_effort.as_str() != Some("none")),
+            );
+        }
+    }
+
+    kwargs
+}
+
 /// Template-visible placeholder tokens per supported modality.
 ///
 /// A `None` token means the loaded model does not support that modality, and
@@ -580,9 +605,10 @@ mod tests {
     use vllm_text::backend::hf::{HfSpecialTokens, NamedSpecialToken};
 
     use super::{ChatTemplateContentFormatOption, HfChatRenderer, MultimodalRenderInfo};
+    use crate::EffortValue;
     use crate::request::{
         ChatContentPart, ChatMessage, ChatRequest, ChatRole, ChatTool, ChatToolChoice,
-        GenerationPromptMode, ReasoningEffort, ResolvedToolContext,
+        GenerationPromptMode, ResolvedToolContext,
     };
     use crate::{AssistantContentBlock, ChatRenderer, Error, Result};
 
@@ -1149,7 +1175,7 @@ mod tests {
     #[test]
     fn chat_template_reasoning_effort_overrides_template_kwargs() {
         let mut request = sample_request(vec![ChatMessage::text(ChatRole::User, "hello")]);
-        request.chat_options.reasoning_effort = Some(ReasoningEffort::Max);
+        request.chat_options.reasoning_effort = Some(EffortValue::from("max"));
         request.chat_options.template_kwargs.insert(
             "reasoning_effort".to_string(),
             Value::String("low".to_string()),
@@ -1181,7 +1207,7 @@ mod tests {
     #[test]
     fn chat_template_reasoning_effort_preserves_request_enable_thinking() {
         let mut request = sample_request(vec![ChatMessage::text(ChatRole::User, "hello")]);
-        request.chat_options.reasoning_effort = Some(ReasoningEffort::None);
+        request.chat_options.reasoning_effort = Some(EffortValue::from("none"));
         request
             .chat_options
             .template_kwargs
@@ -1207,6 +1233,70 @@ mod tests {
         );
     }
 
+    #[test]
+    fn chat_template_preserves_typed_effort_values() {
+        for effort in [
+            serde_json::json!(37),
+            serde_json::json!(0.37),
+            serde_json::json!("custom"),
+        ] {
+            let mut request = sample_request(vec![ChatMessage::user("hello")]);
+            request.chat_options.reasoning_effort =
+                Some(serde_json::from_value(effort.clone()).unwrap());
+            request.chat_options.template_kwargs = [
+                ("reasoning_effort".to_string(), serde_json::json!("high")),
+                ("enable_thinking".to_string(), serde_json::json!(true)),
+            ]
+            .into();
+            let rendered = HfChatRenderer::new(
+                Some("{{ reasoning_effort|tojson }}|{{ enable_thinking }}".to_string()),
+                HashMap::new(),
+                ChatTemplateContentFormatOption::Auto,
+            )
+            .unwrap()
+            .render(&request)
+            .unwrap();
+            assert_eq!(rendered.prompt, Prompt::Text(format!("{effort}|True")));
+            assert_eq!(
+                rendered.effective_template_kwargs["reasoning_effort"],
+                effort
+            );
+        }
+    }
+
+    #[test]
+    fn chat_template_keeps_native_control_conflicts_and_raw_effort_values() {
+        for effort in [
+            serde_json::json!(37),
+            serde_json::json!(null),
+            serde_json::json!("none"),
+        ] {
+            let mut request = sample_request(vec![ChatMessage::user("hello")]);
+            request.chat_options.template_kwargs = [
+                ("thinking".to_string(), serde_json::json!("custom-mode")),
+                ("enable_thinking".to_string(), serde_json::json!(true)),
+                ("reasoning_effort".to_string(), effort),
+            ]
+            .into();
+            let rendered = HfChatRenderer::new(
+                Some(
+                    "{{ thinking }}|{{ enable_thinking }}|{{ reasoning_effort is none }}"
+                        .to_string(),
+                ),
+                HashMap::new(),
+                ChatTemplateContentFormatOption::Auto,
+            )
+            .unwrap()
+            .render(&request)
+            .unwrap();
+            assert!(rendered.prompt.into_text().unwrap().starts_with("custom-mode|True|"));
+            assert_eq!(
+                rendered.effective_template_kwargs,
+                request.chat_options.template_kwargs
+            );
+        }
+    }
+
     #[test]
     fn qwen3_template_omits_reasoning_for_historical_assistant_messages() {
         let request = sample_request(vec![

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

@@ -8,8 +8,10 @@ use thiserror_ext::AsReport as _;
 use vllm_text::Prompt;
 use vllm_text::tokenizer::{DynTokenizer, Tokenizer};
 
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{ChatRenderer, RenderedPrompt};
+use crate::EffortValue;
 use crate::error::{Error, Result};
+use crate::reasoning::ReasoningControl;
 use crate::request::{ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatTool};
 use crate::{AssistantContentBlock, AssistantToolCall};
 
@@ -35,6 +37,7 @@ const MAX_REASONING_EFFORT: f64 = 0.99;
 pub struct InklingChatRenderer {
     tokenizer: DynTokenizer,
     special: InklingSpecialTokenIds,
+    default_template_kwargs: HashMap<String, Value>,
 }
 
 #[derive(Debug, Clone, Copy)]
@@ -78,9 +81,16 @@ impl InklingSpecialTokenIds {
 }
 
 impl InklingChatRenderer {
-    pub fn new(tokenizer: DynTokenizer) -> Result<Self> {
+    pub fn new(
+        tokenizer: DynTokenizer,
+        default_template_kwargs: HashMap<String, Value>,
+    ) -> Result<Self> {
         let special = InklingSpecialTokenIds::resolve(tokenizer.as_ref())?;
-        Ok(Self { tokenizer, special })
+        Ok(Self {
+            tokenizer,
+            special,
+            default_template_kwargs,
+        })
     }
 
     fn write_text_tokens(&self, out: &mut Vec<u32>, text: &str) -> Result<()> {
@@ -150,11 +160,6 @@ impl InklingChatRenderer {
     }
 
     fn write_reasoning_effort(&self, out: &mut Vec<u32>, effort: f64) -> Result<()> {
-        if !(0.0..=MAX_REASONING_EFFORT).contains(&effort) {
-            return Err(Error::ChatTemplate(format!(
-                "Inkling reasoning_effort must be in [0.0, 0.99], got {effort}"
-            )));
-        }
         let formatted = format!("{effort:.2}");
         let effort = formatted.trim_end_matches('0').trim_end_matches('.');
         let effort = if matches!(effort, "0" | "-0") {
@@ -294,11 +299,18 @@ 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 reasoning = resolve_reasoning(ReasoningControl::resolve(
+            request,
+            &self.default_template_kwargs,
+        )?)?;
+        let effective_template_kwargs = reasoning.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"));
+        let mut reasoning_effort = if reasoning.is_enabled() {
+            reasoning.effort().and_then(EffortValue::as_f64)
+        } else {
+            Some(0.0)
+        };
 
         for message in &request.messages {
             if !matches!(
@@ -352,12 +364,33 @@ impl ChatRenderer for InklingChatRenderer {
     }
 }
 
-fn resolve_reasoning_effort(value: Option<&Value>) -> Option<f64> {
+fn resolve_reasoning(control: ReasoningControl) -> Result<ReasoningControl> {
+    let control = control.fallback(ReasoningControl::enabled("high"));
+    if !control.is_enabled() {
+        return Ok(control);
+    }
+    let effort = resolve_reasoning_effort(control.effort());
+    if let Some(effort) = effort {
+        if !(0.0..=MAX_REASONING_EFFORT).contains(&effort) {
+            return Err(Error::InvalidReasoningEffort(format!(
+                "Inkling reasoning_effort must be in [0.0, 0.99], got {effort}"
+            )));
+        }
+        if effort == 0.0 {
+            return Ok(ReasoningControl::Disabled);
+        }
+    }
+    Ok(ReasoningControl::Enabled {
+        effort: effort.map(EffortValue::try_from).transpose()?,
+    })
+}
+
+fn resolve_reasoning_effort(value: Option<&EffortValue>) -> Option<f64> {
     let Some(value) = value else {
         return Some(0.9);
     };
     match value {
-        Value::String(name) => match name.as_str() {
+        EffortValue::String(name) => match name.as_str() {
             "none" => Some(0.0),
             "minimal" => Some(0.1),
             "low" => Some(0.2),
@@ -366,8 +399,7 @@ fn resolve_reasoning_effort(value: Option<&Value>) -> Option<f64> {
             "xhigh" | "max" => Some(0.99),
             _ => None,
         },
-        Value::Number(number) => number.as_f64(),
-        _ => None,
+        EffortValue::Number(number) => number.as_f64(),
     }
 }
 

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

@@ -9,6 +9,7 @@ use serde_json::json;
 use thiserror_ext::AsReport;
 use vllm_text::tokenizer::Tokenizer;
 
+use crate::EffortValue;
 use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
 
 use super::{
@@ -17,7 +18,7 @@ use super::{
     MESSAGE_TOOL, MESSAGE_USER,
 };
 use crate::event::{AssistantContentBlock, AssistantToolCall};
-use crate::request::{ChatMessage, ChatRequest, GenerationPromptMode, ReasoningEffort};
+use crate::request::{ChatMessage, ChatRequest, GenerationPromptMode};
 use crate::{ChatRenderer, Error};
 
 struct FixtureTokenizer;
@@ -87,7 +88,7 @@ impl Tokenizer for FixtureTokenizer {
 }
 
 fn renderer() -> InklingChatRenderer {
-    InklingChatRenderer::new(Arc::new(FixtureTokenizer)).unwrap()
+    InklingChatRenderer::new(Arc::new(FixtureTokenizer), Default::default()).unwrap()
 }
 
 fn render_token_ids(request: &ChatRequest) -> Vec<u32> {
@@ -105,7 +106,7 @@ fn fixture_request(name: &str) -> ChatRequest {
 
 fn inkling_fixture_options() -> FixtureRequestOptions {
     FixtureRequestOptions {
-        enable_thinking: Some(false),
+        enable_thinking: None,
         no_generation_prompt_when_last_assistant: false,
     }
 }
@@ -206,13 +207,13 @@ fn renders_request_and_developer_tools_as_tool_declare() {
 #[test]
 fn renders_named_reasoning_effort_after_tool_declarations() {
     for (effort, expected) in [
-        (ReasoningEffort::None, "0.0"),
-        (ReasoningEffort::Minimal, "0.1"),
-        (ReasoningEffort::Low, "0.2"),
-        (ReasoningEffort::Medium, "0.7"),
-        (ReasoningEffort::High, "0.9"),
-        (ReasoningEffort::XHigh, "0.99"),
-        (ReasoningEffort::Max, "0.99"),
+        (EffortValue::from("none"), "0.0"),
+        (EffortValue::from("minimal"), "0.1"),
+        (EffortValue::from("low"), "0.2"),
+        (EffortValue::from("medium"), "0.7"),
+        (EffortValue::from("high"), "0.9"),
+        (EffortValue::from("xhigh"), "0.99"),
+        (EffortValue::from("max"), "0.99"),
     ] {
         let mut request = fixture_request("tool_declare_input.json");
         request.chat_options.reasoning_effort = Some(effort);
@@ -284,6 +285,43 @@ fn defaults_reasoning_effort_to_high() {
     );
 }
 
+#[test]
+fn normalized_mode_and_scalar_are_reflected_in_prompt_and_kwargs() {
+    for (kwargs, expected, enabled) in [
+        (
+            json!({"enable_thinking": true, "reasoning_effort": "none"}),
+            "0.9",
+            true,
+        ),
+        (
+            json!({"thinking": false, "reasoning_effort": 10}),
+            "0.0",
+            false,
+        ),
+        (json!({"reasoning_effort": 0}), "0.0", false),
+        (json!({"reasoning_effort": null}), "0.9", true),
+        (json!({"reasoning_effort": 0.37}), "0.37", true),
+    ] {
+        let mut request = ChatRequest::for_test();
+        request.chat_options.template_kwargs = serde_json::from_value(kwargs).unwrap();
+        let rendered = renderer().render(&request).unwrap();
+        let prompt = render_symbolic_tokens(&rendered.prompt.into_token_ids().unwrap());
+        assert!(prompt.contains(&format!("Thinking effort level: {expected}<|end_message|>")));
+        assert_eq!(
+            rendered.effective_template_kwargs["enable_thinking"],
+            enabled
+        );
+        assert_eq!(
+            rendered.effective_template_kwargs["reasoning_effort"],
+            if enabled {
+                json!(expected.parse::<f64>().unwrap())
+            } else {
+                json!("none")
+            }
+        );
+    }
+}
+
 #[test]
 fn renders_numeric_reasoning_effort_template_kwarg() {
     let mut request = ChatRequest::for_test();
@@ -301,19 +339,17 @@ 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();
-        request
-            .chat_options
-            .template_kwargs
-            .insert("reasoning_effort".to_string(), value);
+fn ignores_unsupported_reasoning_effort_names() {
+    let mut request = ChatRequest::for_test();
+    request
+        .chat_options
+        .template_kwargs
+        .insert("reasoning_effort".to_string(), json!("invalid"));
 
-        assert_eq!(
-            render_symbolic_tokens(&render_token_ids(&request)),
-            "<|message_user|><|content_text|>test<|end_message|><|message_model|>"
-        );
-    }
+    assert_eq!(
+        render_symbolic_tokens(&render_token_ids(&request)),
+        "<|message_user|><|content_text|>test<|end_message|><|message_model|>"
+    );
 }
 
 #[test]

rust/src/chat/src/renderer/kimi_k3/encoding.rs

@@ -11,7 +11,9 @@ use serde_json::{Map, Value, json};
 use vllm_tokenizer::Tokenizer;
 
 use super::super::MediaPartSource;
+use crate::EffortValue;
 use crate::error::{Error, Result};
+use crate::reasoning::ReasoningControl;
 use crate::request::{
     ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatTool, ChatToolChoice,
 };
@@ -76,9 +78,10 @@ impl<'a> K3TokenWriter<'a> {
 pub(super) fn render_request_with_media_order(
     request: &ChatRequest,
     tokenizer: &dyn Tokenizer,
+    reasoning: &ReasoningControl,
 ) -> Result<(Vec<u32>, Vec<MediaPartSource>)> {
-    let thinking = thinking_enabled(request)?;
-    let thinking_effort = thinking.then(|| thinking_effort(request)).transpose()?;
+    let thinking = reasoning.is_enabled();
+    let thinking_effort = reasoning.effort().and_then(EffortValue::as_str);
     let tools = request_tools(request);
     let mut out = K3TokenWriter::new(tokenizer);
 
@@ -244,58 +247,26 @@ fn request_tools(request: &ChatRequest) -> &[ChatTool] {
     request.initial_tools()
 }
 
-fn thinking_enabled(request: &ChatRequest) -> Result<bool> {
-    if let Some(thinking) = request.parse_template_bool("thinking")? {
-        return Ok(thinking);
-    }
-    if let Some(enable_thinking) = request.parse_template_bool("enable_thinking")? {
-        return Ok(enable_thinking);
-    }
-    if let Some(reasoning_effort) = request.chat_options.reasoning_effort {
-        return Ok(reasoning_effort != crate::request::ReasoningEffort::None);
-    }
-    Ok(request
-        .chat_options
-        .template_kwargs
-        .get("reasoning_effort")
-        .and_then(Value::as_str)
-        != Some("none"))
-}
-
-fn thinking_effort(request: &ChatRequest) -> Result<String> {
-    let effort = if let Some(value) = request.chat_options.template_kwargs.get("thinking_effort") {
-        value.as_str().ok_or_else(|| {
-            Error::ChatTemplate(format!(
-                "template kwarg `thinking_effort` must be a string, got {value}"
-            ))
-        })?
-    } else if let Some(effort) = request.chat_options.reasoning_effort {
-        if effort == crate::request::ReasoningEffort::None {
-            DEFAULT_THINKING_EFFORT
-        } else {
-            effort.as_str()
-        }
-    } else if let Some(value) = request.chat_options.template_kwargs.get("reasoning_effort") {
+/// Resolve standard reasoning controls and validate K3's supported effort grades.
+pub(super) fn resolve_reasoning(
+    request: &ChatRequest,
+    defaults: &HashMap<String, Value>,
+) -> Result<ReasoningControl> {
+    let reasoning = ReasoningControl::resolve(request, defaults)?
+        .fallback(ReasoningControl::enabled(DEFAULT_THINKING_EFFORT));
+    if let Some(value) = reasoning.effort() {
         let effort = value.as_str().ok_or_else(|| {
-            Error::ChatTemplate(format!(
-                "template kwarg `reasoning_effort` must be a string, got {value}"
+            Error::InvalidReasoningEffort(format!(
+                "Kimi K3 reasoning_effort must be a string, got {value}"
             ))
         })?;
-        if effort == "none" {
-            DEFAULT_THINKING_EFFORT
-        } else {
-            effort
+        if !VALID_THINKING_EFFORTS.contains(&effort) {
+            return Err(Error::InvalidReasoningEffort(format!(
+                "unsupported reasoning_effort={effort:?}; supported values are `low`, `high`, and `max`"
+            )));
         }
-    } else {
-        DEFAULT_THINKING_EFFORT
-    };
-
-    if !VALID_THINKING_EFFORTS.contains(&effort) {
-        return Err(Error::ChatTemplate(format!(
-            "unsupported thinking_effort={effort:?}; supported values are `low`, `high`, and `max`"
-        )));
     }
-    Ok(effort.to_string())
+    Ok(reasoning)
 }
 
 fn content_is_empty(content: &ChatContent) -> bool {

rust/src/chat/src/renderer/kimi_k3/mod.rs

@@ -7,10 +7,13 @@ mod encoding;
 #[cfg(test)]
 mod tests;
 
+use std::collections::HashMap;
+
+use serde_json::Value;
 use vllm_text::Prompt;
 use vllm_text::tokenizer::DynTokenizer;
 
-use super::{ChatRenderer, RenderedPrompt, request_template_kwargs};
+use super::{ChatRenderer, RenderedPrompt};
 use crate::Result;
 use crate::request::ChatRequest;
 
@@ -18,25 +21,42 @@ use crate::request::ChatRequest;
 #[derive(Clone)]
 pub struct KimiK3ChatRenderer {
     tokenizer: DynTokenizer,
+    default_template_kwargs: HashMap<String, Value>,
 }
 
 impl KimiK3ChatRenderer {
     /// Create a Kimi K3 renderer.
-    pub fn new(tokenizer: DynTokenizer) -> Self {
-        Self { tokenizer }
+    pub fn new(tokenizer: DynTokenizer, default_template_kwargs: HashMap<String, Value>) -> Self {
+        Self {
+            tokenizer,
+            default_template_kwargs,
+        }
     }
 }
 
 impl ChatRenderer for KimiK3ChatRenderer {
     fn render(&self, request: &ChatRequest) -> Result<RenderedPrompt> {
         request.validate()?;
-        let (token_ids, media_order) =
-            encoding::render_request_with_media_order(request, self.tokenizer.as_ref())?;
+        let reasoning = encoding::resolve_reasoning(request, &self.default_template_kwargs)?;
+        let mut effective_template_kwargs = reasoning.template_kwargs(request);
+        // Export K3's native effort key from the resolved standard controls.
+        // Input `thinking_effort` is ignored; clear it when reasoning is disabled.
+        if let Some(effort) = reasoning.effort() {
+            effective_template_kwargs
+                .insert("thinking_effort".to_string(), serde_json::json!(effort));
+        } else {
+            effective_template_kwargs.remove("thinking_effort");
+        }
+        let (token_ids, media_order) = encoding::render_request_with_media_order(
+            request,
+            self.tokenizer.as_ref(),
+            &reasoning,
+        )?;
 
         Ok(RenderedPrompt {
             prompt: Prompt::TokenIds(token_ids),
             media_order: Some(media_order),
-            effective_template_kwargs: request_template_kwargs(request),
+            effective_template_kwargs,
         })
     }
 }

rust/src/chat/src/renderer/kimi_k3/tests.rs

@@ -15,11 +15,10 @@ use vllm_tokenizer::test_utils::TestTokenizer;
 
 use super::KimiK3ChatRenderer;
 use crate::ChatRenderer;
+use crate::EffortValue;
 use crate::renderer::kimi_k3::encoding::{CLOSE, END_OF_MSG, IMAGE_PLACEHOLDER, OPEN, SEP};
 use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
-use crate::request::{
-    ChatContentPart, ChatMessage, ChatTool, GenerationPromptMode, ReasoningEffort,
-};
+use crate::request::{ChatContentPart, ChatMessage, ChatTool, GenerationPromptMode};
 use crate::{AssistantContentBlock, AssistantToolCall};
 
 const OPEN_ID: u32 = 256;
@@ -37,8 +36,70 @@ fn test_tokenizer() -> TestTokenizer {
         .with_special_token(IMAGE_PLACEHOLDER, MEDIA_ID)
 }
 
+#[test]
+fn standard_effort_controls_prompt_and_exported_native_effort() {
+    let tokenizer = Arc::new(test_tokenizer());
+    for (kwargs, typed, defaults, effort) in [
+        (json!({}), None, json!({}), "max"),
+        (
+            json!({}),
+            Some(EffortValue::from("low")),
+            json!({"reasoning_effort": "high"}),
+            "low",
+        ),
+        (
+            json!({"enable_thinking": true, "reasoning_effort": "none"}),
+            None,
+            json!({"reasoning_effort": "low"}),
+            "low",
+        ),
+        (
+            json!({}),
+            Some(EffortValue::from("none")),
+            json!({"reasoning_effort": "low"}),
+            "none",
+        ),
+        (
+            json!({"thinking": true}),
+            Some(EffortValue::from("none")),
+            json!({}),
+            "max",
+        ),
+    ] {
+        let renderer =
+            KimiK3ChatRenderer::new(tokenizer.clone(), serde_json::from_value(defaults).unwrap());
+        let mut request = crate::ChatRequest::for_test();
+        request.chat_options.reasoning_effort = typed;
+        request.chat_options.template_kwargs = serde_json::from_value(kwargs).unwrap();
+        let rendered = renderer.render(&request).unwrap();
+        let prompt = tokenizer.decode(&rendered.prompt.into_token_ids().unwrap(), false).unwrap();
+        let enabled = effort != "none";
+        assert_eq!(prompt.ends_with("<|open|>think<|sep|>"), enabled);
+        assert_eq!(
+            rendered.effective_template_kwargs["reasoning_effort"],
+            effort
+        );
+        assert_eq!(
+            rendered.effective_template_kwargs["enable_thinking"],
+            enabled
+        );
+        if enabled {
+            assert!(prompt.contains(&format!("thinking_effort={effort}")));
+            assert_eq!(
+                rendered.effective_template_kwargs["thinking_effort"],
+                effort
+            );
+        } else {
+            assert!(!rendered.effective_template_kwargs.contains_key("thinking_effort"));
+        }
+    }
+}
+
 fn render_token_ids(request: &crate::request::ChatRequest, tokenizer: DynTokenizer) -> Vec<u32> {
-    let prompt = KimiK3ChatRenderer::new(tokenizer).render(request).unwrap().prompt;
+    let prompt = KimiK3ChatRenderer::new(tokenizer, Default::default())
+        .render(request)
+        .unwrap()
+        .prompt;
     let Prompt::TokenIds(token_ids) = prompt else {
         panic!("kimi k3 renderer should return token IDs")
     };
@@ -67,7 +128,11 @@ fn kimi_k3_fixture_options() -> FixtureRequestOptions {
 
 fn assert_golden(name: &str) {
     let input_name = format!("{name}_input.json");
-    let request = fixture_chat_request(&fixture_path(&input_name), kimi_k3_fixture_options());
+    let mut request = fixture_chat_request(&fixture_path(&input_name), kimi_k3_fixture_options());
+    // Reference fixtures use encoder-native keys; Rust accepts standard effort input.
+    if let Some(effort) = request.chat_options.template_kwargs.remove("thinking_effort") {
+        request.chat_options.reasoning_effort = Some(EffortValue::try_from(effort).unwrap());
+    }
     let rendered = render_request(&request);
     expect_file![format!("fixtures/{name}_output.txt")].assert_eq(&rendered);
 }
@@ -310,7 +375,7 @@ fn media_order_follows_reordered_tool_results() {
         ),
     ];
     request.chat_options.generation_prompt_mode = GenerationPromptMode::NoGenerationPrompt;
-    let renderer = KimiK3ChatRenderer::new(Arc::new(test_tokenizer()));
+    let renderer = KimiK3ChatRenderer::new(Arc::new(test_tokenizer()), Default::default());
 
     let rendered = renderer.render(&request).unwrap();
     let media_order = rendered.media_order.unwrap();
@@ -371,19 +436,18 @@ fn translates_standard_thinking_kwargs() {
 }
 
 #[test]
-fn native_k3_kwargs_take_precedence() {
+fn thinking_alias_takes_precedence() {
     let mut request = crate::request::ChatRequest::for_test();
     request.chat_options.template_kwargs.extend([
         ("thinking".to_string(), json!(true)),
         ("enable_thinking".to_string(), json!(false)),
-        ("thinking_effort".to_string(), json!("low")),
         ("reasoning_effort".to_string(), json!("high")),
     ]);
 
     let rendered = render_request(&request);
 
-    assert!(rendered.contains("thinking_effort=low"));
-    assert!(!rendered.contains("thinking_effort=high"));
+    assert!(rendered.contains("thinking_effort=high"));
+    assert!(rendered.ends_with("<|open|>think<|sep|>"));
 }
 
 #[test]
@@ -431,7 +495,7 @@ fn enable_thinking_true_overrides_standard_none() {
 #[test]
 fn typed_none_disables_thinking() {
     let mut request = crate::request::ChatRequest::for_test();
-    request.chat_options.reasoning_effort = Some(ReasoningEffort::None);
+    request.chat_options.reasoning_effort = Some(EffortValue::from("none"));
 
     let rendered = render_request(&request);
 
@@ -440,20 +504,17 @@ fn typed_none_disables_thinking() {
 }
 
 #[test]
-fn rejects_removed_medium_thinking_effort() {
+fn rejects_unsupported_standard_reasoning_effort() {
     let mut request = crate::request::ChatRequest::for_test();
-    request
-        .chat_options
-        .template_kwargs
-        .insert("thinking_effort".to_string(), json!("medium"));
+    request.chat_options.reasoning_effort = Some(EffortValue::from("medium"));
 
-    let error = KimiK3ChatRenderer::new(Arc::new(test_tokenizer()))
+    let error = KimiK3ChatRenderer::new(Arc::new(test_tokenizer()), Default::default())
         .render(&request)
         .unwrap_err();
 
     expect![[r#"
-        ChatTemplate(
-            "unsupported thinking_effort=\"medium\"; supported values are `low`, `high`, and `max`",
+        InvalidReasoningEffort(
+            "unsupported reasoning_effort=\"medium\"; supported values are `low`, `high`, and `max`",
         )
     "#]]
     .assert_debug_eq(&error);

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

@@ -4,11 +4,11 @@
 use std::collections::HashMap;
 use std::sync::Arc;
 
-use serde_json::{Value, json};
+use serde_json::Value;
 use vllm_text::Prompt;
 
 use crate::error::Result;
-use crate::request::{ChatRequest, ReasoningEffort};
+use crate::request::ChatRequest;
 
 mod deepseek;
 pub mod deepseek_v32;
@@ -62,33 +62,3 @@ pub trait ChatRenderer: Send + Sync {
 
 /// Shared trait-object form of [`ChatRenderer`].
 pub type DynChatRenderer = Arc<dyn ChatRenderer>;
-
-/// Extract the effective chat-template kwargs visible to the renderer from the request,
-/// using the provided defaults as the base.
-pub(crate) fn effective_template_kwargs(
-    default_template_kwargs: &HashMap<String, Value>,
-    request: &ChatRequest,
-) -> HashMap<String, Value> {
-    let mut kwargs = default_template_kwargs.clone();
-    kwargs.extend(request.chat_options.template_kwargs.clone());
-
-    if let Some(reasoning_effort) = request.chat_options.reasoning_effort {
-        kwargs.insert(
-            "reasoning_effort".to_string(),
-            Value::String(reasoning_effort.as_str().to_string()),
-        );
-        if !request.chat_options.template_kwargs.contains_key("enable_thinking") {
-            kwargs.insert(
-                "enable_thinking".to_string(),
-                json!(reasoning_effort != ReasoningEffort::None),
-            );
-        }
-    }
-
-    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> {
-    effective_template_kwargs(&HashMap::new(), request)
-}

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

@@ -8,10 +8,11 @@ use std::path::Path;
 use serde::Deserialize;
 use serde_json::Value;
 
+use crate::EffortValue;
 use crate::event::{AssistantContentBlock, AssistantToolCall};
 use crate::request::{
     ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatTool, ChatToolChoice,
-    GenerationPromptMode, ReasoningEffort, ResolvedToolContext,
+    GenerationPromptMode, ResolvedToolContext,
 };
 
 /// Options for constructing a [`ChatRequest`] from a fixture file.
@@ -47,7 +48,7 @@ pub(crate) struct FixtureRequest {
     tools: Vec<FixtureTool>,
     messages: Vec<FixtureMessage>,
     add_generation_prompt: Option<bool>,
-    reasoning_effort: Option<ReasoningEffort>,
+    reasoning_effort: Option<EffortValue>,
     /// Standard response format passed to model-specific renderers.
     #[serde(default)]
     response_format: Option<Value>,

rust/src/chat/src/request.rs

@@ -11,9 +11,9 @@ pub use vllm_parser::tool::Tool as ChatTool;
 use vllm_text::TextDecodeOptions;
 pub use vllm_text::{PromptTruncation, SamplingParams};
 
-use crate::AssistantMessageExt;
 use crate::error::{Error, Result};
 use crate::event::{AssistantContentBlock, AssistantMessage};
+use crate::{AssistantMessageExt, EffortValue};
 
 /// Role label for one text-only chat message.
 #[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
@@ -362,33 +362,6 @@ pub enum GenerationPromptMode {
     NoGenerationPrompt,
 }
 
-/// Effort level for reasoning models.
-#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
-#[serde(rename_all = "lowercase")]
-pub enum ReasoningEffort {
-    None,
-    Minimal,
-    Low,
-    Medium,
-    High,
-    XHigh,
-    Max,
-}
-
-impl ReasoningEffort {
-    pub fn as_str(self) -> &'static str {
-        match self {
-            Self::None => "none",
-            Self::Minimal => "minimal",
-            Self::Low => "low",
-            Self::Medium => "medium",
-            Self::High => "high",
-            Self::XHigh => "xhigh",
-            Self::Max => "max",
-        }
-    }
-}
-
 /// Chat-template-related request options.
 ///
 /// These are the small subset of chat controls that currently affect prompt
@@ -404,8 +377,10 @@ pub struct ChatOptions {
     /// used instead of the model's default chat template.
     pub chat_template: Option<String>,
 
-    /// Effort level exposed to chat templates for reasoning models.
-    pub reasoning_effort: Option<ReasoningEffort>,
+    /// Model-specific reasoning effort, typically `none`, `minimal`, `low`,
+    /// `medium`, `high`, `xhigh`, or `max`. Supported names and numeric ranges
+    /// are validated by the selected renderer or HF template.
+    pub reasoning_effort: Option<EffortValue>,
 
     /// Standard response format available to model-specific renderers.
     #[serde(default)]
@@ -673,29 +648,6 @@ impl ChatRequest {
         self.tool_context.parsing_enabled()
     }
 
-    /// Return the request-level thinking toggle when explicitly requested.
-    ///
-    /// We currently accept the two request kwargs `thinking` and
-    /// `enable_thinking`. Both must be booleans when present. If both are
-    /// present, they must have the same value. If neither key is provided,
-    /// return `None`.
-    pub(crate) 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)),
-        }
-    }
-
     pub(crate) fn parse_template_bool(&self, key: &str) -> Result<Option<bool>> {
         match self.chat_options.template_kwargs.get(key) {
             None => Ok(None),
@@ -726,7 +678,7 @@ mod tests {
     use serde_json::{json, to_value};
 
     use super::{
-        ChatContent, ChatContentPart, ChatMessage, ChatRequest, ChatRole, ChatTool, ChatToolChoice,
+        ChatContent, ChatContentPart, ChatMessage, ChatRole, ChatTool, ChatToolChoice,
         ResolvedToolContext,
     };
     use crate::Error;
@@ -919,56 +871,4 @@ mod tests {
             Error::ToolChoiceFunctionNotFound { name } if name == "missing"
         ));
     }
-
-    #[test]
-    fn enable_thinking_is_none_when_no_kwargs_are_present() {
-        let request = ChatRequest::for_test();
-        assert_eq!(request.enable_thinking().unwrap(), None);
-    }
-
-    #[test]
-    fn enable_thinking_accepts_matching_duplicate_kwargs() {
-        let mut request = ChatRequest::for_test();
-        request.chat_options.template_kwargs.insert("thinking".to_string(), json!(true));
-        request
-            .chat_options
-            .template_kwargs
-            .insert("enable_thinking".to_string(), json!(true));
-
-        assert_eq!(request.enable_thinking().unwrap(), Some(true));
-    }
-
-    #[test]
-    fn enable_thinking_rejects_non_boolean_kwargs() {
-        let mut request = ChatRequest::for_test();
-        request
-            .chat_options
-            .template_kwargs
-            .insert("thinking".to_string(), json!("yes"));
-
-        assert!(matches!(
-            request.enable_thinking(),
-            Err(Error::ChatTemplate(message))
-                if message.contains("`thinking` must be a boolean")
-        ));
-    }
-
-    #[test]
-    fn enable_thinking_rejects_conflicting_duplicate_kwargs() {
-        let mut request = ChatRequest::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.enable_thinking(),
-            Err(Error::ChatTemplate(message))
-                if message.contains("`thinking` and `enable_thinking` must match")
-        ));
-    }
 }

rust/src/engine-core-client/src/protocol/mod.rs

@@ -30,6 +30,7 @@ pub mod multimodal;
 pub mod output;
 pub mod request;
 pub mod sampling;
+mod serde_utils;
 pub mod stats;
 pub mod structured_outputs;
 pub mod tensor;

rust/src/engine-core-client/src/protocol/output.rs

@@ -9,7 +9,9 @@ use serde::{Deserialize, Deserializer, Serialize, Serializer};
 use serde_default::DefaultFromSerde;
 use serde_repr::{Deserialize_repr, Serialize_repr};
 use serde_tuple::{Deserialize_tuple, Serialize_tuple};
+use serde_with::serde_as;
 
+use super::serde_utils::AllowTrailingFields;
 use super::utility::UtilityOutput;
 use crate::error::{Error, Result, ext_value_decode};
 use crate::protocol::logprobs::MaybeWireLogprobs;
@@ -156,11 +158,14 @@ impl EngineCoreOutput {
 ///
 /// Original Python definition:
 /// <https://github.com/vllm-project/vllm/blob/f22d6e026798a74e6542a52ef776c054f2de572a/vllm/v1/engine/__init__.py#L186-L214>
+#[serde_as]
 #[derive(Debug, Clone, PartialEq, Serialize_tuple, Deserialize_tuple, DefaultFromSerde)]
 struct WireEngineCoreOutputs {
     #[serde(default)]
     engine_index: u32,
     /// Outputs grouped for this client in the current engine tick.
+    // Match msgspec's append-only schema evolution, including Omni extensions.
+    #[serde_as(deserialize_as = "Vec<AllowTrailingFields>")]
     #[serde(default)]
     outputs: Vec<EngineCoreOutput>,
     #[serde(default)]
@@ -404,6 +409,61 @@ mod tests {
         );
     }
 
+    #[test]
+    fn engine_core_outputs_ignore_appended_fields() {
+        let output = EngineCoreOutput {
+            request_id: "req-1".into(),
+            new_token_ids: vec![42],
+            finish_reason: Some(EngineCoreFinishReason::Length),
+            ..Default::default()
+        };
+        let mut fields = crate::protocol::decode_value(&encode_msgpack(&output).unwrap())
+            .unwrap()
+            .as_array()
+            .unwrap()
+            .clone();
+        // Omni appends multimodal_output, is_segment_finished, and
+        // new_prompt_len_snapshot. Unknown payloads may contain nested values.
+        fields.extend([
+            OpaqueValue::Map(vec![(
+                "audio".into(),
+                OpaqueValue::Array(vec![OpaqueValue::Ext(1, vec![0, 1])]),
+            )]),
+            false.into(),
+            OpaqueValue::Nil,
+        ]);
+        let wire = (0, vec![OpaqueValue::Array(fields); 2]);
+        let frames = [Bytes::from(encode_msgpack(&wire).unwrap())];
+        let decoded = decode_engine_core_outputs(&frames).unwrap();
+        // A second output also checks that the first output's tail was consumed.
+        assert_eq!(decoded.as_request_batch().unwrap().outputs, vec![output; 2]);
+    }
+
+    #[test]
+    fn engine_core_output_requires_valid_known_fields() {
+        for fields in [
+            vec![],
+            vec!["req-1".into()],
+            vec!["req-1".into(), OpaqueValue::Nil],
+            vec!["req-1".into(), OpaqueValue::Array(vec![]), false.into()],
+        ] {
+            let wire = (0, vec![OpaqueValue::Array(fields)]);
+            let frames = [Bytes::from(encode_msgpack(&wire).unwrap())];
+            assert!(decode_engine_core_outputs(&frames).is_err());
+        }
+        let wire = (0, vec![("req-1", vec![42_u32])]);
+        let frames = [Bytes::from(encode_msgpack(&wire).unwrap())];
+        let decoded = decode_engine_core_outputs(&frames).unwrap();
+        assert_eq!(
+            decoded.as_request_batch().unwrap().outputs,
+            vec![EngineCoreOutput {
+                request_id: "req-1".into(),
+                new_token_ids: vec![42],
+                ..Default::default()
+            }]
+        );
+    }
+
     #[test]
     fn engine_core_outputs_classify_request_batch() {
         let outputs = WireEngineCoreOutputs {

rust/src/engine-core-client/src/protocol/serde_utils.rs

@@ -0,0 +1,42 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::fmt;
+use std::marker::PhantomData;
+
+use serde::de::{IgnoredAny, SeqAccess, Visitor, value::SeqAccessDeserializer};
+use serde::{Deserialize, Deserializer};
+use serde_with::DeserializeAs;
+
+/// Decode an array-like struct using its derived field types and defaults,
+/// skipping appended fields as Python's msgspec does. Opt in only for schemas
+/// whose extensions preserve the known field order and semantics.
+pub(super) struct AllowTrailingFields;
+
+impl<'de, T: Deserialize<'de>> DeserializeAs<'de, T> for AllowTrailingFields {
+    fn deserialize_as<D>(deserializer: D) -> Result<T, D::Error>
+    where
+        D: Deserializer<'de>,
+    {
+        struct ArrayVisitor<T>(PhantomData<T>);
+
+        impl<'de, T: Deserialize<'de>> Visitor<'de> for ArrayVisitor<T> {
+            type Value = T;
+
+            fn expecting(&self, formatter: &mut fmt::Formatter<'_>) -> fmt::Result {
+                formatter.write_str("an array-like struct")
+            }
+
+            fn visit_seq<A>(self, mut seq: A) -> Result<T, A::Error>
+            where
+                A: SeqAccess<'de>,
+            {
+                let value = T::deserialize(SeqAccessDeserializer::new(&mut seq))?;
+                while seq.next_element::<IgnoredAny>()?.is_some() {}
+                Ok(value)
+            }
+        }
+
+        deserializer.deserialize_seq(ArrayVisitor(PhantomData))
+    }
+}

rust/src/engine-core-client/src/tests/client.rs

@@ -2641,6 +2641,7 @@ fn python_msgpack_fixtures_match_rust_encoding() {
     let multi_connector_stats_hex =
         lines.next().expect("missing MultiConnector stats fixture line");
     let ready_response_hex = lines.next().expect("missing ready response fixture line");
+    let extended_outputs_hex = lines.next().expect("missing extended outputs fixture line");
 
     let request_bytes = hex::decode(request_hex).unwrap();
     let multimodal_request_bytes = hex::decode(multimodal_request_hex).unwrap();
@@ -2722,6 +2723,14 @@ fn python_msgpack_fixtures_match_rust_encoding() {
     ));
 
     let decoded_outputs: EngineCoreOutputs = rmp_serde::from_slice(&outputs_bytes).unwrap();
+    // Match msgspec's base-schema result for the same extended Python message.
+    let extended_frames = [bytes::Bytes::from(
+        hex::decode(extended_outputs_hex).unwrap(),
+    )];
+    assert_eq!(
+        decode_engine_core_outputs(&extended_frames).unwrap(),
+        decoded_outputs
+    );
     expect_test::expect![[r#"
         RequestBatch(
             RequestBatchOutputs {

rust/src/engine-core-client/src/tests/python_compat.py

@@ -101,6 +101,14 @@ class EngineCoreOutput(
     num_nans_in_logits: int = 0
     mm_cache_miss_hashes: list[str] | None = None
     new_sampling_mask: object | None = None
+    spec_decode_metrics: object | None = None
+
+
+class ExtendedEngineCoreOutput(EngineCoreOutput):
+    # Match Omni's append-only schema without requiring Omni or torch.
+    multimodal_output: dict[str, object] | None = None
+    is_segment_finished: bool = False
+    new_prompt_len_snapshot: int | None = None
 
 
 class EngineCoreOutputs(
@@ -207,6 +215,21 @@ outputs = EngineCoreOutputs(
     finished_requests={"req-1"},
 )
 
+extended_outputs = EngineCoreOutputs(
+    outputs=[
+        ExtendedEngineCoreOutput(
+            **msgspec.structs.asdict(outputs.outputs[0]),
+            multimodal_output={"audio": b"\x00\x01"},
+            is_segment_finished=True,
+            new_prompt_len_snapshot=12,
+        )
+    ],
+    finished_requests=outputs.finished_requests,
+)
+extended_outputs_bytes = msgspec.msgpack.encode(extended_outputs)
+# The ordinary frontend's schema ignores even non-default extension values.
+assert msgspec.msgpack.decode(extended_outputs_bytes, type=EngineCoreOutputs) == outputs
+
 sampling_mask_wire = [
     [
         "<i4",
@@ -509,3 +532,4 @@ print(msgspec.msgpack.encode(nixl_stats).hex())
 print(msgspec.msgpack.encode(mooncake_stats).hex())
 print(msgspec.msgpack.encode(multi_connector_stats).hex())
 print(msgspec.msgpack.encode(ready_response).hex())
+print(extended_outputs_bytes.hex())

rust/src/server/src/error.rs

@@ -175,16 +175,22 @@ mod tests {
     }
 
     #[test]
-    fn invalid_reasoning_effort_maps_to_invalid_request() {
-        let error = vllm_chat::Error::InvalidReasoningEffort(
-            "DeepSeek V4.1 reasoning_effort must be within [1, 100]".to_string(),
-        );
-        let api_error = chat_submit_error("failed to submit chat request", error);
-        assert_eq!(api_error.status_code(), StatusCode::BAD_REQUEST);
-        assert_eq!(
-            api_error.to_error_response().error.error_type,
-            "invalid_request_error"
-        );
+    fn invalid_reasoning_parameters_map_to_invalid_request() {
+        for error in [
+            vllm_chat::Error::InvalidReasoningEffort(
+                "DeepSeek V4.1 reasoning_effort must be within [1, 100]".to_string(),
+            ),
+            vllm_chat::Error::InvalidReasoningControl {
+                message: "template kwarg `thinking` must be a boolean".to_string(),
+            },
+        ] {
+            let api_error = chat_submit_error("failed to submit chat request", error);
+            assert_eq!(api_error.status_code(), StatusCode::BAD_REQUEST);
+            assert_eq!(
+                api_error.to_error_response().error.error_type,
+                "invalid_request_error"
+            );
+        }
     }
 
     #[test]

rust/src/server/src/routes/openai/chat_completions/convert.rs

@@ -180,7 +180,7 @@ pub(super) fn prepare_chat_request(
         chat_options: ChatOptions {
             generation_prompt_mode,
             chat_template: request.chat_template,
-            reasoning_effort: request.reasoning_effort,
+            reasoning_effort: request.reasoning_effort.map(|effort| effort.as_str().into()),
             response_format,
             template_kwargs,
         },
@@ -488,6 +488,50 @@ mod tests {
         }
     }
 
+    #[test]
+    fn chat_http_reasoning_effort_preserves_omission_none_and_kwargs() {
+        for effort in [None, Some(json!(null))].into_iter().chain(
+            ["none", "minimal", "low", "medium", "high", "xhigh", "max"]
+                .map(|name| Some(json!(name))),
+        ) {
+            let kwargs = json!({"reasoning_effort": 37, "thinking": true});
+            let mut value = json!({
+                "messages": [{"role": "user", "content": "hello"}],
+                "chat_template_kwargs": kwargs,
+            });
+            if let Some(effort) = &effort {
+                value["reasoning_effort"] = effort.clone();
+            }
+            let request: ChatCompletionRequest = serde_json::from_value(value).unwrap();
+            let prepared = prepare_chat_request(
+                request,
+                &served(&["test-model"]),
+                request_context(&HeaderMap::new(), None),
+            )
+            .unwrap();
+            let options = prepared.chat_request.chat_options;
+            assert_eq!(
+                serde_json::to_value(options.reasoning_effort).unwrap(),
+                effort.unwrap_or(serde_json::Value::Null)
+            );
+            assert_eq!(
+                serde_json::to_value(options.template_kwargs).unwrap(),
+                kwargs
+            );
+        }
+    }
+
+    #[test]
+    fn chat_http_reasoning_effort_rejects_model_extensions_at_top_level() {
+        for effort in [json!(37), json!("custom")] {
+            let request = json!({
+                "messages": [{"role": "user", "content": "hello"}],
+                "reasoning_effort": effort,
+            });
+            assert!(serde_json::from_value::<ChatCompletionRequest>(request).is_err());
+        }
+    }
+
     #[test]
     fn chat_http_request_defaults_missing_or_null_model() {
         for model in [None, Some(serde_json::Value::Null)] {
@@ -747,7 +791,7 @@ mod tests {
         );
 
         let tokenizer = Arc::new(TestTokenizer::new());
-        let prompt = KimiK3ChatRenderer::new(tokenizer.clone())
+        let prompt = KimiK3ChatRenderer::new(tokenizer.clone(), Default::default())
             .render(&prepared.chat_request)
             .expect("Kimi K3 rendering succeeds")
             .prompt;

rust/src/server/src/routes/openai/chat_completions/types.rs

@@ -9,15 +9,15 @@ use serde::{Deserialize, Serialize};
 use serde_json::Value;
 use serde_with::SerializeDisplay;
 use validator::Validate;
-use vllm_chat::ReasoningEffort;
 use vllm_engine_core_client::protocol::sampling::RepetitionDetectionParams;
 use vllm_text::TruncationSide;
 
 use crate::routes::openai::utils::structured_outputs::ResponseFormat;
 use crate::routes::openai::utils::types::{
-    ChatLogProbs, ChatMessage, Normalizable, PromptLogprobs, StreamOptions, StreamResponseEnvelope,
-    StringOrArray, Tool, ToolCall, ToolCallDelta, ToolChoice, Usage, default_true,
-    deserialize_request_top_k, validate_messages, validate_stop, validate_top_p_value,
+    ChatLogProbs, ChatMessage, Normalizable, PromptLogprobs, ReasoningEffort, StreamOptions,
+    StreamResponseEnvelope, StringOrArray, Tool, ToolCall, ToolCallDelta, ToolChoice, Usage,
+    default_true, deserialize_request_top_k, validate_messages, validate_stop,
+    validate_top_p_value,
 };
 
 /// vLLM-compatible request type for the Chat Completions API.

rust/src/server/src/routes/openai/chat_completions/validate.rs

@@ -125,8 +125,8 @@ fn validate_function_tools(tools: &[Tool], param: &'static str) -> Result<(), Ap
 mod tests {
     use std::collections::HashMap;
 
+    use crate::routes::openai::utils::types::ReasoningEffort;
     use serde_json::json;
-    use vllm_chat::ReasoningEffort;
 
     use super::validate_request_compat;
     use crate::routes::openai::chat_completions::types::ChatCompletionRequest;

rust/src/server/src/routes/openai/utils/types.rs

@@ -39,6 +39,36 @@ where
     }
 }
 
+/// Effort level for reasoning models.
+///
+/// Fixed OpenAI HTTP request grades. Request conversion maps these names to
+/// the model-independent `vllm_chat::EffortValue` before renderer validation.
+#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
+#[serde(rename_all = "lowercase")]
+pub enum ReasoningEffort {
+    None,
+    Minimal,
+    Low,
+    Medium,
+    High,
+    XHigh,
+    Max,
+}
+
+impl ReasoningEffort {
+    pub fn as_str(self) -> &'static str {
+        match self {
+            Self::None => "none",
+            Self::Minimal => "minimal",
+            Self::Low => "low",
+            Self::Medium => "medium",
+            Self::High => "high",
+            Self::XHigh => "xhigh",
+            Self::Max => "max",
+        }
+    }
+}
+
 // ============================================================================
 // String/Array Utilities
 // ============================================================================

tests/kernels/moe/modular_kernel_tools/common.py

@@ -15,6 +15,7 @@ from tests.kernels.quantization.nvfp4_utils import (
     dequantize_nvfp4_to_dtype,
 )
 from tests.kernels.utils import torch_experts
+from vllm._aiter_ops import rocm_aiter_ops
 from vllm.config import VllmConfig
 from vllm.distributed import (
     get_dp_group,
@@ -33,7 +34,9 @@ from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEQuantConfig,
     RoutingMethodType,
 )
+from vllm.model_executor.layers.quantization.utils.fp8_utils import is_fp8
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
+    QuantKey,
     kFp8Dynamic128Sym,
     kFp8DynamicTensorSym,
     kFp8DynamicTokenSym,
@@ -165,31 +168,34 @@ class Config:
 
         return vllm_config, env_dict
 
+    def fp8_quant_key_pair(self) -> tuple[QuantKey, QuantKey]:
+        """Derive the (weight_quant_key, activation_quant_key) pair an FP8
+        quant config of this shape corresponds to (either OCP or FNUZ FP8,
+        see ``current_platform.fp8_dtype()``)."""
+        if self.quant_block_shape is not None:
+            return kFp8Static128BlockSym, kFp8Dynamic128Sym
+        if self.is_per_out_ch_quant:
+            return (
+                kFp8StaticChannelSym,
+                kFp8DynamicTokenSym
+                if self.is_per_act_token_quant
+                else kFp8StaticTensorSym,
+            )
+        return (
+            kFp8StaticTensorSym,
+            kFp8DynamicTensorSym
+            if self.is_per_act_token_quant
+            else kFp8StaticTensorSym,
+        )
+
     def fe_supports_quant_scheme(self) -> bool:
         """Check if the fused experts class supports this quant config.
         See https://github.com/ROCm/aiter/issues/2419 for AITER gaps."""
         if self.quant_config is None or self.quant_dtype is None:
             return True
-        if self.quant_dtype != torch.float8_e4m3fn:
+        if not is_fp8(self.quant_dtype):
             return True
-        # Derive QuantKeys from test config
-        if self.quant_block_shape is not None:
-            w_key = kFp8Static128BlockSym
-            a_key = kFp8Dynamic128Sym
-        elif self.is_per_out_ch_quant:
-            w_key = kFp8StaticChannelSym
-            a_key = (
-                kFp8DynamicTokenSym
-                if self.is_per_act_token_quant
-                else kFp8StaticTensorSym
-            )
-        else:
-            w_key = kFp8StaticTensorSym
-            a_key = (
-                kFp8DynamicTensorSym
-                if self.is_per_act_token_quant
-                else kFp8StaticTensorSym
-            )
+        w_key, a_key = self.fp8_quant_key_pair()
         fe_cls = self.fused_experts_type
         if hasattr(fe_cls, "_supports_quant_scheme"):
             try:
@@ -199,10 +205,7 @@ class Config:
         return True
 
     def is_fp8_block_quantized(self):
-        return (
-            self.quant_dtype == torch.float8_e4m3fn
-            and self.quant_block_shape is not None
-        )
+        return is_fp8(self.quant_dtype) and self.quant_block_shape is not None
 
     def is_batched_prepare_finalize(self):
         info = prepare_finalize_info(self.prepare_finalize_type)
@@ -334,6 +337,11 @@ class Config:
             return False, "Needs Aiter, but Aiter not available."
         if self.needs_mori() and not has_mori():  # noqa: SIM103
             return False, "Needs MoRI, but MoRI not available."
+        if self.needs_mori() and not rocm_aiter_ops.is_fused_moe_enabled():
+            return False, (
+                "Mori requires AITER's fused-moe backend to be enabled "
+                "(VLLM_ROCM_USE_AITER=1 and VLLM_ROCM_USE_AITER_MOE=1)."
+            )
 
         try:
             if not self.fused_experts_type._supports_current_device():
@@ -370,7 +378,7 @@ class WeightTensors:
     def is_quantized(self) -> bool:
         # or w1_scale is not None?
         return (
-            self.w1.dtype == torch.float8_e4m3fn
+            is_fp8(self.w1.dtype)
             or self.w1.dtype == torch.uint8
             or self.w1.dtype == torch.int8
         )

tests/kernels/moe/modular_kernel_tools/mk_objects.py

@@ -86,14 +86,14 @@ MK_FUSED_EXPERT_TYPES: list[mk.FusedMoEExpertsModular] = []
 standard_format = mk.FusedMoEActivationFormat.Standard
 batched_format = mk.FusedMoEActivationFormat.BatchedExperts
 common_float_types: list[torch.dtype | str] = [
-    torch.float8_e4m3fn,
+    current_platform.fp8_dtype(),
     torch.bfloat16,
     torch.float16,
     torch.float32,
 ]
 common_float_and_int_types = common_float_types + [torch.int8]
 nvfp4_types = ["nvfp4"]
-fp8_types = [torch.float8_e4m3fn]
+fp8_types = [current_platform.fp8_dtype()]
 
 
 def register_prepare_and_finalize(
@@ -324,7 +324,9 @@ if has_aiter():
     register_experts(
         AiterExperts,
         standard_format,
-        fp8_types,
+        # AiterExperts also supports the fully-unquantized (None, None)
+        # scheme (see SUPPORTED_W_A in rocm_aiter_moe.py), not just fp8.
+        common_float_types,
         blocked_quantization_support=True,
         needs_aiter=True,
     )
@@ -397,35 +399,35 @@ MK_QUANT_CONFIGS: list[TestMoEQuantConfig | None] = [
     None,
     # per-channel / per-column weights and per-tensor activations
     TestMoEQuantConfig(
-        quant_dtype=torch.float8_e4m3fn,
+        quant_dtype=current_platform.fp8_dtype(),
         per_out_ch_quant=True,
         per_act_token_quant=False,
         block_shape=None,
     ),
     # per-channel / per-column weights and per-token activations
     TestMoEQuantConfig(
-        quant_dtype=torch.float8_e4m3fn,
+        quant_dtype=current_platform.fp8_dtype(),
         per_out_ch_quant=True,
         per_act_token_quant=True,
         block_shape=None,
     ),
     # per-tensor weights and per-tensor activations
     TestMoEQuantConfig(
-        quant_dtype=torch.float8_e4m3fn,
+        quant_dtype=current_platform.fp8_dtype(),
         per_out_ch_quant=False,
         per_act_token_quant=False,
         block_shape=None,
     ),
     # per-tensor weights and per-token activations
     TestMoEQuantConfig(
-        quant_dtype=torch.float8_e4m3fn,
+        quant_dtype=current_platform.fp8_dtype(),
         per_out_ch_quant=False,
         per_act_token_quant=True,
         block_shape=None,
     ),
     # block-quantized weights and 128 block per-token activations
     TestMoEQuantConfig(
-        quant_dtype=torch.float8_e4m3fn,
+        quant_dtype=current_platform.fp8_dtype(),
         per_out_ch_quant=False,
         per_act_token_quant=False,
         block_shape=[128, 128],

tests/kernels/moe/test_gpt_oss_triton_kernels.py

@@ -384,9 +384,8 @@ def test_routing_data_from_sparse_topk_parity(n_tokens, n_experts, topk):
 
     make_routing_data = gptoss_moe.make_routing_data
     routing_data_from_sparse_topk = gptoss_moe.routing_data_from_sparse_topk
-    use_legacy_triton_kernels = gptoss_moe.use_legacy_triton_kernels
 
-    if use_legacy_triton_kernels:
+    if gptoss_moe.triton_kernels_version == "3.5.1":
         pytest.skip("SparseMatrix path requires triton_kernels v3.6.0+")
 
     from triton_kernels.topk import topk as topk_fn

tests/kernels/moe/test_modular_kernel_combinations.py

@@ -37,6 +37,7 @@ from .modular_kernel_tools.parallel_utils import (
     ProcessGroupInfo,
     parallel_launch_with_config,
 )
+from .utils import check_accuracy
 
 has_any_multi_gpu_package = (
     has_deep_ep() or has_deep_gemm() or has_flashinfer_cutlass_fused_moe()
@@ -47,12 +48,6 @@ meets_multi_gpu_requirements = pytest.mark.skipif(
     reason="Requires deep_ep or deep_gemm or flashinfer packages",
 )
 
-if current_platform.is_fp8_fnuz():
-    pytest.skip(
-        "Tests in this file require float8_e4m3fn and platform does not support",
-        allow_module_level=True,
-    )
-
 
 def format_result(verbose, msg, ex=None):
     if ex is not None:
@@ -70,6 +65,28 @@ def format_result(verbose, msg, ex=None):
         print(".", end="")
 
 
+def assert_aiter_quant_scheme_case(config: Config) -> None:
+    """Make the AITER (weight_quant_key, activation_quant_key) pair this
+    config exercises explicit, instead of AiterExperts being reached only
+    indirectly through the general quant-config sweep.
+    See https://github.com/vllm-project/vllm/issues/54966."""
+    fe_cls = config.fused_experts_type
+    if getattr(fe_cls, "__name__", "") != "AiterExperts":
+        return
+
+    if config.quant_config is None:
+        w_key, a_key = None, None
+    else:
+        w_key, a_key = config.fp8_quant_key_pair()
+
+    assert fe_cls._supports_quant_scheme(w_key, a_key), (
+        f"AITER case (weight_key={w_key}, activation_key={a_key}) reached "
+        "the modular-kernel harness, but AiterExperts._supports_quant_scheme "
+        "does not declare it supported."
+    )
+    print(f"[AITER case] weight_key={w_key}, activation_key={a_key}")
+
+
 def rank_worker(
     pgi: ProcessGroupInfo,
     vllm_config: VllmConfig,
@@ -125,6 +142,25 @@ def rank_worker(
                 count -= 1
                 continue
 
+            # Skip unsupported: AITER x DeepEP-HT/Mori dispatch crashes
+            # with an illegal memory access at world_size>1 on gfx942.
+            # https://github.com/vllm-project/vllm/issues/57029
+            if (
+                config.world_size > 1
+                and getattr(config.fused_experts_type, "__name__", "") == "AiterExperts"
+                and getattr(config.prepare_finalize_type, "__name__", "")
+                in ("DeepEPHTPrepareAndFinalize", "MoriPrepareAndFinalize")
+            ):
+                print(
+                    f"Skipping[{pgi.rank}]: m={m}, topk={topk}"
+                    " (AITER x DeepEP-HT/Mori illegal memory access,"
+                    " https://github.com/vllm-project/vllm/issues/57029)"
+                )
+                count -= 1
+                continue
+
+            assert_aiter_quant_scheme_case(config)
+
             # modular kernel out
             mk_out = run_modular_kernel(pgi, vllm_config, config, weights, rank_tensors)
 
@@ -152,32 +188,7 @@ def rank_worker(
                 and config.quant_config is not None
             )
             if is_aiter_fp8:
-                diff = (ref_out - mk_out).abs()
-                n_total = diff.numel()
-                max_diff = diff.max().item()
-                n_exceed = int((diff > atol).sum().item())
-                pct_exceed = n_exceed / n_total * 100
-                # FP8 hw matmul vs f32 reference: up to ~4% of
-                # elements may exceed base tolerance, but max
-                # error should stay within 3x base tolerance.
-                max_pct_allowed = 5.0
-                relaxed_atol = atol * 4
-                print(
-                    f"[AITER FP8 precision] "
-                    f"max_diff={max_diff:.6f}, "
-                    f"exceed_atol={n_exceed}/{n_total} "
-                    f"({pct_exceed:.4f}%), "
-                    f"max_pct_allowed={max_pct_allowed}%, "
-                    f"relaxed_limit={relaxed_atol}"
-                )
-                assert pct_exceed <= max_pct_allowed, (
-                    f"AITER FP8: {pct_exceed:.2f}% elements exceed "
-                    f"atol={atol} (max allowed {max_pct_allowed}%)"
-                )
-                assert max_diff <= relaxed_atol, (
-                    f"AITER FP8: max_diff={max_diff:.6f} exceeds "
-                    f"relaxed limit {relaxed_atol}"
-                )
+                check_accuracy(ref_out, mk_out, atol=atol, rtol=rtol, percent=0.9)
             else:
                 torch.testing.assert_close(ref_out, mk_out, atol=atol, rtol=rtol)
             format_result(verbose, config.describe())

tests/test_config.py

@@ -1520,6 +1520,115 @@ def test_draft_model_enables_async_scheduling_by_default():
     assert cfg.scheduler_config.async_scheduling is True
 
 
+@pytest.mark.skip_global_cleanup
+@pytest.mark.parametrize("tp_size", [1, 2])
+@pytest.mark.parametrize("target_ep", [False, True], ids=["ep-off", "ep-on"])
+@pytest.mark.parametrize(
+    ("method", "draft_is_moe"),
+    [
+        pytest.param("draft_model", False, id="dense-draft"),
+        pytest.param("eagle", False, id="eagle"),
+        pytest.param("eagle3", False, id="eagle3"),
+        pytest.param("draft_model", True, id="moe-draft"),
+        pytest.param("mtp", True, id="mtp"),
+        pytest.param("dspark", True, id="dspark"),
+    ],
+)
+def test_draft_inherits_ep_only_for_moe(
+    tmp_path: Path,
+    tp_size: int,
+    target_ep: bool,
+    method: str,
+    draft_is_moe: bool,
+):
+    """Validate final draft configs without loading weights or mocking validation."""
+    from transformers import LlamaConfig, MixtralConfig
+
+    from vllm.transformers_utils.configs.deepseek_v4 import DeepseekV4Config
+
+    common = dict(
+        hidden_size=128,
+        intermediate_size=256,
+        num_hidden_layers=2,
+        num_attention_heads=4,
+        num_key_value_heads=2,
+        vocab_size=128,
+        max_position_embeddings=2048,
+    )
+    if method in ("mtp", "dspark"):
+        target_hf_config = DeepseekV4Config(
+            architectures=["DeepseekV4ForCausalLM"],
+            n_routed_experts=4,
+            num_experts_per_tok=2,
+            num_nextn_predict_layers=1,
+            compress_ratios=[1, 1],
+            head_dim=32,
+            **common,
+        )
+    else:
+        target_hf_config = MixtralConfig(
+            architectures=["MixtralForCausalLM"], num_local_experts=4, **common
+        )
+    target_path = tmp_path / "target"
+    draft_path = tmp_path / "draft"
+    _write_json(target_path / "config.json", target_hf_config.to_dict())
+    _write_json(
+        draft_path / "config.json",
+        LlamaConfig(architectures=["LlamaForCausalLM"], **common).to_dict(),
+    )
+    target_model_config = ModelConfig(
+        model=str(target_path), tokenizer_mode="skip", max_model_len=2048
+    )
+    target_parallel_config = ParallelConfig(
+        tensor_parallel_size=tp_size,
+        enable_expert_parallel=target_ep,
+        distributed_executor_backend="mp",
+    )
+    target_model_config.verify_with_parallel_config(target_parallel_config)
+    speculative_config = SpeculativeConfig(
+        method=method,
+        model=str(target_path if draft_is_moe else draft_path),
+        num_speculative_tokens=1,
+        target_model_config=target_model_config,
+        target_parallel_config=target_parallel_config,
+    )
+
+    assert speculative_config.method == method
+    assert speculative_config.draft_model_config.is_moe is draft_is_moe
+    assert speculative_config.draft_parallel_config.enable_expert_parallel is (
+        target_ep and draft_is_moe
+    )
+    assert speculative_config.draft_parallel_config.tensor_parallel_size == tp_size
+    assert target_parallel_config.enable_expert_parallel is target_ep
+    assert target_parallel_config.tensor_parallel_size == tp_size
+
+
+@pytest.mark.skip_global_cleanup
+@pytest.mark.parametrize("target_ep", [False, True], ids=["ep-off", "ep-on"])
+@pytest.mark.parametrize("pass_none", [False, True], ids=["omitted", "explicit-none"])
+def test_draft_parallel_config_preserves_ep_without_model(
+    target_ep: bool, pass_none: bool
+):
+    """Legacy callers without draft model information keep EP inheritance."""
+    target_parallel_config = ParallelConfig(
+        tensor_parallel_size=2,
+        enable_expert_parallel=target_ep,
+        distributed_executor_backend="mp",
+    )
+    if pass_none:
+        draft_parallel_config = SpeculativeConfig.create_draft_parallel_config(
+            target_parallel_config, 2, draft_model_config=None
+        )
+    else:
+        draft_parallel_config = SpeculativeConfig.create_draft_parallel_config(
+            target_parallel_config, 2
+        )
+
+    assert draft_parallel_config.enable_expert_parallel is target_ep
+    assert draft_parallel_config.tensor_parallel_size == 2
+    assert target_parallel_config.enable_expert_parallel is target_ep
+
+
 @pytest.mark.parametrize(
     ("method", "parallel_drafting", "expected_slots"),
     [

tests/v1/attention/test_attention_backends.py

@@ -243,14 +243,19 @@ def create_and_prepopulate_kv_cache(
 class MockAttentionLayer:
     """A mock attention layer for testing."""
 
-    def __init__(self, device: torch.device):
+    def __init__(
+        self,
+        device: torch.device,
+        k_scale: float = 1.0,
+        v_scale: float = 1.0,
+    ):
         self._q_scale = torch.tensor(1.0, device=device)
-        self._k_scale = torch.tensor(1.0, device=device)
-        self._v_scale = torch.tensor(1.0, device=device)
+        self._k_scale = torch.tensor(k_scale, device=device)
+        self._v_scale = torch.tensor(v_scale, device=device)
         # Add float versions for flashinfer
         self._q_scale_float = 1.0
-        self._k_scale_float = 1.0
-        self._v_scale_float = 1.0
+        self._k_scale_float = k_scale
+        self._v_scale_float = v_scale
 
 
 def _clone_kv_cache_in_layout(
@@ -285,6 +290,8 @@ def run_attention_backend(
     kv_cache_dtype: str = "auto",
     sinks: torch.Tensor | None = None,
     use_cuda_graph: bool = False,
+    layer_k_scale: float = 1.0,
+    layer_v_scale: float = 1.0,
 ) -> torch.Tensor:
     """Run attention computation using the specified backend's AttentionImpl."""
     use_direct_block_mask = is_torch_equal_or_newer("2.9.0.dev0")
@@ -356,7 +363,7 @@ def run_attention_backend(
         )
 
     # Create mock layer and output buffer
-    mock_layer = MockAttentionLayer(device)
+    mock_layer = MockAttentionLayer(device, layer_k_scale, layer_v_scale)
     output = torch.empty_like(query)
 
     if is_quantized_kv_cache(kv_cache_dtype) and impl.supports_quant_query_input:
@@ -386,9 +393,11 @@ def run_attention_backend(
             )
         graph.replay()
     else:
-        output = impl.forward(
+        backend_output = impl.forward(
             mock_layer, query, key, value, kv_cache, attn_metadata, output=output
         )
+        if backend_output is not None:
+            output = backend_output
 
     return output
 
@@ -413,6 +422,8 @@ def _test_backend_correctness(
     model_dtype: torch.dtype | None = None,
     max_num_seqs: int | None = None,
     max_num_batched_tokens: int | None = None,
+    layer_k_scale: float = 1.0,
+    layer_v_scale: float = 1.0,
 ):
     """Test that all backends produce similar outputs to a reference implementation
     using FlexAttention or an explicit attention-sink reference.
@@ -687,6 +698,8 @@ def _test_backend_correctness(
             kv_cache_dtype=kv_cache_dtype,
             sinks=sinks,
             use_cuda_graph=use_cuda_graph,
+            layer_k_scale=layer_k_scale,
+            layer_v_scale=layer_v_scale,
         )
 
         # Check shape and dtype consistency
@@ -1317,6 +1330,45 @@ def test_sliding_window_backend_correctness(
         )
 
 
+@pytest.mark.skipif(not current_platform.is_rocm(), reason="ROCm-specific test")
+def test_rocm_aiter_fa_unquantized_cache_ignores_kv_scales(
+    default_vllm_config,
+):
+    """Stale scale tensors must not affect an unquantized KV cache."""
+    from vllm._aiter_ops import is_aiter_found_and_supported
+
+    if not is_aiter_found_and_supported():
+        pytest.skip("AITER is required")
+
+    def sliding_window_mask_mod(
+        b: torch.Tensor,
+        h: torch.Tensor,
+        q_idx: torch.Tensor,
+        kv_idx: torch.Tensor,
+        *,
+        context_len: int,
+        sliding_window: int,
+    ):
+        causal_mask = q_idx + context_len >= kv_idx
+        window_mask = q_idx + context_len - kv_idx < sliding_window
+        return causal_mask & window_mask
+
+    model = "microsoft/Phi-tiny-MoE-instruct"
+    batch_spec = BATCH_SPECS["small_decode"]
+    model_config = ModelConfig(model=model, max_model_len=max(batch_spec.seq_lens))
+    sliding_window = model_config.get_sliding_window()
+    assert sliding_window is not None
+
+    _test_backend_correctness(
+        batch_spec,
+        model,
+        [AttentionBackendEnum.ROCM_AITER_FA],
+        partial(sliding_window_mask_mod, sliding_window=sliding_window),
+        layer_k_scale=0.25,
+        layer_v_scale=0.5,
+    )
+
+
 @pytest.mark.parametrize(
     "batch_spec_name",
     [

tests/v1/core/prefix_cache/test_hybrid_prefix_cache_hit_rate.py

@@ -0,0 +1,156 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Replay fixed chat histories and check per-turn cache hits and metrics."""
+
+import asyncio
+import json
+from uuid import uuid4
+
+import httpx
+import pytest
+from openai.types.chat import ChatCompletionMessageParam
+from prometheus_client.parser import text_string_to_metric_families
+
+from tests.utils import RemoteOpenAIServer
+
+MODEL = "Qwen/Qwen3.5-0.8B"
+NUM_CONVERSATIONS = 8
+MAX_TOKENS = 32
+# Long first turns leave room for hybrid cache alignment and MTP recomputation.
+MIN_PROMPT_TOKENS = 2048
+METRICS_TIMEOUT = 30
+
+
+def _user_turns(conversation: int) -> list[str]:
+    records = "\n".join(
+        f"Sensor {conversation}-{i}: temperature {20 + i % 10} C, "
+        f"humidity {40 + i % 20} percent."
+        for i in range(160 + conversation)
+    )
+    return [
+        f"Session {conversation}. Summarize these sensor readings briefly:\n{records}",
+        "Which readings should an operator investigate first? Explain briefly.",
+        "Suggest two checks the operator should perform next.",
+    ]
+
+
+async def _prefix_cache_counters(
+    client: httpx.AsyncClient,
+) -> tuple[float, float]:
+    response = await client.get("/metrics")
+    response.raise_for_status()
+    names = ("vllm:prefix_cache_queries_total", "vllm:prefix_cache_hits_total")
+    counters: dict[str, float] = {}
+    for family in text_string_to_metric_families(response.text):
+        for sample in family.samples:
+            if sample.name in names:
+                counters[sample.name] = counters.get(sample.name, 0) + sample.value
+    assert all(name in counters for name in names), (
+        f"missing prefix-cache metrics: {set(names) - counters.keys()}"
+    )
+    return counters[names[0]], counters[names[1]]
+
+
+async def _wait_for_counters(
+    client: httpx.AsyncClient, expected: tuple[float, float]
+) -> None:
+    """Allow engine stats to reach the frontend after completions are returned."""
+    deadline = asyncio.get_running_loop().time() + METRICS_TIMEOUT
+    while True:
+        actual = await _prefix_cache_counters(client)
+        if actual == expected:
+            return
+        assert all(a <= e for a, e in zip(actual, expected)), (
+            f"prefix-cache counters exceeded usage: {actual=}, {expected=}"
+        )
+        assert asyncio.get_running_loop().time() < deadline, (
+            f"prefix-cache counters did not match usage: {actual=}, {expected=}"
+        )
+        await asyncio.sleep(0.1)
+
+
+@pytest.fixture(scope="module")
+def server(request):
+    args = ["--enable-prompt-tokens-details"]
+    if request.param:
+        args += [
+            "--speculative-config",
+            json.dumps({"method": "mtp", "num_speculative_tokens": 1}),
+        ]
+    with RemoteOpenAIServer(MODEL, args) as remote_server:
+        yield remote_server
+
+
+@pytest.mark.asyncio
+@pytest.mark.parametrize(
+    "server, min_hit_rate",
+    # MTP re-prefills an extra block; both floors reject one more missed block.
+    [pytest.param(False, 0.90, id="base"), pytest.param(True, 0.75, id="mtp")],
+    indirect=["server"],
+)
+async def test_prefix_cache_hit_rate(
+    server: RemoteOpenAIServer, min_hit_rate: float
+) -> None:
+    # Isolate sessions so another conversation cannot hide a cache miss.
+    salts = [uuid4().hex for _ in range(NUM_CONVERSATIONS)]
+    histories: list[list[ChatCompletionMessageParam]] = [
+        [] for _ in range(NUM_CONVERSATIONS)
+    ]
+    prompt_lengths: list[list[int]] = [[] for _ in range(NUM_CONVERSATIONS)]
+
+    async with (
+        server.get_async_client() as client,
+        httpx.AsyncClient(base_url=server.url_root, timeout=METRICS_TIMEOUT) as metrics,
+    ):
+
+        async def run_conversation(session: int, *, replay: bool) -> tuple[int, int]:
+            history = histories[session]
+            queries = hits = 0
+            for turn, user_text in enumerate(_user_turns(session)):
+                if not replay:
+                    history.append({"role": "user", "content": user_text})
+                # Reuse the original replies, even if replay generates different ones.
+                response = await client.chat.completions.create(
+                    model=MODEL,
+                    messages=history[: 2 * turn + 1],
+                    temperature=0,
+                    max_tokens=MAX_TOKENS,
+                    extra_body={"cache_salt": salts[session]},
+                )
+                usage = response.usage
+                assert usage is not None
+                assert usage.prompt_tokens_details is not None
+                prompt = usage.prompt_tokens
+                cached = usage.prompt_tokens_details.cached_tokens
+                context = f"{session=}, {turn=}, {replay=}: {cached=}, {prompt=}"
+                assert cached is not None, context
+                assert 0 <= cached < prompt, context
+                assert prompt >= MIN_PROMPT_TOKENS, context
+                if not replay and turn == 0:
+                    assert cached == 0, context
+                else:
+                    assert cached / prompt >= min_hit_rate, (
+                        f"{context}, {min_hit_rate=}"
+                    )
+                if replay:
+                    assert prompt == prompt_lengths[session][turn], context
+                else:
+                    prompt_lengths[session].append(prompt)
+                    history.append(
+                        {
+                            "role": "assistant",
+                            "content": response.choices[0].message.content or "",
+                        }
+                    )
+                queries += prompt
+                hits += cached
+            return queries, hits
+
+        queries, hits = await _prefix_cache_counters(metrics)
+        for replay in (False, True):
+            counts = await asyncio.gather(
+                *(run_conversation(i, replay=replay) for i in range(NUM_CONVERSATIONS))
+            )
+            queries += sum(prompt for prompt, _ in counts)
+            hits += sum(cached for _, cached in counts)
+            await _wait_for_counters(metrics, (queries, hits))

tests/v1/core/test_scheduler.py

@@ -193,6 +193,22 @@ def test_get_num_unfinished_requests():
         assert scheduler.get_num_unfinished_requests() == len(requests) - i - 1
 
 
+def test_max_num_active_seqs_caps_admission():
+    """Admission is capped by max_num_active_seqs, not max_num_seqs."""
+    scheduler = create_scheduler(max_num_seqs=16, max_num_active_seqs=3)
+    requests = create_requests(num_requests=10)
+    for request in requests:
+        scheduler.add_request(request)
+
+    output = scheduler.schedule()
+
+    assert len(output.scheduled_new_reqs) == 3
+    assert len(scheduler.running) == 3
+    assert len(scheduler.waiting) == 7
+    # Execution capacity is untouched.
+    assert scheduler.max_num_running_reqs == 16
+
+
 def _bind_hisparse_connector(scheduler):
     connector = object.__new__(HiSparseConnector)
     connector.connector_scheduler = HiSparseConnectorScheduler(async_speculative=False)

tests/v1/core/utils.py

@@ -52,6 +52,7 @@ def mock_kv(matched_tokens: int, is_async: bool, num_defers_before_matching: int
 def create_scheduler(
     model: str = "facebook/opt-125m",
     max_num_seqs: int = 16,
+    max_num_active_seqs: int | None = None,
     max_num_batched_tokens: int = 8192,
     enable_chunked_prefill: bool = True,
     enable_prefix_caching: bool = False,
@@ -107,6 +108,7 @@ def create_scheduler(
         max_model_len = max_num_batched_tokens
     scheduler_config = SchedulerConfig(
         max_num_seqs=max_num_seqs,
+        max_num_active_seqs=max_num_active_seqs,
         max_num_batched_tokens=max_num_batched_tokens,
         max_model_len=max_model_len,
         long_prefill_token_threshold=long_prefill_token_threshold,

tests/v1/metrics/test_perf_metrics.py

@@ -133,7 +133,56 @@ def test_base_config_parser():
     assert result.num_attention_heads == 16
     assert result.num_hidden_layers == 24
     assert result.weight_byte_size == 2  # float16 is 2 bytes
-    assert result.activation_byte_size == 2  # default activation size
+    assert result.activation_byte_size == 2  # float16 activations are 2 bytes
+
+
+@pytest.mark.parametrize(
+    "model_dtype,expected_byte_size",
+    [("float16", 2), ("bfloat16", 2), ("float32", 4)],
+)
+def test_base_config_parser_activation_byte_size_follows_dtype(
+    model_dtype: str, expected_byte_size: int
+):
+    """Activation traffic is sized from the model's compute dtype.
+
+    Activations are produced in the model dtype, so pinning them to 2 bytes
+    under-reported memory traffic by half for a float32 model while that same
+    model's weight traffic was sized correctly.
+    """
+    hf_config = Qwen3Config(
+        hidden_size=2048,
+        num_attention_heads=16,
+        num_hidden_layers=24,
+    )
+    vllm_config = create_mock_vllm_config(hf_config, model_dtype=model_dtype)
+
+    result = BaseConfigParser().parse(ParsedArgs(), vllm_config)
+
+    assert result.activation_byte_size == expected_byte_size
+    assert result.weight_byte_size == expected_byte_size
+
+
+def test_attention_read_bytes_scale_with_activation_dtype():
+    """A float32 model reads twice the activation bytes of a bfloat16 one."""
+    hf_config = Qwen3Config(
+        hidden_size=2048,
+        num_attention_heads=16,
+        num_key_value_heads=16,
+        num_hidden_layers=4,
+    )
+    ctx = ExecutionContext.from_single_request(
+        num_tokens=128, context_len=128, is_prefill=True
+    )
+
+    bf16 = AttentionMetrics.from_vllm_config(
+        create_mock_vllm_config(hf_config, model_dtype="bfloat16")
+    ).get_read_bytes_breakdown(ctx)
+    fp32 = AttentionMetrics.from_vllm_config(
+        create_mock_vllm_config(hf_config, model_dtype="float32")
+    ).get_read_bytes_breakdown(ctx)
+
+    # qkv_input is pure activation traffic, so it doubles exactly.
+    assert fp32["qkv_input"] == 2 * bf16["qkv_input"]
 
 
 def test_base_attention_config_parser_with_gqa():

vllm/config/parallel.py

@@ -717,6 +717,7 @@ class ParallelConfig:
                 "allgather_reducescatter",
                 "deepep_high_throughput",
                 "deepep_low_latency",
+                "deepep_v2",
                 "flashinfer_nvlink_one_sided",
                 "mori_high_throughput",
                 "mori_low_latency",

vllm/config/scheduler.py

@@ -67,6 +67,16 @@ class SchedulerConfig:
     In real usage, this should be set in `EngineArgs.create_engine_config`.
     """
 
+    max_num_active_seqs: int | None = Field(default=None, ge=1)
+    """Maximum number of requests the scheduler admits into RUNNING.
+
+    ``max_num_seqs`` sizes the model runner (per-request buffers and CUDA
+    graph capture) and is also the default admission limit. Setting this
+    lowers only the number of requests that may occupy RUNNING, so decode
+    batches stay smaller without shrinking runner or graph capacity. Must
+    be ``<= max_num_seqs``. ``None`` (default) keeps current behavior.
+    """
+
     long_prefill_token_threshold: int = Field(default=0, ge=0)
     """For chunked prefill, a request is considered long if the prompt is
     longer than this number of tokens. 0 disables the cap (default)."""
@@ -310,6 +320,15 @@ class SchedulerConfig:
                 f"({self.max_num_seqs})."
             )
 
+        if (
+            self.max_num_active_seqs is not None
+            and self.max_num_active_seqs > self.max_num_seqs
+        ):
+            raise ValueError(
+                f"max_num_active_seqs ({self.max_num_active_seqs}) cannot be "
+                f"greater than max_num_seqs ({self.max_num_seqs})."
+            )
+
         if self.max_num_batched_tokens > self.max_num_seqs * max_model_len:
             logger.warning(
                 "max_num_batched_tokens (%d) exceeds max_num_seqs "

vllm/config/speculative.py

@@ -1545,9 +1545,13 @@ class SpeculativeConfig:
                     )
                 )
 
+                # Use the final draft config, after any architecture overrides,
+                # so EP agrees with ModelConfig's expert-count validation.
                 self.draft_parallel_config = (
                     SpeculativeConfig.create_draft_parallel_config(
-                        self.target_parallel_config, self.draft_tensor_parallel_size
+                        self.target_parallel_config,
+                        self.draft_tensor_parallel_size,
+                        draft_model_config=self.draft_model_config,
                     )
                 )
 
@@ -1743,15 +1747,21 @@ class SpeculativeConfig:
     def create_draft_parallel_config(
         target_parallel_config: ParallelConfig,
         speculative_draft_tensor_parallel_size: int,
+        draft_model_config: ModelConfig | None = None,
     ) -> ParallelConfig:
         """Create a parallel config for use by the draft worker.
 
-        This is mostly a copy of the target parallel config, except the tp_size.
+        Use the draft TP size and disable inherited EP for known dense drafts.
+        Without a draft model config, preserve the previous EP inheritance.
         """
+        enable_ep = target_parallel_config.enable_expert_parallel
+        if draft_model_config is not None:
+            enable_ep = enable_ep and draft_model_config.is_moe
+
         draft_parallel_config = ParallelConfig(
             pipeline_parallel_size=1,
             tensor_parallel_size=speculative_draft_tensor_parallel_size,
-            enable_expert_parallel=target_parallel_config.enable_expert_parallel,
+            enable_expert_parallel=enable_ep,
             distributed_executor_backend=target_parallel_config.distributed_executor_backend,
             max_parallel_loading_workers=target_parallel_config.max_parallel_loading_workers,
             disable_custom_all_reduce=target_parallel_config.disable_custom_all_reduce,

vllm/engine/arg_utils.py

@@ -556,6 +556,7 @@ class EngineArgs:
     max_num_scheduled_tokens: int | None = None
     long_prefill_token_threshold: int = SchedulerConfig.long_prefill_token_threshold
     max_num_seqs: int | None = None
+    max_num_active_seqs: int | None = SchedulerConfig.max_num_active_seqs
     max_num_queued_reqs: int | None = None
     max_num_queued_tokens: int | None = None
     max_logprobs: int = ModelConfig.max_logprobs
@@ -1589,6 +1590,10 @@ class EngineArgs:
                 "default": None,
             },
         )
+        scheduler_group.add_argument(
+            "--max-num-active-seqs",
+            **scheduler_kwargs["max_num_active_seqs"],
+        )
         scheduler_group.add_argument(
             "--max-num-queued-reqs", **scheduler_kwargs["max_num_queued_reqs"]
         )
@@ -2421,6 +2426,7 @@ class EngineArgs:
             max_num_batched_tokens=self.max_num_batched_tokens,
             max_num_scheduled_tokens=self.max_num_scheduled_tokens,
             max_num_seqs=self.max_num_seqs,
+            max_num_active_seqs=self.max_num_active_seqs,
             max_num_queued_reqs=self.max_num_queued_reqs,
             max_num_queued_tokens=self.max_num_queued_tokens,
             max_model_len=model_config.max_model_len,

vllm/model_executor/layers/fused_moe/all2all_utils.py

@@ -276,6 +276,7 @@ def maybe_make_prepare_finalize(
             num_topk=moe.experts_per_token,
             use_fp8_dispatch=use_fp8_dispatch,
             use_cudagraph=use_cudagraph,
+            sp_size=moe.moe_parallel_config.sp_size,
         )
 
     elif moe.use_mori_kernels:
@@ -284,13 +285,20 @@ def maybe_make_prepare_finalize(
         # Note: We may want to use FP8 dispatch just to reduce
         # data movement.
         use_fp8_dispatch = (
-            quant_config.is_per_act_token or quant_config.is_block_quantized
+            quant_config.is_per_act_token
+            or quant_config.is_block_quantized
+            or quant_config.is_per_tensor
         )
         if use_fp8_dispatch:
-            # For PTPC (per token per channel) quant, scale dim is 1
-            # For 1x128 quant, scale dim is hidden_dim // 128
+            # For PTPC (per token per channel) or per-tensor quant,
+            # scale dim is 1. For 1x128 quant, scale dim is
+            # hidden_dim // 128
             quant_dtype = quant_config.quant_dtype
-            scale_dim = 1 if quant_config.is_per_act_token else moe.hidden_dim // 128
+            scale_dim = (
+                1
+                if (quant_config.is_per_act_token or quant_config.is_per_tensor)
+                else moe.hidden_dim // 128
+            )
         else:
             # Unquantized dispatch (e.g. AITER with defer_input_quant):
             # dispatch raw BF16/FP16 data, no scales needed.

vllm/model_executor/layers/fused_moe/config.py

@@ -17,14 +17,19 @@ from vllm.model_executor.layers.quantization.utils.ocp_mx_utils import (
 )
 from vllm.model_executor.layers.quantization.utils.quant_utils import GroupShape
 from vllm.platforms import current_platform
-from vllm.utils.import_utils import has_triton_kernels
+from vllm.utils.import_utils import get_triton_kernels_version
 from vllm.utils.math_utils import cdiv
 
 logger = init_logger(__name__)
 
-if has_triton_kernels():
+_triton_kernels_version = get_triton_kernels_version()
+
+if _triton_kernels_version is not None:
     try:
-        from triton_kernels.matmul_ogs import PrecisionConfig
+        if _triton_kernels_version == "3.8":
+            from triton_kernels.matmul import PrecisionConfig
+        else:
+            from triton_kernels.matmul_ogs import PrecisionConfig
     except (ImportError, AttributeError) as e:
         logger.error(
             "Failed to import Triton kernels. Please make sure your triton "

vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a16_moe.py

@@ -12,9 +12,6 @@ from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEQuantConfig,
     RoutingMethodType,
 )
-from vllm.model_executor.layers.fused_moe.experts.aiter_mxfp4_w4a8_moe import (
-    patch_gating_output,
-)
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
     QuantKey,
     kMxfp4Static,
@@ -26,17 +23,6 @@ __all__ = [
 ]
 
 
-def _aiter_raw(t):
-    """Unwrap the `triton_kernels.tensor.Tensor` that `_swizzle_mxfp4` returns."""
-    if t is None or isinstance(t, torch.Tensor):
-        return t
-    assert hasattr(t, "storage"), (
-        f"expected a triton_kernels wrapped tensor with a .storage attribute, "
-        f"got {type(t)}"
-    )
-    return t.storage.data
-
-
 def _aiter_w4a16_silu_via_a8w4(
     hidden_states: torch.Tensor,
     w1_data,
@@ -168,11 +154,6 @@ def aiter_triton_kernel_w4a16_moe_forward(
         _routing_mod.is_tdm_avail = lambda: False
     aiter_routing = _routing_mod.routing
 
-    # See context in #50859.
-    # TODO: Remove once https://github.com/ROCm/aiter/pull/4530 is merged,
-    # AITER released, and AITER pin in vLLM increased from 0.1.19.
-    gating_output = patch_gating_output(gating_output, global_num_experts)
-
     if score_mode is not None:
         use_grouped_topk = num_expert_group is not None and num_expert_group > 1
         routing_data, gather_idx, scatter_idx = aiter_routing(
@@ -201,10 +182,14 @@ def aiter_triton_kernel_w4a16_moe_forward(
     assert quant_config.w1_precision is not None
     assert quant_config.w2_precision is not None
 
-    w1_data = _aiter_raw(w1)
-    w2_data = _aiter_raw(w2)
-    w1_wscale = _aiter_raw(quant_config.w1_precision.weight_scale)
-    w2_wscale = _aiter_raw(quant_config.w2_precision.weight_scale)
+    w1_data = w1.storage.data
+    w2_data = w2.storage.data
+    from vllm.model_executor.layers.quantization.utils.mxfp4_utils import (
+        weight_mx_scale,
+    )
+
+    w1_wscale = weight_mx_scale(quant_config.w1_precision).storage.data
+    w2_wscale = weight_mx_scale(quant_config.w2_precision).storage.data
 
     gammas = routing_data.gate_scal if routing_data else None
 

vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a8_moe.py

@@ -5,7 +5,6 @@ import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm._aiter_ops import rocm_aiter_ops
-from vllm.forward_context import get_forward_context, is_forward_context_available
 from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
@@ -25,60 +24,6 @@ __all__ = [
 ]
 
 
-# TODO(akaratza): Remove
-# _get_padding_mask and patch_gating_output after
-# https://github.com/ROCm/aiter/pull/4530 is released and consumed by vLLM.
-def _get_padding_mask() -> torch.Tensor | None:
-    """Retrieves a boolean mask with non-padding (0) and padding (1) tokens.
-
-    `slot_mapping < 0` comes from:
-
-        slot_mapping[num_tokens_unpadded:num_tokens_padded].fill_(-1)
-
-    in gpu_model_runner.py.
-
-    Direct kernel callers do not have model-runner padding metadata, so no
-    mask is needed when there is no forward context.
-    """
-    if not is_forward_context_available():
-        return None
-
-    forward_context = get_forward_context()
-
-    # model runner v2.
-    if forward_context.is_padding is not None:
-        return forward_context.is_padding
-
-    # model runner v1.
-    slot_mapping = forward_context.slot_mapping
-
-    if isinstance(slot_mapping, list):
-        slot_mapping_dict = slot_mapping[0]
-    else:
-        slot_mapping_dict = slot_mapping
-
-    if isinstance(slot_mapping_dict, dict):
-        slot_mapping_sample = next(iter(slot_mapping_dict.values()), None)
-
-    if isinstance(slot_mapping_sample, torch.Tensor):
-        return slot_mapping_sample < 0
-    else:
-        return None
-
-
-def patch_gating_output(
-    gating_output: torch.Tensor, global_num_experts: int
-) -> torch.Tensor:
-    if global_num_experts != 128:
-        is_padding = _get_padding_mask()
-
-        if is_padding is not None:
-            assert is_padding.shape[0] == gating_output.shape[0]
-            gating_output = gating_output.masked_fill(is_padding[:, None], 0.0)
-
-    return gating_output
-
-
 def aiter_triton_kernel_w4a8_moe_forward(
     hidden_states: torch.Tensor,
     w1,  # Tensor or triton_kernels.Tensor
@@ -112,11 +57,6 @@ def aiter_triton_kernel_w4a8_moe_forward(
         _routing_mod.is_tdm_avail = lambda: False
     aiter_routing = _routing_mod.routing
 
-    # See context in #50859.
-    # TODO: Remove once https://github.com/ROCm/aiter/pull/4530 is merged,
-    # AITER released, and AITER pin in vLLM increased from 0.1.19.
-    gating_output = patch_gating_output(gating_output, global_num_experts)
-
     routing_data, gather_idx, scatter_idx = aiter_routing(
         gating_output, topk, sm_first=not renormalize
     )
@@ -191,6 +131,7 @@ def triton_kernel_fused_mxfp4_w4a8_experts(
 
     from vllm.model_executor.layers.quantization.utils.mxfp4_utils import (
         should_use_cdna4_mx_scale_swizzle,
+        weight_mx_scale,
     )
 
     _swizzle_mx_scale = "CDNA4_SCALE" if should_use_cdna4_mx_scale_swizzle() else None
@@ -210,7 +151,7 @@ def triton_kernel_fused_mxfp4_w4a8_experts(
         hidden_states,
         w1.storage.data,
         None,
-        quant_config.w1_precision.weight_scale.storage.data,
+        weight_mx_scale(quant_config.w1_precision).storage.data,
         quant_config.w1_precision.flex_ctx.lhs_data.scale,
         quant_config.w2_precision.flex_ctx.lhs_data.scale,
         quant_config.w1_bias,
@@ -230,7 +171,7 @@ def triton_kernel_fused_mxfp4_w4a8_experts(
         intermediate_cache1,
         w2.storage.data,
         None,
-        quant_config.w2_precision.weight_scale.storage.data,
+        weight_mx_scale(quant_config.w2_precision).storage.data,
         quant_config.w2_precision.flex_ctx.lhs_data.scale,
         None,
         quant_config.w2_bias,

vllm/model_executor/layers/fused_moe/experts/gpt_oss_triton_kernels_moe.py

@@ -26,7 +26,7 @@ from vllm.model_executor.layers.quantization.utils.quant_utils import (
 )
 from vllm.platforms import current_platform
 from vllm.triton_utils import tl, triton
-from vllm.utils.import_utils import has_triton_kernels
+from vllm.utils.import_utils import get_triton_kernels_version, has_triton_kernels
 
 from ..utils import swiglu_limit_func
 
@@ -441,48 +441,59 @@ def _patch_legacy_routing_for_nonpow2_topk() -> None:
     _routing.sort_tokens = _sort_tokens_pow2
 
 
-# Two API generations of triton_kernels are supported:
-#   - v3.5.1 (the version bundled with vLLM): exposes `routing()` and
-#     `routing_from_bitmatrix()` in triton_kernels.routing; the `Bitmatrix`
-#     constructor takes a `scratchpad` argument.
-#   - v3.6.0+: removes the `routing` module in favor of a `SparseMatrix`
-#     based path, and adds a `dtype=BIT` kwarg to `Bitmatrix`. Used only
-#     when the user has triton_kernels installed system-wide at v3.6.0+.
-#
-# `use_legacy_triton_kernels` selects between them at import time based on
-# whether `SparseMatrix` is importable.
-use_legacy_triton_kernels = False
+triton_kernels_version = get_triton_kernels_version()
 
-if has_triton_kernels():
+if triton_kernels_version is not None:
     try:
         import triton_kernels.swiglu
-        from triton_kernels.matmul_ogs import (
-            FnSpecs,
-            FusedActivation,
-            GatherIndx,
-            RoutingData,
-            ScatterIndx,
-            matmul_ogs,
-        )
-        from triton_kernels.tensor import (
-            BIT,
-            Bitmatrix,
-        )
 
+        if triton_kernels_version == "3.8":
+            from triton_kernels.matmul import FnSpecs, FusedActivation, matmul
+            from triton_kernels.matmul_details.opt_flags import (
+                update_opt_flags_constraints,
+            )
+            from triton_kernels.tensor import make_ragged_tensor_metadata
+
+            from vllm.platforms.rocm import on_gfx950
+
+            if on_gfx950():
+                # block_n=64/block_k=256 fit gfx950's 160KB LDS + the CDNA4
+                # scale swizzle; block_m=32 avoids an autotuner block_m=128
+                # padding cliff (~15x) at decode-sized ragged MoE shapes.
+                update_opt_flags_constraints(
+                    {"block_m": 32, "block_n": 64, "block_k": 256}
+                )
+        else:
+            from triton_kernels.matmul_ogs import (
+                FnSpecs,
+                FusedActivation,
+                GatherIndx,
+                RoutingData,
+                ScatterIndx,
+                matmul_ogs,
+            )
+
+            if triton_kernels_version == "3.6":
+                from triton_kernels.tensor import (
+                    SparseMatrix,
+                    make_ragged_tensor_metadata,
+                )
+
+        # BIT/Bitmatrix exist in v3.5.1/v3.6, gone in v3.8.
         try:
             from triton_kernels.tensor import (
-                SparseMatrix,
-                make_ragged_tensor_metadata,
+                BIT,
+                Bitmatrix,
             )
         except ImportError:
-            # TODO(mgoin): drop the v3.5.1 pin and remove this fallback once
-            # the gpt-oss perf regression in v3.6.0+ is resolved upstream.
-            # Tracking: https://github.com/triton-lang/triton/issues/9969
-            use_legacy_triton_kernels = True
-        if not use_legacy_triton_kernels:
+            BIT = None
+            Bitmatrix = None
+
+        # Patch the routing kernels for non-pow2 top_k (DeepSeek-V4 top_k=6);
+        # v3.8 rebuilds routing from topk ids and needs no patch.
+        if triton_kernels_version == "3.6":
             _patch_make_bitmatrix_metadata()
-        else:
-            # Legacy routing fails to compile for non-pow2 top_k (DeepSeek-V4).
+        elif triton_kernels_version == "3.5.1":
             _patch_legacy_routing_for_nonpow2_topk()
     except (AttributeError, ImportError) as e:
         logger.error(
@@ -492,6 +503,21 @@ if has_triton_kernels():
         )
 
 
+if triton_kernels_version == "3.8":
+    # 3.8's matmul dropped matmul_ogs's in-kernel scatter+topk-reduce, so the
+    # forward combines externally via index_add; this is the ragged routing
+    # bundle the (inlined) 3.8 paths pass around, in place of matmul_ogs's.
+    class RoutingData:  # type: ignore[no-redef]
+        __slots__ = ("ragged", "gate_scal", "gather_tok", "n_expts_act", "n_valid")
+
+        def __init__(self, ragged, gate_scal, gather_tok, n_expts_act, n_valid):
+            self.ragged = ragged
+            self.gate_scal = gate_scal
+            self.gather_tok = gather_tok
+            self.n_expts_act = n_expts_act
+            self.n_valid = n_valid
+
+
 @triton.jit
 def pack_bitmatrix(
     bitmatrix,
@@ -562,7 +588,7 @@ def triton_kernel_moe_forward(
     # three separate kernels used by the generic path below.
     # Only available in the legacy (v3.5.1) API; the v3.6.0+ path inlines
     # equivalent logic via SparseMatrix in `make_routing_data`.
-    if use_legacy_triton_kernels and expert_map is None:
+    if triton_kernels_version == "3.5.1" and expert_map is None:
         from triton_kernels.routing import routing as fused_routing
 
         routing_data, gather_idx, scatter_idx = fused_routing(
@@ -662,6 +688,52 @@ def triton_kernel_fused_experts(
     assert activation == MoEActivation.SWIGLUOAI, (
         "Only SWIGLUOAI activation is supported"
     )
+
+    if triton_kernels_version == "3.8":
+        # 3.8 has no in-kernel scatter+topk-reduce, so the two matmuls are
+        # followed by an external index_add combine.
+        assert quant_config is not None
+        M, K = hidden_states.shape[-2:]
+        a_ragged = routing_data.ragged
+        gather_tok = routing_data.gather_tok
+        gammas = routing_data.gate_scal
+        act = FusedActivation(
+            FnSpecs(
+                "swiglu",
+                triton_kernels.swiglu.swiglu_fn,
+                ("alpha", "limit"),
+                reduction_n=2,
+            ),
+            (swiglu_alpha, swiglu_limit),
+        )
+        inter = matmul(
+            hidden_states,
+            w1,
+            quant_config.w1_bias,
+            a_ragged,
+            None,
+            gather_tok,
+            None,
+            quant_config.w1_precision,
+            gammas=gammas if apply_router_weight_on_input else None,
+            fused_activation=act,
+        )
+        y = matmul(
+            inter,
+            w2,
+            quant_config.w2_bias,
+            a_ragged,
+            None,
+            None,
+            None,
+            quant_config.w2_precision,
+            gammas=None if apply_router_weight_on_input else gammas,
+        )
+        acc = torch.zeros((M, K), dtype=torch.float32, device=hidden_states.device)
+        acc.index_add_(0, gather_tok.to(torch.int64), y.to(torch.float32))
+        out = output_tensor.view(M, K)
+        out.copy_(acc.to(out.dtype))
+        return out
     assert quant_config is not None
 
     # type check, uint8 means mxfp4
@@ -704,7 +776,7 @@ def triton_kernel_fused_experts(
             ),
             (swiglu_alpha, swiglu_limit),
         )
-        if not use_legacy_triton_kernels
+        if triton_kernels_version != "3.5.1"
         else FusedActivation(
             FnSpecs("swiglu", triton_kernels.swiglu.swiglu_fn, ("alpha", "limit")),
             (swiglu_alpha, swiglu_limit),
@@ -744,6 +816,24 @@ def make_routing_data(
     topk_weights: torch.Tensor,
     num_local_experts: int,
 ) -> tuple["RoutingData", torch.Tensor, torch.Tensor]:
+    if triton_kernels_version == "3.8":
+        _, topk = topk_ids.shape
+        flat_e = topk_ids.reshape(-1)
+        valid = flat_e >= 0
+        flat_e = flat_e.to(torch.int64).clamp(min=0)
+        flat_w = topk_weights.reshape(-1).to(torch.float32) * valid.to(torch.float32)
+        nrows = flat_e.shape[0]
+        order = torch.argsort(flat_e, stable=True)
+        col_sum = torch.zeros(
+            num_local_experts, dtype=torch.int32, device=flat_e.device
+        )
+        col_sum.scatter_add_(0, flat_e, torch.ones_like(flat_e, dtype=torch.int32))
+        ragged = make_ragged_tensor_metadata(col_sum, nrows)
+        gather_tok = (order // topk).to(torch.int32)
+        gammas = flat_w[order]
+        routing = RoutingData(ragged, gammas, gather_tok, topk, nrows)
+        return routing, gather_tok, gather_tok
+
     topk_ids = topk_ids.to(torch.int16)
     topk_weights = topk_weights.to(torch.bfloat16)
 
@@ -774,7 +864,7 @@ def make_routing_data(
         Bitmatrix(
             bitmatrix, dtype=BIT, shape=bitmatrix_shape, shape_max=bitmatrix_shape_max
         )
-        if not use_legacy_triton_kernels
+        if triton_kernels_version != "3.5.1"
         else Bitmatrix(
             bitmatrix,
             shape=bitmatrix_shape,
@@ -786,7 +876,7 @@ def make_routing_data(
     # matmul_ogs expects invalid topk_weights to be -1s
     topk_weights = torch.where(topk_ids == -1, -1.0, topk_weights)
 
-    if use_legacy_triton_kernels:
+    if triton_kernels_version == "3.5.1":
         from triton_kernels.routing import routing_from_bitmatrix
 
         return routing_from_bitmatrix(
@@ -809,6 +899,10 @@ def routing_data_from_sparse_topk(
     call already produced; when the topk ids are used as-is (no expert-map
     remapping) that work is redundant.
     """
+    if triton_kernels_version == "3.8":
+        # 3.8 routing is rebuilt from raw ids; no SparseMatrix reuse.
+        return make_routing_data(sparse_topk.indx, sparse_topk.vals, n_expts)
+
     dispatch_indx = sparse_topk.mask_metadata.row_sorted_indx
     combine_indx = sparse_topk.mask_metadata.col_sorted_indx
     ragged_batch_metadata = make_ragged_tensor_metadata(
@@ -1197,6 +1291,50 @@ class UnfusedOAITritonExperts(LoRAExpertsMixin, BaseOAITritonExperts):
 
         topk = topk_ids.size(1)
 
+        if triton_kernels_version == "3.8":
+            # Activation and topk-reduce kept unfused from the matmuls (3.8 has
+            # no in-kernel scatter+reduce); combine via external index_add.
+            assert self._lora_context is None, (
+                "tk38 unfused MoE path does not support LoRA yet"
+            )
+            act_out_dim = self.adjust_N_for_activation(w1.shape[2], activation)
+            M, K = hidden_states.shape[-2:]
+            a_ragged = routing_data.ragged
+            gather_tok = routing_data.gather_tok
+            gammas = routing_data.gate_scal
+            inter = matmul(
+                hidden_states,
+                w1,
+                quant_config.w1_bias,
+                a_ragged,
+                None,
+                gather_tok,
+                None,
+                quant_config.w1_precision,
+                gammas=gammas if apply_router_weight_on_input else None,
+            )
+            act_out = torch.empty(
+                (routing_data.n_valid, act_out_dim),
+                dtype=hidden_states.dtype,
+                device=hidden_states.device,
+            )
+            self.activation(activation, act_out, inter)
+            down = matmul(
+                act_out,
+                w2,
+                quant_config.w2_bias,
+                a_ragged,
+                None,
+                None,
+                None,
+                quant_config.w2_precision,
+                gammas=None if apply_router_weight_on_input else gammas,
+            )
+            acc = torch.zeros((M, K), dtype=torch.float32, device=hidden_states.device)
+            acc.index_add_(0, gather_tok.to(torch.int64), down.to(torch.float32))
+            output.view(M, K).copy_(acc.to(output.dtype))
+            return
+
         # type check, uint8 means mxfp4
         assert hidden_states.dtype == torch.bfloat16
         assert (

vllm/model_executor/layers/fused_moe/oracle/mxfp4.py

@@ -45,7 +45,7 @@ from vllm.model_executor.layers.quantization.utils.quant_utils import (
 )
 from vllm.model_executor.layers.quantization.utils.w8a8_utils import all_close_1d
 from vllm.platforms import current_platform
-from vllm.utils.import_utils import has_triton_kernels
+from vllm.utils.import_utils import get_triton_kernels_version
 from vllm.utils.math_utils import round_up
 
 if TYPE_CHECKING:
@@ -54,9 +54,14 @@ if TYPE_CHECKING:
 
 logger = init_logger(__name__)
 
-if has_triton_kernels():
+triton_kernels_version = get_triton_kernels_version()
+
+if triton_kernels_version is not None:
     try:
-        from triton_kernels.matmul_ogs import PrecisionConfig
+        if triton_kernels_version == "3.8":
+            from triton_kernels.matmul import FlexCtx, PrecisionConfig
+        else:
+            from triton_kernels.matmul_ogs import FlexCtx, PrecisionConfig
     except (ImportError, AttributeError) as e:
         logger.error(
             "Failed to import Triton kernels. Please make sure your triton "
@@ -65,6 +70,15 @@ if has_triton_kernels():
         )
 
 
+def _mx_scale_kwargs(scale):
+    """PrecisionConfig weight-scale kwargs: 3.8 uses b_mx_scale/b_microblock_size,
+    3.5.1/3.6 use weight_scale.
+    """
+    if triton_kernels_version == "3.8":
+        return {"b_mx_scale": scale, "b_microblock_size": 32}
+    return {"weight_scale": scale}
+
+
 def _pack_deepgemm_mxfp4_scales(
     w13_weight: torch.Tensor,
     w2_weight: torch.Tensor,
@@ -1182,7 +1196,6 @@ def convert_gpt_oss_weight_to_mxfp4_moe_kernel_format(
 
     elif mxfp4_backend == Mxfp4MoeBackend.AITER_MXFP4_FP8:
         # W4A8: MXFP4 weights + static FP8 activations (triton kernel)
-        from triton_kernels.matmul_ogs import FlexCtx, PrecisionConfig
         from triton_kernels.numerics import InFlexData
 
         if w13_bias is not None:
@@ -1217,11 +1230,11 @@ def convert_gpt_oss_weight_to_mxfp4_moe_kernel_format(
 
         # Create PrecisionConfig with both weight and activation info
         w13_precision_config = PrecisionConfig(
-            weight_scale=w13_scale,
+            **_mx_scale_kwargs(w13_scale),
             flex_ctx=FlexCtx(rhs_data=w13_flex, lhs_data=lhs_data13),
         )
         w2_precision_config = PrecisionConfig(
-            weight_scale=w2_scale,
+            **_mx_scale_kwargs(w2_scale),
             flex_ctx=FlexCtx(rhs_data=w2_flex, lhs_data=lhs_data2),
         )
 
@@ -1238,8 +1251,6 @@ def convert_gpt_oss_weight_to_mxfp4_moe_kernel_format(
         )
 
     elif mxfp4_backend in TRITON_BACKENDS:
-        from triton_kernels.matmul_ogs import FlexCtx, PrecisionConfig
-
         if w13_bias is not None:
             w13_bias = w13_bias.to(torch.float32)
         if w2_bias is not None:
@@ -1255,10 +1266,10 @@ def convert_gpt_oss_weight_to_mxfp4_moe_kernel_format(
         )
 
         w13_precision_config = PrecisionConfig(
-            weight_scale=w13_scale, flex_ctx=FlexCtx(rhs_data=w13_flex)
+            **_mx_scale_kwargs(w13_scale), flex_ctx=FlexCtx(rhs_data=w13_flex)
         )
         w2_precision_config = PrecisionConfig(
-            weight_scale=w2_scale, flex_ctx=FlexCtx(rhs_data=w2_flex)
+            **_mx_scale_kwargs(w2_scale), flex_ctx=FlexCtx(rhs_data=w2_flex)
         )
 
         # The original mxfp4 block scales have been swizzled into the
@@ -1653,8 +1664,6 @@ def convert_weight_to_mxfp4_moe_kernel_format(
     elif mxfp4_backend in TRITON_BACKENDS or (
         mxfp4_backend == Mxfp4MoeBackend.AITER_MXFP4_BF16 and is_gfx1250
     ):
-        from triton_kernels.matmul_ogs import FlexCtx, PrecisionConfig
-
         if mxfp4_backend == Mxfp4MoeBackend.AITER_TRITON_MXFP4_BF16:
             # AITER moe_gemm_a16w4 needs gate/up interleaved
             def interleave_gate_up(w: torch.Tensor) -> torch.Tensor:
@@ -1698,10 +1707,10 @@ def convert_weight_to_mxfp4_moe_kernel_format(
         )
 
         w13_precision_config = PrecisionConfig(
-            weight_scale=w13_scale, flex_ctx=FlexCtx(rhs_data=w13_flex)
+            **_mx_scale_kwargs(w13_scale), flex_ctx=FlexCtx(rhs_data=w13_flex)
         )
         w2_precision_config = PrecisionConfig(
-            weight_scale=w2_scale, flex_ctx=FlexCtx(rhs_data=w2_flex)
+            **_mx_scale_kwargs(w2_scale), flex_ctx=FlexCtx(rhs_data=w2_flex)
         )
 
         # The original mxfp4 block scales have been swizzled into the

vllm/model_executor/layers/fused_moe/prepare_finalize/deepep_v2.py

@@ -18,7 +18,7 @@ from vllm.model_executor.layers.quantization.utils.mxfp8_utils import (
     swizzle_mxfp8_scale,
 )
 from vllm.triton_utils import tl, triton
-from vllm.utils.math_utils import round_up
+from vllm.utils.math_utils import cdiv, round_up
 from vllm.v1.worker.ubatching import (
     dbo_current_ubatch_id,
     dbo_enabled,
@@ -114,6 +114,7 @@ class DeepEPV2PrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):
         num_topk: int,
         use_fp8_dispatch: bool = False,
         use_cudagraph: bool = False,
+        sp_size: int = 1,
     ):
         super().__init__()
         self.buffer = buffer
@@ -124,6 +125,7 @@ class DeepEPV2PrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):
         self.num_topk = num_topk
         self.use_fp8_dispatch = use_fp8_dispatch
         self.use_cudagraph = use_cudagraph
+        self.sp_size = sp_size
 
         # DBO microbatching: one handle slot per micro-batch.
         self.handles: list[deep_ep.EPHandle | None] = [None, None]
@@ -192,12 +194,14 @@ class DeepEPV2PrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):
         # high concurrency). Round up to a power of 2 instead: this bounds the
         # set to ~log2(max_num_batched_tokens) values (compiled once, then
         # cached) while staying small for decode (e.g. 1 token -> 1) and capped
-        # at the buffer's init capacity for prefill.
+        # at the buffer's init capacity for prefill. With sequence parallelism,
+        # each EP rank holds a ceil(n / sp_size) shard of its DP rank's batch.
         num_max_tokens_per_rank = None
         if not do_expand:
             dp_meta = get_forward_context().dp_metadata
             if dp_meta is not None:
                 n = int(dp_meta.num_tokens_across_dp_cpu.max())
+                n = cdiv(n, self.sp_size)
             else:
                 n = tokens.shape[0]
             num_max_tokens_per_rank = 1 << max(n - 1, 0).bit_length()

vllm/model_executor/layers/fused_moe/prepare_finalize/mori.py

@@ -82,6 +82,15 @@ class MoriPrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):
             elif quant_config.is_per_act_token:
                 quant_func = get_hip_quant(QuantType.per_Token)
                 a1, scale = quant_func(a1, quant_dtype=current_platform.fp8_dtype())
+            elif quant_config.is_per_tensor:
+                quant_func = get_hip_quant(QuantType.per_Tensor)
+                a1, scale = quant_func(
+                    a1,
+                    scale=quant_config.a1_scale,
+                    quant_dtype=current_platform.fp8_dtype(),
+                )
+                # mori expects one scale slot per token; broadcast.
+                scale = scale.expand(a1.shape[0], 1).contiguous()
 
         (
             dispatch_a1,

vllm/model_executor/layers/quantization/utils/mxfp4_utils.py

@@ -7,7 +7,7 @@ import torch
 from vllm.logger import init_logger
 from vllm.platforms import current_platform
 from vllm.triton_utils import tl, triton
-from vllm.utils.import_utils import has_triton_kernels
+from vllm.utils.import_utils import get_triton_kernels_version, has_triton_kernels
 from vllm.utils.torch_utils import direct_register_custom_op, is_torch_equal_or_newer
 
 logger = init_logger(__name__)
@@ -27,10 +27,23 @@ def should_use_cdna4_mx_scale_swizzle() -> bool:
     return on_gfx950() and get_tensor_model_parallel_world_size() <= 2
 
 
+def weight_mx_scale(precision_config):
+    """mxfp4 weight scale from a triton_kernels PrecisionConfig: 3.8 exposes it
+    as b_mx_scale, 3.5.1/3.6 as weight_scale.
+    """
+    if get_triton_kernels_version() == "3.8":
+        return precision_config.b_mx_scale
+    return precision_config.weight_scale
+
+
 def _swizzle_mxfp4(quant_tensor, scale, num_warps=8):
     """Weight swizzle for mxfp4 moe, used for OAI mxfp4 kernel."""
     assert has_triton_kernels()
-    import triton_kernels.matmul_ogs_details.opt_flags as opt_flags
+    if get_triton_kernels_version() == "3.8":
+        # 3.8: matmul_ogs_details -> matmul_details
+        import triton_kernels.matmul_details.opt_flags as opt_flags
+    else:
+        import triton_kernels.matmul_ogs_details.opt_flags as opt_flags
     from triton_kernels.numerics import InFlexData
     from triton_kernels.tensor import FP4, convert_layout, wrap_torch_tensor
     from triton_kernels.tensor_details import layout
@@ -51,6 +64,17 @@ def _swizzle_mxfp4(quant_tensor, scale, num_warps=8):
         )
         value_layout = StridedLayout
         scale_layout = StridedLayout
+    elif current_platform.is_rocm() and get_triton_kernels_version() == "3.8":
+        # StridedLayout(-2) is triton_kernels' default CDNA mxfp4 weight layout
+        # (make_default_matmul_mxfp4_w_layout); 3.8 wants layout instances.
+        from triton_kernels.tensor_details.layout import CDNA4MXScaleLayout
+
+        value_layout = StridedLayout(-2)
+        scale_layout = (
+            CDNA4MXScaleLayout()
+            if should_use_cdna4_mx_scale_swizzle()
+            else StridedLayout()
+        )
     elif current_platform.is_rocm():
         value_layout = StridedLayout
         if should_use_cdna4_mx_scale_swizzle():

vllm/utils/import_utils.py

@@ -527,6 +527,26 @@ def has_triton_kernels() -> bool:
     return is_available
 
 
+@cache
+def get_triton_kernels_version() -> str | None:
+    """The triton_kernels MoE-API generation ("3.5.1"/"3.6"/"3.8"), or None.
+
+    Inferred by capability since the package exposes no usable version: 3.8
+    replaced ``matmul_ogs`` with ``matmul``, and 3.5.1 predates ``SparseMatrix``.
+    """
+    if not has_triton_kernels():
+        return None
+    try:
+        import triton_kernels.matmul_ogs  # noqa: F401
+    except ImportError:
+        return "3.8"
+    try:
+        from triton_kernels.tensor import SparseMatrix  # noqa: F401
+    except ImportError:
+        return "3.5.1"
+    return "3.6"
+
+
 @cache
 def has_tilelang() -> bool:
     """Whether the optional `tilelang` package is available.

vllm/v1/attention/backends/rocm_aiter_fa.py

@@ -989,6 +989,20 @@ class AiterFlashAttentionImpl(AttentionImpl):
                 "cross-attention."
             )
 
+    def _get_kv_cache_descales(
+        self,
+        layer: AttentionLayer,
+        num_decodes: int,
+    ) -> tuple[torch.Tensor | None, torch.Tensor | None]:
+        # AITER treats non-null descales as an instruction to dequantize.
+        if not is_quantized_kv_cache(self.kv_cache_dtype):
+            return None, None
+        descale_shape = (num_decodes, self.num_kv_heads)
+        return (
+            layer._k_scale.expand(descale_shape),
+            layer._v_scale.expand(descale_shape),
+        )
+
     def extend_for_sliding_window(
         self,
         attn_metadata: AiterFlashAttentionMetadata,
@@ -1367,6 +1381,9 @@ class AiterFlashAttentionImpl(AttentionImpl):
                     or decode_max_query_len > 1
                     or self.sinks is not None
                 ):
+                    k_descale, v_descale = self._get_kv_cache_descales(
+                        layer, num_decodes
+                    )
                     assert not rocm_aiter_ops.is_shuffle_kv_cache_enabled(), (
                         "Shuffle KV cache layout is not supported with sliding "
                         "window, sinks, or speculative decoding (multi-token decode)."
@@ -1376,7 +1393,6 @@ class AiterFlashAttentionImpl(AttentionImpl):
                             flash_attn_with_kvcache,
                         )
 
-                        descale_shape = (num_decodes, key_cache.shape[2])
                         decode_query = query[:num_decode_tokens].reshape(
                             num_decodes,
                             decode_max_query_len,
@@ -1393,8 +1409,8 @@ class AiterFlashAttentionImpl(AttentionImpl):
                             window_size=self.sliding_window,
                             softcap=self.logits_soft_cap,
                             q_descale=None,
-                            k_descale=layer._k_scale.expand(descale_shape),
-                            v_descale=layer._v_scale.expand(descale_shape),
+                            k_descale=k_descale,
+                            v_descale=v_descale,
                             page_table=attn_metadata.block_table[:num_decodes],
                         )
                         output[:num_decode_tokens].copy_(
@@ -1412,10 +1428,6 @@ class AiterFlashAttentionImpl(AttentionImpl):
                             unified_attention,
                         )
 
-                        descale_shape = (
-                            num_decodes,
-                            key_cache.shape[2],
-                        )
                         unified_attention(
                             q=query[:num_decode_tokens],
                             k=key_cache,
@@ -1434,8 +1446,8 @@ class AiterFlashAttentionImpl(AttentionImpl):
                             block_table=attn_metadata.block_table[:num_decodes],
                             softcap=self.logits_soft_cap,
                             q_descale=None,
-                            k_descale=layer._k_scale.expand(descale_shape),
-                            v_descale=layer._v_scale.expand(descale_shape),
+                            k_descale=k_descale,
+                            v_descale=v_descale,
                             sinks=self.sinks,
                         )
                     return
@@ -1449,6 +1461,9 @@ class AiterFlashAttentionImpl(AttentionImpl):
                 use_unified_attention = self.head_size < _MIN_HEAD_SIZE_FOR_LL4MI
 
                 if use_unified_attention:
+                    k_descale, v_descale = self._get_kv_cache_descales(
+                        layer, num_decodes
+                    )
                     assert not rocm_aiter_ops.is_shuffle_kv_cache_enabled(), (
                         "unified_attention fallback with shuffle layout "
                         "is not supported yet."
@@ -1460,10 +1475,6 @@ class AiterFlashAttentionImpl(AttentionImpl):
                     decode_cu_seqlens_q = attn_metadata.query_start_loc[
                         : num_decodes + 1
                     ]
-                    descale_shape = (
-                        num_decodes,
-                        key_cache.shape[2],
-                    )
                     unified_attention(
                         q=query[:num_decode_tokens],
                         k=key_cache,
@@ -1480,8 +1491,8 @@ class AiterFlashAttentionImpl(AttentionImpl):
                         block_table=attn_metadata.block_table[:num_decodes],
                         softcap=self.logits_soft_cap,
                         q_descale=None,
-                        k_descale=layer._k_scale.expand(descale_shape),
-                        v_descale=layer._v_scale.expand(descale_shape),
+                        k_descale=k_descale,
+                        v_descale=v_descale,
                     )
                 elif rocm_aiter_ops.is_shuffle_kv_cache_enabled():
                     _, num_heads, head_size = query.shape

vllm/v1/core/sched/scheduler.py

@@ -123,6 +123,13 @@ class Scheduler(SchedulerInterface):
 
         # Scheduling constraints.
         self.max_num_running_reqs = self.scheduler_config.max_num_seqs
+        # Admission into RUNNING. Separate from max_num_running_reqs, which
+        # is the model-runner slot count.
+        self.max_num_active_reqs = (
+            self.scheduler_config.max_num_active_seqs
+            if self.scheduler_config.max_num_active_seqs is not None
+            else self.max_num_running_reqs
+        )
         self.max_num_scheduled_tokens = (
             self.scheduler_config.max_num_scheduled_tokens
             if self.scheduler_config.max_num_scheduled_tokens is not None
@@ -850,7 +857,7 @@ class Scheduler(SchedulerInterface):
                 # Paused streaming sessions (WAITING_FOR_STREAMING_REQ) are not
                 # in `running` but still hold a model-runner request slot.
                 num_running = len(self.running) + self.num_waiting_for_streaming_input
-                if num_running >= self.max_num_running_reqs:
+                if num_running >= self.max_num_active_reqs:
                     break
 
                 request_queue = self._select_waiting_queue_for_scheduling()

vllm/v1/metrics/perf.py

@@ -369,9 +369,10 @@ class BaseConfigParser(Parser):
 
         args.weight_byte_size = get_dtype_size(torch_dtype)
 
-        # FIXME: handle this better by parsing whether activations use
-        # bf16, fp32, etc...
-        args.activation_byte_size = 2
+        # Activations are produced in the model's compute dtype. Quantization
+        # overrides weight_byte_size below but leaves activations alone, so
+        # this stays keyed on the model dtype.
+        args.activation_byte_size = get_dtype_size(torch_dtype)
 
         args.dp_size = vllm_config.parallel_config.data_parallel_size
         args.tp_size = vllm_config.parallel_config.tensor_parallel_size