diff --git a/README.md b/README.md
index ca31a8cce9..5681be7612 100644
--- a/README.md
+++ b/README.md
@@ -1,110 +1,5 @@
-
-
-
-
-
-
+# VLLM Backport
-
-Easy, fast, and cheap LLM serving for everyone
-
+A VLLM fork that focuses on running Deepseek V4 Flash 0731 on Ampere at this moment.
-
-| Documentation Blog Paper Twitter/X User Forum Developer Slack
-
-π₯ We have built a vLLM website to help you get started with vLLM. Please visit [vllm.ai](https://vllm.ai) to learn more.
-For events, please visit [vllm.ai/events](https://vllm.ai/events) to join us.
-
----
-
-## About
-
-vLLM is a fast and easy-to-use library for LLM inference and serving.
-
-Originally developed in the [Sky Computing Lab](https://sky.cs.berkeley.edu) at UC Berkeley, vLLM has grown into one of the most active open-source AI projects built and maintained by a diverse community of many dozens of academic institutions and companies from over 2000 contributors.
-
-vLLM is fast with:
-
-- State-of-the-art serving throughput
-- Efficient management of attention key and value memory with [**PagedAttention**](https://blog.vllm.ai/2023/06/20/vllm.html)
-- Continuous batching of incoming requests, chunked prefill, prefix caching
-- Fast and flexible model execution with piecewise and full CUDA/HIP graphs
-- Quantization: FP8, MXFP8/MXFP4, NVFP4, INT8, INT4, GPTQ/AWQ, GGUF, compressed-tensors, ModelOpt, TorchAO, and [more](https://docs.vllm.ai/en/latest/features/quantization/index.html)
-- Optimized attention kernels including FlashAttention, FlashInfer, TRTLLM-GEN, FlashMLA, and Triton
-- Optimized GEMM/MoE kernels for various precisions using CUTLASS, TRTLLM-GEN, CuTeDSL
-- Speculative decoding including n-gram, suffix, EAGLE, DFlash
-- Automatic kernel generation and graph-level transformations using torch.compile
-- Disaggregated prefill, decode, and encode
-
-vLLM is flexible and easy to use with:
-
-- Seamless integration with popular Hugging Face models
-- High-throughput serving with various decoding algorithms, including *parallel sampling*, *beam search*, and more
-- Tensor, pipeline, data, expert, and context parallelism for distributed inference
-- Streaming outputs
-- Generation of structured outputs using xgrammar or guidance
-- Tool calling and reasoning parsers
-- OpenAI-compatible API server, plus Anthropic Messages API and gRPC support
-- Efficient multi-LoRA support for dense and MoE layers
-- Support for NVIDIA GPUs, AMD GPUs, Intel GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
-
-vLLM seamlessly supports 200+ model architectures on Hugging Face, including:
-
-- Decoder-only LLMs (e.g., Llama, Qwen, Gemma)
-- Mixture-of-Expert LLMs (e.g., Mixtral, DeepSeek-V3, Qwen-MoE, GPT-OSS)
-- Hybrid attention and state-space models (e.g., Mamba, Qwen3.5)
-- Multi-modal models (e.g., LLaVA, Qwen-VL, Pixtral)
-- Embedding and retrieval models (e.g., E5-Mistral, GTE, ColBERT)
-- Reward and classification models (e.g., Qwen-Math)
-
-Find the full list of supported models [here](https://docs.vllm.ai/en/latest/models/supported_models.html).
-
-## Getting Started
-
-Install vLLM with [`uv`](https://docs.astral.sh/uv/) (recommended) or `pip`:
-
-```bash
-uv pip install vllm
-```
-
-Or [build from source](https://docs.vllm.ai/en/latest/getting_started/installation/gpu/index.html#build-wheel-from-source) for development.
-
-Visit our [documentation](https://docs.vllm.ai/en/latest/) to learn more.
-
-- [Installation](https://docs.vllm.ai/en/latest/getting_started/installation.html)
-- [Quickstart](https://docs.vllm.ai/en/latest/getting_started/quickstart.html)
-- [List of Supported Models](https://docs.vllm.ai/en/latest/models/supported_models.html)
-
-## Contributing
-
-We welcome and value any contributions and collaborations.
-Please check out [Contributing to vLLM](https://docs.vllm.ai/en/latest/contributing/index.html) for how to get involved.
-
-## Citation
-
-If you use vLLM for your research, please cite our [paper](https://arxiv.org/abs/2309.06180):
-
-```bibtex
-@inproceedings{kwon2023efficient,
- title={Efficient Memory Management for Large Language Model Serving with PagedAttention},
- author={Woosuk Kwon and Zhuohan Li and Siyuan Zhuang and Ying Sheng and Lianmin Zheng and Cody Hao Yu and Joseph E. Gonzalez and Hao Zhang and Ion Stoica},
- booktitle={Proceedings of the ACM SIGOPS 29th Symposium on Operating Systems Principles},
- year={2023}
-}
-```
-
-## Contact Us
-
-
-- For technical questions and feature requests, please use GitHub [Issues](https://github.com/vllm-project/vllm/issues)
-- For discussing with fellow users, please use the [vLLM Forum](https://discuss.vllm.ai)
-- For coordinating contributions and development, please use [Slack](https://slack.vllm.ai)
-- For security disclosures, please use GitHub's [Security Advisories](https://github.com/vllm-project/vllm/security/advisories) feature
-- For collaborations and partnerships, please contact us at [collaboration@vllm.ai](mailto:collaboration@vllm.ai)
-
-
-## Media Kit
-
-- If you wish to use vLLM's logo, please refer to [our media kit repo](https://github.com/vllm-project/media-kit)
+Currently achieving 1783 tps prefill and 355 tps decoding on 8xA6000 and this should also work on A100.
diff --git a/README.vllm.md b/README.vllm.md
new file mode 100644
index 0000000000..ca31a8cce9
--- /dev/null
+++ b/README.vllm.md
@@ -0,0 +1,110 @@
+
+
+
+
+
+
+
+
+Easy, fast, and cheap LLM serving for everyone
+
+
+
+| Documentation Blog Paper Twitter/X User Forum Developer Slack
+
+π₯ We have built a vLLM website to help you get started with vLLM. Please visit [vllm.ai](https://vllm.ai) to learn more.
+For events, please visit [vllm.ai/events](https://vllm.ai/events) to join us.
+
+---
+
+## About
+
+vLLM is a fast and easy-to-use library for LLM inference and serving.
+
+Originally developed in the [Sky Computing Lab](https://sky.cs.berkeley.edu) at UC Berkeley, vLLM has grown into one of the most active open-source AI projects built and maintained by a diverse community of many dozens of academic institutions and companies from over 2000 contributors.
+
+vLLM is fast with:
+
+- State-of-the-art serving throughput
+- Efficient management of attention key and value memory with [**PagedAttention**](https://blog.vllm.ai/2023/06/20/vllm.html)
+- Continuous batching of incoming requests, chunked prefill, prefix caching
+- Fast and flexible model execution with piecewise and full CUDA/HIP graphs
+- Quantization: FP8, MXFP8/MXFP4, NVFP4, INT8, INT4, GPTQ/AWQ, GGUF, compressed-tensors, ModelOpt, TorchAO, and [more](https://docs.vllm.ai/en/latest/features/quantization/index.html)
+- Optimized attention kernels including FlashAttention, FlashInfer, TRTLLM-GEN, FlashMLA, and Triton
+- Optimized GEMM/MoE kernels for various precisions using CUTLASS, TRTLLM-GEN, CuTeDSL
+- Speculative decoding including n-gram, suffix, EAGLE, DFlash
+- Automatic kernel generation and graph-level transformations using torch.compile
+- Disaggregated prefill, decode, and encode
+
+vLLM is flexible and easy to use with:
+
+- Seamless integration with popular Hugging Face models
+- High-throughput serving with various decoding algorithms, including *parallel sampling*, *beam search*, and more
+- Tensor, pipeline, data, expert, and context parallelism for distributed inference
+- Streaming outputs
+- Generation of structured outputs using xgrammar or guidance
+- Tool calling and reasoning parsers
+- OpenAI-compatible API server, plus Anthropic Messages API and gRPC support
+- Efficient multi-LoRA support for dense and MoE layers
+- Support for NVIDIA GPUs, AMD GPUs, Intel GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
+
+vLLM seamlessly supports 200+ model architectures on Hugging Face, including:
+
+- Decoder-only LLMs (e.g., Llama, Qwen, Gemma)
+- Mixture-of-Expert LLMs (e.g., Mixtral, DeepSeek-V3, Qwen-MoE, GPT-OSS)
+- Hybrid attention and state-space models (e.g., Mamba, Qwen3.5)
+- Multi-modal models (e.g., LLaVA, Qwen-VL, Pixtral)
+- Embedding and retrieval models (e.g., E5-Mistral, GTE, ColBERT)
+- Reward and classification models (e.g., Qwen-Math)
+
+Find the full list of supported models [here](https://docs.vllm.ai/en/latest/models/supported_models.html).
+
+## Getting Started
+
+Install vLLM with [`uv`](https://docs.astral.sh/uv/) (recommended) or `pip`:
+
+```bash
+uv pip install vllm
+```
+
+Or [build from source](https://docs.vllm.ai/en/latest/getting_started/installation/gpu/index.html#build-wheel-from-source) for development.
+
+Visit our [documentation](https://docs.vllm.ai/en/latest/) to learn more.
+
+- [Installation](https://docs.vllm.ai/en/latest/getting_started/installation.html)
+- [Quickstart](https://docs.vllm.ai/en/latest/getting_started/quickstart.html)
+- [List of Supported Models](https://docs.vllm.ai/en/latest/models/supported_models.html)
+
+## Contributing
+
+We welcome and value any contributions and collaborations.
+Please check out [Contributing to vLLM](https://docs.vllm.ai/en/latest/contributing/index.html) for how to get involved.
+
+## Citation
+
+If you use vLLM for your research, please cite our [paper](https://arxiv.org/abs/2309.06180):
+
+```bibtex
+@inproceedings{kwon2023efficient,
+ title={Efficient Memory Management for Large Language Model Serving with PagedAttention},
+ author={Woosuk Kwon and Zhuohan Li and Siyuan Zhuang and Ying Sheng and Lianmin Zheng and Cody Hao Yu and Joseph E. Gonzalez and Hao Zhang and Ion Stoica},
+ booktitle={Proceedings of the ACM SIGOPS 29th Symposium on Operating Systems Principles},
+ year={2023}
+}
+```
+
+## Contact Us
+
+
+- For technical questions and feature requests, please use GitHub [Issues](https://github.com/vllm-project/vllm/issues)
+- For discussing with fellow users, please use the [vLLM Forum](https://discuss.vllm.ai)
+- For coordinating contributions and development, please use [Slack](https://slack.vllm.ai)
+- For security disclosures, please use GitHub's [Security Advisories](https://github.com/vllm-project/vllm/security/advisories) feature
+- For collaborations and partnerships, please contact us at [collaboration@vllm.ai](mailto:collaboration@vllm.ai)
+
+
+## Media Kit
+
+- If you wish to use vLLM's logo, please refer to [our media kit repo](https://github.com/vllm-project/media-kit)
diff --git a/requirements/test/cuda.txt b/requirements/test/cuda.txt
index 192f3ac5a6..ca2de20834 100644
--- a/requirements/test/cuda.txt
+++ b/requirements/test/cuda.txt
@@ -498,7 +498,7 @@ mbstrdecoder==1.1.3
# dataproperty
# pytablewriter
# typepy
-mcp==1.28.1
+mcp==2.0.0
# via -r requirements/test/../common.txt
mdurl==0.1.2
# via markdown-it-py
diff --git a/tests/distributed/test_hier_all_reduce.py b/tests/distributed/test_hier_all_reduce.py
new file mode 100644
index 0000000000..a951897a5f
--- /dev/null
+++ b/tests/distributed/test_hier_all_reduce.py
@@ -0,0 +1,78 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import ray
+import torch
+import torch.distributed as dist
+
+from vllm.distributed.device_communicators.hier_all_reduce import (
+ _TWO_SHOT_MIN_ELEMS,
+ HierarchicalAllReduce,
+)
+from vllm.distributed.parallel_state import get_tp_group
+
+from ..utils import init_test_distributed_environment, multi_process_parallel
+
+# Straddle the one-shot/two-shot dispatch threshold in both directions, and
+# include a size that is not a multiple of the CTA tiling to exercise the
+# masked tails.
+TEST_SIZES = [
+ 1024,
+ _TWO_SHOT_MIN_ELEMS // 2,
+ _TWO_SHOT_MIN_ELEMS,
+ _TWO_SHOT_MIN_ELEMS * 2 + 512,
+]
+
+
+@ray.remote(num_gpus=1, max_calls=1)
+def hier_allreduce_matches_nccl(
+ monkeypatch: pytest.MonkeyPatch,
+ tp_size,
+ pp_size,
+ rank,
+ distributed_init_port,
+):
+ with monkeypatch.context() as m:
+ m.delenv("CUDA_VISIBLE_DEVICES", raising=False)
+ m.delenv("HIP_VISIBLE_DEVICES", raising=False)
+ device = torch.device(f"cuda:{rank}")
+ torch.accelerator.set_device_index(device)
+ init_test_distributed_environment(tp_size, pp_size, rank, distributed_init_port)
+
+ tp_group = get_tp_group()
+ group = tp_group.cpu_group
+ # Islands are group-local rank indices; split the TP group in half so
+ # the two halves stand in for the two PCIe islands.
+ half = tp_size // 2
+ islands = [list(range(half)), list(range(half, tp_size))]
+ comm = HierarchicalAllReduce(group, device, islands)
+
+ for numel in TEST_SIZES:
+ inp = torch.randn(numel, dtype=torch.bfloat16, device=device)
+ ref = inp.clone()
+ dist.all_reduce(ref, group=tp_group.device_group)
+ assert comm.should_use(inp)
+ # Run twice: the flag protocol alternates buffer halves by
+ # sequence-token parity, so the second call takes the other half.
+ for _ in range(2):
+ out = comm.all_reduce(inp)
+ torch.cuda.synchronize()
+ # Reduction order differs from NCCL's, so compare within the
+ # dtype's tolerance rather than bit-exactly.
+ torch.testing.assert_close(out, ref, atol=2e-2, rtol=2e-2)
+
+
+@pytest.mark.parametrize("tp_size", [2, 4])
+@pytest.mark.parametrize("pipeline_parallel_size", [1])
+def test_hier_all_reduce(
+ monkeypatch: pytest.MonkeyPatch,
+ tp_size,
+ pipeline_parallel_size,
+):
+ world_size = tp_size * pipeline_parallel_size
+ if world_size > torch.accelerator.device_count():
+ pytest.skip("Not enough GPUs to run the test.")
+ multi_process_parallel(
+ monkeypatch, tp_size, pipeline_parallel_size, hier_allreduce_matches_nccl
+ )
diff --git a/tests/parser/engine/test_deepseek_v32.py b/tests/parser/engine/test_deepseek_v32.py
index 7825f5add8..55ceb0ba41 100644
--- a/tests/parser/engine/test_deepseek_v32.py
+++ b/tests/parser/engine/test_deepseek_v32.py
@@ -22,6 +22,7 @@
DSML_INVOKE_END,
DSML_INVOKE_NAME_END,
DSML_INVOKE_PREFIX,
+ DSML_TOOL_START,
)
from vllm.parser.deepseek_v32 import (
DSML_FUNC_END,
@@ -67,13 +68,7 @@ def _make_tool(name, properties):
)
-@pytest.fixture
-def mock_tokenizer():
- return make_mock_tokenizer({})
-
-
-@pytest.fixture
-def mock_request():
+def _request_without_tools():
from unittest.mock import MagicMock
from vllm.entrypoints.openai.chat_completion.protocol import (
@@ -86,6 +81,16 @@ def mock_request():
return req
+@pytest.fixture
+def mock_tokenizer():
+ return make_mock_tokenizer({})
+
+
+@pytest.fixture
+def mock_request():
+ return _request_without_tools()
+
+
# ββ Non-streaming extraction ββββββββββββββββββββββββββββββββββββββββ
@@ -130,6 +135,67 @@ def test_content_before_tool_call(self, mock_tokenizer, mock_request):
assert result.content is not None
assert "Let me check" in result.content
+ def test_missing_func_start_orphan_invoke(self, mock_tokenizer, mock_request):
+ """Orphan invoke without the <ο½DSMLο½function_calls> wrapper is
+ still parsed as a tool call when the request declared the tool
+ (see gh-48931)."""
+ tool = _make_tool("get_weather", {"city": {"type": "string"}})
+ mock_request.tools = [tool]
+ text = _invoke("get_weather", _param("city", "true", "SF")) + DSML_FUNC_END
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[tool])
+ result = parser.extract_tool_calls(text, mock_request)
+ assert result.tools_called
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"city": "SF"}
+ assert result.content is None
+
+ def test_orphan_invoke_without_declared_tools_stays_content(
+ self, mock_tokenizer, mock_request
+ ):
+ """A request that declared no tools can never accept a recovered
+ name, so the orphan invoke stays plain content."""
+ text = _invoke("get_weather", _param("city", "true", "SF")) + DSML_FUNC_END
+ parser = DeepSeekV32Parser(mock_tokenizer)
+ result = parser.extract_tool_calls(text, mock_request)
+ assert not result.tools_called
+ assert result.tool_calls == []
+ assert result.content == text
+
+ def test_unclosed_foreign_wrapper_then_native_call(
+ self, mock_tokenizer, mock_request
+ ):
+ """A foreign wrapper that never closes must not disable native
+ tool parsing: the token backed function_calls wrapper still
+ wins."""
+ text = (
+ DSML_TOOL_START
+ + "\nStray foreign text.\n"
+ + _func_calls(_invoke("get_weather", _param("city", "true", "SF")))
+ )
+ parser = DeepSeekV32Parser(mock_tokenizer)
+ result = parser.extract_tool_calls(text, mock_request)
+ assert result.tools_called
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"city": "SF"}
+ assert "Stray foreign text." in result.content
+
+ def test_foreign_tool_calls_wrapper_rejected(self, mock_tokenizer, mock_request):
+ """An invoke inside the V4-style tool_calls wrapper stays plain
+ content: the orphan fallback must not fire inside a foreign
+ wrapper."""
+ text = _func_calls(
+ _invoke("get_weather", _param("city", "true", "SF")),
+ ).replace("function_calls", "tool_calls")
+ parser = DeepSeekV32Parser(mock_tokenizer)
+ result = parser.extract_tool_calls(text, mock_request)
+ assert not result.tools_called
+ assert result.tool_calls == []
+ assert result.content == text
+
def test_non_string_params_json_parsed(self, mock_tokenizer, mock_request):
text = _func_calls(
_invoke(
@@ -159,6 +225,418 @@ def test_wrapper_unwrapping(self, mock_tokenizer, mock_request):
assert args == {"location": "Beijing"}
+# ββ Orphan invoke name validation ββββββββββββββββββββββββββββββββββββ
+
+
+class TestOrphanInvokeNameValidation:
+ """Recovered (orphan) invokes must carry a plausible tool name.
+
+ Mirrors the V4 coverage: when the request declares tools, the
+ (CONTENT, INVOKE_PREFIX) recovery path only commits to a tool call
+ if the parsed name is one of the declared functions; otherwise the
+ consumed text is re-emitted as plain content.
+ """
+
+ @pytest.fixture
+ def weather_tool(self):
+ return _make_tool("get_weather", {"city": {"type": "string"}})
+
+ def test_declared_name_recovered(self, mock_tokenizer, mock_request, weather_tool):
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ text = _invoke("get_weather", _param("city", "true", "SF")) + DSML_FUNC_END
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"city": "SF"}
+ assert result.content is None
+
+ def test_undeclared_name_stays_content(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ text = (
+ "Quoting "
+ + DSML_INVOKE_PREFIX
+ + "made_up_tool"
+ + DSML_INVOKE_NAME_END
+ + " literally."
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert not result.tools_called
+ assert result.tool_calls == []
+ assert result.content == text
+
+ def test_char_by_char_undeclared_name_stays_content(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ text = DSML_INVOKE_PREFIX + "made_up_tool" + DSML_INVOKE_NAME_END + " after."
+ results = simulate_tool_streaming(parser, mock_request, list(text))
+ finish_delta = parser.finish_streaming()
+
+ assert collect_function_name(results) is None
+ content = collect_content(results) + (
+ finish_delta.content if finish_delta and finish_delta.content else ""
+ )
+ assert content == text
+
+ def test_quoted_marker_then_wrapped_call_non_streaming(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Prose that quotes the invoke marker and never closes it must
+ not swallow a real wrapped tool call that follows."""
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ text = (
+ "Docs quote "
+ + DSML_INVOKE_PREFIX
+ + " as the marker. "
+ + _func_calls(_invoke("get_weather", _param("city", "true", "SF")))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"city": "SF"}
+ assert DSML_INVOKE_PREFIX in result.content
+
+ def test_quoted_marker_directly_before_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A quoted marker followed immediately by the real wrapper must
+ release the hold and parse the wrapped call."""
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ text = (
+ "See "
+ + DSML_INVOKE_PREFIX
+ + _func_calls(_invoke("get_weather", _param("city", "true", "SF")))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"city": "SF"}
+ assert DSML_INVOKE_PREFIX in result.content
+ # The wrapper token opens the real tool call, so it must be
+ # consumed by the parser rather than left in the content.
+ assert DSML_FUNC_START not in result.content
+
+ def test_streaming_quoted_marker_then_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ chunks = [
+ "Docs quote ",
+ DSML_INVOKE_PREFIX,
+ " as the marker. ",
+ DSML_FUNC_START,
+ _invoke("get_weather", _param("city", "true", "SF")),
+ DSML_FUNC_END,
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_function_name(results) == "get_weather"
+ args = json.loads(collect_tool_arguments(results))
+ assert args == {"city": "SF"}
+ assert DSML_INVOKE_PREFIX in collect_content(results)
+
+ def test_streaming_quoted_marker_prose_released_before_finish(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Prose after a quoted marker must stream out promptly instead
+ of being buffered until the end of the response."""
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ prose = "this marker starts a tool call block in the raw output."
+ chunks = ["Quote: ", DSML_INVOKE_PREFIX, prose, " More prose."]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_function_name(results) is None
+ content = collect_content(results)
+ assert prose in content
+ assert " More prose." in content
+
+ def test_trailing_prose_after_orphan_invoke_is_kept(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A model that drops the opening wrapper often drops the
+ closing one too, which leaves the response ending between
+ invokes. The text after the invoke is real output and must
+ survive as content."""
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ text = _invoke("get_weather", _param("city", "true", "SF")) + "\nThanks!"
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ assert result.content == "\nThanks!"
+
+ def test_streaming_trailing_prose_after_orphan_invoke_is_kept(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ chunks = [_invoke("get_weather", _param("city", "true", "SF")), "\nThanks!"]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_function_name(results) == "get_weather"
+ # Streamed out as it arrives, not buffered until finish.
+ assert collect_content(results) == "\nThanks!"
+
+ def test_whitespace_between_parallel_orphan_invokes_is_ignored(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A response that is only two invokes and the padding between
+ them comes back with no content at all.
+
+ This case is already covered by the parser dropping content that
+ is nothing but whitespace when the response called tools, so it
+ passes whether or not the engine holds the padding back. The
+ test that actually pins the holding back is
+ ``test_padding_between_orphan_invokes_is_dropped_after_prose``.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ text = (
+ _invoke("get_weather", _param("city", "true", "SF"))
+ + "\n \n"
+ + _invoke("get_time", _param("timezone", "true", "EST"))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 2
+ assert result.content is None
+
+ def test_padding_between_orphan_invokes_is_dropped_after_prose(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Padding between two recovered invokes is dropped even when
+ the response already produced real text.
+
+ The prose in front means the content is no longer whitespace
+ only, so the parser's own whitespace dropping does not apply and
+ the engine holding the padding back is the only thing keeping it
+ out. A wrapped call written the same way returns just the
+ prose, and the recovered call has to match it.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ text = (
+ "Some prose "
+ + _invoke("get_weather", _param("city", "true", "SF"))
+ + "\n \n"
+ + _invoke("get_time", _param("timezone", "true", "EST"))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 2
+ assert result.content == "Some prose "
+
+ def test_streaming_padding_between_orphan_invokes_is_dropped_after_prose(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ chunks = [
+ "Some prose ",
+ _invoke("get_weather", _param("city", "true", "SF")),
+ "\n \n",
+ _invoke("get_time", _param("timezone", "true", "EST")),
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_content(results) == "Some prose "
+
+ def test_recovery_does_not_carry_into_a_later_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Once a recovered sequence ends, a later wrapped call in the
+ same response is treated as an ordinary wrapped call.
+
+ Text between the invokes of a wrapped call is dropped, so if the
+ engine still thought it was inside a recovered sequence the
+ stray text below would come back as content.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ text = (
+ _invoke("get_weather", _param("city", "true", "SF"))
+ + DSML_FUNC_END
+ + DSML_FUNC_START
+ + _invoke("get_weather", _param("city", "true", "SF"))
+ + "stray between wrapped invokes"
+ + _invoke("get_time", _param("timezone", "true", "EST"))
+ + DSML_FUNC_END
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 3
+ assert result.content is None
+
+ def test_streaming_recovery_does_not_carry_into_a_later_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ chunks = [
+ _invoke("get_weather", _param("city", "true", "SF")),
+ DSML_FUNC_END,
+ DSML_FUNC_START,
+ _invoke("get_weather", _param("city", "true", "SF")),
+ "stray between wrapped invokes",
+ _invoke("get_time", _param("timezone", "true", "EST")),
+ DSML_FUNC_END,
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_content(results) == ""
+
+ def test_padding_held_before_one_invoke_does_not_reach_a_later_gap(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Padding held before one invoke is dropped when that invoke
+ starts, so it cannot reappear in front of later text.
+
+ The first gap is padding and belongs to nothing. Only the
+ second gap runs into real text, so only that one is content.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ text = (
+ _invoke("get_weather", _param("city", "true", "SF"))
+ + "\n\n"
+ + _invoke("get_time", _param("timezone", "true", "EST"))
+ + " "
+ + "Real text"
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 2
+ assert result.content == " Real text"
+
+ def test_streaming_padding_held_before_one_invoke_does_not_reach_a_later_gap(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ chunks = [
+ _invoke("get_weather", _param("city", "true", "SF")),
+ "\n\n",
+ _invoke("get_time", _param("timezone", "true", "EST")),
+ " ",
+ "Real text",
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_content(results) == " Real text"
+
+ def test_abandoned_recovery_does_not_affect_a_later_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A recovery attempt that turns out not to name a declared tool
+ must leave nothing behind.
+
+ The invoke below is held while its name is read, then given up
+ on because ``get_nothing`` was never declared. The wrapped call
+ after it is ordinary, so the stray text between its invokes is
+ dropped.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+ abandoned = _invoke("get_nothing", _param("city", "true", "SF"))
+ text = (
+ abandoned
+ + DSML_FUNC_START
+ + _invoke("get_weather", _param("city", "true", "SF"))
+ + "stray between wrapped invokes"
+ + _invoke("get_time", _param("timezone", "true", "EST"))
+ + DSML_FUNC_END
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called
+ assert len(result.tool_calls) == 2
+ assert result.content == abandoned
+
+ def test_recovery_does_not_leak_into_the_next_request(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A response that ends part way through a recovered sequence
+ must not leave the engine set up for recovery.
+
+ The engine is reused, so without a clean start the next
+ response would treat an ordinary wrapped call as a recovered
+ one and hand back the text between its invokes as content.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool, time_tool])
+ mock_request.tools = [weather_tool, time_tool]
+
+ first = parser.extract_tool_calls(
+ _invoke("get_weather", _param("city", "true", "SF")), mock_request
+ )
+ assert first.tools_called
+
+ second = parser.extract_tool_calls(
+ DSML_FUNC_START
+ + _invoke("get_weather", _param("city", "true", "SF"))
+ + "stray between wrapped invokes"
+ + _invoke("get_time", _param("timezone", "true", "EST"))
+ + DSML_FUNC_END,
+ mock_request,
+ )
+
+ assert second.tools_called
+ assert len(second.tool_calls) == 2
+ assert second.content is None
+
+ def test_declared_names_do_not_leak_into_the_next_request(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """The engine is reused across requests, so a request that
+ declares no tools must not recover a tool that an earlier
+ request declared."""
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[weather_tool])
+ mock_request.tools = [weather_tool]
+ text = _invoke("get_weather", _param("city", "true", "SF")) + DSML_FUNC_END
+
+ first = parser.extract_tool_calls(text, mock_request)
+ assert first.tools_called
+
+ second = parser.extract_tool_calls(text, _request_without_tools())
+
+ assert not second.tools_called
+ assert second.tool_calls == []
+ assert second.content == text
+
+
# ββ Initial state ββββββββββββββββββββββββββββββββββββββββββββββββββββ
@@ -254,6 +732,20 @@ def test_streaming_wrapper_unwrap_consistency(self, mock_tokenizer, mock_request
assert '"arguments"' not in streamed_args
assert final_args.startswith(streamed_args)
+ def test_missing_func_start_orphan_invoke(self, mock_tokenizer, mock_request):
+ """Orphan invoke without the <ο½DSMLο½function_calls> wrapper is
+ still parsed as a tool call when the request declared the tool
+ (see gh-48931)."""
+ tool = _make_tool("get_weather", {"city": {"type": "string"}})
+ mock_request.tools = [tool]
+ text = _invoke("get_weather", _param("city", "true", "SF")) + DSML_FUNC_END
+ parser = DeepSeekV32Parser(mock_tokenizer, tools=[tool])
+ results = simulate_tool_streaming(parser, mock_request, list(text))
+ assert collect_function_name(results) == "get_weather"
+ args = json.loads(collect_tool_arguments(results))
+ assert args == {"city": "SF"}
+ assert "DSML" not in collect_content(results)
+
def test_missing_invoke_end(self, mock_tokenizer, mock_request):
text = (
f"{DSML_FUNC_START}\n"
diff --git a/tests/parser/engine/test_deepseek_v4.py b/tests/parser/engine/test_deepseek_v4.py
index e5e7b075ba..5ddb01c5b4 100644
--- a/tests/parser/engine/test_deepseek_v4.py
+++ b/tests/parser/engine/test_deepseek_v4.py
@@ -23,6 +23,7 @@
)
from vllm.parser.abstract_parser import DelegatingParser
from vllm.parser.deepseek_v4 import (
+ DSML_FOREIGN_TOOL_START,
DSML_INVOKE_END,
DSML_INVOKE_NAME_END,
DSML_INVOKE_PREFIX,
@@ -231,6 +232,798 @@ def test_streaming_with_trailing_content(self, mock_tokenizer, mock_request):
assert "Done." in collect_content(results)
+# ββ Missing <ο½DSMLο½tool_calls> before <ο½DSMLο½invoke ...> ββββββββββ
+
+
+class TestMissingToolStart:
+ """Orphan invoke blocks are parsed when the START wrapper is missing.
+
+ At long context DeepSeek V4 models intermittently omit the
+ <ο½DSMLο½tool_calls> wrapper while still emitting a well-formed
+ <ο½DSMLο½invoke ...> block. The (CONTENT, INVOKE_PREFIX) transition
+ recovers the tool call instead of leaking raw DSML into content.
+ Recovery only accepts tool names the request declared, so these
+ tests declare the tools they invoke.
+ See https://github.com/vllm-project/vllm/issues/48931.
+ """
+
+ @pytest.fixture
+ def weather_tool(self):
+ return _make_tool("get_weather", {"location": {"type": "string"}})
+
+ def _declared_parser(self, mock_tokenizer, mock_request, *tools):
+ parser = DeepSeekV4Parser(mock_tokenizer, tools=list(tools))
+ mock_request.tools = list(tools)
+ return parser
+
+ def _orphan_invoke(self, with_tool_end: bool = True) -> str:
+ text = _invoke("get_weather", ("location", "true", "NYC"))
+ if with_tool_end:
+ text += DSML_TOOL_END
+ return text
+
+ def test_non_streaming_orphan_invoke(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ result = parser.extract_tool_calls(self._orphan_invoke(), mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"location": "NYC"}
+ assert result.content is None
+
+ def test_non_streaming_orphan_invoke_no_tool_end(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ result = parser.extract_tool_calls(
+ self._orphan_invoke(with_tool_end=False), mock_request
+ )
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"location": "NYC"}
+
+ def test_non_streaming_orphan_matches_wrapped_parse(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """The orphan payload parses identically to its wrapped form."""
+ invoke = _invoke("get_weather", ("location", "true", "NYC"))
+
+ wrapped_parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool
+ )
+ wrapped = wrapped_parser.extract_tool_calls(
+ DSML_TOOL_START + invoke + DSML_TOOL_END, mock_request
+ )
+ orphan_parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool
+ )
+ orphan = orphan_parser.extract_tool_calls(invoke + DSML_TOOL_END, mock_request)
+
+ assert orphan.tools_called is wrapped.tools_called is True
+ assert orphan.tool_calls[0].function.name == wrapped.tool_calls[0].function.name
+ assert (
+ orphan.tool_calls[0].function.arguments
+ == wrapped.tool_calls[0].function.arguments
+ )
+
+ def test_streaming_orphan_invoke_split_marker(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """The invoke marker may arrive split across streaming deltas."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ chunks = [
+ "I'll check the weather.\n",
+ "<ο½DSML",
+ 'ο½invoke name="get_weather">',
+ "\n" + _param("location", "true", "NYC") + "\n",
+ DSML_INVOKE_END,
+ DSML_TOOL_END,
+ ]
+
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_function_name(results) == "get_weather"
+ args = json.loads(collect_tool_arguments(results))
+ assert args == {"location": "NYC"}
+ content = collect_content(results)
+ assert "I'll check the weather." in content
+ assert "DSML" not in content
+
+ def test_streaming_orphan_invoke_char_by_char(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = self._orphan_invoke()
+ results = simulate_tool_streaming(parser, mock_request, list(text))
+
+ assert collect_function_name(results) == "get_weather"
+ args = json.loads(collect_tool_arguments(results))
+ assert args == {"location": "NYC"}
+ assert "DSML" not in collect_content(results)
+
+ def test_orphan_parallel_invokes(self, mock_tokenizer, mock_request, weather_tool):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ text = (
+ _invoke("get_weather", ("location", "true", "NYC"))
+ + "\n"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ + DSML_TOOL_END
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 2
+ assert result.tool_calls[0].function.name == "get_weather"
+ assert result.tool_calls[1].function.name == "get_time"
+
+ def test_plain_content_unaffected(self, mock_tokenizer, mock_request):
+ parser = DeepSeekV4Parser(mock_tokenizer)
+ text = 'Use style tags to call tools.'
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.tool_calls == []
+ assert result.content == text
+
+ def test_partial_marker_mention_stays_content(self, mock_tokenizer, mock_request):
+ """A DSML-like fragment that never completes the invoke marker
+ must be flushed as content, not swallowed."""
+ parser = DeepSeekV4Parser(mock_tokenizer)
+ text = "The prefix <ο½DSMLο½invoke is reserved."
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.content == text
+
+ def test_foreign_function_calls_wrapper_still_rejected(
+ self, mock_tokenizer, mock_request
+ ):
+ """An invoke inside the V3.2-style function_calls wrapper stays
+ plain content: the orphan fallback must not fire inside a
+ foreign wrapper."""
+ parser = DeepSeekV4Parser(mock_tokenizer)
+ text = _tool_calls(
+ _invoke("get_weather", ("location", "true", "NYC")),
+ ).replace("tool_calls", "function_calls")
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.tool_calls == []
+ assert result.content == text
+
+ def test_unclosed_foreign_wrapper_then_native_call(
+ self, mock_tokenizer, mock_request
+ ):
+ """A foreign wrapper that never closes must not disable native
+ tool parsing: the token backed tool_calls wrapper still wins."""
+ parser = DeepSeekV4Parser(mock_tokenizer)
+ text = (
+ DSML_FOREIGN_TOOL_START
+ + "\nStray foreign text.\n"
+ + _tool_calls(_invoke("get_weather", ("location", "true", "NYC")))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"location": "NYC"}
+ assert "Stray foreign text." in result.content
+
+ def test_orphan_invoke_without_declared_tools_stays_content(
+ self, mock_tokenizer, mock_request
+ ):
+ """A request that declared no tools can never accept a recovered
+ name, so the orphan invoke stays plain content."""
+ parser = DeepSeekV4Parser(mock_tokenizer)
+ text = self._orphan_invoke()
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.tool_calls == []
+ assert result.content == text
+
+
+# ββ Orphan invoke name validation ββββββββββββββββββββββββββββββββββββ
+
+
+class TestOrphanInvokeNameValidation:
+ """Recovered (orphan) invokes must carry a plausible tool name.
+
+ The invoke marker has no dedicated special token in the DeepSeek
+ vocab, so prose that literally quotes the marker would otherwise be
+ misparsed as a tool call. The (CONTENT, INVOKE_PREFIX) recovery
+ transition holds its events until the name completes and only
+ commits to a tool call when the name is one of the tools the
+ request declared. The hold ends early once the text seen so far
+ can no longer grow into a declared tool name, so streaming is not
+ stalled by prose that quotes the marker.
+ The wrapped (TOOL_PREAMBLE, INVOKE_PREFIX) path is not validated.
+ """
+
+ @pytest.fixture
+ def weather_tool(self):
+ return _make_tool("get_weather", {"location": {"type": "string"}})
+
+ def _declared_parser(self, mock_tokenizer, mock_request, *tools):
+ parser = DeepSeekV4Parser(mock_tokenizer, tools=list(tools))
+ mock_request.tools = list(tools)
+ return parser
+
+ def test_declared_name_recovered_matches_wrapped(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ invoke = _invoke("get_weather", ("location", "true", "NYC"))
+
+ wrapped_parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool
+ )
+ wrapped = wrapped_parser.extract_tool_calls(
+ DSML_TOOL_START + invoke + DSML_TOOL_END, mock_request
+ )
+ orphan_parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool
+ )
+ orphan = orphan_parser.extract_tool_calls(invoke + DSML_TOOL_END, mock_request)
+
+ assert orphan.tools_called is wrapped.tools_called is True
+ assert (
+ orphan.tool_calls[0].function.name
+ == wrapped.tool_calls[0].function.name
+ == "get_weather"
+ )
+ assert (
+ orphan.tool_calls[0].function.arguments
+ == wrapped.tool_calls[0].function.arguments
+ )
+
+ def test_undeclared_name_stays_content(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = (
+ "The marker "
+ + DSML_INVOKE_PREFIX
+ + "made_up_tool"
+ + DSML_INVOKE_NAME_END
+ + " is reserved syntax."
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.tool_calls == []
+ assert result.content == text
+
+ def test_undeclared_orphan_then_wrapped_call_still_parses(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = (
+ DSML_INVOKE_PREFIX
+ + "made_up_tool"
+ + DSML_INVOKE_NAME_END
+ + " then a real call: "
+ + _tool_calls(_invoke("get_weather", ("location", "true", "NYC")))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"location": "NYC"}
+ assert DSML_INVOKE_PREFIX + "made_up_tool" in result.content
+
+ def test_no_tools_name_with_space_stays_content(self, mock_tokenizer, mock_request):
+ parser = DeepSeekV4Parser(mock_tokenizer)
+ text = DSML_INVOKE_PREFIX + "not a name" + DSML_INVOKE_NAME_END + " more text."
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.content == text
+
+ def test_no_tools_empty_name_stays_content(self, mock_tokenizer, mock_request):
+ parser = DeepSeekV4Parser(mock_tokenizer)
+ text = DSML_INVOKE_PREFIX + DSML_INVOKE_NAME_END + " more text."
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.content == text
+
+ def test_truncated_name_flushes_content(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = "Say " + DSML_INVOKE_PREFIX + "get_wea"
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.tool_calls == []
+ assert result.content == text
+
+ def test_streaming_ends_mid_name_flushes_content(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ chunks = ["Say ", DSML_INVOKE_PREFIX, "get_wea"]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+ finish_delta = parser.finish_streaming()
+
+ assert collect_function_name(results) is None
+ assert finish_delta is not None
+ assert not finish_delta.tool_calls
+ content = collect_content(results) + (finish_delta.content or "")
+ assert content == "Say " + DSML_INVOKE_PREFIX + "get_wea"
+
+ def test_char_by_char_declared_name_recovers(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = _invoke("get_weather", ("location", "true", "NYC")) + DSML_TOOL_END
+ results = simulate_tool_streaming(parser, mock_request, list(text))
+
+ assert collect_function_name(results) == "get_weather"
+ args = json.loads(collect_tool_arguments(results))
+ assert args == {"location": "NYC"}
+ assert "DSML" not in collect_content(results)
+
+ def test_char_by_char_undeclared_name_stays_content(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = DSML_INVOKE_PREFIX + "made_up_tool" + DSML_INVOKE_NAME_END + " after."
+ results = simulate_tool_streaming(parser, mock_request, list(text))
+ finish_delta = parser.finish_streaming()
+
+ assert collect_function_name(results) is None
+ content = collect_content(results) + (
+ finish_delta.content if finish_delta and finish_delta.content else ""
+ )
+ assert content == text
+
+ def test_wrapped_path_not_validated(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """An undeclared name inside the tool_calls wrapper still parses:
+ validation applies only to the orphan recovery path."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = _tool_calls(_invoke("undeclared_fn", ("location", "true", "NYC")))
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "undeclared_fn"
+
+ def test_wrapped_path_never_holds_events(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """TOOL_CALL_START fires immediately on the wrapped path, before
+ the name completes: no hold window and no tool_index rewind."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ simulate_tool_streaming(
+ parser, mock_request, [DSML_TOOL_START, DSML_INVOKE_PREFIX]
+ )
+ engine = parser._engine
+
+ assert engine._hold_active is False
+ assert engine.tool_index == 0
+
+ def test_parallel_orphan_invokes_with_declared_tools(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ text = (
+ _invoke("get_weather", ("location", "true", "NYC"))
+ + "\n"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ + DSML_TOOL_END
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 2
+ assert result.tool_calls[0].function.name == "get_weather"
+ assert result.tool_calls[1].function.name == "get_time"
+
+ def test_quoted_marker_then_wrapped_call_non_streaming(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Prose that quotes the invoke marker and never closes it must
+ not swallow a real wrapped tool call that follows."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = (
+ "Docs quote "
+ + DSML_INVOKE_PREFIX
+ + " as the marker. "
+ + _tool_calls(_invoke("get_weather", ("location", "true", "NYC")))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"location": "NYC"}
+ assert DSML_INVOKE_PREFIX in result.content
+
+ def test_quoted_marker_directly_before_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A quoted marker followed immediately by the real wrapper must
+ release the hold and parse the wrapped call."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = (
+ "See "
+ + DSML_INVOKE_PREFIX
+ + _tool_calls(_invoke("get_weather", ("location", "true", "NYC")))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"location": "NYC"}
+ assert DSML_INVOKE_PREFIX in result.content
+ # The wrapper token opens the real tool call, so it must be
+ # consumed by the parser rather than left in the content.
+ assert DSML_TOOL_START not in result.content
+
+ def test_streaming_quoted_marker_then_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ chunks = [
+ "Docs quote ",
+ DSML_INVOKE_PREFIX,
+ " as the marker. ",
+ DSML_TOOL_START,
+ _invoke("get_weather", ("location", "true", "NYC")),
+ DSML_TOOL_END,
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_function_name(results) == "get_weather"
+ args = json.loads(collect_tool_arguments(results))
+ assert args == {"location": "NYC"}
+ content = collect_content(results)
+ assert DSML_INVOKE_PREFIX in content
+
+ def test_streaming_quoted_marker_prose_released_before_finish(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Prose after a quoted marker must stream out promptly instead
+ of being buffered until the end of the response."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ prose = "this marker starts a tool call block in the raw output."
+ chunks = ["Quote: ", DSML_INVOKE_PREFIX, prose, " More prose."]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_function_name(results) is None
+ content = collect_content(results)
+ assert prose in content
+ assert " More prose." in content
+
+ def test_tool_choice_none_keeps_quoted_invoke_as_content(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """With tool_choice set to none, invoke text must stay in the
+ content instead of being consumed by tool recovery."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ mock_request.tool_choice = "none"
+ text = (
+ "Docs say you write "
+ + _invoke("get_weather", ("location", "true", "NYC"))
+ + " to call it."
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is False
+ assert result.tool_calls == []
+ assert result.content == text
+
+ def test_non_ascii_declared_name_recovered(self, mock_tokenizer, mock_request):
+ """A declared tool name is recoverable even when it contains
+ characters outside the ASCII range."""
+ tool = _make_tool("ζ₯诒倩ζ°", {"location": {"type": "string"}})
+ parser = self._declared_parser(mock_tokenizer, mock_request, tool)
+ text = _invoke("ζ₯诒倩ζ°", ("location", "true", "NYC")) + DSML_TOOL_END
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "ζ₯诒倩ζ°"
+ args = json.loads(result.tool_calls[0].function.arguments)
+ assert args == {"location": "NYC"}
+
+ def test_trailing_prose_after_orphan_invoke_is_kept(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A model that drops the opening wrapper often drops the
+ closing one too, which leaves the response ending between
+ invokes. The text after the invoke is real output and must
+ survive as content."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = _invoke("get_weather", ("location", "true", "NYC")) + "\nThanks!"
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 1
+ assert result.tool_calls[0].function.name == "get_weather"
+ assert result.content == "\nThanks!"
+
+ def test_streaming_trailing_prose_after_orphan_invoke_is_kept(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ chunks = [_invoke("get_weather", ("location", "true", "NYC")), "\nThanks!"]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_function_name(results) == "get_weather"
+ # Streamed out as it arrives, not buffered until finish.
+ assert collect_content(results) == "\nThanks!"
+
+ def test_whitespace_between_parallel_orphan_invokes_is_ignored(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A response that is only two invokes and the padding between
+ them comes back with no content at all.
+
+ This case is already covered by the parser dropping content that
+ is nothing but whitespace when the response called tools, so it
+ passes whether or not the engine holds the padding back. The
+ test that actually pins the holding back is
+ ``test_padding_between_orphan_invokes_is_dropped_after_prose``.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ text = (
+ _invoke("get_weather", ("location", "true", "NYC"))
+ + "\n \n"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 2
+ assert result.content is None
+
+ def test_padding_between_orphan_invokes_is_dropped_after_prose(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Padding between two recovered invokes is dropped even when
+ the response already produced real text.
+
+ The prose in front means the content is no longer whitespace
+ only, so the parser's own whitespace dropping does not apply and
+ the engine holding the padding back is the only thing keeping it
+ out. A wrapped call written the same way returns just the
+ prose, and the recovered call has to match it.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ text = (
+ "Some prose "
+ + _invoke("get_weather", ("location", "true", "NYC"))
+ + "\n \n"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 2
+ assert result.content == "Some prose "
+
+ def test_streaming_padding_between_orphan_invokes_is_dropped_after_prose(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ chunks = [
+ "Some prose ",
+ _invoke("get_weather", ("location", "true", "NYC")),
+ "\n \n",
+ _invoke("get_time", ("timezone", "true", "EST")),
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_content(results) == "Some prose "
+
+ def test_recovery_does_not_carry_into_a_later_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Once a recovered sequence ends, a later wrapped call in the
+ same response is treated as an ordinary wrapped call.
+
+ Text between the invokes of a wrapped call is dropped, so if the
+ engine still thought it was inside a recovered sequence the
+ stray text below would come back as content.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ text = (
+ _invoke("get_weather", ("location", "true", "NYC"))
+ + DSML_TOOL_END
+ + DSML_TOOL_START
+ + _invoke("get_weather", ("location", "true", "NYC"))
+ + "stray between wrapped invokes"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ + DSML_TOOL_END
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 3
+ assert result.content is None
+
+ def test_streaming_recovery_does_not_carry_into_a_later_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ chunks = [
+ _invoke("get_weather", ("location", "true", "NYC")),
+ DSML_TOOL_END,
+ DSML_TOOL_START,
+ _invoke("get_weather", ("location", "true", "NYC")),
+ "stray between wrapped invokes",
+ _invoke("get_time", ("timezone", "true", "EST")),
+ DSML_TOOL_END,
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_content(results) == ""
+
+ def test_padding_held_before_one_invoke_does_not_reach_a_later_gap(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """Padding held before one invoke is dropped when that invoke
+ starts, so it cannot reappear in front of later text.
+
+ The first gap is padding and belongs to nothing. Only the
+ second gap runs into real text, so only that one is content.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ text = (
+ _invoke("get_weather", ("location", "true", "NYC"))
+ + "\n\n"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ + " "
+ + "Real text"
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 2
+ assert result.content == " Real text"
+
+ def test_streaming_padding_held_before_one_invoke_does_not_reach_a_later_gap(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ chunks = [
+ _invoke("get_weather", ("location", "true", "NYC")),
+ "\n\n",
+ _invoke("get_time", ("timezone", "true", "EST")),
+ " ",
+ "Real text",
+ ]
+ results = simulate_tool_streaming(parser, mock_request, chunks)
+
+ assert collect_content(results) == " Real text"
+
+ def test_abandoned_recovery_does_not_affect_a_later_wrapped_call(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A recovery attempt that turns out not to name a declared tool
+ must leave nothing behind.
+
+ The invoke below is held while its name is read, then given up
+ on because ``get_nothing`` was never declared. The wrapped call
+ after it is ordinary, so the stray text between its invokes is
+ dropped.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+ abandoned = _invoke("get_nothing", ("location", "true", "NYC"))
+ text = (
+ abandoned
+ + DSML_TOOL_START
+ + _invoke("get_weather", ("location", "true", "NYC"))
+ + "stray between wrapped invokes"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ + DSML_TOOL_END
+ )
+ result = parser.extract_tool_calls(text, mock_request)
+
+ assert result.tools_called is True
+ assert len(result.tool_calls) == 2
+ assert result.content == abandoned
+
+ def test_recovery_does_not_leak_into_the_next_request(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """A response that ends part way through a recovered sequence
+ must not leave the engine set up for recovery.
+
+ The engine is reused, so without a clean start the next
+ response would treat an ordinary wrapped call as a recovered
+ one and hand back the text between its invokes as content.
+ """
+ time_tool = _make_tool("get_time", {"timezone": {"type": "string"}})
+ parser = self._declared_parser(
+ mock_tokenizer, mock_request, weather_tool, time_tool
+ )
+
+ first = parser.extract_tool_calls(
+ _invoke("get_weather", ("location", "true", "NYC")), mock_request
+ )
+ assert first.tools_called is True
+
+ second = parser.extract_tool_calls(
+ DSML_TOOL_START
+ + _invoke("get_weather", ("location", "true", "NYC"))
+ + "stray between wrapped invokes"
+ + _invoke("get_time", ("timezone", "true", "EST"))
+ + DSML_TOOL_END,
+ mock_request,
+ )
+
+ assert second.tools_called is True
+ assert len(second.tool_calls) == 2
+ assert second.content is None
+
+ def test_declared_names_do_not_leak_into_the_next_request(
+ self, mock_tokenizer, mock_request, weather_tool
+ ):
+ """The engine is reused across requests, so a request that
+ declares no tools must not recover a tool that an earlier
+ request declared."""
+ parser = self._declared_parser(mock_tokenizer, mock_request, weather_tool)
+ text = _invoke("get_weather", ("location", "true", "NYC")) + DSML_TOOL_END
+
+ first = parser.extract_tool_calls(text, mock_request)
+ assert first.tools_called is True
+
+ second = parser.extract_tool_calls(text, _request_without_tools())
+
+ assert second.tools_called is False
+ assert second.tool_calls == []
+ assert second.content == text
+
+
# ββ Thinking mode initial state ββββββββββββββββββββββββββββββββββββββ
@@ -540,6 +1333,20 @@ def _tool_calls(*invokes):
return DSML_TOOL_START + "\n".join(invokes) + DSML_TOOL_END
+def _request_without_tools():
+ from unittest.mock import MagicMock
+
+ from vllm.entrypoints.openai.chat_completion.protocol import ( # noqa: E501
+ ChatCompletionRequest,
+ )
+
+ req = MagicMock(spec=ChatCompletionRequest)
+ req.tools = []
+ req.tool_choice = "auto"
+ req.include_reasoning = True
+ return req
+
+
class TestParallelUnwrapping:
@pytest.fixture
def weather_tool(self):
diff --git a/tests/tokenizers_/test_deepseek_v4.py b/tests/tokenizers_/test_deepseek_v4.py
index 358732eabf..5a1c676e06 100644
--- a/tests/tokenizers_/test_deepseek_v4.py
+++ b/tests/tokenizers_/test_deepseek_v4.py
@@ -288,3 +288,71 @@ def test_deepseek_v4_matches_reference_golden_fixtures(case_id, kwargs):
expected = (FIXTURES_DIR / f"test_output_{case_id}.txt").read_text()
assert prompt == expected
+
+
+def _render(messages, **kwargs):
+ return _tokenizer().apply_chat_template(
+ conversation=messages, messages=messages, tokenize=False, **kwargs
+ )
+
+
+@pytest.mark.parametrize(
+ ("kwargs", "expected_tail"),
+ [({}, ""), ({"thinking": True}, "")],
+)
+def test_deepseek_v4_trailing_system_gets_generation_prompt(kwargs, expected_tail):
+ """A system message after the last user turn must still open an assistant turn.
+
+ Agent frameworks append context/reminder system messages after the user
+ turn. Without the generation prompt the model sees no assistant boundary
+ and continues the prompt as a document instead of answering.
+ """
+ prompt = _render(
+ [
+ {"role": "system", "content": "you are helpful"},
+ {"role": "user", "content": "write the report"},
+ {"role": "system", "content": "Available agent types: ..."},
+ ],
+ **kwargs,
+ )
+
+ assert prompt.endswith("<ο½Assistantο½>" + expected_tail)
+
+
+def test_deepseek_v4_system_only_conversation_gets_generation_prompt():
+ prompt = _render([{"role": "system", "content": "just a system prompt"}])
+
+ assert prompt.endswith("<ο½Assistantο½> ")
+
+
+def test_deepseek_v4_mid_conversation_system_does_not_open_a_turn():
+ """A system message that is not last must not emit a spurious turn marker."""
+ prompt = _render(
+ [
+ {"role": "system", "content": "you are helpful"},
+ {"role": "system", "content": "extra context"},
+ {"role": "user", "content": "hi"},
+ ]
+ )
+
+ assert prompt.count("<ο½Assistantο½>") == 1
+ assert prompt.endswith("<ο½Assistantο½>")
+
+
+def test_deepseek_v4_system_before_latest_reminder_emits_no_turn_marker():
+ """Regression: a non-final system message must not open an assistant turn.
+
+ `latest_reminder` is exempt from the "what may follow" early return, so a
+ system message preceding one reaches the generation-prompt branch. Treating
+ it as a turn boundary injects a stray marker mid-prompt.
+ """
+ prompt = _render(
+ [
+ {"role": "system", "content": "sys"},
+ {"role": "latest_reminder", "content": "2026-08-04"},
+ {"role": "user", "content": "hi"},
+ ]
+ )
+
+ assert prompt.index("<ο½latest_reminderο½>") < prompt.index("<ο½Assistantο½>")
+ assert prompt.count("<ο½Assistantο½>") == 1
diff --git a/vllm/distributed/device_communicators/cuda_communicator.py b/vllm/distributed/device_communicators/cuda_communicator.py
index 9e02463b5d..9e8086c57d 100644
--- a/vllm/distributed/device_communicators/cuda_communicator.py
+++ b/vllm/distributed/device_communicators/cuda_communicator.py
@@ -76,6 +76,9 @@ def __init__(
from vllm.distributed.device_communicators.flashinfer_all_reduce import (
FlashInferAllReduce,
)
+ from vllm.distributed.device_communicators.hier_all_reduce import (
+ HierarchicalAllReduce,
+ )
from vllm.distributed.device_communicators.pynccl import PyNcclCommunicator
from vllm.distributed.device_communicators.quick_all_reduce import (
QuickAllReduce,
@@ -131,6 +134,24 @@ def __init__(
**extra,
)
+ self.hier_ar_comm: HierarchicalAllReduce | None = None
+ if envs.VLLM_HIER_ALL_REDUCE and self.world_size > 1:
+ islands = [
+ [int(r) for r in part.split(",")]
+ for part in envs.VLLM_HIER_ALL_REDUCE.split(";")
+ ]
+ if sorted(r for i in islands for r in i) == list(range(self.world_size)):
+ self.hier_ar_comm = HierarchicalAllReduce(
+ self.cpu_group, self.device, islands
+ )
+ else:
+ logger.warning(
+ "VLLM_HIER_ALL_REDUCE=%s does not cover ranks 0..%d exactly; "
+ "hierarchical allreduce disabled.",
+ envs.VLLM_HIER_ALL_REDUCE,
+ self.world_size - 1,
+ )
+
if use_custom_allreduce and self.world_size > 1 and current_platform.is_rocm():
# Initialize a custom quick all-reduce implementation for AMD.
# Quick reduce is designed as a complement to custom allreduce
@@ -317,6 +338,9 @@ def all_reduce(self, input_):
out = aiter_ar_comm.custom_all_reduce(input_)
assert out is not None
return out
+ hier_ar_comm = self.hier_ar_comm
+ if hier_ar_comm is not None and hier_ar_comm.should_use(input_):
+ return hier_ar_comm.all_reduce(input_)
ca_comm = self.ca_comm
if (
ca_comm is not None
diff --git a/vllm/distributed/device_communicators/hier_all_reduce.py b/vllm/distributed/device_communicators/hier_all_reduce.py
new file mode 100644
index 0000000000..79ad1cc5db
--- /dev/null
+++ b/vllm/distributed/device_communicators/hier_all_reduce.py
@@ -0,0 +1,595 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Island-aware hierarchical allreduce for multi-island PCIe topologies.
+
+Targets boxes like 2x4 A100/A800 PCIe (two PIX islands bridged by the CPU
+interconnect), where the NCCL 8-GPU ring pays the cross-socket latency on
+every hop. Strategy per rank, all inside one Triton kernel launch:
+
+ 1. publish: copy the local input into this rank's IPC-shared slot,
+ release-fence, bump the phase-A flag.
+ 2. island reduce: spin on the island peers' phase-A flags, then read the
+ island's slots directly over P2P and reduce -> island partial, publish
+ to the partial slot with a phase-B flag.
+ 3. cross exchange: spin on the counterpart rank's phase-B flag (the rank
+ with the same island-local index in the other island), read its partial
+ over P2P, add, write the final result. Cross-socket traffic is exactly
+ one message per rank instead of the ring's repeated crossings.
+
+Latency-bound small messages only (decode hidden states); large payloads
+should stay on NCCL. Buffers are IPC-registered once; flags use monotonically
+increasing sequence tokens so no zeroing is needed between calls.
+"""
+
+import ctypes
+import glob
+import os
+from collections.abc import Sequence
+
+import torch
+import torch.distributed as dist
+
+from vllm.logger import init_logger
+from vllm.triton_utils import tl, triton
+
+logger = init_logger(__name__)
+
+_MAX_ELEMS = 256 * 1024 # 512KB bf16 cap; decode messages are ~8-512KB
+_MAX_CTA = 16
+_NUM_WARPS = 8
+# One-shot moves 4n remote bytes in 2 sync rounds, two-shot 7n/4 in 3 and
+# crosses the slow inter-island link with only n/island_size. Latency wins
+# below this size, bandwidth above it; crossover measured on 2x4 A800 PCIe
+# (32KB: 45.6 vs 47.3us one-shot favoured, 48KB: 54.3 vs 49.6 two-shot).
+_TWO_SHOT_MIN_ELEMS = 24 * 1024
+
+
+class _IpcHandle(ctypes.Structure):
+ _fields_ = [("reserved", ctypes.c_char * 64)]
+
+
+_cudart = None
+
+
+def _find_cudart() -> str:
+ """Locate libcudart across install layouts: the CUDA runtime may come
+ from the pip wheel, from torch's bundled libs, or from a system CUDA
+ install (as in most container images)."""
+ # If torch already loaded it, reuse exactly that library.
+ try:
+ with open("/proc/self/maps") as f:
+ for line in f:
+ path = line.split()[-1] if " /" in line else ""
+ if os.path.basename(path).startswith("libcudart.so"):
+ return path
+ except OSError:
+ pass
+ tdir = os.path.dirname(torch.__file__)
+ for pattern in (
+ os.path.join(
+ os.path.dirname(tdir), "nvidia", "cuda_runtime", "lib", "libcudart.so*"
+ ),
+ os.path.join(tdir, "lib", "libcudart*.so*"),
+ ):
+ found = glob.glob(pattern)
+ if found:
+ return found[0]
+ from ctypes.util import find_library
+
+ name = find_library("cudart")
+ if name:
+ return name
+ raise RuntimeError(
+ "hierarchical allreduce could not locate libcudart; it needs the CUDA "
+ "runtime to map peer buffers via cudaIpcOpenMemHandle"
+ )
+
+
+def _get_cudart():
+ """The shared buffers are cudaMalloc'd and IPC-opened via ctypes: torch
+ opens IPC handles inside the OWNER's device context, so the mapping is
+ never made peer-accessible to the importing device's kernels."""
+ global _cudart
+ if _cudart is None:
+ lib = ctypes.CDLL(_find_cudart())
+ lib.cudaMalloc.argtypes = [
+ ctypes.POINTER(ctypes.c_void_p),
+ ctypes.c_size_t,
+ ]
+ lib.cudaMemset.argtypes = [
+ ctypes.c_void_p,
+ ctypes.c_int,
+ ctypes.c_size_t,
+ ]
+ lib.cudaIpcGetMemHandle.argtypes = [
+ ctypes.POINTER(_IpcHandle),
+ ctypes.c_void_p,
+ ]
+ lib.cudaIpcOpenMemHandle.argtypes = [
+ ctypes.POINTER(ctypes.c_void_p),
+ _IpcHandle,
+ ctypes.c_uint,
+ ]
+ lib.cudaIpcCloseMemHandle.argtypes = [ctypes.c_void_p]
+ lib.cudaFree.argtypes = [ctypes.c_void_p]
+ _cudart = lib
+ return _cudart
+
+
+def _check(rc: int, what: str) -> None:
+ if rc != 0:
+ raise RuntimeError(f"{what} failed: cudaError {rc}")
+
+
+@triton.jit
+def _fence_sys(FENCE_LANES: tl.constexpr):
+ """System-scope acq_rel fence, executed by every thread in the CTA.
+
+ PCIe has no native remote atomics, so signalling is a plain remote store
+ bracketed by this fence. Applying the asm to a tensor rather than a
+ scalar is what makes every thread issue it: a fence executed by one
+ thread does not order another thread's stores, so a scalar (single-lane)
+ fence would leave the other lanes' data unordered against the flag.
+ """
+ tl.inline_asm_elementwise(
+ "fence.acq_rel.sys; mov.u32 $0, $1;",
+ "=r,r",
+ [tl.zeros((FENCE_LANES,), dtype=tl.int32)],
+ dtype=tl.int32,
+ is_pure=False,
+ pack=1,
+ )
+
+
+@triton.jit
+def _hier_all_reduce_kernel(
+ inp_ptr,
+ out_ptr,
+ ptrs_ptr, # [world] int64 device pointers to each rank's data slot
+ partial_ptrs_ptr, # [world] int64 pointers to each rank's partial slot
+ flag_ptrs_ptr, # [world] int64 pointers to each rank's flag array
+ rank: tl.constexpr,
+ island_base: tl.constexpr, # first rank of this island
+ island_size: tl.constexpr,
+ counterpart: tl.constexpr, # same-index rank in the other island
+ WORLD: tl.constexpr,
+ numel,
+ token_ptr, # [MAX_CTA] device int32 sequence counters (cudagraph-safe)
+ MAX_ELEMS: tl.constexpr,
+ FENCE_LANES: tl.constexpr,
+ BLOCK: tl.constexpr,
+):
+ # Each CTA owns a disjoint chunk with its own flag row and sequence
+ # counter; chunks synchronize independently, so multiple CTAs keep
+ # multiple PCIe read streams in flight (single-CTA queue depth is the
+ # bandwidth bottleneck at >=32KB).
+ pid = tl.program_id(0)
+ ncta = tl.num_programs(0)
+ token = tl.atomic_add(token_ptr + pid, 1, sem="relaxed") + 1
+ chunk = tl.cdiv(numel, ncta)
+ start = pid * chunk
+ end = tl.minimum(numel, start + chunk)
+ fbase = pid * 2 * WORLD
+
+ # Double-buffer data/partial slots by token parity: a rank can lap a
+ # slow peer by one call, but launch-granularity stream ordering plus
+ # the flag waits make a two-call lap (same-parity reuse) impossible
+ # before the peer's read.
+ buf_off = (token % 2) * MAX_ELEMS
+ my_slot = tl.cast(tl.load(ptrs_ptr + rank), tl.pointer_type(tl.bfloat16)) + buf_off
+ my_partial = (
+ tl.cast(tl.load(partial_ptrs_ptr + rank), tl.pointer_type(tl.bfloat16))
+ + buf_off
+ )
+ my_flags = tl.cast(tl.load(flag_ptrs_ptr + rank), tl.pointer_type(tl.int32))
+
+ # Phase A: publish this chunk locally, then signal island peers.
+ for off in range(start, end, BLOCK):
+ offs = off + tl.arange(0, BLOCK)
+ mask = offs < end
+ x = tl.load(inp_ptr + offs, mask=mask, other=0.0)
+ tl.store(my_slot + offs, x, mask=mask)
+ _fence_sys(FENCE_LANES)
+ tl.debug_barrier()
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ peer_flags = tl.cast(
+ tl.load(flag_ptrs_ptr + peer), tl.pointer_type(tl.int32)
+ )
+ tl.store(peer_flags + fbase + 0 * WORLD + rank, token)
+
+ # Wait for the island peers' phase-A signals for this chunk.
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ while (
+ tl.atomic_add(
+ my_flags + fbase + 0 * WORLD + peer,
+ 0,
+ sem="acquire",
+ scope="sys",
+ )
+ < token
+ ):
+ pass
+ tl.debug_barrier()
+ _fence_sys(FENCE_LANES)
+
+ # Phase B: island reduce over P2P reads, publish partial, signal the
+ # cross-island counterpart. Partials are bf16 to halve cross-island
+ # bytes; accumulation stays fp32.
+ for off in range(start, end, BLOCK):
+ offs = off + tl.arange(0, BLOCK)
+ mask = offs < end
+ acc = tl.load(my_slot + offs, mask=mask, other=0.0).to(tl.float32)
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ peer_slot = (
+ tl.cast(
+ tl.load(ptrs_ptr + peer),
+ tl.pointer_type(tl.bfloat16),
+ )
+ + buf_off
+ )
+ acc += tl.load(peer_slot + offs, mask=mask, other=0.0).to(tl.float32)
+ tl.store(my_partial + offs, acc.to(tl.bfloat16), mask=mask)
+ _fence_sys(FENCE_LANES)
+ tl.debug_barrier()
+ cp_flags = tl.cast(tl.load(flag_ptrs_ptr + counterpart), tl.pointer_type(tl.int32))
+ tl.store(cp_flags + fbase + 1 * WORLD + rank, token)
+
+ # Phase C: wait for the counterpart's phase-B signal for this chunk,
+ # then do the single cross-island exchange.
+ while (
+ tl.atomic_add(
+ my_flags + fbase + 1 * WORLD + counterpart,
+ 0,
+ sem="acquire",
+ scope="sys",
+ )
+ < token
+ ):
+ pass
+ tl.debug_barrier()
+ _fence_sys(FENCE_LANES)
+ cp_partial = (
+ tl.cast(
+ tl.load(partial_ptrs_ptr + counterpart),
+ tl.pointer_type(tl.bfloat16),
+ )
+ + buf_off
+ )
+ for off in range(start, end, BLOCK):
+ offs = off + tl.arange(0, BLOCK)
+ mask = offs < end
+ acc = tl.load(my_partial + offs, mask=mask, other=0.0).to(tl.float32)
+ acc += tl.load(cp_partial + offs, mask=mask, other=0.0).to(tl.float32)
+ tl.store(out_ptr + offs, acc.to(tl.bfloat16), mask=mask)
+
+
+@triton.jit
+def _hier_two_shot_kernel(
+ inp_ptr,
+ out_ptr,
+ ptrs_ptr, # [world] int64 pointers to each rank's data slot
+ partial_ptrs_ptr, # [world] pointers to each rank's island-partial slot
+ gather_ptrs_ptr, # [world] pointers to each rank's global-shard slot
+ flag_ptrs_ptr, # [world] pointers to each rank's flag array
+ rank: tl.constexpr,
+ island_base: tl.constexpr,
+ island_idx: tl.constexpr, # this rank's index inside its island
+ island_size: tl.constexpr,
+ counterpart: tl.constexpr,
+ WORLD: tl.constexpr,
+ numel,
+ token_ptr, # [MAX_CTA] device int32 sequence counters
+ MAX_ELEMS: tl.constexpr,
+ MAX_CTA: tl.constexpr,
+ FENCE_LANES: tl.constexpr,
+ BLOCK: tl.constexpr,
+):
+ """Two-shot variant: reduce-scatter, one cross-island shard exchange,
+ then allgather. Moves 7n/4 remote bytes against the one-shot kernel's
+ 4n, and only n/island_size of that crosses the slow inter-island link
+ (vs n) -- the win grows with message size.
+
+ Partitioning is a stripe: CTA k owns slice k of *every* shard, so a
+ reader CTA always waits on the same CTA index of its peers and no
+ cross-CTA barrier is needed.
+ """
+ pid = tl.program_id(0)
+ ncta = tl.num_programs(0)
+ token = tl.atomic_add(token_ptr + pid, 1, sem="relaxed") + 1
+ buf_off = (token % 2) * MAX_ELEMS
+
+ shard = tl.cdiv(numel, island_size)
+ ss = tl.cdiv(shard, ncta)
+ my_start = island_idx * shard + pid * ss
+ my_end = tl.minimum(tl.minimum(island_idx * shard + shard, my_start + ss), numel)
+
+ my_data = tl.cast(tl.load(ptrs_ptr + rank), tl.pointer_type(tl.bfloat16)) + buf_off
+ my_partial = (
+ tl.cast(tl.load(partial_ptrs_ptr + rank), tl.pointer_type(tl.bfloat16))
+ + buf_off
+ )
+ my_gather = (
+ tl.cast(tl.load(gather_ptrs_ptr + rank), tl.pointer_type(tl.bfloat16)) + buf_off
+ )
+ my_flags = tl.cast(tl.load(flag_ptrs_ptr + rank), tl.pointer_type(tl.int32))
+
+ # Phase A: publish this CTA's stripe of every shard -- exactly the bytes
+ # each island peer reads for the shard it owns.
+ for s in tl.static_range(island_size):
+ start = s * shard + pid * ss
+ end = tl.minimum(tl.minimum(s * shard + shard, start + ss), numel)
+ for off in range(start, end, BLOCK):
+ offs = off + tl.arange(0, BLOCK)
+ mask = offs < end
+ x = tl.load(inp_ptr + offs, mask=mask, other=0.0)
+ tl.store(my_data + offs, x, mask=mask)
+ _fence_sys(FENCE_LANES)
+ tl.debug_barrier()
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ pf = tl.cast(tl.load(flag_ptrs_ptr + peer), tl.pointer_type(tl.int32))
+ tl.store(pf + (0 * WORLD + rank) * MAX_CTA + pid, token)
+
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ while (
+ tl.atomic_add(
+ my_flags + (0 * WORLD + peer) * MAX_CTA + pid,
+ 0,
+ sem="acquire",
+ scope="sys",
+ )
+ < token
+ ):
+ pass
+ tl.debug_barrier()
+ _fence_sys(FENCE_LANES)
+
+ # Phase B: reduce-scatter -- sum the island over the shard this rank owns.
+ for off in range(my_start, my_end, BLOCK):
+ offs = off + tl.arange(0, BLOCK)
+ mask = offs < my_end
+ acc = tl.load(my_data + offs, mask=mask, other=0.0).to(tl.float32)
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ pd = (
+ tl.cast(tl.load(ptrs_ptr + peer), tl.pointer_type(tl.bfloat16))
+ + buf_off
+ )
+ acc += tl.load(pd + offs, mask=mask, other=0.0).to(tl.float32)
+ tl.store(my_partial + offs, acc.to(tl.bfloat16), mask=mask)
+ _fence_sys(FENCE_LANES)
+ tl.debug_barrier()
+ cpf = tl.cast(tl.load(flag_ptrs_ptr + counterpart), tl.pointer_type(tl.int32))
+ tl.store(cpf + (1 * WORLD + rank) * MAX_CTA + pid, token)
+
+ # Phase C: single cross-island exchange, shard-sized. The counterpart
+ # holds the same island-local index, hence the same shard offsets.
+ while (
+ tl.atomic_add(
+ my_flags + (1 * WORLD + counterpart) * MAX_CTA + pid,
+ 0,
+ sem="acquire",
+ scope="sys",
+ )
+ < token
+ ):
+ pass
+ tl.debug_barrier()
+ _fence_sys(FENCE_LANES)
+ cp_partial = (
+ tl.cast(
+ tl.load(partial_ptrs_ptr + counterpart),
+ tl.pointer_type(tl.bfloat16),
+ )
+ + buf_off
+ )
+ for off in range(my_start, my_end, BLOCK):
+ offs = off + tl.arange(0, BLOCK)
+ mask = offs < my_end
+ acc = tl.load(my_partial + offs, mask=mask, other=0.0).to(tl.float32)
+ acc += tl.load(cp_partial + offs, mask=mask, other=0.0).to(tl.float32)
+ tl.store(my_gather + offs, acc.to(tl.bfloat16), mask=mask)
+ _fence_sys(FENCE_LANES)
+ tl.debug_barrier()
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ pf = tl.cast(tl.load(flag_ptrs_ptr + peer), tl.pointer_type(tl.int32))
+ tl.store(pf + (2 * WORLD + rank) * MAX_CTA + pid, token)
+
+ # Phase D: allgather the finished shards from the island.
+ for i in tl.static_range(island_size):
+ peer = island_base + i
+ if peer != rank:
+ while (
+ tl.atomic_add(
+ my_flags + (2 * WORLD + peer) * MAX_CTA + pid,
+ 0,
+ sem="acquire",
+ scope="sys",
+ )
+ < token
+ ):
+ pass
+ tl.debug_barrier()
+ _fence_sys(FENCE_LANES)
+ for s in tl.static_range(island_size):
+ peer = island_base + s
+ pg = (
+ tl.cast(tl.load(gather_ptrs_ptr + peer), tl.pointer_type(tl.bfloat16))
+ + buf_off
+ )
+ start = s * shard + pid * ss
+ end = tl.minimum(tl.minimum(s * shard + shard, start + ss), numel)
+ for off in range(start, end, BLOCK):
+ offs = off + tl.arange(0, BLOCK)
+ mask = offs < end
+ tl.store(
+ out_ptr + offs,
+ tl.load(pg + offs, mask=mask, other=0.0),
+ mask=mask,
+ )
+
+
+class HierarchicalAllReduce:
+ """Two-level island-aware allreduce over IPC-shared buffers.
+
+ Args:
+ group: torch.distributed group covering all ranks on this node.
+ device: this rank's CUDA device.
+ islands: rank partition, e.g. [[0,1,2,3],[4,5,6,7]]. Exactly two
+ islands of equal size are supported.
+ """
+
+ def __init__(
+ self,
+ group: dist.ProcessGroup,
+ device: torch.device,
+ islands: Sequence[Sequence[int]],
+ ) -> None:
+ self.group = group
+ self.device = device
+ self.rank = dist.get_rank(group)
+ self.world_size = dist.get_world_size(group)
+ assert len(islands) == 2 and len(islands[0]) == len(islands[1]), (
+ "HierarchicalAllReduce supports exactly two equal islands"
+ )
+ self.islands = [list(i) for i in islands]
+ me = self.rank
+ self.island_idx = 0 if me in self.islands[0] else 1
+ island = self.islands[self.island_idx]
+ other = self.islands[1 - self.island_idx]
+ self.island_base = min(island)
+ self.island_size = len(island)
+ self.counterpart = other[island.index(me)]
+
+ rt = _get_cudart()
+ _check(rt.cudaSetDevice(ctypes.c_int(device.index)), "cudaSetDevice")
+ self._own = []
+ self._opened = []
+ # data slots (bf16), partial slots (fp32) β double-buffered; flags
+ # [phase(2) x writer(world)] int32
+ data_ptr = self._alloc(rt, 2 * _MAX_ELEMS * 2)
+ partial_ptr = self._alloc(rt, 2 * _MAX_ELEMS * 2)
+ gather_ptr = self._alloc(rt, 2 * _MAX_ELEMS * 2)
+ flags_ptr = self._alloc(rt, _MAX_CTA * 2 * self.world_size * 4)
+ flags2_ptr = self._alloc(rt, 3 * self.world_size * _MAX_CTA * 4)
+ self._token_ctr = torch.zeros(_MAX_CTA, dtype=torch.int32, device=device)
+ self._token_ctr2 = torch.zeros(_MAX_CTA, dtype=torch.int32, device=device)
+
+ self._data_ptrs = self._exchange_ptrs(rt, data_ptr)
+ self._partial_ptrs = self._exchange_ptrs(rt, partial_ptr)
+ self._gather_ptrs = self._exchange_ptrs(rt, gather_ptr)
+ self._flag_ptrs = self._exchange_ptrs(rt, flags_ptr)
+ self._flag2_ptrs = self._exchange_ptrs(rt, flags2_ptr)
+
+ def _alloc(self, rt, nbytes: int) -> int:
+ ptr = ctypes.c_void_p()
+ _check(rt.cudaMalloc(ctypes.byref(ptr), nbytes), "cudaMalloc")
+ _check(rt.cudaMemset(ptr, 0, nbytes), "cudaMemset")
+ self._own.append(ptr.value)
+ return ptr.value
+
+ def _exchange_ptrs(self, rt, local_ptr: int) -> torch.Tensor:
+ """Share a raw device allocation with all ranks via CUDA IPC and
+ return a device tensor of every rank's pointer. Handles are opened
+ with the importing device current so lazy peer access covers our
+ kernels' direct loads/stores."""
+ handle = _IpcHandle()
+ _check(
+ rt.cudaIpcGetMemHandle(ctypes.byref(handle), ctypes.c_void_p(local_ptr)),
+ "cudaIpcGetMemHandle",
+ )
+ objs: list = [None] * self.world_size
+ dist.all_gather_object(objs, (self.rank, bytes(handle)), group=self.group)
+ ptrs = torch.zeros(self.world_size, dtype=torch.int64, device=self.device)
+ for rank, hbytes in objs:
+ if rank == self.rank:
+ ptrs[rank] = local_ptr
+ continue
+ h = _IpcHandle.from_buffer_copy(hbytes)
+ peer_ptr = ctypes.c_void_p()
+ _check(
+ rt.cudaIpcOpenMemHandle(ctypes.byref(peer_ptr), h, ctypes.c_uint(1)),
+ "cudaIpcOpenMemHandle",
+ )
+ self._opened.append(peer_ptr.value)
+ ptrs[rank] = peer_ptr.value
+ return ptrs
+
+ def __del__(self):
+ try:
+ rt = _get_cudart()
+ for ptr in getattr(self, "_opened", []):
+ rt.cudaIpcCloseMemHandle(ctypes.c_void_p(ptr))
+ for ptr in getattr(self, "_own", []):
+ rt.cudaFree(ctypes.c_void_p(ptr))
+ except Exception:
+ pass
+
+ def should_use(self, inp: torch.Tensor) -> bool:
+ return (
+ inp.dtype == torch.bfloat16
+ and inp.is_contiguous()
+ and inp.numel() <= _MAX_ELEMS
+ and inp.numel() % self.island_size == 0
+ )
+
+ def all_reduce(self, inp: torch.Tensor, out: torch.Tensor | None = None):
+ if out is None:
+ out = torch.empty_like(inp)
+ numel = inp.numel()
+ if numel >= _TWO_SHOT_MIN_ELEMS:
+ ncta = min(_MAX_CTA, max(1, numel // (self.island_size * 1024)))
+ _hier_two_shot_kernel[(ncta,)](
+ inp.view(-1),
+ out.view(-1),
+ self._data_ptrs,
+ self._partial_ptrs,
+ self._gather_ptrs,
+ self._flag2_ptrs,
+ rank=self.rank,
+ island_base=self.island_base,
+ island_idx=self.rank - self.island_base,
+ island_size=self.island_size,
+ counterpart=self.counterpart,
+ WORLD=self.world_size,
+ numel=numel,
+ token_ptr=self._token_ctr2,
+ MAX_ELEMS=_MAX_ELEMS,
+ MAX_CTA=_MAX_CTA,
+ FENCE_LANES=_NUM_WARPS * 32,
+ BLOCK=min(4096, triton.next_power_of_2(numel)),
+ num_warps=_NUM_WARPS,
+ )
+ return out
+ ncta = min(_MAX_CTA, max(1, (numel + 4095) // 4096))
+ _hier_all_reduce_kernel[(ncta,)](
+ inp.view(-1),
+ out.view(-1),
+ self._data_ptrs,
+ self._partial_ptrs,
+ self._flag_ptrs,
+ rank=self.rank,
+ island_base=self.island_base,
+ island_size=self.island_size,
+ counterpart=self.counterpart,
+ WORLD=self.world_size,
+ numel=numel,
+ token_ptr=self._token_ctr,
+ MAX_ELEMS=_MAX_ELEMS,
+ FENCE_LANES=_NUM_WARPS * 32,
+ BLOCK=min(8192, triton.next_power_of_2(numel)),
+ num_warps=_NUM_WARPS,
+ )
+ return out
diff --git a/vllm/envs.py b/vllm/envs.py
index b32d7fae95..7e8183f6bb 100755
--- a/vllm/envs.py
+++ b/vllm/envs.py
@@ -294,6 +294,7 @@
VLLM_ELASTIC_EP_SCALE_UP_LAUNCH: bool = False
VLLM_ELASTIC_EP_DRAIN_REQUESTS: bool = False
VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS: bool = True
+ VLLM_HIER_ALL_REDUCE: str = ""
VLLM_NIXL_EP_MAX_NUM_RANKS: int = 32
VLLM_XPU_ENABLE_XPU_GRAPH: bool = False
VLLM_XPU_USE_SAMPLER_KERNEL: bool = True
@@ -2030,6 +2031,9 @@ def _resolve_rust_cli_path() -> str | None:
"VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS": lambda: bool(
int(os.getenv("VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS", "1"))
),
+ # Island partition for the hierarchical allreduce, e.g. "0,1,2,3;4,5,6,7"
+ # for a 2x4 PCIe box. Empty disables it.
+ "VLLM_HIER_ALL_REDUCE": lambda: os.getenv("VLLM_HIER_ALL_REDUCE", ""),
# NIXL EP environment variables
"VLLM_NIXL_EP_MAX_NUM_RANKS": lambda: int(
os.getenv("VLLM_NIXL_EP_MAX_NUM_RANKS", "32")
diff --git a/vllm/parser/deepseek_v32.py b/vllm/parser/deepseek_v32.py
index 0d9ac9f53c..30768d660f 100644
--- a/vllm/parser/deepseek_v32.py
+++ b/vllm/parser/deepseek_v32.py
@@ -26,6 +26,8 @@
DSML_INVOKE_NAME_END,
DSML_INVOKE_PREFIX,
DSML_PARAM_CLOSE,
+ DSML_TOOL_END,
+ DSML_TOOL_START,
_dsml_arg_converter,
_unwrap_wrapper_args,
)
@@ -59,6 +61,8 @@ def deepseek_v32_config() -> ParserEngineConfig:
"INVOKE_NAME_END": DSML_INVOKE_NAME_END,
"INVOKE_END": DSML_INVOKE_END,
"PARAM_CLOSE": DSML_PARAM_CLOSE,
+ "FOREIGN_START": DSML_TOOL_START,
+ "FOREIGN_END": DSML_TOOL_END,
},
token_id_terminals={
"TOOL_START": DSML_FUNC_START,
@@ -69,6 +73,34 @@ def deepseek_v32_config() -> ParserEngineConfig:
ParserState.TOOL_PREAMBLE,
(),
),
+ # Orphan invoke: at long context the model may omit the
+ # <ο½DSMLο½function_calls> wrapper and emit the invoke
+ # directly. The invoke marker has no dedicated special
+ # token, so hold events and validate the parsed name
+ # before committing. Only names the request declared are
+ # accepted.
+ (ParserState.CONTENT, "INVOKE_PREFIX"): Transition(
+ ParserState.TOOL_NAME,
+ (EventType.TOOL_CALL_START,),
+ validate_tool_name=True,
+ ),
+ # V4-style tool_calls wrapper is foreign to V3.2: pass it
+ # and its contents through as plain content
+ (ParserState.CONTENT, "FOREIGN_START"): Transition(
+ ParserState.FOREIGN_BLOCK,
+ (EventType.TEXT_CHUNK,),
+ ),
+ (ParserState.FOREIGN_BLOCK, "FOREIGN_END"): Transition(
+ ParserState.CONTENT,
+ (EventType.TEXT_CHUNK,),
+ ),
+ # The native wrapper always wins over an unclosed foreign
+ # block, so a stray foreign start cannot disable tool
+ # parsing for the rest of the response.
+ (ParserState.FOREIGN_BLOCK, "TOOL_START"): Transition(
+ ParserState.TOOL_PREAMBLE,
+ (),
+ ),
(ParserState.TOOL_PREAMBLE, "INVOKE_PREFIX"): Transition(
ParserState.TOOL_NAME,
(EventType.TOOL_CALL_START,),
@@ -99,6 +131,7 @@ def deepseek_v32_config() -> ParserEngineConfig:
ParserState.CONTENT: EventType.TEXT_CHUNK,
ParserState.TOOL_NAME: EventType.TOOL_NAME,
ParserState.TOOL_ARGS: EventType.ARG_VALUE_CHUNK,
+ ParserState.FOREIGN_BLOCK: EventType.TEXT_CHUNK,
},
arg_converter=_dsml_arg_converter,
arg_structural_chars=frozenset(">"),
diff --git a/vllm/parser/deepseek_v4.py b/vllm/parser/deepseek_v4.py
index c656687082..94c5b83c4f 100644
--- a/vllm/parser/deepseek_v4.py
+++ b/vllm/parser/deepseek_v4.py
@@ -48,6 +48,9 @@
DSML_INVOKE_NAME_END = '">'
DSML_INVOKE_END = f""
DSML_PARAM_CLOSE = f""
+# DeepSeek V3.2-style wrapper, recognized only to reject it as foreign
+DSML_FOREIGN_TOOL_START = f"<{_DSML}function_calls>"
+DSML_FOREIGN_TOOL_END = f""
_ESCAPED_DSML = re.escape(_DSML)
_PARAM_RE = re.compile(
@@ -135,6 +138,8 @@ def deepseek_v4_config(thinking: bool = False) -> ParserEngineConfig:
"INVOKE_NAME_END": DSML_INVOKE_NAME_END,
"INVOKE_END": DSML_INVOKE_END,
"PARAM_CLOSE": DSML_PARAM_CLOSE,
+ "FOREIGN_START": DSML_FOREIGN_TOOL_START,
+ "FOREIGN_END": DSML_FOREIGN_TOOL_END,
},
token_id_terminals={
"THINK_START": DSML_THINK_START,
@@ -170,6 +175,33 @@ def deepseek_v4_config(thinking: bool = False) -> ParserEngineConfig:
ParserState.TOOL_PREAMBLE,
(),
),
+ # Orphan invoke: at long context the model may omit the
+ # <ο½DSMLο½tool_calls> wrapper and emit the invoke directly.
+ # The invoke marker has no dedicated special token, so hold
+ # events and validate the parsed name before committing.
+ # Only names the request declared are accepted.
+ (ParserState.CONTENT, "INVOKE_PREFIX"): Transition(
+ ParserState.TOOL_NAME,
+ (EventType.TOOL_CALL_START,),
+ validate_tool_name=True,
+ ),
+ # V3.2-style function_calls wrapper is foreign to V4: pass
+ # it and its contents through as plain content
+ (ParserState.CONTENT, "FOREIGN_START"): Transition(
+ ParserState.FOREIGN_BLOCK,
+ (EventType.TEXT_CHUNK,),
+ ),
+ (ParserState.FOREIGN_BLOCK, "FOREIGN_END"): Transition(
+ ParserState.CONTENT,
+ (EventType.TEXT_CHUNK,),
+ ),
+ # The native wrapper always wins over an unclosed foreign
+ # block, so a stray foreign start cannot disable tool
+ # parsing for the rest of the response.
+ (ParserState.FOREIGN_BLOCK, "TOOL_START"): Transition(
+ ParserState.TOOL_PREAMBLE,
+ (),
+ ),
(ParserState.TOOL_PREAMBLE, "INVOKE_PREFIX"): Transition(
ParserState.TOOL_NAME,
(EventType.TOOL_CALL_START,),
@@ -201,6 +233,7 @@ def deepseek_v4_config(thinking: bool = False) -> ParserEngineConfig:
ParserState.REASONING: EventType.REASONING_CHUNK,
ParserState.TOOL_NAME: EventType.TOOL_NAME,
ParserState.TOOL_ARGS: EventType.ARG_VALUE_CHUNK,
+ ParserState.FOREIGN_BLOCK: EventType.TEXT_CHUNK,
},
arg_converter=_dsml_arg_converter,
arg_structural_chars=frozenset(">"),
diff --git a/vllm/parser/engine/parser_engine.py b/vllm/parser/engine/parser_engine.py
index 048e714cb4..b58405695f 100644
--- a/vllm/parser/engine/parser_engine.py
+++ b/vllm/parser/engine/parser_engine.py
@@ -29,6 +29,7 @@
from vllm.parser.engine.streaming_parser_engine import StreamingParserEngine
from vllm.tool_parsers.utils import (
coerce_to_schema_type,
+ collect_tool_names,
extract_types_from_schema,
find_tool_name,
find_tool_properties,
@@ -107,6 +108,7 @@ def __init__(
self._engine = StreamingParserEngine(
parser_engine_config, tokenizer, vocab=self.vocab
)
+ self._engine.allowed_tool_names = self._declared_tool_names()
self._has_reasoning = (
"THINK_END" in parser_engine_config.token_id_terminals
@@ -401,6 +403,11 @@ def _accept_tool_name(self, name: str) -> bool:
# ββ Private helpers βββββββββββββββββββββββββββββββββββββββββββββ
+ def _declared_tool_names(self) -> frozenset[str] | None:
+ if not self._tools:
+ return None
+ return collect_tool_names(self._tools) or None
+
def _check_skip_tool_parsing(
self,
request: ChatCompletionRequest | ResponsesRequest,
@@ -408,10 +415,21 @@ def _check_skip_tool_parsing(
tools = getattr(request, "tools", None)
if tools:
self._tools = tools
+ self._engine.allowed_tool_names = self._declared_tool_names()
+ else:
+ # The engine is reused across requests and reset() keeps this
+ # field, so it has to be cleared here. Otherwise a request
+ # that declares no tools would inherit the names of the
+ # previous one and could recover a tool it never asked for.
+ self._engine.allowed_tool_names = None
if not self.skip_tool_parsing and not self._suppress_tool_calls:
tool_choice = getattr(request, "tool_choice", None)
if tool_choice == "none" and tools:
self._suppress_tool_calls = True
+ # The engine needs the suppression state too: recovery
+ # transitions must not consume text that will never be allowed
+ # to become a tool call.
+ self._engine.suppress_tool_calls = self._suppress_tool_calls
def _strip_content_whitespace(
self,
diff --git a/vllm/parser/engine/parser_engine_config.py b/vllm/parser/engine/parser_engine_config.py
index ad83e33149..196256706e 100644
--- a/vllm/parser/engine/parser_engine_config.py
+++ b/vllm/parser/engine/parser_engine_config.py
@@ -33,6 +33,9 @@ class ParserState(Enum):
TOOL_NAME = auto()
TOOL_ARGS = auto()
TOOL_BETWEEN = auto()
+ # Inside a block belonging to a different model format; terminals
+ # matched here pass through as plain content.
+ FOREIGN_BLOCK = auto()
@dataclass(frozen=True, slots=True)
@@ -40,6 +43,11 @@ class Transition:
next_state: ParserState
events: tuple[EventType, ...] = field(default_factory=tuple)
skip_in_token_id_mode: bool = False
+ # Hold this transition's events until the tool name completes, then
+ # validate the name before committing to the tool call. Set on
+ # recovery transitions whose trigger marker has no dedicated special
+ # token, so prose quoting the marker is not misparsed as a tool call.
+ validate_tool_name: bool = False
@dataclass(frozen=True)
diff --git a/vllm/parser/engine/streaming_parser_engine.py b/vllm/parser/engine/streaming_parser_engine.py
index f5c395f76d..dbccb73cda 100644
--- a/vllm/parser/engine/streaming_parser_engine.py
+++ b/vllm/parser/engine/streaming_parser_engine.py
@@ -158,6 +158,17 @@ def __init__(
)
self.skip_tool_parsing = False
+ # Function names declared by the request, or None when unknown.
+ # Consulted only by transitions with ``validate_tool_name``;
+ # set per request by the owning ParserEngine, like
+ # ``skip_tool_parsing`` it survives reset().
+ self.allowed_tool_names: frozenset[str] | None = None
+ # True when the request asked for tool_choice "none". Recovery
+ # transitions are skipped while set, so text that looks like a
+ # recovered tool call stays plain content instead of being
+ # consumed and then suppressed. Set per request by the owning
+ # ParserEngine; survives reset() like ``skip_tool_parsing``.
+ self.suppress_tool_calls = False
self.reset(initial_state=initial_state)
def _reset_args_state(self) -> None:
@@ -187,6 +198,14 @@ def reset(self, initial_state: ParserState | None = None) -> None:
self._lexer.reset()
self._message_header_buffer = ""
self._reset_args_state()
+ self._recovered_tool_call = False
+ self._pending_between_text = ""
+ self._hold_active = False
+ self._held_events: list[SemanticEvent] = []
+ self._held_raw: list[str] = []
+ self._held_name: list[str] = []
+ self._held_prior_state: ParserState = self.state
+ self._held_prior_tool_index: int = -1
def feed(
self,
@@ -240,6 +259,12 @@ def finish(self) -> list[SemanticEvent]:
events.extend(self._process_lex_tokens(self._lexer.flush()))
+ if self._hold_active:
+ # Stream ended before the recovered tool name completed:
+ # the held events never validated, so flush the raw text
+ # as content in the pre-recovery state.
+ events.extend(self._abort_hold("".join(self._held_raw)))
+
if self._args_buffer:
events.append(
SemanticEvent(
@@ -316,6 +341,15 @@ def _on_terminal(self, terminal: str, value: str) -> list[SemanticEvent]:
if transition is None:
if self._has_drops and terminal == DROP_TERMINAL:
return []
+ if self._hold_active and self.state == ParserState.TOOL_NAME:
+ # A terminal with no meaning inside a held tool name,
+ # for example a real tool call start token, ends the
+ # hold: replay the held text as content, then handle
+ # the terminal again in the restored state so it keeps
+ # its normal meaning.
+ events = self._abort_hold("".join(self._held_raw))
+ events.extend(self._on_terminal(terminal, value))
+ return events
return self._emit_for_state(value)
if self.skip_tool_parsing and terminal in self._tool_terminals:
@@ -356,6 +390,23 @@ def _on_terminal(self, terminal: str, value: str) -> list[SemanticEvent]:
return self._apply_transition(transition, value)
def _emit_for_state(self, text: str) -> list[SemanticEvent]:
+ if self._hold_active and self.state == ParserState.TOOL_NAME:
+ candidate = "".join(self._held_name) + text
+ if not self._can_grow_into_declared_name(candidate):
+ # The held text can no longer become a declared tool
+ # name, so holding longer would only stall streaming.
+ # Release everything consumed so far as content.
+ return self._abort_hold("".join(self._held_raw) + text)
+ self._held_raw.append(text)
+ self._held_name.append(text)
+ self._held_events.append(
+ SemanticEvent(
+ EventType.TOOL_NAME,
+ value=text,
+ tool_index=self.tool_index,
+ )
+ )
+ return []
if self.state == ParserState.MESSAGE_HEADER:
self._message_header_buffer += text
return []
@@ -372,6 +423,24 @@ def _emit_for_state(self, text: str) -> list[SemanticEvent]:
content_type = self.config.content_events.get(self.state)
if content_type is not None:
return [SemanticEvent(content_type, value=text, tool_index=self.tool_index)]
+ if self._recovered_tool_call and self.state == ParserState.TOOL_BETWEEN:
+ # A response that lost its opening wrapper usually loses the
+ # closing one too, so text after a recovered invoke is often
+ # the rest of the answer rather than padding before the next
+ # invoke. Whitespace is held back because that is what
+ # padding looks like; as soon as anything else shows up the
+ # whole run is real output and goes out as content.
+ self._pending_between_text += text
+ if self._pending_between_text.strip():
+ held = self._pending_between_text
+ self._pending_between_text = ""
+ return [
+ SemanticEvent(
+ EventType.TEXT_CHUNK,
+ value=held,
+ tool_index=self.tool_index,
+ )
+ ]
return []
def _on_content(self, text: str) -> list[SemanticEvent]:
@@ -383,6 +452,87 @@ def _apply_transition(
self,
transition: Transition,
value: str,
+ ) -> list[SemanticEvent]:
+ if self._hold_active:
+ return self._resolve_hold(transition, value)
+ if transition.validate_tool_name:
+ if self.suppress_tool_calls or self.allowed_tool_names is None:
+ # Recovery could never be accepted for this request, so
+ # the trigger text stays plain content and nothing is
+ # buffered.
+ return self._emit_for_state(value)
+ return self._begin_hold(transition, value)
+ return self._run_transition(transition, value)
+
+ def _begin_hold(
+ self,
+ transition: Transition,
+ value: str,
+ ) -> list[SemanticEvent]:
+ """Apply a ``validate_tool_name`` transition but hold its events.
+
+ The events (and every TOOL_NAME chunk that follows) stay
+ buffered until the name completes and validates, so a false
+ positive can be undone without having emitted anything.
+ """
+ prior_state = self.state
+ prior_tool_index = self.tool_index
+ self._held_events = self._run_transition(transition, value)
+ self._held_raw = [value]
+ self._held_name = []
+ self._held_prior_state = prior_state
+ self._held_prior_tool_index = prior_tool_index
+ self._hold_active = True
+ self._recovered_tool_call = True
+ return []
+
+ def _resolve_hold(
+ self,
+ transition: Transition,
+ value: str,
+ ) -> list[SemanticEvent]:
+ """End the hold window at the name-completing transition."""
+ name = "".join(self._held_name)
+ allowed = self.allowed_tool_names
+ if allowed is not None and name in allowed:
+ events = self._held_events
+ self._clear_hold()
+ events.extend(self._run_transition(transition, value))
+ return events
+ return self._abort_hold("".join(self._held_raw) + value)
+
+ def _abort_hold(self, raw: str) -> list[SemanticEvent]:
+ """Discard held events and re-emit the raw text as content."""
+ self.state = self._held_prior_state
+ self.tool_index = self._held_prior_tool_index
+ self._recovered_tool_call = self._held_prior_state in self._TOOL_STATES
+ self._clear_hold()
+ return self._emit_for_state(raw)
+
+ def _clear_hold(self) -> None:
+ self._hold_active = False
+ self._held_events = []
+ self._held_raw = []
+ self._held_name = []
+
+ def _can_grow_into_declared_name(self, candidate: str) -> bool:
+ """Return True when *candidate* is a prefix of a declared tool name.
+
+ Consulted while a recovery hold is active. Membership in the
+ declared set is the only way a held name can validate, so once
+ the text seen so far stops being a prefix of any declared name
+ the caller aborts the hold. This also bounds how much text a
+ hold can buffer to the length of the longest declared name.
+ """
+ allowed = self.allowed_tool_names
+ if allowed is None:
+ return False
+ return any(name.startswith(candidate) for name in allowed)
+
+ def _run_transition(
+ self,
+ transition: Transition,
+ value: str,
) -> list[SemanticEvent]:
events: list[SemanticEvent] = []
previous_state = self.state
@@ -402,10 +552,17 @@ def _apply_transition(
)
self._args_buffer = ""
+ # Whatever is still held between invokes is whitespace padding,
+ # which the wrapped path drops too.
+ self._pending_between_text = ""
+
if previous_state == ParserState.MESSAGE_HEADER:
message_header = self._message_header_buffer
self._message_header_buffer = ""
+ if transition.next_state not in self._TOOL_STATES:
+ self._recovered_tool_call = False
+
self.state = transition.next_state
for event_type in transition.events:
diff --git a/vllm/tokenizers/deepseek_v4_encoding.py b/vllm/tokenizers/deepseek_v4_encoding.py
index 16bfa1a99a..52268e2197 100644
--- a/vllm/tokenizers/deepseek_v4_encoding.py
+++ b/vllm/tokenizers/deepseek_v4_encoding.py
@@ -351,8 +351,21 @@ def render_message(index: int, messages: List[Dict[str, Any]], thinking_mode: st
prompt += thinking_end_token if thinking_mode != "thinking" else thinking_start_token
prompt += task_sp_token
- elif messages[index].get("role") in ["user", "developer"]:
- # Normal generation: append Assistant + thinking token
+ elif messages[index].get("role") in ["user", "developer"] or (
+ messages[index].get("role") == "system" and index + 1 == len(messages)
+ ):
+ # Normal generation: append Assistant + thinking token.
+ #
+ # For user/developer this token is both the inter-turn separator and the
+ # generation prompt. A trailing `system` message needs only the latter:
+ # agent frameworks routinely append a context/reminder system message
+ # after the last user turn, and without the generation prompt the model
+ # sees no assistant boundary and continues the prompt as a document
+ # instead of answering. The `index + 1 == len(messages)` guard keeps a
+ # non-final system message from emitting a separator it does not own --
+ # a system message followed by `latest_reminder` reaches this branch
+ # (that role is exempt from the early return above), and unconditionally
+ # appending there corrupts the reference-golden layouts.
prompt += ASSISTANT_SP_TOKEN
if not drop_thinking and thinking_mode == "thinking":
prompt += thinking_start_token
diff --git a/vllm/tool_parsers/utils.py b/vllm/tool_parsers/utils.py
index a11d4a9eec..3900bb0fe7 100644
--- a/vllm/tool_parsers/utils.py
+++ b/vllm/tool_parsers/utils.py
@@ -289,6 +289,23 @@ def find_tool_properties(
return {}
+def collect_tool_names(tools: list[Tool] | None) -> frozenset[str]:
+ """Collect the names of all declared function tools."""
+ if not tools:
+ return frozenset()
+ names: set[str] = set()
+ for tool in tools:
+ if isinstance(tool, (FunctionTool, NamespaceTool)):
+ for name, _ in iter_response_function_tool_info(tool):
+ names.add(name)
+ continue
+ if not _is_function_tool(tool):
+ continue
+ name, _ = _extract_tool_info(tool)
+ names.add(name)
+ return frozenset(names)
+
+
def find_tool_name(
tools: list[Tool] | None,
tool_name: str,
-