main
py 572 lines 22.9 KB
Raw
1 # SPDX-License-Identifier: Apache-2.0
2 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3 """MTP draft model for DeepSeek V4 (internal codename: DeepseekV4).
4
5 Split from ``deepseek_mtp.py`` because the V4 architecture introduces several
6 pieces that have no analogue in V3/V32:
7 * separate ``e_proj`` / ``h_proj`` with fp8 linear quantization (instead of
8 the fused ``eh_proj``);
9 * ``hc_head`` hypercompressed vocab projection applied in ``compute_logits``;
10 * ``DeepseekV4DecoderLayer`` with its own aux-stream management;
11 * V4-specific checkpoint weight-name remapping in ``load_weights``.
12 """
13
14 import typing
15 from collections.abc import Callable, Iterable
16
17 import regex as re
18 import torch
19 import torch.nn as nn
20
21 import vllm.envs as envs
22 from vllm.config import VllmConfig
23 from vllm.config.kernel import MEGA_MOE_BACKENDS
24 from vllm.distributed import (
25 get_tensor_model_parallel_rank,
26 get_tensor_model_parallel_world_size,
27 )
28 from vllm.forward_context import get_forward_context, is_forward_context_available
29 from vllm.logger import init_logger
30 from vllm.model_executor.kernels.mhc.tilelang import (
31 hc_head_fused_kernel_tilelang,
32 mhc_post_tilelang,
33 )
34 from vllm.model_executor.layers.fused_moe import (
35 fused_moe_make_expert_params_mapping,
36 )
37 from vllm.model_executor.layers.layernorm import RMSNorm
38 from vllm.model_executor.layers.linear import ReplicatedLinear
39 from vllm.model_executor.layers.logits_processor import LogitsProcessor
40 from vllm.model_executor.layers.vocab_parallel_embedding import (
41 VocabParallelEmbedding,
42 )
43 from vllm.model_executor.model_loader.mtp_validation import (
44 is_mtp_completeness_check_enabled,
45 )
46 from vllm.model_executor.model_loader.weight_utils import default_weight_loader
47 from vllm.model_executor.models.deepseek_mtp import SharedHead
48 from vllm.model_executor.models.deepseek_v2 import get_spec_layer_idx_from_weight_name
49 from vllm.model_executor.models.utils import maybe_prefix
50 from vllm.models.common.ops.sequence_parallel import (
51 sp_all_gather,
52 sp_padding_mask,
53 sp_shard,
54 )
55 from vllm.models.deepseek_v4.common.mm_preprocess import IMAGE_SENTINEL_BASE_ID
56 from vllm.models.deepseek_v4.common.ops.fused_mtp_input_rmsnorm import (
57 _FUSED_MTP_INPUT_RMSNORM_KERNEL,
58 _MTP_SHARED_HEAD_RMSNORM_KERNEL,
59 )
60 from vllm.sequence import IntermediateTensors
61
62 from .model import (
63 DeepseekV4DecoderLayer,
64 DeepseekV4Model,
65 _use_sequence_parallel,
66 make_deepseek_v4_expert_params_mapping,
67 prepare_mega_gate_routing_metadata,
68 )
69
70 logger = init_logger(__name__)
71
72 # MoE expert scales are fused into per-layer w13/w2 tensors. The exact
73 # parameter suffix depends on which FusedMoEFactory method handles the experts:
74 # - fp4 experts (Mxfp4MoEMethod) register ``w{1,2,3}_weight_scale``;
75 # - fp8 experts (Fp8MoEMethod with block_quant=True) register
76 # ``w{1,2,3}_weight_scale_inv``.
77 # Other FP8 linear scales (including shared experts) always use
78 # ``.weight_scale_inv``. Mirrors the per-instance mapper built by
79 # ``_make_deepseek_v4_weights_mapper`` in deepseek_v4.py.
80 _EXPERT_SCALE_RE = re.compile(r"\.experts\.\d+\.w[123]\.scale$")
81
82
83 class DeepSeekV4MultiTokenPredictorLayer(nn.Module):
84 def __init__(
85 self,
86 vllm_config: VllmConfig,
87 topk_indices_buffer: torch.Tensor,
88 prefix: str,
89 aux_stream_list: list[torch.cuda.Stream] | None = None,
90 ) -> None:
91 super().__init__()
92
93 assert vllm_config.speculative_config is not None
94 config = vllm_config.speculative_config.draft_model_config.hf_config
95 self.config = config
96 quant_config = vllm_config.quant_config
97 self.rms_norm_eps = config.rms_norm_eps
98 self.use_mega_moe = vllm_config.kernel_config.moe_backend in MEGA_MOE_BACKENDS
99
100 self.enorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
101 self.hnorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
102
103 # V4 keeps e_ and h_ proj separate (with fp8 linear quant) rather than
104 # fusing them the way V3 does with eh_proj.
105 self.e_proj = ReplicatedLinear(
106 config.hidden_size,
107 config.hidden_size,
108 bias=False,
109 return_bias=False,
110 quant_config=quant_config,
111 prefix=f"{prefix}.e_proj",
112 )
113 self.h_proj = ReplicatedLinear(
114 config.hidden_size,
115 config.hidden_size,
116 bias=False,
117 return_bias=False,
118 quant_config=quant_config,
119 prefix=f"{prefix}.h_proj",
120 )
121
122 self.hc_eps = config.hc_eps
123 self.hc_mult = config.hc_mult
124 self.hc_dim = self.hc_mult * config.hidden_size
125 self.hc_head_fn = nn.Parameter(
126 torch.empty(self.hc_mult, self.hc_dim, dtype=torch.float32),
127 requires_grad=False,
128 )
129 self.hc_head_base = nn.Parameter(
130 torch.empty(self.hc_mult, dtype=torch.float32),
131 requires_grad=False,
132 )
133 self.hc_head_scale = nn.Parameter(
134 torch.empty(1, dtype=torch.float32),
135 requires_grad=False,
136 )
137
138 self.shared_head = SharedHead(
139 config=config, prefix=prefix, quant_config=quant_config
140 )
141 self.mtp_block = DeepseekV4DecoderLayer(
142 vllm_config,
143 prefix,
144 topk_indices_buffer=topk_indices_buffer,
145 aux_stream_list=aux_stream_list,
146 )
147
148 if vllm_config.kernel_config.enable_jit_warmup:
149 _FUSED_MTP_INPUT_RMSNORM_KERNEL.register_warmup()
150 _MTP_SHARED_HEAD_RMSNORM_KERNEL.register_warmup()
151
152 def forward(
153 self,
154 input_ids: torch.Tensor,
155 positions: torch.Tensor,
156 previous_hidden_states: torch.Tensor,
157 inputs_embeds: torch.Tensor | None = None,
158 spec_step_index: int = 0,
159 ) -> torch.Tensor:
160 assert inputs_embeds is not None
161 # Target stashes pre-hc_head residual as flat (T, hc_mult * D);
162 # reshape to (T, hc_mult, D) — the training-time layout — before
163 # the fused norm pass so both inputs are 3D-friendly.
164 previous_hidden_states = previous_hidden_states.view(
165 -1, self.hc_mult, self.config.hidden_size
166 )
167 # Fused: mask inputs at position 0 (not needed by MTP), enorm, hnorm.
168 inputs_embeds, previous_hidden_states = _FUSED_MTP_INPUT_RMSNORM_KERNEL(
169 inputs_embeds,
170 positions,
171 previous_hidden_states,
172 self.enorm.weight.data,
173 self.hnorm.weight.data,
174 self.enorm.variance_epsilon,
175 self.hc_mult,
176 )
177 if self.mtp_block.use_sequence_parallel:
178 if envs.VLLM_MOE_SKIP_PADDING and is_forward_context_available():
179 forward_context = get_forward_context()
180 forward_context.is_padding = sp_padding_mask(
181 forward_context.is_padding, inputs_embeds
182 )
183 inputs_embeds = sp_shard(inputs_embeds)
184 previous_hidden_states = sp_shard(previous_hidden_states)
185 hidden_states = self.h_proj(previous_hidden_states) + self.e_proj(
186 inputs_embeds
187 ).unsqueeze(-2)
188 mega_gate_metadata = None
189 if self.use_mega_moe:
190 routing_input_ids = input_ids
191 if self.mtp_block.use_sequence_parallel:
192 routing_input_ids = sp_shard(routing_input_ids)
193 mega_gate_metadata = prepare_mega_gate_routing_metadata(
194 routing_input_ids,
195 has_hash_routing=False,
196 image_sentinel_base_id=IMAGE_SENTINEL_BASE_ID
197 if getattr(self.config, "vision_n_layers", 0) > 0
198 else None,
199 )
200 hidden_states, residual, post_mix, res_mix = self.mtp_block(
201 positions=positions,
202 x=hidden_states,
203 input_ids=input_ids,
204 mega_gate_metadata=mega_gate_metadata,
205 )
206 hidden_states = mhc_post_tilelang(hidden_states, residual, post_mix, res_mix)
207 if self.mtp_block.use_sequence_parallel:
208 hidden_states = sp_all_gather(hidden_states)[: positions.shape[0]]
209 # Return the flat pre-hc_head residual so it can be re-fed as the
210 # next spec step's `previous_hidden_states` when
211 # num_speculative_tokens > 1. hc_head is deferred to compute_logits.
212 return hidden_states.flatten(1)
213
214
215 class DeepSeekV4MultiTokenPredictor(nn.Module):
216 def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
217 super().__init__()
218 config = vllm_config.model_config.hf_config
219 self.mtp_start_layer_idx = config.num_hidden_layers
220 self.num_mtp_layers = config.num_nextn_predict_layers
221
222 topk_tokens = config.index_topk
223 self.topk_indices_buffer = torch.empty(
224 vllm_config.scheduler_config.max_num_batched_tokens,
225 topk_tokens,
226 dtype=torch.int32,
227 )
228
229 # Three aux streams shared across all MTP layers, mirroring DeepseekV4Model.
230 aux_stream_list = [torch.cuda.Stream() for _ in range(3)]
231
232 # to map the exact layer index from weights
233 self.layers = torch.nn.ModuleDict(
234 {
235 str(idx): DeepSeekV4MultiTokenPredictorLayer(
236 vllm_config,
237 self.topk_indices_buffer,
238 f"{prefix}.layers.{idx}",
239 aux_stream_list=aux_stream_list,
240 )
241 for idx in range(
242 self.mtp_start_layer_idx,
243 self.mtp_start_layer_idx + self.num_mtp_layers,
244 )
245 }
246 )
247 self.embed_tokens = VocabParallelEmbedding(
248 config.vocab_size,
249 config.hidden_size,
250 prefix=maybe_prefix(prefix, "embed_tokens"),
251 )
252 self.logits_processor = LogitsProcessor(config.vocab_size)
253
254 def embed_input_ids(self, input_ids: torch.Tensor) -> torch.Tensor:
255 return self.embed_tokens(input_ids)
256
257 def forward(
258 self,
259 input_ids: torch.Tensor,
260 positions: torch.Tensor,
261 previous_hidden_states: torch.Tensor,
262 inputs_embeds: torch.Tensor | None = None,
263 spec_step_idx: int = 0,
264 ) -> torch.Tensor:
265 if inputs_embeds is None:
266 inputs_embeds = self.embed_tokens(input_ids)
267 current_step_idx = spec_step_idx % self.num_mtp_layers
268 return self.layers[str(self.mtp_start_layer_idx + current_step_idx)](
269 input_ids,
270 positions,
271 previous_hidden_states,
272 inputs_embeds,
273 current_step_idx,
274 )
275
276 def compute_logits(
277 self,
278 hidden_states: torch.Tensor,
279 spec_step_idx: int = 0,
280 ) -> torch.Tensor:
281 current_step_idx = spec_step_idx % self.num_mtp_layers
282 mtp_layer = self.layers[str(self.mtp_start_layer_idx + current_step_idx)]
283 # MTP forward returns the pre-hc_head residual (T, hc_mult * D); apply
284 # hc_head here so logits are computed from the dense hidden state.
285 hidden_states = hidden_states.view(
286 -1, mtp_layer.hc_mult, mtp_layer.config.hidden_size
287 )
288 hidden_states = hc_head_fused_kernel_tilelang(
289 hidden_states,
290 mtp_layer.hc_head_fn,
291 mtp_layer.hc_head_scale,
292 mtp_layer.hc_head_base,
293 mtp_layer.rms_norm_eps,
294 mtp_layer.hc_eps,
295 )
296 hidden_states = _MTP_SHARED_HEAD_RMSNORM_KERNEL(
297 hidden_states,
298 mtp_layer.shared_head.norm.weight.data,
299 mtp_layer.shared_head.norm.variance_epsilon,
300 )
301 logits = self.logits_processor(mtp_layer.shared_head.head, hidden_states)
302 return logits
303
304
305 class DeepSeekV4MTP(nn.Module):
306 def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
307 super().__init__()
308 self.config = vllm_config.model_config.hf_config
309 self.quant_config = vllm_config.quant_config
310 self.pad_shared_expert = getattr(
311 self.quant_config, "weight_block_size", None
312 ) is not None and not _use_sequence_parallel(vllm_config)
313 self.model = DeepSeekV4MultiTokenPredictor(
314 vllm_config=vllm_config, prefix=maybe_prefix(prefix, "model")
315 )
316
317 def embed_input_ids(self, input_ids: torch.Tensor) -> torch.Tensor:
318 return self.model.embed_input_ids(input_ids)
319
320 def forward(
321 self,
322 input_ids: torch.Tensor | None,
323 positions: torch.Tensor,
324 hidden_states: torch.Tensor,
325 intermediate_tensors: IntermediateTensors | None = None,
326 inputs_embeds: torch.Tensor | None = None,
327 spec_step_idx: int = 0,
328 ) -> torch.Tensor:
329 hidden_states = self.model(
330 input_ids, positions, hidden_states, inputs_embeds, spec_step_idx
331 )
332 return hidden_states
333
334 def compute_logits(
335 self,
336 hidden_states: torch.Tensor,
337 spec_step_idx: int = 0,
338 ) -> torch.Tensor | None:
339 return self.model.compute_logits(hidden_states, spec_step_idx)
340
341 def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
342 # Weight name remapping for checkpoint compatibility.
343 # Maps checkpoint weight paths to model parameter paths.
344 WEIGHT_NAME_REMAPPING: dict[str, str] = {
345 ".emb.tok_emb.weight": ".embed_tokens.weight",
346 ".head.weight": ".shared_head.head.weight",
347 ".norm.weight": ".shared_head.norm.weight",
348 }
349
350 def _remap_weight_name(name: str) -> str:
351 """Remap checkpoint weight names to model parameter names."""
352 for old_pattern, new_pattern in WEIGHT_NAME_REMAPPING.items():
353 if old_pattern in name:
354 name = name.replace(old_pattern, new_pattern)
355 return name
356
357 def _find_mtp_layer_idx(name: str) -> int:
358 subnames = name.split(".")
359 for subname in subnames:
360 try:
361 # we return the first encountered integer
362 return int(subname)
363 except ValueError:
364 continue
365 return 0
366
367 stacked_params_mapping = [
368 # (param_name, shard_name, shard_id)
369 ("gate_up_proj", "w1", 0),
370 ("gate_up_proj", "w3", 1),
371 ("attn.fused_wqa_wkv", "attn.wq_a", 0),
372 ("attn.fused_wqa_wkv", "attn.wkv", 1),
373 ]
374 params_dict = dict(self.named_parameters())
375 loaded_params: set[str] = set()
376
377 # TP for attention
378 tp_size = get_tensor_model_parallel_world_size()
379 tp_rank = get_tensor_model_parallel_rank()
380 n_head = self.config.num_attention_heads
381 n_local_head = n_head // tp_size
382 head_rank_start = n_local_head * tp_rank
383 head_rank_end = n_local_head * (tp_rank + 1)
384
385 # Pre-compute expert mapping ONCE.
386 first_layer = next(iter(self.model.layers.values()))
387 if first_layer.mtp_block.ffn.use_mega_moe:
388 expert_mapping = make_deepseek_v4_expert_params_mapping(
389 self.config.n_routed_experts
390 )
391 else:
392 expert_mapping = fused_moe_make_expert_params_mapping(
393 self,
394 ckpt_gate_proj_name="w1",
395 ckpt_down_proj_name="w2",
396 ckpt_up_proj_name="w3",
397 num_experts=self.config.n_routed_experts,
398 )
399
400 # FP8 experts register ``..._weight_scale_inv`` (block_quant) while
401 # FP4/MXFP4 experts register ``..._weight_scale``. Choose the suffix
402 # for the rename below based on the model's expert dtype.
403 expert_scale_suffix = (
404 ".weight_scale"
405 if getattr(self.config, "expert_dtype", "fp4") == "fp4"
406 else ".weight_scale_inv"
407 )
408
409 for name, loaded_weight in weights:
410 mtp_layer_idx = _find_mtp_layer_idx(name)
411 # V4 checkpoints store MTP weights as `mtp.{i}.*`; remap to
412 # `model.layers.{num_hidden_layers + i}.*` so that
413 # get_spec_layer_idx_from_weight_name can identify them.
414 name = name.replace(
415 f"mtp.{mtp_layer_idx}.",
416 f"model.layers.{self.config.num_hidden_layers + mtp_layer_idx}.",
417 )
418
419 spec_layer = get_spec_layer_idx_from_weight_name(self.config, name)
420 if spec_layer is None:
421 continue
422
423 name = _remap_weight_name(name)
424 name = self._rewrite_spec_layer_name(spec_layer, name)
425
426 if spec_layer != self.model.mtp_start_layer_idx and ".layers" not in name:
427 continue
428 if name.endswith(".scale"):
429 suffix = (
430 expert_scale_suffix
431 if _EXPERT_SCALE_RE.search(name)
432 else ".weight_scale_inv"
433 )
434 name = name.removesuffix(".scale") + suffix
435 if ".shared_experts.w2" in name:
436 name = name.replace(".shared_experts.w2", ".shared_experts.down_proj")
437 if self.pad_shared_expert and ".shared_experts." in name:
438 loaded_weight = DeepseekV4Model._pad_shared_expert_weight(
439 self.quant_config, name, loaded_weight
440 )
441 for param_name, weight_name, shard_id in stacked_params_mapping:
442 # Skip non-stacked layers and experts (experts handled below).
443 if ".experts." in name:
444 continue
445 if weight_name not in name:
446 continue
447 name = name.replace(weight_name, param_name)
448
449 param = params_dict[name]
450 weight_loader = param.weight_loader
451 weight_loader(param, loaded_weight, shard_id)
452 loaded_params.add(name)
453 break
454 else:
455 if ".experts." in name:
456 # Reinterpret E8M0 scales as uint8 to preserve raw
457 # exponent bytes; numeric copy_() would zero them.
458 # Mirrors the main DeepseekV4 loader.
459 if (
460 "weight_scale" in name
461 and loaded_weight.dtype == torch.float8_e8m0fnu
462 ):
463 loaded_weight = loaded_weight.view(torch.uint8)
464 for mapping in expert_mapping:
465 param_name, weight_name, expert_id, expert_shard_id = mapping
466 if weight_name not in name:
467 continue
468 name_mapped = name.replace(weight_name, param_name)
469 param = params_dict[name_mapped]
470 # We should ask the weight loader to return success or not
471 # here since otherwise we may skip experts with other
472 # available replicas.
473 weight_loader = typing.cast(
474 Callable[..., bool], param.weight_loader
475 )
476 success = weight_loader(
477 param,
478 loaded_weight,
479 name_mapped,
480 shard_id=expert_shard_id,
481 expert_id=expert_id,
482 return_success=True,
483 )
484 if success:
485 name = name_mapped
486 loaded_params.add(name_mapped)
487 break
488 continue
489 elif "attn_sink" in name:
490 narrow_weight = loaded_weight[head_rank_start:head_rank_end]
491 n = narrow_weight.shape[0]
492 params_dict[name][:n].copy_(narrow_weight)
493 loaded_params.add(name)
494 continue
495 else:
496 if name.endswith(".ffn.gate.bias"):
497 # ``e_score_correction_bias`` lives on the gate
498 # under a different attribute name.
499 name = name.replace(
500 ".ffn.gate.bias",
501 ".ffn.gate.e_score_correction_bias",
502 )
503 param = params_dict[name]
504 weight_loader = getattr(
505 param, "weight_loader", default_weight_loader
506 )
507 weight_loader(param, loaded_weight)
508 loaded_params.add(name)
509 continue
510
511 loaded_layers: set[int] = set()
512 for param_name in loaded_params:
513 spec_layer = get_spec_layer_idx_from_weight_name(self.config, param_name)
514 if spec_layer is not None:
515 loaded_layers.add(spec_layer)
516 for layer_idx in range(
517 self.model.mtp_start_layer_idx,
518 self.model.mtp_start_layer_idx + self.model.num_mtp_layers,
519 ):
520 if layer_idx not in loaded_layers and is_mtp_completeness_check_enabled():
521 raise ValueError(
522 f"MTP speculative decoding layer {layer_idx} weights "
523 f"missing from checkpoint. The checkpoint may have "
524 f"been quantized without including the MTP layers. "
525 f"Use a checkpoint that includes MTP layer weights, "
526 f"or disable speculative decoding."
527 )
528 self.process_weights_after_loading()
529 logger.info_once("MTP draft model loaded: %d params", len(loaded_params))
530 return loaded_params
531
532 def finalize_mega_moe_weights(self) -> None:
533 for layer in self.model.layers.values():
534 layer.mtp_block.ffn.finalize_mega_moe_weights()
535
536 def process_weights_after_loading(self) -> None:
537 self.finalize_mega_moe_weights()
538
539 def _rewrite_spec_layer_name(self, spec_layer: int, name: str) -> str:
540 """Rewrite the weight name to match the format of the original model.
541 Add .mtp_block for modules in transformer layer block for spec layer
542 and rename shared layer weights to be top level.
543 """
544 spec_layer_weight_names = [
545 "embed_tokens",
546 "enorm",
547 "hnorm",
548 "h_proj",
549 "e_proj",
550 "shared_head",
551 "hc_head_fn",
552 "hc_head_base",
553 "hc_head_scale",
554 ]
555 shared_weight_names = ["embed_tokens"]
556 spec_layer_weight = False
557 shared_weight = False
558 for weight_name in spec_layer_weight_names:
559 if weight_name in name:
560 spec_layer_weight = True
561 if weight_name in shared_weight_names:
562 shared_weight = True
563 break
564 if not spec_layer_weight:
565 # treat rest weights as weights for transformer layer block
566 name = name.replace(
567 f"model.layers.{spec_layer}.", f"model.layers.{spec_layer}.mtp_block."
568 )
569 elif shared_weight:
570 # treat shared weights as top level weights
571 name = name.replace(f"model.layers.{spec_layer}.", "model.")
572 return name