draft_trunc.py 6.9 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159
  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 commit, the image's VLLM_REF. On any other they would put back old code
  26. # over new, so the engine stops instead.
  27. VLLM_COMMIT = "ddd6fbca"
  28. if f"+g{VLLM_COMMIT}" not in vllm.__version__:
  29. raise RuntimeError(f"experimental/draft-trunc copies vLLM {VLLM_COMMIT}'s files, but this image runs "
  30. f"vLLM {vllm.__version__}. Refresh the copies from this image, or remove the "
  31. "draft-trunc mounts from experimental/compose/adaptive-k.yaml.")
  32. _TAU = float(os.environ.get("VLLM_DRAFT_TRUNC_TAU") or 0)
  33. ENABLED = bool(os.environ.get("VLLM_DRAFT_TRUNC_CONTROL")) or _TAU > 0
  34. # Measurement: per verify step, each request's draft confidences (the
  35. # acceptance estimator's, computed when the drafts were made) and how many
  36. # tokens it sampled, one JSON line per step, on tensor-parallel rank 0.
  37. LOG = os.environ.get("VLLM_DRAFT_TRUNC_LOG", "")
  38. _LOG_MAX = int(os.environ.get("VLLM_DRAFT_TRUNC_LOG_STEPS", "20000"))
  39. _logged = 0
  40. _control = os.environ.get("VLLM_DRAFT_TRUNC_CONTROL", "")
  41. _mtime = 0.0
  42. _mode: tuple[str, float] | None = None
  43. _EXPLORE = int(os.environ.get("VLLM_DRAFT_TRUNC_EXPLORE") or 4)
  44. _steps = 0
  45. _cut = False
  46. def _read_control() -> tuple[str, float] | None:
  47. """("live", N), ("tau", X), or None to truncate nothing."""
  48. global _mtime, _mode
  49. default = ("tau", _TAU) if _TAU > 0 else None
  50. try:
  51. mtime = os.stat(_control).st_mtime if _control else None
  52. except OSError:
  53. mtime = None
  54. if mtime is None:
  55. _mtime, _mode = 0.0, default
  56. return default
  57. if mtime != _mtime:
  58. _mtime = mtime
  59. words = open(_control).read().split()
  60. _mode = (words[0], float(words[1])) if len(words) == 2 and words[0] in ("live", "tau") else None
  61. return _mode
  62. def live_drafts() -> int | None:
  63. """The control file's N, or None."""
  64. mode = _read_control()
  65. return int(mode[1]) if mode and mode[0] == "live" else None
  66. @triton.jit
  67. def _mark_dead_kernel(is_padding_ptr, query_start_loc_ptr, cu_num_logits_ptr, live, BLOCK: tl.constexpr):
  68. b = tl.program_id(0)
  69. num_drafts = tl.load(cu_num_logits_ptr + b + 1) - tl.load(cu_num_logits_ptr + b) - 1
  70. end = tl.load(query_start_loc_ptr + b + 1)
  71. j = tl.arange(0, BLOCK)
  72. # A request's drafts are its last num_drafts tokens.
  73. tl.store(is_padding_ptr + end - num_drafts + j, tl.full((BLOCK,), 1, tl.int1),
  74. mask=(j >= live) & (j < num_drafts))
  75. @triton.jit
  76. def _mark_cut_kernel(is_padding_ptr, query_start_loc_ptr, cu_num_logits_ptr, pred_ptr, pred_stride,
  77. idx_mapping_ptr, tau, BLOCK: tl.constexpr):
  78. b = tl.program_id(0)
  79. num_drafts = tl.load(cu_num_logits_ptr + b + 1) - tl.load(cu_num_logits_ptr + b) - 1
  80. end = tl.load(query_start_loc_ptr + b + 1)
  81. state = tl.load(idx_mapping_ptr + b).to(tl.int64)
  82. j = tl.arange(0, BLOCK)
  83. in_range = j < num_drafts
  84. p = tl.load(pred_ptr + state * pred_stride + j, mask=in_range & (state >= 0), other=1.0)
  85. # Survival only falls along a request, so the dead drafts are a suffix.
  86. dead = (tl.cumprod(p, axis=0) < tau) & in_range
  87. tl.store(is_padding_ptr + end - num_drafts + j, tl.full((BLOCK,), 1, tl.int1), mask=dead)
  88. def mark_dead(is_padding: torch.Tensor, query_start_loc: torch.Tensor, cu_num_logits: torch.Tensor,
  89. num_reqs: int, max_drafts: int, predictions: torch.Tensor | None = None,
  90. idx_mapping: torch.Tensor | None = None) -> None:
  91. """Set is_padding on each request's dead drafts, per the control file."""
  92. global _steps, _cut
  93. _cut = False
  94. mode = _read_control()
  95. if mode is None or num_reqs == 0 or max_drafts == 0:
  96. return
  97. block = triton.next_power_of_2(max_drafts)
  98. if mode[0] == "live":
  99. _mark_dead_kernel[(num_reqs,)](is_padding, query_start_loc, cu_num_logits, int(mode[1]), BLOCK=block)
  100. _cut = True
  101. return
  102. if predictions is None or idx_mapping is None:
  103. return
  104. _steps += 1
  105. if _EXPLORE > 0 and _steps % _EXPLORE == 0:
  106. return
  107. _mark_cut_kernel[(num_reqs,)](is_padding, query_start_loc, cu_num_logits, predictions, predictions.stride(0),
  108. idx_mapping, mode[1], BLOCK=block)
  109. _cut = True
  110. def graded() -> bool:
  111. """Whether the last step verified every draft, so the estimator may learn from it."""
  112. return not _cut
  113. def log_step(confidence: torch.Tensor, cu_num_logits: torch.Tensor, num_sampled: torch.Tensor,
  114. estimator=None, idx_mapping: torch.Tensor | None = None) -> None:
  115. """Append this step's (drafts verified, sampled, confidences) per request to LOG,
  116. with the estimator's features and coefficients when it has one."""
  117. global _logged
  118. if _logged >= _LOG_MAX:
  119. return
  120. from vllm.distributed import get_tensor_model_parallel_rank
  121. if get_tensor_model_parallel_rank() != 0:
  122. return
  123. n = num_sampled.shape[0]
  124. drafts = (cu_num_logits[1:n + 1] - cu_num_logits[:n] - 1).tolist()
  125. feats = (estimator.features[idx_mapping].tolist() if estimator is not None and idx_mapping is not None
  126. else [None] * n)
  127. rows = [{"k": k, "sampled": s, "conf": [round(c, 4) for c in conf[:k]],
  128. "feat": None if f is None else [round(x, 3) for x in f[:k]]}
  129. for k, s, conf, f in zip(drafts, num_sampled.tolist(), confidence.tolist(), feats) if k > 0]
  130. if rows and estimator is not None and _logged % 200 == 0:
  131. rows.append({"slope": estimator.slope.tolist(), "intercepts": estimator.intercepts.tolist(),
  132. "refits": estimator._refits})
  133. if rows:
  134. import json
  135. with open(LOG, "a") as f:
  136. f.write(json.dumps(rows) + "\n")
  137. _logged += 1