Repository navigation
perf(spec-decode): fuse the unfiltered per-row-noise draw into two launches (ATOM_SPEC_BONUS_FUSED) - #2914
perf(spec-decode): fuse the unfiltered per-row-noise draw into two launches (ATOM_SPEC_BONUS_FUSED)#2914chelnnexy wants to merge 8 commits into
Conversation
🏷️ CI GuideRuns automatically on every eligible PR before approval:
Heavy model tests:
|
Validation on MI355X (gfx950)Ran on an 8×MI355X host ( Unit tests
GPU bit-exact (real Triton kernels on CUDA, not the interpreter)Compared PASS (24/24 rows) — greedy rows bit-exact, sampling rows matched with no near-tie exceptions needed at these seeds. End-to-end (both decode bands)Full GSM8K 1319, 5-shot, 64 concurrency,
Important nuance: the fused path is gated on The decisive correctness evidence is the GPU bit-exact 24/24 and the 3200/3200 agreement above, which run the actual fused path. |
|
@valar-cop review this pr |
valar-cop
left a comment
There was a problem hiding this comment.
some changes needed..
e2e with the fused path armed (temperature>0 or n>1)? and a gfx942 run for the new kernel?
| logits_ptr + r.to(tl.int64) * stride_l + offs, mask=mask, other=-float("inf") | ||
| ).to(tl.float32) | ||
| temp = tl.load(temps_ptr + r).to(tl.float32) | ||
| inv_t = 1.0 / tl.maximum(temp, 1e-5) |
There was a problem hiding this comment.
all -inf tile here gives -inf - -inf = NaN.. reachable with required tool_choice, please guard m == -inf and add a test
There was a problem hiding this comment.
Confirmed on MI355X (gfx950) and MI300X (gfx942): the current kernel gives the same NaN result on both. The fix follows @yitingw1's version: guard the tile max and remove the clamp. Full results are in the top-level comment.
Before the fix (PR head 680d7359), gpu_check.py:
[1] NaN partial keys on a forced row: 78 / 80 (unpatched: >0 expected = bug present; patched: must be 0)
[2] forced rows wrong: 0 / 7
[3] masked rows: 0 unexplained mismatches, 0 near-ties, 0 out-of-range ids, over 1600 rows
[4] all -inf rows out of range: 0 / 4 [163839, 163839, 0, 163839]
FAIL (1 check(s))
There are NaN keys in 78 of 80 tiles for a forced row. The forced id still wins only because tl.max drops NaN. Separately, on a row that is -inf throughout, the sampling result is 163839: that is the vocab sentinel after tl.minimum(tok, vocab - 1), which shows the clamp is what's hiding it.
Fix:
@@ _bonus_tile_kernel
- key = tl.where(mask, tl.exp(xs - m) / e, -1.0)
+ m_safe = tl.where(m == -float("inf"), 0.0, m)
+ key = tl.where(mask, tl.exp(xs - m_safe) / e, -1.0)
@@ _bonus_pick_kernel
mm = tl.max(m, axis=0)
+ mm = tl.where(mm == -float("inf"), 0.0, mm)
ks = tl.where(tm, k * tl.exp(m - mm), -1.0)
@@
- tok = tl.minimum(tok, vocab - 1)- Cost: one scalar
whereper program in each kernel. The per-element work is unchanged. - Rows with no
-inflogits: results are unchanged. Checked against the current kernel on 80 rows at vocab 1000 to 163840. - All-
-inftile: keys every id at 0, so it loses to any tile with a finite logit, and itsidxstays in range. - Pick kernel: the
mmguard covers a row that is-infthroughout. That row now returns id 0 instead of the clamped sentinel. - The clamp: with the guard in place it can no longer change a result, so it's removed. Any future escape will now show up as an out-of-range id instead of a plausible but wrong token.
Tests in tests/test_bonus_sample.py, with RuntimeWarning turned into an error:
test_a_forced_channel_row_samples_the_forced_id: forced id at 5 / 100000 / 163837, for both T=1.0 and greedy rows.test_masked_ids_never_win_and_every_id_is_in_range: whole-inftiles plus scattered-infids, compared against the float64 race, plus one row that is-infthroughout.
Results under the interpreter: the current kernel fails 4 of these, and with the fix the file gives 13 passed, 1 skipped. black and ruff report nothing.
After the fix, gpu_check.py:
MI355X (gfx950):
[1] NaN partial keys on a forced row: 0 / 80 (unpatched: >0 expected = bug present; patched: must be 0)
[2] forced rows wrong: 0 / 7
[3] masked rows: 0 unexplained mismatches, 0 near-ties, 0 out-of-range ids, over 1600 rows
[4] all -inf rows out of range: 0 / 4 [0, 0, 0, 0]
PASS
MI300X (gfx942):
[1] NaN partial keys on a forced row: 0 / 80 (unpatched: >0 expected = bug present; patched: must be 0)
[2] forced rows wrong: 0 / 7
[3] masked rows: 0 unexplained mismatches, 0 near-ties, 0 out-of-range ids, over 1600 rows
[4] all -inf rows out of range: 0 / 4 [0, 0, 0, 0]
PASS
Pushed as e34da58. The branch is also rebased onto main, which brings in the loader fix from #2936.
| # Rows with top-k/top-p, and batches without per-row noise, keep the aiter | ||
| # path. Read when the `Sampler` is built, so it must be set before the | ||
| # server starts. | ||
| "ATOM_SPEC_BONUS_FUSED": lambda: os.getenv("ATOM_SPEC_BONUS_FUSED", "0") == "1", |
There was a problem hiding this comment.
why default off? the sibling ATOM_DSPARK_FUSED_MARKOV_SAMPLE is on by default..
There was a problem hiding this comment.
It stays off on purpose. The reason is now in the flag's comment and in environment_variables.md (c70d5bf):
- Not bit-exact: sampling rows can differ from the aiter kernel on fp32 near-ties (~1e-7 relative).
- The gain shrinks with verify rows: on MI355X it's about −25% up to 128 rows, but nothing at 256 (about Kimi-K3 DSpark at concurrency 64).
ATOM_DSPARK_FUSED_MARKOV_SAMPLE doesn't have that trade-off. The tuned aiter version (see the aiter thread) is faster at 256 rows too; switching to it is the point to revisit the default.
|
+1 on valar-cop's NaN finding — it is reachable exactly as described. Two refinements that change what the fix should be, plus a few things not yet raised. 1. The NaN is real and reachable, but it does not produce a wrong token today. Worth knowing before you fix it, so you don't change a result that is currently correct. The greedy half is genuinely safe and needs no guard: 2. So the thing to remove is the clamp, not just the NaN. m = tl.max(xs, axis=0)
m_safe = tl.where(m == -float("inf"), 0.0, m)
key = tl.where(mask, tl.exp(xs - m_safe) / e, -1.0)An all- 3. No test uses a 4. The microbench percentages have a cost excluded from the denominator that is paid on the same critical path every step. Excluding the 5. 6. The aiter side finding will get lost where it is. " 7. On "why default off" — the answer is already in the PR body, it just isn't wired to the question. Sampling rows are not bit-exact (the ~1e-7 near-tie divergence), unlike Things I checked and found correct, so nobody re-derives them: the two-stage rescale identity ( |
…unches (ATOM_SPEC_BONUS_FUSED) The speculative bonus token, and every draft-conditioned verification row of a non-greedy batch, is an unfiltered draw with per-row noise. It ran as mixed_sample_outer_exponential + topk_select + <= + where. With ATOM_SPEC_BONUS_FUSED=1 (default off) it runs as two vocab-tiled Triton launches: a (row, tile) kernel keeping each tile's best race key and raw-logit argmax, and a (row,) kernel combining the tiles. Same exponential_ noise, same race argmax_i p_i / (E_i + eps), so the distribution is unchanged; greedy rows are bit-identical (lowest id on ties). Rows with top-k/top-p and batches without per-row noise keep the aiter path. Tested on MI350X (gfx950) and on CPU under the Triton interpreter; Sampler routing checked against a fake aiter.
Add two interpreter tests to test_bonus_sample.py: - same (logits, temps, noise) returns the same token every call (the kernels hold no state of their own); - a dominating race key in the final vocab tile beats tile 0, exercising the per-tile running max and the (row,) kernel's k * exp(m - mm) rescale.
With tool_choice="required" the bonus and verification rows are masked to -inf everywhere but the forced id, so 39 of 40 vocab tiles have a tile max of -inf and exp(-inf - -inf) put NaN into the race keys. The forced id only won because tl.max drops NaN. Shift by 0 when the tile max (and the pick kernel's global max) is -inf, so no NaN reaches a reduction. Remove tl.minimum(tok, vocab - 1): with the guard it can no longer change a result, and it turned an escaped sentinel into a valid-looking token id. Add tests for forced-channel rows and for partly and fully masked rows.
680d735 to
e34da58
Compare
MI355X (gfx950) and MI300X (gfx942) results for the review asksSummary
All runs have the NaN fix applied, now pushed as MI355X (gfx950)Image: End-to-end run that actually reaches the fused path (@valar-cop)Kimi-K3 + DSpark, Note: before the rebase, the branch didn't start in this image. The weight loader failed with Kernel:
|
_temperature_sample and the ATOM_SPEC_BONUS_FUSED branch of forward each spelled the (num_tokens, vocab_size) exponential_(1) draw out by hand. Both now call get_per_row_exponential, so the two paths cannot drift apart and keep consuming the RNG identically.
| import pytest | ||
| import torch | ||
|
|
||
| pytest.importorskip("triton") |
There was a problem hiding this comment.
stub bonus_sample in the fixture like aiter, drop importorskip so this runs on cpu?
There was a problem hiding this comment.
Done in a8347e8. The fixture now fakes atom.model_ops.bonus_sample the same way it fakes aiter, and the importorskip is gone. A new test checks that the flag wires Sampler to bonus_temperature_sample.
With Triton blocked (as on a CPU runner) all 7 tests pass. Before, the whole file was skipped.
| n_tiles = triton.cdiv(vocab, _BONUS_BLOCK_V) | ||
| f32 = {"dtype": torch.float32, "device": logits.device} | ||
| i32 = {"dtype": torch.int32, "device": logits.device} | ||
| part_m = torch.empty((n_tiles, rows), **f32) |
There was a problem hiding this comment.
still allocate 5 buffers per call on the hot path.. cache them keyed by (n_tiles, rows)?
There was a problem hiding this comment.
I tried it and measured it, and kept torch.empty.
A cached buffer, keyed by (device, stream) and grown on demand, made the draw slower: about +2.2 µs per call at every row count on MI355X. The main path stayed within ±0.4 µs between the two runs, so this is not run-to-run noise. The cost is in Python: looking up the current stream builds a Stream object on every call, and slicing views out of the buffer creates as many tensor objects as torch.empty does. The memory itself is not allocated per call. PyTorch's caching allocator already hands back the same blocks, so there is nothing to save there.
MI355X, median µs of 300 calls:
| rows | main path | fused, torch.empty |
fused, cached buffer |
|---|---|---|---|
| 1 | 60.08 → 59.64 | 44.76 | 46.88 (+2.12) |
| 4 | 61.42 → 61.08 | 44.60 | 46.80 (+2.20) |
| 8 | 62.36 → 62.12 | 44.68 | 46.96 (+2.28) |
| 16 | 62.60 → 62.24 | 44.64 | 46.66 (+2.02) |
| 32 | 64.08 → 64.12 | 45.04 | 47.36 (+2.32) |
| 64 | 66.88 → 66.52 | 43.16 | 45.36 (+2.20) |
| 128 | 70.84 → 70.66 | 54.06 | 57.08 (+3.02) |
| 256 | 78.28 → 78.28 | 78.80 | 81.44 (+2.64) |
So the five torch.empty calls stay.
|
|
||
|
|
||
| @triton.jit | ||
| def _bonus_tile_kernel( |
There was a problem hiding this comment.
this is a second impl of aiter's mixed_sample_outer_exponential.. why not put it in aiter and call here?
There was a problem hiding this comment.
Agreed, and it's moved: ROCm/aiter#7416 adds it to aiter as fused_mixed_sample (Triton, in aiter/ops/triton/fusions/, with configs tuned per arch, a unit test and a benchmark).
It uses the same race and the same noise as mixed_sample_outer_exponential (that's why the agreement check is exact), but it isn't the same kernel: aiter's kernel runs one workgroup per row over the whole vocab, while this one tiles each row across the grid and folds in the greedy override (topk_select + <= + where), so the draw takes 2 launches instead of 5. In the aiter version, greedy rows also skip the race and the noise read.
With the tuned configs, the aiter version is faster than the HIP path at every row count we measured, including 256 rows (−33% on MI355X, −27% on MI300X), where the kernel here is break-even.
#2914 doesn't need to wait for it: the kernel here stays behind the default-off flag. Once a release with fused_mixed_sample is in ATOM's image, a small follow-up switches sampler.py to it and deletes atom/model_ops/bonus_sample.py.
Related: the existing aiter kernel's out-of-range ids at small vocab sizes, from the PR description, were fixed today in ROCm/aiter#7373.
| # The same noise draw as `_temperature_sample`, then the race | ||
| # and the greedy override in two vocab-tiled launches. | ||
| num_tokens, vocab_size = logits.shape | ||
| exponential = torch.empty( |
There was a problem hiding this comment.
same exponential_ draw spelled twice, fused path relies on same rng.. one helper for both?
There was a problem hiding this comment.
Done in 6a9bc8f. get_per_row_exponential(num_tokens, vocab_size, device) in sampler.py is now the only place the per-row draw happens. _temperature_sample and the fused branch both call it, so they can't drift apart and they consume the RNG identically.
pytest.importorskip("triton") skipped the whole file on a CPU CI runner
without Triton, although no test here launches a kernel. The fixture now
fakes atom.model_ops.bonus_sample the way it already fakes aiter, so all of
it runs on any CPU runner, and a new test checks that the flag wires
Sampler to bonus_temperature_sample.
Sampling rows are not bit-exact with the aiter kernel (fp32 near-ties), and on MI355X the gain shrinks with the verify-row count: about -25% up to 128 rows but none at 256 (about Kimi-K3 DSpark at concurrency 64). State that in the flag's comment and in environment_variables.md.
Follow-up to the review (3 commits on top of
|
| Review comment | Answer |
|---|---|
Stub bonus_sample in the wiring fixture, drop importorskip |
a8347e8 |
exponential_ draw written twice |
6a9bc8f: get_per_row_exponential |
| Why default off | c70d5bf: reason documented |
| Five buffers allocated per call | Measured a cached buffer: +2.2 µs per call on MI355X, so torch.empty stays (numbers in the thread) |
| Move into aiter | ROCm/aiter#7416 (fused_mixed_sample, tuned for gfx950 and gfx942). #2914 doesn't depend on it |
@yitingw1's points
- move from private repo to ROCm #1–update the link to the repo #3 (the NaN is real; remove the clamp; add a
-inftest): done ine34da58, see the reply on the NaN thread. - Update README.md and pyproject.toml #4–fix the example code cmd #5 (include the noise-draw cost; add a rows=256 point): in the results comment above. The draw here is break-even at 256 rows on MI355X, which is also why the flag stays off. The tuned aiter version is 33% faster than the HIP path at 256 rows.
- update readme for deepseek #6 (the aiter out-of-range ids at vocab 1000 / 4097): fixed upstream in [HIP] [Bugfix] sample_kernels: partial-vector and prologue-tail defects aiter#7373, thanks.
- support gpt oss #7 (state the default-off reason in the env doc):
c70d5bf.
Validation
bonus_sample.py and its kernel tests are unchanged from e34da58 (GPU results in the comment above). The three commits change sampler.py, the wiring tests and docs.
- CPU: the full CI unit suite (
.github/scripts/run_unit_tests.sh) passes on the new head, andblackandruffare clean. - MI355X, no
LD_PRELOAD: the server starts on the rebased branch (main's fix(loader): resolve OpenMP symbol through torch #2936 loader fix is in). The end-to-end run with the follow-up passes with the flag OFF and ON:
metric off on
wall_s 64.5 68.2
sampling_requests 200 200
sampling_errors 0 0
sampling_wrong_n 0 0
sampling_choices 300 300
sampling_empty 0 0
sampling_correct 300 300
required_requests 64 64
required_errors 0 0
required_broken 0 0
mtp_acceptance 0.648 0.6493
hard_failures 0 0
Kernel checks:
[1] NaN partial keys on a forced row: 0 / 80 (unpatched: >0 expected = bug present; patched: must be 0)
[2] forced rows wrong: 0 / 7
[3] masked rows: 0 unexplained mismatches, 0 near-ties, 0 out-of-range ids, over 1600 rows
[4] all -inf rows out of range: 0 / 4 [0, 0, 0, 0]
PASS
That GPU run also included the cached-buffer commit, which was then dropped. The draw it measured is the same draw.
|
The previous round is fully addressed, as far as I can tell — the Two notes from re-reading it with fresh eyes, neither blocking. Coverage reads thinner than it is. A row-count gate would answer the second of the two default-off reasons. The crossover is now pinned on both arches (~256 rows), so Nit: |
…default Per review: a row-count gate (fused at or below the crossover, aiter above) would remove the speed reason for keeping the flag off, and aiter's fused_mixed_sample (ROCm/aiter#7416), already faster at 256 rows, would remove it without one. Say so in the env var doc and the envs.py comment. The test module docstring now says why a plain pytest run reports most tests as skipped: they run in the interpreter child, which fails if any of them fails or if it collects nothing.
|
Thanks for re-reading it, and for running the interpreter child and the two mutations.
|
|
Thanks — the three points are answered, and I'll withdraw the You were also right that Having read both, I think the scope question comes before anything else here. By line count, 410 of this PR's 592 additions are the part ROCm/aiter#7416 now covers: bonus_temperature_sample(logits, temperatures, exponential, eps, greedy_eps)
fused_mixed_sample(input, exponentials, temperatures, eps=…, greedy_eps=…, out=…)Two things I found while comparing them argue the same way — the in-tree copy isn't just a duplicate, it's the weaker of the two. 1. The
They agree exactly at T = 2.0 / 1.0 / 0.5, where 2. Stage 1 here computes both paths for every row and stores 5 partials; #7416 branches on → Change: land the kernel in #7416 and reduce this PR to the ATOM side — One thing I couldn't settle: #7416's |
…ample Per review: the kernel now lives in aiter as fused_mixed_sample (ROCm/aiter#7416), so ATOM keeps only its side. With ATOM_SPEC_BONUS_FUSED=1 the sampler calls fused_mixed_sample(logits, noise, temperatures, eps=..., greedy_eps=...) on the same per-row exponential draw. Remove the in-tree bonus_sample.py, its kernel tests and tests/triton_interpreter.py. aiter's op subtracts the max before scaling by 1/T, so a sampling row differs from mixed_sample_outer_exponential only on a near tie of a few 1e-7 relative. The in-tree kernel scaled first, and that window grew like 1/T. An aiter without fused_mixed_sample keeps the aiter race and warns once. On a GPU aiter has no tuned config for (it ships gfx942 and gfx950), the op raises FileNotFoundError before launching; the sampler then runs the aiter race on the noise it already drew, so the RNG stream matches the flag-off path, and keeps that path for later steps. The wiring tests cover both fallbacks and the argument order.
|
Thanks — agreed on all of it.
Depends on #7416 landing in ATOM's aiter; until then the flag just warns. Description updated. |
Summary
Speculative decoding draws the bonus token, plus every draft-conditioned verify row of a non-greedy batch (
needs_independent_noise=True), with no top-k/top-p filter. On main that draw is five launches on the critical path of every verify step:mixed_sample_outer_exponential(one program per row walking the whole 163840 vocab) →topk_select(2 launches) →<=→whereWith
ATOM_SPEC_BONUS_FUSED=1the sampler calls aiter'sfused_mixed_sample(ROCm/aiter#7416) instead: the same race and greedy override in two vocab-tiled Triton launches, on the sameexponential_noise. Default off. With the flag off, the kernels and their order are exactly as on main.Depends on ROCm/aiter#7416. Until ATOM's aiter has
fused_mixed_sample, setting the flag only warns once and keeps the aiter path, so this PR is safe to merge before or after it; the flag does nothing useful before.What changes
atom/model_ops/sampler.py__init__reads the flag once and importsaiter.ops.triton.fusions.fused_mixed_sample. An aiter without it: one warning, aiter path.forwardcallsfused_mixed_sample(logits, noise, temperatures, eps=self.eps, greedy_eps=SAMPLER_EPS)only for an unfiltered batch with per-row noise. Top-k/top-p, shared-noise and all-greedy batches keep the aiter path.FileNotFoundErrorbefore launching elsewhere. The sampler then runs the aiter race on the noise it already drew (so the RNG stream is the flag-off one) and keeps the aiter path for later steps, with one warning.get_per_row_exponential; the aiter race takes it through_race.atom/utils/envs.py,docs/environment_variables.md,docs/model_ops_guide.md: the flag, its dependency, and why it is off by default.tests/test_bonus_sample_wiring.py(CPU, no AITER or Triton needed).The kernel, its tests and its tuned configs are in ROCm/aiter#7416, not here.
Measurements
Image
rocm/atom-dev:nightly_202611011508-rocm10.0(ROCm 10.0.0, torch 2.11.0+rocm10.0.0, triton 3.8.0), aiter0dd271600with #7416's 6 files at6074827, this head0a4e7ec. MI355X (gfx950) and MI300X (gfx942).Tests. aiter's
test_fused_mixed_sample.py: 39 passed on each arch. ATOM's wiring and sampler tests: 71 passed.ATOM's
Sampler, flag off vs on, on the same noise (sampler_check.py: 20 seeds × 256 rows, vocab 163840, bf16 logits). Greedy rows that differ: 0 on each arch. Sampling rows that differ, per temperature:Sampler.forward, noise draw included, median µs at T = 0.7:Kernel-level numbers (stage by stage, tuned configs) are in ROCm/aiter#7416.
End to end with the fused path armed: Kimi-K3 + DSpark on MI355X,
recipes/Agentic-Kimi-K3.md§1 at CONC=4 (DCP1) without--spec-decode-acceptance-length,ATOM_SPEC_BONUS_FUSEDtoggled alone. Each run: 200 sampling requests (T 0.7 / 1.0, half withn=2) and 64tool_choice="required"requests (T 0.7).n, broken tool calls)With the flag on, aiter's
fused_mixed_sampleran 17,080 times during the requests (1 to 49 rows,BLOCK_V1024 at 1 row and 2048 above, from the gfx950 config), with no fallback. The acceptance difference is within run-to-run noise (earlier pairs: 63.1 / 64.7%, 64.8 / 65.2%).Why default off
Sampling rows are not bit-exact with the aiter kernel: they can differ when a row's two best race keys are within a few 1e-7 relative (fp32 rounding order), unlike
ATOM_DSPARK_FUSED_MARKOV_SAMPLE. It also needs an aiter that hasfused_mixed_sample.Testing
.github/scripts/run_unit_tests.sh): 11124 passed, 0 failed.pytest tests/test_bonus_sample_wiring.py→ 9 passed: flag off keeps the aiter ops; the flag wires aiter'sfused_mixed_sample; the argument order; filtered, shared-noise and all-greedy batches never reach it; an aiter without the op, and a GPU without a config, fall back (the latter on the same noise, and for later steps).blackandruffare clean on the changed files.Samplercheck and (MI355X) the end-to-end run above.Notes
sampler.py,envs.py). There is no semantic overlap: feat(spec-decode): probability-ratio rejection sampling behind draft_probs #2582 changes the verify rule, this PR changes the unfiltered draw.mixed_sample_outer_exponentialat vocab 1000 and 4097 (vocab not a multiple of 4, or below 4096) are fixed upstream in [HIP] [Bugfix] sample_kernels: partial-vector and prologue-tail defects aiter#7373. Vocab 163840 (production) was never affected.