model_state.py 4.9 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127
  1. # SPDX-License-Identifier: Apache-2.0
  2. # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
  3. from typing import Any
  4. import torch
  5. import torch.nn as nn
  6. from vllm.config import VllmConfig
  7. from vllm.models.deepseek_v4_1.common.engram import (
  8. Engram,
  9. EngramDiskStager,
  10. NgramHashState,
  11. )
  12. from vllm.triton_utils import tl, triton
  13. from vllm.v1.worker.gpu.input_batch import InputBatch
  14. from vllm.v1.worker.gpu.mm.encoder_cache import EncoderCache
  15. from vllm.v1.worker.gpu.model_states.default import DefaultModelState
  16. from vllm.v1.worker.gpu.states import RequestState
  17. @triton.jit
  18. def _gather_lookback_kernel(
  19. lookback_ptr,
  20. idx_mapping_ptr,
  21. num_computed_tokens_ptr,
  22. all_token_ids_ptr,
  23. all_token_ids_stride,
  24. num_reqs,
  25. DEPTH: tl.constexpr,
  26. BLOCK_DEPTH: tl.constexpr,
  27. ):
  28. # One program per lookback row; rows past the batch are filled with -1.
  29. batch_idx = tl.program_id(0)
  30. in_batch = batch_idx < num_reqs
  31. req_state_idx = tl.load(idx_mapping_ptr + batch_idx, mask=in_batch, other=0)
  32. num_computed = tl.load(num_computed_tokens_ptr + req_state_idx)
  33. offs = tl.arange(0, BLOCK_DEPTH)
  34. pos = num_computed - 1 - offs
  35. valid = in_batch & (offs < DEPTH) & (pos >= 0)
  36. ids = tl.load(
  37. all_token_ids_ptr + req_state_idx * all_token_ids_stride + pos,
  38. mask=valid,
  39. other=-1,
  40. )
  41. tl.store(lookback_ptr + batch_idx * DEPTH + offs, ids, mask=offs < DEPTH)
  42. class DeepseekV41ModelState(DefaultModelState):
  43. """DefaultModelState plus the engram lookback window.
  44. The engram n-gram hash needs the ids of the ``depth`` tokens preceding
  45. each request's chunk start (see ``common/engram.py``). The runner keeps
  46. the full token history on device, so the window is gathered there every
  47. step: exact for prompt and generated tokens alike, whatever instance
  48. produced their KV.
  49. """
  50. def __init__(
  51. self,
  52. vllm_config: VllmConfig,
  53. model: nn.Module,
  54. encoder_cache: EncoderCache | None,
  55. device: torch.device,
  56. ):
  57. super().__init__(vllm_config, model, encoder_cache, device)
  58. depth = model.token_lookback_depth
  59. self.lookback_token_ids: torch.Tensor | None = None
  60. if depth > 0:
  61. # Persistent so a captured graph can read it on replay.
  62. self.lookback_token_ids = torch.full(
  63. (self.max_num_reqs, depth), -1, dtype=torch.int32, device=device
  64. )
  65. # Tech2Wild/Kai 2026-09-10 (CUDA graphs): DISK-backed Engram rows are
  66. # staged in prepare_inputs, outside the (possibly captured) forward.
  67. self.engram_stager: EngramDiskStager | None = None
  68. engrams = [m for m in model.modules() if isinstance(m, Engram) and m.prestage]
  69. if engrams:
  70. hash_states = [m for m in model.modules() if isinstance(m, NgramHashState)]
  71. assert len(hash_states) == 1, (
  72. f"expected one NgramHashState, found {len(hash_states)}"
  73. )
  74. self.engram_stager = EngramDiskStager(hash_states[0], engrams)
  75. def prepare_inputs(
  76. self, input_batch: InputBatch, req_states: RequestState
  77. ) -> dict[str, torch.Tensor | None]:
  78. model_inputs = super().prepare_inputs(input_batch, req_states)
  79. window = self.lookback_token_ids
  80. if window is None:
  81. return model_inputs
  82. all_token_ids = req_states.all_token_ids.gpu
  83. depth = window.shape[1]
  84. _gather_lookback_kernel[(window.shape[0],)](
  85. window,
  86. input_batch.idx_mapping,
  87. req_states.num_computed_tokens.gpu,
  88. all_token_ids,
  89. all_token_ids.stride(0),
  90. input_batch.idx_mapping.shape[0],
  91. DEPTH=depth,
  92. BLOCK_DEPTH=triton.next_power_of_2(depth),
  93. )
  94. model_inputs["lookback_token_ids"] = window
  95. if self.engram_stager is not None and input_batch.input_ids is not None:
  96. # After the runner wrote this step's ids/positions/query_start_loc
  97. # and the lookback window above; before the forward / graph replay.
  98. positions = model_inputs.get("positions")
  99. self.engram_stager.stage(
  100. input_batch.input_ids,
  101. positions if positions is not None else input_batch.positions,
  102. input_batch.query_start_loc[: input_batch.num_reqs + 1],
  103. window,
  104. input_batch.num_tokens,
  105. input_batch=input_batch,
  106. req_states=req_states,
  107. )
  108. return model_inputs
  109. def prepare_dummy_inputs(self, num_reqs: int, num_tokens: int) -> dict[str, Any]:
  110. model_inputs = super().prepare_dummy_inputs(num_reqs, num_tokens)
  111. if self.lookback_token_ids is not None:
  112. # The captured graph reads this buffer; replays refill it in place.
  113. self.lookback_token_ids.fill_(-1)
  114. model_inputs["lookback_token_ids"] = self.lookback_token_ids
  115. return model_inputs