Repository navigation
[ROCm][Perf] Triton port of the fused GDN MTP decode kernel - #61855
Conversation
The fused GDN MTP decode kernel (fused_gdn_decode_post_conv_mtp) only builds for CUDA, so on ROCm the speculative-decode step of every flat-layout Qwen GDN layer (Qwen3.5, Qwen3.8-27B, Qwen3.8-Flash-Next) falls back to the unfused path. Add a Triton port with the same signature and use it on ROCm; the existing kernel test now covers both implementations. Co-authored-by: Claude <noreply@anthropic.com> Signed-off-by: iseeyuan <myuan@meta.com>
e57aeb1 to
7d56761
Compare
|
|
||
|
|
||
| @triton.jit | ||
| def _fused_gdn_decode_post_conv_mtp_kernel( |
There was a problem hiding this comment.
This eliminates the copy/cat/fill overhead, but drops the BV=32 V-tiling that fused_sigmoid_gating_delta_rule_update_kernel uses.
Since only the final RMSNorm needs the full V reduced, would tiling the state update over V like the existing kernel, with the norm done as a small epilogue over the accumulated sum-of-squares, be faster?
There was a problem hiding this comment.
Good call, measured it: on MI350X at the Qwen3.5-35B-A3B shape (fp32 state), V-tiling with BV=32 plus a small norm epilogue is 33% faster per call at 1 request and 15% at 32, with bit-identical outputs and states. I'll switch the PR to that layout.
simondanielsson
left a comment
There was a problem hiding this comment.
Thanks for the fix, this looks good
Suggestion: Can we please attach
- gpqa and niah accuracy results vs main
- vllm bench serve with 64k/1k ISL/OSL
Suggestion: Can we also check the perf of this on gfx942, or add a 950 gate?
Address review: like fused_sigmoid_gating_delta_rule_update, split each value head's state update over the value dimension, and move the gated RMSNorm, the only step that needs the whole value dimension, into a small epilogue kernel that combines per-tile sums of squares. The tile size follows the grid: 32 rows per program while there are fewer (request, value head) pairs than CUs, 64 up to 4x the CUs, else one tile per value head. Tuned with a tile-size sweep on MI300X and MI350X; it is within 8% of the best tile size at every measured batch size on both. States are bit-identical to the single-kernel version; outputs are too on MI350X and match within the rounding of the norm's sum on MI300X. Per call (Qwen3.5-35B-A3B shape, FP32 state, 4 tokens per request): MI350X 18.7 -> 11.5 us at 1 request and 71.9 -> 64.3 us at 32; MI300X 22.3 -> 14.4 us and 97.0 -> 90.1 us. Co-authored-by: Claude <noreply@anthropic.com> Signed-off-by: iseeyuan <myuan@meta.com>
|
Thanks, good call. Switched to that layout in 66eb7ec: the state update is tiled over V like Per call (Qwen3.5-35B-A3B shape, FP32 state, 4 tokens per request), previous -> now: MI350X 18.7 -> 11.5 us at 1 request and 71.9 -> 64.3 us at 32; MI300X 21.7 -> 13.3 us and 96.8 -> 88.8 us. End to end on MI350X at batch 1, decode went from 2.19 to 2.08 ms/token (unfused fallback: 2.39). The description has the refreshed numbers. |
Review nit: one condition instead of an early return in forward_hip; the routing is unchanged (speculative decode still takes forward_cuda). Co-authored-by: Claude <noreply@anthropic.com> Signed-off-by: iseeyuan <myuan@meta.com>
|
@simondanielsson Thanks for the suggestions; here are the results (methodology and full tables in the updated description).
GPQA, NIAH and 64K/1K ran on MI350X with vLLM nightly 77154ca + aiter main and shared-expert fusion, main vs this PR side by side. |
Review round 3: - Use vllm.utils.platform_utils.num_compute_units for the tile rule instead of a local device-properties cache. - qwen_triton_warmup compiles the ROCm Triton MTP decode kernels when the layer takes the fused spec-decode path: once per V tile size the request count can select (up to max_num_seqs), with the layer's packed QKVZ/BA layouts so the stride specializations match serving. - The wrapper asserts the layouts the kernels index with, as the CUDA op checks them. - test_gdn_fused_mtp.py now runs on ROCm against the Triton entry point (Triton GDN prefill there, as cutedsl is CUDA-only), and the MTP kernel test only applies the compute-capability gate on CUDA. Co-authored-by: Claude <noreply@anthropic.com> Signed-off-by: iseeyuan <myuan@meta.com>
|
✅ @iseeyuan, CI is now available for this PR.
|
simondanielsson
left a comment
There was a problem hiding this comment.
Thanks for the changes! LGTM
|
/ci run |
|
✅ Triggered Buildkite CI #96005 for commit |
Review suggestion: `supports_fused_mtp_decode()` and `fused_mtp_decode_impl()` hold the "Triton port on ROCm, CUDA op elsewhere" logic that the init-time eligibility check, the decode-time eligibility check and the dispatch each repeated; the model-path test patches the impl instead of choosing a target per platform. Co-authored-by: Claude <noreply@anthropic.com> Signed-off-by: iseeyuan <myuan@meta.com>
…iton Signed-off-by: iseeyuan <myuan@meta.com>
…cm-fused-gdn-mtp-triton Signed-off-by: iseeyuan <myuan@meta.com>
|
/ci run |
|
❌ This PR is 61 commits behind upstream |
|
/ci run |
|
✅ Triggered Buildkite CI #96663 for commit |
|
@simondanielsson since all tests pass, could we merge the pr to main to have the improvement landed, and to unblock the deployment by default? |
|
@iseeyuan you can join vllm slack => ping maintainers in |
…ject#61855) Signed-off-by: iseeyuan <myuan@meta.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Simon Danielsson <70206058+simondanielsson@users.noreply.github.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Overview
GDN layers with the flat qkvz layout (Qwen3.5 family) fall back to the unfused speculative-decode path on ROCm because
fused_gdn_decode_post_conv_mtp(#51674) only builds for CUDA. This adds a Triton port with the same signature and uses it on ROCm: a state-update kernel tiled over the value dimension likefused_sigmoid_gating_delta_rule_update(tile size chosen from the grid size) plus a small gated-RMSNorm epilogue. It covers every model built onQwenGatedDeltaNetAttentionwith the flat layout: the Qwen3.5 family, Qwen3.8-27B and Qwen3.8-Flash-Next (table below). Qwen3-Next (interleaved layout) stays ineligible, as on CUDA.Claims
77154ca) with aiter main. The first version of this PR (one program per value head, no V tiling) gave 2.19 ms/token in the same side-by-side run.vllm bench serve, 2K/256).tests/evals/gsm8k/configs/Qwen3.5-35B-A3B-DEP2.yaml.Validation
Kernel test (MI350X and MI300X, ROCm 7.2.3; the CUDA-op cases skip because the op is not built on ROCm):
End to end: Qwen/Qwen3.5-35B-A3B (BF16), 1x MI350X,
--speculative-config '{"method":"mtp","num_speculative_tokens":3}' --attention-backend ROCM_AITER_FA --compilation-config '{"cudagraph_mode":"FULL_AND_PIECEWISE"}' --async-scheduling,VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS=1. Workload: 104 image-to-video prompt-rewriting requests (one image + 9.5K-token system prompt, ~180 output tokens, temperature 0.6 / top_p 0.95 / top_k 3), one at a time, first pass on a fresh server.The three servers ran side by side on three GPUs of the same host. The I2V, profile, GSM8K, GPQA and NIAH runs predate the tile-size rule; at one request it picks the same 32-row tiles, and it leaves states and outputs unchanged in the checks above.
aiter main is used because aiter <= v0.1.25 sends the multi-token verify batch to an unsplit gfx950 attention kernel (fixed by ROCm/aiter#6336, not yet released); without it attention dominates the step and hides this change.
Per decode step GPU time (torch profiler, 39 steps, same servers):
vllm bench serve(base and PR servers side by side; random dataset, 2048 input / 256 output tokens,--ignore-eos, 4x concurrency prompts, min 32). Acceptance differs between runs because the outputs differ. With MTP each streamed chunk is one decode step, so median ITL is the per-step time independent of acceptance; TPOT x acceptance is shown too:Long context,
vllm bench servewith 64K input / 1K output tokens (random dataset,--ignore-eos, 8 / 32 / 64 prompts at concurrency 1 / 8 / 32, both servers with--max-model-len 69632and prompts neither had cached):The 64K prefill dominates these runs, and acceptance on random prompts varied from 2.49 to 2.89, which moves throughput more than this change does. The decode step itself is 3-12% faster and prefill is unchanged; at 32 concurrent requests, time to first token mostly reflects queueing behind other prefills and varied by about 30% between passes. A pass with another prompt seed gave -12% / -7% / -6% median ITL.
GSM8K (
tests/evals/gsm8k/gsm8k_eval.py, 1319 questions, 5-shot, greedy,--max-concurrency 16, so MTP verify batches of up to 16 requests go through the new kernel):GPQA Diamond and needle-in-a-haystack, main vs this PR side by side (both servers with
--max-model-len 69632):GPQA Diamond uses the OpenAI simple-evals prompt and answer pattern: 198 questions x 2 repeats with shuffled choices, thinking on, temperature 0.6 / top_p 0.95 / top_k 20, up to 24K output tokens, 64 concurrent requests. NIAH follows RULER's
niah_single_1: a 7-digit number hidden at 11 depths (0-100%) in the noise haystack, 3 samples per depth, thinking off, greedy.Reproduce (on a ROCm image, from the vLLM repo root):
Details
Models that take the fused path on ROCm with this change (eligibility is unchanged from the CUDA op: flat qkvz layout, K = V = 128, value/key head ratio 1/2/3/4/8, BF16 activations, at most 8 tokens per verify):
qwen3_5.pyqwen3_5.pyqwen3_5.pyqwen3_5.pyqwen3_5.py(Qwen3_5ForConditionalGeneration)vllm/models/qwen4_exp(same GDN layer, flat layout)qwen3_next.py(interleaved layout)Root cause. On ROCm the layer logs
Falling back to the Triton GDN spec decode path: torch.ops._C.fused_gdn_decode_post_conv_mtp is not built. The AITER fused GDN path only covers non-speculative decode, so every GDN layer's verify step runs ~8 small kernels around the recurrent kernel: contiguous copies ofa/b, three copies plus acatto unpack q/k/v, a zero-fill, an output copy, and the gated RMSNorm.Kernels. Two launches. The update kernel runs one program per (request, value head, tile of the value dimension). It reads q/k/v and
a/bin place from the projections and runs the gated delta-rule update over the request's tokens with the same state reads and writes asfused_sigmoid_gating_delta_rule_update: start from the state at the last accepted token, store the state after every token, slot index <= 0 means no state. Each tile writes its rows of the output and their sum of squares per token. The epilogue, one program per (request, value head, token), adds the partial sums and applies the gated RMSNorm, the only step that needs the whole value dimension, then writes theout_projinput. Recurrent state layout is[slot, HV, V, K], BF16 or FP32.Tile size. 32 rows per program while there are fewer (request, value head) pairs than CUs, 64 up to 4x the CUs, else one tile per value head. A sweep of 32/64/128 rows on MI300X and MI350X (both head shapes, both state dtypes, 1-32 requests) put this rule within 8% of the best tile size everywhere; a fixed 32 was up to 36% slower at 16-32 requests on MI300X.
Wiring. Same keyword signature as the CUDA op; the layer picks the Triton port on ROCm.
forward_hiproutes speculative-decode configs to the packed fused-norm path (forward_cuda); non-speculative ROCm decode keeps the AITER path. Opt out with the existingVLLM_GDN_DECODE_KERNEL=triton.Per call. Qwen3.5-35B-A3B shape (16/32 K/V heads), 4 tokens per request, CUDA-graph replay, current ROCm nightly image; "first version" is this PR before the review (one program per value head):
With FP32 state, the Qwen3.5 and Qwen3.8 default (
mamba_ssm_dtype: float32), this PR is faster than the first version at every batch size on both GPUs. With BF16 state at 32 requests on MI350X it is 8-9% slower than the first version (one tile per head plus the epilogue launch), still well ahead of the unfused path. The 16/48 head shape behaves the same way.Coverage. End to end this was run on Qwen3.5-35B-A3B only. The other models in the table are covered by the kernel test, which runs every eligible head ratio (1/2/3/4/8) at TP 1-16, but were not run end to end.
CI. The parametrized test runs the Triton port on any GPU (NVIDIA too), so existing CUDA CI covers its numerics; AMD CI covers the ROCm path where it runs
tests/kernels/fla.Related open PRs (not duplicates).
VLLM_ROCM_QWEN_MTP_GDN), shape-specialized ROCm kernel that also fuses the conv. It accepts only 16/48 K/V heads (Qwen3.8-27B), TP=1 and verify widths up to 4, and raises for any other shape, including Qwen3.5-35B-A3B (16/32). This PR makes the default ROCm path fused under the same contract and eligibility as the upstream CUDA op (head ratios 1/2/3/4/8, any TP, up to 8 tokens). The two can coexist, with [Perf][ROCm][Kernel] Fuse Qwen3.8 MTP verify conv, GDN, and RMSNorm #61675 taking over when it is enabled.mamba_cache_mode="none"). On ROCm this PR replaces that path whenever the fused path is eligible; the deferred-store idea could be applied to this kernel in a follow-up.Provenance. The first version of the kernel was written and validated in Meta's internal vLLM fork; the V tiling was added in this PR after review.
AI assistance. Drafted with Claude Code (port, V tiling and tile-size rule after review, test parametrization, benchmarks); reviewed and run by the submitter.
Pull Request Checklist
I used vLLM's
/pr-checklistskill. (Mandatory for agents, optional for humans).AI assistance was used during the creation of this PR.
Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.
Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.
Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.
Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.