diff --git a/.buildkite/scripts/crcr-report.sh b/.buildkite/scripts/crcr-report.sh index e19146fddb..3b12d451cd 100644 --- a/.buildkite/scripts/crcr-report.sh +++ b/.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 diff --git a/.buildkite/test-amd.yaml b/.buildkite/test-amd.yaml index 78130fa3c7..4b72256a3f 100644 --- a/.buildkite/test-amd.yaml +++ b/.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 diff --git a/docker/Dockerfile.rock b/docker/Dockerfile.rock index 14ba523b07..32623ff1a9 100644 --- a/docker/Dockerfile.rock +++ b/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 diff --git a/rust/src/chat/src/backend/hf.rs b/rust/src/chat/src/backend/hf.rs index 5b5b8772c5..af9af8b51c 100644 --- a/rust/src/chat/src/backend/hf.rs +++ b/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("")); + 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("")); + 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( diff --git a/rust/src/chat/src/backend/mod.rs b/rust/src/chat/src/backend/mod.rs index eed4e95cbe..af4579159e 100644 --- a/rust/src/chat/src/backend/mod.rs +++ b/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, - /// 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, /// Maximum number of input items allowed per prompt for each modality. /// Unspecified modalities are unlimited. diff --git a/rust/src/chat/src/error.rs b/rust/src/chat/src/error.rs index a8304f17aa..db985df58d 100644 --- a/rust/src/chat/src/error.rs +++ b/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, diff --git a/rust/src/chat/src/lib.rs b/rust/src/chat/src/lib.rs index 9c59d14028..69d39ca78e 100644 --- a/rust/src/chat/src/lib.rs +++ b/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; diff --git a/rust/src/chat/src/reasoning.rs b/rust/src/chat/src/reasoning.rs new file mode 100644 index 0000000000..b5a5b17684 --- /dev/null +++ b/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 for EffortValue { + type Error = crate::error::Error; + + fn try_from(value: Value) -> Result { + 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 for EffortValue { + type Error = crate::error::Error; + + /// Construct a numeric effort, rejecting NaN and infinities. + fn try_from(value: f64) -> Result { + 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 { + 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 }, +} + +impl ReasoningControl { + /// Parse request controls; the typed effort field takes precedence over kwargs effort. + pub(crate) fn from_request(request: &ChatRequest) -> Result { + 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) -> Result { + Self::from_kwargs(kwargs, None) + } + + /// Select controls within one source, then derive its mode and optional effort. + fn from_kwargs( + kwargs: &HashMap, + typed_effort: Option<&EffortValue>, + ) -> Result { + // 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, + ) -> Result { + 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) -> 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 { + 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 { + 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") + ); + } + } +} diff --git a/rust/src/chat/src/renderer/deepseek.rs b/rust/src/chat/src/renderer/deepseek.rs index 01576a7c43..7892fd2497 100644 --- a/rust/src/chat/src/renderer/deepseek.rs +++ b/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 { - render_request_with_media_order(request, dialect).map(|(prompt, _)| prompt) +pub(super) fn render_request( + request: &ChatRequest, + dialect: DsDialect, + reasoning: &ReasoningControl, +) -> Result { + 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)> { - 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 { + 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" + )), } } diff --git a/rust/src/chat/src/renderer/deepseek_v32/encoding.rs b/rust/src/chat/src/renderer/deepseek_v32/encoding.rs index 4fc5f26530..74c7e4cbf2 100644 --- a/rust/src/chat/src/renderer/deepseek_v32/encoding.rs +++ b/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 { - let thinking_mode = match request.enable_thinking()?.unwrap_or(false) { +pub(super) fn render_request( + request: &ChatRequest, + reasoning: &ReasoningControl, +) -> Result { + let thinking_mode = match reasoning.is_enabled() { true => ThinkingMode::Thinking, false => ThinkingMode::Chat, }; diff --git a/rust/src/chat/src/renderer/deepseek_v32/mod.rs b/rust/src/chat/src/renderer/deepseek_v32/mod.rs index 03c2765d19..b4ac1bae11 100644 --- a/rust/src/chat/src/renderer/deepseek_v32/mod.rs +++ b/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, +} impl DeepSeekV32ChatRenderer { /// Create the dedicated DeepSeek V3.2 renderer. - pub fn new() -> Self { - Self + pub fn new(default_template_kwargs: HashMap) -> Self { + Self { + default_template_kwargs, + } } } impl ChatRenderer for DeepSeekV32ChatRenderer { fn render(&self, request: &ChatRequest) -> Result { 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), }) } } diff --git a/rust/src/chat/src/renderer/deepseek_v32/tests.rs b/rust/src/chat/src/renderer/deepseek_v32/tests.rs index 56c03a92ca..dd9ec13ddc 100644 --- a/rust/src/chat/src/renderer/deepseek_v32/tests.rs +++ b/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(""), + enabled + ); + assert_eq!( + rendered.effective_template_kwargs["enable_thinking"], + enabled + ); + } +} + fn render_result(request: &ChatRequest) -> Result { - 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, diff --git a/rust/src/chat/src/renderer/deepseek_v4/mod.rs b/rust/src/chat/src/renderer/deepseek_v4/mod.rs index 54b111e505..136938fbc0 100644 --- a/rust/src/chat/src/renderer/deepseek_v4/mod.rs +++ b/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, +} impl DeepSeekV4ChatRenderer { - pub fn new() -> Self { - Self + pub fn new(default_template_kwargs: HashMap) -> Self { + Self { + default_template_kwargs, + } } } impl ChatRenderer for DeepSeekV4ChatRenderer { fn render(&self, request: &ChatRequest) -> Result { 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), }) } } diff --git a/rust/src/chat/src/renderer/deepseek_v4/tests.rs b/rust/src/chat/src/renderer/deepseek_v4/tests.rs index 9404b54296..21a33ec2d5 100644 --- a/rust/src/chat/src/renderer/deepseek_v4/tests.rs +++ b/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) -> 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(""), 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｜>")); } #[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); diff --git a/rust/src/chat/src/renderer/deepseek_v41/mod.rs b/rust/src/chat/src/renderer/deepseek_v41/mod.rs index 7737f7ebda..54544ec383 100644 --- a/rust/src/chat/src/renderer/deepseek_v41/mod.rs +++ b/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, +} impl DeepSeekV41ChatRenderer { - pub fn new() -> Self { - Self + pub fn new(default_template_kwargs: HashMap) -> Self { + Self { + default_template_kwargs, + } } } impl ChatRenderer for DeepSeekV41ChatRenderer { fn render(&self, request: &ChatRequest) -> Result { 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), }) } } diff --git a/rust/src/chat/src/renderer/deepseek_v41/tests.rs b/rust/src/chat/src/renderer/deepseek_v41/tests.rs index 3d621c2246..62cf6d2024 100644 --- a/rust/src/chat/src/renderer/deepseek_v41/tests.rs +++ b/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("")); + 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("")); + 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｜>"]; 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)); } diff --git a/rust/src/chat/src/renderer/harmony/mod.rs b/rust/src/chat/src/renderer/harmony/mod.rs index 71431cc92b..83b59e9198 100644 --- a/rust/src/chat/src/renderer/harmony/mod.rs +++ b/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, } 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 { + pub fn new(default_template_kwargs: HashMap) -> Result { 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, system_start_date: impl Into, use_system_instructions: bool, ) -> Result { 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> { + fn render_token_ids( + &self, + request: &ChatRequest, + effort: HarmonyReasoningEffort, + ) -> Result> { 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 { + 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> { +fn to_harmony_messages( + request: &ChatRequest, + options: &Options, + effort: HarmonyReasoningEffort, +) -> Result> { 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, leading_developer_tools: Option<&[ChatTool]>, options: &Options, + effort: HarmonyReasoningEffort, ) -> Result> { 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, + reasoning_effort: Option, system_start_date: &str, ) -> Result { 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 { .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 { - 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")) ))), } } diff --git a/rust/src/chat/src/renderer/harmony/tests.rs b/rust/src/chat/src/renderer/harmony/tests.rs index 4849c93223..609c1f1159 100644 --- a/rust/src/chat/src/renderer/harmony/tests.rs +++ b/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 { @@ -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 { diff --git a/rust/src/chat/src/renderer/hf/mod.rs b/rust/src/chat/src/renderer/hf/mod.rs index ba2b76d100..911a73928a 100644 --- a/rust/src/chat/src/renderer/hf/mod.rs +++ b/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, + request: &ChatRequest, +) -> HashMap { + 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![ diff --git a/rust/src/chat/src/renderer/inkling/mod.rs b/rust/src/chat/src/renderer/inkling/mod.rs index e8afe66396..9feee371aa 100644 --- a/rust/src/chat/src/renderer/inkling/mod.rs +++ b/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, } #[derive(Debug, Clone, Copy)] @@ -78,9 +81,16 @@ impl InklingSpecialTokenIds { } impl InklingChatRenderer { - pub fn new(tokenizer: DynTokenizer) -> Result { + pub fn new( + tokenizer: DynTokenizer, + default_template_kwargs: HashMap, + ) -> Result { 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, text: &str) -> Result<()> { @@ -150,11 +160,6 @@ impl InklingChatRenderer { } fn write_reasoning_effort(&self, out: &mut Vec, 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 { +fn resolve_reasoning(control: ReasoningControl) -> Result { + 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 { 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 { "xhigh" | "max" => Some(0.99), _ => None, }, - Value::Number(number) => number.as_f64(), - _ => None, + EffortValue::Number(number) => number.as_f64(), } } diff --git a/rust/src/chat/src/renderer/inkling/tests.rs b/rust/src/chat/src/renderer/inkling/tests.rs index 1c4cd7236d..d0dd3055e4 100644 --- a/rust/src/chat/src/renderer/inkling/tests.rs +++ b/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 { @@ -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::().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] diff --git a/rust/src/chat/src/renderer/kimi_k3/encoding.rs b/rust/src/chat/src/renderer/kimi_k3/encoding.rs index 88b4e3ae61..b434d72c53 100644 --- a/rust/src/chat/src/renderer/kimi_k3/encoding.rs +++ b/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, Vec)> { - 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 { - 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 { - 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, +) -> Result { + 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 { diff --git a/rust/src/chat/src/renderer/kimi_k3/mod.rs b/rust/src/chat/src/renderer/kimi_k3/mod.rs index 90750eeac4..c39ea2985a 100644 --- a/rust/src/chat/src/renderer/kimi_k3/mod.rs +++ b/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, } impl KimiK3ChatRenderer { /// Create a Kimi K3 renderer. - pub fn new(tokenizer: DynTokenizer) -> Self { - Self { tokenizer } + pub fn new(tokenizer: DynTokenizer, default_template_kwargs: HashMap) -> Self { + Self { + tokenizer, + default_template_kwargs, + } } } impl ChatRenderer for KimiK3ChatRenderer { fn render(&self, request: &ChatRequest) -> Result { 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, }) } } diff --git a/rust/src/chat/src/renderer/kimi_k3/tests.rs b/rust/src/chat/src/renderer/kimi_k3/tests.rs index d0ecbe44d3..5aadb3d403 100644 --- a/rust/src/chat/src/renderer/kimi_k3/tests.rs +++ b/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 { - 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); diff --git a/rust/src/chat/src/renderer/mod.rs b/rust/src/chat/src/renderer/mod.rs index bc1831c375..347ec43d03 100644 --- a/rust/src/chat/src/renderer/mod.rs +++ b/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; - -/// 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, - request: &ChatRequest, -) -> HashMap { - 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 { - effective_template_kwargs(&HashMap::new(), request) -} diff --git a/rust/src/chat/src/renderer/test_utils.rs b/rust/src/chat/src/renderer/test_utils.rs index 1090a4039f..2c59ed7a22 100644 --- a/rust/src/chat/src/renderer/test_utils.rs +++ b/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, messages: Vec, add_generation_prompt: Option, - reasoning_effort: Option, + reasoning_effort: Option, /// Standard response format passed to model-specific renderers. #[serde(default)] response_format: Option, diff --git a/rust/src/chat/src/request.rs b/rust/src/chat/src/request.rs index 50e00bbc43..e3bbcaa2aa 100644 --- a/rust/src/chat/src/request.rs +++ b/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, - /// Effort level exposed to chat templates for reasoning models. - pub reasoning_effort: Option, + /// 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, /// 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> { - 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> { 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") - )); - } } diff --git a/rust/src/engine-core-client/src/protocol/mod.rs b/rust/src/engine-core-client/src/protocol/mod.rs index 23bd09d8d8..8f4ee06127 100644 --- a/rust/src/engine-core-client/src/protocol/mod.rs +++ b/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; diff --git a/rust/src/engine-core-client/src/protocol/output.rs b/rust/src/engine-core-client/src/protocol/output.rs index 85de8c4e39..c8a8705435 100644 --- a/rust/src/engine-core-client/src/protocol/output.rs +++ b/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: /// +#[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")] #[serde(default)] outputs: Vec, #[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 { diff --git a/rust/src/engine-core-client/src/protocol/serde_utils.rs b/rust/src/engine-core-client/src/protocol/serde_utils.rs new file mode 100644 index 0000000000..708353e42e --- /dev/null +++ b/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(deserializer: D) -> Result + where + D: Deserializer<'de>, + { + struct ArrayVisitor(PhantomData); + + impl<'de, T: Deserialize<'de>> Visitor<'de> for ArrayVisitor { + type Value = T; + + fn expecting(&self, formatter: &mut fmt::Formatter<'_>) -> fmt::Result { + formatter.write_str("an array-like struct") + } + + fn visit_seq(self, mut seq: A) -> Result + where + A: SeqAccess<'de>, + { + let value = T::deserialize(SeqAccessDeserializer::new(&mut seq))?; + while seq.next_element::()?.is_some() {} + Ok(value) + } + } + + deserializer.deserialize_seq(ArrayVisitor(PhantomData)) + } +} diff --git a/rust/src/engine-core-client/src/tests/client.rs b/rust/src/engine-core-client/src/tests/client.rs index 3db573d2aa..739c34d6b7 100644 --- a/rust/src/engine-core-client/src/tests/client.rs +++ b/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 { diff --git a/rust/src/engine-core-client/src/tests/python_compat.py b/rust/src/engine-core-client/src/tests/python_compat.py index 0cebca972d..5013f9e2a5 100755 --- a/rust/src/engine-core-client/src/tests/python_compat.py +++ b/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 = [ [ "(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; diff --git a/rust/src/server/src/routes/openai/chat_completions/types.rs b/rust/src/server/src/routes/openai/chat_completions/types.rs index a371462694..13c7c7233f 100644 --- a/rust/src/server/src/routes/openai/chat_completions/types.rs +++ b/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. diff --git a/rust/src/server/src/routes/openai/chat_completions/validate.rs b/rust/src/server/src/routes/openai/chat_completions/validate.rs index 7260155799..1bfc1d9a09 100644 --- a/rust/src/server/src/routes/openai/chat_completions/validate.rs +++ b/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; diff --git a/rust/src/server/src/routes/openai/utils/types.rs b/rust/src/server/src/routes/openai/utils/types.rs index e9aab2d2b2..5554b5d366 100644 --- a/rust/src/server/src/routes/openai/utils/types.rs +++ b/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 // ============================================================================ diff --git a/tests/kernels/moe/modular_kernel_tools/common.py b/tests/kernels/moe/modular_kernel_tools/common.py index 8adfa6588c..3ec522c8ec 100644 --- a/tests/kernels/moe/modular_kernel_tools/common.py +++ b/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 ) diff --git a/tests/kernels/moe/modular_kernel_tools/mk_objects.py b/tests/kernels/moe/modular_kernel_tools/mk_objects.py index 7f1924bac5..e0f37f5436 100644 --- a/tests/kernels/moe/modular_kernel_tools/mk_objects.py +++ b/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], diff --git a/tests/kernels/moe/test_gpt_oss_triton_kernels.py b/tests/kernels/moe/test_gpt_oss_triton_kernels.py index 4afcba706c..86028a7feb 100644 --- a/tests/kernels/moe/test_gpt_oss_triton_kernels.py +++ b/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 diff --git a/tests/kernels/moe/test_modular_kernel_combinations.py b/tests/kernels/moe/test_modular_kernel_combinations.py index 0c0e1d61f9..fc1d8ac07b 100644 --- a/tests/kernels/moe/test_modular_kernel_combinations.py +++ b/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()) diff --git a/tests/test_config.py b/tests/test_config.py index ce2580934a..3d6a17f7a2 100644 --- a/tests/test_config.py +++ b/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"), [ diff --git a/tests/v1/attention/test_attention_backends.py b/tests/v1/attention/test_attention_backends.py index 66a1c5c088..b02c85f682 100644 --- a/tests/v1/attention/test_attention_backends.py +++ b/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", [ diff --git a/tests/v1/core/prefix_cache/test_hybrid_prefix_cache_hit_rate.py b/tests/v1/core/prefix_cache/test_hybrid_prefix_cache_hit_rate.py new file mode 100644 index 0000000000..c75591b697 --- /dev/null +++ b/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)) diff --git a/tests/v1/core/test_scheduler.py b/tests/v1/core/test_scheduler.py index 3239ec25a0..3add7beda0 100644 --- a/tests/v1/core/test_scheduler.py +++ b/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) diff --git a/tests/v1/core/utils.py b/tests/v1/core/utils.py index 28d1f71ac4..f69b10ced3 100644 --- a/tests/v1/core/utils.py +++ b/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, diff --git a/tests/v1/metrics/test_perf_metrics.py b/tests/v1/metrics/test_perf_metrics.py index 3868cc772e..c0997ec459 100644 --- a/tests/v1/metrics/test_perf_metrics.py +++ b/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(): diff --git a/vllm/config/parallel.py b/vllm/config/parallel.py index f66af2f30b..f0a2b7e457 100644 --- a/vllm/config/parallel.py +++ b/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", diff --git a/vllm/config/scheduler.py b/vllm/config/scheduler.py index d3264b2d80..321c76fe30 100644 --- a/vllm/config/scheduler.py +++ b/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 " diff --git a/vllm/config/speculative.py b/vllm/config/speculative.py index 63f729a832..ef9a16f903 100644 --- a/vllm/config/speculative.py +++ b/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, diff --git a/vllm/engine/arg_utils.py b/vllm/engine/arg_utils.py index ccb30a2597..d637391cd3 100644 --- a/vllm/engine/arg_utils.py +++ b/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, diff --git a/vllm/model_executor/layers/fused_moe/all2all_utils.py b/vllm/model_executor/layers/fused_moe/all2all_utils.py index 75ea91bfca..775f595bfa 100644 --- a/vllm/model_executor/layers/fused_moe/all2all_utils.py +++ b/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. diff --git a/vllm/model_executor/layers/fused_moe/config.py b/vllm/model_executor/layers/fused_moe/config.py index ccca24cc6f..9a2f8229b1 100644 --- a/vllm/model_executor/layers/fused_moe/config.py +++ b/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 " diff --git a/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a16_moe.py b/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a16_moe.py index 757d1d8295..f6afd6fbfa 100644 --- a/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a16_moe.py +++ b/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 diff --git a/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a8_moe.py b/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a8_moe.py index 0691d9fed4..b327cbeb59 100644 --- a/vllm/model_executor/layers/fused_moe/experts/aiter_mxfp4_w4a8_moe.py +++ b/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, diff --git a/vllm/model_executor/layers/fused_moe/experts/gpt_oss_triton_kernels_moe.py b/vllm/model_executor/layers/fused_moe/experts/gpt_oss_triton_kernels_moe.py index e2b4987bd0..27fdcf8bb2 100644 --- a/vllm/model_executor/layers/fused_moe/experts/gpt_oss_triton_kernels_moe.py +++ b/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 ( diff --git a/vllm/model_executor/layers/fused_moe/oracle/mxfp4.py b/vllm/model_executor/layers/fused_moe/oracle/mxfp4.py index 912b713918..ce19ba39bf 100644 --- a/vllm/model_executor/layers/fused_moe/oracle/mxfp4.py +++ b/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 diff --git a/vllm/model_executor/layers/fused_moe/prepare_finalize/deepep_v2.py b/vllm/model_executor/layers/fused_moe/prepare_finalize/deepep_v2.py index 8bd215423d..a584ab5f52 100644 --- a/vllm/model_executor/layers/fused_moe/prepare_finalize/deepep_v2.py +++ b/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() diff --git a/vllm/model_executor/layers/fused_moe/prepare_finalize/mori.py b/vllm/model_executor/layers/fused_moe/prepare_finalize/mori.py index 35e15325b4..75b97c69aa 100644 --- a/vllm/model_executor/layers/fused_moe/prepare_finalize/mori.py +++ b/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, diff --git a/vllm/model_executor/layers/quantization/utils/mxfp4_utils.py b/vllm/model_executor/layers/quantization/utils/mxfp4_utils.py index 1bd6b3f931..a1d0501775 100644 --- a/vllm/model_executor/layers/quantization/utils/mxfp4_utils.py +++ b/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(): diff --git a/vllm/utils/import_utils.py b/vllm/utils/import_utils.py index d6a2aabbb4..4e90367a3e 100644 --- a/vllm/utils/import_utils.py +++ b/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. diff --git a/vllm/v1/attention/backends/rocm_aiter_fa.py b/vllm/v1/attention/backends/rocm_aiter_fa.py index 1615f883bc..4ce8b23d31 100644 --- a/vllm/v1/attention/backends/rocm_aiter_fa.py +++ b/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 diff --git a/vllm/v1/core/sched/scheduler.py b/vllm/v1/core/sched/scheduler.py index 19ea411aa1..22c231cc43 100644 --- a/vllm/v1/core/sched/scheduler.py +++ b/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() diff --git a/vllm/v1/metrics/perf.py b/vllm/v1/metrics/perf.py index dd046a8250..1cd8fa79f1 100644 --- a/vllm/v1/metrics/perf.py +++ b/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