qwen3_vl.py · tduncan/Nvidia-Optimization-Stack-VLLM

@tduncan / Nvidia-Optimization-Stack-VLLM

Code Commits Issues Pulls Cloud Agent CI/CD

main

py 3,170 lines 120 KB

Raw

1	# SPDX-License-Identifier: Apache-2.0
2	# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3
4	# Copyright 2025 The vLLM team.
5	# Copyright 2025 The Qwen Team.
6	# Copyright 2025 The HuggingFace Inc. team.
7	# All rights reserved.
8	#
9	# This code is based on EleutherAI's GPT-NeoX library and the GPT-NeoX
10	# and OPT implementations in this library. It has been modified from its
11	# original forms to accommodate minor architectural differences compared
12	# to GPT-NeoX and OPT used by the Meta AI team that trained the model.
13	#
14	# Licensed under the Apache License, Version 2.0 (the "License");
15	# you may not use this file except in compliance with the License.
16	# You may obtain a copy of the License at
17	#
18	# http://www.apache.org/licenses/LICENSE-2.0
19	#
20	# Unless required by applicable law or agreed to in writing, software
21	# distributed under the License is distributed on an "AS IS" BASIS,
22	# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
23	# See the License for the specific language governing permissions and
24	# limitations under the License.
25	"""Inference-only Qwen3VL model compatible with HuggingFace weights."""
26
27	from collections.abc import Callable, Hashable, Iterable, Iterator, Mapping, Sequence
28	from functools import lru_cache, partial
29	from itertools import islice
30	from typing import Any, Literal
31
32	import numpy as np
33	import torch
34	import torch.nn as nn
35	import torch.nn.functional as F
36	from transformers import BatchFeature, ProcessorMixin
37	from transformers.models.qwen2_vl import Qwen2VLImageProcessor
38	from transformers.models.qwen2_vl.image_processing_qwen2_vl import (
39	smart_resize as image_smart_resize,
40	)
41	from transformers.models.qwen3_vl import Qwen3VLProcessor, Qwen3VLVideoProcessor
42	from transformers.models.qwen3_vl.configuration_qwen3_vl import (
43	Qwen3VLConfig,
44	Qwen3VLVisionConfig,
45	)
46	from transformers.models.qwen3_vl.video_processing_qwen3_vl import (
47	smart_resize as video_smart_resize,
48	)
49	from transformers.video_utils import VideoMetadata
50
51	from vllm.compilation.decorators import (
52	should_torch_compile_mm_encoder,
53	support_torch_compile,
54	)
55	from vllm.config import VllmConfig
56	from vllm.config.multimodal import (
57	MultiModalConfig,
58	MultiModalDummyOptions,
59	VideoDummyOptions,
60	VideoPruningMethod,
61	)
62	from vllm.distributed import get_pp_group, parallel_state
63	from vllm.inputs import MultiModalDataDict
64	from vllm.logger import init_logger
65	from vllm.model_executor.layers.activation import _ACTIVATION_REGISTRY
66	from vllm.model_executor.layers.attention.mm_encoder_attention import (
67	MMEncoderAttention,
68	)
69	from vllm.model_executor.layers.conv import Conv3dLayer
70	from vllm.model_executor.layers.fusion.mm_input_norm import (
71	IdentityInputNorm,
72	build_mm_input_norm,
73	)
74	from vllm.model_executor.layers.linear import (
75	ColumnParallelLinear,
76	RowParallelLinear,
77	)
78	from vllm.model_executor.layers.logits_processor import LogitsProcessor
79	from vllm.model_executor.layers.quantization import QuantizationConfig
80	from vllm.model_executor.layers.rotary_embedding import get_rope
81	from vllm.model_executor.layers.vocab_parallel_embedding import ParallelLMHead
82	from vllm.model_executor.models.module_mapping import MultiModelKeys
83	from vllm.multimodal import MULTIMODAL_REGISTRY
84	from vllm.multimodal.inputs import (
85	MultiModalFeatureSpec,
86	MultiModalFieldConfig,
87	MultiModalFieldElem,
88	MultiModalKwargsItem,
89	MultiModalKwargsItems,
90	PlaceholderRange,
91	VideoItem,
92	)
93	from vllm.multimodal.parse import ImageSize, MultiModalDataItems
94	from vllm.multimodal.processing import (
95	BaseDummyInputsBuilder,
96	BaseMultiModalProcessor,
97	PromptReplacement,
98	PromptUpdate,
99	PromptUpdateDetails,
100	)
101	from vllm.multimodal.video_prune.evs import (
102	compute_mrope_for_media,
103	compute_retained_tokens_count,
104	compute_retention_mask,
105	recompute_mrope_positions,
106	)
107	from vllm.multimodal.video_prune.vidcom2 import (
108	compute_retained_tokens_count as vidcom2_compute_retained_tokens_count,
109	)
110	from vllm.multimodal.video_prune.vidcom2 import (
111	compute_retention_mask as vidcom2_compute_retention_mask,
112	)
113	from vllm.sequence import IntermediateTensors
114	from vllm.tokenizers.protocol import TokenizerLike
115	from vllm.tokenizers.registry import cached_tokenizer_from_config
116	from vllm.triton_utils import HAS_TRITON, tl, triton
117	from vllm.utils.cache import LRUCache
118	from vllm.utils.collection_utils import is_list_of
119	from vllm.utils.math_utils import cdiv, round_up
120	from vllm.utils.torch_utils import PIN_MEMORY
121	from vllm.v1.worker.encoder_cudagraph_defs import EncoderCudaGraphReplayBuffers
122
123	from ...utils.gpu_sync_debug import gpu_sync_allowed
124	from ...utils.torch_utils import async_tensor_h2d
125	from .interfaces import (
126	MultiModalEmbeddings,
127	SupportsEagle,
128	SupportsEagle3,
129	SupportsEncoderCudaGraph,
130	SupportsLoRA,
131	SupportsMRoPE,
132	SupportsMultiModal,
133	SupportsMultiModalPruning,
134	SupportsPP,
135	_require_is_multimodal,
136	)
137	from .qwen2_5_vl import (
138	Qwen2_5_VisionAttention,
139	Qwen2_5_VLImageEmbeddingInputs,
140	Qwen2_5_VLImageInputs,
141	Qwen2_5_VLImagePixelInputs,
142	Qwen2_5_VLVideoEmbeddingInputs,
143	Qwen2_5_VLVideoInputs,
144	Qwen2_5_VLVideoPixelInputs,
145	)
146	from .qwen2_vl import (
147	Qwen2VLMultiModalDataParser,
148	Qwen2VLProcessingInfo,
149	_create_qwen2vl_field_factory,
150	)
151	from .qwen3 import Qwen3ForCausalLM, Qwen3Model
152	from .utils import (
153	AutoWeightsLoader,
154	PPMissingLayer,
155	StageMissingLayer,
156	WeightsMapper,
157	_merge_multimodal_embeddings,
158	maybe_prefix,
159	)
160	from .vision import (
161	get_fp8_padded_hidden_size,
162	get_vit_attn_backend,
163	is_vit_use_data_parallel,
164	run_dp_sharded_mrope_vision_model,
165	)
166
167	logger = init_logger(__name__)
168
169	# We use 2048 dummy video frames that would generate vision embeddings
170	# of the maximum size.
171	DUMMY_VIDEO_NUM_FRAMES = 2048
172
173	# ---------------------------------------------------------------------------
174	# Triton kernel: fused bilinear position-embedding interpolation
175	# ---------------------------------------------------------------------------
176	# Replaces many small eager-mode CUDA kernels with a single launch.
177	# The spatial-merge reorder is baked into the index math so the output
178	# is ready to be added to the patch embeddings directly.
179	# ---------------------------------------------------------------------------
180
181	if HAS_TRITON:
182
183	@triton.jit
184	def _bilinear_pos_embed_kernel(
185	embed_ptr,
186	output_ptr,
187	H,
188	W,
189	h_scale,
190	w_scale,
191	NUM_GRID: tl.constexpr,
192	M_SIZE: tl.constexpr,
193	HIDDEN_DIM: tl.constexpr,
194	BLOCK_D: tl.constexpr,
195	):
196	"""Fused bilinear pos-embed interpolation with spatial-merge reorder."""
197	pid = tl.program_id(0)
198	total_spatial = H * W
199	spatial_idx = pid % total_spatial
200
201	num_blocks_w = W // M_SIZE
202	block_idx = spatial_idx // (M_SIZE * M_SIZE)
203	local_idx = spatial_idx % (M_SIZE * M_SIZE)
204	br = block_idx // num_blocks_w
205	bc = block_idx % num_blocks_w
206	lr = local_idx // M_SIZE
207	lc = local_idx % M_SIZE
208	row = br * M_SIZE + lr
209	col = bc * M_SIZE + lc
210
211	h_frac = row.to(tl.float32) * h_scale
212	w_frac = col.to(tl.float32) * w_scale
213
214	hf = tl.math.floor(h_frac).to(tl.int32)
215	wf = tl.math.floor(w_frac).to(tl.int32)
216	hc = tl.minimum(hf + 1, NUM_GRID - 1)
217	wc = tl.minimum(wf + 1, NUM_GRID - 1)
218
219	dh = h_frac - hf.to(tl.float32)
220	dw = w_frac - wf.to(tl.float32)
221	w11 = dh * dw
222	w10 = dh - w11
223	w01 = dw - w11
224	w00 = 1.0 - dh - w01
225
226	off00 = (hf * NUM_GRID + wf) * HIDDEN_DIM
227	off01 = (hf * NUM_GRID + wc) * HIDDEN_DIM
228	off10 = (hc * NUM_GRID + wf) * HIDDEN_DIM
229	off11 = (hc * NUM_GRID + wc) * HIDDEN_DIM
230	out_off = pid * HIDDEN_DIM
231
232	# Cast weights to output dtype so the multiply-accumulate stays
233	# in the same precision as the native PyTorch implementation.
234	out_dtype = output_ptr.dtype.element_ty
235	w00_c = w00.to(out_dtype)
236	w01_c = w01.to(out_dtype)
237	w10_c = w10.to(out_dtype)
238	w11_c = w11.to(out_dtype)
239
240	for d in tl.range(0, HIDDEN_DIM, BLOCK_D):
241	cols = d + tl.arange(0, BLOCK_D)
242	mask = cols < HIDDEN_DIM
243
244	e00 = tl.load(embed_ptr + off00 + cols, mask=mask)
245	e01 = tl.load(embed_ptr + off01 + cols, mask=mask)
246	e10 = tl.load(embed_ptr + off10 + cols, mask=mask)
247	e11 = tl.load(embed_ptr + off11 + cols, mask=mask)
248
249	val = w00_c * e00 + w01_c * e01 + w10_c * e10 + w11_c * e11
250
251	tl.store(output_ptr + out_off + cols, val, mask=mask)
252
253	def triton_pos_embed_interpolate(
254	embed_weight: torch.Tensor,
255	t: int,
256	h: int,
257	w: int,
258	num_grid_per_side: int,
259	m_size: int,
260	dtype: torch.dtype,
261	) -> torch.Tensor:
262	"""Launch the fused Triton kernel for one (t,h,w) grid.
263
264	Returns a tensor of shape ``(t * h * w, hidden_dim)`` with the
265	bilinearly-interpolated position embeddings in spatial-merge order.
266	"""
267	assert h % m_size == 0 and w % m_size == 0, (
268	f"h={h} and w={w} must be divisible by m_size={m_size}"
269	)
270	hidden_dim = embed_weight.shape[1]
271	total_out = t * h * w
272	output = torch.empty(
273	total_out,
274	hidden_dim,
275	device=embed_weight.device,
276	dtype=dtype,
277	)
278
279	h_scale = float(num_grid_per_side - 1) / float(h - 1) if h > 1 else 0.0
280	w_scale = float(num_grid_per_side - 1) / float(w - 1) if w > 1 else 0.0
281
282	BLOCK_D = triton.next_power_of_2(hidden_dim)
283
284	_bilinear_pos_embed_kernel[(total_out,)](
285	embed_weight,
286	output,
287	h,
288	w,
289	h_scale,
290	w_scale,
291	num_grid_per_side,
292	m_size,
293	hidden_dim,
294	BLOCK_D,
295	)
296	return output
297
298
299	def pos_embed_interpolate_native(
300	embed_weight: torch.Tensor,
301	t: int,
302	h: int,
303	w: int,
304	num_grid_per_side: int,
305	m_size: int,
306	dtype: torch.dtype,
307	) -> torch.Tensor:
308	"""Eager PyTorch bilinear position-embedding interpolation.
309
310	Returns a tensor of shape ``(t * h * w, hidden_dim)`` with the
311	bilinearly-interpolated position embeddings in spatial-merge order.
312	"""
313	assert h % m_size == 0 and w % m_size == 0, (
314	f"h={h} and w={w} must be divisible by m_size={m_size}"
315	)
316	hidden_dim = embed_weight.shape[1]
317	device = embed_weight.device
318
319	h_idxs = torch.linspace(
320	0,
321	num_grid_per_side - 1,
322	h,
323	dtype=torch.float32,
324	device=device,
325	)
326	w_idxs = torch.linspace(
327	0,
328	num_grid_per_side - 1,
329	w,
330	dtype=torch.float32,
331	device=device,
332	)
333
334	h_floor = h_idxs.to(torch.long)
335	w_floor = w_idxs.to(torch.long)
336	h_ceil = torch.clamp(h_floor + 1, max=num_grid_per_side - 1)
337	w_ceil = torch.clamp(w_floor + 1, max=num_grid_per_side - 1)
338
339	dh = h_idxs - h_floor
340	dw = w_idxs - w_floor
341
342	dh_grid, dw_grid = torch.meshgrid(dh, dw, indexing="ij")
343	h_floor_grid, w_floor_grid = torch.meshgrid(h_floor, w_floor, indexing="ij")
344	h_ceil_grid, w_ceil_grid = torch.meshgrid(h_ceil, w_ceil, indexing="ij")
345
346	w11 = dh_grid * dw_grid
347	w10 = dh_grid - w11
348	w01 = dw_grid - w11
349	w00 = 1 - dh_grid - w01
350
351	h_grid = torch.stack([h_floor_grid, h_floor_grid, h_ceil_grid, h_ceil_grid])
352	w_grid = torch.stack([w_floor_grid, w_ceil_grid, w_floor_grid, w_ceil_grid])
353	h_grid_idx = h_grid * num_grid_per_side
354
355	indices = (h_grid_idx + w_grid).reshape(4, -1)
356	weights = torch.stack([w00, w01, w10, w11], dim=0).reshape(4, -1, 1)
357	weights = weights.to(dtype=dtype)
358
359	embeds = embed_weight[indices]
360	embeds *= weights
361	combined = embeds.sum(dim=0)
362
363	combined = combined.reshape(h // m_size, m_size, w // m_size, m_size, hidden_dim)
364	combined = combined.permute(0, 2, 1, 3, 4).reshape(1, -1, hidden_dim)
365	repeated = combined.expand(t, -1, -1).reshape(-1, hidden_dim)
366	return repeated.to(dtype=dtype)
367
368
369	@support_torch_compile(
370	dynamic_arg_dims={
371	"x": 0,
372	},
373	enable_if=should_torch_compile_mm_encoder,
374	is_encoder=True,
375	)
376	class Qwen3_VisionPatchEmbed(nn.Module):
377	def __init__(
378	self,
379	patch_size: int = 14,
380	temporal_patch_size: int = 2,
381	in_channels: int = 3,
382	hidden_size: int = 1152,
383	) -> None:
384	super().__init__()
385	self.patch_size = patch_size
386	self.temporal_patch_size = temporal_patch_size
387	self.hidden_size = hidden_size
388
389	kernel_size = (temporal_patch_size, patch_size, patch_size)
390	self.proj = Conv3dLayer(
391	in_channels,
392	hidden_size,
393	kernel_size=kernel_size,
394	stride=kernel_size,
395	bias=True,
396	)
397
398	def forward(self, x: torch.Tensor) -> torch.Tensor:
399	L, C = x.shape
400	x = x.view(L, -1, self.temporal_patch_size, self.patch_size, self.patch_size)
401	x = self.proj(x).view(L, self.hidden_size)
402	return x
403
404
405	class Qwen3_VisionMLP(nn.Module):
406	def __init__(
407	self,
408	in_features: int,
409	hidden_features: int,
410	bias: bool = False,
411	act_fn: Callable[[torch.Tensor], torch.Tensor] = F.silu,
412	quant_config: QuantizationConfig \| None = None,
413	prefix: str = "",
414	):
415	super().__init__()
416	use_data_parallel = is_vit_use_data_parallel()
417	self.linear_fc1 = ColumnParallelLinear(
418	in_features,
419	hidden_features,
420	bias=bias,
421	quant_config=quant_config,
422	return_bias=False,
423	prefix=f"{prefix}.linear_fc1",
424	disable_tp=use_data_parallel,
425	)
426	self.linear_fc2 = RowParallelLinear(
427	hidden_features,
428	in_features,
429	bias=bias,
430	quant_config=quant_config,
431	return_bias=False,
432	prefix=f"{prefix}.linear_fc2",
433	disable_tp=use_data_parallel,
434	)
435	self.act_fn = act_fn
436
437	def forward(self, x: torch.Tensor):
438	mlp_output = self.linear_fc2(self.act_fn(self.linear_fc1(x)))
439	return mlp_output
440
441
442	@support_torch_compile(
443	dynamic_arg_dims={
444	"x": 0,
445	"cu_seqlens": 0,
446	"rotary_pos_emb_cos": 0,
447	"rotary_pos_emb_sin": 0,
448	"sequence_lengths": 0,
449	},
450	enable_if=should_torch_compile_mm_encoder,
451	is_encoder=True,
452	)
453	class Qwen3_VisionBlock(nn.Module):
454	def __init__(
455	self,
456	dim: int,
457	num_heads: int,
458	mlp_hidden_dim: int,
459	act_fn: Callable[[torch.Tensor], torch.Tensor] = F.silu,
460	norm_layer: Callable[[int], nn.Module] \| None = None,
461	quant_config: QuantizationConfig \| None = None,
462	prefix: str = "",
463	) -> None:
464	super().__init__()
465	if norm_layer is None:
466	norm_layer = partial(nn.LayerNorm, eps=1e-6)
467	self.norm1 = norm_layer(dim)
468	self.norm2 = norm_layer(dim)
469	self.attn = Qwen2_5_VisionAttention(
470	embed_dim=dim,
471	num_heads=num_heads,
472	projection_size=dim,
473	quant_config=quant_config,
474	prefix=f"{prefix}.attn",
475	)
476	self.mlp = Qwen3_VisionMLP(
477	dim,
478	mlp_hidden_dim,
479	act_fn=act_fn,
480	bias=True,
481	quant_config=quant_config,
482	prefix=f"{prefix}.mlp",
483	)
484
485	def forward(
486	self,
487	x: torch.Tensor,
488	cu_seqlens: torch.Tensor,
489	rotary_pos_emb_cos: torch.Tensor,
490	rotary_pos_emb_sin: torch.Tensor,
491	max_seqlen: torch.Tensor, # Only used for Flash Attention
492	sequence_lengths: torch.Tensor, # Only used for FlashInfer CuDNN backend
493	) -> torch.Tensor:
494	x = x + self.attn(
495	self.norm1(x),
496	cu_seqlens=cu_seqlens,
497	rotary_pos_emb_cos=rotary_pos_emb_cos,
498	rotary_pos_emb_sin=rotary_pos_emb_sin,
499	max_seqlen=max_seqlen,
500	sequence_lengths=sequence_lengths,
501	)
502
503	x = x + self.mlp(self.norm2(x))
504	return x
505
506
507	class Qwen3_VisionPatchMerger(nn.Module):
508	def __init__(
509	self,
510	d_model: int,
511	context_dim: int,
512	norm_layer: Callable[[int], nn.Module] \| None = None,
513	spatial_merge_size: int = 2,
514	use_postshuffle_norm: bool = False,
515	quant_config: QuantizationConfig \| None = None,
516	prefix: str = "",
517	) -> None:
518	super().__init__()
519	use_data_parallel = is_vit_use_data_parallel()
520	self.hidden_size = context_dim * (spatial_merge_size**2)
521
522	self.use_postshuffle_norm = use_postshuffle_norm
523	if self.use_postshuffle_norm:
524	context_dim = self.hidden_size
525
526	if norm_layer is None:
527	norm_layer = partial(nn.LayerNorm, eps=1e-6)
528	self.norm = norm_layer(context_dim)
529	self.linear_fc1 = ColumnParallelLinear(
530	self.hidden_size,
531	self.hidden_size,
532	bias=True,
533	quant_config=quant_config,
534	prefix=f"{prefix}.linear_fc1",
535	disable_tp=use_data_parallel,
536	)
537	self.act_fn = nn.GELU()
538	self.linear_fc2 = RowParallelLinear(
539	self.hidden_size,
540	d_model,
541	bias=True,
542	quant_config=quant_config,
543	prefix=f"{prefix}.linear_fc2",
544	disable_tp=use_data_parallel,
545	)
546
547	def forward(self, x: torch.Tensor) -> torch.Tensor:
548	if self.use_postshuffle_norm:
549	x = self.norm(x.view(-1, self.hidden_size))
550	else:
551	x = self.norm(x).view(-1, self.hidden_size)
552
553	x_parallel, _ = self.linear_fc1(x)
554	x_parallel = self.act_fn(x_parallel)
555	out, _ = self.linear_fc2(x_parallel)
556	return out
557
558
559	class Qwen3_VisionTransformer(nn.Module):
560	hf_to_vllm_mapper = WeightsMapper(
561	orig_to_new_stacked={
562	"attn.q.": ("attn.qkv.", "q"),
563	"attn.k.": ("attn.qkv.", "k"),
564	"attn.v.": ("attn.qkv.", "v"),
565	}
566	)
567
568	def __init__(
569	self,
570	vision_config: Qwen3VLVisionConfig,
571	norm_eps: float = 1e-6,
572	quant_config: QuantizationConfig \| None = None,
573	input_norm: nn.Module \| None = None,
574	prefix: str = "",
575	) -> None:
576	super().__init__()
577	self.hidden_size = vision_config.hidden_size
578	self.num_heads = vision_config.num_heads
579	self.num_position_embeddings = vision_config.num_position_embeddings
580	self.patch_size = vision_config.patch_size
581	self.spatial_merge_size = vision_config.spatial_merge_size
582	self.spatial_merge_unit = self.spatial_merge_size**2
583	self.temporal_patch_size = vision_config.temporal_patch_size
584	self.deepstack_visual_indexes = (
585	vision_config.deepstack_visual_indexes
586	if hasattr(vision_config, "deepstack_visual_indexes")
587	else []
588	)
589	self.num_grid_per_side = int(self.num_position_embeddings**0.5)
590	self._rot_pos_ids_cache: LRUCache[tuple[int, int, int], torch.Tensor] = (
591	LRUCache(capacity=1024)
592	)
593
594	use_data_parallel = is_vit_use_data_parallel()
595	self.tp_size = (
596	1
597	if use_data_parallel
598	else parallel_state.get_tensor_model_parallel_world_size()
599	)
600
601	# NOTE: This is used for creating empty tensor for all_gather for
602	# DP ViT. Here out_hidden_size is enlarged due to deepstack
603	self.out_hidden_size = vision_config.out_hidden_size * (
604	1 + len(self.deepstack_visual_indexes)
605	)
606
607	self.patch_embed = Qwen3_VisionPatchEmbed(
608	patch_size=self.patch_size,
609	temporal_patch_size=self.temporal_patch_size,
610	in_channels=vision_config.in_channels,
611	hidden_size=self.hidden_size,
612	)
613	self.input_norm = input_norm if input_norm is not None else IdentityInputNorm()
614
615	self.pos_embed = nn.Embedding(self.num_position_embeddings, self.hidden_size)
616
617	norm_layer = partial(nn.LayerNorm, eps=norm_eps)
618	head_dim = self.hidden_size // self.num_heads
619
620	# FP8 attention: Q/K/V become independent contiguous tensors
621	# after quantization, so cu_seqlens uses uniform stride (no 3x V).
622	self.fp8_padded_hidden_size = get_fp8_padded_hidden_size(
623	self.num_heads, head_dim
624	)
625
626	self.rotary_pos_emb = get_rope(
627	head_size=head_dim,
628	max_position=8192,
629	is_neox_style=True,
630	rope_parameters={"partial_rotary_factor": 0.5},
631	)
632
633	self.merger = Qwen3_VisionPatchMerger(
634	d_model=vision_config.out_hidden_size,
635	context_dim=self.hidden_size,
636	norm_layer=norm_layer,
637	spatial_merge_size=self.spatial_merge_size,
638	quant_config=quant_config,
639	prefix=f"{prefix}.merger",
640	)
641
642	self.deepstack_merger_list = nn.ModuleList(
643	[
644	Qwen3_VisionPatchMerger(
645	d_model=vision_config.out_hidden_size,
646	context_dim=self.hidden_size,
647	spatial_merge_size=self.spatial_merge_size,
648	use_postshuffle_norm=True,
649	norm_layer=norm_layer,
650	quant_config=quant_config,
651	prefix=f"{prefix}.deepstack_merger_list.{layer_idx}",
652	)
653	for layer_idx in range(len(self.deepstack_visual_indexes))
654	]
655	)
656
657	self.attn_backend = get_vit_attn_backend(
658	head_size=head_dim,
659	dtype=torch.get_default_dtype(),
660	)
661
662	self.blocks = nn.ModuleList(
663	[
664	Qwen3_VisionBlock(
665	dim=self.hidden_size,
666	num_heads=self.num_heads,
667	mlp_hidden_dim=vision_config.intermediate_size,
668	act_fn=_ACTIVATION_REGISTRY[vision_config.hidden_act],
669	norm_layer=norm_layer,
670	quant_config=quant_config,
671	prefix=f"{prefix}.blocks.{layer_idx}",
672	)
673	for layer_idx in range(vision_config.depth)
674	]
675	)
676
677	@property
678	def dtype(self) -> torch.dtype:
679	return self.patch_embed.proj.weight.dtype
680
681	@property
682	def device(self) -> torch.device:
683	return self.patch_embed.proj.weight.device
684
685	def rot_pos_ids(self, h: int, w: int, spatial_merge_size: int) -> torch.Tensor:
686	cache_key = (h, w, spatial_merge_size)
687	if cache_key in self._rot_pos_ids_cache:
688	return self._rot_pos_ids_cache[cache_key]
689
690	hpos_ids = np.broadcast_to(np.arange(h).reshape(h, 1), (h, w))
691	h_div = h // spatial_merge_size
692	w_div = w // spatial_merge_size
693	hpos_ids = hpos_ids.reshape(
694	h_div,
695	spatial_merge_size,
696	w_div,
697	spatial_merge_size,
698	)
699	hpos_ids = hpos_ids.transpose(0, 2, 1, 3)
700	hpos_ids = hpos_ids.flatten()
701
702	wpos_ids = np.broadcast_to(np.arange(w).reshape(1, w), (h, w))
703	wpos_ids = wpos_ids.reshape(
704	h_div,
705	spatial_merge_size,
706	w_div,
707	spatial_merge_size,
708	)
709	wpos_ids = wpos_ids.transpose(0, 2, 1, 3)
710	wpos_ids = wpos_ids.flatten()
711
712	result = torch.from_numpy(np.stack([hpos_ids, wpos_ids], axis=-1))
713	self._rot_pos_ids_cache[cache_key] = result
714	return result
715
716	def rot_pos_emb(self, grid_thw: Sequence[Sequence[int]]):
717	max_grid_size = max(max(h, w) for _, h, w in grid_thw)
718	pos_ids = [
719	self.rot_pos_ids(h, w, self.spatial_merge_size)
720	if t == 1
721	else self.rot_pos_ids(h, w, self.spatial_merge_size).repeat(t, 1)
722	for t, h, w in grid_thw
723	]
724	num_pos = sum(p.shape[0] for p in pos_ids)
725	pinned = torch.empty(
726	(num_pos, pos_ids[0].shape[1]),
727	dtype=pos_ids[0].dtype,
728	pin_memory=PIN_MEMORY,
729	)
730	pos_ids = torch.cat(pos_ids, dim=0, out=pinned).to(
731	self.device, non_blocking=True
732	)
733
734	# Use pre-computed cos_sin_cache from RotaryEmbedding
735	cos, sin = self.rotary_pos_emb.get_cos_sin(max_grid_size)
736
737	cos_combined = cos[pos_ids].flatten(1)
738	sin_combined = sin[pos_ids].flatten(1)
739
740	return cos_combined, sin_combined
741
742	def fast_pos_embed_interpolate(
743	self, grid_thw: Sequence[Sequence[int]]
744	) -> torch.Tensor:
745	interpolate_fn = (
746	triton_pos_embed_interpolate if HAS_TRITON else pos_embed_interpolate_native
747	)
748	outputs = []
749	for t, h, w in grid_thw:
750	outputs.append(
751	interpolate_fn(
752	self.pos_embed.weight,
753	t,
754	h,
755	w,
756	self.num_grid_per_side,
757	self.spatial_merge_size,
758	self.dtype,
759	)
760	)
761	return torch.cat(outputs, dim=0)
762
763	def prepare_encoder_metadata(
764	self,
765	grid_thw_list: Sequence[Sequence[int]],
766	*,
767	max_batch_size: int \| None = None,
768	max_frames_per_batch: int \| None = None,
769	max_seqlen_override: int \| None = None,
770	device: torch.device \| None = None,
771	) -> dict[str, torch.Tensor \| None]:
772	"""Compute encoder metadata from grid_thw_list.
773
774	Shared by the eager forward path, CUDA graph capture, and
775	CUDA graph replay to avoid duplicated implementation.
776
777	Args:
778	grid_thw_list: Grid configurations as list of [t, h, w].
779	max_batch_size: If set, pad cu_seqlens to this size
780	(needed for CUDA graph capture/replay).
781	max_frames_per_batch: If set, overrides max_batch_size for
782	cu_seqlens padding. For video inputs each item contributes
783	T attention sequences (frames); this sizes the buffer to
784	the total frame budget so video replays never overflow.
785	max_seqlen_override: If set, use this value for max_seqlen
786	instead of computing from cu_seqlens (needed for CUDA
787	graph capture to cover worst-case replay scenarios).
788	device: Device to place tensors on. Defaults to self.device.
789
790	"""
791	if device is None:
792	device = self.device
793
794	metadata: dict[str, torch.Tensor \| None] = {}
795
796	# Positional embeddings
797	metadata["pos_embeds"] = self.fast_pos_embed_interpolate(grid_thw_list)
798	rotary_cos, rotary_sin = self.rot_pos_emb(grid_thw_list)
799	metadata["rotary_pos_emb_cos"] = rotary_cos
800	metadata["rotary_pos_emb_sin"] = rotary_sin
801
802	# cu_seqlens from grid_thw
803	grid_thw_np = np.array(grid_thw_list, dtype=np.int32)
804	patches_per_frame = grid_thw_np[:, 1] * grid_thw_np[:, 2]
805	cu_seqlens = np.repeat(patches_per_frame, grid_thw_np[:, 0]).cumsum(
806	dtype=np.int32
807	)
808	cu_seqlens = np.concatenate([np.zeros(1, dtype=np.int32), cu_seqlens])
809
810	# Pad cu_seqlens to the required number of sequences.
811	# For videos each item contributes T frames = T attention sequences,
812	# so the total can exceed max_batch_size. max_frames_per_batch
813	# overrides the pad target when set.
814	pad_to = (
815	max_frames_per_batch if max_frames_per_batch is not None else max_batch_size
816	)
817	if pad_to is not None:
818	num_seqs = len(cu_seqlens) - 1
819	if num_seqs < pad_to:
820	cu_seqlens = np.concatenate(
821	[
822	cu_seqlens,
823	np.full(
824	pad_to - num_seqs,
825	cu_seqlens[-1],
826	dtype=np.int32,
827	),
828	]
829	)
830
831	# sequence_lengths (backend-specific)
832	metadata["sequence_lengths"] = MMEncoderAttention.maybe_compute_seq_lens(
833	self.attn_backend, cu_seqlens, device
834	)
835
836	# max_seqlen
837	if max_seqlen_override is not None:
838	max_seqlen_val = max_seqlen_override
839	else:
840	max_seqlen_val = MMEncoderAttention.compute_max_seqlen(
841	self.attn_backend, cu_seqlens
842	)
843	# Keep max_seqlen on CPU: attention wrappers call .item() on it,
844	# and having it on GPU would capture a wasteful D2H copy in CUDA
845	# graphs without changing behavior (the scalar is baked at capture).
846	metadata["max_seqlen"] = torch.tensor(max_seqlen_val, dtype=torch.int32)
847
848	# Recompute cu_seqlens (backend-specific transformation)
849	metadata["cu_seqlens"] = MMEncoderAttention.maybe_recompute_cu_seqlens(
850	self.attn_backend,
851	cu_seqlens,
852	self.hidden_size,
853	self.tp_size,
854	device,
855	fp8_padded_hidden_size=self.fp8_padded_hidden_size,
856	)
857
858	return metadata
859
860	def forward(
861	self,
862	x: torch.Tensor,
863	grid_thw: torch.Tensor \| list[list[int]],
864	*,
865	encoder_metadata: dict[str, torch.Tensor] \| None = None,
866	) -> torch.Tensor:
867	hidden_states = self.input_norm(
868	x.to(device=self.device, non_blocking=True), self.dtype
869	)
870	hidden_states = self.patch_embed(hidden_states)
871
872	if encoder_metadata is None:
873	if isinstance(grid_thw, list):
874	grid_thw_list = grid_thw
875	else:
876	grid_thw_list = grid_thw.tolist()
877	encoder_metadata = self.prepare_encoder_metadata(grid_thw_list)
878
879	pos_embeds = encoder_metadata["pos_embeds"]
880	hidden_states = hidden_states + pos_embeds
881	hidden_states = hidden_states.unsqueeze(1)
882
883	deepstack_feature_lists = []
884	for layer_num, blk in enumerate(self.blocks):
885	hidden_states = blk(
886	hidden_states,
887	cu_seqlens=encoder_metadata["cu_seqlens"],
888	rotary_pos_emb_cos=encoder_metadata["rotary_pos_emb_cos"],
889	rotary_pos_emb_sin=encoder_metadata["rotary_pos_emb_sin"],
890	max_seqlen=encoder_metadata["max_seqlen"],
891	sequence_lengths=encoder_metadata.get("sequence_lengths"),
892	)
893	if layer_num in self.deepstack_visual_indexes:
894	deepstack_merger_idx = self.deepstack_visual_indexes.index(layer_num)
895	deepstack_feature = self.deepstack_merger_list[deepstack_merger_idx](
896	hidden_states
897	)
898	deepstack_feature_lists.append(deepstack_feature)
899	hidden_states = self.merger(hidden_states)
900	hidden_states = torch.cat(
901	[hidden_states] + deepstack_feature_lists, dim=1
902	) # [seq_len, hidden_size * (1 + depth_of_deepstack)]
903	return hidden_states
904
905	def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
906	loader = AutoWeightsLoader(self)
907	return loader.load_weights(weights, mapper=self.hf_to_vllm_mapper)
908
909
910	class Qwen3VLMultiModalDataParser(Qwen2VLMultiModalDataParser):
911	# Timestamps are part of the prompt replacement (each frame group's
912	# "<X.X seconds>" text), so they size the placeholder range and must
913	# travel with the grid when embeddings are delivered out of band.
914	embedding_fields = {
915	**Qwen2VLMultiModalDataParser.embedding_fields,
916	"video": {
917	**Qwen2VLMultiModalDataParser.embedding_fields["video"],
918	"timestamps": "metadata",
919	},
920	}
921
922
923	class Qwen3VLProcessingInfo(Qwen2VLProcessingInfo):
924	def get_hf_config(self):
925	return self.ctx.get_hf_config(Qwen3VLConfig)
926
927	def get_hf_processor(self, **kwargs: object) -> Qwen3VLProcessor:
928	return self.ctx.get_hf_processor(
929	Qwen3VLProcessor,
930	**kwargs,
931	)
932
933	def get_image_processor(self, **kwargs: object) -> Qwen2VLImageProcessor:
934	return self.get_hf_processor(**kwargs).image_processor
935
936	def get_video_processor(self, **kwargs: object) -> Qwen3VLVideoProcessor:
937	return self.get_hf_processor(**kwargs).video_processor
938
939	def get_data_parser(self):
940	return Qwen3VLMultiModalDataParser(
941	self.get_hf_config().vision_config.spatial_merge_size,
942	video_needs_metadata=True,
943	expected_hidden_size=self._get_expected_hidden_size(),
944	allow_missing_mm_embeddings=self.allow_missing_mm_embeddings,
945	)
946
947	def get_supported_mm_processor_kwargs(self) -> dict[str, set[str]]:
948	supported = super().get_supported_mm_processor_kwargs()
949
950	# vLLM has historically accepted `min_pixels` / `max_pixels` for
951	# Qwen3-VL video. Include them in the routing schema even though
952	# Qwen3VLVideoProcessor only accepts the equivalent `size`; they are
953	# folded into `size` before calling HF.
954	supported["videos_kwargs"].update({"min_pixels", "max_pixels"})
955	return supported
956
957	def _get_vision_info(
958	self,
959	*,
960	image_width: int,
961	image_height: int,
962	num_frames: int = 2,
963	do_resize: bool = True,
964	image_processor: Qwen2VLImageProcessor \| Qwen3VLVideoProcessor,
965	mm_kwargs: Mapping[str, object],
966	modality: Literal["image", "video"] \| None = None,
967	) -> tuple[ImageSize, int]:
968	is_video = isinstance(image_processor, Qwen3VLVideoProcessor)
969
970	hf_config = self.get_hf_config()
971	vision_config = hf_config.vision_config
972	patch_size = vision_config.patch_size
973	merge_size = vision_config.spatial_merge_size
974	temporal_patch_size = vision_config.temporal_patch_size
975
976	if modality is None:
977	modality = "video" if is_video else "image"
978
979	merged_mm_kwargs = self._merge_and_resolve_mm_processor_kwargs(mm_kwargs)
980	scope = "videos_kwargs" if modality == "video" else "images_kwargs"
981	scoped_mm_kwargs = merged_mm_kwargs.get(scope, {})
982	size = self._get_vision_size(
983	scoped_mm_kwargs,
984	default_size=image_processor.size,
985	)
986
987	if do_resize:
988	if is_video:
989	smart_resize = video_smart_resize
990	extra_kwargs = {
991	"num_frames": num_frames,
992	"temporal_factor": temporal_patch_size,
993	}
994	else:
995	smart_resize = image_smart_resize
996	extra_kwargs = {}
997
998	resized_height, resized_width = smart_resize(
999	height=image_height,
1000	width=image_width,
1001	factor=patch_size * merge_size,
1002	min_pixels=size["shortest_edge"],
1003	max_pixels=size["longest_edge"],
1004	**extra_kwargs,
1005	)
1006	preprocessed_size = ImageSize(width=resized_width, height=resized_height)
1007	else:
1008	preprocessed_size = ImageSize(width=image_width, height=image_height)
1009
1010	padded_num_frames = round_up(num_frames, temporal_patch_size)
1011
1012	grid_t = max(padded_num_frames // temporal_patch_size, 1)
1013	grid_h = preprocessed_size.height // patch_size
1014	grid_w = preprocessed_size.width // patch_size
1015
1016	num_patches = grid_t * grid_h * grid_w
1017	num_vision_tokens = num_patches // (merge_size**2)
1018
1019	return preprocessed_size, num_vision_tokens
1020
1021	def _get_max_video_frames(self, max_tokens: int, start_num_frames: int = 2) -> int:
1022	return super()._get_max_video_frames(
1023	max_tokens, start_num_frames=start_num_frames
1024	)
1025
1026	def get_num_frames_with_most_features(
1027	self,
1028	seq_len: int,
1029	mm_counts: Mapping[str, int],
1030	max_frames_per_video: int = DUMMY_VIDEO_NUM_FRAMES,
1031	) -> int:
1032	return super().get_num_frames_with_most_features(
1033	seq_len, mm_counts, max_frames_per_video=max_frames_per_video
1034	)
1035
1036	def get_max_video_tokens(
1037	self,
1038	seq_len: int,
1039	mm_counts: Mapping[str, int],
1040	) -> int:
1041	video_processor = self.get_video_processor()
1042
1043	merged_mm_kwargs = self._merge_and_resolve_mm_processor_kwargs({})
1044	video_mm_kwargs = merged_mm_kwargs.get("videos_kwargs", {})
1045	video_size = self._get_vision_size(
1046	video_mm_kwargs,
1047	default_size=video_processor.size,
1048	)
1049	temporal_patch_size = video_mm_kwargs.get(
1050	"temporal_patch_size", video_processor.temporal_patch_size
1051	)
1052
1053	# video_max_pixels contains the temporal compression factor,
1054	# so we divide by 2 to get the maximum number of image pixels.
1055	video_max_pixels = video_size["longest_edge"]
1056	target_width, target_height = self.get_image_size_with_most_features(
1057	max_pixels=video_max_pixels // temporal_patch_size
1058	)
1059	num_video_soft_tokens = self.get_num_video_tokens(
1060	image_width=target_width,
1061	image_height=target_height,
1062	num_frames=2,
1063	image_processor=video_processor,
1064	mm_kwargs={},
1065	)
1066	return num_video_soft_tokens
1067
1068	def _calculate_timestamps(
1069	self, indices: list[int] \| torch.Tensor, video_fps: float, merge_size: int
1070	):
1071	if not isinstance(indices, list):
1072	indices = indices.tolist()
1073	if len(indices) % merge_size != 0:
1074	# don't update metadata's frames_indices directly
1075	indices = indices + [indices[-1]] * (merge_size - len(indices) % merge_size)
1076	timestamps = [idx / video_fps for idx in indices]
1077	timestamps = [
1078	(timestamps[i] + timestamps[i + merge_size - 1]) / 2
1079	for i in range(0, len(timestamps), merge_size)
1080	]
1081	return timestamps
1082
1083	def _get_video_second_idx(
1084	self,
1085	metadata: dict[str, Any],
1086	do_sample_frames: bool \| None = None,
1087	sampled_fps: float \| None = None,
1088	sampled_num_frames: int \| None = None,
1089	) -> list[int]:
1090	video_processor = self.get_video_processor()
1091	temporal_patch_size = video_processor.temporal_patch_size
1092	indices = metadata["frames_indices"]
1093
1094	# metadata["fps"] refers to the true fps of the input video.
1095	video_fps = metadata["fps"]
1096	if do_sample_frames is None:
1097	do_sample_frames = metadata.get("do_sample_frames", False)
1098
1099	# If video frames are sampled in HF processor (instead of vLLM
1100	# video loader), we need to re-calculate the indices from original
1101	# metadata.
1102	if do_sample_frames:
1103	total_num_frames = metadata["total_num_frames"]
1104
1105	# When num_frames is explicitly provided, use it directly
1106	# instead of computing from fps. This mirrors the behavior of
1107	# HF's Qwen3VLVideoProcessor.sample_frames where num_frames
1108	# and fps are mutually exclusive.
1109	if sampled_num_frames is not None:
1110	num_frames = sampled_num_frames
1111	else:
1112	# here video_fps is the fps of the sampled video, and
1113	# metadata["fps"] refers to the fps of the original video.
1114	sampled_fps = sampled_fps if sampled_fps else video_processor.fps
1115	num_frames = int(total_num_frames / metadata["fps"] * sampled_fps)
1116
1117	num_frames = min(
1118	min(
1119	max(num_frames, video_processor.min_frames),
1120	video_processor.max_frames,
1121	),
1122	total_num_frames,
1123	)
1124	indices = (
1125	np.linspace(0, total_num_frames - 1, num_frames)
1126	.round()
1127	.astype(int)
1128	.tolist()
1129	)
1130	timestamps = self._calculate_timestamps(indices, video_fps, temporal_patch_size)
1131	return timestamps
1132
1133
1134	class Qwen3VLDummyInputsBuilder(BaseDummyInputsBuilder[Qwen3VLProcessingInfo]):
1135	def get_dummy_text(self, mm_counts: Mapping[str, int]) -> str:
1136	num_images = mm_counts.get("image", 0)
1137	num_videos = mm_counts.get("video", 0)
1138
1139	image_token = "<\|vision_start\|><\|image_pad\|><\|vision_end\|>"
1140	video_token = "<\|vision_start\|><\|video_pad\|><\|vision_end\|>"
1141
1142	return image_token * num_images + video_token * num_videos
1143
1144	def get_dummy_mm_data(
1145	self,
1146	seq_len: int,
1147	mm_counts: Mapping[str, int],
1148	mm_options: MultiModalDummyOptions,
1149	) -> MultiModalDataDict:
1150	video_overrides = mm_options.get("video")
1151
1152	target_image_width, target_image_height = (
1153	self.info.get_image_size_with_most_features()
1154	)
1155
1156	# treat videos as special images
1157	target_num_frames = 2
1158	if video_overrides:
1159	num_frames_override = video_overrides.num_frames
1160	if num_frames_override:
1161	if num_frames_override > target_num_frames:
1162	logger.warning(
1163	"video.num_frames override (%d) exceeds model's "
1164	"maximum number of frames (%d), will be ignored",
1165	num_frames_override,
1166	target_num_frames,
1167	)
1168	if num_frames_override < 2:
1169	logger.warning(
1170	"video.num_frames override (%d) cannot be less "
1171	"than 2, will be ignored",
1172	num_frames_override,
1173	)
1174	target_num_frames = min(target_num_frames, num_frames_override)
1175	target_num_frames = max(target_num_frames, 2)
1176
1177	video_processor = self.info.get_video_processor()
1178
1179	merged_mm_kwargs = self.info._merge_and_resolve_mm_processor_kwargs({})
1180	video_mm_kwargs = merged_mm_kwargs.get("videos_kwargs", {})
1181	video_size = self.info._get_vision_size(
1182	video_mm_kwargs,
1183	default_size=video_processor.size,
1184	)
1185	temporal_patch_size = video_mm_kwargs.get(
1186	"temporal_patch_size", video_processor.temporal_patch_size
1187	)
1188
1189	# video_max_pixels contains the temporal compression factor,
1190	# so we divide by 2 to get the maximum number of image pixels.
1191	video_max_pixels = video_size["longest_edge"]
1192
1193	# With the HF processor's per-frame pixel cap enabled
1194	# (cap_pixels_per_frame, huggingface/transformers#48071), a 2-frame
1195	# dummy is processed at only 2 * cap pixels, so memory profiling
1196	# underestimates the largest possible item: a fully sampled video
1197	# still reaches the whole longest_edge budget. Spread the budget
1198	# over enough frames that the cap is not binding for the dummy.
1199	# This takes precedence over a shrinking num_frames override, which
1200	# would otherwise reintroduce the under-profiling.
1201	if video_mm_kwargs.get("cap_pixels_per_frame"):
1202	max_video_tokens = video_mm_kwargs.get(
1203	"max_video_tokens",
1204	getattr(video_processor, "max_video_tokens", 768),
1205	)
1206	patch_size = video_mm_kwargs.get("patch_size", video_processor.patch_size)
1207	merge_size = video_mm_kwargs.get("merge_size", video_processor.merge_size)
1208	per_frame_cap = max_video_tokens * (patch_size * merge_size) ** 2
1209	target_num_frames = max(
1210	target_num_frames,
1211	min(video_processor.max_frames, cdiv(video_max_pixels, per_frame_cap)),
1212	)
1213
1214	target_video_width, target_video_height = (
1215	self.info.get_image_size_with_most_features(
1216	max_pixels=video_max_pixels // temporal_patch_size
1217	)
1218	)
1219	target_video_size, _ = self.info._get_vision_info(
1220	image_width=target_video_width,
1221	image_height=target_video_height,
1222	num_frames=target_num_frames,
1223	image_processor=video_processor,
1224	mm_kwargs={},
1225	)
1226	# NOTE: we need to do this check here since Qwen3-VL resizes video
1227	# frames depending on how many frames there are.
1228	target_video_width, target_video_height = (
1229	target_video_size.width,
1230	target_video_size.height,
1231	)
1232	if video_overrides:
1233	width_override = video_overrides.width
1234	if width_override:
1235	if width_override > target_video_width:
1236	logger.warning(
1237	"video.width override (%d) exceeds model's "
1238	"maximum width (%d), will be ignored",
1239	width_override,
1240	target_video_width,
1241	)
1242	target_video_width = min(target_video_width, width_override)
1243	height_override = video_overrides.height
1244	if height_override:
1245	if height_override > target_video_height:
1246	logger.warning(
1247	"video.height override (%d) exceeds model's "
1248	"maximum height (%d), will be ignored",
1249	height_override,
1250	target_video_height,
1251	)
1252	target_video_height = min(target_video_height, height_override)
1253
1254	return {
1255	"image": self._get_dummy_images(
1256	width=target_image_width,
1257	height=target_image_height,
1258	num_images=mm_counts.get("image", 0),
1259	overrides=mm_options.get("image"),
1260	),
1261	"video": self._get_dummy_videos(
1262	width=target_video_width,
1263	height=target_video_height,
1264	num_frames=target_num_frames,
1265	num_videos=mm_counts.get("video", 0),
1266	),
1267	}
1268
1269	def _get_dummy_videos(
1270	self,
1271	*,
1272	width: int,
1273	height: int,
1274	num_frames: int,
1275	num_videos: int,
1276	overrides: VideoDummyOptions \| None = None,
1277	) -> list[VideoItem]:
1278	videos = super()._get_dummy_videos(
1279	width=width,
1280	height=height,
1281	num_frames=num_frames,
1282	num_videos=num_videos,
1283	overrides=overrides,
1284	)
1285	videos = [v.copy() for v in videos]
1286
1287	video_items: list[VideoItem] = []
1288	for video in videos:
1289	video_num_frames = video.shape[0]
1290	video_metadata = {
1291	"fps": 2.0,
1292	"duration": video_num_frames / 2.0,
1293	"total_num_frames": video_num_frames,
1294	"frames_indices": list(range(video_num_frames)),
1295	"video_backend": "opencv",
1296	"do_sample_frames": False,
1297	}
1298	video_items.append((video, video_metadata))
1299
1300	return video_items
1301
1302
1303	def _replace_video_token_placeholders(
1304	prompt_ids: list[int],
1305	target: list[int],
1306	replacements: list[list[int]],
1307	) -> list[int]:
1308	"""Replace each video placeholder with its expanded sequence.
1309
1310	Args:
1311	prompt_ids: Token IDs of the original (unexpanded) prompt.
1312	target: Token-ID sequence to search for. New-style
1313	(transformers>=5.10) processors expand only the bare
1314	`[video_pad_id]`; older ones expand the full
1315	`[vision_start_id, video_pad_id, vision_end_id]` triplet.
1316	replacements: Per-video expanded token sequences, in prompt order.
1317
1318	Returns:
1319	Token IDs with every placeholder replaced.
1320
1321	"""
1322	result: list[int] = []
1323	repl_idx = 0
1324	i = 0
1325	n = len(prompt_ids)
1326	k = len(target)
1327	num_repl = len(replacements)
1328
1329	while i < n:
1330	if prompt_ids[i : i + k] == target:
1331	result.extend(replacements[repl_idx])
1332	repl_idx += 1
1333	i += k
1334	else:
1335	result.append(prompt_ids[i])
1336	i += 1
1337
1338	assert repl_idx == num_repl, (
1339	f"Found {repl_idx} video placeholders but expected {num_repl}"
1340	)
1341	return result
1342
1343
1344	class Qwen3VLMultiModalProcessor(BaseMultiModalProcessor[Qwen3VLProcessingInfo]):
1345	@staticmethod
1346	def _expands_only_video_token(hf_processor: ProcessorMixin) -> bool:
1347	"""Transformers>=5.10 processors override `replace_video_token`
1348	to expand only the bare video token, keeping the prompt's outer
1349	`<\|vision_start\|>`/`<\|vision_end\|>` markers."""
1350	mixin_impl = getattr(ProcessorMixin, "replace_video_token", None)
1351	proc_impl = getattr(type(hf_processor), "replace_video_token", None)
1352	return proc_impl is not None and proc_impl is not mixin_impl
1353
1354	def _get_hf_mm_text(self, mm_counts: Mapping[str, int]) -> str:
1355	return self.dummy_inputs.get_dummy_text(mm_counts)
1356
1357	def _apply_hf_processor_main(
1358	self,
1359	mm_items: MultiModalDataItems,
1360	hf_kwargs: Mapping[str, object],
1361	) -> BatchFeature:
1362	hf_data, hf_kwargs, passthrough_data = self._get_hf_mm_inputs(
1363	mm_items, hf_kwargs
1364	)
1365
1366	if not hf_data:
1367	return self._finalize_hf_mm_data(hf_data, hf_kwargs, passthrough_data)
1368
1369	# Merge and resolve the multimodal processor kwargs once so every HF
1370	# call below starts from the same engine/request state and scoped form.
1371	merged_mm_kwargs = self.info._merge_and_resolve_mm_processor_kwargs(hf_kwargs)
1372
1373	if "images_kwargs" in merged_mm_kwargs:
1374	image_scoped_kwargs = dict(merged_mm_kwargs["images_kwargs"])
1375
1376	# Qwen2VLImageProcessor requires both `size` edges. If a partial size
1377	# override is passed, fill missing edges from the processor defaults
1378	# before any HF call.
1379	if image_scoped_kwargs.keys() & {"size", "min_pixels", "max_pixels"}:
1380	default_image_size = self.info.get_image_processor().size
1381	image_scoped_kwargs["size"] = self.info._get_vision_size(
1382	image_scoped_kwargs,
1383	default_size=default_image_size,
1384	)
1385
1386	merged_mm_kwargs["images_kwargs"] = image_scoped_kwargs
1387
1388	if "videos_kwargs" in merged_mm_kwargs:
1389	video_scoped_kwargs = dict(merged_mm_kwargs["videos_kwargs"])
1390
1391	# Qwen3VLVideoProcessor requires both `size` edges. If a partial size
1392	# override is passed, fill missing edges from the processor defaults
1393	# before any HF call.
1394	if video_scoped_kwargs.keys() & {"size", "min_pixels", "max_pixels"}:
1395	default_video_size = self.info.get_video_processor().size
1396	video_scoped_kwargs["size"] = self.info._get_vision_size(
1397	video_scoped_kwargs,
1398	default_size=default_video_size,
1399	)
1400
1401	# Qwen3VLVideoProcessor does not accept vLLM's min/max pixel
1402	# compatibility kwargs. They have already been folded into `size`,
1403	# so remove them before any HF call, including calls that contain
1404	# only text/images.
1405	video_scoped_kwargs.pop("min_pixels", None)
1406	video_scoped_kwargs.pop("max_pixels", None)
1407
1408	merged_mm_kwargs["videos_kwargs"] = video_scoped_kwargs
1409
1410	# Separate video processing from image processing. Because the videos
1411	# are processed into several image patches
1412	video_input_ids_lst: list[list[int]] = []
1413	if videos := hf_data.pop("videos", []):
1414	assert isinstance(videos, list)
1415	video_grid_thw_lst = []
1416	pixel_values_videos_lst = []
1417	timestamps_per_video = []
1418
1419	hf_config = self.info.get_hf_config()
1420	tokenizer = self.info.get_tokenizer()
1421	merge_size = hf_config.vision_config.spatial_merge_size
1422	pruning_spec = self.info.ctx.get_mm_config().get_video_pruning_spec()
1423	vision_start_token_id = hf_config.vision_start_token_id
1424	vision_end_token_id = hf_config.vision_end_token_id
1425	video_token_id = hf_config.video_token_id
1426
1427	for item_idx, item in enumerate(videos):
1428	video_array, metadata = item
1429
1430	# NOTE: @JJJYmmm new attr metadata.frames_indices indicates
1431	# the sampled frames indices of pre-sampled videos, which is
1432	# used to calculate the timestamps. Make sure that
1433	# do_sample_frames in mm_kwargs is false for presampled videos.
1434
1435	# The root copy is shallow, so give each item its own video
1436	# scope before modifying per-video values below.
1437	video_mm_kwargs = dict(merged_mm_kwargs)
1438	video_scoped_kwargs = dict(video_mm_kwargs.get("videos_kwargs", {}))
1439	video_mm_kwargs["videos_kwargs"] = video_scoped_kwargs
1440
1441	sampled_fps = video_scoped_kwargs.get("fps")
1442	if is_list_of(sampled_fps, float):
1443	video_scoped_kwargs["fps"] = sampled_fps[item_idx]
1444	sampled_num_frames = video_scoped_kwargs.get("num_frames")
1445	if is_list_of(sampled_num_frames, int):
1446	video_scoped_kwargs["num_frames"] = sampled_num_frames[item_idx]
1447	if "do_sample_frames" not in video_scoped_kwargs:
1448	# qwen_vl_utils already has "do_sample_frames" in
1449	# mm_kwargs, don't overwrite it.
1450	video_scoped_kwargs["do_sample_frames"] = metadata.get(
1451	"do_sample_frames", False
1452	)
1453
1454	metadata = VideoMetadata(
1455	**{k: metadata[k] for k in metadata if k != "do_sample_frames"}
1456	)
1457
1458	# Compute timestamps here where we have access to metadata
1459	timestamps = self.info._get_video_second_idx(
1460	metadata=metadata,
1461	do_sample_frames=video_scoped_kwargs["do_sample_frames"],
1462	sampled_fps=video_scoped_kwargs.get("fps"),
1463	sampled_num_frames=video_scoped_kwargs.get("num_frames"),
1464	)
1465	timestamps_per_video.append(timestamps)
1466
1467	video_mm_data = dict()
1468	video_mm_data["videos"] = [[video_array]]
1469	video_mm_data["video_metadata"] = [[metadata]]
1470
1471	# When num_frames is specified, explicitly set fps=None
1472	# to prevent HF's BaseVideoProcessor.preprocess() from
1473	# filling in the class default (fps=2) via setdefault(),
1474	# which would conflict with num_frames (mutually exclusive).
1475	if (
1476	"num_frames" in video_scoped_kwargs
1477	and "fps" not in video_scoped_kwargs
1478	):
1479	video_scoped_kwargs["fps"] = None
1480
1481	video_outputs = self.info.ctx.call_hf_processor(
1482	self.info.get_hf_processor(**video_mm_kwargs),
1483	dict(
1484	text="<\|vision_start\|><\|video_pad\|><\|vision_end\|>",
1485	**video_mm_data,
1486	),
1487	video_mm_kwargs,
1488	)
1489
1490	# Discard HF output input_ids — we use get_video_repl below
1491	# to generate the correct (EVS-adjusted) token sequence.
1492	video_outputs.pop("input_ids", None)
1493
1494	video_grid_thw = video_outputs["video_grid_thw"]
1495	num_frames = int(video_grid_thw[0, 0])
1496	tokens_per_frame_base = int(video_grid_thw[0, 1:].prod()) // (
1497	merge_size**2
1498	)
1499
1500	# Apply video pruning (EVS or VidCom2) if enabled.
1501	if pruning_spec is not None:
1502	method, prune_q = pruning_spec
1503	count_fn = (
1504	vidcom2_compute_retained_tokens_count
1505	if method == "vidcom2"
1506	else compute_retained_tokens_count
1507	)
1508	num_tokens = count_fn(
1509	tokens_per_frame=tokens_per_frame_base,
1510	num_frames=num_frames,
1511	q=prune_q,
1512	)
1513	tokens_per_frame = [num_tokens] + [0] * (num_frames - 1)
1514	select_token_id = False
1515	else:
1516	tokens_per_frame = [tokens_per_frame_base] * num_frames
1517	select_token_id = True
1518
1519	video_repl = Qwen3VLMultiModalProcessor.get_video_repl(
1520	tokens_per_frame=tokens_per_frame,
1521	timestamps=timestamps,
1522	tokenizer=tokenizer,
1523	vision_start_token_id=vision_start_token_id,
1524	vision_end_token_id=vision_end_token_id,
1525	video_token_id=video_token_id,
1526	select_token_id=select_token_id,
1527	)
1528	video_input_ids_lst.append(list(video_repl.full))
1529
1530	video_grid_thw_lst.append(video_outputs["video_grid_thw"])
1531	pixel_values_videos_lst.append(video_outputs["pixel_values_videos"])
1532	video_outputs = dict(
1533	pixel_values_videos=torch.cat(pixel_values_videos_lst),
1534	video_grid_thw=torch.cat(video_grid_thw_lst),
1535	timestamps=timestamps_per_video,
1536	)
1537	else:
1538	video_outputs = dict()
1539
1540	# The root copy is shallow, so copy the video scope before removing
1541	# per-video kwargs to avoid mutating `merged_mm_kwargs`.
1542	non_video_mm_kwargs = dict(merged_mm_kwargs)
1543	if "videos_kwargs" in non_video_mm_kwargs:
1544	non_video_scoped_kwargs = dict(non_video_mm_kwargs["videos_kwargs"])
1545	non_video_scoped_kwargs.pop("fps", None)
1546	non_video_scoped_kwargs.pop("num_frames", None)
1547	non_video_mm_kwargs["videos_kwargs"] = non_video_scoped_kwargs
1548	processed_data = self.info.ctx.call_hf_processor(
1549	self.info.get_hf_processor(**non_video_mm_kwargs),
1550	hf_data,
1551	non_video_mm_kwargs,
1552	)
1553
1554	# Replace each placeholder with pre-computed video tokens.
1555	if video_input_ids_lst:
1556	hf_config = self.info.get_hf_config()
1557	if self._expands_only_video_token(self.info.get_hf_processor()):
1558	video_target = [hf_config.video_token_id]
1559	else:
1560	video_target = [
1561	hf_config.vision_start_token_id,
1562	hf_config.video_token_id,
1563	hf_config.vision_end_token_id,
1564	]
1565	input_ids = processed_data.pop("input_ids")
1566	if not isinstance(input_ids, list):
1567	input_ids = input_ids.tolist()
1568	(prompt_ids,) = input_ids
1569	expanded_ids = _replace_video_token_placeholders(
1570	prompt_ids, video_target, video_input_ids_lst
1571	)
1572	processed_data["input_ids"] = [expanded_ids]
1573
1574	processed_data.update(video_outputs)
1575	return self._finalize_hf_mm_data(
1576	hf_data, hf_kwargs, passthrough_data, processed_data
1577	)
1578
1579	def _get_mm_fields_config(
1580	self,
1581	hf_inputs: BatchFeature,
1582	hf_processor_mm_kwargs: Mapping[str, object],
1583	) -> Mapping[str, MultiModalFieldConfig]:
1584	return _create_qwen2vl_field_factory(
1585	self.info.get_hf_config().vision_config.spatial_merge_size
1586	)(hf_inputs)
1587
1588	def _get_prompt_updates(
1589	self,
1590	mm_items: MultiModalDataItems,
1591	hf_processor_mm_kwargs: Mapping[str, Any],
1592	out_mm_kwargs: MultiModalKwargsItems,
1593	) -> Sequence[PromptUpdate]:
1594	hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs)
1595	image_processor = self.info.get_image_processor(**hf_processor_mm_kwargs)
1596	tokenizer = self.info.get_tokenizer()
1597	hf_config = self.info.get_hf_config()
1598
1599	video_token_id = hf_config.video_token_id
1600	vision_start_token_id = hf_config.vision_start_token_id
1601	vision_end_token_id = hf_config.vision_end_token_id
1602
1603	merge_length = image_processor.merge_size**2
1604
1605	def get_image_replacement_qwen3vl(item_idx: int):
1606	out_item = out_mm_kwargs["image"][item_idx]
1607	grid_thw = out_item["image_grid_thw"].data
1608	assert isinstance(grid_thw, torch.Tensor)
1609
1610	num_tokens = int(grid_thw.prod()) // merge_length
1611	return [hf_processor.image_token_id] * num_tokens
1612
1613	def get_video_replacement_qwen3vl(item_idx: int):
1614	out_item = out_mm_kwargs["video"][item_idx]
1615	grid_thw = out_item["video_grid_thw"].data
1616	assert isinstance(grid_thw, torch.Tensor)
1617
1618	sampled_fps = hf_processor_mm_kwargs.get("fps")
1619	if is_list_of(sampled_fps, float):
1620	sampled_fps = sampled_fps[item_idx]
1621
1622	timestamps = out_item["timestamps"].data
1623	assert isinstance(timestamps, (list, torch.Tensor))
1624	assert len(timestamps) == grid_thw[0], (
1625	f"The timestamps length({len(timestamps)}) should be equal "
1626	f"video length ({grid_thw[0]})."
1627	)
1628
1629	# Compute tokens per frame, with EVS / VidCom2 support
1630	num_frames = int(grid_thw[0])
1631	tokens_per_frame_base = int(grid_thw[1:].prod()) // merge_length
1632
1633	pruning_spec = self.info.ctx.get_mm_config().get_video_pruning_spec()
1634	if pruning_spec is not None:
1635	method, prune_q = pruning_spec
1636	count_fn = (
1637	vidcom2_compute_retained_tokens_count
1638	if method == "vidcom2"
1639	else compute_retained_tokens_count
1640	)
1641	num_tokens = count_fn(
1642	tokens_per_frame=tokens_per_frame_base,
1643	num_frames=num_frames,
1644	q=prune_q,
1645	)
1646	tokens_per_frame = [num_tokens] + [0] * (num_frames - 1)
1647	select_token_id = False
1648	else:
1649	tokens_per_frame = [tokens_per_frame_base] * num_frames
1650	select_token_id = True
1651
1652	return Qwen3VLMultiModalProcessor.get_video_repl(
1653	tokens_per_frame=tokens_per_frame,
1654	timestamps=timestamps,
1655	tokenizer=tokenizer,
1656	vision_start_token_id=vision_start_token_id,
1657	vision_end_token_id=vision_end_token_id,
1658	video_token_id=video_token_id,
1659	select_token_id=select_token_id,
1660	)
1661
1662	if self._expands_only_video_token(hf_processor):
1663	# transformers>=5.10 expands only the bare video_token
1664	video_target = [video_token_id]
1665	else:
1666	# Old-style processors expand the full placeholder
1667	video_target = [
1668	vision_start_token_id,
1669	video_token_id,
1670	vision_end_token_id,
1671	]
1672
1673	return [
1674	PromptReplacement(
1675	modality="image",
1676	target=[hf_processor.image_token_id],
1677	replacement=get_image_replacement_qwen3vl,
1678	),
1679	PromptReplacement(
1680	modality="video",
1681	target=video_target,
1682	replacement=get_video_replacement_qwen3vl,
1683	),
1684	]
1685
1686	@staticmethod
1687	def get_video_repl(
1688	*,
1689	tokens_per_frame: list[int],
1690	timestamps: Sequence[float] \| torch.Tensor,
1691	tokenizer: TokenizerLike,
1692	vision_start_token_id: int,
1693	vision_end_token_id: int,
1694	video_token_id: int,
1695	select_token_id: bool = False,
1696	) -> PromptUpdateDetails:
1697	"""Build prompt replacement for a video in Qwen3VL format.
1698
1699	The replacement structure for each frame is:
1700	timestamp_tokens + vision_start_token + video_tokens + vision_end_token
1701
1702	Args:
1703	tokens_per_frame: Number of video tokens per frame (can vary per frame for
1704	EVS).
1705	timestamps: List of timestamps in seconds for each frame
1706	tokenizer: Tokenizer to encode timestamp strings
1707	vision_start_token_id: Token ID for vision start marker
1708	vision_end_token_id: Token ID for vision end marker
1709	video_token_id: Token ID for video content
1710	select_token_id: Whether to return token ids instead of text
1711
1712	Returns:
1713	PromptUpdateDetails with full token sequence
1714
1715	"""
1716	assert len(timestamps) == len(tokens_per_frame), (
1717	"timestamps and tokens_per_frame must have the same length"
1718	)
1719
1720	# Tokenize timestamp strings independently to avoid tokenizer merging
1721	# tokens across boundaries.
1722	timestamp_token_ids = [
1723	tokenizer.encode(f"<{timestamp:.1f} seconds>", add_special_tokens=False)
1724	for timestamp in timestamps
1725	]
1726
1727	# Build the full token sequence
1728	all_token_ids = []
1729	for frame_timestamp_ids, num_tokens in zip(
1730	timestamp_token_ids, tokens_per_frame
1731	):
1732	# Add timestamp tokens
1733	all_token_ids.extend(frame_timestamp_ids)
1734
1735	# Add vision tokens: vision_start + video_tokens + vision_end
1736	all_token_ids.append(vision_start_token_id)
1737	all_token_ids.extend([video_token_id] * num_tokens)
1738	all_token_ids.append(vision_end_token_id)
1739
1740	if select_token_id:
1741	return PromptUpdateDetails.select_token_id(all_token_ids, video_token_id)
1742
1743	# NOTE: we use `from_seq` instead of `select_token_id` because we want all
1744	# tokens in the placeholder to be initially marked as candidates. Then
1745	# in `get_input_embeddings``, we refine the mask to only replace
1746	# `video_token_id` / `image_token_id`` positions with video/image embeddings,
1747	# keeping text embeddings for timestamps and structural tokens.
1748	return PromptUpdateDetails.from_seq(all_token_ids)
1749
1750
1751	@support_torch_compile(
1752	dynamic_arg_dims={
1753	"input_ids": 0,
1754	# positions is of shape (3, seq_len) if mrope is enabled for qwen2-vl,
1755	# otherwise (seq_len, ).
1756	"positions": -1,
1757	"intermediate_tensors": 0,
1758	"inputs_embeds": 0,
1759	# the same shape as input_embeds
1760	"deepstack_input_embeds": 0,
1761	}
1762	)
1763	class Qwen3LLMModel(Qwen3Model):
1764	def forward(
1765	self,
1766	input_ids: torch.Tensor \| None,
1767	positions: torch.Tensor,
1768	intermediate_tensors: IntermediateTensors \| None = None,
1769	inputs_embeds: torch.Tensor \| None = None,
1770	# args for deepstack
1771	deepstack_input_embeds: IntermediateTensors \| None = None,
1772	) -> torch.Tensor \| IntermediateTensors:
1773	if get_pp_group().is_first_rank:
1774	if inputs_embeds is not None:
1775	hidden_states = inputs_embeds
1776	else:
1777	hidden_states = self.embed_input_ids(input_ids)
1778	residual = None
1779	else:
1780	assert intermediate_tensors is not None
1781	hidden_states = intermediate_tensors["hidden_states"]
1782	residual = intermediate_tensors["residual"]
1783
1784	aux_hidden_states = self._maybe_add_hidden_state([], 0, hidden_states, residual)
1785	for layer_idx, layer in islice(
1786	enumerate(self.layers), self.start_layer, self.end_layer
1787	):
1788	hidden_states, residual = layer(
1789	positions,
1790	hidden_states,
1791	residual,
1792	)
1793
1794	if deepstack_input_embeds is not None and layer_idx in range(
1795	0, len(deepstack_input_embeds)
1796	):
1797	hidden_states = (
1798	hidden_states
1799	+ deepstack_input_embeds[f"deepstack_input_embeds_{layer_idx}"]
1800	)
1801	self._maybe_add_hidden_state(
1802	aux_hidden_states, layer_idx + 1, hidden_states, residual
1803	)
1804
1805	if not get_pp_group().is_last_rank:
1806	return IntermediateTensors(
1807	{"hidden_states": hidden_states, "residual": residual}
1808	)
1809	hidden_states, _ = self.norm(hidden_states, residual)
1810
1811	if len(aux_hidden_states) > 0:
1812	return hidden_states, aux_hidden_states
1813	return hidden_states
1814
1815
1816	class Qwen3LLMForCausalLM(Qwen3ForCausalLM):
1817	def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
1818	super(Qwen3ForCausalLM, self).__init__()
1819	config = vllm_config.model_config.hf_config
1820	quant_config = vllm_config.quant_config
1821
1822	self.config = config
1823
1824	self.quant_config = quant_config
1825	self.model = Qwen3LLMModel(
1826	vllm_config=vllm_config, prefix=maybe_prefix(prefix, "model")
1827	)
1828
1829	if get_pp_group().is_last_rank:
1830	self.lm_head = ParallelLMHead(
1831	config.vocab_size,
1832	config.hidden_size,
1833	quant_config=quant_config,
1834	prefix="lm_head",
1835	)
1836	if config.tie_word_embeddings:
1837	self.lm_head = self.lm_head.tie_weights(self.model.embed_tokens)
1838	else:
1839	self.lm_head = PPMissingLayer()
1840
1841	self.logits_processor = LogitsProcessor(config.vocab_size)
1842
1843	self.make_empty_intermediate_tensors = (
1844	self.model.make_empty_intermediate_tensors
1845	)
1846
1847
1848	@MULTIMODAL_REGISTRY.register_processor(
1849	Qwen3VLMultiModalProcessor,
1850	info=Qwen3VLProcessingInfo,
1851	dummy_inputs=Qwen3VLDummyInputsBuilder,
1852	)
1853	class Qwen3VLForConditionalGeneration(
1854	nn.Module,
1855	SupportsMultiModal,
1856	SupportsEncoderCudaGraph,
1857	SupportsLoRA,
1858	SupportsPP,
1859	SupportsMRoPE,
1860	SupportsEagle,
1861	SupportsEagle3,
1862	SupportsMultiModalPruning,
1863	):
1864	packed_modules_mapping: dict[str, list[str]] = {
1865	"qkv_proj": [
1866	"q_proj",
1867	"k_proj",
1868	"v_proj",
1869	],
1870	"gate_up_proj": [
1871	"gate_proj",
1872	"up_proj",
1873	],
1874	"qkv": ["qkv"], # For vision tower's already-packed QKV
1875	}
1876
1877	supports_encoder_tp_data = True
1878	supports_tower_connector_lora = True
1879	supports_mm_device_do_normalize = True
1880
1881	supported_video_pruning_methods = ("evs", "vidcom2")
1882
1883	# To ensure correct weight loading and mapping.
1884	hf_to_vllm_mapper = WeightsMapper(
1885	orig_to_new_prefix={
1886	"model.visual.": "visual.",
1887	"lm_head.": "language_model.lm_head.",
1888	"model.language_model.": "language_model.model.",
1889	}
1890	)
1891
1892	@classmethod
1893	def get_placeholder_str(cls, modality: str, i: int) -> str \| None:
1894	if modality.startswith("image"):
1895	return "<\|vision_start\|><\|image_pad\|><\|vision_end\|>"
1896	if modality.startswith("video"):
1897	return "<\|vision_start\|><\|video_pad\|><\|vision_end\|>"
1898
1899	raise ValueError("Only image or video modality is supported")
1900
1901	def _init_video_pruning(self, multimodal_config: MultiModalConfig) -> None:
1902	pruning_spec = multimodal_config.get_video_pruning_spec()
1903	if pruning_spec is None:
1904	self.video_pruning_method: VideoPruningMethod \| None = None
1905	self.video_pruning_rate = multimodal_config.video_pruning_rate
1906	else:
1907	self.video_pruning_method, self.video_pruning_rate = pruning_spec
1908	self.is_multimodal_pruning_enabled = (
1909	multimodal_config.is_multimodal_pruning_enabled()
1910	)
1911
1912	def __init__(self, *, vllm_config: VllmConfig, prefix: str = "model"):
1913	super().__init__()
1914	config: Qwen3VLConfig = vllm_config.model_config.hf_config
1915	quant_config = vllm_config.quant_config
1916	multimodal_config = vllm_config.model_config.get_multimodal_config()
1917
1918	self.config = config
1919	self.model_config = vllm_config.model_config
1920	self._tokenizer = cached_tokenizer_from_config(vllm_config.model_config)
1921	self.multimodal_config = multimodal_config
1922	self.use_data_parallel = multimodal_config.mm_encoder_tp_mode == "data"
1923	self._init_video_pruning(multimodal_config)
1924
1925	with self._mark_tower_model(vllm_config, {"image", "video"}):
1926	self.visual = Qwen3_VisionTransformer(
1927	config.vision_config,
1928	norm_eps=getattr(config, "rms_norm_eps", 1e-6),
1929	quant_config=quant_config,
1930	input_norm=build_mm_input_norm(self.model_config),
1931	prefix=maybe_prefix(prefix, "visual"),
1932	)
1933
1934	self.use_deepstack = hasattr(
1935	config.vision_config, "deepstack_visual_indexes"
1936	) and not isinstance(self.visual, StageMissingLayer)
1937	self.deepstack_num_level = (
1938	len(config.vision_config.deepstack_visual_indexes)
1939	if self.use_deepstack
1940	else 0
1941	)
1942	self.visual_dim = config.vision_config.out_hidden_size
1943	self.multiscale_dim = self.visual_dim * self.deepstack_num_level
1944
1945	if self.use_deepstack:
1946	self.deepstack_input_embeds = [
1947	torch.zeros(
1948	vllm_config.scheduler_config.max_num_batched_tokens,
1949	config.text_config.hidden_size,
1950	)
1951	for _ in range(self.deepstack_num_level)
1952	]
1953	# Tracks the valid token span currently stored in the buffer.
1954	# Zero means there is no active deepstack payload to consume.
1955	self.deepstack_input_embeds_num_tokens = 0
1956
1957	with self._mark_language_model(vllm_config):
1958	self.language_model = Qwen3LLMForCausalLM(
1959	vllm_config=vllm_config.with_hf_config(
1960	config.text_config,
1961	architectures=["Qwen3ForCausalLM"],
1962	),
1963	prefix=maybe_prefix(prefix, "language_model"),
1964	)
1965
1966	if not get_pp_group().is_first_rank and hasattr(
1967	config.vision_config, "deepstack_visual_indexes"
1968	):
1969	assert self.language_model.model.start_layer >= len(
1970	config.vision_config.deepstack_visual_indexes
1971	), (
1972	"start_layer should be greater than or equal to "
1973	"len(deepstack_visual_indexes)"
1974	)
1975
1976	self.make_empty_intermediate_tensors = (
1977	self.language_model.make_empty_intermediate_tensors
1978	)
1979
1980	def _get_deepstack_input_embeds(
1981	self,
1982	num_tokens: int,
1983	) -> IntermediateTensors \| None:
1984	if not getattr(self, "deepstack_input_embeds", None):
1985	return None # If vision tower is skipped
1986	if num_tokens > self.deepstack_input_embeds[0].size(0):
1987	self._resize_deepstack_input_embeds(num_tokens)
1988
1989	# get deepstack_input_embeds from buffer, and clear the buffer
1990	return IntermediateTensors(
1991	{
1992	f"deepstack_input_embeds_{idx}": self.deepstack_input_embeds[idx][
1993	:num_tokens
1994	]
1995	for idx in range(self.deepstack_num_level)
1996	}
1997	)
1998
1999	def _resize_deepstack_input_embeds(self, num_tokens: int) -> None:
2000	self.deepstack_input_embeds = [
2001	torch.zeros(
2002	num_tokens,
2003	self.config.text_config.hidden_size,
2004	device=self.deepstack_input_embeds[0].device,
2005	dtype=self.deepstack_input_embeds[0].dtype,
2006	)
2007	for _ in range(self.deepstack_num_level)
2008	]
2009
2010	def _set_deepstack_input_embeds(self, deepstack_input_embeds: torch.Tensor) -> None:
2011	if not getattr(self, "deepstack_input_embeds", None):
2012	return
2013
2014	# set deepstack_input_embeds to buffer
2015	num_tokens = deepstack_input_embeds.size(1)
2016	if num_tokens > self.deepstack_input_embeds[0].size(0):
2017	self._resize_deepstack_input_embeds(num_tokens)
2018	for idx in range(self.deepstack_num_level):
2019	self.deepstack_input_embeds[idx][:num_tokens].copy_(
2020	deepstack_input_embeds[idx]
2021	)
2022	self.deepstack_input_embeds_num_tokens = num_tokens
2023
2024	def _clear_deepstack_input_embeds(self, num_tokens: int) -> None:
2025	if not getattr(self, "deepstack_input_embeds", None):
2026	return
2027	if getattr(self, "deepstack_input_embeds_num_tokens", 0) == 0:
2028	return
2029
2030	# clear deepstack_input_embeds in buffer
2031	if num_tokens > 0:
2032	for idx in range(self.deepstack_num_level):
2033	self.deepstack_input_embeds[idx][:num_tokens].zero_()
2034	self.deepstack_input_embeds_num_tokens = 0
2035
2036	# -- SupportsEncoderCudaGraph protocol methods --
2037
2038	def get_encoder_cudagraph_config(self):
2039	from vllm.v1.worker.encoder_cudagraph_defs import (
2040	EncoderCudaGraphConfig,
2041	)
2042
2043	# When video pruning is enabled, embed_multimodal post-processes both
2044	# image and video embeddings (mrope positions are appended for image,
2045	# prune+append for video). The encoder CUDA graph path bypasses that
2046	# post-process, producing inconsistent embedding formats vs eager. So
2047	# disable CUDA graph for all modalities when pruning is on.
2048	modalities = [] if self.is_multimodal_pruning_enabled else ["image", "video"]
2049
2050	# Compute max_frames_per_video for budget sizing.
2051	max_frames = self.get_max_frames_per_video() if "video" in modalities else 1
2052
2053	return EncoderCudaGraphConfig(
2054	modalities=modalities,
2055	buffer_keys=[
2056	"pixel_values",
2057	"pos_embeds",
2058	"rotary_pos_emb_cos",
2059	"rotary_pos_emb_sin",
2060	"cu_seqlens",
2061	"max_seqlen",
2062	"sequence_lengths",
2063	],
2064	out_hidden_size=self.visual.out_hidden_size,
2065	max_frames_per_video=max_frames,
2066	)
2067
2068	def get_input_modality(
2069	self,
2070	mm_kwargs: dict[str, Any],
2071	) -> str:
2072	if "image_grid_thw" in mm_kwargs:
2073	return "image"
2074	elif "video_grid_thw" in mm_kwargs:
2075	return "video"
2076	raise AssertionError("This line should be unreachable.")
2077
2078	def get_max_frames_per_video(self) -> int:
2079	mm_registry = MULTIMODAL_REGISTRY
2080	info = mm_registry.get_processing_info(self.model_config)
2081	assert isinstance(info, Qwen3VLProcessingInfo)
2082	max_frames_per_video = info.get_num_frames_with_most_features(
2083	seq_len=self.model_config.max_model_len,
2084	mm_counts={"video": self.multimodal_config.get_limit_per_prompt("video")},
2085	)
2086	return max_frames_per_video
2087
2088	def get_encoder_cudagraph_budget_range(
2089	self,
2090	vllm_config,
2091	) -> tuple[int, int]:
2092	# Min: estimated smallest possible encoder input.
2093	# 224x224 image → 16x16 patches (patch_size=14)
2094	# spatial_merge_size=2 → 8x8 = 64 tokens
2095	min_budget = 64
2096	# Max: capped by max_num_batched_tokens
2097	# TODO(shen-shanshan): the max_budget auto-infer needs to be optimized later.
2098	max_budget = min(
2099	vllm_config.scheduler_config.max_num_batched_tokens,
2100	self.model_config.max_model_len,
2101	)
2102	return (min_budget, max_budget)
2103
2104	def _get_pixel_values_by_modality(
2105	self,
2106	mm_kwargs: dict[str, Any],
2107	) -> torch.Tensor:
2108	modality = self.get_input_modality(mm_kwargs)
2109	if modality == "image":
2110	return mm_kwargs["pixel_values"]
2111	elif modality == "video":
2112	return mm_kwargs["pixel_values_videos"]
2113	raise AssertionError("This line should be unreachable.")
2114
2115	def _get_grid_thw_by_modality(
2116	self,
2117	mm_kwargs: dict[str, Any],
2118	) -> list[tuple[int, int, int]]:
2119	grid_thw_key = f"{self.get_input_modality(mm_kwargs)}_grid_thw"
2120	grid_thw = mm_kwargs[grid_thw_key]
2121	if not isinstance(grid_thw, list):
2122	grid_thw = grid_thw.tolist()
2123	return grid_thw
2124
2125	def get_encoder_cudagraph_item_specs(
2126	self,
2127	mm_kwargs: dict[str, Any],
2128	):
2129	from vllm.v1.worker.encoder_cudagraph_defs import EncoderItemSpec
2130
2131	m = self.visual.spatial_merge_size
2132	grid_thw = self._get_grid_thw_by_modality(mm_kwargs)
2133	return [
2134	EncoderItemSpec(
2135	input_size=t * h * w,
2136	output_tokens=t * (h // m) * (w // m),
2137	)
2138	for t, h, w in grid_thw
2139	]
2140
2141	def select_encoder_cudagraph_items(
2142	self,
2143	mm_kwargs: dict[str, Any],
2144	indices: list[int],
2145	) -> dict[str, Any]:
2146	grid_thw = self._get_grid_thw_by_modality(mm_kwargs)
2147	pixel_values = self._get_pixel_values_by_modality(mm_kwargs)
2148
2149	if len(indices) == 0:
2150	if self.get_input_modality(mm_kwargs) == "image":
2151	return {
2152	"pixel_values": pixel_values[:0],
2153	"image_grid_thw": [],
2154	}
2155	else:
2156	return {
2157	"pixel_values_videos": pixel_values[:0],
2158	"video_grid_thw": [],
2159	}
2160
2161	# Compute cumulative patch offsets for slicing pixel_values
2162	patches_per_item = [t * h * w for t, h, w in grid_thw]
2163	cum_patches = [0]
2164	for p in patches_per_item:
2165	cum_patches.append(cum_patches[-1] + p)
2166
2167	selected_pv = torch.cat(
2168	[pixel_values[cum_patches[i] : cum_patches[i + 1]] for i in indices]
2169	)
2170	selected_grid = [grid_thw[i] for i in indices]
2171
2172	if self.get_input_modality(mm_kwargs) == "image":
2173	return {
2174	"pixel_values": selected_pv,
2175	"image_grid_thw": selected_grid,
2176	}
2177	else:
2178	return {
2179	"pixel_values_videos": selected_pv,
2180	"video_grid_thw": selected_grid,
2181	}
2182
2183	def prepare_encoder_cudagraph_capture_inputs(
2184	self,
2185	token_budget: int,
2186	max_batch_size: int,
2187	max_frames_per_batch: int,
2188	device: torch.device,
2189	dtype: torch.dtype,
2190	path: str = "default",
2191	axis_keys: tuple[Hashable, ...] \| None = None,
2192	):
2193	from vllm.v1.worker.encoder_cudagraph_defs import (
2194	EncoderCudaGraphCaptureInputs,
2195	)
2196
2197	spatial_merge_size = self.visual.spatial_merge_size
2198	# Ceil so the buffer fits the worst case of one item using the full
2199	# budget. Floor under-allocates when budget is not a multiple of
2200	# max_batch_size.
2201	per_mm_item_output = (token_budget + max_batch_size - 1) // max_batch_size
2202
2203	frames_per_item = max_frames_per_batch // max_batch_size
2204	if frames_per_item > 1:
2205	# Build the capture grid using a video-format layout so that
2206	# cu_seqlens is sized for video replays from the start.
2207	# cu_seqlens has one entry per attention sequence (one per frame),
2208	# so using T > 1 per item makes the buffer large enough without
2209	# relying solely on padding.
2210	# Ceiling ensures frames_per_item * tokens_per_frame >= per_mm_item_output
2211	# so the pixel_values buffer covers any valid single-item replay.
2212	tokens_per_frame = (
2213	per_mm_item_output + frames_per_item - 1
2214	) // frames_per_item
2215	# Video-format grid_config (T=frames_per_item).
2216	grid_config = [
2217	[
2218	frames_per_item,
2219	spatial_merge_size,
2220	tokens_per_frame * spatial_merge_size,
2221	]
2222	for _ in range(max_batch_size)
2223	]
2224	else:
2225	# Image-format grid_config (T=1).
2226	grid_config = [
2227	[1, spatial_merge_size, per_mm_item_output * spatial_merge_size]
2228	for _ in range(max_batch_size)
2229	]
2230
2231	# Create dummy pixel_values
2232	patch_embed = self.visual.patch_embed
2233	in_channels = patch_embed.proj.in_channels
2234	patch_size = patch_embed.patch_size
2235	temporal_patch_size = patch_embed.temporal_patch_size
2236	total_patches = sum(t * h * w for t, h, w in grid_config)
2237	flattened_patch_size = (
2238	in_channels * temporal_patch_size * patch_size * patch_size
2239	)
2240	dummy_pixel_values = torch.zeros(
2241	total_patches,
2242	flattened_patch_size,
2243	device=device,
2244	dtype=self.visual.input_norm.input_dtype or dtype,
2245	)
2246
2247	# Override max_seqlen with a safe upper bound for capture.
2248	# max_seqlen.item() gets baked into the CUDA graph (not replayed),
2249	# so the capture value must cover any replay scenario.
2250	# Worst case: 1 item consuming the full budget ->
2251	# seq_len = token_budget * spatial_merge_size^2.
2252	metadata = self.visual.prepare_encoder_metadata(
2253	grid_config,
2254	max_batch_size=max_batch_size,
2255	max_frames_per_batch=max_frames_per_batch,
2256	max_seqlen_override=token_budget * (spatial_merge_size**2),
2257	device=device,
2258	)
2259
2260	# Just use image-modality dummy input_buffer for capturing, since it's also
2261	# compatible for video inputs (has the same shape: [num_patches, CTP*P]).
2262	values = metadata \| {
2263	"pixel_values": dummy_pixel_values,
2264	}
2265
2266	return EncoderCudaGraphCaptureInputs(
2267	values=values,
2268	)
2269
2270	def prepare_encoder_cudagraph_replay_buffers(
2271	self,
2272	mm_kwargs: dict[str, Any],
2273	max_batch_size: int,
2274	max_frames_per_batch: int,
2275	path: str = "default",
2276	):
2277	modality = self.get_input_modality(mm_kwargs)
2278	grid_thw_list = self._get_grid_thw_by_modality(mm_kwargs)
2279
2280	if modality == "image":
2281	metadata = self.visual.prepare_encoder_metadata(
2282	grid_thw_list,
2283	max_batch_size=max_batch_size,
2284	)
2285	elif modality == "video":
2286	metadata = self.visual.prepare_encoder_metadata(
2287	grid_thw_list,
2288	max_frames_per_batch=max_frames_per_batch,
2289	)
2290	else:
2291	raise AssertionError("This line should be unreachable.")
2292
2293	values = metadata \| {
2294	"pixel_values": self._get_pixel_values_by_modality(mm_kwargs),
2295	}
2296	return EncoderCudaGraphReplayBuffers(values=values)
2297
2298	def encoder_cudagraph_forward(
2299	self,
2300	values: dict[str, torch.Tensor],
2301	path: str = "default",
2302	) -> torch.Tensor:
2303	pixel_values = values.pop("pixel_values")
2304	metadata = values
2305	return self.visual(pixel_values, None, encoder_metadata=metadata)
2306
2307	def encoder_eager_forward(
2308	self,
2309	mm_kwargs: dict[str, Any],
2310	path: str = "default",
2311	) -> torch.Tensor:
2312	pixel_values = self._get_pixel_values_by_modality(mm_kwargs)
2313	grid_thw = self._get_grid_thw_by_modality(mm_kwargs)
2314	return self.visual(pixel_values, grid_thw)
2315
2316	def _parse_and_validate_image_input(
2317	self, **kwargs: object
2318	) -> Qwen2_5_VLImageInputs \| None:
2319	pixel_values = kwargs.pop("pixel_values", None)
2320	image_embeds = kwargs.pop("image_embeds", None)
2321	image_grid_thw = kwargs.pop("image_grid_thw", None)
2322
2323	if pixel_values is None and image_embeds is None:
2324	return None
2325
2326	if pixel_values is not None:
2327	return Qwen2_5_VLImagePixelInputs(
2328	type="pixel_values",
2329	pixel_values=pixel_values,
2330	image_grid_thw=image_grid_thw,
2331	)
2332
2333	else:
2334	return Qwen2_5_VLImageEmbeddingInputs(
2335	type="image_embeds",
2336	image_embeds=image_embeds,
2337	image_grid_thw=image_grid_thw,
2338	)
2339
2340	def _parse_and_validate_video_input(
2341	self, **kwargs: object
2342	) -> Qwen2_5_VLVideoInputs \| None:
2343	pixel_values_videos = kwargs.pop("pixel_values_videos", None)
2344	video_embeds = kwargs.pop("video_embeds", None)
2345	video_grid_thw = kwargs.pop("video_grid_thw", None)
2346	second_per_grid_ts = kwargs.pop("second_per_grid_ts", None)
2347	timestamps = kwargs.pop("timestamps", None)
2348
2349	if pixel_values_videos is None and video_embeds is None:
2350	return None
2351
2352	if pixel_values_videos is not None:
2353	return Qwen2_5_VLVideoPixelInputs(
2354	type="pixel_values_videos",
2355	pixel_values_videos=pixel_values_videos,
2356	video_grid_thw=video_grid_thw,
2357	second_per_grid_ts=second_per_grid_ts,
2358	timestamps=timestamps,
2359	)
2360
2361	else:
2362	return Qwen2_5_VLVideoEmbeddingInputs(
2363	type="video_embeds",
2364	video_embeds=video_embeds,
2365	video_grid_thw=video_grid_thw,
2366	timestamps=timestamps,
2367	)
2368
2369	def _process_image_input(
2370	self, image_input: Qwen2_5_VLImageInputs
2371	) -> tuple[torch.Tensor, ...]:
2372	grid_thw = image_input["image_grid_thw"]
2373	assert grid_thw.ndim == 2
2374
2375	if image_input["type"] == "image_embeds":
2376	image_embeds = image_input["image_embeds"].type(self.visual.dtype)
2377	else:
2378	pixel_values = image_input["pixel_values"]
2379	if self.use_data_parallel:
2380	return run_dp_sharded_mrope_vision_model(
2381	self.visual, pixel_values, grid_thw.tolist(), rope_type="rope_3d"
2382	)
2383	else:
2384	image_embeds = self.visual(pixel_values, grid_thw=grid_thw)
2385
2386	# Split concatenated embeddings for each image item.
2387	merge_size = self.visual.spatial_merge_size
2388	sizes = (grid_thw.prod(-1) // merge_size // merge_size).tolist()
2389	return image_embeds.split(sizes)
2390
2391	def _process_video_input(
2392	self, video_input: Qwen2_5_VLVideoInputs
2393	) -> tuple[torch.Tensor, ...]:
2394	grid_thw = video_input["video_grid_thw"]
2395	assert grid_thw.ndim == 2
2396
2397	if video_input["type"] == "video_embeds":
2398	video_embeds = video_input["video_embeds"].type(self.visual.dtype)
2399	else:
2400	pixel_values_videos = video_input["pixel_values_videos"]
2401	if self.use_data_parallel:
2402	grid_thw_list = grid_thw.tolist()
2403	return run_dp_sharded_mrope_vision_model(
2404	self.visual, pixel_values_videos, grid_thw_list, rope_type="rope_3d"
2405	)
2406	else:
2407	video_embeds = self.visual(pixel_values_videos, grid_thw=grid_thw)
2408
2409	# Split concatenated embeddings for each video item.
2410	merge_size = self.visual.spatial_merge_size
2411	sizes = (grid_thw.prod(-1) // merge_size // merge_size).tolist()
2412	return video_embeds.split(sizes)
2413
2414	def _postprocess_image_embeds_evs(
2415	self,
2416	image_embeds_split: tuple[torch.Tensor, ...],
2417	image_input: Qwen2_5_VLImageInputs,
2418	) -> tuple[torch.Tensor, ...]:
2419	"""Append mrope positions for each for images.
2420	This is necessary to recover correct mrope
2421	positions after video pruning
2422
2423	Args:
2424	image_embeds_split: Tuple of image embeddings for
2425	each image item.
2426	image_input: Image input data.
2427
2428	Returns:
2429	Tuple of image embeddings for each image item.
2430	Resulting embeddings will have extra 5 channels for
2431	computed mrope positions, consistent with video embeddings.
2432
2433	"""
2434	if self.is_multimodal_pruning_enabled:
2435	merge_size = self.visual.spatial_merge_size
2436	grid_thw = image_input["image_grid_thw"]
2437	grid_thw_list = grid_thw.tolist()
2438	image_embeds_out = []
2439	for emb, size in zip(image_embeds_split, grid_thw_list):
2440	positions = async_tensor_h2d(
2441	compute_mrope_for_media(size, merge_size), emb.device
2442	)
2443	positions = torch.cat(
2444	[
2445	positions,
2446	torch.zeros_like(
2447	positions[:, 0:1]
2448	), # Dummy extra fifth channel
2449	],
2450	dim=1,
2451	)
2452	emb = torch.cat([emb, positions], dim=1)
2453	image_embeds_out.append(emb)
2454	image_embeds_split = tuple(image_embeds_out)
2455	return image_embeds_split
2456
2457	def _postprocess_video_embeds_evs(
2458	self,
2459	video_embeds_split: tuple[torch.Tensor, ...],
2460	video_input: Qwen2_5_VLVideoInputs,
2461	) -> tuple[torch.Tensor, ...]:
2462	"""Prunes video embeddings via Efficient Video Sampling (EVS)
2463	and then appends mrope positions for each retained embeddings
2464
2465	Args:
2466	video_embeds_split: Tuple of video embeddings for each video item.
2467	video_input: Video input data.
2468
2469	Returns:
2470	Tuple of video embeddings for each video item.
2471	Resulting embeddings will have extra 5 channels for computed mrope
2472	positions, and whether the index corresponds to a video embedding.
2473
2474	"""
2475	grid_thw = video_input["video_grid_thw"]
2476	assert grid_thw.ndim == 2
2477	grid_thw_list = grid_thw.tolist()
2478	merge_size = self.visual.spatial_merge_size
2479
2480	# Apply EVS to each video.
2481	video_embeds_out = []
2482	assert video_input.timestamps is not None
2483	for video_idx, (emb, size) in enumerate(zip(video_embeds_split, grid_thw_list)):
2484	# Compute positions.
2485	timestamps = video_input.timestamps[video_idx]
2486	num_frames = len(timestamps)
2487
2488	t, h, w = size
2489	if self.is_multimodal_pruning_enabled:
2490	assert self.video_pruning_rate is not None
2491	# Compute the retention mask for each video (EVS or VidCom2).
2492	if self.video_pruning_method == "vidcom2":
2493	mask_fn = vidcom2_compute_retention_mask
2494	else:
2495	mask_fn = compute_retention_mask
2496	retention_mask = mask_fn(
2497	emb,
2498	size,
2499	spatial_merge_size=self.visual.spatial_merge_size,
2500	q=self.video_pruning_rate,
2501	)
2502
2503	with gpu_sync_allowed():
2504	# Apply retention mask.
2505	emb = emb[retention_mask]
2506
2507	# Calculate the actual number of retained tokens per frame.
2508	num_frames, rows, cols = t, h // merge_size, w // merge_size
2509	retention_mask_thw = retention_mask.reshape(num_frames, rows, cols)
2510	num_tokens_per_frame = (
2511	retention_mask_thw.sum(dim=(1, 2)).long().tolist()
2512	)
2513	else:
2514	feature_size = emb.shape[0] // num_frames
2515	num_tokens_per_frame = [feature_size] * num_frames
2516	retention_mask = None
2517
2518	emb = self._create_final_video_embeddings(
2519	video_embeddings=emb,
2520	num_tokens_per_frame=num_tokens_per_frame,
2521	timestamps=timestamps,
2522	video_grid_thw=size,
2523	retention_mask=retention_mask,
2524	)
2525
2526	video_embeds_out.append(emb)
2527
2528	return tuple(video_embeds_out)
2529
2530	def _create_final_video_embeddings(
2531	self,
2532	video_embeddings: torch.Tensor,
2533	num_tokens_per_frame: list[int],
2534	timestamps: list[float],
2535	video_grid_thw: list[int],
2536	retention_mask: torch.Tensor,
2537	) -> torch.Tensor:
2538	"""Create final embeddings that combine video embeddings with
2539	text embeddings of indicator tokens.
2540
2541	These final embeddings contain:
2542	- Actual video embeddings in positions corresponding to video content
2543	- Text embeddings for indicator tokens (<img>, </img>, and
2544	frame separation text) in their respective positions
2545
2546	These embeddings will replace the placeholder embeddings to create
2547	input_embeds for the LLM.
2548	"""
2549	device = video_embeddings.device
2550
2551	# Generate video replacement token IDs using get_video_repl
2552	# This tokenizes each frame separator independently, then uses pre-tokenized
2553	# special tokens to ensure consistent tokenization regardless of
2554	# num_tokens_per_frame values.
2555	video_repl = Qwen3VLMultiModalProcessor.get_video_repl(
2556	tokens_per_frame=num_tokens_per_frame,
2557	tokenizer=self._tokenizer,
2558	timestamps=timestamps,
2559	vision_start_token_id=self.config.vision_start_token_id,
2560	vision_end_token_id=self.config.vision_end_token_id,
2561	video_token_id=self.config.video_token_id,
2562	select_token_id=self.is_multimodal_pruning_enabled,
2563	)
2564
2565	repl_token_ids = torch.tensor(video_repl.full, device=device)
2566	embed_token_id = _cached_tensor(self.config.video_token_id, device=device)
2567	is_video_embed = torch.isin(repl_token_ids, embed_token_id)
2568
2569	# Get text embeddings for indicator tokens (has only `visual_dim``).
2570	text_embeddings = self.get_language_model().embed_input_ids(repl_token_ids)
2571
2572	if self.use_deepstack:
2573	(
2574	deepstack_input_embeds,
2575	multimodal_embeddings,
2576	) = self._compute_deepstack_embeds(
2577	inputs_embeds=text_embeddings,
2578	multimodal_embeddings=[video_embeddings],
2579	is_multimodal=is_video_embed,
2580	)
2581	else:
2582	deepstack_input_embeds = None
2583	multimodal_embeddings = [video_embeddings]
2584
2585	merged_embeddings = _merge_multimodal_embeddings(
2586	inputs_embeds=text_embeddings,
2587	multimodal_embeddings=multimodal_embeddings,
2588	is_multimodal=is_video_embed,
2589	)
2590
2591	to_concat = [merged_embeddings]
2592	if deepstack_input_embeds is not None:
2593	to_concat.append(
2594	deepstack_input_embeds.permute(1, 0, 2).reshape(
2595	deepstack_input_embeds.shape[1], -1
2596	)
2597	)
2598
2599	expanded_positions = None
2600	if self.is_multimodal_pruning_enabled:
2601	is_vision_start = repl_token_ids.eq(self.config.vision_start_token_id)
2602	expanded_positions = self._get_expanded_positions(
2603	device=merged_embeddings.device,
2604	seq_len=merged_embeddings.shape[0],
2605	video_grid_thw=video_grid_thw,
2606	num_tokens_per_frame=num_tokens_per_frame,
2607	timestamps=timestamps,
2608	is_video_embed=is_video_embed,
2609	is_vision_start=is_vision_start,
2610	retention_mask=retention_mask,
2611	)
2612	to_concat.append(expanded_positions)
2613
2614	final_video_embeddings = torch.cat(to_concat, dim=-1)
2615
2616	return final_video_embeddings
2617
2618	def _get_expanded_positions(
2619	self,
2620	device,
2621	seq_len,
2622	video_grid_thw,
2623	num_tokens_per_frame,
2624	timestamps,
2625	is_video_embed,
2626	is_vision_start,
2627	retention_mask,
2628	):
2629	embed_token_id = _cached_tensor(self.config.video_token_id, device=device)
2630
2631	# Expand positions to match the full sequence length
2632	# (includes both video tokens and indicator tokens)
2633	# Shape: [full_length, 5] where positions are filled for video tokens
2634	# and zeros for indicator tokens.
2635	# Channel 3 flags VISION_START tokens so that
2636	# recompute_mrope_positions can reliably count timestamp tokens
2637	# (even when early frames have all video tokens pruned).
2638	# Channel 4 flags video-embedding tokens.
2639	expanded_positions = torch.zeros(
2640	seq_len,
2641	5, # [t_index, h_index, w_index, is_vision_start, is_video]
2642	device=device,
2643	dtype=torch.long,
2644	)
2645	_, h, w = video_grid_thw
2646	merge_size = self.visual.spatial_merge_size
2647	num_frames = len(num_tokens_per_frame)
2648	unpruned_token_ids = Qwen3VLMultiModalProcessor.get_video_repl(
2649	tokens_per_frame=[(h // merge_size) * (w // merge_size)] * num_frames,
2650	tokenizer=self._tokenizer,
2651	timestamps=timestamps,
2652	vision_start_token_id=self.config.vision_start_token_id,
2653	vision_end_token_id=self.config.vision_end_token_id,
2654	video_token_id=self.config.video_token_id,
2655	).full
2656	unpruned_token_ids_tensor = torch.tensor(unpruned_token_ids, device=device)
2657	mm_feature = MultiModalFeatureSpec(
2658	data=MultiModalKwargsItem(
2659	{
2660	"video_grid_thw": MultiModalFieldElem(
2661	data=torch.tensor(video_grid_thw),
2662	field=MultiModalFieldConfig.batched("video").field,
2663	),
2664	}
2665	),
2666	modality="video",
2667	identifier="DUMMY",
2668	mm_position=PlaceholderRange(offset=0, length=len(unpruned_token_ids)),
2669	)
2670	original_mrope_cpu = self.get_mrope_input_positions(
2671	input_tokens=unpruned_token_ids,
2672	mm_features=[mm_feature],
2673	)[0]
2674	original_mrope = async_tensor_h2d(original_mrope_cpu, device).permute(1, 0)
2675	full_is_video_embed = unpruned_token_ids_tensor == embed_token_id
2676
2677	with gpu_sync_allowed():
2678	expanded_positions[is_video_embed, :3] = original_mrope[
2679	full_is_video_embed
2680	][retention_mask]
2681	expanded_positions[~is_video_embed, :3] = original_mrope[
2682	~full_is_video_embed
2683	]
2684	expanded_positions[..., 3] = is_vision_start
2685	expanded_positions[..., 4] = is_video_embed
2686
2687	return expanded_positions
2688
2689	def _parse_and_validate_multimodal_inputs(self, **kwargs: object) -> dict:
2690	mm_input_by_modality: dict[
2691	str, Qwen2_5_VLImageInputs \| Qwen2_5_VLVideoInputs \| None
2692	] = {}
2693	for input_key in kwargs:
2694	if (
2695	input_key in ("pixel_values", "image_embeds")
2696	and "image" not in mm_input_by_modality
2697	):
2698	mm_input_by_modality["image"] = self._parse_and_validate_image_input(
2699	**kwargs
2700	)
2701	if (
2702	input_key in ("pixel_values_videos", "video_embeds")
2703	and "video" not in mm_input_by_modality
2704	):
2705	mm_input_by_modality["video"] = self._parse_and_validate_video_input(
2706	**kwargs
2707	)
2708	return mm_input_by_modality
2709
2710	@staticmethod
2711	def _iter_mm_grid_hw(
2712	input_tokens: list[int],
2713	mm_features: list[MultiModalFeatureSpec],
2714	video_token_id: int,
2715	vision_start_token_id: int,
2716	vision_end_token_id: int,
2717	spatial_merge_size: int,
2718	) -> Iterator[tuple[int, int, int, int]]:
2719	"""Iterate over multimodal features and yield position info.
2720
2721	Args:
2722	input_tokens: List of token IDs in the input sequence.
2723	mm_features: List of multimodal feature specifications containing
2724	image/video data and position information.
2725	video_token_id: Token ID used for video tokens.
2726	vision_start_token_id: Token ID marking the start of a vision sequence.
2727	vision_end_token_id: Token ID marking the end of a vision sequence.
2728	spatial_merge_size: Size of the spatial merge operation used to
2729	compute logical grid dimensions from the original feature grid.
2730
2731	Yields:
2732	offset: Position of the first video/image token in the sequence.
2733	llm_grid_h: Logical grid height (may not match actual token count with EVS).
2734	llm_grid_w: Logical grid width (may not match actual token count with EVS).
2735	actual_num_tokens: Actual number of video/image tokens in the placeholder.
2736
2737	"""
2738	for mm_feature in sorted(mm_features, key=lambda f: f.mm_position.offset):
2739	offset = mm_feature.mm_position.offset
2740	data = mm_feature.data
2741	assert data is not None
2742	if mm_feature.modality == "image":
2743	grid = data["image_grid_thw"].data
2744	assert isinstance(grid, torch.Tensor)
2745	t, h, w = grid.tolist()
2746	assert t == 1, f"Image must have 1 frame, got {t}"
2747	llm_grid_h = h // spatial_merge_size
2748	llm_grid_w = w // spatial_merge_size
2749	yield offset, llm_grid_h, llm_grid_w, llm_grid_h * llm_grid_w
2750	elif mm_feature.modality == "video":
2751	grid = data["video_grid_thw"].data
2752	assert isinstance(grid, torch.Tensor)
2753	t, h, w = grid.tolist()
2754	llm_grid_h = h // spatial_merge_size
2755	llm_grid_w = w // spatial_merge_size
2756
2757	for _ in range(t):
2758	# When EVS is enabled, some frames may have 0 video tokens in the
2759	# placeholder. We use `vision_start_token_id` to locate each frame
2760	# since it is always present for every frame.
2761	# We then look for the first `video_token_id` after
2762	# `vision_start_token_id` and before `vision_end_token_id`.
2763	offset = input_tokens.index(vision_start_token_id, offset)
2764	vision_end_offset = input_tokens.index(vision_end_token_id, offset)
2765
2766	try:
2767	actual_num_tokens = 0
2768	video_offset = input_tokens.index(
2769	video_token_id, offset, vision_end_offset
2770	)
2771	# NOTE: looking at the
2772	# `Qwen3VLMultiModalProcessor.get_video_repl` code, we can
2773	# see that we can use the below formula to get the token
2774	# count, since everything in between `video_offset` and
2775	# `vision_end_offset` is populated as `video_token_id`.
2776	# This saves us from manually counting the number tokens
2777	# that match `video_token_id` in between.
2778	actual_num_tokens += vision_end_offset - video_offset
2779	except ValueError:
2780	# No `video_token_id` in this frame (EVS with 0 tokens for
2781	# this frame) -> use `offset + 1`` to move past
2782	# `vision_start_token_id`.
2783	video_offset = offset + 1
2784
2785	yield video_offset, llm_grid_h, llm_grid_w, actual_num_tokens
2786	# Move offset past this frame for next iteration.
2787	offset = vision_end_offset + 1
2788	else:
2789	raise ValueError(f"Unsupported modality: {mm_feature.modality}")
2790
2791	def get_mrope_input_positions(
2792	self,
2793	input_tokens: list[int],
2794	mm_features: list[MultiModalFeatureSpec],
2795	) -> tuple[torch.Tensor, int]:
2796	return self._get_mrope_input_positions(
2797	input_tokens=input_tokens,
2798	mm_features=mm_features,
2799	config=self.config,
2800	)
2801
2802	@staticmethod
2803	def _get_mrope_input_positions(
2804	input_tokens: list[int],
2805	mm_features: list[MultiModalFeatureSpec],
2806	config: Qwen3VLConfig,
2807	):
2808	llm_pos_ids_list: list[np.ndarray] = []
2809	st = 0
2810	for (
2811	offset,
2812	llm_grid_h,
2813	llm_grid_w,
2814	actual_num_tokens,
2815	) in Qwen3VLForConditionalGeneration._iter_mm_grid_hw(
2816	input_tokens,
2817	mm_features,
2818	video_token_id=config.video_token_id,
2819	vision_start_token_id=config.vision_start_token_id,
2820	vision_end_token_id=config.vision_end_token_id,
2821	spatial_merge_size=config.vision_config.spatial_merge_size,
2822	):
2823	# Skip frames with 0 tokens (EVS placeholder with tokens lumped elsewhere)
2824	if actual_num_tokens == 0:
2825	continue
2826
2827	text_len = offset - st
2828	st_idx = llm_pos_ids_list[-1].max() + 1 if len(llm_pos_ids_list) > 0 else 0
2829	llm_pos_ids_list.append(
2830	np.broadcast_to(np.arange(text_len), (3, text_len)) + st_idx
2831	)
2832
2833	# Check if this is a "lumped placeholder" (all tokens from multiple frames
2834	# assigned to the 0-th frame - see
2835	# `Qwen3VLMultiModalProcessor.get_video_repl`.
2836	expected_tokens_per_frame = llm_grid_h * llm_grid_w
2837	if actual_num_tokens > expected_tokens_per_frame:
2838	# Lumped placeholder: create grid positions for all "logical" frames
2839	# represented.
2840	num_logical_frames = actual_num_tokens // expected_tokens_per_frame
2841	remainder = actual_num_tokens % expected_tokens_per_frame
2842
2843	# Create positions for complete frames.
2844	for _ in range(num_logical_frames):
2845	grid_indices = np.indices((1, llm_grid_h, llm_grid_w)).reshape(
2846	3, -1
2847	)
2848	llm_pos_ids_list.append(grid_indices + text_len + st_idx)
2849	st_idx = llm_pos_ids_list[-1].max() + 1
2850	text_len = 0 # No text between frames within the lump
2851
2852	# Handle remainder tokens if any (partial frame).
2853	# NOTE: this should never be the case. Should we have an assert?
2854	if remainder > 0:
2855	# Create a partial grid - take first 'remainder' positions
2856	full_grid = np.indices((1, llm_grid_h, llm_grid_w)).reshape(3, -1)
2857	grid_indices = full_grid[:, :remainder]
2858	llm_pos_ids_list.append(grid_indices + text_len + st_idx)
2859	else:
2860	# Normal case: frame has exactly the expected tokens (after actual EVS
2861	# pruning).
2862	grid_indices = np.indices((1, llm_grid_h, llm_grid_w)).reshape(3, -1)
2863	llm_pos_ids_list.append(grid_indices + text_len + st_idx)
2864
2865	st = offset + actual_num_tokens
2866
2867	if st < len(input_tokens):
2868	st_idx = llm_pos_ids_list[-1].max() + 1 if len(llm_pos_ids_list) > 0 else 0
2869	text_len = len(input_tokens) - st
2870	llm_pos_ids_list.append(
2871	np.broadcast_to(np.arange(text_len), (3, text_len)) + st_idx
2872	)
2873
2874	llm_positions = np.concatenate(llm_pos_ids_list, axis=1).reshape(3, -1)
2875	mrope_position_delta = (llm_positions.max() + 1 - len(input_tokens)).item()
2876	return torch.from_numpy(llm_positions), mrope_position_delta
2877
2878	def recompute_mrope_positions(
2879	self,
2880	input_ids: list[int] \| torch.Tensor,
2881	multimodal_embeddings: Sequence[torch.Tensor],
2882	mrope_positions: torch.LongTensor,
2883	num_computed_tokens: int,
2884	) -> tuple[Sequence[torch.Tensor], torch.Tensor, int]:
2885	"""Update part of input mrope positions (starting with
2886	num_computed_tokens index). Original mrope_positions are computed
2887	for unpruned sequence and becomes incorrect once pruning occurs,
2888	so once we prune media tokens we should reflect this in the
2889	mrope_positions before we feed it to LLM.
2890
2891	Args:
2892	input_ids: (N,) All input tokens of the prompt containing
2893	entire sequence.
2894	multimodal_embeddings: Sequence of multimodal embeddings that
2895	fits into the prefill chunk that is being processed.
2896	mrope_positions: Existing mrope positions (3, N) for entire
2897	sequence
2898	num_computed_tokens: A number of computed tokens so far.
2899
2900	Returns:
2901	Tuple of (multimodal_embeddings, mrope_positions,
2902	mrope_position_delta).
2903
2904	"""
2905	return self._recompute_mrope_positions(
2906	input_ids=input_ids,
2907	multimodal_embeddings=multimodal_embeddings,
2908	mrope_positions=mrope_positions,
2909	num_computed_tokens=num_computed_tokens,
2910	image_token_id=self.config.image_token_id,
2911	video_token_id=self.config.video_token_id,
2912	vision_start_token_id=self.config.vision_start_token_id,
2913	)
2914
2915	@staticmethod
2916	def _recompute_mrope_positions(
2917	input_ids: list[int] \| torch.Tensor,
2918	multimodal_embeddings: Sequence[torch.Tensor],
2919	mrope_positions: torch.LongTensor,
2920	num_computed_tokens: int,
2921	vision_start_token_id: int,
2922	image_token_id: int,
2923	video_token_id: int,
2924	) -> tuple[Sequence[torch.Tensor], torch.Tensor, int]:
2925	# Device
2926	device = (
2927	multimodal_embeddings[0].device
2928	if len(multimodal_embeddings)
2929	else mrope_positions.device
2930	)
2931
2932	# Tensors. input_ids may already be a (device-side) tensor.
2933	if isinstance(input_ids, torch.Tensor):
2934	assert input_ids.device == device
2935	input_ids_t = input_ids.to(dtype=torch.long)
2936	else:
2937	input_ids_t = async_tensor_h2d(input_ids, device=device, dtype=torch.long)
2938
2939	mm_embeddings_out = []
2940	mm_embeddings_pos = []
2941	# Strip position information from embeddings (last 5 channels)
2942	# For Qwen3 VL, handle potentially empty frames (from unpacking)
2943	for mm in multimodal_embeddings:
2944	if mm.shape[0] > 0: # Only process non-empty frames
2945	mm_embeddings_out.append(mm[:, :-5])
2946	mm_embeddings_pos.append(mm[:, -5:].permute(1, 0).long())
2947	else:
2948	# Empty frame - keep as is
2949	mm_embeddings_out.append(mm)
2950	# Create empty position tensor with correct shape
2951	mm_embeddings_pos.append(
2952	torch.empty(5, 0, device=device, dtype=torch.long)
2953	)
2954
2955	with gpu_sync_allowed():
2956	positions, mrope_positions_delta = recompute_mrope_positions(
2957	input_ids_t,
2958	mm_embeddings_pos,
2959	mrope_positions,
2960	num_computed_tokens,
2961	vision_start_token_id,
2962	image_token_id,
2963	video_token_id,
2964	)
2965
2966	return mm_embeddings_out, positions, mrope_positions_delta
2967
2968	def embed_multimodal(self, **kwargs: object) -> MultiModalEmbeddings \| None:
2969	mm_input_by_modality = self._parse_and_validate_multimodal_inputs(**kwargs)
2970	if not mm_input_by_modality:
2971	return None
2972
2973	# The result multimodal_embeddings is tuple of tensors, with each
2974	# tensor corresponding to a multimodal data item (image or video).
2975	multimodal_embeddings: list[torch.Tensor] = []
2976
2977	# NOTE: It is important to iterate over the keys in this dictionary
2978	# to preserve the order of the modalities.
2979	for modality in mm_input_by_modality:
2980	multimodal_input = mm_input_by_modality[modality]
2981	if modality == "image":
2982	image_embeddings = self._process_image_input(multimodal_input)
2983	image_embeddings = self._postprocess_image_embeds_evs(
2984	image_embeddings, multimodal_input
2985	)
2986	multimodal_embeddings.extend(image_embeddings)
2987	if modality == "video":
2988	video_embeddings = self._process_video_input(multimodal_input)
2989	if self.is_multimodal_pruning_enabled:
2990	video_embeddings = self._postprocess_video_embeds_evs(
2991	video_embeddings, multimodal_input
2992	)
2993	multimodal_embeddings.extend(video_embeddings)
2994
2995	embeddings_tuple = tuple(multimodal_embeddings)
2996	return embeddings_tuple
2997
2998	def _compute_deepstack_embeds(
2999	self,
3000	inputs_embeds: torch.Tensor,
3001	multimodal_embeddings: MultiModalEmbeddings,
3002	is_multimodal: torch.Tensor,
3003	) -> tuple[torch.Tensor, MultiModalEmbeddings]:
3004	visual_lens = [len(x) for x in multimodal_embeddings]
3005	multimodal_embeddings_cat = torch.cat(multimodal_embeddings, dim=0)
3006
3007	(
3008	multimodal_embeddings_main,
3009	multimodal_embeddings_multiscale,
3010	) = torch.split(
3011	multimodal_embeddings_cat,
3012	[self.visual_dim, self.multiscale_dim],
3013	dim=-1,
3014	)
3015
3016	multimodal_embeddings = torch.split(
3017	multimodal_embeddings_main, visual_lens, dim=0
3018	)
3019	multimodal_embeddings_multiscale = torch.split(
3020	multimodal_embeddings_multiscale, visual_lens, dim=0
3021	)
3022
3023	deepstack_input_embeds = inputs_embeds.new_zeros(
3024	inputs_embeds.size(0), self.deepstack_num_level * inputs_embeds.size(1)
3025	)
3026
3027	deepstack_input_embeds = _merge_multimodal_embeddings(
3028	inputs_embeds=deepstack_input_embeds,
3029	multimodal_embeddings=multimodal_embeddings_multiscale,
3030	is_multimodal=is_multimodal,
3031	)
3032	deepstack_input_embeds = deepstack_input_embeds.view(
3033	inputs_embeds.shape[0], self.deepstack_num_level, self.visual_dim
3034	)
3035	deepstack_input_embeds = deepstack_input_embeds.permute(1, 0, 2)
3036
3037	return deepstack_input_embeds, multimodal_embeddings
3038
3039	def embed_input_ids(
3040	self,
3041	input_ids: torch.Tensor,
3042	multimodal_embeddings: MultiModalEmbeddings \| None = None,
3043	*,
3044	is_multimodal: torch.Tensor \| None = None,
3045	) -> torch.Tensor:
3046	inputs_embeds = self._embed_text_input_ids(
3047	input_ids,
3048	self.language_model.embed_input_ids,
3049	is_multimodal=is_multimodal,
3050	)
3051
3052	if multimodal_embeddings is None or len(multimodal_embeddings) == 0:
3053	return inputs_embeds
3054
3055	is_multimodal = _require_is_multimodal(is_multimodal)
3056
3057	if self.use_deepstack:
3058	(
3059	deepstack_input_embeds,
3060	multimodal_embeddings,
3061	) = self._compute_deepstack_embeds(
3062	inputs_embeds=inputs_embeds,
3063	multimodal_embeddings=multimodal_embeddings,
3064	is_multimodal=is_multimodal,
3065	)
3066	else:
3067	deepstack_input_embeds = None
3068
3069	inputs_embeds = _merge_multimodal_embeddings(
3070	inputs_embeds=inputs_embeds,
3071	multimodal_embeddings=multimodal_embeddings,
3072	is_multimodal=is_multimodal,
3073	)
3074
3075	if deepstack_input_embeds is not None:
3076	self._set_deepstack_input_embeds(deepstack_input_embeds)
3077
3078	return inputs_embeds
3079
3080	def forward(
3081	self,
3082	input_ids: torch.Tensor \| None,
3083	positions: torch.Tensor,
3084	intermediate_tensors: IntermediateTensors \| None = None,
3085	inputs_embeds: torch.Tensor \| None = None,
3086	**kwargs: object,
3087	) -> torch.Tensor \| IntermediateTensors:
3088	"""Run forward pass for Qwen3VL.
3089
3090	Args:
3091	input_ids: Flattened (concatenated) input_ids corresponding to a
3092	batch.
3093	positions: Flattened (concatenated) position ids corresponding to a
3094	batch.
3095	NOTE: If mrope is enabled (default setting for Qwen3VL
3096	opensource models), the shape will be `(3, seq_len)`,
3097	otherwise it will be `(seq_len,).
3098	intermediate_tensors: Intermediate tensors from previous pipeline
3099	stages.
3100	inputs_embeds: Pre-computed input embeddings.
3101	**kwargs: Additional keyword arguments including:
3102	- pixel_values: Pixel values to be fed to a model.
3103	`None` if no images are passed.
3104	- image_grid_thw: Tensor `(n_images, 3)` of image 3D grid in
3105	LLM. `None` if no images are passed.
3106	- pixel_values_videos: Pixel values of videos to be fed to a
3107	model. `None` if no videos are passed.
3108	- video_grid_thw: Tensor `(n_videos, 3)` of video 3D grid in
3109	LLM. `None` if no videos are passed.
3110
3111	"""
3112	if intermediate_tensors is not None:
3113	inputs_embeds = None
3114
3115	if inputs_embeds is not None and get_pp_group().is_first_rank:
3116	deepstack_input_embeds = self._get_deepstack_input_embeds(
3117	inputs_embeds.size(0)
3118	)
3119	else:
3120	deepstack_input_embeds = None
3121
3122	hidden_states = self.language_model.model(
3123	input_ids=input_ids,
3124	positions=positions,
3125	intermediate_tensors=intermediate_tensors,
3126	inputs_embeds=inputs_embeds,
3127	# args for deepstack
3128	deepstack_input_embeds=deepstack_input_embeds,
3129	)
3130
3131	if inputs_embeds is not None and get_pp_group().is_first_rank:
3132	self._clear_deepstack_input_embeds(inputs_embeds.size(0))
3133
3134	return hidden_states
3135
3136	def compute_logits(
3137	self,
3138	hidden_states: torch.Tensor,
3139	) -> torch.Tensor \| None:
3140	return self.language_model.compute_logits(hidden_states)
3141
3142	def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
3143	loader = AutoWeightsLoader(self)
3144	return loader.load_weights(weights, mapper=self.hf_to_vllm_mapper)
3145
3146	def get_mm_mapping(self) -> MultiModelKeys:
3147	"""Get the module prefix in multimodal models."""
3148	return MultiModelKeys.from_string_field(
3149	language_model="language_model",
3150	connector=["visual.merger", "visual.deepstack_merger_list"],
3151	tower_model="visual.",
3152	)
3153
3154	def get_mm_lora_token_counts(
3155	self,
3156	*,
3157	modality: str,
3158	mm_kwargs: MultiModalKwargsItem \| None,
3159	num_mm_embeds: int,
3160	) -> tuple[int, int \| None]:
3161	del modality, mm_kwargs
3162	hf_config = self.config
3163	vision_config = hf_config.vision_config
3164	merge_size = vision_config.spatial_merge_size
3165	return num_mm_embeds * merge_size**2, num_mm_embeds
3166
3167
3168	@lru_cache
3169	def _cached_tensor(x, device) -> torch.Tensor:
3170	return torch.tensor(x, device=device)