rejection_sampler.py 13 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340
  1. # SPDX-License-Identifier: Apache-2.0
  2. # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
  3. from collections.abc import Iterable, Iterator
  4. from typing import Any
  5. import numpy as np
  6. import torch
  7. from vllm.v1.worker.gpu.spec_decode import draft_trunc
  8. from vllm.config import SpeculativeConfig
  9. from vllm.config.model import PROCESSED_LOGPROBS_MODES
  10. from vllm.triton_utils import tl, triton
  11. from vllm.v1.outputs import LogprobsTensors
  12. from vllm.v1.spec_decode.utils import unconditional_to_conditional_rates
  13. from vllm.v1.watermarking.gpu_sampler import GPUWatermarkSampler
  14. from vllm.v1.worker.gpu.input_batch import (
  15. InputBatch,
  16. get_num_sampled_and_rejected,
  17. )
  18. from vllm.v1.worker.gpu.metrics.logits import get_num_nans
  19. from vllm.v1.worker.gpu.sample.logprob import compute_topk_scores
  20. from vllm.v1.worker.gpu.sample.output import SamplerOutput
  21. from vllm.v1.worker.gpu.sample.sampler import Sampler
  22. from vllm.v1.worker.gpu.sample.states import NO_LOGPROBS
  23. from vllm.v1.worker.gpu.spec_decode.rejection_sampler_utils import (
  24. rejection_sample,
  25. )
  26. # Cap on the FP32 target-logits buffer materialized by apply_sampling_params.
  27. # TODO(mgoin): Chunking is a workaround. The rejection kernels already upcast
  28. # per vocab block on load and apply ops like temperature and gumbel, so folding
  29. # sampling-param application into those kernels would remove this buffer and
  30. # its traffic entirely.
  31. MAX_CHUNK_BYTES = 2**30 # 1GB
  32. _FP32_BYTES = 4
  33. def get_max_chunk_logits(vocab_size: int) -> int:
  34. """Largest number of logits rows one verification chunk may hold."""
  35. return max(1, MAX_CHUNK_BYTES // (vocab_size * _FP32_BYTES))
  36. def _iter_request_chunks(
  37. cu_num_logits: np.ndarray, max_chunk_logits: int
  38. ) -> Iterator[tuple[int, int]]:
  39. """Yield maximally packed request ranges without splitting requests."""
  40. assert max_chunk_logits > 0
  41. num_reqs = cu_num_logits.size - 1
  42. start = 0
  43. while start < num_reqs:
  44. max_logit = int(cu_num_logits[start]) + max_chunk_logits
  45. end = int(np.searchsorted(cu_num_logits, max_logit, side="right") - 1)
  46. end = min(num_reqs, max(start + 1, end))
  47. yield start, end
  48. start = end
  49. @triton.jit
  50. def _flatten_sampled_kernel(
  51. # [num_logits]
  52. flat_sampled_ptr,
  53. # [num_reqs, num_speculative_steps + 1]
  54. sampled_ptr,
  55. sampled_stride,
  56. # [num_reqs]
  57. num_sampled_ptr,
  58. # [num_reqs + 1]
  59. cu_num_logits_ptr,
  60. ):
  61. req_idx = tl.program_id(0)
  62. start_idx = tl.load(cu_num_logits_ptr + req_idx)
  63. num_sampled = tl.load(num_sampled_ptr + req_idx)
  64. for i in range(num_sampled):
  65. token_id = tl.load(sampled_ptr + req_idx * sampled_stride + i)
  66. tl.store(flat_sampled_ptr + start_idx + i, token_id)
  67. class RejectionSampler:
  68. def __init__(
  69. self,
  70. sampler: Sampler,
  71. spec_config: SpeculativeConfig,
  72. device: torch.device,
  73. *,
  74. watermark_key: int | None = None,
  75. ):
  76. self.sampler = sampler
  77. self.watermark_key = watermark_key
  78. if watermark_key is not None:
  79. assert isinstance(sampler, GPUWatermarkSampler)
  80. self.num_speculative_steps = spec_config.num_speculative_tokens
  81. self.enable_adaptive_verification = spec_config.enable_adaptive_verification
  82. rejection_sample_method = spec_config.rejection_sample_method
  83. self.use_block_verification: bool = False
  84. self.synthetic_conditional_rates: torch.Tensor | None = None
  85. if rejection_sample_method == "synthetic":
  86. assert spec_config.synthetic_acceptance_rates is not None
  87. self.synthetic_conditional_rates = torch.tensor(
  88. unconditional_to_conditional_rates(
  89. spec_config.synthetic_acceptance_rates
  90. ),
  91. dtype=torch.float32,
  92. device=device,
  93. )
  94. elif rejection_sample_method == "block":
  95. self.use_block_verification = True
  96. def _get_logprobs_tensors(
  97. self,
  98. sampled: torch.Tensor,
  99. num_sampled: torch.Tensor,
  100. logits: torch.Tensor,
  101. cu_num_logits: torch.Tensor,
  102. cu_num_logits_np: np.ndarray,
  103. max_num_logprobs: int,
  104. ) -> LogprobsTensors | None:
  105. if max_num_logprobs == NO_LOGPROBS:
  106. return None
  107. num_reqs = cu_num_logits.shape[0] - 1
  108. num_logits = logits.shape[0]
  109. flat_sampled = torch.zeros(
  110. num_logits, dtype=sampled.dtype, device=sampled.device
  111. )
  112. _flatten_sampled_kernel[(num_reqs,)](
  113. flat_sampled,
  114. sampled,
  115. sampled.stride(0),
  116. num_sampled,
  117. cu_num_logits,
  118. num_warps=1,
  119. )
  120. expanded_logits = num_logits != num_reqs
  121. cu_num_generated_tokens: list[int] | torch.Tensor | None = None
  122. if expanded_logits:
  123. if self.enable_adaptive_verification:
  124. # Adaptive verification keeps the true per-request boundaries
  125. # on device only; cu_num_logits_np holds the pre-compacted layout.
  126. cu_num_generated_tokens = cu_num_logits.clone()
  127. else:
  128. cu_num_generated_tokens = cu_num_logits_np.tolist()
  129. return compute_topk_scores(
  130. logits,
  131. max_num_logprobs,
  132. flat_sampled,
  133. cu_num_generated_tokens,
  134. logits_mode=self.sampler.logprobs_mode
  135. in ("raw_logits", "processed_logits"),
  136. )
  137. def _watermarking_kwargs(
  138. self,
  139. draft_sampled: torch.Tensor,
  140. expanded_idx_mapping: torch.Tensor,
  141. expanded_local_pos: torch.Tensor,
  142. ) -> dict[str, Any]:
  143. if self.watermark_key is None:
  144. return {}
  145. assert isinstance(self.sampler, GPUWatermarkSampler)
  146. return {
  147. "contexts": self.sampler._get_contexts(
  148. expanded_idx_mapping, expanded_local_pos, draft_sampled
  149. ),
  150. "watermarking": self.sampler.watermarking.gpu,
  151. "watermark_key": self.watermark_key,
  152. }
  153. def _verify(
  154. self,
  155. logits: torch.Tensor,
  156. draft_logits: torch.Tensor | None,
  157. draft_sampled: torch.Tensor,
  158. accept_drafts: torch.Tensor,
  159. pos: torch.Tensor,
  160. cu_num_logits: torch.Tensor,
  161. idx_mapping: torch.Tensor,
  162. idx_mapping_np: np.ndarray,
  163. expanded_idx_mapping: torch.Tensor,
  164. expanded_local_pos: torch.Tensor,
  165. seq_lens_upper_bound_np: np.ndarray,
  166. ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
  167. processed_logits = self.sampler.apply_sampling_params(
  168. logits,
  169. expanded_idx_mapping,
  170. idx_mapping,
  171. idx_mapping_np,
  172. pos,
  173. draft_sampled,
  174. expanded_local_pos,
  175. seq_lens_upper_bound_np,
  176. )
  177. sampled, num_sampled = rejection_sample(
  178. processed_logits,
  179. draft_logits,
  180. accept_drafts,
  181. cu_num_logits,
  182. pos,
  183. idx_mapping,
  184. expanded_idx_mapping,
  185. expanded_local_pos,
  186. self.sampler.sampling_states.temperature.gpu,
  187. self.sampler.sampling_states.seeds.gpu,
  188. self.num_speculative_steps,
  189. self.synthetic_conditional_rates,
  190. use_fp64=self.sampler.use_fp64_gumbel,
  191. use_block_verification=self.use_block_verification,
  192. **self._watermarking_kwargs(
  193. draft_sampled, expanded_idx_mapping, expanded_local_pos
  194. ),
  195. )
  196. return processed_logits, sampled, num_sampled
  197. def _verify_in_chunks(
  198. self,
  199. logits: torch.Tensor,
  200. input_batch: InputBatch,
  201. draft_logits: torch.Tensor | None,
  202. draft_sampled: torch.Tensor,
  203. pos: torch.Tensor,
  204. max_chunk_logits: int,
  205. max_num_logprobs: int,
  206. accept_drafts: torch.Tensor | None = None,
  207. ) -> tuple[torch.Tensor, torch.Tensor, LogprobsTensors | None]:
  208. cu_num_logits_np = input_batch.cu_num_logits_np
  209. use_processed_logits = self.sampler.logprobs_mode in PROCESSED_LOGPROBS_MODES
  210. num_reqs = input_batch.num_reqs
  211. if logits.shape[0] <= max_chunk_logits:
  212. # One chunk covers the batch. Adaptive verification compacts the logits
  213. # without updating cu_num_logits_np (it keeps the pre-compacted layout),
  214. # so the stale sums must not pick chunk boundaries; its budget cap
  215. # guarantees the compacted batch always lands here.
  216. request_chunks: Iterable[tuple[int, int]] = ((0, num_reqs),)
  217. else:
  218. assert not self.enable_adaptive_verification
  219. request_chunks = _iter_request_chunks(cu_num_logits_np, max_chunk_logits)
  220. sampled_chunks: list[torch.Tensor] = []
  221. num_sampled_chunks: list[torch.Tensor] = []
  222. logprobs_chunks: list[LogprobsTensors] = []
  223. for start, end in request_chunks:
  224. lo = int(cu_num_logits_np[start])
  225. hi = int(cu_num_logits_np[end])
  226. chunk_cu_num_logits_np = cu_num_logits_np[start : end + 1] - lo
  227. chunk_cu_num_logits = input_batch.cu_num_logits[start : end + 1] - lo
  228. # draft_logits uses persistent request-state indices and stays global.
  229. processed_logits, sampled, num_sampled = self._verify(
  230. logits[lo:hi],
  231. draft_logits,
  232. draft_sampled[lo:hi],
  233. (draft_sampled if accept_drafts is None else accept_drafts)[lo:hi],
  234. pos[lo:hi],
  235. chunk_cu_num_logits,
  236. input_batch.idx_mapping[start:end],
  237. input_batch.idx_mapping_np[start:end],
  238. input_batch.expanded_idx_mapping[lo:hi],
  239. input_batch.expanded_local_pos[lo:hi],
  240. input_batch.seq_lens_cpu_upper_bound.numpy()[start:end],
  241. )
  242. chunk_logprobs = self._get_logprobs_tensors(
  243. sampled,
  244. num_sampled,
  245. processed_logits if use_processed_logits else logits[lo:hi],
  246. chunk_cu_num_logits,
  247. chunk_cu_num_logits_np,
  248. max_num_logprobs,
  249. )
  250. if chunk_logprobs is not None:
  251. logprobs_chunks.append(chunk_logprobs)
  252. del processed_logits
  253. sampled_chunks.append(sampled)
  254. num_sampled_chunks.append(num_sampled)
  255. if len(sampled_chunks) == 1:
  256. logprobs_tensors = logprobs_chunks[0] if logprobs_chunks else None
  257. return sampled_chunks[0], num_sampled_chunks[0], logprobs_tensors
  258. logprobs_tensors = None
  259. if logprobs_chunks:
  260. expanded_logits = logits.shape[0] != input_batch.num_reqs
  261. logprobs_tensors = LogprobsTensors.cat(
  262. logprobs_chunks,
  263. cu_num_generated_tokens=(
  264. cu_num_logits_np.tolist() if expanded_logits else None
  265. ),
  266. )
  267. sampled = torch.cat(sampled_chunks)
  268. num_sampled = torch.cat(num_sampled_chunks)
  269. return sampled, num_sampled, logprobs_tensors
  270. def __call__(
  271. self,
  272. logits: torch.Tensor,
  273. input_batch: InputBatch,
  274. draft_logits: torch.Tensor | None = None,
  275. ) -> SamplerOutput:
  276. # NOTE(woosuk): We intentionally compute num_nans before sampling to make clear
  277. # that num_nans is computed before applying penalties and temperature.
  278. num_nans = get_num_nans(logits) if self.sampler.compute_nans else None
  279. draft_sampled = input_batch.input_ids[input_batch.logits_indices]
  280. # A dead draft (draft_trunc) is -1 for acceptance only, so it never matches.
  281. accept_drafts = draft_sampled
  282. if draft_trunc.ENABLED:
  283. accept_drafts = draft_sampled.masked_fill(input_batch.is_padding[input_batch.logits_indices], -1)
  284. pos = input_batch.positions[input_batch.logits_indices]
  285. max_num_logprobs = self.sampler.sampling_states.max_num_logprobs(
  286. input_batch.idx_mapping_np
  287. )
  288. chunk_logit_limit = get_max_chunk_logits(logits.shape[1])
  289. sampled, num_sampled, logprobs_tensors = self._verify_in_chunks(
  290. logits,
  291. input_batch,
  292. draft_logits,
  293. draft_sampled,
  294. pos,
  295. chunk_logit_limit,
  296. max_num_logprobs,
  297. accept_drafts=accept_drafts,
  298. )
  299. num_sampled, num_rejected = get_num_sampled_and_rejected(
  300. num_sampled,
  301. input_batch.seq_lens,
  302. input_batch.cu_num_logits,
  303. input_batch.idx_mapping,
  304. self.sampler.req_states.prefill_len.gpu,
  305. )
  306. return SamplerOutput(
  307. sampled_token_ids=sampled,
  308. logprobs_tensors=logprobs_tensors,
  309. num_nans=num_nans,
  310. num_sampled=num_sampled,
  311. num_rejected=num_rejected,
  312. )