Skip to content

[Triton/Gluon] [gfx950] unified_attention: split-KV 3D configs for head_size <= 128 keyed on max_seqlen_q - #6983

Merged
Boss2002n merged 3 commits into
ROCm:mainfrom
lishuoshuo-amd:perf/gfx950-ua-3d-decode-q2
Nov 3, 2026
Merged

Boss2002n merged 3 commits into
ROCm:mainfrom
lishuoshuo-amd:perf/gfx950-ua-3d-decode-q2

Conversation

@lishuoshuo-amd

@lishuoshuo-amd lishuoshuo-amd commented Oct 27, 2026 •

Copy link
Copy Markdown
Contributor

Summary

On gfx950, a triton unified-attention call with head_size <= 128 that takes the split-KV (3D) path resolves to attn_3d.any (BLOCK_M 16) and kv_split.D_LEQ_128 (SEGMENTS_PER_CU 4). For multi-token queries at long context, as in speculative-decode verify, the grid is too small: a 16q/1kv head, 23 sequences, q=4, context p50 114K gets 16 KV segments and takes 0.774 ms per call; with this change it takes 0.419 ms.

This PR adds a Q axis to the attn_3d and kv_split schemas and three entries:

section key values applies to
attn_3d D_LEQ_128.Q_GEQ_4 BLOCK_M 32, rest as any head_size <= 128, q >= 4
kv_split D_LEQ_128.Q_GEQ_2 SEGMENTS_PER_CU 16, SEGMENT_TILE_MIN/MAX 256, rest as D_LEQ_128 64 < head_size <= 128, q >= 2
kv_split D_LEQ_64 same values as D_LEQ_128 on main keeps head_size <= 64 on main's split

max_seqlen_q == 1 resolves to the same configs as on main, so decode does not change. In SGLang with speculative decoding, total throughput goes up 11-22% on Llama-3.1-8B and Qwen3-14B-FP8.

It also masks padded query rows when the 2D and 3D triton kernels load qq_bias (see "qq_bias padded query rows" below). That out-of-bounds read already exists on main for query lengths that are not a multiple of BLOCK_Q. BLOCK_M 32 makes it reachable at q = 4 as well.

Change

Lookup scan with unified_attention_utils._lookup over head_size {32, 64, 96, 128, 192, 256, 512} x Q {1, 2, 3, 4, 8, 16, 64, 300, 4096} x q/kv dtype {bf16/fp8, bf16/bf16, fp8/fp8} x sliding window {off, on}: the resolved config changes only for attn_3d at head_size <= 128 with Q >= 4, and for kv_split at 64 < head_size <= 128 with Q >= 2. Q = 1, head_size > 128, attn_2d and reduce resolve to the same configs as on main. head_size 96 falls in the changed range; I did not measure a head_size 96 model.

How the values were picked: a sweep of 57 variants (BLOCK_M 16/32 x SEGMENTS_PER_CU 4/8/16/32 x SEGMENT_TILE 64/128/256/512 x MAX_SEGMENTS 128/256) over 143 q >= 2 shapes, two GPU replicas (median disagreement 0.4%), choosing per keyable cell (head_size, q = 2 vs q >= 4). Findings that shaped the entries:

  • The gain comes from BLOCK_M 32 at q >= 4. At q = 2 it halves the query blocks without filling them, so q = 2 keeps 16.
  • 32 segments per CU over-split short KV ranges (worst 0.63x); 16 with a 256-token floor keeps most of the gain.
  • MAX_SEGMENTS 256 made batch-1 shapes slower (down to 0.58x).
  • head_size 64 (gpt-oss) does best with BLOCK_M 32 and main's split.
  • Applying BLOCK_M 32 and 32 segments per CU to every q >= 2 call instead gives a 1.17x geomean with worst case 0.63x; the entries here give 1.23x with worst case 0.81x.

When these entries are read on gfx950:

  1. Triton, not gluon. unified_attention takes the gluon kernel when _gfx950_gluon_supported holds: q and kv dtypes match (bf16/bf16 or fp8/fp8), head_size is a power of two <= 256, and there is no softcap, alibi or qq_bias. Otherwise it reads this triton table, for example with bf16 q and an fp8 KV cache, a softcap, or head_size 96/192.
  2. 3D, not 2D. use_2d_kernel takes the 2D kernel when there is a sliding window, when max_seqlen_k <= 512, or when the 2D grid (num_2d_prgms, about (tokens / BLOCK_Q + seqs) x kv heads) exceeds 4 x CU (1024 on MI355X). So the 3D entries serve long-context calls with a small or medium batch; for Llama-8B (8 kv heads) at q = 4 that is up to about 64 sequences, and gpt-oss at 64 sequences already takes 2D.
  3. Q >= 2. q = 1 resolves to main's entries.

In SGLang (aiter backend, fp8 KV) this is the speculative-decode verify/extend step; q = 1 decode quantizes q to fp8 and goes to gluon.

I measured only gfx950.

qq_bias padded query rows

Both triton kernels compute query_pos = q_block_local_idx * BLOCK_Q + offs_m // num_queries_per_kv and load qq_bias_ptr + query_pos[:, None] * qq_bias_stride_0 + key_rel_pos[None, :] with mask=is_query_key[None, :]. Rows past the sequence's query length are not masked. When the query length is not a multiple of BLOCK_Q, the last query block therefore reads qq_bias rows that do not exist. On main this happens with BLOCK_M 16 and 4 query heads per KV head at q = 3, 5, 6 and 7. With attn_3d BLOCK_M 32, BLOCK_Q becomes 8 and q = 4 reads past the tensor too; this is the case Copilot flagged. The padded rows are already -inf and are not stored, so valid outputs do not change, but the read itself is out of bounds.

The fix adds query_mask_0[:, None] to the three qq_bias loads: two in the 2D kernel and one in the 3D kernel. USE_QQ_BIAS is a tl.constexpr, so calls without qq_bias compile to the same code, and none of the measurements below change.

To show the read, a probe places qq_bias at the end of a mapped HIP VMM page whose next page is reserved but unmapped. Any read past qq_bias then faults. The probe uses 32q/8kv, D 128, bf16, page 16, 1 sequence, and compares the output against a torch reference with the bias:

tree 3D, q=3, kv 4096 3D, q=4, kv 4096 2D, q=3, kv 256
main 036a6ce memory fault ok memory fault
main + these configs, no mask memory fault memory fault memory fault
this PR ok, max abs err 4.1e-4 ok, 3.5e-4 ok, 3.2e-3

The new unit test checks qq_bias results with padded rows on both paths. It cannot detect the read itself, because the read does not change the output.

Results (MI355X, gfx950, ROCm 10.0.0, Triton 3.8.0)

Micro, aiter main be0237f source, CUDA-graph timing over distinct KV "layers", speedup = main / this PR, geo-mean of two GPU replicas:

shape family q speedup geomean [min, max]
Llama-8B verify, 32q/8kv D128, fp8 KV, page 1, ctx 4K-64K, batch 1-64 4 1.16x [0.86, 1.41]
16q/1kv D128 (MiniMax-M3 TP4), page 1, ctx 16K-150K 4 1.39x [0.81, 2.09]
16q/1kv D128, page 16, ctx 16K-150K 2 / 4 / 8 1.05x / 1.35x / 1.48x
16q/1kv D128, 23 seqs, lognormal ctx p50 114K 2 / 4 / 8 1.71x / 1.85x / 1.96x
Llama-8B, page 16, batch 1-2, ctx 16K-64K 16 / 64 / 300 1.19x / 1.51x / 1.28x
16q/1kv D128, page 16, batch 1-2, ctx 16K-64K 16 / 64 / 300 1.01x / 1.16x / 1.58x
gpt-oss full layers, 64q/8kv D64, sinks 4 1.08x [0.93, 1.44]
all 143 q >= 2 shapes 2-300 1.23x [0.81, 2.09], 16 shapes below 0.97x
Llama-8B / 16q/1kv / gpt-oss 1 1.00x

13 of the 16 slower shapes run 1 or 2 sequences. The worst are 16q/1kv q=4 at batch 1 (0.81-0.86x at 64K-150K): BLOCK_M 32 halves the query blocks and the split is already at MAX_SEGMENTS, so the grid shrinks. The others are Llama q=4 at 4K with 2-4 sequences (0.86-0.87x) and 16q/1kv q=2 with 16 sequences (0.89-0.91x). These calls take 10-96 us. The config lookup has no batch-size axis, so a table entry cannot separate batch 1 from batch 16; a batch-aware key would be a separate change to the lookup.

All 16 shapes below 0.97x. Times are per call, averaged over the two GPU replicas, and main / PR is the geomean of the two per-GPU ratios. Page size 1 is the SGLang layout and 16 the vLLM layout. The last column is the 3D kernel's NUM_SEGMENTS and BLOCK_M on each side:

model q / kv heads head size q / KV dtype sinks page size batch kv_len q_len main (us) this PR (us) main / PR segments, BLOCK_M: main -> PR
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 1 1 150000 4 39.4 48.4 0.81x 128, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 1 150000 4 40.2 48.5 0.83x 128, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 1 65536 4 23.4 27.5 0.85x 128, 16 -> 128, 32
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 1 4 4000 4 18.8 21.9 0.86x 16, 16 -> 64, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 1 1 64000 4 22.1 25.7 0.86x 128, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 2 16384 16 26.4 30.7 0.86x 32, 16 -> 128, 32
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 1 2 4000 4 13.2 15.2 0.87x 32, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 16 16000 2 30.7 34.4 0.89x 32, 16 -> 64, 16
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 1 16000 8 15.3 17.0 0.90x 128, 16 -> 64, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 16 65536 2 87.4 95.9 0.91x 32, 16 -> 128, 16
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 16 1 16384 16 28.8 31.6 0.91x 32, 16 -> 128, 32
gpt-oss-120b 64 / 8 64 bf16 / fp8 yes 16 1 4096 4 10.0 10.8 0.93x 64, 16 -> 64, 32
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 1 1 4000 4 11.0 11.6 0.95x 64, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 1 1 16000 4 12.8 13.6 0.95x 128, 16 -> 128, 32
gpt-oss-120b 64 / 8 64 bf16 / fp8 yes 1 1 4000 4 10.8 11.2 0.96x 64, 16 -> 64, 32
gpt-oss-120b 64 / 8 64 bf16 / fp8 yes 16 1 16384 4 17.3 18.0 0.96x 64, 16 -> 64, 32

End to end, SGLang 52a6f70 (--attention-backend aiter --kv-cache-dtype fp8_e4m3, TP1, CUDA graphs on), real-text prompts, OSL 1000. For models without sliding window, SGLang's aiter backend takes this kernel only with SGLANG_USE_AITER_UNIFIED_ATTN=1 (not the default), which these runs set. Speculative decoding is either an EAGLE3 draft (--speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4) or NGRAM as a linear chain (--speculative-ngram-max-bfs-breadth 1 --speculative-num-draft-tokens 4), so verify runs q = 4 in both. The arms differ only in this DEFAULT.json, swapped through the aiter config directory on the image's aiter build. Each arm: 4 consecutive runs, a fresh server for each run, run 1 discarded, median of runs 2-4; order main -> this PR -> main repeat. A per-process probe confirmed that verify calls resolve to the new entries and that prefill stays on the 2D kernel.

model, speculative method ISL, conc total tok/s main this PR change main repeat drift mean TPOT mean TTFT
Qwen3-14B-FP8 (40q/8kv), NGRAM 32K, 16 12795 15360 +20.1% -0.2% -18.6% -0.3%
Qwen3-14B-FP8, NGRAM 8K, 32 12588 15291 +21.5% -0.4% -18.8% -0.4%
Llama-3.1-8B (32q/8kv), NGRAM 32K, 16 28649 31979 +11.6% +1.8% -10.7% +0.9%
Llama-3.1-8B, EAGLE3 32K, 16 21047 23269 +10.6% +0.4% -10.6% +2.8%
gpt-oss-120b (D 64, sinks), NGRAM 32K, 16 25418 26857 +5.7% +0.5% -7.2% +0.2%
gemma-4-26B-A4B (D 256/512, no match), none 8K, 32 21305 21249 -0.3% - +0.3% -0.1%

Within-arm spread of total tok/s is 0.2-3.7% (gpt-oss up to 4.4%). Speculative acceptance length is the same in both arms of every row (Qwen 1.5 / 1.4, Llama NGRAM 1.8, EAGLE3 1.1, gpt-oss 1.5), so the gain is attention time in verify. With the low acceptance on these prompts, speculative decoding is itself slower than plain decode here; every row compares the two configs with the same speculative setting. The gemma-4 row compares main's table with that all-q >= 2 variant on the same SGLang commit; head_size 256/512 reaches neither. MiniMax-M3 TP4 + EAGLE3, measured with the same variant on SGLang 3ea7e210, was flat: only 4 of its 61 attention call sites per step take this kernel.

Accuracy, GSM8K with the Llama-8B + EAGLE3 server, chat API, greedy, 1319 questions: main 0.8021, this PR 0.8044, main repeat 0.7953.

Test plan

On gfx950, rebased on main 036a6ce:

  • test_unified_attention.py::test_triton_unified_attn, triton backend ids (2048): 1024 passed, 1024 skipped (skip shuffled_kv_cache, 2d gluon not available), the same counts as before the rebase and as with main's table.
  • test_unified_attention.py::test_triton_unified_attn_3d: 64 passed, 32 skipped.
  • New test_triton_unified_attn_qq_bias_padded_query_rows: 5 passed. It covers 32q/8kv at q = 3 and 4 and 16q/1kv at q = 3 with KV 4096 on the 3D path, and 32q/8kv at q = 3 and 5 with KV 256 on the 2D path. ref_paged_attn gains an optional qq_bias.
  • Guard-page probe above: memory fault on main and without the mask, clean with this PR.
  • Lookup scan above: q = 1 configs unchanged for every head_size, dtype and sliding-window setting.

The micro and end-to-end numbers were measured on main be0237f. Between that commit and 036a6ce, this table, the lookup and the triton kernels are unchanged. The wrapper changed in two ways: #6581 casts NUM_SEGMENTS and TILE_SIZE to Python ints, which gives the same launch values, and #6896 is gfx1250 only. The new mask only compiles in when qq_bias is passed.

Not covered: vLLM end to end (no vLLM environment was available), and head_size 96, which falls in the changed range.

Powered by Hyperloom, with human review, writing and testing.

@github-actions

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:gfx1250-ffm-triton Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x Run an additional Triton test job on MI300X in PRs (added automatically when gfx942 configs change); main branch always runs both MI35X and MI300X
ci:triton-355 Run the full Triton test suite on MI35X, not only the tests the change affects
ci:windows Windows JIT op smoke test on the gfx1201 ROCm runner
ci:multi-gpu Aiter multi-GPU tests on the 8-GPU runner
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: DeepSeek-V4-Pro, MiniMax-M3-MXFP4
ci:all All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 6983 --add-label <label>

One backend per PR:
A PR changes one kernel backend: [Triton/Gluon] (Triton and Gluon count as one), [HIP], [ASM], [CK], [OPUS] or [FlyDSL]. If the title ends up with two backend tags, split the PR -- as stacked pull requests when one part cannot merge without the other.

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to stop the title rewrites; labels stay in sync either way.

@lishuoshuo-amd

Copy link
Copy Markdown
Contributor Author

@cagrikymk @vgokhale Could you review? It only touches the gfx950 triton unified-attention DEFAULT.json; q = 1 resolves to the same configs as main (lookup scan in the description). Fork PR, so the workflow needs a maintainer approval to run. Thanks!

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

The wider query tile exposes out-of-bounds query-bias reads because the 3D kernel’s bias load does not mask padded query rows.

1 open issue
What changed in this PR

This PR tunes gfx950 Triton unified attention for multi-token queries on the split-KV path, while preserving single-token decode settings.

Changes:

  • Adds query-length (Q) matching to the 3D attention and KV-split schemas.
  • Increases query tile size for head sizes ≤128 when Q≥4.
  • Adjusts KV splitting for head sizes 65–128 when Q≥2, preserving existing splits for smaller heads.
File Description
aiter/​ops/​triton/​configs/​gfx950/​triton/​attention/​unified_attention/​DEFAULT.json Adds query-length-specific tile and split settings for gfx950.

Review effort: Balanced

@cagrikymk

Copy link
Copy Markdown
Contributor

@cagrikymk @vgokhale Could you review? It only touches the gfx950 triton unified-attention DEFAULT.json; q = 1 resolves to the same configs as main (lookup scan in the description). Fork PR, so the workflow needs a maintainer approval to run. Thanks!

Thank you for the contribution!
Could you add a bit more data (maybe a table) about your regression case? With dtype, num q/kv heads, kv_len, batch size etc.
I see that the regressions are relatively small compared to the upside but having a table about them would make it clear where it regresses (time in us for both arms)

@lishuoshuo-amd
lishuoshuo-amd force-pushed the perf/gfx950-ua-3d-decode-q2 branch from 603961c to 9d1fd50 Compare October 28, 2026 03:11
@lishuoshuo-amd

lishuoshuo-amd commented Oct 28, 2026 •

Copy link
Copy Markdown
Contributor Author

Could you add a bit more data (maybe a table) about your regression case? With dtype, num q/kv heads, kv_len, batch size etc.
I see that the regressions are relatively small compared to the upside but having a table about them would make it clear where it regresses (time in us for both arms)

@cagrikymk Thanks for the review! Below are all the q >= 2 shapes that get slower: 16 of the 143 measured are below 0.97x.

How the numbers were taken:

  • MI355X, aiter main be0237f, triton 3D kernel.
  • CUDA-graph timing per call, averaged over two GPUs. main / PR is the geomean of the two per-GPU ratios.
  • Page size 1 is the SGLang KV layout; 16 is the vLLM layout.
  • The last column is the 3D kernel's NUM_SEGMENTS and BLOCK_M on each side.
model q / kv heads head size q / KV dtype sinks page size batch kv_len q_len main (us) this PR (us) main / PR segments, BLOCK_M: main -> PR
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 1 1 150000 4 39.4 48.4 0.81x 128, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 1 150000 4 40.2 48.5 0.83x 128, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 1 65536 4 23.4 27.5 0.85x 128, 16 -> 128, 32
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 1 4 4000 4 18.8 21.9 0.86x 16, 16 -> 64, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 1 1 64000 4 22.1 25.7 0.86x 128, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 2 16384 16 26.4 30.7 0.86x 32, 16 -> 128, 32
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 1 2 4000 4 13.2 15.2 0.87x 32, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 16 16000 2 30.7 34.4 0.89x 32, 16 -> 64, 16
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 1 16000 8 15.3 17.0 0.90x 128, 16 -> 64, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 16 16 65536 2 87.4 95.9 0.91x 32, 16 -> 128, 16
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 16 1 16384 16 28.8 31.6 0.91x 32, 16 -> 128, 32
gpt-oss-120b 64 / 8 64 bf16 / fp8 yes 16 1 4096 4 10.0 10.8 0.93x 64, 16 -> 64, 32
Llama-3.1-8B 32 / 8 128 bf16 / fp8 no 1 1 4000 4 11.0 11.6 0.95x 64, 16 -> 128, 32
MiniMax-M3 TP4 16 / 1 128 bf16 / fp8 no 1 1 16000 4 12.8 13.6 0.95x 128, 16 -> 128, 32
gpt-oss-120b 64 / 8 64 bf16 / fp8 yes 1 1 4000 4 10.8 11.2 0.96x 64, 16 -> 64, 32
gpt-oss-120b 64 / 8 64 bf16 / fp8 yes 16 1 16384 4 17.3 18.0 0.96x 64, 16 -> 64, 32

Where it regresses:

  • 13 of the 16 run 1 or 2 sequences.
  • The worst rows are 16q/1kv at q=4 with batch 1 and long context. There BLOCK_M 32 halves the number of query blocks while main's split is already at MAX_SEGMENTS 128, so the grid shrinks.
  • The Llama-3.1-8B rows at q=4, 4K context, 2-4 sequences use the SGLang layout. SGLang passes the page-table width (70000 in this benchmark) as max_seqlen_k, so the split is sized for 70K rather than the 4K context. With the new kv_split entry that becomes 64-128 segments, and most of them cover no keys.
  • The lookup has no batch-size axis, so the table cannot separate these batch-1/2 shapes from the batch-16+ shapes that gain. Doing that would need a batch-aware key in the lookup, which would be a separate change.

The push also fixes the out-of-bounds qq_bias read Copilot flagged; I replied in that thread.

@github-actions github-actions Bot changed the title [Triton/Gluon] [Config] [gfx950] unified_attention: split-KV 3D configs for head_size <= 128 keyed on max_seqlen_q [Triton/Gluon] [gfx950] unified_attention: split-KV 3D configs for head_size <= 128 keyed on max_seqlen_q Oct 28, 2026
@github-actions github-actions Bot removed the Config label Oct 28, 2026
…_size <= 128 keyed on max_seqlen_q

Add a Q axis to the attn_3d and kv_split schemas of the gfx950 triton
unified-attention table:

- attn_3d D_LEQ_128.Q_GEQ_4: BLOCK_M 32.
- kv_split D_LEQ_128.Q_GEQ_2: 16 segments per CU with a 256-token segment
  floor.
- kv_split D_LEQ_64: main's D_LEQ_128 split, so head_size <= 64 keeps it.

max_seqlen_q == 1 resolves to the same configs as before.

Mask padded query rows when loading qq_bias in the 2D and 3D triton
kernels. A query length that is not a multiple of BLOCK_Q read qq_bias
rows past the tensor; the wider BLOCK_M makes that reachable at q = 4.
Add a qq_bias test with padded query rows on both paths.
@lishuoshuo-amd
lishuoshuo-amd force-pushed the perf/gfx950-ua-3d-decode-q2 branch from 9d1fd50 to 935301c Compare October 29, 2026 02:26
@lishuoshuo-amd

Copy link
Copy Markdown
Contributor Author

@cagrikymk The gfx1250 skip is in 935301c (thread above). Could you take another look when you have time? The workflows also still need a maintainer approval to run. Thanks!

@aiter-bot

aiter-bot commented Oct 30, 2026 •

Copy link
Copy Markdown

NOT READY FOR HUMAN REVIEW

PR Self-Review

Revision b4da6fcad69e · mode shadow

Waiting for the PR author. The readiness report will appear after the self-review answers are complete for this revision.

Questions to answer in the PR description

  1. What behavior does this PR change, and why is the existing behavior insufficient?
  2. What did you validate, and what remains unverified? Include commands, hardware/configuration and results.
  3. Which exact workload keys consume these configuration changes, and what evidence shows that unrelated shapes keep their previous selection?
  4. For every changed op_tests/ file, what behavior does it cover and why is it not redundant? Write None if no op test changed.

Edit the PR Self-Review section in the PR description. Include the current revision SHA; this comment refreshes automatically.

Missing or invalid answers

  • FAIL — missing ### Reviewed revision in the PR description
  • FAIL — missing ### Behavior and intent in the PR description
  • FAIL — missing ### Validation and unverified scope in the PR description
  • FAIL — missing ### PR-specific question in the PR description
  • FAIL — missing ### Changed op tests in the PR description
  • FAIL — Reviewed revision must contain at least the first 7 characters of the current head SHA
  • FAIL — Behavior and intent has no substantive answer
  • FAIL — Validation and unverified scope has no substantive answer
  • FAIL — PR-specific question has no substantive answer
  • FAIL — changed op test not explained in Changed op tests: op_tests/triton_tests/attention/test_unified_attention.py
  • FAIL — Changed op tests must explain why each changed op test is necessary and non-redundant

Shadow mode reports what would block but keeps the status successful. Set AITER_PR_SELF_REVIEW_MODE=enforce only after reviewing real false positives.

@cagrikymk
cagrikymk self-requested a review October 30, 2026 02:03
@cagrikymk

Copy link
Copy Markdown
Contributor

@lishuoshuo-amd LGTM! Triggered the CI

@lishuoshuo-amd

Copy link
Copy Markdown
Contributor Author

@cagrikymk Thanks for the approval! Triton Test and OPUS Test passed. The only failure is one Aiter Test shard (Standard Tests (1 GPU) ... 8, 7), where the self-hosted runner lost communication before the tests started, so the gate failed with it. Could you re-run the failed jobs?

@lishuoshuo-amd

Copy link
Copy Markdown
Contributor Author

@azaidy @cagrikymk @Boss2002n This PR has been approved; could you please re-trigger the CI? Thanks!

@cagrikymk

Copy link
Copy Markdown
Contributor

@azaidy @cagrikymk @Boss2002n This PR has been approved; could you please re-trigger the CI? Thanks!

@lishuoshuo-amd Triggered the CI and merged main to the branch

@lishuoshuo-amd

Copy link
Copy Markdown
Contributor Author

@cagrikymk The CI has finished running; could you please help with the merge? Thanks!

@Boss2002n
Boss2002n merged commit 0df10f8 into ROCm:main Nov 3, 2026
86 of 108 checks passed
yuychang pushed a commit to yuychang/aiter that referenced this pull request Nov 4, 2026
- Resolve conflicts with ROCm#6889 (window-aware split) and ROCm#6983 (Q axis in the attn_3d and kv_split schemas).
- Kernels: keep main's split_window_tiles carving next to the length-aware per-sequence segment count.
- compute_segment_params: keep the num_segments helper and apply main's SPLIT_FIT_BUDGET / SPLIT_SEQS_PER_TILE in it.
- Config: keep kv_split D_EQ_64.DT_bf16_fp8 next to main's new D_LEQ_64 and D_LEQ_128.Q_GEQ_2 entries.
- Config: add SPLIT_RUNS_PER_XCD to the new attn_3d D_LEQ_128.Q_GEQ_4 entry so q>=4 length-aware calls do not assert.
yuychang pushed a commit to yuychang/aiter that referenced this pull request Nov 4, 2026
- Main's attn_3d D_LEQ_128.Q_GEQ_4 entry (ROCm#6983) sets BLOCK_M 32, which pads q4 verify blocks to twice their rows.
- A new SPLIT_BLOCK_M key cuts BLOCK_M to max(SPLIT_BLOCK_M, query rows) on the length-aware launch only.
- The eligibility check still uses the tuned BLOCK_M, and the uniform split keeps BLOCK_M 32.
- Test: a q=8 skewed case and q=1/4/8 config-key checks in test_length_aware_split_scope.
yuychang pushed a commit to yuychang/aiter that referenced this pull request Nov 4, 2026
…usal

- Pick up ROCm#7088's merge with main (ROCm#6983 Q-axis configs through ROCm#7302) so this stacked PR is conflict-free again.
- Pick up ROCm#7088's SPLIT_BLOCK_M, which keeps q4 draft verify blocks at 16 rows under main's BLOCK_M 32 entry.
- No textual conflicts; the CAUSAL constexpr composes with the merged split code.
yuychang added a commit to yuychang/aiter that referenced this pull request Nov 4, 2026
…usal

- Pick up ROCm#7088's merge with main (ROCm#6983 Q-axis configs through ROCm#7302) so this stacked PR is conflict-free again.
- Pick up ROCm#7088's SPLIT_BLOCK_M, which keeps q4 draft verify blocks at 16 rows under main's BLOCK_M 32 entry.
- No textual conflicts; the CAUSAL constexpr composes with the merged split code.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants