dflash2_speculator.py 8.7 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265
  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. from vllm.config import VllmConfig
  6. from vllm.config.compilation import CUDAGraphMode
  7. from vllm.triton_utils import tl, triton
  8. from vllm.v1.worker.gpu.sample.gumbel import gumbel_noised_argmax
  9. from vllm.v1.worker.gpu.spec_decode.dflash.speculator import DFlashSpeculator
  10. @triton.jit
  11. def _selector_walk_kernel(
  12. scores_ptr,
  13. candidate_ptr,
  14. sample_pos_ptr,
  15. req_state_ptr,
  16. temperature_ptr,
  17. seeds_ptr,
  18. tokens_ptr,
  19. realized_scores_ptr,
  20. num_steps: tl.constexpr,
  21. top_k: tl.constexpr,
  22. BLOCK_K: tl.constexpr,
  23. SAMPLE_PROBABILISTIC: tl.constexpr,
  24. USE_FP64: tl.constexpr,
  25. ):
  26. row = tl.program_id(0)
  27. offsets = tl.arange(0, BLOCK_K)
  28. mask = offsets < top_k
  29. req_state = tl.load(req_state_ptr + row * num_steps)
  30. valid = req_state >= 0
  31. temperature = tl.load(temperature_ptr + req_state, mask=valid, other=0.0)
  32. seed = tl.load(seeds_ptr + req_state, mask=valid, other=0)
  33. previous = 0
  34. for step in range(num_steps):
  35. flat = row * num_steps + step
  36. score_base = (flat * top_k + previous) * top_k
  37. scores = tl.load(
  38. scores_ptr + score_base + offsets,
  39. mask=mask & valid,
  40. other=float("-inf"),
  41. ).to(tl.float64 if USE_FP64 else tl.float32)
  42. candidate_base = flat * top_k
  43. candidates = tl.load(
  44. candidate_ptr + candidate_base + offsets,
  45. mask=mask & valid,
  46. other=0,
  47. )
  48. # sample_pos is the predicted token's position P. Sampling keys a draw
  49. # by the position before the sampled token, P-1.
  50. sample_pos = tl.load(sample_pos_ptr + flat) - 1
  51. _, index = gumbel_noised_argmax(
  52. scores,
  53. candidates,
  54. mask & valid,
  55. seed,
  56. sample_pos,
  57. temperature if SAMPLE_PROBABILISTIC else 0.0,
  58. IS_DRAFTING=True,
  59. USE_FP64=USE_FP64,
  60. )
  61. tl.store(
  62. realized_scores_ptr + candidate_base + offsets,
  63. scores,
  64. mask=mask & valid,
  65. )
  66. token = tl.load(candidate_ptr + candidate_base + index, mask=valid, other=0)
  67. tl.store(tokens_ptr + flat, token, mask=valid)
  68. previous = index
  69. @triton.jit
  70. def _cache_draft_logits_kernel(
  71. draft_logits_ptr,
  72. cached_candidate_ptr,
  73. candidate_ptr,
  74. scores_ptr,
  75. req_state_ptr,
  76. draft_logits_stride_0,
  77. draft_logits_stride_1,
  78. num_steps: tl.constexpr,
  79. top_k: tl.constexpr,
  80. BLOCK_K: tl.constexpr,
  81. ):
  82. flat = tl.program_id(0)
  83. req_state = tl.load(req_state_ptr + flat)
  84. step = flat % num_steps
  85. offsets = tl.arange(0, BLOCK_K)
  86. mask = (req_state >= 0) & (offsets < top_k)
  87. candidate_base = flat * top_k
  88. cache_base = (req_state * num_steps + step) * top_k
  89. old_token_ids = tl.load(cached_candidate_ptr + cache_base + offsets, mask=mask)
  90. logits_base = (
  91. draft_logits_ptr
  92. + req_state * draft_logits_stride_0
  93. + step * draft_logits_stride_1
  94. )
  95. tl.store(logits_base + old_token_ids, -float("inf"), mask=mask)
  96. token_ids = tl.load(candidate_ptr + candidate_base + offsets, mask=mask)
  97. scores = tl.load(scores_ptr + candidate_base + offsets, mask=mask)
  98. tl.store(logits_base + token_ids, scores, mask=mask)
  99. tl.store(cached_candidate_ptr + cache_base + offsets, token_ids, mask=mask)
  100. class CandidateSampler:
  101. """The shared DFlash2/LiLiCorr candidate walk and realized proposal cache."""
  102. def __init__(
  103. self, max_num_reqs: int, num_steps: int, top_k: int, device: torch.device
  104. ):
  105. self.num_steps = num_steps
  106. self.top_k = top_k
  107. self.scores = torch.empty(
  108. max_num_reqs, num_steps, top_k, dtype=torch.float32, device=device
  109. )
  110. self.cached_candidate_ids = torch.zeros(
  111. self.scores.shape, dtype=torch.int64, device=device
  112. )
  113. def sample(
  114. self,
  115. candidate_ids: torch.Tensor,
  116. scores: torch.Tensor,
  117. num_reqs: int,
  118. sample_pos: torch.Tensor,
  119. idx_mapping: torch.Tensor,
  120. temperature: torch.Tensor,
  121. seeds: torch.Tensor,
  122. draft_tokens: torch.Tensor,
  123. draft_logits: torch.Tensor | None,
  124. use_fp64: bool,
  125. ) -> None:
  126. candidate_ids = candidate_ids.contiguous()
  127. sample_pos = sample_pos.contiguous()
  128. idx_mapping = idx_mapping.contiguous()
  129. block_k = triton.next_power_of_2(self.top_k)
  130. _selector_walk_kernel[(num_reqs,)](
  131. scores.contiguous(),
  132. candidate_ids.contiguous(),
  133. sample_pos,
  134. idx_mapping,
  135. temperature,
  136. seeds,
  137. draft_tokens,
  138. self.scores,
  139. num_steps=self.num_steps,
  140. top_k=self.top_k,
  141. BLOCK_K=block_k,
  142. SAMPLE_PROBABILISTIC=draft_logits is not None,
  143. USE_FP64=use_fp64,
  144. num_warps=1,
  145. )
  146. if draft_logits is not None:
  147. self._cache_draft_logits(
  148. candidate_ids, num_reqs * self.num_steps, idx_mapping, draft_logits
  149. )
  150. def _cache_draft_logits(
  151. self,
  152. candidate_ids: torch.Tensor,
  153. num_sample: int,
  154. idx_mapping: torch.Tensor,
  155. draft_logits: torch.Tensor,
  156. ) -> None:
  157. block_k = triton.next_power_of_2(self.top_k)
  158. _cache_draft_logits_kernel[(num_sample,)](
  159. draft_logits,
  160. self.cached_candidate_ids,
  161. candidate_ids,
  162. self.scores,
  163. idx_mapping,
  164. draft_logits.stride(0),
  165. draft_logits.stride(1),
  166. num_steps=self.num_steps,
  167. top_k=self.top_k,
  168. BLOCK_K=block_k,
  169. num_warps=1,
  170. )
  171. class DFlash2Speculator(DFlashSpeculator):
  172. _speculator_name = "DFlash2"
  173. _candidate_top_k_key = "selector_top_k"
  174. def __init__(self, vllm_config: VllmConfig, device: torch.device):
  175. super().__init__(vllm_config, device)
  176. draft_config = self.draft_model_config.hf_config.dflash_config
  177. self.top_k = int(draft_config[self._candidate_top_k_key])
  178. self.candidate_sampler = CandidateSampler(
  179. self.max_num_reqs, self.num_speculative_steps, self.top_k, device
  180. )
  181. def draft_logits_spec(self, vllm_config: VllmConfig) -> tuple[torch.dtype, float]:
  182. # fp32 so the walk and the rejection that checks it read the same
  183. # distribution; -inf because the cache kernel writes only the K
  184. # candidates.
  185. return torch.float32, -float("inf")
  186. def _generate_draft(
  187. self,
  188. num_reqs: int,
  189. num_tokens_padded: int,
  190. attn_metadata: dict[str, Any] | None,
  191. slot_mappings: dict[str, torch.Tensor] | None,
  192. num_tokens_across_dp: torch.Tensor | None,
  193. cudagraph_runtime_mode: CUDAGraphMode = CUDAGraphMode.NONE,
  194. ) -> None:
  195. last_hidden_states = self._run_model(
  196. num_tokens_padded,
  197. attn_metadata,
  198. slot_mappings,
  199. num_tokens_across_dp,
  200. cudagraph_runtime_mode,
  201. )
  202. num_sample = num_reqs * self.num_speculative_steps
  203. hidden_states = last_hidden_states[self.sample_indices[:num_sample]].view(
  204. num_reqs, self.num_speculative_steps, -1
  205. )
  206. candidate_ids, unary_logits = self.model.compute_candidates(
  207. hidden_states.flatten(0, 1)
  208. )
  209. candidate_ids = candidate_ids.view(
  210. num_reqs, self.num_speculative_steps, self.top_k
  211. )
  212. unary_logits = unary_logits.view_as(candidate_ids)
  213. scores = self._score_candidates(candidate_ids, unary_logits, hidden_states)
  214. self.candidate_sampler.sample(
  215. candidate_ids,
  216. scores,
  217. num_reqs,
  218. self.sample_pos,
  219. self.sample_idx_mapping,
  220. self.temperature,
  221. self.seeds,
  222. self.draft_tokens,
  223. self.draft_logits,
  224. self.use_fp64_gumbel,
  225. )
  226. if self.acceptance_estimator is not None:
  227. self._maybe_predict_acceptance(
  228. self.candidate_sampler.scores[:num_reqs].flatten(0, 1),
  229. self.sample_idx_mapping[:num_sample],
  230. self.sample_col[:num_sample],
  231. )
  232. def _score_candidates(
  233. self,
  234. candidate_ids: torch.Tensor,
  235. unary_logits: torch.Tensor,
  236. hidden_states: torch.Tensor,
  237. ) -> torch.Tensor:
  238. num_reqs = candidate_ids.shape[0]
  239. anchor_token_ids = self.input_buffers.input_ids[
  240. : num_reqs * self.num_query_per_req : self.num_query_per_req
  241. ]
  242. return self.model.model.candidate_selector(
  243. candidate_ids, unary_logits, hidden_states, anchor_token_ids
  244. )