[ROCm] EAGLE: a malformed draft probability must reject, not accept

[JohnQinAMD]

Follow-up to #58, which landed before this commit was pushed. #58 guards q in the two residual passes; this guards it in the accept test, which is the site where a corrupt row does the most damage.

Problem

q = tl.load(DraftProbs + offset_draft)
if coin * q < p:      # accept

coin * q < p is -inf < p for an -inf q and 0 < p for a zero one. Either way the step accepts unconditionally and commits the draft head's token instead of the target's. Zero also passes the (q >= 0) & (q <= 1) range test #58 added to the residual passes, so that guard does not cover this site.

The failure is invisible from outside: acceptance climbing toward 1.0 reads as a good draft model, not as a broken test.

Fix

X was sampled from q, so q(X) has to be strictly positive. Anything else means the row is not the distribution X came from — reject, which sends the step through the residual path and resamples from the target.

q_is_prob = (q > 0.0) & (q <= 1.0)
if q_is_prob & (coin * q < p):

Scope

This code is only reached when EAGLE verify actually samples, which on ROCm requires sgl-project#37134 (_verify_uses_greedy). On this repo's main, which does not carry that patch, _is_hip in eagle_utils.py still forces the argmax branch unconditionally and the chain sampler is dead code — so this lands ahead of the path it protects, same as the rest of #58.

Why it matters

Measured on a tree that does carry sgl-project#37134 (GLM-5.2-FP8, 8x MI355X, TP8, NEXTN steps=1 draft=2, GPQA-Diamond 198q, temperature=1.0 top_p=0.95, max_tokens=98304):

	verify takes argmax (no sgl-project#37134)	sgl-project#37134 + #55 + #58 + this
requests that reach a final answer	144/198 = 72.7%	190/198 = 96.0%
responses that loop until the token cap	59 (29.8%)	0
tokens burned on capped requests	~5.3M	~0.8M
GPQA-Diamond	0.7222	0.8939

The same benchmark with speculative decoding switched off entirely scores 0.904 at 96.0% termination, so the fixed speculative path is back to parity with no speculation.

Cost, single request, 512 tokens with ignore_eos, same hardware: -2.2% throughput at temperature=1.0 (125.98 vs 128.76 tok/s), and nothing measurable at temperature=0 (135.28 vs 135.66) where is_all_greedy sends verify down the argmax path regardless.

🤖 Generated with Claude Code

---

CI States

Latest PR Test (Base): ❌ Run #34720143350
Latest PR Test (Extra): ❌ Run #34720143222
Latest PR Test (AMD ROCm 10): ❌ Run #34720143421