diff --git a/perf_knowledge/expert_skills/index.yaml b/perf_knowledge/expert_skills/index.yaml index a8093a317..60013b338 100644 --- a/perf_knowledge/expert_skills/index.yaml +++ b/perf_knowledge/expert_skills/index.yaml @@ -120,6 +120,33 @@ skills: e2e_delta_min_pct: 1.0 parity: required validation_status: validated +- id: flydsl_fused_attention_backward + file: skills/flydsl_fused_attention_backward/skill.md + scope: kernel + match: + operator: + - attention_prefill_fmha + - gqa_mqa_attention + - mla_attention + arch_class: + - '*' + gens: + - gfx942 + - gfx950 + dtypes: + - bf16 + - fp16 + regimes: + - training + from_backend: '' + to_backend: flydsl + profile_signature: + op_name_regex: fmha_bwd|mha_bwd|_bwd_kernel|flash.*bwd|attn.*backward + min_pct_gpu: 15.0 + expects: + isolated_speedup_min: 1.1 + parity: required + validation_status: draft - id: flydsl_prefill_moe_stage2_fp8partial file: skills/flydsl_prefill_moe_stage2_fp8partial/skill.md scope: kernel diff --git a/perf_knowledge/expert_skills/skills/flydsl_fused_attention_backward/skill.md b/perf_knowledge/expert_skills/skills/flydsl_fused_attention_backward/skill.md new file mode 100644 index 000000000..3c1d156f9 --- /dev/null +++ b/perf_knowledge/expert_skills/skills/flydsl_fused_attention_backward/skill.md @@ -0,0 +1,167 @@ +--- +id: flydsl_fused_attention_backward +title: "Attention backward: author a fused multi-GEMM FlyDSL kernel (CDNA3/CDNA4)" +kind: expert_skill +authors: [GEAK Team] +scope: kernel +# ---- selector: the workflow matches these against the live bottleneck ---- +match: + operator: [attention_prefill_fmha, gqa_mqa_attention, mla_attention] + arch_class: ['*'] + gens: [gfx942, gfx950] + dtypes: [bf16, fp16] + regimes: [training] # training-only: no inference kernel can match this selector + from_backend: "" # any — asm/ck/triton/fa_rocm backward, or no backward at all + to_backend: flydsl + profile_signature: + op_name_regex: fmha_bwd|mha_bwd|_bwd_kernel|flash.*bwd|attn.*backward + min_pct_gpu: 15.0 +# ---- expected effect: the validation gate's pass criteria ---- +expects: + isolated_speedup_min: 1.10 # general authoring bar; the measured instance reached 1.38 (see Sources) + parity: required # cos >= 0.999 on every gradient vs an fp32 reference +# ---- validation: AUTO-FILLED by validate_skill.py — do NOT hand-edit ---- +validation: + status: draft + last_verified: "" + gpu: "" + model: "" + measured: {isolated: "", e2e_pct: "", parity: ""} + artifact: "" +role: advisory_prior +supersedes: [] +--- + +## When to use +An attention **backward** kernel is a top GPU-time entry in a training or fine-tuning run and the live +path is a poor fit for the shape. The two situations that make authoring worthwhile: + +- **The vendor fast path does not accept your head dims and pads to reach them.** Padding costs work + proportional to the dim ratio, and it can also cost you the fast *forward*: the aiter gfx942 v3 + backward dispatch is gated on `hdim_q == hdim_v` while its ASM forward requires `hdim_v == 128`, so an + asymmetric-head-dim model has no configuration where both kernels are fast. +- **No tuned backward exists for the variant at all** (GQA/MQA group ratios, small head dims, unusual + causal alignment), leaving a generic CK-tile or Triton fallback. + +**Do not use it when a tuned asm/CK backward dispatches natively for your exact shape** — that path is +strong and this recipe is unlikely to beat it. Check the dispatch conditions first; that check is +minutes and decides whether the rest of the work has any headroom. + +## Mechanism +The performance comes from structure, not from instruction-level tuning, and the binding constraint is +the **256 arch-VGPR cap** rather than the tile size: a fused backward carries several long-lived operand +sets *plus* accumulators across the whole loop, so most classical "fewer instructions" optimizations +lose to spill. The four structural decisions, in the order they must be made: + +1. **Fuse all the GEMMs into one kernel.** Splitting one out forces the scores and their gradient to be + recomputed there, because those intermediates are too large to spill. Per score element the fused + chain is `2·(HD_QK + HD_V + HD_V + HD_QK + HD_QK)` FLOPs; evaluate the split's overhead for your head + dims before choosing (~+33% at symmetric 128, more as the dims diverge). +2. **Choose the operand-swapped MFMA fragment orientation** so the score comes out as `S[qrow, kpos]` — + already the A-operand layout that both dV and dK need, since both contract over q rows. The + probabilities and their gradient then **never leave registers and never transit LDS**. Free if chosen + first, a rewrite if retrofitted. +3. **Loop KV-outer.** KV rows become exclusive to a block, so dK/dV are stored directly with no atomics + and no workspace, and only dQ accumulates. The q-outer orientation puts the larger tensors on atomics + instead — compare the fp32 workspace volumes for your shape rather than assuming. +4. **Pick the wave count against the register cap, not by occupancy intuition.** Fewer waves gives each + wave more registers but doubles its slice of the split axis, hence doubles the operand set that must + live in arch VGPRs — against a cap that does **not** move with occupancy. This is why the + fewer-waves direction can lose here while it wins on GEMM. + +Full derivations, the geometry constraints and the anti-pattern table are in +[[languages/flydsl/authoring_attention_levers]]; this skill is the regulated procedure that applies them. + +## Procedure +Build in dependency order; each phase gates the next. Add the structural optimisations **one at a time** +(step 5) or an unexplained regression becomes unattributable. + +1. **fp32 reference first.** A varlen shape descriptor (seqlens, nheads, `HD_QK`, `HD_V`, softmax scale, + `cu_seqlens`), an fp32 forward returning `(out, lse)` and backward returning every gradient plus `D`, + and a `compare` reporting cosine / `max_err/scale` / `rms/scale`. Fix the **causal convention here** + (bottom-right aligned: `delta = seqlen_k − seqlen_q`, keep `j <= i + delta`) — top-left alignment is a + different problem that silently produces high-cosine garbage. Validate the reference itself against an + independent forward + autograd; a subtly wrong reference certifies a matching bug in the kernel. +2. **Harness before kernel.** Build *all* variants up front, then time them round-robin in one process + keeping the per-variant minimum ([[profiling/benchmarking_methodology]] §single-kernel variant + sweeps). Include a smoke shape with at least one non-tile-multiple length so tail handling runs every + time. Add an **ISA probe** now, not later — register allocation is the dominant failure mode and is + invisible from timings. +3. **Preprocess kernel:** `D = rowsum(dO ⊙ O)` into a **head-major** `[nheads, seqlen]` fp32 tensor. The + main kernel must index `D` and `lse` the same way; transposing this still yields plausible-looking + gradients on uniform shapes, so gate it explicitly. +4. **Main kernel.** Derive all geometry from the tile shape and assert the divisibility constraints at + build time (see the levers card). `lse` is natural-log with the scale folded in, as a + Flash-Attention forward emits it, so the kernel computes `p = exp2(s·scale·log2e − lse·log2e)`. + Required structure: + - **Three LDS tiles, not five** — read the second orientation of each staged tile out of the single + copy instead of materialising a transposed one ([[optimization/lds_and_bank_conflicts]]). + - **The causal cost axis on the SLOWEST grid axis** (longest-job-first) + ([[optimization/xcd_l2_locality]]). + - **One scheduling region, two barriers** — double-buffer the staging and issue the next tile's loads + inside this tile's MFMA chain with no barrier between them; keep the staging branch-free by clamping + rather than guarding ([[optimization/memory_pipelining]]). + - **Re-partition the waves** for the output that contracts over the split axis, instead of reducing + partials across waves. +5. **Bring-up ladder**, layout → scheduling → dispatch, measuring each step, and dumping the ISA at + every one. A step that improves instruction count but spills is a regression. +6. **Head-to-head** against the live path interleaved in one process — and assert the baseline is + actually computing gradients rather than silently no-op'ing on an unsupported configuration. + +## Knobs & pitfalls +Expose `qk_head_dim`, `v_head_dim`, `block_m`, `block_n`, `num_waves`, `dq_groups`, `transposed_tiles`, +`pipeline`, `schedule`, `sched_steps`, and the grid-order flag, with defaults at the measured optimum. +Keep the three ablation paths (`transposed_tiles=True`, `pipeline=False`, grid-order off) alive +permanently — they are what make the step-5 ladder re-checkable in one interleaved process on new +hardware, and they let a regression be bisected against a mechanism rather than a commit. + +Gates, in order: parity on every gradient (`cos ≥ 0.999`, relative bound ~3e-2) → resolved-geometry echo +→ **arch VGPR / AGPR / `v_accvgpr` move count** → the dispatch-order ablation ratio against its scheduling +model → the MFMA counter matching the decomposition exactly. + +**Judge the wave count from the ISA, not the clock.** In the measured instance the best 4-wave build +reported *lower* `private_seg_size` than the 8-wave build and was still 15% slower, because the spill +relocated into the accumulator file where it does not appear as scratch. Count `v_accvgpr` moves +([[optimization/occupancy_and_registers]]). + +**Dead ends, built and measured — do not re-derive** (mechanisms in +[[languages/flydsl/authoring_attention_levers]] §anti-patterns): growing the KV tile beyond the point +where the operand set fits the arch-VGPR cap; halving the q tile to relieve it; fanning dQ atomics across +more slots; `sched_group_barrier` hints once the region is already branch-free and single; hoisting +invariant staging address math out of the hot loop; folding the loop to a single barrier by deferring the +dQ phase; splitting the longest blocks to collect a low occupancy figure; coalescing a once-per-block +prologue. + +## Do-no-harm notes +- **Training-only selector.** `regimes: [training]` means no prefill or decode kernel can match this + skill, so nothing on a serving path is affected by its presence. It also lands `validation.status: + draft`, so it is not auto-applied even when `use_expert_skills` is on (itself default OFF). +- **Symmetric head dims are usually not worth authoring.** Where `hdim_q == hdim_v`, the tuned asm + backward dispatches natively and is strong; the recipe should decline rather than compete. +- **The speedup bar is generic; the measured instance is not.** `isolated_speedup_min: 1.10` reflects the + shape-independent structural wins. The 1.38× (per unit of useful work) in the instance below also + includes the baseline's padding tax, which exists only on gfx942 with asymmetric head dims — aiter + supports native asymmetric dims on gfx950. Do not carry the ratio across gens. +- **Re-derive the wave count per gen.** gfx950's larger register file is the one place the fewer-waves + schedule may win; the gfx942 result is not evidence about gfx950 either way. +- **Confirm the compute partition mode (SPX/NPS1) before trusting any timing.** The grid-order lever + depends on workgroups reaching CUs in x-fastest order with static XCD assignment; CPX changes that and + the device-scope atomic behaviour. A dispatch-order ablation that comes out neutral when the model says + otherwise is usually a partition-mode problem, not a kernel problem. +- **Non-causal removes the cost gradient** that longest-job-first exploits — re-derive the grid order + rather than assuming it still wins. +- Budget the output-cast overhead as this recipe's own cost if the epilogue leaves gradients in fp32; the + vendor path typically folds that cast into an existing pass. + +## Sources +- Levers, derivations and anti-patterns (general): [[languages/flydsl/authoring_attention_levers]]. +- Worked instance with per-step deltas, counter expectations and the full dead-end ledger — fused + five-GEMM MLA backward, `HD_QK=192`/`HD_V=128`, bf16, causal, varlen, gfx942 MI300X/MI325X SPX/NPS1, + ROCm 7.1.0: [[operators/mla_attention/backends/flydsl]]. Measured 715–745 µs + vs 984 µs for the padded-ASM chain (−27.3% wall while executing 15.4% less work, so 1.38× per unit of + useful work), `cos = 0.999999` on all three gradients, and 1.49×/1.79× e2e against the two aiter + configurations. +- Wave-count counter-measurement against HipKittens §3.3.2: [[optimization/mfma_scheduling]] · + [[languages/hipkittens/perf_findings]]. +- FlyDSL authoring surface: [[languages/flydsl/authoring_optimization]] (register-allocation gate) · + [[languages/flydsl/debugging]]. diff --git a/perf_knowledge/index/capability_index.yaml b/perf_knowledge/index/capability_index.yaml index 06ec9244a..590beecbf 100644 --- a/perf_knowledge/index/capability_index.yaml +++ b/perf_knowledge/index/capability_index.yaml @@ -649,6 +649,15 @@ candidates: - https://github.com/Dao-AILab/flash-attention - https://vllm.ai/blog/2026-02-27-rocm-attention-backend - ROCm/aiter@a6bb499375849eec45d68c5ccaebc8865fd422c0:aiter/mla.py + - operator: mla_attention + backend: flydsl + gens: [gfx942] + dtypes: [bf16] + regimes: [training] + card: ../operators/mla_attention/backends/flydsl.md + sources: + - Attention-Kernels/geak_trans_py2flydsl/fmha_backward/FMHA_BWD_FlyDSL_Skills.md + - https://arxiv.org/abs/2511.08083 - operator: mla_attention backend: hip gens: [gfx942, gfx950] diff --git a/perf_knowledge/index/changelog.md b/perf_knowledge/index/changelog.md index 58789915b..902d16465 100644 --- a/perf_knowledge/index/changelog.md +++ b/perf_knowledge/index/changelog.md @@ -78,3 +78,14 @@ - **No operator cards and no new operator.** The index is untouched by content: same 54 operators, same 225 cards, no `taxonomy.md` change. Only `expert_skills/index.yaml` gains the new skill. - Deliberately **no SOTA cards** in this ingest. No Gluon attention kernel exists, forward or backward, published or measured, so any card would have been a route sketch rather than a result — and a `sota_card` whose measured column is empty invites being read as evidence. Cards land when there is a Gluon number to anchor them; until then the skill carries the mechanics and `languages/gluon/gemm_cookbook.md` carries the one real Gluon measurement AMD has published (near-peak GEMM on gfx950). - The skill lands as `validation.status: draft` (no on-box Gluon A/B), so it is **not** auto-applied; `use_expert_skills` remains default-OFF and the OFF-identical regression is unaffected. Upstream stays the SSOT; this is a one-way snapshot recorded in `## Sources`. + +## FlyDSL attention-authoring levers + gfx942 wave-count correction (2026-08-12) +- Ingested `Attention-Kernels/geak_trans_py2flydsl/fmha_backward/FMHA_BWD_FlyDSL_Skills.md` (882-line first-party build playbook) as a **four-tier merge**, not a copy: the dimension-agnostic levers became a new language card, the regulated procedure an expert skill, the numbers a SOTA card, and the transferable mechanisms amendments to existing `optimization/` cards. Measured on gfx942 MI300X/MI325X, ROCm 7.1.0, SPX/NPS1: a fused five-GEMM MLA 192/128 causal varlen backward at **715–745 µs** vs **984 µs** for aiter's ASM backward with V padded 128→192, and **1.49× / 1.79×** e2e against the two aiter configurations. +- **New**: `languages/flydsl/authoring_attention_levers.md` — the sibling `authoring_gemm_levers.md` had no counterpart for **fused multi-GEMM** kernels, which is where FlyDSL authoring actually differs: the binding constraint is the per-wave resident *operand* set against the 256 arch-VGPR cap, not the tile size, so a GEMM's register intuitions invert. Carries the three pre-code decisions (fusion boundary as a symbolic FLOP ratio, MFMA fragment orientation, which output goes on atomics), the geometry derivations with their divisibility constraints, six ordered levers, five verification gates and the anti-pattern table — all in terms of `HD_QK`/`HD_V`/`BM`/`BN`/`NW`, with the 192/128 build cited as one resolution rather than a default. Linked from `overview.md`'s deep-dive map. +- **New**: `operators/mla_attention/backends/flydsl.md` (`status: sota`, `gens: [gfx942]`, `regimes: [training]`) — the first FlyDSL card for any attention operator and the first backward/training card in the attention family, framed as the measured instance of the levers card. `mla_attention/overview.md` gains `training` to its regimes and a "Training / backward" shape regime documenting the dispatch trap: aiter's gfx942 ASM bwd is gated on `hdim_q == hdim_v` and its ASM *forward* on `hdim_v == 128`, so **no aiter configuration on gfx942 gets both fast**. Registries regenerated from card frontmatter (226 cards; note `_gen_registry.py` writes LF while the committed indices are CRLF — restore CRLF after running it, or the whole file shows as changed). +- **New**: `expert_skills/skills/flydsl_fused_attention_backward/` (`scope: kernel`, `operator: [attention_prefill_fmha, gqa_mqa_attention, mla_attention]`, `regimes: [training]`, `gens: [gfx942, gfx950]`, `to_backend: flydsl`, `isolated_speedup_min: 1.10`). Deliberately **not** an MLA-192/128 selector: the procedure is the same for any fused attention backward, so the selector is the attention family restricted to `training`, and the bar is the shape-independent structural win — the instance's 1.38×/unit-work also collects the baseline's padding tax, which exists only on gfx942 with asymmetric head dims. `regimes: [training]` is what guarantees no prefill/decode kernel can match it, so nothing on a serving path is touched. The manifest measures **two** shape classes (symmetric head dims where a tuned backward dispatches natively, as the generality/do-no-harm run, and the asymmetric instance at `expect_min: 1.25`) and marks gfx950 and non-causal `reverify_knobs` rather than carrying the ratio over. Lands **`draft`** — first-party numbers, but the manifest has not been re-run through `validate_skill.py`, so it is not auto-applied. +- **Correction, not just an addition — and scoped so no GEMM guidance moves.** `optimization/mfma_scheduling.md` previously carried "reach for 8-wave ping-pong or 4-wave interleave" as an unqualified gfx942/gfx950 prior, and `operators/gqa_mqa_attention/backends/hipkittens.md` said "4-wave interleave often wins the backward" — both inherited from HipKittens §3.3.2, author-reported on MI355X and never re-measured. The amendment states the **decision rule** (is the per-wave resident set accumulator-bound and retileable, or operand-bound and scaling with the work slice?) with a two-row table for which way each resolves, and explicitly records that **the GEMM prior is unchanged** — dense/scaled GEMM keeps both starting points and their documented ranking. Only the attention-backward half is corrected, with the measured counter-example: every 4-wave build 15% slower, all shuttling AGPRs, because halving the wave count doubles the per-wave K/V/Kᵀ set against a **256 arch-VGPR cap that does not move with occupancy**. `languages/hipkittens/primitives.md` and the GQA card now scope the paper's backward row to its own shape (symmetric `hdim=128`, non-causal, MI355X). Per the sourcing doctrine the on-box measurement outranks the vendor-labeled claim; HK's own numbers are left intact. +- **Transferable mechanisms** amended into the cross-cutting cards, each with its measured delta and its mechanism: `lds_and_bank_conflicts.md` §5 — **read an operand in the awkward orientation rather than materialising the transpose** (−19%; conflict-free by construction, no XOR swizzle, since the feared conflict lives on the scatter side of a materialised transpose); `occupancy_and_registers.md` — **count `v_accvgpr` moves, not scratch** (AGPR shuttling never appears as `private_seg_size`, so a build can report *less* scratch and be 15% slower), plus the corollary that dropping wave count does not relieve arch-VGPR pressure; `xcd_l2_locality.md` — dispatch order as a **makespan** lever (longest-job-first grid axis, −20%, matching 114/144), model **8 static XCD queues of ~38 CUs** rather than one CU pool, and empty workgroups are free appended but **+29% interleaved**; `languages/flydsl/authoring_optimization.md` — the register-allocation gate as a per-step check, with the two measured spill regressions (−8.6% hoisting address math, −36% folding to one barrier) and why `s_barrier` ending a scheduling region makes branch-free staging worth 7.3%. +- **One doctrine conflict resolved rather than papered over.** `profiling/benchmarking_methodology.md` mandated median-of-7-with-spread; the playbook mandated the minimum, because sustained replays throttle and the median of an *unchanged* kernel drifted 438 → 534 µs. Added as a **scoped exception**: minimum-over-interleaved-replays for isolated single-kernel variant sweeps (build all variants up front, time round-robin in one process, no CUDA graphs), median+spread unchanged as the e2e rule. +- **Transferable mechanisms are written rule-first**, with the 192/128 build demoted to a labelled measured instance in each: the LDS rule is stated as "an MFMA operand does not need its contraction-axis elements to arrive in one load", the register rule as "scratch is not the whole spill surface", the dispatch rule as "when per-block cost varies systematically along one grid axis, put that axis on the slowest one". Each holds for any kernel with the stated property; the attention backward is just where they were measured. +- **No new operator, and nothing existing changes behaviour.** Backward is a `regime` in `taxonomy.md`, not an operator, so the taxonomy is untouched: same 54 operators, 226 cards. No existing card's recommendation is reversed, no default or tuned config is altered, and the one prior that is qualified (wave count) is qualified only for operand-bound kernels. The new skill is `draft` and `regimes: [training]`, so it cannot fire on any current inference path even with `use_expert_skills` on (itself default OFF). diff --git a/perf_knowledge/index/sota_matrix.md b/perf_knowledge/index/sota_matrix.md index 3b0190954..c38189a13 100644 --- a/perf_knowledge/index/sota_matrix.md +++ b/perf_knowledge/index/sota_matrix.md @@ -3,7 +3,7 @@ AUTO-GENERATED from per-card frontmatter (`index/_gen_registry.py`). Each cell links to the SOTA card. Legend: 🟢 sota · 🟡 competitive · 🧪 experimental · 🟤 legacy · ⚪ na · `·` no card. -Coverage: **54 operators**, **225 backend cards**. +Coverage: **54 operators**, **226 backend cards**. ## GEMM | operator | triton | flydsl | hip | ck | asm | tilelang | gluon | hipkittens | rocwmma | aiter | hipblaslt | @@ -21,7 +21,7 @@ Coverage: **54 operators**, **225 backend cards**. |---|---|---|---|---|---|---|---|---|---|---|---| | [attention_prefill_fmha](../operators/attention_prefill_fmha/overview.md) | [🟡](../operators/attention_prefill_fmha/backends/triton.md) | · | · | [🟢](../operators/attention_prefill_fmha/backends/ck.md) | [🟢](../operators/attention_prefill_fmha/backends/asm.md) | [🟡](../operators/attention_prefill_fmha/backends/tilelang.md) | [🟢](../operators/attention_prefill_fmha/backends/hipkittens.md) | [🟢](../operators/attention_prefill_fmha/backends/aiter.md) | [🟡](../operators/attention_prefill_fmha/backends/fa_rocm.md) | · | · | | [attention_decode_paged](../operators/attention_decode_paged/overview.md) | [🟡](../operators/attention_decode_paged/backends/triton.md) | · | [🟢](../operators/attention_decode_paged/backends/hip.md) | [🟡](../operators/attention_decode_paged/backends/ck.md) | · | · | · | [🟢](../operators/attention_decode_paged/backends/aiter.md) | [🟡](../operators/attention_decode_paged/backends/fa_rocm.md) | · | [🟢](../operators/attention_decode_paged/backends/vllm_kernels.md) | -| [mla_attention](../operators/mla_attention/overview.md) | [🟡](../operators/mla_attention/backends/triton.md) | · | [🟡](../operators/mla_attention/backends/hip.md) | [🟡](../operators/mla_attention/backends/ck.md) | · | · | · | [🟢](../operators/mla_attention/backends/aiter.md) | [⚪](../operators/mla_attention/backends/fa_rocm.md) | · | · | +| [mla_attention](../operators/mla_attention/overview.md) | [🟡](../operators/mla_attention/backends/triton.md) | [🟢](../operators/mla_attention/backends/flydsl.md) | [🟡](../operators/mla_attention/backends/hip.md) | [🟡](../operators/mla_attention/backends/ck.md) | · | · | · | [🟢](../operators/mla_attention/backends/aiter.md) | [⚪](../operators/mla_attention/backends/fa_rocm.md) | · | · | | [gqa_mqa_attention](../operators/gqa_mqa_attention/overview.md) | [🟡](../operators/gqa_mqa_attention/backends/triton.md) | · | · | [🟡](../operators/gqa_mqa_attention/backends/ck.md) | · | · | [🟢](../operators/gqa_mqa_attention/backends/hipkittens.md) | [🟢](../operators/gqa_mqa_attention/backends/aiter.md) | [🟡](../operators/gqa_mqa_attention/backends/fa_rocm.md) | · | · | | [sliding_window_attention](../operators/sliding_window_attention/overview.md) | [🟡](../operators/sliding_window_attention/backends/triton.md) | · | · | [🟢](../operators/sliding_window_attention/backends/ck.md) | · | · | · | [🟡](../operators/sliding_window_attention/backends/aiter.md) | [🟡](../operators/sliding_window_attention/backends/fa_rocm.md) | · | · | | [chunked_prefill](../operators/chunked_prefill/overview.md) | [🟢](../operators/chunked_prefill/backends/triton.md) | · | · | · | · | · | · | [🟢](../operators/chunked_prefill/backends/aiter.md) | · | [🟢](../operators/chunked_prefill/backends/sglang_kernels.md) | [🟢](../operators/chunked_prefill/backends/vllm_kernels.md) | diff --git a/perf_knowledge/index/sota_registry.yaml b/perf_knowledge/index/sota_registry.yaml index a2946a923..f706fa9aa 100644 --- a/perf_knowledge/index/sota_registry.yaml +++ b/perf_knowledge/index/sota_registry.yaml @@ -714,6 +714,16 @@ entries: - https://github.com/Dao-AILab/flash-attention - https://vllm.ai/blog/2026-02-27-rocm-attention-backend - ROCm/aiter@a6bb499375849eec45d68c5ccaebc8865fd422c0:aiter/mla.py + - operator: mla_attention + backend: flydsl + status: sota + gens: [gfx942] + dtypes: [bf16] + regimes: [training] + card: ../operators/mla_attention/backends/flydsl.md + sources: + - Attention-Kernels/geak_trans_py2flydsl/fmha_backward/FMHA_BWD_FlyDSL_Skills.md + - https://arxiv.org/abs/2511.08083 - operator: mla_attention backend: hip status: competitive diff --git a/perf_knowledge/index/sources_index.md b/perf_knowledge/index/sources_index.md index 19d3d3ba8..0e397ad0a 100644 --- a/perf_knowledge/index/sources_index.md +++ b/perf_knowledge/index/sources_index.md @@ -4,7 +4,7 @@ kind: reference updated: 2026-06-09 --- -# Sources index — 503 unique URLs across 649 docs +# Sources index — 503 unique URLs across 656 docs Auto-generated union of every `## Sources` / inline URL (run `index/_gen_sources.py`). Each doc keeps its own inline `## Sources`. diff --git a/perf_knowledge/languages/flydsl/authoring_attention_levers.md b/perf_knowledge/languages/flydsl/authoring_attention_levers.md new file mode 100644 index 000000000..27e232dab --- /dev/null +++ b/perf_knowledge/languages/flydsl/authoring_attention_levers.md @@ -0,0 +1,172 @@ +--- +title: "FlyDSL — attention authoring levers (fused multi-GEMM kernels)" +kind: language +gens: [gfx942] +dtypes: [bf16, fp16] +regimes: [prefill, training, both] +updated: 2026-08-12 +sources: + - AMD-AGI/GEAK@c0a1f937:src/minisweagent/skills/flydsl/docs/flydsl_optimization.md + - https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-mi300-cdna3-instruction-set-architecture.pdf + - https://arxiv.org/abs/2511.08083 +--- + +> **Reference (how-to), not a verdict.** Attention-specific follow-on to +> [`authoring_optimization.md`](authoring_optimization.md), and the sibling of +> [`authoring_gemm_levers.md`](authoring_gemm_levers.md). For *which backend to use* on a given +> attention operator (library vs author) see the operator cards under +> [`../../operators/`](../../operators/) — authoring is usually the wrong answer when a tuned +> asm/CK path already dispatches for your shape. + +# FlyDSL attention authoring levers + +## Overview +Use this document when the kernel is an attention **fused multi-GEMM** — several matmuls chained +through register- or LDS-resident intermediates, with a softmax (or its derivative) in between — rather +than a single GEMM. The canonical cases are flash-attention forward (2 GEMMs) and attention backward +(up to 5: score, dP, dV, dK, dQ). + +These kernels behave differently from GEMM in one decisive way: **the per-wave resident operand set, +not the tile size, is what binds.** A GEMM's registers are dominated by one accumulator you can shrink +by tiling; a fused attention kernel carries several long-lived operand sets *plus* accumulators across +the whole loop, so the 256 arch-VGPR cap is reached by the kernel's *structure* and most classical +"reduce instructions" optimizations lose to spill. Read +[`../../optimization/occupancy_and_registers.md`](../../optimization/occupancy_and_registers.md) +before pulling any lever here. + +## Three decisions to make before writing code +Each of these is a rewrite if changed later, and none of them is a tuning knob. + +### 1. The fusion boundary +Splitting a GEMM out of the chain forces its inputs to be **recomputed**, because the intermediates +(scores, probabilities) are too large to spill to HBM. Cost the split symbolically before choosing: +for a backward pass, per score element, the fused chain costs +`2·(HD_QK + HD_V + HD_V + HD_QK + HD_QK)` FLOPs while splitting the dQ GEMM into its own kernel adds a +second score + dP recomputation. Evaluate that ratio for *your* head dims — at symmetric +`HD_QK = HD_V = 128` it is ~+33%, and it rises as the head dims diverge. A split that adds >20% of +total work rarely wins back the register pressure it relieves. + +### 2. The MFMA fragment orientation +For lane `l`, with `grp = l/16`, `col = l%16`, `g4 = 4·grp`, the two useful conventions differ in which +index lands on `col`: + +``` +standard: A[m=g4+v, k=col] B[k=col, n=g4+v] C[m=col, n=g4+v] +operand-swapped: A[m=col, k=g4+v] B[k=g4+v, n=col] C[m=g4+v, n=col] +``` + +Pick the one that makes an intermediate come out of its producing MFMA **already in the layout its +consumer needs**. In a backward pass the operand-swapped form yields `S[qrow, kpos]`, which is exactly +the A-operand layout that both the dV and dK GEMMs want (both contract over q rows), so the +probabilities and their gradient **never leave registers and never transit LDS**. The transpose you +avoid this way is free; the one you retrofit costs LDS traffic and a barrier. + +### 3. Which axis is outer (and therefore what goes on atomics) +In a fused backward exactly one output contracts over the axis you parallelize, so exactly one output +must be accumulated across blocks. Choose the orientation that puts the **smaller** tensor on atomics: +a KV-outer loop makes KV rows exclusive to a block (dK/dV stored directly, no atomics, no workspace) +and accumulates dQ; a q-outer loop does the reverse. Compare the workspace volumes explicitly — +`fp32` accumulator bytes are `tokens · nheads · head_dim · 4`, so with typical `nheads ≫ 1` on the q +side this is usually a several-fold difference, not a marginal one. Decide once, then stop revisiting. + +## Geometry derivation +Derive every count from the tile shape rather than hard-coding it, and **assert the divisibility +constraints at build time** so an invalid configuration fails loudly instead of computing plausible +garbage. With `WARP=64`, MFMA `16×16×16`, `FRAG=4`, `VEC=8` (one 128-bit bf16 load): + +| symbol | derivation | meaning | +|---|---|---| +| `KPW` | `BN / NW` | KV rows per wave — the split axis | +| `NT_Q` | `BM / 16` | q-row tiles; also the dV/dK contraction steps | +| `KS_QK`, `KS_V` | `HD_QK / 16`, `HD_V / 16` | score and dP contraction steps | +| `NT_V`, `NT_HD` | `HD_V / 16`, `HD_QK / 16` | dV and dK output tiles | +| `TPW` | `NT_Q · NT_HD / NW` | per-wave share of the re-partitioned output | +| staging counts | `ceil((BM · HD_x / VEC) / BLOCK_THREADS)` | per-thread global loads per tile | + +Required: `BN % (NW·16) == 0`, and `NT_Q·NT_HD` divisible by `NW` with a wave's share inside one q-row +tile. A hard-coded staging count silently drops part of a tile at any other geometry — this is the +most common cause of a kernel that is correct at one `BM` and wrong at another. + +## The levers, in the order they pay +1. **Wave count — decide it against the register cap, not by occupancy intuition.** More waves means a + smaller per-wave slice of the split axis, hence a *smaller* resident operand set; fewer waves gives + each wave more registers but doubles what it must hold. Since the arch-VGPR cap is 256 **regardless + of occupancy** (only the accumulator file grows as you drop to 1 wave/SIMD, and an MFMA cannot take + an AGPR as an A or B operand), the fewer-waves direction often *loses* on these kernels even though + it wins on GEMM. Measure both, and judge by the ISA — see + [`../../optimization/mfma_scheduling.md`](../../optimization/mfma_scheduling.md). +2. **Read an operand in its awkward orientation instead of materialising the transpose.** Q and dO are + each needed in two orientations (the score GEMM contracts over the head dim, dV/dK over q rows). + Assembling a fragment from four 2-byte LDS reads down a *column* is the same operand and is + conflict-free by construction — much cheaper than a second global read plus a second LDS copy. See + [`../../optimization/lds_and_bank_conflicts.md`](../../optimization/lds_and_bank_conflicts.md). +3. **Keep the hot loop one scheduling region.** `s_barrier` ends an instruction scheduling region, so a + global load and the MFMA chain you want it hidden behind must sit in the *same* region with no + barrier between them. Double-buffer the staged tiles and issue tile `i+1`'s loads inside tile `i`'s + MFMA chain. Keep the staging **branch-free** — clamp the tail index rather than guarding it, because + a guard issues the same instructions under exec mask but ends the basic block and splits the region. + The same reasoning explains why removing bounds-check VALU behind a block-uniform `scf.If` can + measure *slower*. See [`../../optimization/memory_pipelining.md`](../../optimization/memory_pipelining.md). +4. **Re-partition the waves for the output that contracts over the split axis** instead of reducing + partials. Publish the intermediate to LDS once, then divide the output tiles among the waves so each + wave contracts over the whole axis alone. This converts a cross-wave reduction into a single barrier + and collapses that accumulator's register cost by roughly `NW`. +5. **Order the grid axes by cost, not by convention.** Under causal masking a block's work falls + monotonically with its KV-tile index, so putting that index on the *slowest* axis dispatches the + expensive blocks first (longest-job-first). This is usually a few SALU and a permuted launch grid, + and it is frequently the largest single win available. See + [`../../optimization/xcd_l2_locality.md`](../../optimization/xcd_l2_locality.md). +6. **Emit the output dtype the caller wants.** Leaving gradients in fp32 for a bf16 pipeline buys a + separate cast kernel per tensor and doubles the store traffic. Casting in the epilogue is + register-neutral. + +## Verification gates +Cheap, ordered, and each catching a distinct failure class: + +1. **Parity** against an fp32 reference on *every* output (`cos ≥ 0.999` with a relative-error bound is + a reasonable bf16 gate). One output passing while another fails localizes the bug: a good dV with a + bad dQ points at the transpose or the re-partition indices, not at the masking. +2. **Echo the resolved geometry** (LDS bytes, tile shape, wave count, knob states) and treat it as part + of the gate — a build that silently fell back to different parameters invalidates every timing. +3. **Register allocation**: arch VGPR, AGPR, **and `v_accvgpr` move count**. See the register-allocation + gate in [`authoring_optimization.md`](authoring_optimization.md). +4. **MFMA count identity.** Compute the expected MFMA issues from the decomposition and compare against + the profiled counter — it should match exactly. One counter pins the whole work decomposition and + catches entire classes of masking and range bugs. +5. **Keep the ablation paths alive** behind knobs, permanently. Structural claims are only re-checkable + if you can build both sides in one interleaved process (see + [`../../profiling/benchmarking_methodology.md`](../../profiling/benchmarking_methodology.md) + §single-kernel variant sweeps), and it lets a regression be bisected against a mechanism rather than + a commit. + +## Anti-patterns +| attempt | why it loses | +|---|---| +| Growing the KV/N tile to cut atomic traffic | the resident operand set scales with it and overruns the arch-VGPR cap | +| Hoisting invariant address math out of the hot loop | lengthens live ranges; the spill costs more than the arithmetic saved | +| Merging phases to remove a barrier | same cause — the merged region's live set exceeds the cap | +| Fanning atomics across more slots to spread contention | device-scope atomics are L2-uncached by construction, so this only multiplies DRAM traffic | +| Splitting the longest blocks to raise a low occupancy figure | if the block count is pinned by the register cap, total time is unchanged | +| Trusting `scratch`/`private_seg_size` as the spill measure | AGPR shuttling never appears there | +| Materialising a second orientation of a staged tile | see lever 2 | + +**Before any change that keeps more values live, predict its effect on the live set and read the ISA.** +It is a two-minute check, and on these kernels it is the difference between a lever and a regression. + +## Worked instance +An end-to-end application of every lever above, with per-step measured deltas, counter expectations and +a dead-end ledger, is recorded for a fused five-GEMM MLA backward with asymmetric head dims on gfx942: +[`../../operators/mla_attention/backends/flydsl.md`](../../operators/mla_attention/backends/flydsl.md) +and the skill +[`../../expert_skills/skills/flydsl_fused_attention_backward`](../../expert_skills/skills/flydsl_fused_attention_backward/skill.md). +Treat its constants as one instance of the derivations above, not as defaults to copy. + +## Sources +- Generic authoring workflow this specializes: [`authoring_optimization.md`](authoring_optimization.md) + (origin `AMD-AGI/GEAK@c0a1f937:src/minisweagent/skills/flydsl/docs/flydsl_optimization.md`). +- MFMA fragment layouts, AGPR/arch-VGPR split, `s_barrier` scheduling semantics: AMD CDNA3 ISA reference. +- Wave-scheduling patterns for AMD attention kernels (8-wave ping-pong / 4-wave interleave, and the + register-allocation reason wave specialization loses on CDNA): HipKittens, arXiv 2511.08083 — + see [`../hipkittens/primitives.md`](../hipkittens/primitives.md). +- Levers 1–6 and the anti-pattern table are each measured on-box; the instance, deltas and mechanisms + are cited in the operator card above. diff --git a/perf_knowledge/languages/flydsl/authoring_optimization.md b/perf_knowledge/languages/flydsl/authoring_optimization.md index 819a8ff53..bcca65d6d 100644 --- a/perf_knowledge/languages/flydsl/authoring_optimization.md +++ b/perf_knowledge/languages/flydsl/authoring_optimization.md @@ -216,6 +216,33 @@ Always verify after each patch — violations often cause silent corruption: 6. If the generated form did not change, assume the optimization did not land yet, even if the Python source looks right 7. If speedup is marginal or absolute time regresses, move to the next structural strategy rather than re-tuning the same approach +### The register-allocation gate (run this at every step, not at the end) + +On register-heavy kernels this is the dominant failure mode and it is **invisible from timings**. Make a +one-config ISA probe part of the harness from the start, then check three numbers together: + +```bash +FLYDSL_DUMP_IR=1 FLYDSL_DUMP_DIR=/tmp/isa python isa_probe.py "" +grep -E 'vgpr_count|agpr_count|private_seg_size' /tmp/isa//21_final_isa.s +grep -c v_accvgpr /tmp/isa//21_final_isa.s # the discriminating number +``` + +- **Count `v_accvgpr` moves, not scratch.** When arch VGPRs run out the compiler first parks values in the + accumulator file, which costs issue slots but never shows up as `private_seg_size`/`scratch_`. A build + can report *lower* scratch than its rival and be 15% slower. See + [[optimization/occupancy_and_registers]]. +- A small non-zero `private_seg_size` (e.g. 16 B) is often a fixed prologue slot — confirm there is no + `scratch_`/`buffer_store` targeting it before calling it spill. +- **A step that cuts instruction count but spills is a regression.** Before any change that lengthens a + live range, predict its effect on the live set and check the ISA — measured examples where it lost: + hoisting invariant address math out of the hot loop (**−8.6%**, `private_seg_size` 16 → 48 B), and + merging a second compute phase into the main region to remove a barrier (**−36%**, 16 → 60 B). +- Related: `s_barrier` **ends an instruction scheduling region**, so a global load and the MFMA chain you + want it hidden behind must sit in the same region with no barrier between them. Keep staging + branch-free (clamp indices instead of guarding) — a guard issues the same instructions under exec mask + but ends the basic block and splits the region, which is why removing bounds-check VALU by splitting on + a block-uniform condition can measure *slower*. See [[optimization/memory_pipelining]]. + ## Key FlyDSL APIs - Device kernel: `@flyc.kernel` | Host launcher: `@flyc.jit` @@ -232,4 +259,8 @@ Always verify after each patch — violations often cause silent corruption: - AMD CDNA3 ISA (LDS banks, `s_waitcnt`/`lgkmcnt`, MFMA variants): https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-mi300-cdna3-instruction-set-architecture.pdf - AMD CDNA4 ISA (gfx950 LDS/bank changes, scaled MFMA): https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-cdna4-instruction-set-architecture.pdf - Matrix Cores on CDNA (MFMA programming): https://rocm.blogs.amd.com/software-tools-optimization/matrix-cores-cdna/README.html +- Register-allocation gate (`v_accvgpr`-vs-scratch, the two measured spill regressions, barrier-ends-region): + first-party on-box gfx942 MI300X, ROCm 7.1.0, mirrored in + [[operators/mla_attention/backends/flydsl]] and + [`expert_skills/skills/flydsl_fused_attention_backward`](../../expert_skills/skills/flydsl_fused_attention_backward/skill.md). - Cross-refs: [`authoring_gemm_levers.md`](authoring_gemm_levers.md) · [`debugging.md`](debugging.md) · [`knobs.md`](knobs.md) diff --git a/perf_knowledge/languages/flydsl/debugging.md b/perf_knowledge/languages/flydsl/debugging.md index bb6da85e5..ad4de1453 100644 --- a/perf_knowledge/languages/flydsl/debugging.md +++ b/perf_knowledge/languages/flydsl/debugging.md @@ -17,8 +17,10 @@ sources: This reference covers correctness, stability, and hang triage on runnable FlyDSL kernels. **Scope**: execution debugging, not performance triage. For kernel-speed work see -[`authoring_optimization.md`](authoring_optimization.md) (and [`authoring_gemm_levers.md`](authoring_gemm_levers.md) -for GEMM). Here, "tracing" refers to FlyDSL frontend tracing, not ATT collection. +[`authoring_optimization.md`](authoring_optimization.md) — in particular its **register-allocation gate**, +which is the dominant silent failure mode on register-heavy kernels and is invisible from timings — and +[`authoring_gemm_levers.md`](authoring_gemm_levers.md) for GEMM. Here, "tracing" refers to FlyDSL frontend +tracing, not ATT collection. ## Step 0: Decide whether cache is involved diff --git a/perf_knowledge/languages/flydsl/overview.md b/perf_knowledge/languages/flydsl/overview.md index 9d36b2cde..fd4b5aa41 100644 --- a/perf_knowledge/languages/flydsl/overview.md +++ b/perf_knowledge/languages/flydsl/overview.md @@ -82,6 +82,7 @@ gfx (`..._gfx942`). Key arch behavior baked into aiter's wrappers: - [authoring_tile_programming.md](authoring_tile_programming.md) — write a first correct kernel (CuTe-style tile model, the 4 patterns, MFMA reference). - [authoring_optimization.md](authoring_optimization.md) — structure-first optimization workflow (fusion → LDS → MFMA-loop → tuning). - [authoring_gemm_levers.md](authoring_gemm_levers.md) — GEMM-specific levers (tiling / LDS staging / swizzle / epilogue). +- [authoring_attention_levers.md](authoring_attention_levers.md) — fused multi-GEMM attention levers (fusion boundary, MFMA fragment orientation, which output goes on atomics, wave count vs the register cap). - [debugging.md](debugging.md) — correctness/stability/hang triage (NaN / zeros / mismatch / compile / hang). ## Sources diff --git a/perf_knowledge/languages/hipkittens/primitives.md b/perf_knowledge/languages/hipkittens/primitives.md index 8f2c7b176..b0d17ebfd 100644 --- a/perf_knowledge/languages/hipkittens/primitives.md +++ b/perf_knowledge/languages/hipkittens/primitives.md @@ -65,6 +65,13 @@ arithmetic intensity (only ~80% of peak BF16 GEMM). The two performant AMD patte buys ~22% on backward at ~3× the code.) **8-wave ping-pong is the default** — it already reaches SOTA for GEMM and attention-forward on MI355X. +> **The backward row does not generalise.** It is symmetric `hdim=128` non-causal on MI355X. A gfx942 +> counter-measurement on a register-heavier backward (fused five-GEMM causal, asymmetric head dims) has every 4-wave +> build **15% slower** than 8-wave, all of them shuttling AGPRs — halving the waves doubles the per-wave +> resident operand set against a 256 arch-VGPR cap that does not move with occupancy. HK's gain also relies +> on hand-staggered issue, which a compiler DSL can only approximate with `sched_group_barrier` hints. +> See [[optimization/mfma_scheduling]] and [[operators/mla_attention/backends/flydsl]]. + ## XCD-aware grid swizzle (chiplet scheduling) MI355X = 256 CUs across **8 XCDs** (32 CUs each); each XCD has a private 4 MB L2, all share an LLC before HBM (miss ~300 ns L2, ~500 ns LLC). Blocks are assigned round-robin to XCDs. HK's Algorithm 1: diff --git a/perf_knowledge/operators/gqa_mqa_attention/backends/hipkittens.md b/perf_knowledge/operators/gqa_mqa_attention/backends/hipkittens.md index 3a81ab7c0..9a7288169 100644 --- a/perf_knowledge/operators/gqa_mqa_attention/backends/hipkittens.md +++ b/perf_knowledge/operators/gqa_mqa_attention/backends/hipkittens.md @@ -41,8 +41,14 @@ backward). See [[languages/hipkittens]] and arXiv 2511.08083 Table 1. | HK GQA/MQA **forward**, d=64 | arXiv 2511.08083v1 | gfx950; bf16/fp16 | beats baselines (fwd in the 1.0–2.1× vs AITER class); **d=64 a strength** | small-head GQA fwd | ## Config space / knobs (backend-specific) -- **Schedule**: 4-wave interleave often wins the backward (full register budget eases dQ/dK/dV pressure); - 8-wave ping-pong for compact forward. +- **Schedule**: 4-wave interleave wins **this** backward (symmetric `hdim`, non-causal, MI355X — the full + per-wave register budget eases dQ/dK/dV pressure); 8-wave ping-pong for compact forward. **Do not carry + the 4-wave conclusion to a register-heavier backward.** On a fused five-GEMM *causal* backward with + asymmetric head dims on **gfx942**, every 4-wave build measured **15% slower** than 8-wave and all of them + shuttled AGPRs, because halving the waves doubles the per-wave K/V/Kᵀ set against a 256 arch-VGPR cap that + does not move — [[optimization/mfma_scheduling]] has the accumulator-bound vs operand-bound split that + decides which way this goes, and [[operators/mla_attention/backends/flydsl]] the instance. Pick the wave + count by measurement, and read the ISA (`v_accvgpr` move count), not the clock. - **Pinned register tiles**: the decisive lever for backward — pin registers, feed AGPRs to MFMA, avoid `v_accvgpr_read`. Sharp tool: bypasses the allocator; validate parity. - **Head dim** d=64 vs 128 (d=64 underserved by tuned asm → HK win); **GQA group ratio** (KV-head sharing). diff --git a/perf_knowledge/operators/mla_attention/backends/flydsl.md b/perf_knowledge/operators/mla_attention/backends/flydsl.md new file mode 100644 index 000000000..30453e47e --- /dev/null +++ b/perf_knowledge/operators/mla_attention/backends/flydsl.md @@ -0,0 +1,158 @@ +--- +title: mla_attention backward on FlyDSL — SOTA card +kind: sota_card +operator: mla_attention +backend: flydsl +gens: [gfx942] +dtypes: [bf16] +regimes: [training] +status: sota +updated: 2026-08-12 +sources: + - https://arxiv.org/abs/2511.08083 +--- + +# mla_attention (backward) × FlyDSL + +> **This card is the measured instance.** The dimension-agnostic authoring levers it applies — fusion +> boundary, MFMA fragment orientation, which output goes on atomics, wave count vs the register cap, +> geometry derivation, verification gates — are in +> [`languages/flydsl/authoring_attention_levers.md`](../../../languages/flydsl/authoring_attention_levers.md). +> Read that for the *rules*; read this for the numbers that back them. The constants here are one +> resolution of those derivations, not defaults to copy. + +## TL;DR +For the **unabsorbed MLA training backward** — `qk_head_dim=192`, `v_head_dim=128`, bf16, causal, varlen — +a hand-authored FlyDSL three-kernel chain is the **measured fastest option on gfx942**: **715–745 µs** +(~230 TFLOP/s of useful work) against **984 µs** for aiter's ASM backward with V padded 128→192, at +`cos ≥ 0.999999` on all three gradients. The structural point is that **no aiter configuration on gfx942 +gets both the forward and the backward fast**: the ASM bwd dispatch is gated on `hdim_q == hdim_v` (so V +must be padded to 192) while the ASM *forward* requires `hdim_v == 128`. End to end that makes this +**1.49×** faster than the padded/ASM-bwd config and **1.79×** faster than the unpadded/CK-tile-bwd config. +This is a first-party on-box measurement, not a vendor number — and it is the source of GEAK's gfx942 +counter-measurement to HipKittens §3.3.2 on wave counts (see +[[optimization/mfma_scheduling]]). + +## SOTA implementation +Three device-side steps per backward call: zero the fp32 dQ accumulator, a preprocess kernel computing +`D = rowsum(dO ⊙ O)` into a **head-major** `[nheads, seqlen]` fp32 tensor, then one main kernel running +**all five backward GEMMs fused**. Fusing is not an optimisation to retrofit — splitting dQ into its own +kernel forces S and dP to be recomputed there, raising work from 1664 to 2304 FLOPs per score element +(+38%), which loses to the ASM baseline on its own. + +Four structural decisions carry the performance — the general form of each is a lever in +[[languages/flydsl/authoring_attention_levers]]; here is how they resolved at these head dims. Two are +layout choices that are painful to retrofit: + +- **Operand-swapped MFMA fragments** (`v_mfma_f32_16x16x16_bf16_1k`) make the score land as + `S[qrow, kpos]`, which is already the A-operand layout both dV and dK want. **P and dS never leave + registers and never transit LDS** — free, if chosen first. +- **KV-outer schedule**: KV rows are exclusive to a block, so dK/dV are stored directly, no atomics. Only + dQ accumulates, via global atomics into a 606 MB fp32 workspace. The q-outer alternative puts dK/dV on + atomics at 4.04 GB. +- **`kv_tile` on the slowest grid axis** (worth **20%**) — longest-job-first, because a block's cost falls + monotonically with `kv_tile` under causal masking. See [[optimization/xcd_l2_locality]]. +- **Read the second operand orientation from LDS instead of materialising it** (worth **19%**) — see + [[optimization/lds_and_bank_conflicts]]. + +| impl | source | gens/dtypes | measured perf | when best | +|---|---|---|---|---| +| FlyDSL fused 5-GEMM bwd chain (native 192/128) | 13 varlen causal seqs, 32768 tokens, 2 heads | gfx942 (MI300X/MI325X, **SPX/NPS1**); bf16 | **715–745 µs, ~230 TFLOP/s useful** — vs padded-ASM 984 µs (−27.3% wall, −15.4% executed work); e2e **1008.6 µs/iter** vs 1497.9 (B) / 1806.9 (A) | asymmetric MLA head dims on gfx942 training | +| FlyDSL preprocess (`D = Σ dO·O`) | same | gfx942; bf16→fp32 | **3.9 µs** vs ASM `odo` 22.2 µs (**5.7×**) | always — smallest, easiest win | +| (ref) aiter ASM bwd, V padded to 192 | aiter v3 gfx942 dispatch | gfx942; bf16 | 984 µs, 196.9 TFLOP/s executed / 170.7 useful | symmetric head dims, where no padding is needed | +| (ref) CK-tile bwd, V unpadded | aiter fallback | gfx942; bf16 | 1560.8 µs | keeps the fast ASM forward, loses the backward | + +## Config space / knobs +Defaults are the measured optimum; the last three exist only to reproduce the ablation ladder and should +be kept anyway, so the comparisons stay checkable in one interleaved process on new hardware. + +| param | range / source | effect | default | +|---|---|---|---| +| `block_m` | 16 / 32 / 64 | q rows per tile; 64 is over LDS budget once double-buffered, 16 is rejected at 8 waves (12 dQ tiles don't divide by 8) | **32** | +| `block_n` | 64 / 128 / 192 / 256 | KV rows per block; ≥192 is 3–6× slower (K/V/Kᵀ operand set overruns the 256 arch-VGPR cap) | **128** | +| `num_waves` | 4 / 8 | 8 → 512 threads, 2 waves/SIMD. **4 is 15% slower on this kernel** — see Pitfalls | **8** | +| `dq_groups` | 1 / 2 / 4 / 8 | dQ atomic fan-out slots; monotonically worse (1187/1216/1259/1339 µs) — atomics are L2-uncached regardless | **1** | +| `kv_major` | bool | `kv_tile` on the slow axis (longest-job-first) | **True** (−20%) | +| `pipeline` | bool | double-buffered staging, 2 barriers instead of 3 | **True** (−5.4%) | +| `transposed_tiles` | bool | materialise the second orientation instead of reading it | **False** (True costs +19%) | +| `schedule` / `sched_steps` | `sched_group_barrier` hints | neutral once the region is branch-free and single | off | + +Resolved geometry at the default: `LDS 50948 B BLOCK_M=32 BLOCK_N=128 waves=8 dq_groups=1 pipeline=True +kv_major=True`. Print it and treat it as part of the gate — anything else means no timing is comparable. + +## Numerics / parity +fp32 accumulate throughout; gradients come back fp32 and the caller casts to bf16. `lse` is +`[nheads, total_q]` fp32 in natural log with the softmax scale already folded in, exactly as a +Flash-Attention forward emits it, because the kernel computes `p = exp2(s·scale·log2e − lse·log2e)`. +Causal convention is **bottom-right aligned** (`delta = seqlen_k − seqlen_q`, keep `j <= i + delta`); +top-left alignment is a different problem that silently produces high-cosine garbage. Masking is +branch-free and a zeroed P also zeroes dS, so no separate dS mask is needed. + +Measured `cos = 0.999999` on dq/dk/dv with `max_err/scale ≈ 1.6–2.0e-03`; gate at `cos_tol=0.999`, +`rel_tol=3e-2`. Accuracy is also *better* than both aiter configs (dq max abs err 1.193e-02 vs 2.546e-02 +unpadded-CK and 3.627e-02 padded-ASM). + +## Counter expectations (rocprofiler-compute) +These confirm the design is behaving as specified rather than accidentally working: + +| claim | counter | +|---|---| +| 8 waves/CU, 2 waves/SIMD, 1 block/CU | 128 arch VGPR + 128 AGPR, 112 SGPR, LDS 51200 B | +| no spill despite a full register file | scratch 16 B, **0** spill/stack instructions | +| the column-wise transpose read is conflict-free | 0.03 bank conflicts/access, 0 address conflicts | +| branch-free staging keeps the exec mask full | 63.97 of 64 VALU active threads | +| reading the transpose beats materialising it | LDS instructions 12.7M → 42.4M (3.3×), but vL1D read requests 322.6M → 66.0M (−80%) and VMEM instructions −54% | +| the decomposition is right | MFMA count exactly 21,582,080 (104/wave-iteration × 207,520), 100% BF16 | + +**Nothing is saturated**: MFMA 16.1% of peak, VALU 19.7% of issue slots, LDS 22.8%, vL1D 26.1%, HBM 33%. +Per resident wave-cycle: 36.1% active, **40.4% waiting on a dependency**. The kernel is latency-bound at +2 waves/SIMD, and 2 waves/SIMD is fixed by the register cap — do not read the 51% occupancy figure as +free headroom (splitting the longest blocks to collect it was implemented and produced nothing). + +## Pitfalls & anti-patterns +- **`num_waves=4` is 15% slower, and scratch will lie to you about why.** The best 4-wave build reports + `private_seg_size 0 / agpr_count 123 / vgpr_count 379` with **75 `v_accvgpr` moves** — *lower* scratch + than the 8-wave build, which has 256 VGPR / 0 AGPR / 0 moves. The spill relocated into the accumulator + file. **Count `v_accvgpr` moves, not scratch** ([[optimization/occupancy_and_registers]]). +- **SPX/NPS1 is part of the specification.** The 20% grid-mapping win is built on static XCD round-robin; + CPX changes both dispatch and atomic behaviour. If the `kv_major` ablation ratio comes out near 1.0 + instead of 0.79, check the partition mode before suspecting the kernel. +- **Optimisations that keep more values live lose to the register cap**: hoisting invariant staging + address math out of the q-loop is **−8.6%** (triples spill), and deferring the dQ phase to fold to a + single barrier is **−36%**. Predict the effect on the live set and check the ISA first — a two-minute + check that pre-empts both. +- **Splitting the region on a block-uniform condition measures slower** even when it removes ~200 VALU of + bounds checks, because the `scf.If` boundary stops the scheduler hoisting global loads into the previous + iteration's MFMA chain. Overlap is worth more than instruction count here. +- Do not design around L2 residency for the dQ atomics — every one is uncached by construction. +- Budget ~49 µs of `bfloat16_copy_kernel` as this path's own overhead (it leaves gradients in fp32 where + the ASM path folds the dQ cast into its convert pass). Emitting bf16 from the epilogue is + register-neutral and the most accessible remaining win. + +## How to verify +```bash +rocm-smi --showproductname # confirm gfx942 + expected CU count; check SPX/NPS1 +FLYDSL_DUMP_IR=1 FLYDSL_DUMP_DIR=/tmp/isa python isa_probe.py "num_waves=8" +grep -E 'vgpr_count|agpr_count|private_seg_size' /tmp/isa/fmha_bwd_main/21_final_isa.s +grep -c v_accvgpr /tmp/isa/fmha_bwd_main/21_final_isa.s # must be 0 +# then: build BOTH sides in ONE interleaved sweep (min over replays), never across processes +``` + +## Alternatives / cross-links +[[operators/mla_attention/backends/aiter]] (sota decode/prefill; the ASM bwd baseline here) · +[[operators/mla_attention/backends/ck]] (CK-tile bwd fallback) · +[[operators/gqa_mqa_attention/backends/hipkittens]] (the other attention-backward card; note its 4-wave +recommendation does **not** transfer to this shape) · +[`expert_skills/skills/flydsl_fused_attention_backward`](../../../expert_skills/skills/flydsl_fused_attention_backward/skill.md) +(the regulated procedure) · +[`languages/flydsl/authoring_attention_levers.md`](../../../languages/flydsl/authoring_attention_levers.md) +(the general levers) · +[[languages/flydsl/authoring_optimization]] · [[languages/flydsl/debugging]]. + +## Sources +- gfx942 MI300X/MI325X, SPX/NPS1, 13 varlen causal sequences of + 472–3638 tokens (32768 total), 2 heads, bf16, softmax scale 0.08838834764831843, bottom-right causal. +- aiter gfx942 v3 bwd `hdim_q == hdim_v` dispatch gate and ASM-forward `hdim_v == 128` requirement: same + document §8.3–§8.4 (measured against an aiter source checkout). +- HipKittens §3.3.2 wave-schedule patterns that this kernel's 4-wave result contradicts on gfx942: + https://arxiv.org/abs/2511.08083 diff --git a/perf_knowledge/operators/mla_attention/overview.md b/perf_knowledge/operators/mla_attention/overview.md index 6cd5fb8d4..ea7fdb21d 100644 --- a/perf_knowledge/operators/mla_attention/overview.md +++ b/perf_knowledge/operators/mla_attention/overview.md @@ -4,13 +4,14 @@ kind: operator_overview operator: mla_attention gens: [gfx942, gfx950] dtypes: [bf16, fp16, fp8_e4m3_fnuz, fp8_e4m3] -regimes: [prefill, decode] -updated: 2026-06-08 +regimes: [prefill, decode, training] +updated: 2026-08-12 sources: - ROCm/aiter@a6bb499375849eec45d68c5ccaebc8865fd422c0:aiter/mla.py - https://rocm.blogs.amd.com/software-tools-optimization/aiter-mla/README.html - https://vllm.ai/blog/2026-02-27-rocm-attention-backend - https://arxiv.org/abs/2405.04434 + - Attention-Kernels/geak_trans_py2flydsl/fmha_backward/FMHA_BWD_FlyDSL_Skills.md --- # mla_attention (DeepSeek Multi-head Latent Attention) @@ -44,6 +45,10 @@ backends give **1.2–1.6× faster TPOT** vs Triton MLA. ## Shape regimes - **Decode** (`sq=1`): the headline — bandwidth-bound MQA over the latent; `mla_decode_fwd` + splitKV. - **Prefill** (long sq): `mla_prefill_fwd` (and a persistent variant `mla_prefill_ps_fwd`); GEMM-bound. +- **Training / backward**: the unabsorbed form with **asymmetric head dims** `qk_head_dim=192` + (`qk_nope 128 + qk_rope 64`), `v_head_dim=128`. On gfx942 this asymmetry is a dispatch trap — aiter's + ASM backward is gated on `hdim_q == hdim_v` and its ASM forward on `hdim_v == 128`, so **no aiter + configuration gets both fast**. See [backends/flydsl.md](backends/flydsl.md). - DeepSeek shapes: `num_heads` 16/64/128, `kv_lora_rank=512`, `qk_rope_head_dim=64`, `v_head_dim=128`. ## Where it matters (Amdahl) @@ -56,6 +61,7 @@ broadcasts a real KV head) — MLA attends a *compressed latent*. | backend | status | card | |---|---|---| | aiter | 🟢 sota (asm `mla_decode_fwd`, 17×) | [backends/aiter.md](backends/aiter.md) | +| flydsl | 🟢 sota **backward only**, gfx942 192/128 (715–745 µs vs padded-ASM 984) | [backends/flydsl.md](backends/flydsl.md) | | triton | 🟡 (reference + fallback; `mla_decode.py`) | [backends/triton.md](backends/triton.md) | | ck | 🟡 (CK-Tile MLA; from-source) | [backends/ck.md](backends/ck.md) | | hip | 🟡 (vLLM custom; mostly routes to AITER MLA) | [backends/hip.md](backends/hip.md) | diff --git a/perf_knowledge/optimization/lds_and_bank_conflicts.md b/perf_knowledge/optimization/lds_and_bank_conflicts.md index b2c06c9cc..4ba7aef6a 100644 --- a/perf_knowledge/optimization/lds_and_bank_conflicts.md +++ b/perf_knowledge/optimization/lds_and_bank_conflicts.md @@ -59,6 +59,34 @@ footprint — feasible for many tiles on CDNA3's 64 KB, much easier on CDNA4's 1 Pre-permuting B into the MFMA-native layout (aiter `bpreshuffle`, `[[operators/dense_gemm/tuning.md]]`) moves the shuffle off the hot path so the LDS read pattern is already conflict-free. +### 5. When an operand is needed in two orientations: read the awkward one, don't materialise it +A kernel that contracts the same staged tile over two different axes (classic case: attention backward, +where the score GEMM contracts over the head dim while dV/dK contract over q rows) appears to need **two** +LDS copies of that tile — the convenient orientation plus a transposed one, written by a second global +read. **Prefer reading the awkward orientation out of the single copy.** + +The assumption worth dropping is that an MFMA operand needs its contraction-axis elements to arrive in one +load. It does not: a fragment assembled from four 2-byte `ds_read`s walking a **column** of a +row-major tile is the same operand. Make it conflict-free by construction rather than swizzling — if +consecutive lanes differ only in the MFMA column and the four lane groups sit `4·pad` elements apart, the +64 lanes cover 32 distinct banks and the two lanes sharing a bank also share its dword. **No XOR swizzle +is needed**, because the feared many-way conflict lives on the *scatter* side of a materialised transpose, +which this removes entirely. + +Measured instance (first-party, gfx942 MI300X, fused attention backward with asymmetric head dims; +authoring guidance in +[`languages/flydsl/authoring_attention_levers.md`](../languages/flydsl/authoring_attention_levers.md)): +worth **−19% wall clock**. It trades ~160 +extra `ds_read` per loop iteration against a whole second global read of both tiles and **23 KB of LDS** — +LDS instructions went 12.7M → 42.4M (3.3×) while vL1D read requests fell 322.6M → 66.0M (**−80%**) and +VMEM instructions −54%, with MFMA count identical. + +> **On CDNA3, LDS read instructions are cheap enough that reading an operand in the awkward orientation +> beats materialising it in the convenient one, by a wide margin, as long as the awkward read is +> conflict-free.** Check this before spending LDS on a transposed copy — and note the freed LDS usually +> buys a second *staging* buffer (`§3`), not a second resident block, since the block count is normally +> pinned by registers anyway (`[[optimization/occupancy_and_registers.md]]`). + ## Sizing budget (capacity → tile) LDS bytes per stage ≈ `(BM·BK + BK·BN) · sizeof(dtype) · num_stages · (2 if double-buffer)`. On CDNA3 (64 KB) a bf16 256×64 + 64×128 double-buffered tile is already tight; CDNA4 (160 KB) lets you diff --git a/perf_knowledge/optimization/mfma_scheduling.md b/perf_knowledge/optimization/mfma_scheduling.md index 57a55d8d5..6cea12f86 100644 --- a/perf_knowledge/optimization/mfma_scheduling.md +++ b/perf_knowledge/optimization/mfma_scheduling.md @@ -49,6 +49,30 @@ Practical prior: on gfx942/gfx950, reach for 8-wave ping-pong or 4-wave interlea compute, full register budget, no producer/consumer split). See `[[languages/hipkittens]]`, `[[operators/dense_gemm]]`, `[[operators/scaled_quant_gemm/tuning.md]]`. +### Which wave count: decide it against the register cap, not from the prior +**Rule: pick the wave count by what each wave must keep resident, and judge the choice from the ISA +rather than the clock.** Fewer waves gives each wave more of the register file but a *larger* slice of the +split axis, so whatever operands it holds across the loop scale up — against an **arch-VGPR cap of 256 +that does not move with occupancy**. Only the accumulator file effectively grows as you approach +1 wave/SIMD, and an MFMA cannot take an AGPR as an A or B operand, so a kernel whose *operands* (not just +accumulators) grow with the slice gets no relief from dropping waves. Which way this resolves depends on +one property of the kernel: + +| the per-wave resident set is | fewer waves | typical case | +| --- | --- | --- | +| dominated by accumulators you can retile | often wins — more registers per wave, deeper unroll | GEMM, GEMM-epilogue fusions | +| dominated by long-lived **operands** that scale with the slice | often loses — the operand set outgrows the 256 cap | fused multi-GEMM attention backward | + +The GEMM prior above is unchanged: 4-wave interleave and 8-wave ping-pong both remain the right starting +points, and for dense/scaled GEMM the ranking between them stands as documented. +Note the 4-wave builds report *lower* scratch than the 8-wave build while being slower — the spill moved +into the accumulator file, where `private_seg_size` cannot see it. Count `v_accvgpr` moves +(`[[optimization/occupancy_and_registers.md]]`). HK's gain also comes from **hand-staggered instruction +issue**, available in a compiler DSL only as `sched_group_barrier` hints, which measure neutral at both +wave counts once the loop body is a single branch-free region. Authoring guidance: +[`languages/flydsl/authoring_attention_levers.md`](../languages/flydsl/authoring_attention_levers.md); +instance: `[[operators/mla_attention/backends/flydsl]]`. + ## Concepts (the hardware) - **MFMA instruction**: `D = A·B + C`, one instruction per wave processes a fixed M×N×K block. Common CDNA3 shapes: `mfma_16x16x16` and `mfma_32x32x8` (bf16/fp16); fp8 variants pack 2× K; diff --git a/perf_knowledge/optimization/occupancy_and_registers.md b/perf_knowledge/optimization/occupancy_and_registers.md index 9457ed8d3..b74c6b826 100644 --- a/perf_knowledge/optimization/occupancy_and_registers.md +++ b/perf_knowledge/optimization/occupancy_and_registers.md @@ -76,14 +76,50 @@ report and the profiler: (`[[profiling/]]`, `[[hardware/cdna3_mi300/occupancy.md]]`). - Rule of thumb: prefer **2 waves/EU with no spills** over 3 waves/EU that spill, for GEMM-class kernels. +### Count `v_accvgpr` moves, not scratch — scratch under-reports the spill +Scratch is not the whole spill surface. When arch VGPRs run out, the compiler's *first* move is to park +values in the **accumulator file** and shuttle them with `v_accvgpr_read/write`, which costs issue slots +and serialises against MFMA **but never appears as `private_seg_size`, `scratch_` or `buffer_store`**. +A build can therefore report *less* scratch than its rival and be materially slower. + +Measured instance (gfx942 MI300X, fused attention backward, same kernel, one interleaved process): + +| waves | `private_seg_size` | vgpr | agpr | `v_accvgpr` moves | µs | +| ---: | ---: | ---: | ---: | ---: | ---: | +| 8 | 16 B | 256 | 0 | **0** | **902** | +| 4 | **0 B** | 379 | 123 | 75 | 1055 | + +The 4-wave build has *zero* scratch and is 15% slower. Read the arch-VGPR count, the AGPR count **and** +the `v_accvgpr` move count together: + +```bash +grep -E 'vgpr_count|agpr_count|private_seg_size' .s +grep -c v_accvgpr .s # the number that actually discriminates +``` + +A small non-zero `private_seg_size` (e.g. 16 B) is often a fixed prologue slot rather than spill — confirm +by checking there is no `scratch_`/`buffer_store` targeting it. + +**Corollary for occupancy hunting:** the arch-VGPR cap is **256 per wave regardless of occupancy** — only +the accumulator file effectively grows as you drop to 1 wave/SIMD, and an MFMA cannot take an AGPR as an A +or B operand. So dropping the wave count does *not* relieve pressure on operands that must be arch VGPRs; +in a kernel whose per-wave resident operand set scales with its work slice, halving the waves *doubles* +that set against a fixed cap. `[[optimization/mfma_scheduling.md]]` has the accumulator-bound vs +operand-bound split and which way each resolves. + ## Pitfalls - Treating "more occupancy = faster" as universal. MFMA-bound kernels run great at 1–2 waves/EU. - Forgetting AGPRs count against the 512 budget — a fat accumulator silently caps occupancy. - Setting `waves_per_eu` high without checking the ISA dump for spills. +- **Judging spill by scratch alone** — AGPR shuttling is invisible there (see the cliff section above). +- Reading a low total-wave-cycle "occupancy %" as free headroom. If the block count is pinned by the + register cap, the usual way to collect it (split the work into more, smaller blocks) changes nothing: + in the measured case a 2-way split of the longest blocks landed at 747 vs 743 µs. - Assuming CUDA "blocks/SM" math; CDNA granularity is **16 VGPR**, slots are **8/SIMD**, wave is **64**. ## Verify -- ISA/asm: confirm VGPR/AGPR counts and zero scratch (`amdgpu-arch` dump; triton `TRITON_CACHE`/`AMDGCN`). +- ISA/asm: confirm VGPR/AGPR counts, zero scratch **and zero `v_accvgpr` moves** (`amdgpu-arch` dump; + triton `TRITON_CACHE`/`AMDGCN`; FlyDSL `FLYDSL_DUMP_IR=1` — see `[[languages/flydsl/debugging.md]]`). - Profiler: occupancy and `VALUBusy` from Omniperf; compare across `waves_per_eu` settings. - A/B: sweep `waves_per_eu ∈ {1,2,3,4}` and `num_warps ∈ {4,8}`, keep the lowest latency with no spill. @@ -91,3 +127,7 @@ report and the profiler: - 512 VGPR/EU, 16-granule, worked 170→176→2-waves example, `waves_per_eu` hint: ROCm MI300X workload guide. - 256 architected + 256 AGPR pools, allocation granularity, `v_accvgpr_*`: AMD CDNA3 (MI300) ISA reference. - Register-pressure / occupancy reasoning (CDNA lab notes): AMD GPUOpen register-pressure note. +- `v_accvgpr`-vs-scratch table, the 256-arch-VGPR-cap-regardless-of-occupancy corollary, and the + occupancy-%-is-not-headroom result: first-party on-box gfx942 MI300X, ROCm 7.1.0 — + `Attention-Kernels/geak_trans_py2flydsl/fmha_backward/FMHA_BWD_FlyDSL_Skills.md` §9–§12, mirrored in + `[[operators/mla_attention/backends/flydsl]]`. diff --git a/perf_knowledge/optimization/xcd_l2_locality.md b/perf_knowledge/optimization/xcd_l2_locality.md index 92e52d636..29b009ba6 100644 --- a/perf_knowledge/optimization/xcd_l2_locality.md +++ b/perf_knowledge/optimization/xcd_l2_locality.md @@ -53,12 +53,46 @@ For a tiled GEMM with `T` tiles on 8 XCDs, instead of `xcd = pid % 8` (scatters `group = pid / tiles_per_xcd; xcd = group; local = pid % tiles_per_xcd` so a contiguous run of data-sharing tiles stays on one XCD's L2. Tune `tiles_per_xcd` to the L2 working-set size. +## Dispatch order is also a *makespan* lever, not only an L2 lever +The round-robin assignment is **static**: under SPX, workgroups reach the CUs in **x-fastest linear order** +and block `i` is handed to XCD `i % 8` with no migration — a block cannot move to whichever XCD happens to +be idle. Three consequences that are easy to miss if you only think about L2: + +- **Put the expensive axis on the SLOWEST grid axis (longest-job-first).** When per-block cost varies + systematically along one grid axis — the standard case is a causal attention kernel with a KV-outer + schedule, where a block's q-tile count falls monotonically with its KV tile index — having that axis + *fast* means the whole cost range of one sequence is dispatched before the next starts, so a maximal + block can begin past the halfway point and become the tail. Moving it to the slow axis dispatches all + the expensive blocks first. Measured **−20%** on a gfx942 fused attention backward (754 blocks ranging + 0–114 q-tile iterations), matching the greedy-list-scheduling prediction of 114/144 = **0.79**. Cost: + two SALU to decode the block index plus a permuted grid in the launcher. +- **Model 8 static queues of ~38 CUs, not one global CU pool.** A single-pool model reproduces headline + critical paths but mispredicts anything that changes grid *composition*, because the busiest XCD sets + the makespan. +- **Empty workgroups are free at the end of the grid and expensive in the middle.** Appending 5504 wholly + empty blocks cost nothing measurable; interleaving the *same* number between live blocks cost **+29%**, + because inserting blocks repartitions every later block across the XCDs (XCD load went from a balanced + 3193–3293 to a skewed 2243–4259 iteration units). Early-exit blocks are cheap, but do not scatter them + through the grid. This also qualifies the "≥1024 workgroups" lever above: pad by *extending* the grid, + never by interleaving no-op blocks among live ones. + +**Confirm the partition mode before trusting any of this** — SPX/NPS1 is what makes the mapping +predictable, and CPX changes both dispatch and device-scope atomic behaviour. If a dispatch-order ablation +comes out neutral when the model says it should be ~0.79, suspect the partition mode before the kernel. + ## Pitfalls - Default linear pid mapping ⇒ B/A panels re-fetched from HBM/foreign L2 instead of local L2 hits. - Tile count not a multiple of 8 ⇒ one or more XCDs finish early (load imbalance, tail latency). - <1024 workgroups on prefill ⇒ idle CUs/XCDs (`[[operators/dense_gemm/tuning.md]]`). - Over-grouping (too many tiles pinned to one XCD) ⇒ L2 thrash; size groups to the L2 capacity. - Assuming a unified L2 (it is **partitioned per XCD**). +- Optimizing the loop body before checking the grid axis order on a kernel with a structural cost gradient + across one axis — the axis order was the single largest and cheapest win in the measured case. +- **Designing around L2 residency for device-scope atomics.** A private L2 per XCD means such an atomic + must carry the uncached bit and execute past the fabric, so it misses L2 by construction (counted as + *Write and Atomic (Uncached)*, with equal atomic counts at the L1→L2 and L2→fabric interfaces). Fanning + atomics out across more slots to "spread contention" therefore only multiplies the DRAM footprint — + measured monotonically worse at 1/2/4/8 slots. ## Verify - Omniperf: L2 hit rate per channel / cross-XCD traffic, HBM read BW; XCD-aware order should *raise* L2 @@ -70,3 +104,6 @@ data-sharing tiles stays on one XCD's L2. Tune `tiles_per_xcd` to the L2 working - 8 XCDs, partitioned L2, round-robin dispatch: AMD CDNA3 whitepaper + MI300X Hot Chips 2024 architecture deck. - ≥1024 WGs, 8-multiple tiles, XCD/L2 placement levers: ROCm MI300X workload optimization guide. - Swizzle mechanics: `[[hardware/shared/l2_xcd_swizzle.md]]`. +- Longest-job-first grid axis (−20%, 114/144 model), static-XCD-queue modelling, empty-workgroup + placement (+29% interleaved vs free appended), and uncached device-scope atomics: first-party on-box + gfx942 MI300X SPX/NPS1, ROCm 7.1.0 diff --git a/perf_knowledge/profiling/benchmarking_methodology.md b/perf_knowledge/profiling/benchmarking_methodology.md index 6ee64795d..f3d1326e8 100644 --- a/perf_knowledge/profiling/benchmarking_methodology.md +++ b/perf_knowledge/profiling/benchmarking_methodology.md @@ -15,7 +15,8 @@ A trustworthy MI-GPU measurement is: **warm** (discard cold runs), **repeated** perf_knowledge e2e standard is **REPEATS=7**), inside a **noise band** (accept a change only if it clears the **~0.5%** e2e band), with **clocks controlled** (or at least monitored), and done as a **same-session, non-overlapping A/B** (ref vs candidate back-to-back). If the delta is inside the noise band, it is not -a result. Profiling perturbs timing, so measure in a *separate, untraced* pass from your counter/trace +a result. **One scoped exception**: an isolated single-kernel sweep over many build variants should report +the **minimum** over interleaved replays, not the median — see §scope exception below. Profiling perturbs timing, so measure in a *separate, untraced* pass from your counter/trace diagnosis ([`trace_analysis.md`](trace_analysis.md), [`rocprofv3_counters.md`](rocprofv3_counters.md)). ## Why MI300X is noisy (what you're fighting) @@ -41,6 +42,31 @@ Net: compute achieved TFLOP/s from *measured time*, never from assumed clock. near-zero host overhead, both to *get* the real GPU-bound time and as a perf technique when the trace shows host-launch gaps ([`trace_analysis.md`](trace_analysis.md)). +## Scope exception: single-kernel variant sweeps use the MINIMUM, not the median +The median-with-spread rule above is the **e2e** standard and stays the rule for anything you report as an +e2e delta. It is the wrong statistic for an **isolated single-kernel sweep over many build variants**, +where the failure mode is different: sustained back-to-back replays throttle the clocks, so the median +drifts *within* a run. A measured instance (gfx942 MI300X, attention backward): the median of an +**unchanged** kernel drifted **438 → 534 µs** across runs — larger than most effects you would be trying +to measure. The minimum over replays reflects unthrottled capability and is stable enough to rank variants. + +Two rules for this scope, and neither is optional: +1. **Report the minimum over replays**, never the mean or median. +2. **Build every variant up front, then time them round-robin in one process.** Clock drift then hits all + variants equally. Numbers from separate processes are not comparable, and a table you produce is only + internally comparable within one invocation — including the ablation rows, so re-derive any single + comparison by building both sides in **one** interleaved sweep rather than quoting two old runs. + +Print correctness per variant alongside the time, so a fast-but-wrong build is obvious, and print each +variant's delta against the first — that is the number to reason about. Keep ablation code paths alive +permanently for this reason: it is what makes a claim re-checkable on new hardware, and lets a regression +be bisected against a mechanism rather than a commit. + +**Do not use graphs for this sweep**: capturing one graph per variant and replaying round-robin faults, and +at ~1 ms per call launch overhead is far under the noise floor anyway. Graphs remain correct (and +necessary) for a *separate* small-kernel harness — a 4 µs kernel is otherwise swamped by ~40 µs of Python +dispatch per launch, so without graphs you are measuring the launch, not the kernel. + ## Per-leg vs 2-launch A/B For an e2e serving change, prefer a **2-launch A/B** (full ref launch vs full candidate launch) over summing **per-leg** microbenchmarks: per-leg sums miss overlap, caching, and dispatch interactions and @@ -59,6 +85,8 @@ repeats, with spread; never present theoretical peak as achievable. - Cold-cache first run counted in the median; clock not yet ramped. - Comparing across sessions/days — clocks, thermals, and background load differ. - Trusting summed per-leg microbenchmarks over a real e2e 2-launch A/B. +- Using the median for a throttling single-kernel variant sweep, or comparing two variants timed in + separate processes (§scope exception). ## Verify A real win clears the 0.5% band across REPEATS=7, reproduces on a re-run of the same A/B, and is