main
py 524 lines 21.1 KB
Raw
1 # SPDX-License-Identifier: Apache-2.0
2 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3 """DSpark draft model for DeepSeek-V4 on ROCm/AMD (gfx950).
4
5 ROCm port of ``nvidia/dspark.py``. Follows the same nvidia->amd recipe used for
6 ``amd/mtp.py``:
7
8 * import ``DeepseekV4DecoderLayer`` from the AMD ``.model`` (aiter/triton
9 attention + MHC CustomOp path) instead of the nvidia one;
10 * route the MHC head through the ``HCHeadOp`` CustomOp dispatcher (aiter /
11 tilelang / triton / torch) instead of calling the tilelang kernels directly,
12 and gate the trailing ``mhc_post`` on ``use_fused_mhc`` (True when AITER
13 or TileLang fused MHC is available; False only on the torch fallback);
14 * drop the mega-MoE weight path (``make_deepseek_v4_expert_params_mapping`` /
15 ``use_mega_moe`` / ``finalize_mega_moe_weights`` do not exist in amd/model.py).
16
17 Everything else — the semi-autoregressive drafting hooks, the Markov head, the
18 sliding-window context-KV insert, and the checkpoint ``mtp.*`` weight remap — is
19 pure torch / Triton and shared with the nvidia implementation unchanged.
20 """
21
22 from collections.abc import Iterable
23
24 import regex as re
25 import torch
26 import torch.nn as nn
27
28 from vllm.config import VllmConfig, get_current_vllm_config
29 from vllm.distributed import (
30 get_tensor_model_parallel_rank,
31 get_tensor_model_parallel_world_size,
32 )
33 from vllm.logger import init_logger
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.mhc import HCHeadOp
41 from vllm.model_executor.layers.vocab_parallel_embedding import (
42 ParallelLMHead,
43 VocabParallelEmbedding,
44 )
45 from vllm.model_executor.model_loader.weight_utils import default_weight_loader
46 from vllm.model_executor.models.qwen3_dspark import (
47 DSparkConfidenceHead,
48 DSparkMarkovHead,
49 )
50 from vllm.model_executor.models.utils import maybe_prefix
51
52 from .model import (
53 DeepseekV4DecoderLayer,
54 )
55
56 logger = init_logger(__name__)
57
58 # MoE expert scale suffix differs by expert dtype (mirrors deepseek_v4 loaders):
59 # fp4 experts register ``.weight_scale``; block-fp8 experts ``.weight_scale_inv``.
60 _EXPERT_SCALE_RE = re.compile(r"\.experts\.\d+\.w[123]\.scale$")
61
62
63 class DSparkDeepseekV4Model(nn.Module):
64 def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None:
65 super().__init__()
66 assert vllm_config.speculative_config is not None
67 config = vllm_config.speculative_config.draft_model_config.hf_config
68 self.config = config
69 self.hidden_size = config.hidden_size
70 self.hc_mult = config.hc_mult
71 self.hc_eps = config.hc_eps
72 self.rms_norm_eps = config.rms_norm_eps
73 self.num_hidden_layers = config.num_hidden_layers
74 self.target_layer_ids = tuple(config.dspark_target_layer_ids)
75
76 self.num_dspark_layers = getattr(config, "n_mtp_layers", None) or 3
77
78 # Shared with the target (aliased by the speculator's loading utility).
79 self.embed_tokens = VocabParallelEmbedding(
80 config.vocab_size,
81 config.hidden_size,
82 prefix=maybe_prefix(prefix, "embed_tokens"),
83 )
84
85 self.main_proj = ReplicatedLinear(
86 config.hidden_size * len(self.target_layer_ids),
87 config.hidden_size,
88 bias=False,
89 return_bias=False,
90 quant_config=vllm_config.quant_config,
91 prefix=maybe_prefix(prefix, "main_proj"),
92 )
93 self.main_norm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
94
95 current_vllm_config = get_current_vllm_config()
96 self.layers = nn.ModuleList(
97 [
98 DeepseekV4DecoderLayer(
99 current_vllm_config,
100 prefix=maybe_prefix(prefix, f"layers.{self.num_hidden_layers + i}"),
101 )
102 for i in range(self.num_dspark_layers)
103 ]
104 )
105
106 # Heads: final norm + hc_head, and the Markov + confidence heads
107 # Loaded from the "final" MTP layer weights (mtp.*) in the target checkpoint
108 self.norm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
109 hc_dim = self.hc_mult * config.hidden_size
110 self.hc_head_fn = nn.Parameter(
111 torch.empty(self.hc_mult, hc_dim, dtype=torch.float32),
112 requires_grad=False,
113 )
114 self.hc_head_base = nn.Parameter(
115 torch.empty(self.hc_mult, dtype=torch.float32), requires_grad=False
116 )
117 self.hc_head_scale = nn.Parameter(
118 torch.empty(1, dtype=torch.float32), requires_grad=False
119 )
120 draft_vocab_size = (
121 getattr(config, "draft_vocab_size", None) or config.vocab_size
122 )
123 self.markov_head = DSparkMarkovHead(
124 config.vocab_size,
125 draft_vocab_size,
126 config.dspark_markov_rank,
127 prefix=maybe_prefix(prefix, "markov_head"),
128 )
129 self.confidence_head: DSparkConfidenceHead | None = None
130 if getattr(config, "enable_confidence_head", True):
131 self.confidence_head = DSparkConfidenceHead(
132 config.hidden_size + config.dspark_markov_rank,
133 prefix=maybe_prefix(prefix, "confidence_head"),
134 )
135
136 # MHC head CustomOp dispatcher (aiter / tilelang / triton / torch),
137 # replacing the direct nvidia tilelang kernel call.
138 self.hc_head_op = HCHeadOp()
139
140 def embed_input_ids(self, input_ids: torch.Tensor) -> torch.Tensor:
141 return self.embed_tokens(input_ids)
142
143 def combine_hidden_states(self, aux_hidden_states: torch.Tensor) -> torch.Tensor:
144 """main_x = main_norm(main_proj(concat of target aux hidden states)).
145
146 ``aux_hidden_states`` is [T, hidden_size * len(target_layer_ids)].
147 """
148 return self.main_norm(self.main_proj(aux_hidden_states))
149
150 @torch.inference_mode()
151 def precompute_and_store_context_kv(
152 self,
153 main_x: torch.Tensor,
154 context_positions: torch.Tensor,
155 context_slot_mappings: list[torch.Tensor | None] | None = None,
156 ) -> None:
157 """Insert the sliding-window context KV for every draft layer.
158
159 Mirrors the reference DSparkAttention: each layer derives its context KV
160 from the SAME projected target hidden ``main_x``, via that layer's own
161 ``wkv`` + ``kv_norm`` + RoPE + quant, then writes it at the
162 layer's context slots.
163
164 ``context_slot_mappings`` is a per-layer list (each entry is the context
165 slot mapping for that layer's kv-cache group, since the hybrid manager may
166 place draft layers in different groups). ``None`` (or a ``None`` entry)
167 runs the projection to reserve workspace but writes nothing (profiling).
168 """
169 for i, layer in enumerate(self.layers):
170 slot_mapping = (
171 None if context_slot_mappings is None else context_slot_mappings[i]
172 )
173 attn = layer.attn
174 # Optimized DSV4 MLA path: wkv part of the fused wq_a|wkv projection
175 # (q_lora part discarded), then RoPE/quant/insert via the fused op.
176 qr_kv, _ = attn.fused_wqa_wkv(main_x)
177 kv = qr_kv[..., attn.q_lora_rank :]
178 kv = attn.kv_norm(kv)
179 if slot_mapping is None:
180 continue
181 _insert_context_kv(attn, kv, context_positions, slot_mapping)
182
183 def forward(
184 self,
185 input_ids: torch.Tensor,
186 positions: torch.Tensor,
187 inputs_embeds: torch.Tensor | None = None,
188 ) -> torch.Tensor:
189 if inputs_embeds is None:
190 inputs_embeds = self.embed_input_ids(input_ids)
191 # Expand to hc_mult copies for hyper-connections ([T, H] -> [T, hc, H]).
192 hidden_states = inputs_embeds.unsqueeze(-2).repeat(1, self.hc_mult, 1)
193
194 residual = post_mix = res_mix = None
195 for layer in self.layers:
196 hidden_states, residual, post_mix, res_mix = layer(
197 hidden_states,
198 positions,
199 input_ids,
200 post_mix,
201 res_mix,
202 residual,
203 )
204 # On the fused-MHC path the trailing hc_post must be applied here; on the
205 # aiter unfused path (ROCm default) the decoder layer already applied
206 # hc_post in-layer and returned None mixes, so this is skipped. Mirrors
207 # amd/mtp.py.
208 last_layer = self.layers[-1]
209 if last_layer.use_fused_mhc:
210 hidden_states = last_layer.hc_post(
211 hidden_states, residual, post_mix, res_mix
212 )
213 # hc_head reduces the hc copies; return the PRE-norm head hidden.
214 hidden_states = self.hc_head_op(
215 hidden_states,
216 self.hc_head_fn,
217 self.hc_head_scale,
218 self.hc_head_base,
219 self.rms_norm_eps,
220 self.hc_eps,
221 )
222 return hidden_states
223
224
225 def _insert_context_kv(
226 attn: nn.Module,
227 kv: torch.Tensor,
228 positions: torch.Tensor,
229 slot_mapping: torch.Tensor,
230 ) -> None:
231 """RoPE + quant + paged-cache insert of (already kv_norm'd) context KV.
232
233 Reuses the DSV4 fused insert ops (which also process a query; we pass a dummy
234 query and discard it, since context tokens have no query). Mirrors
235 ``DeepseekV4Attention._fused_qnorm_rope_kv_insert``.
236 """
237 swa_cache = attn.swa_cache_layer.kv_cache
238 block_size = attn.swa_cache_layer.block_size
239 cos_sin_cache = attn.rotary_emb.cos_sin_cache
240 cache_dtype = swa_cache.dtype
241 n_ctx = kv.shape[0]
242 dummy_q = torch.zeros(
243 (n_ctx, attn.n_local_heads, attn.head_dim),
244 dtype=kv.dtype,
245 device=kv.device,
246 )
247 if cache_dtype == torch.uint8:
248 # fp8_ds_mla UE8M0 paged layout (the gfx950 aiter SWA cache path).
249 swa_2d = swa_cache.view(swa_cache.shape[0], -1)
250 torch.ops._C.fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert(
251 dummy_q,
252 kv,
253 swa_2d,
254 slot_mapping,
255 positions,
256 cos_sin_cache,
257 attn.padded_heads,
258 attn.eps,
259 block_size,
260 )
261 elif cache_dtype == torch.bfloat16:
262 swa_3d = swa_cache.view(-1, block_size, attn.head_dim)
263 torch.ops._C.fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_bf16_insert(
264 dummy_q,
265 kv,
266 swa_3d,
267 slot_mapping,
268 positions,
269 cos_sin_cache,
270 attn.eps,
271 block_size,
272 )
273 else: # per-tensor fp8 (torch.float8_e4m3fn)
274 # NOTE(rocm): unreachable on ROCm/aiter, where the SWA cache dtype is
275 # uint8 (fp8_ds_mla) or bfloat16. This branch relies on FlashInfer-only
276 # attributes (``_flashinfer_fp8_*``) that the aiter attention layer does
277 # not define; kept for parity with the nvidia path.
278 swa_3d = swa_cache.view(-1, block_size, attn.head_dim)
279 dummy_q_fp8 = torch.zeros_like(dummy_q, dtype=torch.float8_e4m3fn)
280 torch.ops._C.fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert(
281 dummy_q,
282 kv,
283 dummy_q_fp8,
284 swa_3d,
285 slot_mapping,
286 positions,
287 cos_sin_cache,
288 attn._flashinfer_fp8_kv_scale,
289 attn._flashinfer_fp8_q_scale_inv,
290 attn.eps,
291 block_size,
292 )
293
294
295 class DSparkDeepseekV4ForCausalLM(nn.Module):
296 # Draft weights ship in the target checkpoint (mtp.*) without embed/head, so
297 # load_dspark_model always aliases the target's.
298 has_own_embed_tokens = False
299 has_own_lm_head = False
300 # Full-vocab draft: draft ids are target ids, no remapping needed.
301 draft_id_to_target_id = None
302
303 def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None:
304 super().__init__()
305 assert vllm_config.speculative_config is not None
306 self.draft_model_config = vllm_config.speculative_config.draft_model_config
307 self.config = self.draft_model_config.hf_config
308 self.model = DSparkDeepseekV4Model(
309 vllm_config=vllm_config, prefix=maybe_prefix(prefix, "model")
310 )
311 # Shared with the target (aliased by the speculator's load utility).
312 self.lm_head = ParallelLMHead(
313 self.config.vocab_size,
314 self.config.hidden_size,
315 prefix=maybe_prefix(prefix, "lm_head"),
316 )
317 self.logits_processor = LogitsProcessor(self.config.vocab_size)
318
319 # --- Hooks used by the speculator -------------------------------------
320
321 def embed_input_ids(self, input_ids: torch.Tensor) -> torch.Tensor:
322 return self.model.embed_input_ids(input_ids)
323
324 def combine_hidden_states(self, aux_hidden_states: torch.Tensor) -> torch.Tensor:
325 return self.model.combine_hidden_states(aux_hidden_states)
326
327 def get_draft_kv_cache_layer_names(self) -> list[str]:
328 # DSV4 MLA path: each draft layer's sliding-window cache is a separate
329 # layer, named by its prefix.
330 return [layer.attn.swa_cache_layer.prefix for layer in self.model.layers]
331
332 def precompute_and_store_context_kv(
333 self,
334 context_states: torch.Tensor,
335 context_positions: torch.Tensor,
336 context_slot_mappings: list[torch.Tensor | None] | None = None,
337 ) -> None:
338 self.model.precompute_and_store_context_kv(
339 context_states, context_positions, context_slot_mappings
340 )
341
342 def forward(
343 self,
344 input_ids: torch.Tensor,
345 positions: torch.Tensor,
346 inputs_embeds: torch.Tensor | None = None,
347 ) -> torch.Tensor:
348 # Returns the pre-norm hc_head hidden ([T, hidden_size]).
349 return self.model(input_ids, positions, inputs_embeds)
350
351 def compute_logits(self, hidden_states: torch.Tensor) -> torch.Tensor:
352 """Base logits U_k = lm_head(norm(head_hidden))."""
353 return self.logits_processor(self.lm_head, self.model.norm(hidden_states))
354
355 def compute_draft_logits(self, hidden_states: torch.Tensor) -> torch.Tensor:
356 # Full-vocab draft: base logits, no d2t scatter.
357 return self.compute_logits(hidden_states)
358
359 def map_draft_to_target(self, draft_ids: torch.Tensor) -> torch.Tensor:
360 return draft_ids # full-vocab: draft ids are target ids
361
362 def markov_embed(self, token_ids: torch.Tensor) -> torch.Tensor:
363 return self.model.markov_head.embed(token_ids)
364
365 def markov_bias(self, markov_embed: torch.Tensor) -> torch.Tensor:
366 return self.model.markov_head.bias(markov_embed, self.logits_processor)
367
368 def compute_confidence(
369 self, head_hidden: torch.Tensor, markov_embed: torch.Tensor
370 ) -> torch.Tensor:
371 """Per-position acceptance probability for each drafted token."""
372 if self.model.confidence_head is None:
373 raise RuntimeError(
374 "compute_confidence() requires a confidence head, but the "
375 "checkpoint did not provide confidence_head weights."
376 )
377 return torch.sigmoid(self.model.confidence_head(head_hidden, markov_embed))
378
379 # --- Weight loading ----------------------------------------------------
380
381 def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
382 """Load the ``mtp.{0,1,2}.*`` draft weights from the target checkpoint.
383
384 Non-mtp weights (embed/head/main layers) belong to the target model and
385 are skipped here. ``embed_tokens``/``lm_head`` are aliased from the target.
386 """
387 # AMD DeepseekV4MoE has no mega-MoE path; always use the standard
388 # per-expert fused-MoE mapping (mirrors amd/mtp.py).
389 expert_mapping = fused_moe_make_expert_params_mapping(
390 self,
391 ckpt_gate_proj_name="w1",
392 ckpt_down_proj_name="w2",
393 ckpt_up_proj_name="w3",
394 num_experts=self.config.n_routed_experts,
395 )
396 expert_scale_suffix = (
397 ".weight_scale"
398 if getattr(self.config, "expert_dtype", "fp4") == "fp4"
399 else ".weight_scale_inv"
400 )
401
402 # (param_name, ckpt_shard_name, shard_id) for non-expert stacked params.
403 stacked_params_mapping = [
404 ("gate_up_proj", "w1", 0),
405 ("gate_up_proj", "w3", 1),
406 ("attn.fused_wqa_wkv", "attn.wq_a", 0),
407 ("attn.fused_wqa_wkv", "attn.wkv", 1),
408 ]
409
410 params_dict = dict(self.named_parameters())
411 loaded_params: set[str] = set()
412 loaded_confidence_head = False
413
414 tp_size = get_tensor_model_parallel_world_size()
415 tp_rank = get_tensor_model_parallel_rank()
416 n_local_head = self.config.num_attention_heads // tp_size
417 head_start = n_local_head * tp_rank
418 head_end = n_local_head * (tp_rank + 1)
419
420 for name, loaded_weight in weights:
421 mapped = self._remap_dspark_name(name)
422 if mapped is None:
423 continue
424 name = mapped
425 if "confidence_head." in name:
426 loaded_confidence_head = True
427
428 # ``.scale`` -> per-method scale suffix.
429 if name.endswith(".scale"):
430 suffix = (
431 expert_scale_suffix
432 if _EXPERT_SCALE_RE.search(name)
433 else ".weight_scale_inv"
434 )
435 name = name.removesuffix(".scale") + suffix
436
437 # E8M0 expert scales: keep raw exponent bytes.
438 if ".experts." in name:
439 if (
440 "weight_scale" in name
441 and loaded_weight.dtype == torch.float8_e8m0fnu
442 ):
443 loaded_weight = loaded_weight.view(torch.uint8)
444 for param_name, weight_name, expert_id, shard_id in expert_mapping:
445 if weight_name not in name:
446 continue
447 name_mapped = name.replace(weight_name, param_name)
448 param = params_dict[name_mapped]
449 success = param.weight_loader(
450 param,
451 loaded_weight,
452 name_mapped,
453 shard_id=shard_id,
454 expert_id=expert_id,
455 return_success=True,
456 )
457 if success:
458 loaded_params.add(name_mapped)
459 break
460 continue
461
462 # Stacked rules only apply to decoder-layer weights. Head-stack params
463 # (main_proj/norm/hc_head/markov_head/confidence_head) load directly —
464 # otherwise e.g. "markov_w1" would collide with the "w1" shard rule.
465 is_layer_param = name.startswith("model.layers.")
466 for param_name, weight_name, stacked_shard_id in stacked_params_mapping:
467 if not is_layer_param or weight_name not in name:
468 continue
469 name = name.replace(weight_name, param_name)
470 param = params_dict[name]
471 param.weight_loader(param, loaded_weight, stacked_shard_id)
472 loaded_params.add(name)
473 break
474 else:
475 if "attn_sink" in name:
476 narrow = loaded_weight[head_start:head_end]
477 params_dict[name][: narrow.shape[0]].copy_(narrow)
478 loaded_params.add(name)
479 continue
480 if ".shared_experts.w2" in name:
481 name = name.replace(
482 ".shared_experts.w2", ".shared_experts.down_proj"
483 )
484 if name.endswith(".ffn.gate.bias"):
485 name = name.replace(
486 ".ffn.gate.bias", ".ffn.gate.e_score_correction_bias"
487 )
488 param = params_dict[name]
489 weight_loader = getattr(param, "weight_loader", default_weight_loader)
490 weight_loader(param, loaded_weight)
491 loaded_params.add(name)
492
493 if self.model.confidence_head is not None and not loaded_confidence_head:
494 self.model.confidence_head = None
495 logger.info_once("DSpark draft model loaded: %d params", len(loaded_params))
496 return loaded_params
497
498 def _remap_dspark_name(self, name: str) -> str | None:
499 """Map a checkpoint ``mtp.{i}.*`` name to this model's parameter path.
500
501 Returns None for non-mtp weights (owned by the target model).
502 """
503 m = re.match(r"mtp\.(\d+)\.(.*)", name)
504 if m is None:
505 return None
506 stage = int(m.group(1))
507 rest = m.group(2)
508 if rest.startswith("confidence_head.") and self.model.confidence_head is None:
509 return None
510 # Head-stack params live at model level (mtp.last), context combiner at
511 # model level (mtp.0); everything else is a per-layer decoder block.
512 head_prefixes = (
513 "norm.",
514 "hc_head_fn",
515 "hc_head_base",
516 "hc_head_scale",
517 "markov_head.",
518 "confidence_head.",
519 )
520 if rest.startswith(("main_proj.", "main_norm.")) or rest.startswith(
521 head_prefixes
522 ):
523 return f"model.{rest}"
524 return f"model.layers.{stage}.{rest}"