Repository navigation
[AMD] Add Work-Centric (Lean) Attention: a persistent-CTA decode kernel for long-context serving - #33576
Conversation
This adds a persistent-CTA, work-stealing Triton decode-attention kernel for the ROCm/AMD Triton backend. It targets the long-context, low-batch decode regime where the standard flash-decode kernel leaves GPU compute units idle. The kernel launches a fixed persistent grid sized to the device CU count and derives its own tile schedule on-device from kv_indptr, keeping every CU busy via work-stealing. Decode-step latency stays nearly flat as context grows. Lean is opt-in and gated: numerically identical to the standard kernel (cosine similarity 1.0000), and a cheap host-side seqlen gate keeps it off in the short-context regime where the standard kernel is faster. Benchmarked on Qwen2.5-7B (GQA) / MI355X: kernel-level Lean vs SplitK up to 5.6x at 128K; end-to-end decode ITL 1.14x@32K -> 2.21x@128K. Changes: - python/sglang/kernels/ops/attention/decode_attention.py: Add Lean kernel and gate helpers; decode_attention_fwd gains enable_lean parameter. - python/sglang/srt/layers/attention/triton_backend.py: Read enable_lean_attention from server args; resolve it in forward_decode (explicit override, else host-side auto-gate, else kill-switch); pass through to decode_attention_fwd. - python/sglang/srt/server_args.py: Add --enable-lean-attention flag. - python/sglang/srt/environ.py: Add SGLANG_DISABLE_LEAN_ATTENTION kill-switch. - benchmark/lean_kernel_qwen_gqa.py: Kernel-level comparison benchmark. Signed-off-by: Valerie Chen <valechen@amd.com>
- Add missing _LEAN_BLOCK_M and _NUM_CU global constants before _lean_num_cus - Update benchmark import from triton_ops to kernels.ops (matches file move) Signed-off-by: Valerie Chen <valechen@amd.com>
Signed-off-by: Valerie Chen <valechen@amd.com>
Follow-up to the initial Lean (Work-Centric) Attention decode kernel: - decode_attention.py: kernel/launch/gate refinements (persistent-grid sizing, head-dim guard, gate docstring + batch-aware thresholds). - triton_backend.py: resolve enable_lean per forward (explicit override → auto-gate → SGLANG_DISABLE_LEAN_ATTENTION kill-switch) and allocate the Lean scratch buffers (Mp/Lp/Op/locks) for both eager and CUDA-graph paths. - benchmark/lean_kernel_qwen_gqa.py: allocate Lean scratch buffers, clearer attn_logits/attn_lse naming. Validated on MI355X / triton 3.7.0: bit-identical output (cos=1.0000) across 60 kernel configs; up to 4.36x kernel and 1.33x E2E throughput for GQA at long context / low batch. Signed-off-by: Valerie Chen <valechen@amd.com>
Per-commit correctness gate on the MI35x AMD suite
(stage-b-test-1-gpu-small-amd-mi35x), where Lean actually runs:
- Parity: Lean output is numerically identical to the standard SplitK
grouped kernel (cosine sim > 0.999, finite) across the two GQA head
shapes (Qwen 28Q/4KV, Llama 32Q/8KV) x batch {1,8} x context {8K,32K}.
- Gate logic: lean_decode_seqlen_gate enables Lean for long-context /
low-batch, stays off for short context, stays off for MLA, and its
threshold relaxes (not tightens) with batch.
Correctness is triton-version-independent (only perf varies with the
build), so this is a robust per-commit gate. Validated on MI355X:
5 tests / 8 subtests pass.
Signed-off-by: Valerie Chen <valechen@amd.com>
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
The Lean decode path crashed on fp8 KV cache (--kv-cache-dtype fp8_e4m3): tl.dot(q, k) fed a bf16 q and an fp8 k into the MMA -> "Unsupported rhs dtype fp8e4nv". The standard SplitK kernel already handles fp8; Lean now adopts the same dtype/scale handling: - Dispatch folds k_scale into sm_scale and passes v_scale to the Lean path (previously dropped), matching the standard grouped call. - Kernel casts q -> K_Buffer.dtype for the main dot (fp8xfp8 MMA), dequantizes the rope-split kpe -> q dtype, and applies v_scale at the single output-write. - For bf16/fp16 KV, k_scale=v_scale=1 and the casts are no-ops, so existing parity is preserved. Adds an fp8 KV parity test (Lean vs standard, cos>0.99) alongside the bf16 case. Validated on MI355X: 6 tests / 16 subtests pass. Phase 2 of the Lean attention work; unblocks fp8-KV GQA and removes the DeepSeek-R1-MXFP4 fp8 crash. Does not change the MLA auto-gate. Signed-off-by: Valerie Chen <valechen@amd.com>
The seq-len auto-gate (lean_decode_seqlen_gate) is a no-op under CUDA graphs: decode-graph capture fills seq_lens with the fill value (1), so avg_len is 1 and the gate always bakes the standard kernel into every captured graph. With --enable-lean-attention unset (auto) and CUDA graphs on (the default), Lean therefore never activated. Lean's fixed 512-CTA persistent grid still adapts to per-step raggedness on-device at replay (work-stealing from kv_indptr), so the fix is to key the capture-time bake decision on capture-time-known signals (batch, head-tile parallelism, is_mla) rather than seq_lens. - Add lean_capture_policy(num_q_heads, kv_group_num, batch, is_mla), used during CUDA-graph capture (detected via get_is_capture_mode). Calibrated on MI355X (triton 3.7.0, CUDA graphs on): bake at GQA batch>=16 and MLA batch>=8; never for heavy TP shards (tiles<4). - Drop the blanket is_mla->off in lean_decode_seqlen_gate; MLA is now gated on batch (batch 1 is a hard loss as the many-query-head decode already saturates the CUs, while batch>=8 is parity on uniform batches and a win on ragged ones). Both gate functions share the new _lean_head_tiles helper. - forward_decode resolves auto mode via the capture policy during capture and the seq-len gate in eager decode. - Tests: add TestLeanCapturePolicy and update the MLA gate test to assert the batch-threshold behavior. Signed-off-by: Valerie Chen <valechen@amd.com>
The Lean decode path assumed a contiguous [N, head, dim] KV buffer and addressed each slot as kv_loc * stride, so it could not read a paged [num_pages, page_size, head, dim] buffer and was skipped whenever page_size > 1. Add the same page-aware address math the standard grouped kernel already uses: - _lean_attention_decode_kernel takes page/token strides and a PAGE_SIZE constexpr; the K, rope-split KPE, and V loads split kv_loc into (kv_loc // PAGE_SIZE, kv_loc % PAGE_SIZE) when PAGE_SIZE > 1, and keep the contiguous fast path at PAGE_SIZE == 1 (Triton specializes the dead branch away). - _decode_lean_attention_fwd derives the strides via _extract_kv_strides (which handles both 3-D and 4-D buffers) and reads head_num from dim -2 so the 4-D paged layout is addressed correctly. - decode_attention_fwd drops the page_size == 1 guard on the Lean dispatch and threads page_size through. Adds a paged-KV parity test (4-D buffer, scattered slot indices, cos > 0.999) that exercises the page_id/tok_in_p math. All 10 tests / 20 subtests pass. Signed-off-by: Valerie Chen <valechen@amd.com>
The Lean persistent grid was hard-coded to 2x the device CU count. Kernel and end-to-end A/B sweeps on MI300X (gfx942) and MI355X (gfx950) found one CTA per CU (1x) beats 2x across uniform and ragged decode: 2x oversubscribed the CUs and regressed mid/high-batch decode, while 1x maximizes KV work-tiles per CTA and minimizes the cross-CTA combine/atomic reduction. - environ.py: add SGLANG_FORCE_LEAN_GRID_CU_MULT (EnvFloat, default 1.0); grid size = round(device_CU_count * mult), a rebuild-free A/B knob. - decode_attention.py: size the persistent grid from the multiplier. - benchmark: replace lean_kernel_qwen_gqa.py with lean_kernel_sweep.py, a full GQA batch x context kernel-vs-SplitK sweep. Signed-off-by: Valerie Chen <valechen@amd.com>
The eager seq-len gate computed avg_len = seq_lens_sum / batch unconditionally. Under EAGLE3 speculative decoding the draft runner issues decode forwards without a CPU length mirror, so seq_lens_sum is None and the division raised TypeError, crashing the scheduler. Fall back to the standard kernel when seq_lens_sum is None (no length signal to gate on), matching the have_cpu_mirror handling elsewhere in the triton backend. Add a regression test for the None case. Signed-off-by: Valerie Chen <valechen@amd.com>
Replace `tl.static_range(max_output_tile_cnt + 1)` with a `while iter < cta_end_tile_gid` loop. Triton's `tl.static_range` requires the loop bound to be fully resolvable at compile time, but `max_output_tile_cnt` is computed from runtime inputs (head_num, batch, total_programs) and then passed as a constexpr parameter. Even though each kernel specialization sees a constant value, Triton's MLIR Coalesce pass cannot statically analyze the bound during compilation, causing a crash in `mlir::triton::gpu::CoalescePass`. The work-stealing loop terminates when `iter >= cta_end_tile_gid` (the actual work boundary computed on-device), so the artificial iteration limit `max_output_tile_cnt + 1` was only ever a safety upper bound. The `while` formulation preserves the intended semantics while removing the Triton compile-time constraint. This fixes CI timeouts in stage-b-test-1-gpu-large-amd where the scheduler watchdog would fire after 300s waiting for the compiler to finish. Signed-off-by: Valerie Chen <valechen@amd.com>
|
@amd-bot ci-status |
CI Status for PR #33576Merge verdict: Caution The changed code is not exercised by any PR-CI test that actually ran. Caution AMD PR CI is incomplete. Changed files: Executed CI failure attribution: AMD: 5 failures (0 clearly PR-caused, 2 possibly) · NPU: 1 failure (0 related) · Others: gate/finish cascades only (no executed test failures) AMD Executed Failures
Other Executed Failures
Details / what to do before merge
Generated by amd-bot using Claude Code CLI |
Motivation
The standard flash-decode ("SplitK") attention kernel parallelizes decode by splitting each sequence's KV cache into chunks and assigning chunks to GPU compute units (CUs). While effective at short context, it leaves performance on the table in two common serving regimes:
batch × head-tiles × num_splitswith a fixed split count, so a single (or few) long sequence cannot generate enough KV-split work to fill all CUs: many CUs sit idle while the few launched programs each grind through a long serial KV loop, and decode-step latency grows roughly linearly with context, making the workload occupancy-starved, not bandwidth-bound. As context windows reach 64K–128K+, attention dominates thedecode step here.
Work-Centric Attention (WCA / "Lean") launches a fixed persistent grid sized to the device CU count and derives its KV-tile schedule on-device from
kv_indptr, cooperatively splitting long sequences across CTAs (work-stealing) so no CU idles. Decode-attention latency stays nearly flat as context grows (1.06–1.30× E2E throughput even on uniform batches) and absorbs raggedness for free (up to 1.52× throughput, 3.62× lower ITL). Lean is numerically identical to the standard kerneland ships behind a conservative auto-gate that enables it only where it wins — including by default under CUDA graphs — so existing workloads are unaffected.
Validated on AMD MI355X (gfx950) across GQA (Qwen2.5-7B, Llama-3.1-8B) and MLA (DeepSeek-R1-MXFP4, TP=4, fp8 KV). Also re-validated on AMD MI300X (gfx942) at the same 1×CU grid across GQA (Qwen2.5-7B, Llama-3.1-8B) and MLA (proxied with DeepSeek-V2-Lite, fp8 KV, since fp4/MXFP4 GEMM is gfx950-only).
Modifications
A minimal, mostly-additive change (6 files, ~+1540 lines) on the ROCm/Triton backend. Almost all of it is new code; the only edits to existing paths are the
forward_decodedispatch hook and thedecode_attention_fwdsignature. The standard decode path is untouched when Lean is inactive.python/sglang/kernels/ops/attention/decode_attention.py(+822)_lean_attention_decode_kernel) and its launch/dispatch helpers.decode_attention_fwdgains anenable_leanparameter and early dispatch to Lean. Supports bf16/fp16 and fp8 KV (foldk_scaleintosm_scale, castq → K.dtypefor the fp8×fp8 MMA, dequantize rope-splitkpe, applyv_scale), and both contiguous and paged KV (kv_loc // page_size/% page_sizeaddressing). Auto-gate:lean_decode_seqlen_gate(eager) andlean_capture_policy(CUDA-graph capture).test/registered/kernels/test_lean_attention.py(+440)stage-b-test-1-gpu-small-amd-mi35x): parity vs standard SplitK for bf16 (cos > 0.999), fp8 (cos > 0.99), and paged KV (cos > 0.999), plus gate/policy logic.python/sglang/srt/layers/attention/triton_backend.py(+119)enable_lean_attentioninforward_decode(explicit override → auto-gate → kill-switch); allocate the Lean scratch buffers (Mp/Lp/Op/locks) for the eager and CUDA-graph paths.benchmark/lean_kernel_sweep.py(+150)python/sglang/srt/server_args.py(+4)--enable-lean-attentionflag (Optional[bool], defaultNone= auto-gate).python/sglang/srt/environ.py(+11)SGLANG_DISABLE_LEAN_ATTENTIONkill-switch (force the standard kernel) andSGLANG_FORCE_LEAN_GRID_CU_MULT(persistent-grid size multiplier, default1.0= one CTA per CU; a knob for grid A/B without a rebuild).Design:
total_programs = device CU countCTAs — one CTA per CU (rounded to whole XCDs under the ROCm XCD remap). One CTA per CU maximizes the KV work-tiles handled per CTA, so each CTA's cross-CTA combine/atomic reduction amortizes more useful work; oversubscribing (e.g. 2 CTAs/CU) halves tiles-per-CTA and inflates that reduction, regressing high-batch decode (measured on MI300X/gfx942 and confirmed on MI355X/gfx950). The grid is independent of batch×heads×splits, so it is CUDA-graph capturable and a single captured grid replays across batch sizes.kv_indptr; long sequences are cooperatively split across CTAs. The schedule is derived fromkv_indptrat replay, so a captured graph adapts to the real per-step sequence lengths and raggedness — no host sync, no per-step relaunch.None) resolves to a concrete decision before launch; explicitTrue/Falseoverrides it. Because Lean is a win only in specific regimes, the gate operates in two modes:seq_lensare known:lean_decode_seqlen_gatekeys a length threshold on query-head parallelism (tiles = ceil(num_q_heads / min(16, kv_group)), i.e. how well the standard kernel already fills the CUs), relaxed as batch grows.lean_capture_policymakes the bake decision from capture-time-known signals (batch,tiles,is_mla) — bake GQA/MHA atbatch ≥ 16, MLA atbatch ≥ 8, and never for heavy TP shards (tiles < 4). Lean's fixed grid then adapts to the actual lengths at replay.qk_head_dim != v_head_dim); bf16/fp16/fp8 KV; contiguous and paged KV. Falls back to the standard kernel for unimplemented features (logit capping, attention sinks, xAI temperature, score modification), so results are never silently wrong.Usage:
None): auto-gate — Lean turns on automatically where it wins (including by default under CUDA graphs). No flag needed.--enable-lean-attention(with--attention-backend triton).--enable-lean-attention=False, orSGLANG_DISABLE_LEAN_ATTENTION=1.Accuracy Tests
Test configuration:
Command (the registered CI gating test):
Results. Lean matches the standard kernel across every tested configuration:
Every measured configuration is bit-identical in bf16 (the one unmeasured cell, Llama batch-32 × 128K, exceeds the microbenchmark's single contiguous KV tensor — it is a harness limit, not a kernel result). The committed CI gating test (
test_lean_attention.py) reproduces the bf16, fp8, and paged parity checks on the MI35x per-commit suite. Because output is numerically equivalent, downstream generation quality is unchanged.Speed Tests and Profiling
Unless noted otherwise, the tables below are on 1× MI355X (gfx950); the dedicated MI300X (gfx942) subsection re-runs the same suites on that arch. Both: ROCm 7.2.0, PyTorch2.9.1+rocm7.2.0, Triton 3.7.0 (AMD wheel; the production ROCm wheel — the stock rocm7.2 image's Triton 3.5.1 inverts the crossover and must not be used for Lean benchmarking). The baseline is the Triton backend's own standard decode kernel, not AITER: both arms use
--attention-backend tritonand dispatch through the same
decode_attention_fwd, differing only in theenable_leanbranch, so this is a clean Lean-vs-standard A/B.Kernel-level (isolated decode attention, speedup = standard ÷ lean)
Direct kernel call, fp16, random Q/K/V,
page_size=1; batch ∈ {1,2,4,8,16,32} × context ∈ {8K,16K,32K,64K,128K}. Command:SGLANG_FORCE_LEAN_GRID_CU_MULT=1.0 python3 benchmark/lean_kernel_sweep.pyQwen2.5-7B (28Q/4KV):
Llama-3.1-8B (32Q/8KV):
Bold = Lean is faster at the kernel level (speedup ≥ 1×). Peak 5.33× (Qwen) / 5.92× (Llama) at batch=1 / 128K; speedup grows superlinearly with context (SplitK's idle-CU penalty scales with length, Lean's stays flat) and the crossover falls with batch. The production auto-gate turns Lean on around the end-to-end break-even (later than the isolated-kernel crossover), so it enables the long-context / larger-batch cells and keeps the standard kernel for short-context low-batch. (The batch-32 × 128K Llama cell exceeds the microbenchmark's single contiguous KV tensor and is not
measured; real serving uses a paged KV pool and is unaffected — see the E2E tables.)
End-to-end serving — GQA
Full-model serving, closed-loop / fixed-batch (
num-prompts == max-concurrency == reqs,--request-rate inf), CUDA graphs on (--cuda-graph-bs 1 2 4 8 16 32),--seed 42(both arms see the identical batch).--random-range-ratio rsamples per-request prompt and output lengths in[r·ctx, ctx], soratio < 1= ragged. Rows include batch 1 to show the low-batch long-context win and the auto-gate's enabled region (GQAbatch ≥ 16); below b16 the standard kernel runs by default.Columns: reqs = concurrent requests; KV mean / KV range = mean and min–max prompt (KV) length across the batch; dec/req = mean decode (output) tokens per request; dec iters = total decode tokens measured; tps× = Lean ÷ baseline output throughput; ITL× / p99× = baseline ÷ Lean median / p99 inter-token latency (>1 = Lean faster).
Lean wins across every measured cell — all 26 have
tps× ≥ 1.06, none regress. The win grows with raggedness (lower ratio) and context: peak 1.52× throughput (Llama b32/16K r0.1) and 3.62× / 2.91× lower median / p99 ITL (Llama b32/65K r0.1). Uniform batches are also faster at every size (1.06–1.30× throughput, 1.28–1.93× lower ITL) — e.g. b1 uniform 65K: Qwen 1.27× tps / 1.63× ITL, Llama 1.24× / 1.59×; b32 uniform 65K: Qwen 1.07× / 1.66×, Llama 1.16× / 1.93×.Command (per arm; identical
--seed 42):End-to-end serving — MLA (DeepSeek-R1-MXFP4, TP=4, fp8 KV)
Same closed-loop methodology,
--tp-size 4 --kv-cache-dtype fp8_e4m3. Rows include batch 1 to show the gate boundary and where Lean helps vs regresses. The auto-gate enables MLA atbatch ≥ 8(below it, the standard kernel runs by default).Lean wins at batch ≥ 8 (the gate threshold): 1.05–1.13× tps, 1.33–2.41× lower ITL, growing with raggedness (peak b32 ragged: 1.13× tps, 2.41× ITL). At batch 1 it regresses (0.83× tps, 0.60× ITL), confirming the gate is correct — MLA's splitK is faster at low batch, so the auto-gate keeps it off there.
Command:
Default path (auto-gate, no flag)
The tables above are Lean-vs-baseline in the region the auto-gate enables. Confirming the default path (no
--enable-lean-attention) reproduces them under CUDA graphs — 3 arms, Qwen GQA, 64K, raggedratio 0.1,--seed 42:auto(no flag) tracksforcedwhere the capture policy bakes Lean (b≥16) and stays at the standard kernel below the threshold:So the throughput/latency gains are delivered by default, with no configuration change. (At b8 the
forcedarm is already faster than baseline — 88.83 vs 79.55, ITL 1.98× — so the b≥16 GQA capture threshold is deliberately conservative and leaves some b8 headroom on the table; tightening it is future work.)MI300X (gfx942, 304 CUs)
All tests: 1× AMD Instinct MI300X (304 CUs, gfx942, 192 GB HBM3), ROCm 7.2.0, PyTorch2.9.1+rocm7.2.0, Triton 3.7.0 (AMD wheel), CUDA graphs on,
--seed 42, identical harnesses, commands, and auto-gate as above. Parity is arch-independent and passes unchanged (bf16 cos 1.0000, fp8 >0.99, paged >0.999). Results are at the default 1×CU grid(
SGLANG_FORCE_LEAN_GRID_CU_MULTdefault1.0, one CTA per CU).Kernel-level (speedup = standard ÷ lean; fp16,
page_size=1)Qwen2.5-7B (28Q/4KV):
Llama-3.1-8B (32Q/8KV):
Peak 3.72× (Qwen) / 3.34× (Llama) at batch=1 / 128K; high batch now wins across context. Short-context low-batch stays below 1× and is kept on the standard kernel by the gate.
End-to-end serving — GQA (auto-gate enabled region,
batch ≥ 16)Same closed-loop methodology and columns as the MI355X table above. KV mean / KV range / dec-per-req / dec iters are identical to the MI355X rows by construction (same
--seed 42, model,random-input-len,random-range-ratio); onlytps× / ITL× / p99×are re-measured on MI300X.At
1×CUthe enabled region wins across the board on MI300X — including uniform batches (neutral-to-positive throughput, better tail latency). Peak 1.14× throughput and 2.14× / 1.83× lower median / p99 ITL (Llama b32×64K ragged). The MI355X auto-gate thresholds transfer unchanged.End-to-end serving — MLA (proxy: DeepSeek-V2-Lite, fp8 KV)
amd/DeepSeek-R1-MXFP4cannot run on MI300X: fp4/MXFP4 matmul is a CDNA4 (gfx950) feature — gfx942 has no fp4 GEMM (0AFP4WFP4configs;aitermarks the OCP fp4 dtypes "gfx950+"). MLA is therefore proxied on MI300X with DeepSeek-V2-Lite (16B MoE, MLA, fp8 KV, TP=1), same ragged sweep.b1 loses (0.31–0.42× — Lean's fixed grid is pure overhead at V2-Lite's 16-head batch-1 decode); b8 is borderline (~parity to 1.04×); b16/b32 win (up to 1.14× tps, 1.75× ITL). The clean-win region on this proxy is b ≥ 16 (b8 borderline). Note this differs from R1 on MI355X, where the 128-head decode makes batch 1 a win at 1×CU — so the exact
is_mlabatch threshold is a function of head count and should be set from the R1 calibration, not this proxy. (b1 rows omitted from the table.)Conclusion
Lean Attention is a flat-latency, work-stealing decode kernel for long-context and ragged serving on ROCm — up to ~5.9× at the kernel level, 1.52× end-to-end throughput and up to 3.6× lower median / 2.9× lower p99 ITL on GQA, and 1.13× tps / 2.41× lower ITL on MLA at batch ≥ 8 — at zero accuracy cost (bit-identical bf16; fp8/paged parity within quantization noise). It supports MHA/GQA/MLA, bf16/fp16/fp8 KV, and contiguous or paged KV, and runs behind a conservative, CUDA-graph-aware auto-gate that turns it on only where it wins (delivering the gains on the default path) and leaves the standard kernel in place everywhere else — so there is no regression risk.
Reproduce: kernel sweep —
SGLANG_FORCE_LEAN_GRID_CU_MULT=1.0 python3 benchmark/lean_kernel_sweep.py;CI parity (bf16/fp8/paged) —
pytest test/registered/kernels/test_lean_attention.py; end-to-end — thesglang.launch_server+sglang.bench_servingcommands in each subsection above.CI States
Latest PR Test (Base): ✅ Run #32810803823
Latest PR Test (Extra): ❌ Run #32810803658
Latest PR Test (AMD ROCm 7.2): ❌ Run #32810803793