draft_trunc.py 7.1 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161
  1. """Per-step draft truncation.
  2. The scheduler picks how many drafts each request verifies before the drafts
  3. exist, so the verify shape is fixed. Truncation leaves the shape alone and
  4. marks a request's later drafts dead on the device: their MoE routes go to
  5. expert -1, which the MoE kernels skip without reading those experts' weights,
  6. and the rejection sampler sees them as -1, so acceptance stops before them.
  7. They come after the live tokens, so causal attention leaves the live tokens
  8. exact.
  9. By default (VLLM_DRAFT_TRUNC_TAU, 0.3 in adaptive-k.yaml) each request keeps
  10. its leading drafts whose predicted survival, the running product of the
  11. acceptance estimator's per-draft probabilities, is at least tau. A cut step
  12. grades only the drafts the estimator already expected to survive, so fitting
  13. on it would bias the estimator, and a position it always kills could never
  14. recover. Every VLLM_DRAFT_TRUNC_EXPLORE-th step (4 by default) therefore skips
  15. the cut, and only those steps are folded into the estimator. The file named by
  16. VLLM_DRAFT_TRUNC_CONTROL overrides that while it exists: "tau X" sets the
  17. threshold, "live N" keeps each request's first N drafts, and anything else
  18. truncates nothing.
  19. """
  20. import os
  21. import torch
  22. import vllm
  23. from vllm.triton_utils import tl, triton
  24. # The other files in this directory are whole copies of vLLM's, with the cut added. They
  25. # match one vLLM version, the image's VLLM_REF. On any other they would put back old code
  26. # over new, so the engine stops instead.
  27. # The v0.31.0 base is a RELEASE build: its version string is the bare "0.31.0"
  28. # (the nightly bases carried a +g<sha> scm suffix, e.g. +gddd6fbca).
  29. VLLM_COMMIT = "0.31.0" # v0.31.0 release (tag commit db9527a4)
  30. if vllm.__version__ != VLLM_COMMIT:
  31. raise RuntimeError(f"experimental/draft-trunc copies vLLM {VLLM_COMMIT}'s files, but this image runs "
  32. f"vLLM {vllm.__version__}. Refresh the copies from this image, or remove the "
  33. "draft-trunc mounts from experimental/compose/adaptive-k.yaml.")
  34. _TAU = float(os.environ.get("VLLM_DRAFT_TRUNC_TAU") or 0)
  35. ENABLED = bool(os.environ.get("VLLM_DRAFT_TRUNC_CONTROL")) or _TAU > 0
  36. # Measurement: per verify step, each request's draft confidences (the
  37. # acceptance estimator's, computed when the drafts were made) and how many
  38. # tokens it sampled, one JSON line per step, on tensor-parallel rank 0.
  39. LOG = os.environ.get("VLLM_DRAFT_TRUNC_LOG", "")
  40. _LOG_MAX = int(os.environ.get("VLLM_DRAFT_TRUNC_LOG_STEPS", "20000"))
  41. _logged = 0
  42. _control = os.environ.get("VLLM_DRAFT_TRUNC_CONTROL", "")
  43. _mtime = 0.0
  44. _mode: tuple[str, float] | None = None
  45. _EXPLORE = int(os.environ.get("VLLM_DRAFT_TRUNC_EXPLORE") or 4)
  46. _steps = 0
  47. _cut = False
  48. def _read_control() -> tuple[str, float] | None:
  49. """("live", N), ("tau", X), or None to truncate nothing."""
  50. global _mtime, _mode
  51. default = ("tau", _TAU) if _TAU > 0 else None
  52. try:
  53. mtime = os.stat(_control).st_mtime if _control else None
  54. except OSError:
  55. mtime = None
  56. if mtime is None:
  57. _mtime, _mode = 0.0, default
  58. return default
  59. if mtime != _mtime:
  60. _mtime = mtime
  61. words = open(_control).read().split()
  62. _mode = (words[0], float(words[1])) if len(words) == 2 and words[0] in ("live", "tau") else None
  63. return _mode
  64. def live_drafts() -> int | None:
  65. """The control file's N, or None."""
  66. mode = _read_control()
  67. return int(mode[1]) if mode and mode[0] == "live" else None
  68. @triton.jit
  69. def _mark_dead_kernel(is_padding_ptr, query_start_loc_ptr, cu_num_logits_ptr, live, BLOCK: tl.constexpr):
  70. b = tl.program_id(0)
  71. num_drafts = tl.load(cu_num_logits_ptr + b + 1) - tl.load(cu_num_logits_ptr + b) - 1
  72. end = tl.load(query_start_loc_ptr + b + 1)
  73. j = tl.arange(0, BLOCK)
  74. # A request's drafts are its last num_drafts tokens.
  75. tl.store(is_padding_ptr + end - num_drafts + j, tl.full((BLOCK,), 1, tl.int1),
  76. mask=(j >= live) & (j < num_drafts))
  77. @triton.jit
  78. def _mark_cut_kernel(is_padding_ptr, query_start_loc_ptr, cu_num_logits_ptr, pred_ptr, pred_stride,
  79. idx_mapping_ptr, tau, BLOCK: tl.constexpr):
  80. b = tl.program_id(0)
  81. num_drafts = tl.load(cu_num_logits_ptr + b + 1) - tl.load(cu_num_logits_ptr + b) - 1
  82. end = tl.load(query_start_loc_ptr + b + 1)
  83. state = tl.load(idx_mapping_ptr + b).to(tl.int64)
  84. j = tl.arange(0, BLOCK)
  85. in_range = j < num_drafts
  86. p = tl.load(pred_ptr + state * pred_stride + j, mask=in_range & (state >= 0), other=1.0)
  87. # Survival only falls along a request, so the dead drafts are a suffix.
  88. dead = (tl.cumprod(p, axis=0) < tau) & in_range
  89. tl.store(is_padding_ptr + end - num_drafts + j, tl.full((BLOCK,), 1, tl.int1), mask=dead)
  90. def mark_dead(is_padding: torch.Tensor, query_start_loc: torch.Tensor, cu_num_logits: torch.Tensor,
  91. num_reqs: int, max_drafts: int, predictions: torch.Tensor | None = None,
  92. idx_mapping: torch.Tensor | None = None) -> None:
  93. """Set is_padding on each request's dead drafts, per the control file."""
  94. global _steps, _cut
  95. _cut = False
  96. mode = _read_control()
  97. if mode is None or num_reqs == 0 or max_drafts == 0:
  98. return
  99. block = triton.next_power_of_2(max_drafts)
  100. if mode[0] == "live":
  101. _mark_dead_kernel[(num_reqs,)](is_padding, query_start_loc, cu_num_logits, int(mode[1]), BLOCK=block)
  102. _cut = True
  103. return
  104. if predictions is None or idx_mapping is None:
  105. return
  106. _steps += 1
  107. if _EXPLORE > 0 and _steps % _EXPLORE == 0:
  108. return
  109. _mark_cut_kernel[(num_reqs,)](is_padding, query_start_loc, cu_num_logits, predictions, predictions.stride(0),
  110. idx_mapping, mode[1], BLOCK=block)
  111. _cut = True
  112. def graded() -> bool:
  113. """Whether the last step verified every draft, so the estimator may learn from it."""
  114. return not _cut
  115. def log_step(confidence: torch.Tensor, cu_num_logits: torch.Tensor, num_sampled: torch.Tensor,
  116. estimator=None, idx_mapping: torch.Tensor | None = None) -> None:
  117. """Append this step's (drafts verified, sampled, confidences) per request to LOG,
  118. with the estimator's features and coefficients when it has one."""
  119. global _logged
  120. if _logged >= _LOG_MAX:
  121. return
  122. from vllm.distributed import get_tensor_model_parallel_rank
  123. if get_tensor_model_parallel_rank() != 0:
  124. return
  125. n = num_sampled.shape[0]
  126. drafts = (cu_num_logits[1:n + 1] - cu_num_logits[:n] - 1).tolist()
  127. feats = (estimator.features[idx_mapping].tolist() if estimator is not None and idx_mapping is not None
  128. else [None] * n)
  129. rows = [{"k": k, "sampled": s, "conf": [round(c, 4) for c in conf[:k]],
  130. "feat": None if f is None else [round(x, 3) for x in f[:k]]}
  131. for k, s, conf, f in zip(drafts, num_sampled.tolist(), confidence.tolist(), feats) if k > 0]
  132. if rows and estimator is not None and _logged % 200 == 0:
  133. rows.append({"slope": estimator.slope.tolist(), "intercepts": estimator.intercepts.tolist(),
  134. "refits": estimator._refits})
  135. if rows:
  136. import json
  137. with open(LOG, "a") as f:
  138. f.write(json.dumps(rows) + "\n")
  139. _logged += 1