File size: 11,682 Bytes
e1acc15 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179 180 181 182 183 184 185 186 187 188 189 190 191 192 193 194 195 196 197 198 199 200 201 202 203 204 205 206 207 208 209 210 211 212 213 214 215 216 217 218 219 220 221 222 223 224 225 226 227 228 229 230 231 232 233 234 235 236 237 238 239 240 241 242 243 244 245 246 247 248 249 250 251 252 253 254 255 256 257 258 259 260 261 262 263 264 265 266 267 268 269 270 271 272 | """Multi-round speculative generation through the image's real verifier, on CPU (AUD1 item 1).
Each trial is one request: a prompt token, then rounds of (deterministic drafts -> the image's
rejection_sample() -> commit) until `length` tokens are generated, exactly as the V2 model runner
commits them (num_computed_tokens += accepted + 1; the next round's first row is the last committed
token; row i is keyed by position = start + i, model_runner/input_batch.prepare_pos_seq_lens).
Production's remote-DSpark shape: draft_logits=None (one-hot q), standard verification, T=1,
top_p=1, an unseeded request seed (random int64), a per-round draft window (0047 adaptive window)
and depth-policy dead rows passed as -1 (model_runner.sample -> invalid_drafts).
The target is a context-dependent toy (a Markov chain on a small vocabulary) so the exact joint
law of the generated tokens is known; the drafter is a deterministic function of the committed
context. Only synthetic numbers; no model, prompt or generated text from production.
"""
from __future__ import annotations
import math
import multiprocessing as mp
import os
import random
from dataclasses import dataclass, field
from pathlib import Path
from typing import Any
import torch
from kernel_loader import DEFAULT_SRC, load
# vLLM #61161 hunk (rejection_sampler_utils.py, _rejection_kernel), applied to a copy for the
# positive control only. spec_step_seed is the PR's gumbel.py helper, defined in the copy.
_PRE = " u = tl_rand32(seed, pos, includes_zero=False)\n"
_POST = (
" step_seed = spec_step_seed(seed, i) if USE_BLOCK_VERIFICATION else seed\n"
" u = tl_rand32(step_seed, pos, includes_zero=False)\n"
)
_HELPER = (
"\n\n@triton.jit\n"
"def spec_step_seed(seed, step):\n"
" return seed ^ (tl.cast(step, tl.int64) << 32)\n"
)
_ANCHOR = "from vllm.v1.worker.gpu.sample.watermark import philox_gumbel_block_argmax\n"
def write_pr61161_variant(dest: Path, src_root: Path = DEFAULT_SRC) -> Path:
"""Copy the image's rejection_sampler_utils.py with #61161's kernel change (control only)."""
text = (src_root / "vllm/v1/worker/gpu/spec_decode/rejection_sampler_utils.py").read_text()
if text.count(_PRE) != 1 or text.count(_ANCHOR) != 1:
raise RuntimeError("image kernel differs from the audited pre-#61161 source")
out = dest / "rejection_sampler_utils_pr61161.py"
out.write_text(text.replace(_ANCHOR, _ANCHOR + _HELPER).replace(_PRE, _POST))
return out
@dataclass(frozen=True)
class Toy:
"""Context-dependent target p(x | prev), deterministic drafter, per-round window/depth."""
target: tuple[tuple[float, ...], ...] # rows: prev token -> distribution over the vocab
draft_next: tuple[int, ...] # deterministic proposal after token x
max_k: int # num_speculative_tokens (production: 5 at 1-2 streams, 4 at 3-4)
length: int # generated tokens per trial
vary_window: bool = True # 0047-style per-round window from committed history
dead_rows: bool = True # depth-policy rows past the depth -> -1
prompt_token: int = 0
prompt_pos: int = 37 # position of the prompt's last token (any offset)
@property
def vocab(self) -> int:
return len(self.target)
def joint(self) -> dict[tuple[int, ...], float]:
"""Exact law of the first `length` generated tokens."""
out: dict[tuple[int, ...], float] = {(): 1.0}
for _ in range(self.length):
nxt: dict[tuple[int, ...], float] = {}
for seq, pr in out.items():
prev = seq[-1] if seq else self.prompt_token
for x, px in enumerate(self.target[prev]):
if px > 0:
nxt[seq + (x,)] = pr * px
out = nxt
return out
# The review's first case (#61161 validation): context-free p=(.5,.5), drafter always 0, k=2.
PR_CASE = Toy(target=((0.5, 0.5), (0.5, 0.5)), draft_next=(0, 0), max_k=2, length=6,
vary_window=False, dead_rows=False)
# Production-shaped case: context-dependent target, a drafter that is right about half the time,
# windows 1..5 and dead rows chosen from committed history only.
PROD_CASE = Toy(target=((0.5, 0.3, 0.2), (0.2, 0.5, 0.3), (0.3, 0.2, 0.5)),
draft_next=(0, 2, 2), max_k=5, length=4)
@dataclass
class Result:
counts: dict[tuple[int, ...], int] = field(default_factory=dict)
trials: int = 0
rounds: int = 0
accepted: int = 0
# acceptance-uniform keys already consulted in an earlier round for a position that was not
# committed (the #61161 condition), per the kernel's read set. The key is (seed, position),
# or (seed ^ row << 32, position) for block rounds in the #61161-fixed copy.
reread_keys: int = 0
def merge(self, other: Result) -> None:
for key, value in other.counts.items():
self.counts[key] = self.counts.get(key, 0) + value
self.trials += other.trials
self.rounds += other.rounds
self.accepted += other.accepted
self.reread_keys += other.reread_keys
def _window(toy: Toy, seq: list[int], rnd: int) -> tuple[int, int]:
"""(window k, depth d <= k) from committed history only, like 0047 / the depth policy."""
if not toy.vary_window:
return toy.max_k, toy.max_k
h = (sum((i + 1) * t for i, t in enumerate(seq)) * 2654435761 + rnd * 40503) & 0xFFFF
k = 1 + h % toy.max_k
d = k if not toy.dead_rows else 1 + (h >> 4) % k
return k, d
def _run_chunk(args: tuple[Toy, str, int, int, str | None]) -> Result:
toy, mode, n, chunk_seed, variant = args
module = load(Path(variant) if variant else None)
step_keyed = variant is not None and mode == "block"
rng = random.Random(chunk_seed)
seeds = [rng.getrandbits(63) - (1 << 62) for _ in range(n)]
log_target = torch.tensor(toy.target, dtype=torch.float32).log()
seqs: list[list[int]] = [[toy.prompt_token] for _ in range(n)]
consulted: list[set[tuple[int, int]]] = [set() for _ in range(n)]
res = Result(trials=n)
rnd = 0
while True:
active = [r for r in range(n) if len(seqs[r]) - 1 < toy.length]
if not active:
break
rnd += 1
logits_rows: list[torch.Tensor] = []
draft_sampled: list[int] = []
pos: list[int] = []
cu = [0]
local_pos: list[int] = []
expanded: list[int] = []
for b, r in enumerate(active):
seq = seqs[r]
k, d = _window(toy, seq, rnd)
start = toy.prompt_pos + len(seq) - 1
row_tokens = [seq[-1]]
for _ in range(k):
row_tokens.append(toy.draft_next[row_tokens[-1]])
for i, tok in enumerate(row_tokens):
logits_rows.append(log_target[tok])
# rows past the depth are invalid drafts (-1), as model_runner.sample passes them
draft_sampled.append(-1 if (i > d) else tok)
pos.append(start + i)
local_pos.append(i)
expanded.append(b)
cu.append(cu[-1] + k + 1)
nb = len(active)
sampled, num_sampled = module.rejection_sample(
torch.stack(logits_rows).contiguous(),
None, # remote DSpark capacity: draft_logits None -> one-hot q
torch.tensor(draft_sampled, dtype=torch.int64),
torch.tensor(cu, dtype=torch.int32),
torch.tensor(pos, dtype=torch.int64),
torch.arange(nb, dtype=torch.int32),
torch.tensor(expanded, dtype=torch.int32),
torch.tensor(local_pos, dtype=torch.int32),
torch.ones(nb, dtype=torch.float32), # T=1
torch.tensor([seeds[r] for r in active], dtype=torch.int64),
toy.max_k,
use_block_verification=(mode == "block"),
)
for b, r in enumerate(active):
m = int(num_sampled[b])
k = cu[b + 1] - cu[b] - 1
start = pos[cu[b]]
# Rows whose acceptance uniform the kernel reads (rejection_sampler_utils.py
# _rejection_kernel): standard stops after the first rejection or -1; block reads every
# row up to the first -1. Rows are keyed by (seed, start + i).
valid = 0
while valid < k and draft_sampled[cu[b] + valid + 1] >= 0:
valid += 1
read_rows = valid if mode == "block" else min(m, valid)
for i in range(read_rows):
key = (i if step_keyed else 0, start + i)
if key in consulted[r]:
res.reread_keys += 1
consulted[r].add(key)
new = [int(t) for t in sampled[b, :m]]
seqs[r].extend(new)
res.accepted += m - 1
# committed positions can never be consulted again; drop them from the watch set
committed_upto = start + m - 1
consulted[r] = {c for c in consulted[r] if c[1] > committed_upto}
res.rounds += 1
for seq in seqs:
gen = tuple(seq[1 : 1 + toy.length])
res.counts[gen] = res.counts.get(gen, 0) + 1
return res
def simulate(toy: Toy, mode: str, trials: int, seed: int, workers: int = 8,
variant: Path | None = None) -> Result:
if mode not in ("standard", "block"):
raise ValueError(mode)
chunks = max(1, min(workers, trials))
sizes = [trials // chunks + (1 if i < trials % chunks else 0) for i in range(chunks)]
jobs = [(toy, mode, s, seed * 1000 + i, str(variant) if variant else None)
for i, s in enumerate(sizes) if s]
total = Result()
if len(jobs) == 1:
total.merge(_run_chunk(jobs[0]))
return total
ctx = mp.get_context("spawn")
env_ok = os.environ.get("TRITON_INTERPRET") == "1" and os.environ.get("CUDA_VISIBLE_DEVICES") == ""
if not env_ok:
raise RuntimeError("TRITON_INTERPRET=1 and CUDA_VISIBLE_DEVICES= are required")
with ctx.Pool(len(jobs)) as pool:
for part in pool.map(_run_chunk, jobs):
total.merge(part)
return total
def chi2_threshold(dof: int, alpha: float) -> float:
"""Wilson-Hilferty upper quantile of chi-square(dof)."""
z = _norm_ppf(1.0 - alpha)
c = 2.0 / (9.0 * dof)
return dof * (1.0 - c + z * math.sqrt(c)) ** 3
def _norm_ppf(p: float) -> float:
lo, hi = -10.0, 10.0
for _ in range(200):
mid = (lo + hi) / 2
if 0.5 * math.erfc(-mid / math.sqrt(2)) < p:
lo = mid
else:
hi = mid
return (lo + hi) / 2
def chi2_joint(toy: Toy, res: Result) -> tuple[float, int]:
exact = toy.joint()
stat = 0.0
for key, pr in exact.items():
e = pr * res.trials
stat += (res.counts.get(key, 0) - e) ** 2 / e
if any(key not in exact for key in res.counts):
stat = math.inf
return stat, len(exact) - 1
def marginals(toy: Toy, res: Result, token: int) -> list[tuple[float, float]]:
"""(observed P(x_t = token), exact) per generated position t."""
exact = toy.joint()
out = []
for t in range(toy.length):
obs = sum(c for k, c in res.counts.items() if k[t] == token) / res.trials
ex = sum(p for k, p in exact.items() if k[t] == token)
out.append((obs, ex))
return out
def summary(toy: Toy, res: Result) -> dict[str, Any]:
stat, dof = chi2_joint(toy, res)
return {"trials": res.trials, "rounds": res.rounds, "accepted_drafts": res.accepted,
"reread_keys": res.reread_keys, "chi2": round(stat, 2), "dof": dof,
"marginals_token0": [(round(o, 4), round(e, 4)) for o, e in marginals(toy, res, 0)]}
|