Script 'mail_helper' called by obssrc Hello community, here is the log from the commit of package python-vllm for openSUSE:Factory checked in at 2026-10-12 22:07:34 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Comparing /work/SRC/openSUSE:Factory/python-vllm (Old) and /work/SRC/openSUSE:Factory/.python-vllm.new.3698736 (New) ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
Package is "python-vllm" Mon Oct 12 22:07:34 2026 rev:11 rq:1384276 version:0.30.0 Changes: -------- --- /work/SRC/openSUSE:Factory/python-vllm/python-vllm.changes 2026-09-23 14:36:35.030398193 +0200 +++ /work/SRC/openSUSE:Factory/.python-vllm.new.3698736/python-vllm.changes 2026-10-12 22:07:35.095219508 +0200 @@ -1,0 +2,17 @@ +Mon Oct 12 08:32:09 UTC 2026 - Martin Pluskal <[email protected]> + +- CVE-2026-105775: Mamba "all"-mode block size not a multiple + of the chunk size corrupts the state write-back and takes + EngineCore down (boo#1284279) + * fix-CVE-2026-105775.patch +- CVE-2026-105922: prompt-embeds positions fed stale token IDs + into the penalty kernels, tripping a device-side assert in + EngineCore (boo#1284365) + * fix-CVE-2026-105922.patch + +------------------------------------------------------------------- +Thu Sep 24 07:28:25 UTC 2026 - Andreas Schwab <[email protected]> + +- Enable build on riscv64 + +------------------------------------------------------------------- New: ---- fix-CVE-2026-105775.patch fix-CVE-2026-105922.patch ----------(New B)---------- New: EngineCore down (boo#1284279) * fix-CVE-2026-105775.patch - CVE-2026-105922: prompt-embeds positions fed stale token IDs New: EngineCore (boo#1284365) * fix-CVE-2026-105922.patch ----------(New E)---------- ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Other differences: ------------------ ++++++ python-vllm.spec ++++++ --- /var/tmp/diff_new_pack.IMQAbZ/_old 2026-10-12 22:07:36.307269539 +0200 +++ /var/tmp/diff_new_pack.IMQAbZ/_new 2026-10-12 22:07:36.309269622 +0200 @@ -70,6 +70,10 @@ Patch1: vllm-cpu-disable-rust-frontend.patch # PATCH-FIX-UPSTREAM vllm-thinking-budget-scan-cursor.patch boo#1280982 [email protected] -- CVE-2026-92365: resume the thinking-budget marker search from a cursor so a generation is not quadratic in its own length Patch2: vllm-thinking-budget-scan-cursor.patch +# PATCH-FIX-UPSTREAM fix-CVE-2026-105775.patch https://github.com/vllm-project/vllm/pull/57278 (unmerged as of 2026-10-12, rebased onto 0.30.0) -- CVE-2026-105775: keep the mamba block size a multiple of the chunk size so the "all" write-back reads states off chunk ends; DROP as obsolete when updating past the upstream removal of mamba_cache_mode "all" (#58997) +Patch3: fix-CVE-2026-105775.patch +# PATCH-FIX-UPSTREAM fix-CVE-2026-105922.patch https://github.com/vllm-project/vllm/pull/57762 (open/unmerged as of 2026-10-12, rebased onto 0.30.0) -- CVE-2026-105922: treat prompt-embedding positions as invalid token IDs so penalty kernels do not index stale CPU token buffer contents +Patch4: fix-CVE-2026-105922.patch BuildRequires: %{python_module Jinja2} # Runtime modules imported eagerly by "import vllm" (needed for the %%check smoke test). BuildRequires: %{python_module cachetools} @@ -184,7 +188,7 @@ Requires: python-watchfiles Requires: python-xgrammar >= 0.2.1 # Limited to the arches python-torch is built for. -ExclusiveArch: x86_64 aarch64 +ExclusiveArch: x86_64 aarch64 riscv64 %if %{with cpu_kernels} # The compiled kernels need a C++ toolchain and torch's CMake package config, # which ships in the -devel subpackage (find_package(Torch)). @@ -255,6 +259,9 @@ %else export FETCHCONTENT_SOURCE_DIR_ONEDNN="$(readlink -f ../oneDNN-%{onednn_x86_version})" %endif +%ifarch riscv64 +export CMAKE_ARGS=-DVLLM_RVV_VLEN=256 +%endif export CMAKE_GENERATOR=Ninja export MAX_JOBS=%{?jobs:%{jobs}}%{!?jobs:4} # cpu_extension.cmake does a bare find_library(OPEN_MP NAMES gomp REQUIRED), @@ -307,6 +314,13 @@ # Import smoke test validating the package and its metadata load. # Full model/serving tests need model weights and are out of scope here. %python_expand $python -c "import vllm; print(vllm.__version__)" +# Regression test for CVE-2026-105775 (CPU-marked, needs only pytest+torch): +# run it when pytest is already in the build root; no test-only +# BuildRequires are added for it. +%python_expand if $python -c "import pytest" 2>/dev/null; then $python -m pytest tests/v1/core/test_mamba_block_size_alignment.py; fi +# The CUDA penalties test from fix-CVE-2026-105922 +# (tests/v1/worker/test_gpu_penalties.py) is carried-not-run: it needs a +# CUDA device, which OBS build roots do not have. %pre # Migrate an install that still carries the old update-alternatives symlink. ++++++ fix-CVE-2026-105775.patch ++++++ From: Tianyao Wu <[email protected]> Subject: [PATCH] [Bugfix][Mamba] Keep the mamba block size a multiple of the chunk size References: CVE-2026-105775, boo#1284279 Upstream-PR: https://github.com/vllm-project/vllm/pull/57278 (CLOSED UNMERGED, needs-rebase; unmerged as of 2026-10-12, rebased onto 0.30.0) With mamba_cache_mode "all", _align_hybrid_block_size drops the model's mamba chunk size from the block alignment as soon as --mamba-block-size is given. That size has to stay a multiple of the chunk size: the "all" write-back reads every state it caches off a chunk end, walking the intermediate states with a stride of mamba_block_size // chunk_size chunks per block, and chunk ends sit at absolute sequence positions that are multiples of the chunk size. Once the block size is not a multiple of the chunk size that stride no longer matches the block spacing, so a slot is filled from a chunk end at a different offset or from a neighbouring sequence's chunks, and nothing raises. Walking past the end of the intermediate states yields an empty slice rather than an error, and the state write then fails with a shape mismatch inside EngineCore, which does not recover. Fold the chunk size back into the alignment: an explicit --mamba-block-size widens it instead of replacing it, so every configuration that is correct today derives the same block size as before. On Falcon-H1-Tiny-90M-Instruct (chunk size 128), --mamba-block-size 96 derived 288 and now derives 384. Check the invariant once where the write-back's chunk size is resolved, so that a future reintroduction fails at start-up instead of corrupting state. NOTE: upstream later removed mamba_cache_mode "all" entirely (#58997). That removal is NOT backported here (behavior break); this patch keeps "all" working correctly. DROP this patch as obsolete when updating past the removal. Co-authored-by: Claude <[email protected]> Signed-off-by: Tianyao Wu <[email protected]> --- .../core/test_mamba_block_size_alignment.py | 115 ++++++++++++++++++ vllm/config/cache.py | 7 +- vllm/platforms/interface.py | 12 +- vllm/v1/attention/backends/mamba2_attn.py | 8 ++ 4 files changed, 139 insertions(+), 3 deletions(-) create mode 100644 tests/v1/core/test_mamba_block_size_alignment.py diff --git a/tests/v1/core/test_mamba_block_size_alignment.py b/tests/v1/core/test_mamba_block_size_alignment.py new file mode 100644 index 000000000000..29389940a700 --- /dev/null +++ b/tests/v1/core/test_mamba_block_size_alignment.py @@ -0,0 +1,115 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Mamba "all" mode: the block size has to stay a multiple of the chunk size. + +In "all" mode the cached states are read off chunk ends: the write-back in +``MambaMixer2.conv_ssm_forward`` walks the intermediate states with a stride of +``mamba_block_size // chunk_size`` chunks per block. A block size that is not a +multiple of the chunk size makes that stride drift away from the block spacing, +so a slot is filled from a chunk end at the wrong offset, and when the walk +leaves the intermediate states the slice is empty rather than an error -- the +state write then dies with a shape mismatch and takes EngineCore with it +(#57266). + +``--mamba-block-size`` used to pick the alignment by itself, which is how a +block size of 288 tokens against a chunk size of 128 was reachable. +""" + +from types import SimpleNamespace + +import pytest +import torch + +from vllm.config import CacheConfig, VllmConfig +from vllm.platforms.interface import Platform +from vllm.v1.attention.backend import MultipleOf + +pytestmark = pytest.mark.cpu_test + +# Falcon-H1-Tiny-90M-Instruct, the smallest model that reproduces #57266. +MAMBA_CHUNK_SIZE = 128 +CONV_STATE_SHAPE = (896, 3) +SSM_STATE_SHAPE = (24, 32, 64) +NUM_KV_HEADS = 2 +HEAD_SIZE = 64 + +# (--mamba-block-size, derived mamba block size). The mamba state is 103680 +# bytes and a token of attention is 512, so 203 tokens of attention pay for one +# state and the derivation rounds that up to the alignment. Requests that are +# not a multiple of the chunk size round up to a common multiple, which can be +# well above the 256 the constraint alone would need: 176 -> 1408. +DERIVED_BLOCK_SIZES = [ + (None, 256), + (16, 256), + (32, 256), + (64, 256), + (96, 384), + (128, 256), + (176, 1408), + (256, 256), +] + + +class _FakeModel: + @classmethod + def get_mamba_state_shape_from_config( + cls, vllm_config: VllmConfig + ) -> tuple[tuple[int, ...], ...]: + return (CONV_STATE_SHAPE, SSM_STATE_SHAPE) + + @classmethod + def get_mamba_state_dtype_from_config( + cls, vllm_config: VllmConfig + ) -> tuple[torch.dtype, ...]: + return (torch.bfloat16, torch.bfloat16) + + +class _FakeBackend: + @staticmethod + def customize_spec(spec): + return spec + + @staticmethod + def get_supported_kernel_block_sizes() -> list[MultipleOf]: + return [MultipleOf(16)] + + +def _hybrid_config( + monkeypatch: pytest.MonkeyPatch, mamba_block_size: int | None +) -> VllmConfig: + cache_config = CacheConfig( + enable_prefix_caching=True, mamba_block_size=mamba_block_size + ) + cache_config.mamba_cache_mode = "all" + vllm_config = VllmConfig(cache_config=cache_config) + vllm_config.model_config = SimpleNamespace( + architecture="FalconH1ForCausalLM", + dtype=torch.bfloat16, + use_mla=False, + get_num_kv_heads=lambda parallel_config: NUM_KV_HEADS, + get_head_size=lambda: HEAD_SIZE, + get_mamba_chunk_size=lambda: MAMBA_CHUNK_SIZE, + ) + monkeypatch.setattr( + "vllm.model_executor.models.ModelRegistry.resolve_model_cls", + lambda architecture, model_config: (_FakeModel, architecture), + ) + # The public entry point needs backend discovery, so drive the aligner that + # owns the derivation directly. + return vllm_config + + [email protected]("requested,expected", DERIVED_BLOCK_SIZES) +def test_mamba_all_mode_block_size_is_chunk_aligned( + monkeypatch: pytest.MonkeyPatch, requested: int | None, expected: int +) -> None: + vllm_config = _hybrid_config(monkeypatch, requested) + + Platform._align_hybrid_block_size(vllm_config, _FakeBackend) + + block_size = vllm_config.cache_config.mamba_block_size + assert block_size % MAMBA_CHUNK_SIZE == 0, ( + f"--mamba-block-size {requested} produced {block_size}, " + f"which no chunk end lands on" + ) + assert block_size == expected diff --git a/vllm/config/cache.py b/vllm/config/cache.py index 37ad5b2b34fe..e2b6dd3d1bd9 100644 --- a/vllm/config/cache.py +++ b/vllm/config/cache.py @@ -166,7 +166,12 @@ class CacheConfig: mamba_block_size: int | None = Field(default=None, gt=0) """Size of a contiguous cache block in number of tokens for mamba cache. Can be set only when prefix caching is enabled. - Value must be a multiple of 8 to align with causal_conv1d kernel.""" + Value must be a multiple of 8 to align with causal_conv1d kernel. + With `mamba_cache_mode` "all" the effective block size is a multiple of + both this value and the model's mamba chunk size, since that is where the + cached states are read from, and it also sets the attention block size. + With "align" it is ignored: the mamba block size follows the attention + block size.""" mamba_cache_dtype: MambaDType = "auto" """The data type to use for the Mamba cache (both the conv as well as the ssm state). If set to 'auto', the data type will be inferred from the model diff --git a/vllm/platforms/interface.py b/vllm/platforms/interface.py index 13e5b55402c6..596031b719a6 100644 --- a/vllm/platforms/interface.py +++ b/vllm/platforms/interface.py @@ -896,8 +896,16 @@ def _align_hybrid_block_size( # TODO(tdoublep): this constraint can be relaxed fairly # easily by changing the way we layout chunks in the # mamba2 kernels. - base_chunk_size = mamba_block_size or model_config.get_mamba_chunk_size() - assert base_chunk_size is not None + model_chunk_size = model_config.get_mamba_chunk_size() + # The "all" write-back reads the states it caches off chunk ends, + # so a block boundary that is not also a chunk boundary has no + # state to read. An explicit --mamba-block-size therefore only + # widens the alignment; it cannot replace the chunk size. + base_chunk_size = ( + lcm(mamba_block_size, model_chunk_size) + if mamba_block_size is not None + else model_chunk_size + ) attn_tokens_per_mamba_state = cdiv(mamba_page_size, attn_page_size_1_token) chunk_size = lcm(base_chunk_size, kernel_block_alignment_size) attn_block_size = chunk_size * cdiv(attn_tokens_per_mamba_state, chunk_size) diff --git a/vllm/v1/attention/backends/mamba2_attn.py b/vllm/v1/attention/backends/mamba2_attn.py index 0a4584632dca..82261e3c065d 100644 --- a/vllm/v1/attention/backends/mamba2_attn.py +++ b/vllm/v1/attention/backends/mamba2_attn.py @@ -128,6 +128,14 @@ def __init__( "chunk_size needs to be set in the model config for Mamba2 models" ) self.chunk_size: int = chunk_size + if vllm_config.cache_config.mamba_cache_mode == "all": + # "all" mode caches states read off chunk ends, so a block boundary + # that is not also a chunk boundary has no state to read from. + assert kv_cache_spec.block_size % chunk_size == 0, ( + f"mamba_cache_mode 'all' needs the mamba block size " + f"({kv_cache_spec.block_size}) to be a multiple of the mamba " + f"chunk size ({chunk_size})" + ) def build( self, ++++++ fix-CVE-2026-105922.patch ++++++ From: cross2pro <[email protected]> Subject: [PATCH] Fix penalties with prompt embeddings References: CVE-2026-105922, boo#1284365 Upstream-PR: https://github.com/vllm-project/vllm/pull/57762 (OPEN, blocked on rebase; open/unmerged as of 2026-10-12, rebased onto 0.30.0) Treat prompt embedding positions as invalid token IDs when building sampling metadata. This prevents penalty kernels from indexing stale CPU token buffer contents. MRV2 writes placeholder zeros into all_token_ids at prompt-embedding positions (vllm/v1/request.py), so the penalty bincount kernel counted token 0 as a prompt token for every prompt-embeds request, wrongly applying the repetition penalty to its logit. Consult the per-position is-token-ids mask that PromptEmbedsState already maintains on device: skip embed positions for mixed prompts and the whole prompt for pure prompt-embeds requests. Rebased onto 0.30.0: the MRV2 sampler/penalties plumbing in this tree still uses RequestState directly (no LogitsProcRequestState yet), so the PenaltiesState/Sampler signatures are extended in place and the model_runner passes model_state.prompt_embeds_state through; the triton bincount masking and the gpu_input_batch CPU-tensor pad are upstream's. Co-authored-by: Nick Hill <[email protected]> Signed-off-by: cross2pro <[email protected]> --- diff -Naur a/tests/v1/worker/test_gpu_input_batch.py b/tests/v1/worker/test_gpu_input_batch.py --- a/tests/v1/worker/test_gpu_input_batch.py 2026-09-22 07:00:46.000000000 +0200 +++ b/tests/v1/worker/test_gpu_input_batch.py 2026-10-12 10:29:33.473825773 +0200 @@ -511,3 +511,35 @@ assert metadata.get_prompt_token_ids_cpu()[0].tolist() == req.prompt_token_ids else: assert metadata.prompt_token_ids_cpu is None + + +def test_sampling_metadata_prompt_embeds_use_vocab_size_as_token_id_pad(): + input_batch = InputBatch( + max_num_reqs=2, + max_model_len=8, + max_num_batched_tokens=8, + device=torch.device("cpu"), + vocab_size=VOCAB_SIZE, + block_sizes=[16], + kernel_block_sizes=[16], + max_num_blocks_per_req=[1], + ) + input_batch.token_ids_cpu[:].fill(VOCAB_SIZE + 1) + req = CachedRequestState( + req_id="prompt-embeds", + prompt_token_ids=None, + prompt_embeds=torch.zeros((3, 4)), + mm_features=[], + sampling_params=SamplingParams(presence_penalty=0.1), + pooling_params=None, + block_ids=([],), + generator=None, + num_computed_tokens=0, + output_token_ids=[], + ) + input_batch.add_request(req) + + metadata = input_batch._make_sampling_metadata() + + assert metadata.prompt_token_ids is not None + assert metadata.prompt_token_ids[0].tolist() == [VOCAB_SIZE] * 3 diff -Naur a/tests/v1/worker/test_gpu_penalties.py b/tests/v1/worker/test_gpu_penalties.py --- a/tests/v1/worker/test_gpu_penalties.py 1970-01-01 01:00:00.000000000 +0100 +++ b/tests/v1/worker/test_gpu_penalties.py 2026-10-12 10:29:44.923112854 +0200 @@ -0,0 +1,93 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +import numpy as np +import pytest +import torch + +pytest.importorskip("triton") +if not torch.cuda.is_available(): + pytest.skip( + "CUDA required for Model Runner V2 penalties tests", + allow_module_level=True, + ) + +from vllm.sampling_params import SamplingParams +from vllm.v1.core.sched.output import NewRequestData +from vllm.v1.worker.gpu.model_states.prompt_embeds import PromptEmbedsState +from vllm.v1.worker.gpu.sample.penalties import PenaltiesState +from vllm.v1.worker.gpu.states import RequestState + +DEVICE = torch.device("cuda") +VOCAB_SIZE = 128 +HIDDEN_SIZE = 8 + + +def _new_req_data( + req_id: str, + embeds_len: int = 0, + prompt_is_token_ids: list[bool] | None = None, +) -> NewRequestData: + return NewRequestData( + req_id=req_id, + prompt_token_ids=None, + mm_features=[], + sampling_params=None, + pooling_params=None, + block_ids=([],), + num_computed_tokens=0, + lora_request=None, + prompt_embeds=(torch.zeros(embeds_len, HIDDEN_SIZE) if embeds_len else None), + prompt_is_token_ids=prompt_is_token_ids, + ) + + +def _prompt_bin_tokens(penalties: PenaltiesState, req_idx: int) -> set[int]: + words = penalties.prompt_bin_mask[req_idx].cpu().numpy().astype(np.uint32) + return {t for t in range(VOCAB_SIZE) if words[t // 32] & (1 << (t % 32))} + + +def test_bincount_ignores_prompt_embeds_positions(): + """Prompt-embedding positions must not contribute to the prompt bin mask. + + `all_token_ids` holds placeholder zeros at prompt-embedding positions, so + without masking the bincount would count token 0 as a prompt token and the + repetition penalty would wrongly scale its logit. + """ + req_states = RequestState( + max_num_reqs=3, + max_model_len=64, + max_num_batched_tokens=64, + num_speculative_steps=1, + vocab_size=VOCAB_SIZE, + device=DEVICE, + ) + pe_state = PromptEmbedsState(3, HIDDEN_SIZE, torch.bfloat16, DEVICE) + + # Ordinary token prompt. + req_states.add_request("tok", 3, [5, 6, 7], 0, 16) + # Pure prompt-embeds prompt: all positions are placeholder zeros. + req_states.add_request("embeds", 3, [0, 0, 0], 0, 16) + # Mixed prompt: position 1 is an embedding (placeholder zero). + req_states.add_request("mixed", 3, [9, 0, 10], 0, 16) + pe_state.add_request(req_states.req_id_to_index["tok"], _new_req_data("tok")) + pe_state.add_request( + req_states.req_id_to_index["embeds"], _new_req_data("embeds", embeds_len=3) + ) + pe_state.add_request( + req_states.req_id_to_index["mixed"], + _new_req_data("mixed", embeds_len=3, prompt_is_token_ids=[True, False, True]), + ) + req_states.apply_staged_writes() + pe_state.apply_staged_writes() + + penalties = PenaltiesState(req_states, pe_state) + params = SamplingParams(repetition_penalty=1.1) + for req_idx in range(3): + penalties.add_request(req_idx, params) + penalties.apply_staged_writes() + + idx = req_states.req_id_to_index + assert _prompt_bin_tokens(penalties, idx["tok"]) == {5, 6, 7} + assert _prompt_bin_tokens(penalties, idx["embeds"]) == set() + assert _prompt_bin_tokens(penalties, idx["mixed"]) == {9, 10} diff -Naur a/vllm/v1/worker/gpu/model_runner.py b/vllm/v1/worker/gpu/model_runner.py --- a/vllm/v1/worker/gpu/model_runner.py 2026-09-22 07:00:46.000000000 +0200 +++ b/vllm/v1/worker/gpu/model_runner.py 2026-10-12 10:29:09.075082970 +0200 @@ -469,6 +469,7 @@ "enable_trace_replay": self.model_config.enable_trace_replay, "reasoning_config": self.vllm_config.reasoning_config, "return_sampling_mask": self.model_config.return_sampling_mask, + "prompt_embeds_state": self.model_state.prompt_embeds_state, } if self.vllm_config.watermark_config is None: self.sampler = Sampler(**sampler_kwargs) diff -Naur a/vllm/v1/worker/gpu/model_states/interface.py b/vllm/v1/worker/gpu/model_states/interface.py --- a/vllm/v1/worker/gpu/model_states/interface.py 2026-09-22 06:51:49.000000000 +0200 +++ b/vllm/v1/worker/gpu/model_states/interface.py 2026-10-12 10:29:21.748507663 +0200 @@ -1,7 +1,7 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project from abc import ABC, abstractmethod -from typing import Any, ClassVar, cast +from typing import TYPE_CHECKING, Any, ClassVar, cast import torch import torch.nn as nn @@ -23,6 +23,9 @@ from vllm.v1.worker.gpu.states import RequestState from vllm.v1.worker.utils import AttentionGroup +if TYPE_CHECKING: + from vllm.v1.worker.gpu.model_states.prompt_embeds import PromptEmbedsState + class ModelSpecificAttnMetadata: """Base class for model-specific attention metadata.""" @@ -46,6 +49,10 @@ supports_prompt_embeds: ClassVar[bool] = False """Whether this state implements user-provided prompt embeddings.""" + prompt_embeds_state: "PromptEmbedsState | None" = None + """GPU state for prompt-embeds overlays; only set by states that support + user-provided prompt embeddings.""" + def __init__( self, vllm_config: VllmConfig, diff -Naur a/vllm/v1/worker/gpu/sample/penalties.py b/vllm/v1/worker/gpu/sample/penalties.py --- a/vllm/v1/worker/gpu/sample/penalties.py 2026-09-22 07:00:46.000000000 +0200 +++ b/vllm/v1/worker/gpu/sample/penalties.py 2026-10-12 10:28:54.045393396 +0200 @@ -1,5 +1,7 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +from typing import TYPE_CHECKING + import numpy as np import torch @@ -10,10 +12,18 @@ from vllm.v1.worker.gpu.buffer_utils import UvaBackedTensor from vllm.v1.worker.gpu.states import RequestState +if TYPE_CHECKING: + from vllm.v1.worker.gpu.model_states.prompt_embeds import PromptEmbedsState + class PenaltiesState: - def __init__(self, req_states: RequestState): + def __init__( + self, + req_states: RequestState, + prompt_embeds_state: "PromptEmbedsState | None" = None, + ): self.req_states = req_states + self.prompt_embeds_state = prompt_embeds_state max_num_reqs = req_states.max_num_reqs self.vocab_size = req_states.vocab_size @@ -63,6 +73,7 @@ prefill_lens = self.req_states.prefill_len.np[self._new_penalties_reqs] max_prefill_len = int(prefill_lens.max()) + pe_state = self.prompt_embeds_state bincount( idx_mapping, self.req_states.all_token_ids.gpu, @@ -71,6 +82,8 @@ self.prompt_bin_mask, self.output_bin_counts, max_prefill_len, + pe_state.embeds_lens.gpu if pe_state is not None else None, + pe_state.mask_ptrs.gpu if pe_state is not None else None, ) self._new_penalties_reqs.clear() @@ -222,10 +235,13 @@ all_token_ids_stride, prompt_len_ptr, prefill_len_ptr, + prompt_embeds_lens_ptr, + prompt_embeds_mask_ptrs_ptr, prompt_bin_mask_ptr, prompt_bin_mask_stride, output_bin_counts_ptr, output_bin_counts_stride, + HAS_PROMPT_EMBEDS: tl.constexpr, BLOCK_SIZE: tl.constexpr, ): token_idx = tl.program_id(0) @@ -240,6 +256,19 @@ block = block_idx * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) if block_idx * BLOCK_SIZE < prompt_len: mask = block < prompt_len + if HAS_PROMPT_EMBEDS and tl.load(prompt_embeds_lens_ptr + req_state_idx) > 0: + # Prompt-embedding positions hold placeholder zeros in + # all_token_ids; exclude them so token 0 is not spuriously + # counted as a prompt token. + mask_ptr_int = tl.load(prompt_embeds_mask_ptrs_ptr + req_state_idx) + if mask_ptr_int == 0: + # The whole prompt is embeddings: no token ids to count. + mask = tl.full([BLOCK_SIZE], 0, tl.int1) + else: + # Mixed prompt: the mask marks real token-id positions. + is_tok_ptr = mask_ptr_int.to(tl.pointer_type(tl.uint8)) + is_tok = tl.load(is_tok_ptr + block, mask=mask, other=0) + mask = mask & (is_tok != 0) prompt_tokens = tl.load( all_token_ids_ptr + req_state_idx * all_token_ids_stride + block, mask=mask ) @@ -275,6 +304,8 @@ prompt_bin_mask: torch.Tensor, output_bin_counts: torch.Tensor, max_prefill_len: int, + prompt_embeds_lens: torch.Tensor | None = None, + prompt_embeds_mask_ptrs: torch.Tensor | None = None, ) -> None: # Use index_fill_ instead of `tensor[idx] = 0` to avoid sync. idx_long = expanded_idx_mapping.long() @@ -283,16 +314,20 @@ num_tokens = expanded_idx_mapping.shape[0] BLOCK_SIZE = 1024 num_blocks = triton.cdiv(max_prefill_len, BLOCK_SIZE) + has_prompt_embeds = prompt_embeds_lens is not None _bincount_kernel[(num_tokens, num_blocks)]( expanded_idx_mapping, all_token_ids, all_token_ids.stride(0), prompt_len, prefill_len, + prompt_embeds_lens if has_prompt_embeds else expanded_idx_mapping, + prompt_embeds_mask_ptrs if has_prompt_embeds else expanded_idx_mapping, prompt_bin_mask, prompt_bin_mask.stride(0), output_bin_counts, output_bin_counts.stride(0), + HAS_PROMPT_EMBEDS=has_prompt_embeds, BLOCK_SIZE=BLOCK_SIZE, ) diff -Naur a/vllm/v1/worker/gpu/sample/sampler.py b/vllm/v1/worker/gpu/sample/sampler.py --- a/vllm/v1/worker/gpu/sample/sampler.py 2026-09-22 07:00:46.000000000 +0200 +++ b/vllm/v1/worker/gpu/sample/sampler.py 2026-10-12 10:30:14.494437137 +0200 @@ -1,6 +1,8 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +from typing import TYPE_CHECKING + import numpy as np import torch @@ -29,6 +31,9 @@ from vllm.v1.worker.gpu.sample.trace_replay import TraceReplayState from vllm.v1.worker.gpu.states import RequestState +if TYPE_CHECKING: + from vllm.v1.worker.gpu.model_states.prompt_embeds import PromptEmbedsState + class Sampler: def __init__( @@ -43,6 +48,7 @@ enable_trace_replay: bool = False, reasoning_config: ReasoningConfig | None = None, return_sampling_mask: bool = False, + prompt_embeds_state: "PromptEmbedsState | None" = None, ): self.logprobs_mode = logprobs_mode self.compute_nans = envs.VLLM_COMPUTE_NANS_IN_LOGITS # False by default. @@ -50,7 +56,9 @@ self.req_states = req_states self.sampling_states = SamplingStates(max_num_reqs, vocab_size) - self.penalties_state = PenaltiesState(req_states) + self.penalties_state = PenaltiesState( + req_states, prompt_embeds_state + ) self.logit_bias_state = LogitBiasState(max_num_reqs, device) self.bad_words_state = BadWordsState(req_states) self.logprob_token_ids_state = LogprobTokenIdsState(max_num_reqs, device) diff -Naur a/vllm/v1/worker/gpu_input_batch.py b/vllm/v1/worker/gpu_input_batch.py --- a/vllm/v1/worker/gpu_input_batch.py 2026-09-22 07:00:46.000000000 +0200 +++ b/vllm/v1/worker/gpu_input_batch.py 2026-10-12 10:28:35.011253658 +0200 @@ -996,6 +996,9 @@ # token_id of this value. for i in range(num_reqs): prompt_token_ids[i, self.num_prompt_tokens[i] :] = self.vocab_size + prompt_token_ids[~self.is_token_ids[:num_reqs, :max_prompt_len]] = ( + self.vocab_size + ) return prompt_token_ids_cpu_tensor def make_lora_inputs(
