Commit e3bb5f5
[Feature] add synthetic mode for speculative decode (vllm-project#12256)
What this PR does / why we need it
Ports upstream vLLM's rejection_sample_method="synthetic"
(controlled-acceptance-rate spec-decode benchmarking) to vLLM-Ascend.
Bug fixed: AscendRejectionSampler never passed spec_config/device to the
base class, so synthetic_mode was never set — the synthetic config was
silently ignored and standard rejection sampling ran
instead.
Changes (additive only — non-synthetic paths byte-for-byte unchanged)
- worker/model_runner_v1.py: pass spec_config, device into
AscendRejectionSampler so base init populates
synthetic_mode/synthetic_conditional_rates (fp32, on-device, via
upstream's
unconditional_to_conditional_rates).
- sample/rejection_sampler.py: forward synthetic_* through
init/forward/rejection_sample; generate uniform_probs_for_greedy
(fp64→fp32 for Ascend's fp32-only triton); add synthetic branches to the
3
pyTorch fallbacks; disable block_verify under synthetic with
warning_once (its joint-cumprod verification conflicts with synthetic's
per-token first-rejection Bernoulli, corrupting the acceptance
distribution).
- ops/triton/reject_sample.py: add SYNTHETIC_MODE: tl.constexpr branches
to the greedy (spec_len=1 + general) and random kernels; thread the new
params through the triton wrapper. block_verify kernel
untouched.
Correctness gotchas: fp64→fp32 at triton entry; int32/int64 mismatch
handled (draft int32 vs target_argmax int64) via separate stores /
tl.where with cast; recovered_token_ids reused as global ids under
enable_reduce_sample.
User-facing change
Opt-in only. Set rejection_sample_method="synthetic" (+
synthetic_acceptance_length/rates) in --speculative-config. No default
change — synthetic_mode=False makes all new branches dead
(tl.constexpr-eliminated). Only side effect: modified triton kernels
recompile once (output identical). block_verify auto-disabled with
warning; reduce_sample/entropy_verify silently bypassed per-token.
Testing
- Static: py_compile on all 3 files; AST undefined-name check; git apply
--check clean; call-site audit.
- Runtime (NPU, TP=16): synthetic_acceptance_length=3,
num_speculative_tokens=3 → metrics reported Mean acceptance length:
3.00, per-position 1.000, 1.000, 0.000, draft rate 66.7% — matches
config
([1,1,0], (3−1)/3). Reproduced & fixed the int32/int64 triton ternary
compile error.
- Pending: synthetic × reduce_sample / entropy_verify / block_verify,
mixed greedy+random batch, pyTorch fallback path, and synthetic-off
regression matrix.
vllm serve <model> --speculative-config \
'{"method":"eagle","model":"","num_speculative_tokens":3,"rejection_sample_method":"synthetic","synthetic_acceptance_length":3.0}'
- vLLM version: v0.25.1
- vLLM main:
vllm-project/vllm@54503ec
---------
Signed-off-by: chengduxiaowu <wuchengtai@huawei-partners.com>
Co-authored-by: chengduxiaowu <wuchengtai@huawei-partners.com>1 parent 3d077f9 commit e3bb5f5
6 files changed
Lines changed: 529 additions & 34 deletions
File tree
- .github/workflows/configs
- docs/source/user_guide/feature_guide
- tests/e2e/nightly/single_node/ops/singlecard_ops/triton
- vllm_ascend
- ops/triton
- sample
- worker
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
31 | 31 | | |
32 | 32 | | |
33 | 33 | | |
| 34 | + | |
| 35 | + | |
| 36 | + | |
34 | 37 | | |
35 | 38 | | |
36 | 39 | | |
| |||
Lines changed: 54 additions & 1 deletion
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
187 | 187 | | |
188 | 188 | | |
189 | 189 | | |
190 | | - | |
| 190 | + | |
191 | 191 | | |
192 | 192 | | |
193 | 193 | | |
| |||
337 | 337 | | |
338 | 338 | | |
339 | 339 | | |
| 340 | + | |
| 341 | + | |
| 342 | + | |
| 343 | + | |
| 344 | + | |
| 345 | + | |
| 346 | + | |
| 347 | + | |
| 348 | + | |
| 349 | + | |
| 350 | + | |
| 351 | + | |
| 352 | + | |
| 353 | + | |
| 354 | + | |
| 355 | + | |
| 356 | + | |
| 357 | + | |
| 358 | + | |
| 359 | + | |
| 360 | + | |
| 361 | + | |
| 362 | + | |
| 363 | + | |
| 364 | + | |
| 365 | + | |
| 366 | + | |
| 367 | + | |
| 368 | + | |
| 369 | + | |
| 370 | + | |
| 371 | + | |
| 372 | + | |
| 373 | + | |
| 374 | + | |
| 375 | + | |
| 376 | + | |
| 377 | + | |
| 378 | + | |
| 379 | + | |
| 380 | + | |
| 381 | + | |
| 382 | + | |
| 383 | + | |
| 384 | + | |
| 385 | + | |
| 386 | + | |
| 387 | + | |
| 388 | + | |
| 389 | + | |
| 390 | + | |
| 391 | + | |
| 392 | + | |
0 commit comments