verify_cap.py 14 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324
  1. """Per-step verify length for DSpark without changing the verify layout. Default OFF.
  2. The target verify runs the fixed [bs, 6] rectangle (anchor + 5 drafts). With DSV41_VERIFY_CAP set,
  3. each request gets a live length L = 1 + k (k = drafts actually verified):
  4. * acceptance is capped at k drafts through the engine's own cutoff (CapCorrectLen + bonus from
  5. row k), so the committed tokens are exactly those of a verify with k drafts;
  6. * the router ids of the dead rows (k+1..5) are replaced by the anchor row's ids, so the dead rows
  7. add no experts to the per-layer set the MoE streams (the saving; MoE is bandwidth-bound);
  8. * dead rows are never committed (commit <= k + 1), and live rows only attend to earlier rows, so
  9. their wrong MoE output cannot reach a committed token.
  10. k comes from DSV41_VERIFY_CAP:
  11. "1".."5" fixed k for every step (5 = all live, measures the machinery's own cost)
  12. "conf:T" k = number of leading drafts whose running product of the draft confidence head's
  13. per-position survival stays >= T (at least DSV41_VERIFY_CAP_MIN, default 1).
  14. The head's value for position i is computed from drafts before i only, so the
  15. cut is a stopping rule: it never looks at the token it drops (exact for sampling).
  16. Sampled rows go through block_verify, which treats dead positions as q = 0, p(X) = 0.
  17. """
  18. import os
  19. import torch
  20. SPEC = os.environ.get("DSV41_VERIFY_CAP", "").strip()
  21. ENABLED = SPEC not in ("", "0", "off")
  22. # verify rows per request = DSPARK block size + 1 (anchor); 6 for the stock block of 5
  23. STRIDE = int(os.environ.get("DSV41_VERIFY_CAP_STRIDE") or int(os.environ.get("DSPARK_BLOCK_SIZE", "5")) + 1)
  24. MAX_BS = 64
  25. E_TARGET, K_TARGET = 384, 6
  26. _state = {"live": None, "in_verify": False, "conf": None, "fixed": None, "thr": None,
  27. "kmin": int(os.environ.get("DSV41_VERIFY_CAP_MIN", "1")), "logged": False}
  28. if ENABLED:
  29. if SPEC.startswith("conf:"):
  30. # one threshold for every draft, or one per draft 2..5 ("conf:0.25,0.25,0.03,0.06");
  31. # draft 1 is always verified (DSV41_VERIFY_CAP_MIN, default 1)
  32. vals = [float(v) for v in SPEC.split(":", 1)[1].split(",")]
  33. _state["thr"] = vals[0] if len(vals) == 1 else [0.0] + vals
  34. else:
  35. _state["fixed"] = int(SPEC)
  36. def _capturing():
  37. return torch.cuda.is_available() and torch.cuda.is_current_stream_capturing()
  38. def live_buf(device=None):
  39. buf = _state["live"]
  40. if buf is None:
  41. assert not _capturing(), "verify_cap: live buffer must exist before capture"
  42. buf = torch.full((MAX_BS,), STRIDE, dtype=torch.int64,
  43. device=device if device is not None else torch.cuda.current_device())
  44. if _state["fixed"] is not None:
  45. buf.fill_(1 + _state["fixed"])
  46. _state["live"] = buf
  47. return buf
  48. def cutoff(bs):
  49. return live_buf()[:bs] if ENABLED else None
  50. def set_live_from_confidence(confidence, bs):
  51. """confidence [bs, 5] per-position survival (position i given drafts < i)."""
  52. if _state["thr"] is None or confidence is None:
  53. return
  54. buf = live_buf(confidence.device)
  55. _state["last_conf"] = confidence
  56. cum = torch.cumprod(confidence.float().clamp(0, 1), dim=1)
  57. thr = _state["thr"]
  58. if isinstance(thr, list):
  59. if _state.get("thr_t") is None or _state["thr_t"].device != cum.device:
  60. t = (thr + [thr[-1]] * cum.shape[1])[:cum.shape[1]]
  61. _state["thr_t"] = torch.tensor(t, dtype=torch.float32, device=cum.device)
  62. thr = _state["thr_t"]
  63. k = (cum >= thr).to(torch.int64).cumprod(dim=1).sum(dim=1).clamp(min=_state["kmin"], max=STRIDE - 1)
  64. buf[:bs].copy_(k + 1)
  65. # Rank-invariant by construction: the head's input comes out of an all-reduce that may differ
  66. # in the last bits between ranks, and a live length that differs across ranks would give one
  67. # verify row two expert sets. Rank 0 decides, as for the engine's own accept (SpecTpSync).
  68. group = _state.get("tp_group")
  69. if group is None:
  70. from sglang.srt.distributed import get_tp_group
  71. group = _state["tp_group"] = get_tp_group()
  72. if group.world_size > 1:
  73. group.broadcast(buf[:bs], src=0)
  74. def _src_rows(m, device):
  75. """Row each verify row takes its router output from (itself if live, else its anchor).
  76. Computed once per forward and reused by all 43 layers (inside the graph: once per replay)."""
  77. src = _state.get("src")
  78. if src is not None and src.shape[0] == m:
  79. return src
  80. live = live_buf(device)
  81. rows = torch.arange(m, device=device)
  82. req, pos = rows // STRIDE, rows % STRIDE
  83. dead = pos >= live[req.clamp(max=MAX_BS - 1)]
  84. src = torch.where(dead, req * STRIDE, rows)
  85. _state["src"] = src
  86. if _capturing() and m not in _state.setdefault("captured", set()):
  87. _state["captured"].add(m)
  88. print(f"[verify_cap] remap captured into the verify graph for M={m}", flush=True)
  89. return src
  90. try:
  91. import triton
  92. import triton.language as tl
  93. @triton.jit
  94. def _remap_kernel(W, I, LIVE, M, SW: tl.constexpr, SI: tl.constexpr, STRIDE_R: tl.constexpr,
  95. K: tl.constexpr, BLOCK: tl.constexpr):
  96. r = tl.arange(0, BLOCK)
  97. rmask = r < M
  98. req = r // STRIDE_R
  99. pos = r % STRIDE_R
  100. live = tl.load(LIVE + req, rmask, STRIDE_R)
  101. dead = rmask & (pos >= live)
  102. anchor = req * STRIDE_R
  103. for k in tl.static_range(K):
  104. w = tl.load(W + anchor * SW + k, dead, 0.0)
  105. i = tl.load(I + anchor * SI + k, dead, 0)
  106. tl.store(W + r * SW + k, w, dead)
  107. tl.store(I + r * SI + k, i, dead)
  108. except Exception: # no triton: fall back to index_select
  109. _remap_kernel = None
  110. def remap_dead_rows(weights, indices):
  111. """weights/indices [M, 6] for M = bs * 6 verify rows (request-major). In place when possible:
  112. one kernel per layer copies the anchor row's ids and weights over the dead rows."""
  113. m = indices.shape[0]
  114. if (_remap_kernel is not None and weights.is_contiguous() and indices.is_contiguous()
  115. and weights.shape[1] == K_TARGET and m <= 1024):
  116. if _capturing() and m not in _state.setdefault("captured", set()):
  117. _state["captured"].add(m)
  118. print(f"[verify_cap] in-place remap captured into the verify graph for M={m}", flush=True)
  119. _remap_kernel[(1,)](weights, indices, live_buf(indices.device), m, weights.stride(0),
  120. indices.stride(0), STRIDE, K_TARGET, triton.next_power_of_2(m))
  121. return weights, indices
  122. src = _src_rows(m, indices.device)
  123. return weights.index_select(0, src), indices.index_select(0, src)
  124. def install_gate(module):
  125. """sglang.kernels.ops.moe.moe_fused_gate"""
  126. if not ENABLED or getattr(module, "_dsv41_verify_cap", False):
  127. return
  128. module._dsv41_verify_cap = True
  129. original = module.moe_fused_gate
  130. fused = None
  131. if (os.environ.get("DSV41_LAUNCHER") == "tp4"
  132. and os.environ.get("DSV41_ROUTER_LIVE", "0").strip() not in ("0", "", "off", "false")):
  133. # DSV41_ROUTER_LIVE: the router loads the anchor's scores for dead rows (no remap launch)
  134. from router_live import build
  135. fused = build(module)
  136. print("[verify_cap] dead-row remap folded into the router kernel", flush=True)
  137. def wrapped(scores, bias, topk, *a, **kw):
  138. m = scores.shape[0] if scores.dim() == 2 else 0
  139. verify = (_state["in_verify"] and scores.shape[-1] == E_TARGET and topk == K_TARGET
  140. and m and m % STRIDE == 0 and m // STRIDE <= MAX_BS)
  141. if verify and fused is not None:
  142. return fused(scores, bias, topk, *a, live=live_buf(scores.device), live_stride=STRIDE, **kw)
  143. weights, indices = original(scores, bias, topk, *a, **kw)
  144. if verify:
  145. return remap_dead_rows(weights, indices)
  146. return weights, indices
  147. module.moe_fused_gate = wrapped
  148. print(f"[verify_cap] armed ({SPEC}): dead verify rows route to the anchor's experts", flush=True)
  149. def install_model(dsv4_module):
  150. """sglang.srt.models.deepseek_v4: mark target-verify forwards (eager and graph capture)."""
  151. if not ENABLED:
  152. return
  153. cls = dsv4_module.DeepseekV4ForCausalLM
  154. if getattr(cls, "_dsv41_verify_cap", False):
  155. return
  156. cls._dsv41_verify_cap = True
  157. orig = cls.forward
  158. def forward(self, input_ids, positions, forward_batch, *a, **kw):
  159. if _state["live"] is None and not _capturing():
  160. live_buf(input_ids.device)
  161. prev = _state["in_verify"]
  162. _state["src"] = None
  163. _state["in_verify"] = (type(self).__name__ == "DeepseekV4ForCausalLM"
  164. and forward_batch.forward_mode.is_target_verify())
  165. try:
  166. return orig(self, input_ids, positions, forward_batch, *a, **kw)
  167. finally:
  168. _state["in_verify"] = prev
  169. cls.forward = forward
  170. class _Cut:
  171. def __init__(self, verify_lens):
  172. self.verify_lens = verify_lens
  173. def install_verify(verify_module):
  174. """sglang.srt.speculative.dspark_components.dspark_verify"""
  175. if not ENABLED:
  176. return
  177. ep = verify_module.DsparkVerifyEpilogue
  178. if getattr(ep, "_dsv41_verify_cap", False):
  179. return
  180. ep._dsv41_verify_cap = True
  181. orig_accept = ep._accept
  182. def _accept(self, *, candidates, logits, draft_tokens, seq_lens, cutoff_verify_lens=None):
  183. if cutoff_verify_lens is None:
  184. cutoff_verify_lens = cutoff(candidates.shape[0])
  185. return orig_accept(self, candidates=candidates, logits=logits, draft_tokens=draft_tokens,
  186. seq_lens=seq_lens, cutoff_verify_lens=cutoff_verify_lens)
  187. ep._accept = _accept
  188. orig_adt = verify_module.accept_draft_tokens
  189. def accept_draft_tokens(*, candidates, cutoff_layout=None, **kw):
  190. if cutoff_layout is None:
  191. cutoff_layout = _Cut(cutoff(candidates.shape[0]))
  192. return orig_adt(candidates=candidates, cutoff_layout=cutoff_layout, **kw)
  193. verify_module.accept_draft_tokens = accept_draft_tokens
  194. ex = verify_module.TargetVerifyExecutor
  195. orig_run = ex.run_non_compact
  196. def run_non_compact(self, *, batch, draft_input, verify_ids_2d, **kw):
  197. bs = verify_ids_2d.shape[0]
  198. if _state["thr"] is not None:
  199. set_live_from_confidence(_state["conf"], bs)
  200. if not _state["logged"]:
  201. _state["logged"] = True
  202. print(f"[verify_cap] confidence {'present' if _state['conf'] is not None else 'MISSING'}",
  203. flush=True)
  204. _state["conf"] = None
  205. return orig_run(self, batch=batch, draft_input=draft_input, verify_ids_2d=verify_ids_2d, **kw)
  206. ex.run_non_compact = run_non_compact
  207. # DSV41_VERIFY_CAP_LOG=/state/vcap-log.bin (rank 0): per verify step and request, the confidence
  208. # vector, the live length used and the accepted drafts. Costs a device sync per step: measurement only.
  209. log_path = os.environ.get("DSV41_VERIFY_CAP_LOG", "")
  210. if log_path and _state["thr"] is not None:
  211. orig_aaf = ex.accept_and_finalize
  212. _state["log"] = None
  213. def accept_and_finalize(self, *a, **kw):
  214. outs = orig_aaf(self, *a, **kw)
  215. conf = _state.get("last_conf")
  216. if conf is not None:
  217. import numpy as np
  218. if _state["log"] is None:
  219. from sglang.srt.distributed import get_tp_group
  220. _state["log"] = open(log_path, "ab") if get_tp_group().rank_in_group == 0 else False
  221. if _state["log"]:
  222. bs = conf.shape[0]
  223. rec = torch.cat([conf.float(), live_buf()[:bs].float().unsqueeze(1),
  224. outs.correct_len[:bs].float().unsqueeze(1)], dim=1)
  225. _state["log"].write(rec.cpu().numpy().astype(np.float32).tobytes())
  226. _state["log"].flush()
  227. return outs
  228. ex.accept_and_finalize = accept_and_finalize
  229. def install_draft(draft_module):
  230. """sglang.srt.speculative.dspark_components.dspark_draft: keep a folded proposal's confidence."""
  231. if not ENABLED or _state["thr"] is None:
  232. return
  233. cls = draft_module.DraftBlockProposer
  234. orig = cls.propose
  235. def propose(self, *a, **kw):
  236. out = orig(self, *a, **kw)
  237. if getattr(out, "confidence", None) is not None:
  238. _state["conf"] = out.confidence
  239. return out
  240. cls.propose = propose
  241. def install_planner(planner_module):
  242. """sglang.srt.speculative.dspark_components.dspark_planner: keep the step's confidence."""
  243. if not ENABLED or _state["thr"] is None:
  244. return
  245. cls = planner_module.DSparkVerifyPlanner
  246. orig = cls.compute_confidence_tensor
  247. def compute_confidence_tensor(self, **kw):
  248. out = orig(self, **kw)
  249. _state["conf"] = out
  250. return out
  251. cls.compute_confidence_tensor = compute_confidence_tensor
  252. def install_dspark(dspark_module):
  253. """sglang.srt.models.deepseek_v4_dspark: the engine builds the draft's confidence head only in
  254. the ragged-verify modes; build it in static mode too when the conf policy needs it."""
  255. if not ENABLED or _state["thr"] is None:
  256. return
  257. orig = dspark_module.build_dspark_v4_confidence_head
  258. mode_fn = dspark_module.read_ragged_verify_mode
  259. from sglang.srt.speculative.ragged_verify import RaggedVerifyMode
  260. def build(*a, **kw):
  261. dspark_module.read_ragged_verify_mode = lambda: RaggedVerifyMode.CAP_ACCEPT
  262. try:
  263. head = orig(*a, **kw)
  264. finally:
  265. dspark_module.read_ragged_verify_mode = mode_fn
  266. print(f"[verify_cap] draft confidence head built: {head is not None}", flush=True)
  267. return head
  268. dspark_module.build_dspark_v4_confidence_head = build