@tduncan / Nvidia-Optimization-Stack-VLLM / commits / 9dd969da09

[Model][ROCm] Enable DeepSeek V4 Vision (#55107)

Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Signed-off-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: OpenAI Codex <codex@openai.com>

Andreas Karatzas committed Sep 11, 2026 at 11:17 UTC 9dd969da096e37256ee37e24f6a4689d860f39ce
17 files changed +1055 -347
.buildkite/test_areas/models_basic.yaml
+4 -2
@@ -48,8 +48,9 @@ steps:
48 - tests/models/test_terratorch.py
49 - tests/models/transformers/test_backend.py
50 - tests/models/test_registry.py
51 + - tests/models/test_deepseek_v4_vl_rocm.py
52 commands:
52 - - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
53 + - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py models/test_deepseek_v4_vl_rocm.py
54 mirror:
55 amd:
56 label: ":amd: (MI300) Basic Models (Other)"
@@ -59,9 +60,10 @@ steps:
60 depends_on:
61 - image-build-amd
62 source_file_dependencies:
63 + - tests/models/test_deepseek_v4_vl_rocm.py
64 - tests/models/test_hyv4_rocm.py
65 commands:
64 - - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
66 + - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py models/test_deepseek_v4_vl_rocm.py
67 - VLLM_ROCM_USE_AITER=1 pytest -v -s models/test_hyv4_rocm.py -m 'not distributed'
68
69 - label: ":nvidia: (B200) Inkling"
tests/kernels/attention/test_rocm_triton_attn_dsv4.py
+95
@@ -336,6 +336,101 @@ def test_compute_global_topk_ragged_indices_and_indptr() -> None:
336 torch.testing.assert_close(actual_lens, expected_lens)
337
338
339 +@torch.inference_mode()
340 +def test_combine_topk_swa_indices_adds_image_visibility() -> None:
341 + from vllm.models.deepseek_v4.amd.rocm import combine_topk_swa_indices
342 +
343 + device = torch.device("cuda")
344 + num_tokens = 8
345 + topk_indices = torch.full((num_tokens, 1), -1, dtype=torch.int32, device=device)
346 + query_start_loc = torch.tensor([0, num_tokens], dtype=torch.int32, device=device)
347 + seq_lens = torch.tensor([num_tokens], dtype=torch.int32, device=device)
348 + gather_lens = seq_lens.clone()
349 + left_visible = torch.tensor(
350 + [0, 0, 0, 1, 2, 3, 4, 0], dtype=torch.int32, device=device
351 + )
352 + right_visible = torch.tensor(
353 + [0, 0, 4, 3, 2, 1, 0, 0], dtype=torch.int32, device=device
354 + )
355 +
356 + indices, lens = combine_topk_swa_indices(
357 + topk_indices,
358 + query_start_loc,
359 + seq_lens,
360 + gather_lens,
361 + window_size=4,
362 + compress_ratio=1,
363 + topk=0,
364 + M=16,
365 + N=0,
366 + max_image_tokens=5,
367 + left_visible=left_visible,
368 + right_visible=right_visible,
369 + )
370 +
371 + expected_rows = [
372 + [0],
373 + [0, 1],
374 + [0, 1, 2, 3, 4, 5, 6],
375 + [0, 1, 2, 3, 4, 5, 6],
376 + [1, 2, 3, 4, 5, 6],
377 + [2, 3, 4, 5, 6],
378 + [2, 3, 4, 5, 6],
379 + [4, 5, 6, 7],
380 + ]
381 + for token_idx, expected in enumerate(expected_rows):
382 + actual = indices[token_idx, : lens[token_idx]].cpu().tolist()
383 + assert actual == expected
384 +
385 +
386 +@torch.inference_mode()
387 +def test_combine_topk_swa_indices_apc_hit_inside_image() -> None:
388 + from vllm.models.deepseek_v4.amd.rocm import combine_topk_swa_indices
389 +
390 + device = torch.device("cuda")
391 + indices, lens = combine_topk_swa_indices(
392 + torch.full((2, 1), -1, dtype=torch.int32, device=device),
393 + torch.tensor([0, 2], dtype=torch.int32, device=device),
394 + torch.tensor([10], dtype=torch.int32, device=device),
395 + # Only positions [5, 10) exist in the gathered SWA workspace.
396 + torch.tensor([5], dtype=torch.int32, device=device),
397 + window_size=4,
398 + compress_ratio=1,
399 + topk=0,
400 + M=10,
401 + N=0,
402 + max_image_tokens=10,
403 + left_visible=torch.tensor([8, 9], dtype=torch.int32, device=device),
404 + # Deliberately extends beyond seq_len to exercise the upper clamp too.
405 + right_visible=torch.tensor([5, 5], dtype=torch.int32, device=device),
406 + )
407 +
408 + assert lens.cpu().tolist() == [5, 5]
409 + assert indices[:, :5].cpu().tolist() == [list(range(5)), list(range(5))]
410 +
411 +
412 +@torch.inference_mode()
413 +def test_combine_topk_swa_indices_keeps_vision_row_width_without_images() -> None:
414 + from vllm.models.deepseek_v4.amd.rocm import combine_topk_swa_indices
415 +
416 + device = torch.device("cuda")
417 + indices, lens = combine_topk_swa_indices(
418 + torch.full((1, 1), -1, dtype=torch.int32, device=device),
419 + torch.tensor([0, 1], dtype=torch.int32, device=device),
420 + torch.tensor([1], dtype=torch.int32, device=device),
421 + torch.tensor([1], dtype=torch.int32, device=device),
422 + window_size=120,
423 + compress_ratio=1,
424 + topk=0,
425 + M=256,
426 + N=0,
427 + max_image_tokens=16,
428 + )
429 +
430 + assert indices.shape == (1, 256)
431 + assert lens.item() == 1
432 +
433 +
434 def test_extra_cache_nan_free_provenance_gate(monkeypatch) -> None:
435 from vllm.models.deepseek_v4.amd import rocm as mod
436
tests/models/multimodal/processing/test_tensor_schema.py
+4 -3
@@ -164,10 +164,11 @@ def test_model_tensor_schema(model_id: str):
164 "Kimi-K2.5's offline inference has issues about vision chunks. Fix later."
165 )
166
167 - if model_id == "deepseek-ai/DeepSeek-V4-Flash-Vision-Exp" and not (
168 - current_platform.is_cuda()
167 + if (
168 + model_id == "deepseek-ai/DeepSeek-V4-Flash-Vision-Exp"
169 + and not current_platform.is_cuda_alike()
170 ):
170 - pytest.skip("Deepseek V4 is only supported on CUDA")
171 + pytest.skip("Deepseek V4 vision is only supported on CUDA and ROCm")
172
173 if model_id == "zai-org/GLM-5.3-Flash" and (current_platform.is_xpu()):
174 pytest.skip("GLM-5.3-Flash is not supported on XPU")
tests/models/test_deepseek_v4_vl_rocm.py new
+458
@@ -0,0 +1,458 @@
1 +# SPDX-License-Identifier: Apache-2.0
2 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3 +
4 +from types import SimpleNamespace
5 +
6 +import pytest
7 +import torch
8 +from torch import nn
9 +
10 +from vllm.model_executor.models.utils import WeightsMapper
11 +from vllm.platforms import current_platform
12 +
13 +pytestmark = pytest.mark.skipif(
14 + not current_platform.is_rocm(), reason="ROCm-specific tests"
15 +)
16 +
17 +
18 +def _preshuffled_fp8_linear(
19 + holder: str = "quant_method", weight_shape: tuple[int, int] = (256, 128)
20 +) -> nn.Module:
21 + pytest.importorskip("aiter")
22 + from vllm.model_executor.kernels.linear.scaled_mm.aiter import (
23 + AiterPreshuffledFp8BlockScaledMMKernel,
24 + )
25 + from vllm.model_executor.kernels.linear.scaled_mm.ScaledMMLinearKernel import (
26 + FP8ScaledMMLinearLayerConfig,
27 + )
28 + from vllm.model_executor.layers.quantization.utils.quant_utils import (
29 + kFp8Dynamic128Sym,
30 + kFp8Static128BlockSym,
31 + )
32 +
33 + config = FP8ScaledMMLinearLayerConfig(
34 + weight_quant_key=kFp8Static128BlockSym,
35 + activation_quant_key=kFp8Dynamic128Sym,
36 + input_dtype=torch.bfloat16,
37 + out_dtype=torch.bfloat16,
38 + weight_shape=weight_shape,
39 + )
40 + layer = nn.Module()
41 + layer.weight = nn.Parameter(
42 + torch.randn(*weight_shape, device="cuda").to(current_platform.fp8_dtype()),
43 + requires_grad=False,
44 + )
45 + layer.register_parameter(
46 + "weight_scale" if holder == "scheme" else "weight_scale_inv",
47 + nn.Parameter(
48 + torch.ones(*(dim // 128 for dim in weight_shape), device="cuda"),
49 + requires_grad=False,
50 + ),
51 + )
52 + setattr(
53 + layer,
54 + holder,
55 + SimpleNamespace(fp8_linear=AiterPreshuffledFp8BlockScaledMMKernel(config)),
56 + )
57 + return layer
58 +
59 +
60 +@pytest.mark.parametrize("holder", ["quant_method", "scheme"])
61 +def test_rocm_wo_a_keeps_row_major_weights(holder: str, default_vllm_config) -> None:
62 + from vllm.v1.attention.ops.rocm_aiter_mla_sparse import _get_cached_wo_a_bf16
63 +
64 + layer = _preshuffled_fp8_linear(holder)
65 + layer.is_bmm = True
66 + original = layer.weight.detach().clone()
67 + kernel = getattr(layer, holder).fp8_linear
68 + kernel.process_weights_after_loading(layer)
69 +
70 + actual = _get_cached_wo_a_bf16(layer, 2, 128, 128)
71 + torch.testing.assert_close(actual, original.to(torch.bfloat16).view(2, 128, 128))
72 +
73 +
74 +@pytest.mark.parametrize("holder", ["quant_method", "scheme"])
75 +def test_rocm_gateup_shuffles_once_and_down_proj_keeps_preshuffled_backend(
76 + holder: str, monkeypatch: pytest.MonkeyPatch, default_vllm_config
77 +) -> None:
78 + from vllm._aiter_ops import rocm_aiter_ops
79 + from vllm.models.deepseek_v4.amd import model as rocm_model
80 +
81 + gate_up, down = _preshuffled_fp8_linear(holder), _preshuffled_fp8_linear(holder)
82 + original_gate_up = gate_up.weight.detach().clone()
83 + original_down = down.weight.detach().clone()
84 + monkeypatch.setattr(
85 + rocm_model, "MergedColumnParallelLinear", lambda *a, **k: gate_up
86 + )
87 + monkeypatch.setattr(rocm_model, "RowParallelLinear", lambda *a, **k: down)
88 + monkeypatch.setattr(rocm_aiter_ops, "is_enabled", lambda: True)
89 + model = rocm_model.DeepseekV4MLP(128, 128, "silu")
90 + for linear in (gate_up, down):
91 + getattr(linear, holder).fp8_linear.process_weights_after_loading(linear)
92 + model.prepare_gateup_preshuffle()
93 +
94 + for linear, original in ((gate_up, original_gate_up), (down, original_down)):
95 + expected = rocm_aiter_ops.shuffle_weight(original, layout=(16, 16))
96 + torch.testing.assert_close(
97 + linear.weight.view(torch.uint8), expected.view(torch.uint8)
98 + )
99 +
100 +
101 +def test_rocm_fused_qkv_quant_matches_preshuffled_gemm_scale_layout(
102 + monkeypatch: pytest.MonkeyPatch, default_vllm_config
103 +) -> None:
104 + from vllm._aiter_ops import rocm_aiter_ops
105 + from vllm.models.deepseek_v4.amd.rocm import (
106 + DeepseekV4ROCMAiterMLAAttention,
107 + apply_pre_quantized_block_scaled_mm,
108 + )
109 +
110 + if not rocm_aiter_ops.is_blockscale_bpreshuffle_tuned(1024, 256):
111 + pytest.skip("Requires a tuned AITER preshuffled FP8 GEMM")
112 + linear = _preshuffled_fp8_linear(weight_shape=(1024, 256))
113 + original = linear.weight.detach().clone()
114 + linear.quant_method.fp8_linear.process_weights_after_loading(linear)
115 + attention = DeepseekV4ROCMAiterMLAAttention.__new__(DeepseekV4ROCMAiterMLAAttention)
116 + nn.Module.__init__(attention)
117 + attention.wq_b = linear
118 + attention.indexer = None
119 + attention.q_lora_rank, attention.head_dim = 256, 128
120 + attention.eps = 1e-5
121 + attention.q_norm = SimpleNamespace(
122 + weight=torch.ones(256, dtype=torch.bfloat16, device="cuda")
123 + )
124 + attention.kv_norm = SimpleNamespace(
125 + weight=torch.ones(128, dtype=torch.bfloat16, device="cuda")
126 + )
127 + monkeypatch.setattr(rocm_aiter_ops, "is_linear_fp8_enabled", lambda: True)
128 + # Multiple rows and groups with different ranges expose byte-order errors.
129 + inputs = torch.randn(16, 384, dtype=torch.bfloat16, device="cuda")
130 + inputs[:, :128] *= 4
131 + quantized, scales, _ = attention._split_qkv_and_norm(inputs)
132 + assert scales is not None
133 + output = apply_pre_quantized_block_scaled_mm(linear, quantized, scales)
134 + q = inputs[:, :256].float()
135 + normalized = q * torch.rsqrt(q.square().mean(dim=-1, keepdim=True) + attention.eps)
136 + reference_scales = normalized.view(16, 2, 128).abs().amax(dim=-1)
137 + reference_scales /= torch.finfo(current_platform.fp8_dtype()).max
138 + dequantized = quantized.float() * reference_scales.repeat_interleave(128, dim=1)
139 + expected = dequantized @ original.float().T
140 + torch.testing.assert_close(output.float(), expected, atol=0.125, rtol=0.01)
141 +
142 +
143 +def test_rocm_packed_kv_cache_auto_uses_ds_mla_layout() -> None:
144 + from vllm.config import CacheConfig
145 + from vllm.models.deepseek_v4.attention import _resolve_dsv4_kv_cache_dtype
146 +
147 + cache_config = CacheConfig()
148 +
149 + resolved_dtype, torch_dtype = _resolve_dsv4_kv_cache_dtype(
150 + use_fp8_ds_mla_layout=True,
151 + kv_cache_dtype=cache_config.cache_dtype,
152 + cache_config=cache_config,
153 + )
154 +
155 + assert resolved_dtype == "fp8_ds_mla"
156 + assert torch_dtype is torch.uint8
157 + assert cache_config.cache_dtype == "fp8_ds_mla"
158 +
159 +
160 +def test_rocm_packed_kv_cache_rejects_unquantized_dtype() -> None:
161 + from vllm.config import CacheConfig
162 + from vllm.models.deepseek_v4.attention import _resolve_dsv4_kv_cache_dtype
163 +
164 + cache_config = CacheConfig(cache_dtype="bfloat16")
165 +
166 + with pytest.raises(ValueError, match="only supports fp8 kv-cache"):
167 + _resolve_dsv4_kv_cache_dtype(
168 + use_fp8_ds_mla_layout=True,
169 + kv_cache_dtype=cache_config.cache_dtype,
170 + cache_config=cache_config,
171 + )
172 +
173 +
174 +def test_vl_mapper_preserves_rocm_weight_mapping() -> None:
175 + from vllm.models.deepseek_v4.amd.model import _make_deepseek_v4_weights_mapper
176 + from vllm.models.deepseek_v4.common.vl_model import (
177 + _make_deepseek_v4_vl_weights_mapper,
178 + )
179 +
180 + text_mapper = _make_deepseek_v4_weights_mapper("fp4", fuse_shared_experts=True)
181 + mapper = _make_deepseek_v4_vl_weights_mapper(text_mapper, image_enabled=True)
182 +
183 + assert mapper._map_name("layers.3.attn.wq_a.input_scale") == (
184 + "language_model.model.layers.3.attn.wq_a.input_scale_2"
185 + )
186 + assert mapper._map_name("layers.3.ffn.shared_experts.w2.weight") == (
187 + "language_model.model.layers.3.ffn.shared_experts.w2.weight"
188 + )
189 + assert mapper._map_name("head.weight") == "language_model.lm_head.weight"
190 +
191 +
192 +def test_rocm_moe_wires_vision_routing_on_hash_and_regular_layers(
193 + monkeypatch: pytest.MonkeyPatch,
194 +) -> None:
195 + from vllm.models.deepseek_v4.amd import model as rocm_model
196 + from vllm.models.deepseek_v4.common.mm_preprocess import IMAGE_SENTINEL_BASE_ID
197 +
198 + captured: list[dict] = []
199 +
200 + class FakeGate(nn.Module):
201 + def __init__(self, **kwargs) -> None:
202 + super().__init__()
203 +
204 + def fake_factory(**kwargs):
205 + captured.append(kwargs)
206 + return nn.Identity()
207 +
208 + monkeypatch.setattr(rocm_model, "GateLinear", FakeGate)
209 + monkeypatch.setattr(rocm_model, "FusedMoEFactory", fake_factory)
210 + monkeypatch.setattr(rocm_model, "get_tensor_model_parallel_world_size", lambda: 1)
211 + monkeypatch.setattr(rocm_model, "get_tensor_model_parallel_rank", lambda: 0)
212 +
213 + config = SimpleNamespace(
214 + hidden_size=16,
215 + n_routed_experts=4,
216 + num_experts_per_tok=2,
217 + moe_intermediate_size=8,
218 + swiglu_limit=None,
219 + norm_topk_prob=True,
220 + scoring_func="sqrtsoftplus",
221 + num_hash_layers=1,
222 + vocab_size=32,
223 + topk_method="noaux_tc",
224 + vision_n_layers=1,
225 + n_shared_experts=None,
226 + )
227 + vllm_config = SimpleNamespace(
228 + model_config=SimpleNamespace(hf_config=config), quant_config=None
229 + )
230 +
231 + hash_moe = rocm_model.DeepseekV4MoE(vllm_config, prefix="model.layers.0.ffn")
232 + regular_moe = rocm_model.DeepseekV4MoE(vllm_config, prefix="model.layers.1.ffn")
233 +
234 + assert hash_moe.gate.tid2eid is not None
235 + assert regular_moe.gate.tid2eid is None
236 + for moe, factory_kwargs in zip((hash_moe, regular_moe), captured, strict=True):
237 + assert moe.gate.e_score_correction_bias is not None
238 + assert moe.gate.bias_vl is not None
239 + assert factory_kwargs["bias_vl"] is moe.gate.bias_vl
240 + assert factory_kwargs["image_sentinel_lo"] == IMAGE_SENTINEL_BASE_ID
241 +
242 +
243 +def test_rocm_mtp_forwards_input_ids_for_vision_routing(
244 + monkeypatch: pytest.MonkeyPatch,
245 +) -> None:
246 + from vllm.models.deepseek_v4.amd import mtp as rocm_mtp
247 +
248 + hidden_size = 4
249 + hc_mult = 2
250 +
251 + class FakeNorm(nn.Module):
252 + def __init__(self) -> None:
253 + super().__init__()
254 + self.weight = nn.Parameter(torch.ones(hidden_size))
255 + self.variance_epsilon = 1e-6
256 +
257 + class FakeMTPBlock(nn.Module):
258 + use_fused_mhc = False
259 +
260 + def __init__(self) -> None:
261 + super().__init__()
262 + self.seen_input_ids: torch.Tensor | None = None
263 +
264 + def forward(
265 + self,
266 + *,
267 + positions: torch.Tensor,
268 + x: torch.Tensor,
269 + input_ids: torch.Tensor | None,
270 + ):
271 + self.seen_input_ids = input_ids
272 + return x, None, None, None
273 +
274 + def passthrough_mtp_input(
275 + inputs_embeds: torch.Tensor,
276 + positions: torch.Tensor,
277 + previous_hidden_states: torch.Tensor,
278 + *args,
279 + ) -> tuple[torch.Tensor, torch.Tensor]:
280 + return inputs_embeds, previous_hidden_states
281 +
282 + monkeypatch.setattr(
283 + rocm_mtp, "_FUSED_MTP_INPUT_RMSNORM_KERNEL", passthrough_mtp_input
284 + )
285 +
286 + layer = object.__new__(rocm_mtp.DeepSeekV4MultiTokenPredictorLayer)
287 + nn.Module.__init__(layer)
288 + layer.config = SimpleNamespace(hidden_size=hidden_size)
289 + layer.hc_mult = hc_mult
290 + layer.enorm = FakeNorm()
291 + layer.hnorm = FakeNorm()
292 + layer.e_proj = nn.Identity()
293 + layer.h_proj = nn.Identity()
294 + layer.mtp_block = FakeMTPBlock()
295 +
296 + input_ids = torch.tensor([11, 12])
297 + positions = torch.tensor([3, 4])
298 + inputs_embeds = torch.arange(8, dtype=torch.float32).view(2, hidden_size)
299 + previous_hidden_states = torch.arange(16, dtype=torch.float32).view(2, -1)
300 +
301 + output = layer(
302 + input_ids,
303 + positions,
304 + previous_hidden_states,
305 + inputs_embeds,
306 + )
307 +
308 + assert layer.mtp_block.seen_input_ids is input_ids
309 + expected = previous_hidden_states.view(2, hc_mult, hidden_size)
310 + expected = expected + inputs_embeds.unsqueeze(-2)
311 + torch.testing.assert_close(output, expected.flatten(1))
312 +
313 +
314 +def test_rocm_compute_logits_local_skips_gather() -> None:
315 + from vllm.models.deepseek_v4.amd.model import DeepseekV4ForCausalLM
316 +
317 + calls: list[tuple[nn.Module, torch.Tensor, bool]] = []
318 +
319 + def logits_processor(
320 + lm_head: nn.Module, hidden_states: torch.Tensor, *, skip_gather: bool = False
321 + ) -> torch.Tensor:
322 + calls.append((lm_head, hidden_states, skip_gather))
323 + return hidden_states + 1
324 +
325 + model = object.__new__(DeepseekV4ForCausalLM)
326 + nn.Module.__init__(model)
327 + model.lm_head = nn.Identity()
328 + model.logits_processor = logits_processor
329 + hidden_states = torch.tensor([4.0])
330 +
331 + result = model.compute_logits_local(hidden_states)
332 +
333 + assert torch.equal(result, torch.tensor([5.0]))
334 + assert calls == [(model.lm_head, hidden_states, True)]
335 +
336 +
337 +class _FakeLanguageModel(nn.Module):
338 + finalizes_weights_during_load = False
339 +
340 + def __init__(self) -> None:
341 + super().__init__()
342 + self.tensor_a = nn.Parameter(torch.zeros(1))
343 + self.tensor_c = nn.Parameter(torch.zeros(1))
344 + self.finalized_values: list[tuple[float, float]] = []
345 +
346 + def process_weights_after_loading(self) -> None:
347 + self.finalized_values.append((self.tensor_a.item(), self.tensor_c.item()))
348 +
349 + def compute_logits_local(self, hidden_states: torch.Tensor) -> torch.Tensor:
350 + return hidden_states + 1
351 +
352 +
353 +def test_vl_wrapper_streams_then_delegates_finalization() -> None:
354 + from vllm.models.deepseek_v4.common.vl_model import (
355 + DeepseekV4ForConditionalGeneration,
356 + )
357 +
358 + model = object.__new__(DeepseekV4ForConditionalGeneration)
359 + nn.Module.__init__(model)
360 + model.language_model = _FakeLanguageModel()
361 + model.vision = nn.Module()
362 + model.vision.tensor_b = nn.Parameter(torch.zeros(1))
363 + model.hf_to_vllm_mapper = WeightsMapper()
364 +
365 + def interleaved_weights():
366 + yield "language_model.tensor_a", torch.tensor([1.0])
367 + assert model.language_model.tensor_a.item() == 1.0
368 + yield "vision.tensor_b", torch.tensor([2.0])
369 + assert model.vision.tensor_b.item() == 2.0
370 + yield "language_model.tensor_c", torch.tensor([3.0])
371 +
372 + loaded = model.load_weights(interleaved_weights())
373 +
374 + assert loaded == {
375 + "language_model.tensor_a",
376 + "vision.tensor_b",
377 + "language_model.tensor_c",
378 + }
379 + assert model.language_model.finalized_values == []
380 +
381 + model.process_weights_after_loading()
382 +
383 + assert model.language_model.finalized_values == [(1.0, 3.0)]
384 + assert torch.equal(
385 + model.compute_logits_local(torch.tensor([4.0])), torch.tensor([5.0])
386 + )
387 + model.process_weights_after_loading()
388 + assert model.language_model.finalized_values == [(1.0, 3.0)]
389 +
390 +
391 +class _FakeFinalizingLanguageModel(_FakeLanguageModel):
392 + finalizes_weights_during_load = True
393 +
394 + def __init__(self) -> None:
395 + super().__init__()
396 + self.load_calls = 0
397 +
398 + def load_weights(self, weights) -> set[str]:
399 + self.load_calls += 1
400 + loaded = set()
401 + for name, value in weights:
402 + getattr(self, name).data.copy_(value)
403 + loaded.add(name)
404 + self.process_weights_after_loading()
405 + return loaded
406 +
407 +
408 +def test_vl_wrapper_groups_child_that_finalizes_during_load() -> None:
409 + from vllm.models.deepseek_v4.common.vl_model import (
410 + DeepseekV4ForConditionalGeneration,
411 + )
412 +
413 + model = object.__new__(DeepseekV4ForConditionalGeneration)
414 + nn.Module.__init__(model)
415 + model.language_model = _FakeFinalizingLanguageModel()
416 + model.vision = nn.Module()
417 + model.vision.tensor_b = nn.Parameter(torch.zeros(1))
418 + model.hf_to_vllm_mapper = WeightsMapper()
419 +
420 + loaded = model.load_weights(
421 + iter(
422 + (
423 + ("language_model.tensor_a", torch.tensor([1.0])),
424 + ("vision.tensor_b", torch.tensor([2.0])),
425 + ("language_model.tensor_c", torch.tensor([3.0])),
426 + )
427 + )
428 + )
429 +
430 + assert loaded == {
431 + "language_model.tensor_a",
432 + "vision.tensor_b",
433 + "language_model.tensor_c",
434 + }
435 + assert model.language_model.load_calls == 1
436 + assert model.language_model.finalized_values == [(1.0, 3.0)]
437 +
438 + # The framework's later model-level hook must not double-finalize a child
439 + # which already completed this work in load_weights.
440 + model.process_weights_after_loading()
441 + assert model.language_model.finalized_values == [(1.0, 3.0)]
442 +
443 +
444 +def test_vl_wrapper_dummy_load_delegates_finalization() -> None:
445 + from vllm.models.deepseek_v4.common.vl_model import (
446 + DeepseekV4ForConditionalGeneration,
447 + )
448 +
449 + model = object.__new__(DeepseekV4ForConditionalGeneration)
450 + nn.Module.__init__(model)
451 + model.language_model = _FakeFinalizingLanguageModel()
452 +
453 + # DummyModelLoader bypasses model.load_weights(), so no finalized marker
454 + # exists and the framework-level hook must still delegate to the child.
455 + model.process_weights_after_loading()
456 + assert model.language_model.finalized_values == [(0.0, 0.0)]
457 + model.process_weights_after_loading()
458 + assert model.language_model.finalized_values == [(0.0, 0.0)]
tests/models/test_initialization.py
+2 -2
@@ -135,8 +135,8 @@ def can_initialize(
135 if model_arch == "DeepseekV4ForConditionalGeneration":
136 from vllm.platforms import current_platform
137
138 - if not current_platform.is_cuda():
139 - pytest.skip("Deepseek V4 is only supported on CUDA")
138 + if not (current_platform.is_cuda() or current_platform.is_rocm()):
139 + pytest.skip("Deepseek V4 vision is only supported on CUDA and ROCm")
140
141 with (
142 patch.object(V1EngineCore, "_initialize_kv_caches", _initialize_kv_caches_v1),
tests/models/test_registry.py
+4 -3
@@ -64,10 +64,11 @@ def test_registry_imports(model_arch):
64 ):
65 pytest.skip("HY V4 is only supported on CUDA and ROCm")
66
67 - if model_arch == "DeepseekV4ForConditionalGeneration" and not (
68 - current_platform.is_cuda()
67 + if (
68 + model_arch == "DeepseekV4ForConditionalGeneration"
69 + and not current_platform.is_cuda_alike()
70 ):
70 - pytest.skip("Deepseek V4 is only supported on CUDA")
71 + pytest.skip("Deepseek V4 vision is only supported on CUDA and ROCm")
72
73 # Ensure all model classes can be imported successfully
74 model_cls = ModelRegistry._try_load_model_cls(model_arch)
tests/test_config.py
+39
@@ -57,6 +57,45 @@ def _write_json(path: Path, value: object) -> None:
57 path.write_text(json.dumps(value), encoding="utf-8")
58
59
60 +@pytest.mark.skipif(not current_platform.is_rocm(), reason="ROCm-specific test")
61 +@pytest.mark.parametrize(
62 + ("is_mm_prefix_lm", "is_multimodal_model", "expected"),
63 + [
64 + pytest.param(True, True, True, id="multimodal-prefix-lm"),
65 + pytest.param(False, True, False, id="multimodal-causal"),
66 + pytest.param(True, False, False, id="text-prefix-lm"),
67 + pytest.param(None, True, False, id="missing-model-config"),
68 + ],
69 +)
70 +def test_rocm_mm_prefix_lm_disables_chunked_mm_input(
71 + is_mm_prefix_lm: bool | None,
72 + is_multimodal_model: bool,
73 + expected: bool,
74 +) -> None:
75 + from vllm.platforms.rocm import RocmPlatform
76 +
77 + config = SimpleNamespace(
78 + compilation_config=SimpleNamespace(cudagraph_mode=CUDAGraphMode.NONE),
79 + parallel_config=SimpleNamespace(
80 + prefill_context_parallel_size=1,
81 + worker_cls="test-worker",
82 + ),
83 + model_config=(
84 + None
85 + if is_mm_prefix_lm is None
86 + else SimpleNamespace(is_mm_prefix_lm=is_mm_prefix_lm)
87 + ),
88 + scheduler_config=SimpleNamespace(
89 + is_multimodal_model=is_multimodal_model,
90 + disable_chunked_mm_input=False,
91 + ),
92 + )
93 +
94 + RocmPlatform.check_and_update_config(config)
95 +
96 + assert config.scheduler_config.disable_chunked_mm_input is expected
97 +
98 +
99 def test_kda_recoverssm_derivation_is_revalidated():
100 config = SimpleNamespace(
101 cache_config=SimpleNamespace(
tests/v1/spec_decode/test_adaptive_verification.py
+1
@@ -42,6 +42,7 @@ def make_manager(
42 ("mode", "piecewise_capture_available", "expected"),
43 [
44 ("FULL_DECODE_ONLY", True, "FULL_DECODE_ONLY"),
45 + ("FULL_DECODE_ONLY", False, "FULL_DECODE_ONLY"),
46 ("FULL", True, "FULL_AND_PIECEWISE"),
47 ("FULL", False, "FULL_DECODE_ONLY"),
48 ("FULL_AND_PIECEWISE", True, "FULL_AND_PIECEWISE"),
vllm/models/deepseek_v4/__init__.py
+4 -4
@@ -20,7 +20,7 @@ if current_platform.is_rocm():
20 )
21 from .amd.model import DeepseekV4ForCausalLM
22 from .amd.mtp import DeepSeekV4MTP
23 - from .vl_stub import ( # type: ignore[assignment]
23 + from .common.vl_model import ( # type: ignore[assignment]
24 DeepseekV4ForConditionalGeneration,
25 )
26 elif current_platform.is_xpu():
@@ -38,14 +38,14 @@ elif current_platform.is_cpu():
38 DeepseekV4ForConditionalGeneration,
39 )
40 else:
41 + from .common.vl_model import ( # type: ignore[assignment]
42 + DeepseekV4ForConditionalGeneration,
43 + )
44 from .nvidia.dspark import ( # type: ignore[assignment]
45 DSparkDeepseekV4ForCausalLM,
46 )
47 from .nvidia.model import DeepseekV4ForCausalLM # type: ignore[assignment]
48 from .nvidia.mtp import DeepSeekV4MTP # type: ignore[assignment]
46 - from .nvidia.vl_model import ( # type: ignore[assignment]
47 - DeepseekV4ForConditionalGeneration,
48 - )
49
50 __all__ = [
51 "DSparkDeepseekV4ForCausalLM",
vllm/models/deepseek_v4/amd/model.py
+23 -1
@@ -79,6 +79,8 @@ from vllm.platforms import current_platform
79 from vllm.platforms.rocm import on_gfx950
80 from vllm.sequence import IntermediateTensors
81
82 +from ..common.mm_preprocess import IMAGE_SENTINEL_BASE_ID
83 +
84 logger = init_logger(__name__)
85
86
@@ -549,6 +551,10 @@ class DeepseekV4MoE(nn.Module):
551
552 self.gate.e_score_correction_bias = None
553 self.gate.tid2eid = None
554 + self.gate.bias_vl = None
555 + self.image_sentinel_lo = (
556 + IMAGE_SENTINEL_BASE_ID if getattr(config, "vision_n_layers", 0) > 0 else 0
557 + )
558 is_hash_moe = extract_layer_index(prefix) < config.num_hash_layers
559 self.hash_indices_dtype = torch.int32
560 if is_hash_moe:
@@ -564,12 +570,20 @@ class DeepseekV4MoE(nn.Module):
570 ),
571 requires_grad=False,
572 )
567 - elif getattr(config, "topk_method", None) == "noaux_tc":
573 + if getattr(config, "topk_method", None) == "noaux_tc" and (
574 + not is_hash_moe or getattr(config, "vision_n_layers", 0) > 0
575 + ):
576 self.gate.e_score_correction_bias = nn.Parameter(
577 torch.empty(config.n_routed_experts, dtype=torch.float32),
578 requires_grad=False,
579 )
580
581 + if getattr(config, "vision_n_layers", 0) > 0:
582 + self.gate.bias_vl = nn.Parameter(
583 + torch.empty(config.n_routed_experts, dtype=torch.float32),
584 + requires_grad=False,
585 + )
586 +
587 self.n_shared_experts = config.n_shared_experts
588
589 # TODO: Historically, only `VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS=1`
@@ -638,6 +652,8 @@ class DeepseekV4MoE(nn.Module):
652 routed_scaling_factor=self.routed_scaling_factor,
653 e_score_correction_bias=self.gate.e_score_correction_bias,
654 hash_indices_table=self.gate.tid2eid,
655 + bias_vl=self.gate.bias_vl,
656 + image_sentinel_lo=self.image_sentinel_lo,
657 swiglu_limit=self.swiglu_limit,
658 router_logits_dtype=torch.float32,
659 routed_experts_cls=(
@@ -660,6 +676,8 @@ class DeepseekV4MoE(nn.Module):
676 ) -> torch.Tensor:
677 if self.gate.tid2eid is not None and input_ids is None:
678 raise ValueError("DeepSeek V4 hash MoE routing requires input_ids.")
679 + if self.gate.bias_vl is not None and input_ids is None:
680 + raise ValueError("DeepSeek V4 vision MoE routing requires input_ids.")
681
682 org_shape = hidden_states.shape
683 final_hidden_states = self.experts(
@@ -1351,6 +1369,7 @@ def _make_deepseek_v4_weights_mapper(
1369
1370 class DeepseekV4ForCausalLM(nn.Module, SupportsPP, SupportsEagle3):
1371 model_cls = DeepseekV4Model
1372 + finalizes_weights_during_load = False
1373
1374 # Default mapper assumes the original FP4-expert checkpoint layout.
1375 # Overridden per-instance in __init__ when expert_dtype != "fp4".
@@ -1397,6 +1416,9 @@ class DeepseekV4ForCausalLM(nn.Module, SupportsPP, SupportsEagle3):
1416 logits = self.logits_processor(self.lm_head, hidden_states)
1417 return logits
1418
1419 + def compute_logits_local(self, hidden_states: torch.Tensor) -> torch.Tensor:
1420 + return self.logits_processor(self.lm_head, hidden_states, skip_gather=True)
1421 +
1422 def forward(
1423 self,
1424 input_ids: torch.Tensor,
vllm/models/deepseek_v4/amd/mtp.py
+1 -1
@@ -160,7 +160,7 @@ class DeepSeekV4MultiTokenPredictorLayer(nn.Module):
160 inputs_embeds
161 ).unsqueeze(-2)
162 hidden_states, residual, post_mix, res_mix = self.mtp_block(
163 - positions=positions, x=hidden_states, input_ids=None
163 + positions=positions, x=hidden_states, input_ids=input_ids
164 )
165 if self.mtp_block.use_fused_mhc:
166 hidden_states = self.mtp_block.hc_post(
vllm/models/deepseek_v4/amd/rocm.py
+61 -7
@@ -67,8 +67,13 @@ def weight_already_preshuffled(linear: torch.nn.Module) -> bool:
67 The hand-shuffles below (fused_wqa_wkv, wo_b, gate_up_proj) must be skipped
68 for those, since shuffle_weight is a permutation rather than an involution.
69 """
70 - kernel = getattr(getattr(linear, "quant_method", None), "fp8_linear", None)
71 - return bool(getattr(kernel, "preshuffles_weight", False))
70 + return any(
71 + getattr(getattr(method, "fp8_linear", None), "preshuffles_weight", False)
72 + for method in (
73 + getattr(linear, "quant_method", None),
74 + getattr(linear, "scheme", None),
75 + )
76 + )
77
78
79 def apply_pre_quantized_block_scaled_mm(
@@ -116,13 +121,18 @@ def _combine_topk_swa_indices_kernel(
121 query_start_loc_ptr,
122 seq_lens_ptr,
123 gather_lens_ptr,
124 + left_visible_ptr,
125 + right_visible_ptr,
126 M,
127 N,
128 TOP_K: tl.constexpr,
129 COMPRESS_RATIO: tl.constexpr,
130 WINDOW_SIZE: tl.constexpr,
131 + SWA_WIDTH: tl.constexpr,
132 TOPK_WIDTH: tl.constexpr,
133 PADDED_TOP_K: tl.constexpr,
134 + PADDED_SWA_WIDTH: tl.constexpr,
135 + HAS_IMAGE: tl.constexpr,
136 ):
137 batch_idx = tl.program_id(0)
138 worker_id = tl.program_id(1)
@@ -141,7 +151,20 @@ def _combine_topk_swa_indices_kernel(
151 token_idx_in_query = token_idx - query_start
152 pos = start_pos + token_idx_in_query
153 topk_len = tl.minimum((pos + 1) // COMPRESS_RATIO, TOP_K)
144 - swa_len = tl.minimum(pos + 1, WINDOW_SIZE)
154 + if HAS_IMAGE:
155 + left = tl.load(left_visible_ptr + token_idx)
156 + right = tl.load(right_visible_ptr + token_idx)
157 + else:
158 + left = 0
159 + right = 0
160 + left_add = tl.maximum(left - (WINDOW_SIZE - 1), 0)
161 + # Prefix caching can resume inside an image span. Do not generate
162 + # indices outside the SWA rows present in the gathered workspace.
163 + swa_start = tl.maximum(
164 + tl.maximum(pos - (WINDOW_SIZE - 1) - left_add, 0), gather_start
165 + )
166 + swa_end = tl.minimum(pos + right + 1, seq_len)
167 + swa_len = tl.maximum(swa_end - swa_start, 0)
168
169 topk_offset = tl.arange(0, PADDED_TOP_K)
170 topk_mask = topk_offset < topk_len
@@ -159,14 +182,14 @@ def _combine_topk_swa_indices_kernel(
182 mask=topk_mask,
183 )
184
162 - swa_offset = tl.arange(0, WINDOW_SIZE)
185 + swa_offset = tl.arange(0, PADDED_SWA_WIDTH)
186 tl.store(
187 combined_indices_ptr
188 + token_idx * combined_indices_stride
189 + topk_len
190 + swa_offset,
168 - M * batch_idx + N + swa_offset + pos - swa_len + 1 - gather_start,
169 - mask=swa_offset < swa_len,
191 + M * batch_idx + N + swa_offset + swa_start - gather_start,
192 + mask=(swa_offset < swa_len) & (swa_offset < SWA_WIDTH),
193 )
194
195 tl.store(combined_lens_ptr + token_idx, topk_len + swa_len)
@@ -182,12 +205,21 @@ def combine_topk_swa_indices(
205 topk: int,
206 M: int,
207 N: int,
208 + max_image_tokens: int = 0,
209 + left_visible: torch.Tensor | None = None,
210 + right_visible: torch.Tensor | None = None,
211 ) -> tuple[torch.Tensor, torch.Tensor]:
212 + if (left_visible is None) != (right_visible is None):
213 + raise ValueError("left_visible and right_visible must be provided together")
214 topk_indices = topk_indices.reshape(topk_indices.shape[0], -1).contiguous()
215 num_tokens = topk_indices.shape[0]
216 num_reqs = seq_lens.shape[0]
217 + has_image = left_visible is not None
218 + # Keep the row shape fixed for a vision model even when a particular batch
219 + # has no image.
220 + swa_width = window_size + max_image_tokens
221 combined_topk = (
190 - (topk + window_size + _SPARSE_PREFILL_TOPK_ALIGNMENT - 1)
222 + (topk + swa_width + _SPARSE_PREFILL_TOPK_ALIGNMENT - 1)
223 // _SPARSE_PREFILL_TOPK_ALIGNMENT
224 * _SPARSE_PREFILL_TOPK_ALIGNMENT
225 )
@@ -211,13 +243,18 @@ def combine_topk_swa_indices(
243 query_start_loc,
244 seq_lens,
245 gather_lens,
246 + left_visible if left_visible is not None else topk_indices,
247 + right_visible if right_visible is not None else topk_indices,
248 M,
249 N,
250 TOP_K=topk,
251 COMPRESS_RATIO=compress_ratio,
252 WINDOW_SIZE=window_size,
253 + SWA_WIDTH=swa_width,
254 TOPK_WIDTH=topk_indices.shape[-1],
255 PADDED_TOP_K=triton.next_power_of_2(topk_indices.shape[-1]),
256 + PADDED_SWA_WIDTH=triton.next_power_of_2(swa_width),
257 + HAS_IMAGE=has_image,
258 )
259 return combined_indices, combined_lens
260
@@ -948,6 +985,12 @@ class DeepseekV4ROCMAiterMLAAttention(DeepseekV4Attention):
985 assert query_start_loc_cpu is not None
986 assert query_start_loc is not None
987 prefill_token_base = query_start_loc_cpu[num_decodes]
988 + left_visible = swa_metadata.prefill_left_visible
989 + right_visible = swa_metadata.prefill_right_visible
990 + if left_visible is not None:
991 + left_visible = left_visible[num_decode_tokens:]
992 + assert right_visible is not None
993 + right_visible = right_visible[num_decode_tokens:]
994
995 if not swa_only:
996 if self.compress_ratio == 4:
@@ -1026,6 +1069,17 @@ class DeepseekV4ROCMAiterMLAAttention(DeepseekV4Attention):
1069 top_k,
1070 M,
1071 N,
1072 + max_image_tokens=self.max_image_tokens,
1073 + left_visible=(
1074 + left_visible[query_start:query_end]
1075 + if left_visible is not None
1076 + else None
1077 + ),
1078 + right_visible=(
1079 + right_visible[query_start:query_end]
1080 + if right_visible is not None
1081 + else None
1082 + ),
1083 )
1084 rocm_sparse_attn_prefill(
1085 q=q[query_start:query_end],
vllm/models/deepseek_v4/common/vl_model.py new
+331
@@ -0,0 +1,331 @@
1 +# SPDX-License-Identifier: Apache-2.0
2 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3 +"""DeepSeek-V4 vision variant (e.g. DeepSeek-V4-Flash-Vision-Exp).
4 +
5 +Thin multimodal wrapper around the text-only ``DeepseekV4ForCausalLM``:
6 +
7 +- ``vision`` ViT + ``aligner`` produce per-image embeddings for the IMAGE
8 + sentinel positions; four learned vectors (``image_start`` / ``image_pad`` /
9 + ``image_newline`` / ``image_end``) fill the remaining sentinel positions.
10 +- Image placeholders (``<｜deepseek_image｜>``) are expanded by the processor
11 + in ``common/mm_preprocess.py`` into sentinel blocks borrowing reserved
12 + in-vocab tokens ``<|place_holder_mm_span_0431|>``..``_0435|>``
13 + (see ``common/vision.py`` for the tower itself).
14 +- Merged embeddings enter the text model via ``inputs_embeds``, i.e. before
15 + its hyper-connection stream expansion. Raw ``input_ids`` still flow into
16 + the model so the MoE router can apply ``bias_vl`` to image tokens
17 + (``requires_raw_input_tokens``).
18 +"""
19 +
20 +from collections.abc import Iterable
21 +
22 +import torch
23 +from torch import nn
24 +
25 +from vllm.model_executor.models.interfaces import (
26 + MultiModalEmbeddings,
27 + SupportsEagle3,
28 + SupportsMultiModal,
29 + SupportsPP,
30 +)
31 +from vllm.model_executor.models.utils import (
32 + AutoWeightsLoader,
33 + WeightsMapper,
34 + init_vllm_registered_model,
35 + maybe_prefix,
36 +)
37 +from vllm.multimodal import MULTIMODAL_REGISTRY
38 +
39 +from .mm_preprocess import (
40 + IMAGE_PLACEHOLDER,
41 + IMAGE_SENTINEL_BASE_ID,
42 + DeepseekV4VLDummyInputsBuilder,
43 + DeepseekV4VLMultiModalProcessor,
44 + DeepseekV4VLProcessingInfo,
45 + image_sentinel_mask,
46 +)
47 +from .vision import DeepseekV4Aligner, DeepseekV4ViT
48 +
49 +
50 +def _make_deepseek_v4_vl_weights_mapper(
51 + text_mapper: WeightsMapper, image_enabled: bool
52 +) -> WeightsMapper:
53 + """Text-checkpoint mapping rules re-rooted under ``language_model.``."""
54 + orig_to_new_prefix = {
55 + src: None if dst is None else f"language_model.{dst}"
56 + for src, dst in text_mapper.orig_to_new_prefix.items()
57 + }
58 + if not image_enabled:
59 + orig_to_new_prefix.update({"vision.": None, "aligner.": None, "image_": None})
60 + return WeightsMapper(
61 + orig_to_new_renaming=text_mapper.orig_to_new_renaming,
62 + orig_to_new_prefix=orig_to_new_prefix,
63 + orig_to_new_regex=text_mapper.orig_to_new_regex,
64 + orig_to_new_stacked=text_mapper.orig_to_new_stacked,
65 + orig_to_new_suffix={
66 + **text_mapper.orig_to_new_suffix,
67 + "head.weight": "language_model.lm_head.weight",
68 + },
69 + orig_to_new_substr={
70 + **text_mapper.orig_to_new_substr,
71 + # Draft models load the checkpoint's MTP weights separately from
72 + # the target model.
73 + "mtp.": None,
74 + },
75 + )
76 +
77 +
78 +@MULTIMODAL_REGISTRY.register_processor(
79 + DeepseekV4VLMultiModalProcessor,
80 + info=DeepseekV4VLProcessingInfo,
81 + dummy_inputs=DeepseekV4VLDummyInputsBuilder,
82 +)
83 +class DeepseekV4ForConditionalGeneration(
84 + nn.Module, SupportsMultiModal, SupportsPP, SupportsEagle3
85 +):
86 + """Multimodal entry point for DeepSeek-V4 checkpoints with a vision tower.
87 +
88 + ``SupportsEagle3`` (aux hidden-state plumbing for MTP/DSpark drafters)
89 + delegates through ``language_model`` via the protocol defaults.
90 + """
91 +
92 + # The MoE router needs raw token ids to detect image sentinel tokens
93 + # (borrowed reserved ids, see common/mm_preprocess.py) and apply bias_vl.
94 + requires_raw_input_tokens = True
95 +
96 + @classmethod
97 + def get_placeholder_str(cls, modality: str, i: int) -> str | None:
98 + if modality == "image":
99 + return IMAGE_PLACEHOLDER
100 + raise ValueError(f"Unsupported modality: {modality!r}")
101 +
102 + def __init__(self, *, vllm_config, prefix: str = "") -> None:
103 + super().__init__()
104 + model_config = vllm_config.model_config
105 + config = model_config.hf_config
106 + self.config = config
107 + self.multimodal_config = model_config.multimodal_config
108 + assert self.multimodal_config is not None
109 +
110 + image_enabled = (
111 + config.vision_n_layers > 0
112 + and self.multimodal_config.get_limit_per_prompt("image") > 0
113 + )
114 + with self._mark_tower_model(vllm_config, {"image"}):
115 + self.vision: DeepseekV4ViT | None = None
116 + self.aligner: DeepseekV4Aligner | None = None
117 + self.image_start: nn.Parameter | None = None
118 + self.image_end: nn.Parameter | None = None
119 + self.image_newline: nn.Parameter | None = None
120 + self.image_pad: nn.Parameter | None = None
121 + if image_enabled:
122 + self.vision = DeepseekV4ViT(config)
123 + self.aligner = DeepseekV4Aligner(config)
124 + for name in (
125 + "image_start",
126 + "image_end",
127 + "image_newline",
128 + "image_pad",
129 + ):
130 + setattr(
131 + self,
132 + name,
133 + nn.Parameter(
134 + torch.empty(config.hidden_size, dtype=torch.float32)
135 + ),
136 + )
137 + self.vision.to(dtype=model_config.dtype)
138 + self.aligner.to(dtype=model_config.dtype)
139 +
140 + with self._mark_language_model(vllm_config):
141 + # The arch convertor routes any config with a vision tower to
142 + # this wrapper class; mark the copy handed to the text backbone
143 + # so it resolves to DeepseekV4ForCausalLM instead of recursing
144 + # (with_hf_config deepcopies the config, the marker survives).
145 + config._dsv4_vl_inner = True # type: ignore[attr-defined]
146 + try:
147 + self.language_model = init_vllm_registered_model(
148 + vllm_config=vllm_config,
149 + hf_config=config,
150 + prefix=maybe_prefix(prefix, "language_model"),
151 + architectures=["DeepseekV4ForCausalLM"],
152 + )
153 + finally:
154 + del config._dsv4_vl_inner # type: ignore[attr-defined]
155 + # The outer mapper (see load_weights) fully resolves HF names into
156 + # this wrapper's namespace before AutoWeightsLoader strips the
157 + # "language_model." prefix and delegates to the child's load_weights,
158 + # so the child's own mapper must be a no-op. Its suffix rules are not
159 + # idempotent (e.g. "lm_head.weight".endswith("head.weight") would
160 + # re-fire "head.weight" -> "lm_head.weight").
161 + text_mapper = self.language_model.hf_to_vllm_mapper
162 + self.language_model.hf_to_vllm_mapper = WeightsMapper()
163 + self.make_empty_intermediate_tensors = ( # type: ignore[method-assign]
164 + self.language_model.make_empty_intermediate_tensors
165 + )
166 +
167 + self.hf_to_vllm_mapper = _make_deepseek_v4_vl_weights_mapper(
168 + text_mapper, image_enabled
169 + )
170 + self._weights_finalized = False
171 +
172 + def _parse_and_validate_image_input(self, **kwargs: object) -> dict | None:
173 + patches = kwargs.pop("patches", None)
174 + if patches is None:
175 + return None
176 + vit_grid = kwargs.pop("vit_grid", None)
177 + llm_grid = kwargs.pop("llm_grid", None)
178 + perm = kwargs.pop("perm", None)
179 + assert vit_grid is not None and llm_grid is not None and perm is not None
180 + return {
181 + "patches": patches,
182 + "vit_grid": vit_grid,
183 + "llm_grid": llm_grid,
184 + "perm": perm,
185 + }
186 +
187 + def _encode_image(
188 + self,
189 + patches: torch.Tensor,
190 + n_vit_h: int,
191 + n_vit_w: int,
192 + perm: torch.Tensor,
193 + ) -> torch.Tensor:
194 + assert self.vision is not None and self.aligner is not None
195 + image_embeds = self.aligner(
196 + self.vision(patches, n_vit_h, n_vit_w), n_vit_h, n_vit_w
197 + )
198 + # Reorder into the N-layout block order used in the prompt.
199 + return image_embeds[perm.to(image_embeds.device)]
200 +
201 + def _process_image_input(
202 + self,
203 + patches: torch.Tensor,
204 + vit_grid: torch.Tensor,
205 + llm_grid: torch.Tensor,
206 + perm: torch.Tensor,
207 + ) -> tuple[torch.Tensor, ...]:
208 + assert self.vision is not None and self.aligner is not None
209 + patches = patches.to(self.aligner.w1.weight.dtype)
210 +
211 + embeds: list[torch.Tensor] = []
212 + vit_offset = 0
213 + llm_offset = 0
214 + for (n_vit_h, n_vit_w), (n_llm_h, n_llm_w) in zip(
215 + vit_grid.tolist(), llm_grid.tolist(), strict=True
216 + ):
217 + n_vit = n_vit_h * n_vit_w
218 + n_llm = n_llm_h * n_llm_w
219 + embeds.append(
220 + self._encode_image(
221 + patches[vit_offset : vit_offset + n_vit],
222 + n_vit_h,
223 + n_vit_w,
224 + perm[llm_offset : llm_offset + n_llm],
225 + )
226 + )
227 + vit_offset += n_vit
228 + llm_offset += n_llm
229 + return tuple(embeds)
230 +
231 + def embed_multimodal(self, **kwargs: object) -> MultiModalEmbeddings:
232 + image_input = self._parse_and_validate_image_input(**kwargs)
233 + if image_input is None or self.vision is None:
234 + return []
235 + return self._process_image_input(
236 + image_input["patches"],
237 + image_input["vit_grid"],
238 + image_input["llm_grid"],
239 + image_input["perm"],
240 + )
241 +
242 + def embed_input_ids(
243 + self,
244 + input_ids: torch.Tensor,
245 + multimodal_embeddings: MultiModalEmbeddings | None = None,
246 + *,
247 + is_multimodal: torch.Tensor | None = None,
248 + ) -> torch.Tensor:
249 + from vllm.model_executor.models.utils import _merge_multimodal_embeddings
250 +
251 + # All ids are in-vocab here: image-block sentinels are borrowed
252 + # reserved tokens (their embedding rows are always overwritten below).
253 + inputs_embeds = self.language_model.embed_input_ids(input_ids)
254 +
255 + if self.image_start is not None:
256 + # Branch-free sentinel overwrite: safe inside compiled/captured
257 + # regions (no data-dependent control flow).
258 + sentinel_mask = image_sentinel_mask(input_ids)
259 + if is_multimodal is not None:
260 + # IMAGE positions get vision embeddings via the merge below.
261 + sentinel_mask = sentinel_mask & ~is_multimodal.to(input_ids.device)
262 + table = torch.stack(
263 + [
264 + self.image_start,
265 + self.image_pad,
266 + self.image_pad,
267 + self.image_newline,
268 + self.image_end,
269 + ]
270 + ).to(inputs_embeds.dtype)
271 + idx = (input_ids - IMAGE_SENTINEL_BASE_ID).clamp(0, 4)
272 + inputs_embeds = torch.where(
273 + sentinel_mask.unsqueeze(-1), table[idx], inputs_embeds
274 + )
275 +
276 + if multimodal_embeddings is None or len(multimodal_embeddings) == 0:
277 + return inputs_embeds
278 +
279 + assert is_multimodal is not None
280 + return _merge_multimodal_embeddings(
281 + inputs_embeds=inputs_embeds,
282 + multimodal_embeddings=multimodal_embeddings,
283 + is_multimodal=is_multimodal,
284 + )
285 +
286 + def forward(
287 + self,
288 + input_ids: torch.Tensor,
289 + positions: torch.Tensor,
290 + intermediate_tensors=None,
291 + inputs_embeds: torch.Tensor | None = None,
292 + **kwargs,
293 + ) -> torch.Tensor:
294 + return self.language_model(
295 + input_ids, positions, intermediate_tensors, inputs_embeds
296 + )
297 +
298 + def compute_logits(self, hidden_states: torch.Tensor) -> torch.Tensor | None:
299 + return self.language_model.compute_logits(hidden_states)
300 +
301 + def compute_logits_local(self, hidden_states: torch.Tensor) -> torch.Tensor:
302 + return self.language_model.compute_logits_local(hidden_states)
303 +
304 + def get_expert_mapping(self) -> list[tuple[str, str, int, str]]:
305 + return self.language_model.get_expert_mapping()
306 +
307 + def get_mtp_target_hidden_states(self) -> torch.Tensor | None:
308 + """Pre-hc_head residual stream buffer for the MTP/DSpark draft model."""
309 + return self.language_model.get_mtp_target_hidden_states()
310 +
311 + def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
312 + child_finalizes = getattr(
313 + self.language_model, "finalizes_weights_during_load", True
314 + )
315 + mapped = self.hf_to_vllm_mapper.apply(weights)
316 + if child_finalizes:
317 + # A child which finalizes inside load_weights must see all of its
318 + # weights in one contiguous delegation from AutoWeightsLoader.
319 + mapped = iter(sorted(mapped, key=lambda x: x[0]))
320 + loader = AutoWeightsLoader(self)
321 + loaded_params = loader.load_weights(mapped)
322 + self._weights_finalized = child_finalizes
323 + return loaded_params
324 +
325 + def process_weights_after_loading(self) -> None:
326 + # Backbones such as the ROCm implementation require this to run only
327 + # after the loader's generic per-layer quantization finalization.
328 + if getattr(self, "_weights_finalized", False):
329 + return
330 + self.language_model.process_weights_after_loading()
331 + self._weights_finalized = True
vllm/models/deepseek_v4/nvidia/model.py
+1
@@ -1790,6 +1790,7 @@ class DeepseekV4ForCausalLM(
1790 SupportsLoRA,
1791 DeepseekV4MixtureOfExperts,
1792 ):
1793 + finalizes_weights_during_load = True
1794 model_cls = DeepseekV4Model
1795
1796 # Default mapper assumes the original FP4-expert checkpoint layout.
vllm/models/deepseek_v4/nvidia/vl_model.py
+10 -322
@@ -1,333 +1,21 @@
1 # SPDX-License-Identifier: Apache-2.0
2 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3 -"""DeepSeek-V4 vision variant (e.g. DeepSeek-V4-Flash-Vision-Exp).
3 +"""Compatibility imports for the platform-neutral DeepSeek-V4 vision model."""
4
5 -Thin multimodal wrapper around the text-only ``DeepseekV4ForCausalLM``:
6 -
7 -- ``vision`` ViT + ``aligner`` produce per-image embeddings for the IMAGE
8 - sentinel positions; four learned vectors (``image_start`` / ``image_pad`` /
9 - ``image_newline`` / ``image_end``) fill the remaining sentinel positions.
10 -- Image placeholders (``<｜deepseek_image｜>``) are expanded by the processor
11 - in ``common/mm_preprocess.py`` into sentinel blocks borrowing reserved
12 - in-vocab tokens ``<|place_holder_mm_span_0431|>``..``_0435|>``
13 - (see ``common/vision.py`` for the tower itself).
14 -- Merged embeddings enter the text model via ``inputs_embeds``, i.e. before
15 - its hyper-connection stream expansion. Raw ``input_ids`` still flow into
16 - the model so the MoE router can apply ``bias_vl`` to image tokens
17 - (``requires_raw_input_tokens``).
18 -"""
19 -
20 -from collections.abc import Iterable
21 -
22 -import torch
23 -from torch import nn
24 -
25 -from vllm.model_executor.models.interfaces import (
26 - MultiModalEmbeddings,
27 - SupportsEagle3,
28 - SupportsMultiModal,
29 - SupportsPP,
5 +from ..common.vl_model import (
6 + DeepseekV4ForConditionalGeneration,
7 )
31 -from vllm.model_executor.models.utils import (
32 - AutoWeightsLoader,
33 - WeightsMapper,
34 - init_vllm_registered_model,
35 - maybe_prefix,
8 +from ..common.vl_model import (
9 + _make_deepseek_v4_vl_weights_mapper as _make_common_vl_weights_mapper,
10 )
37 -from vllm.multimodal import MULTIMODAL_REGISTRY
38 -
39 -from ..common.mm_preprocess import (
40 - IMAGE_PLACEHOLDER,
41 - IMAGE_SENTINEL_BASE_ID,
42 - DeepseekV4VLDummyInputsBuilder,
43 - DeepseekV4VLMultiModalProcessor,
44 - DeepseekV4VLProcessingInfo,
45 - image_sentinel_mask,
46 -)
47 -from ..common.vision import DeepseekV4Aligner, DeepseekV4ViT
11 from .model import _make_deepseek_v4_weights_mapper
12
13
51 -def _make_deepseek_v4_vl_weights_mapper(
52 - expert_dtype: str, image_enabled: bool
53 -) -> WeightsMapper:
54 - """Text-checkpoint mapping rules re-rooted under ``language_model.``."""
55 - base = _make_deepseek_v4_weights_mapper(expert_dtype)
56 - orig_to_new_prefix: dict[str, str | None] = {
57 - "layers.": "language_model.model.layers.",
58 - "embed.": "language_model.model.embed.",
59 - "norm.": "language_model.model.norm.",
60 - "hc_head": "language_model.model.hc_head",
61 - "mtp.": "language_model.model.mtp.",
62 - }
63 - if not image_enabled:
64 - orig_to_new_prefix.update({"vision.": None, "aligner.": None, "image_": None})
65 - return WeightsMapper(
66 - orig_to_new_prefix=orig_to_new_prefix,
67 - orig_to_new_regex=base.orig_to_new_regex,
68 - orig_to_new_suffix={
69 - "head.weight": "language_model.lm_head.weight",
70 - "embed.weight": "embed_tokens.weight",
71 - ".ffn.gate.bias": ".ffn.gate.e_score_correction_bias",
72 - },
73 - orig_to_new_substr={
74 - ".shared_experts.w2": ".shared_experts.down_proj",
75 - # The MTP/DSpark draft heads are not supported for the vision
76 - # variant; drop their weights.
77 - "mtp.": None,
78 - },
14 +def _make_deepseek_v4_vl_weights_mapper(expert_dtype: str, image_enabled: bool):
15 + """Retain the original NVIDIA helper signature for downstream imports."""
16 + return _make_common_vl_weights_mapper(
17 + _make_deepseek_v4_weights_mapper(expert_dtype), image_enabled
18 )
19
20
82 -@MULTIMODAL_REGISTRY.register_processor(
83 - DeepseekV4VLMultiModalProcessor,
84 - info=DeepseekV4VLProcessingInfo,
85 - dummy_inputs=DeepseekV4VLDummyInputsBuilder,
86 -)
87 -class DeepseekV4ForConditionalGeneration(
88 - nn.Module, SupportsMultiModal, SupportsPP, SupportsEagle3
89 -):
90 - """Multimodal entry point for DeepSeek-V4 checkpoints with a vision tower.
91 -
92 - ``SupportsEagle3`` (aux hidden-state plumbing for MTP/DSpark drafters)
93 - delegates through ``language_model`` via the protocol defaults.
94 - """
95 -
96 - # The MoE router needs raw token ids to detect image sentinel tokens
97 - # (borrowed reserved ids, see common/mm_preprocess.py) and apply bias_vl.
98 - requires_raw_input_tokens = True
99 -
100 - @classmethod
101 - def get_placeholder_str(cls, modality: str, i: int) -> str | None:
102 - if modality == "image":
103 - return IMAGE_PLACEHOLDER
104 - raise ValueError(f"Unsupported modality: {modality!r}")
105 -
106 - def __init__(self, *, vllm_config, prefix: str = "") -> None:
107 - super().__init__()
108 - model_config = vllm_config.model_config
109 - config = model_config.hf_config
110 - self.config = config
111 - self.multimodal_config = model_config.multimodal_config
112 - assert self.multimodal_config is not None
113 -
114 - image_enabled = (
115 - config.vision_n_layers > 0
116 - and self.multimodal_config.get_limit_per_prompt("image") > 0
117 - )
118 - with self._mark_tower_model(vllm_config, {"image"}):
119 - self.vision: DeepseekV4ViT | None = None
120 - self.aligner: DeepseekV4Aligner | None = None
121 - self.image_start: nn.Parameter | None = None
122 - self.image_end: nn.Parameter | None = None
123 - self.image_newline: nn.Parameter | None = None
124 - self.image_pad: nn.Parameter | None = None
125 - if image_enabled:
126 - self.vision = DeepseekV4ViT(config)
127 - self.aligner = DeepseekV4Aligner(config)
128 - for name in (
129 - "image_start",
130 - "image_end",
131 - "image_newline",
132 - "image_pad",
133 - ):
134 - setattr(
135 - self,
136 - name,
137 - nn.Parameter(
138 - torch.empty(config.hidden_size, dtype=torch.float32)
139 - ),
140 - )
141 - self.vision.to(dtype=model_config.dtype)
142 - self.aligner.to(dtype=model_config.dtype)
143 -
144 - with self._mark_language_model(vllm_config):
145 - # The arch convertor routes any config with a vision tower to
146 - # this wrapper class; mark the copy handed to the text backbone
147 - # so it resolves to DeepseekV4ForCausalLM instead of recursing
148 - # (with_hf_config deepcopies the config, the marker survives).
149 - config._dsv4_vl_inner = True # type: ignore[attr-defined]
150 - try:
151 - self.language_model = init_vllm_registered_model(
152 - vllm_config=vllm_config,
153 - hf_config=config,
154 - prefix=maybe_prefix(prefix, "language_model"),
155 - architectures=["DeepseekV4ForCausalLM"],
156 - )
157 - finally:
158 - del config._dsv4_vl_inner # type: ignore[attr-defined]
159 - # The outer mapper (see load_weights) fully resolves HF names into
160 - # this wrapper's namespace before AutoWeightsLoader strips the
161 - # "language_model." prefix and delegates to the child's load_weights,
162 - # so the child's own mapper must be a no-op. Its suffix rules are not
163 - # idempotent (e.g. "lm_head.weight".endswith("head.weight") would
164 - # re-fire "head.weight" -> "lm_head.weight").
165 - self.language_model.hf_to_vllm_mapper = WeightsMapper()
166 - self.make_empty_intermediate_tensors = ( # type: ignore[method-assign]
167 - self.language_model.make_empty_intermediate_tensors
168 - )
169 -
170 - expert_dtype = getattr(config, "expert_dtype", "fp4")
171 - self.hf_to_vllm_mapper = _make_deepseek_v4_vl_weights_mapper(
172 - expert_dtype, image_enabled
173 - )
174 -
175 - def _parse_and_validate_image_input(self, **kwargs: object) -> dict | None:
176 - patches = kwargs.pop("patches", None)
177 - if patches is None:
178 - return None
179 - vit_grid = kwargs.pop("vit_grid", None)
180 - llm_grid = kwargs.pop("llm_grid", None)
181 - perm = kwargs.pop("perm", None)
182 - assert vit_grid is not None and llm_grid is not None and perm is not None
183 - return {
184 - "patches": patches,
185 - "vit_grid": vit_grid,
186 - "llm_grid": llm_grid,
187 - "perm": perm,
188 - }
189 -
190 - def _encode_image(
191 - self,
192 - patches: torch.Tensor,
193 - n_vit_h: int,
194 - n_vit_w: int,
195 - perm: torch.Tensor,
196 - ) -> torch.Tensor:
197 - assert self.vision is not None and self.aligner is not None
198 - image_embeds = self.aligner(
199 - self.vision(patches, n_vit_h, n_vit_w), n_vit_h, n_vit_w
200 - )
201 - # Reorder into the N-layout block order used in the prompt.
202 - return image_embeds[perm.to(image_embeds.device)]
203 -
204 - def _process_image_input(
205 - self,
206 - patches: torch.Tensor,
207 - vit_grid: torch.Tensor,
208 - llm_grid: torch.Tensor,
209 - perm: torch.Tensor,
210 - ) -> tuple[torch.Tensor, ...]:
211 - assert self.vision is not None and self.aligner is not None
212 - patches = patches.to(self.aligner.w1.weight.dtype)
213 -
214 - embeds: list[torch.Tensor] = []
215 - vit_offset = 0
216 - llm_offset = 0
217 - for (n_vit_h, n_vit_w), (n_llm_h, n_llm_w) in zip(
218 - vit_grid.tolist(), llm_grid.tolist(), strict=True
219 - ):
220 - n_vit = n_vit_h * n_vit_w
221 - n_llm = n_llm_h * n_llm_w
222 - embeds.append(
223 - self._encode_image(
224 - patches[vit_offset : vit_offset + n_vit],
225 - n_vit_h,
226 - n_vit_w,
227 - perm[llm_offset : llm_offset + n_llm],
228 - )
229 - )
230 - vit_offset += n_vit
231 - llm_offset += n_llm
232 - return tuple(embeds)
233 -
234 - def embed_multimodal(self, **kwargs: object) -> MultiModalEmbeddings:
235 - image_input = self._parse_and_validate_image_input(**kwargs)
236 - if image_input is None or self.vision is None:
237 - return []
238 - return self._process_image_input(
239 - image_input["patches"],
240 - image_input["vit_grid"],
241 - image_input["llm_grid"],
242 - image_input["perm"],
243 - )
244 -
245 - def embed_input_ids(
246 - self,
247 - input_ids: torch.Tensor,
248 - multimodal_embeddings: MultiModalEmbeddings | None = None,
249 - *,
250 - is_multimodal: torch.Tensor | None = None,
251 - ) -> torch.Tensor:
252 - from vllm.model_executor.models.utils import _merge_multimodal_embeddings
253 -
254 - # All ids are in-vocab here: image-block sentinels are borrowed
255 - # reserved tokens (their embedding rows are always overwritten below).
256 - inputs_embeds = self.language_model.embed_input_ids(input_ids)
257 -
258 - if self.image_start is not None:
259 - # Branch-free sentinel overwrite: safe inside compiled/captured
260 - # regions (no data-dependent control flow).
261 - sentinel_mask = image_sentinel_mask(input_ids)
262 - if is_multimodal is not None:
263 - # IMAGE positions get vision embeddings via the merge below.
264 - sentinel_mask = sentinel_mask & ~is_multimodal.to(input_ids.device)
265 - table = torch.stack(
266 - [
267 - self.image_start,
268 - self.image_pad,
269 - self.image_pad,
270 - self.image_newline,
271 - self.image_end,
272 - ]
273 - ).to(inputs_embeds.dtype)
274 - idx = (input_ids - IMAGE_SENTINEL_BASE_ID).clamp(0, 4)
275 - inputs_embeds = torch.where(
276 - sentinel_mask.unsqueeze(-1), table[idx], inputs_embeds
277 - )
278 -
279 - if multimodal_embeddings is None or len(multimodal_embeddings) == 0:
280 - return inputs_embeds
281 -
282 - assert is_multimodal is not None
283 - return _merge_multimodal_embeddings(
284 - inputs_embeds=inputs_embeds,
285 - multimodal_embeddings=multimodal_embeddings,
286 - is_multimodal=is_multimodal,
287 - )
288 -
289 - def forward(
290 - self,
291 - input_ids: torch.Tensor,
292 - positions: torch.Tensor,
293 - intermediate_tensors=None,
294 - inputs_embeds: torch.Tensor | None = None,
295 - **kwargs,
296 - ) -> torch.Tensor:
297 - return self.language_model(
298 - input_ids, positions, intermediate_tensors, inputs_embeds
299 - )
300 -
301 - def compute_logits(self, hidden_states: torch.Tensor) -> torch.Tensor | None:
302 - return self.language_model.compute_logits(hidden_states)
303 -
304 - def compute_logits_local(self, hidden_states: torch.Tensor) -> torch.Tensor:
305 - return self.language_model.compute_logits_local(hidden_states)
306 -
307 - def get_expert_mapping(self) -> list[tuple[str, str, int, str]]:
308 - return self.language_model.get_expert_mapping()
309 -
310 - def get_mtp_target_hidden_states(self) -> torch.Tensor | None:
311 - """Pre-hc_head residual stream buffer for the MTP/DSpark draft model."""
312 - return self.language_model.get_mtp_target_hidden_states()
313 -
314 - def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
315 - # Map HF names into this wrapper's namespace up front and sort, so
316 - # the "language_model." group reaches the child loader as one
317 - # contiguous block (AutoWeightsLoader delegates per contiguous group,
318 - # and the child's load_weights finalizes fused expert weights, which
319 - # must not run on a partially loaded model).
320 - mapped = sorted(self.hf_to_vllm_mapper.apply(weights), key=lambda x: x[0])
321 - loader = AutoWeightsLoader(self)
322 - loaded_params = loader.load_weights(mapped)
323 - # The child's load_weights already ran its post-load finalization.
324 - self._weights_finalized = True
325 - return loaded_params
326 -
327 - def process_weights_after_loading(self) -> None:
328 - # Model-level post-load hook (called by the loader after any load
329 - # format). Under DummyModelLoader the child's load_weights — and
330 - # hence its finalize step — is bypassed, so run it here instead.
331 - if getattr(self, "_weights_finalized", False):
332 - return
333 - self.language_model.process_weights_after_loading()
21 +__all__ = ["DeepseekV4ForConditionalGeneration"]
vllm/models/deepseek_v4/vl_stub.py
+2 -2
@@ -9,6 +9,6 @@ class DeepseekV4ForConditionalGeneration(nn.Module):
9 def __init__(self, *, vllm_config, prefix: str = ""):
10 super().__init__()
11 raise NotImplementedError(
12 - "DeepSeek-V4 vision (DeepseekV4ForConditionalGeneration) is only "
13 - "supported on NVIDIA GPUs for now."
12 + "DeepSeek-V4 vision (DeepseekV4ForConditionalGeneration) is not "
13 + "supported on the current platform; use CUDA or ROCm."
14 )
vllm/platforms/rocm.py
+15
@@ -943,6 +943,21 @@ class RocmPlatform(Platform):
943 if parallel_config.worker_cls == "auto":
944 parallel_config.worker_cls = "vllm.v1.worker.gpu_worker.Worker"
945
946 + model_config = vllm_config.model_config
947 + scheduler_config = vllm_config.scheduler_config
948 + # Note: model_config may be None during testing
949 + if (
950 + model_config is not None
951 + and model_config.is_mm_prefix_lm
952 + and scheduler_config.is_multimodal_model
953 + and not scheduler_config.disable_chunked_mm_input
954 + ):
955 + logger.warning_once(
956 + "Forcing --disable_chunked_mm_input for models "
957 + "with multimodal-bidirectional attention."
958 + )
959 + scheduler_config.disable_chunked_mm_input = True
960 +
961 @classmethod
962 def verify_model_arch(cls, model_arch: str) -> None:
963 if model_arch in _ROCM_UNSUPPORTED_MODELS: