Skip to content

[Triton/Gluon] [Config] Add a tuned gfx1101 MHA config, split small_head/default - #4493

Merged
azaidy merged 2 commits into
ROCm:mainfrom
Ragua1:mha-config-gfx1101
Sep 28, 2026
Merged

azaidy merged 2 commits into
ROCm:mainfrom
Ragua1:mha-config-gfx1101

Conversation

@Ragua1

@Ragua1 Ragua1 commented Jul 31, 2026 •

Copy link
Copy Markdown
Contributor

Rewritten on 2026-09-10, after the branch was force-pushed from c27cabd7 to ec3af272. The earlier description covered the flat-layout file and values that no longer hold; that revision and the review anchored to it are unchanged in the history. What moved and why is in the comment below.

Motivation

aiter.ops.triton.attention.mha.flash_attn_func cannot run at all on gfx1101 (RDNA3). _get_config resolves configs/{arch}/triton/attention/mha/DEFAULT.json, and with no per-arch file and no fallback the call raises

FileNotFoundError: Required config file doesn't exist:
  aiter/ops/triton/configs/gfx1101/triton/attention/mha/DEFAULT.json

before any kernel is compiled. gfx1101 is in RDNA_ARCHS, so the architecture is declared supported while the default MHA path is unreachable on it. This PR is the narrow remedy for the one architecture I can measure on.

What changed since the approval on 2026-08-06

This PR was approved at c27cabd7, and that revision no longer holds, so the branch has been rewritten rather than merged as approved. Two things moved underneath it: the config layout (#5018 / #5019 deleted the flat path without a fallback) and the kernel (#3936, #4414).

Re-measured against today's main, the approved fwd/default is 1.274x slower on head_dim 128, and on head_dim 64 it no longer establishes a gain — the intervals overlap the donor's. What moved is the donor, not the candidate: the inherited gfx1151 entry became faster on today's kernel while the tuned one stood still. Numerics did not move at all, so a correctness re-verification passes with digit-identical residuals — re-verifying rather than re-measuring would have merged the regression.

Rather than withdraw the tuning, this revision splits it in two, which removes the regression instead of trading one shape against another.

Technical Details

One new file, aiter/ops/triton/configs/gfx1101/triton/attention/mha/DEFAULT.json, in the nested layout from #5019. 101 lines added, nothing removed, no code touched.

Nine of the eleven entries are byte-identical to gfx1151/triton/attention/mha/DEFAULT.json, the nearest tuned architecture. Two forward entries are tuned on gfx1101 instead of inherited:

fwd/default      BLOCK_M 64 -> 128, num_warps 4 -> 8, num_stages 2 -> 3   (vs. the gfx1151 donor)
fwd/small_head   BLOCK_M 128, num_warps 4, num_stages 1                   (new key, no donor equivalent)

fwd/small_head is the bucket added in #4414 for 16 < head_dim_v <= 64. Because it is opt-in per architecture by the presence of the key, this remains a data-only change that cannot affect any other architecture.

The split is needed because the two head_dim ranges want a different num_warps, and _get_config already discriminates on exactly that axis. Two independent coarse sweeps, one per range, selected these two entries separately — the numbers are below.

📌 The comment on that bucket does not describe this architecture, and these values should not be read as agreeing with it. It states that 16 < d <= 64 suffers a num_stages=1 pipelining pathology which num_stages=3 cures. On gfx1101 the ordering is the opposite: on the tuned M128 / N32 / w4 tile, num_stages 1 / 2 / 3 measured 2.346 / 2.429 / 2.450 ms. If that comment is meant as an invariant rather than as an observation about the architectures tuned so far, it changes the answer here, and I would rather be told than assume.

⚠️ num_stages is Triton-version-dependent on this architecture, so these values should not be assumed to carry backwards — click for why

Under the Triton used here, use_async_copy defaults to arch in ["gfx950", "gfx1250"] and is_pingpong_schedule_enabled admits only gfx942 and gfx950 (triton/backends/amd/compiler.py:43 and :33), so gfx1101 gets neither. num_stages therefore means classic software pipelining into registers only, which is why it can trade against BLOCK_M and why the answer differs between the two head-dim ranges. Under Triton 3.7.1 the same knob had a different gate — use_async_copy read the env knob without testing the arch — so these numbers should not be assumed to carry backwards.

fwd/dropout_or_fp32, both pe entries and both backward sections are inherited unchanged; only the forward path was exercised, so no claim is made about backward.

Test Plan

Windows native ROCm, RX 7800 XT (gfx1101, RDNA3, 60 CU), AITER checkout on PYTHONPATH with no install, AITER_TRITON_ONLY implicit on win32.

Every number below was measured on one software stack: torch 2.15.0a0+rocm10.1.0a20260822, HIP 7.16.26332, triton-windows 3.8.0.post28, a Windows build of Triton 3.8.0 and the newest release available. Older stacks are deliberately not reported, so that no row here describes a configuration that is no longer current.

fp16, seed 424242. Each point is the median of 20 timed iterations after 5 warmups with torch.cuda.synchronize() per iteration. Sweeps are single-shot rankings over a grid; every claimed ratio comes from a series of 10 independent repeats, and is claimed only where the [min, max] intervals across those repeats are disjoint. Where they overlap this is stated and nothing is claimed. Ratios are against the gfx1151 donor entry, since inheriting is the alternative to this file.

Two shape sets. Diffusion attention: SDXL self-attention (2, 4096, 4096, 10, 64), head_dim 64, and a Flux joint-attention shape (1, 4608, 4608, 24, 128), head_dim 128. LLM prefill, causal=True, GQA, batch 1, head dims from op_tests/op_benchmarks/triton/utils/model_configs.json (hidden_size / num_attention_heads); the sequence lengths are not in that file and are my choice, 4096 and 16384.

⚠️ The upstream MHA test and benchmark harness cannot run on this host, so its shape set is used with my own measuring path — click for the chain

op_tests/triton_tests/attention/test_mha.py and bench_mha.py both reach aiter.dtypes, which aiter/__init__.py binds only in the non-AITER_TRITON_ONLY branch, and the JIT path below it ends in RuntimeError("ROCm version file not found"). The upstream shape set is therefore used with my own measuring path rather than the upstream harness. This is independent of the present change, and I preferred not to introduce a workaround into the verification being reported.

Test Result

The defect is still present on the current base

With the new file removed from the checkout, on main @ 12620102 and the stack above:

FileNotFoundError: Required config file doesn't exist:
  .../aiter/ops/triton/configs/gfx1101/triton/attention/mha/DEFAULT.json
  _triton_kernels/attention/mha.py:955  ->  utils/config_utils.py:50

It is raised on the first call and on every later call in the same process. On that base the architectures that do have an MHA config are gfx942, gfx950, gfx1151 and gfx1250.

Two escape hatches work on the same host and the same shape: passing config= explicitly, and mha_set_impl("dao_ai"). Both return max abs 6.470e-05 against the fp32 reference, against 6.013e-05 for torch's own fp16 SDPA measured in the same process. The kernel therefore runs correctly on this architecture, and it is only the config lookup that makes the default path unreachable.

Routing and correctness

Verified on main @ 12620102 with this file applied, each entry in a separate process because _get_config memoises. head_dim 64 routing to small_head and head_dim 128 routing to default is the behaviour the split depends on.

All six forward entries resolve as intended, max abs 6.469e-05 to 7.008e-05 — click for the per-entry output
v_head_dim 128, fp16            -> default              max abs vs fp32 reference 7.008e-05
v_head_dim  64, fp16            -> small_head           max abs vs fp32 reference 6.469e-05
v_head_dim  64, fp16 + dropout  -> dropout_or_fp32      finite, dropout active
v_head_dim  64, fp32            -> dropout_or_fp32      max abs 8.643e-07
v_head_dim  64, pe              -> pe                   max abs 6.685e-05
v_head_dim  64, pe + dropout    -> pe_dropout_or_fp32   finite, dropout active

Every configuration measured in this campaign produced an identical maximum absolute difference per shape, the inherited one included, so tile parameters have no numerical effect on this kernel and the choice is purely a performance decision.

How the two entries were selected

One coarse sweep per head-dim range, 72 candidates each over BLOCK_M x BLOCK_N x num_warps x num_stages, ranked on that range's shape. The two sweeps are independent and each selected the entry this PR ships for its own range.

sweep shape best of 72 ms vs. donor approved M128 w4 s1 on the same shape
head_dim 64 M128 N32 w4 s1 = this PR's small_head 2.346 0.974 it is that entry
head_dim 128 M128 N32 w8 s3 = this PR's default 6.444 0.933 8.619, i.e. 1.25

The separating knob is num_warps: at head_dim 128 the same tile measures 8.619 ms at w4 s1 against 6.657 ms at w8 s1, while at head_dim 64 the ordering reverses. A refine stage over PRELOAD_V and waves_per_eu on the three best candidates produced no cell that beat M128 N32 w8 s3 pv0 we1 on head_dim 128.

Three of the 72 candidates on the head_dim 128 shape are unusable on this architecture, and the failure is the expected one: out of resource: shared memory, Required: 69632 (resp. 73728), Hardware limit: 65536.

head_dim 128 — the regression and its removal

10 repeats per configuration, and the whole series run twice. [min, max] over the repeats in ms, ratio of medians against the donor in brackets.

series donor M64 w4 s2 approved M128 w4 s1 this PR's default, M128 w8 s3
run 1 [6.805, 7.028] [8.575, 8.776] (1.273) ❌ [6.437, 6.503] (0.944) ✅
run 2 [6.787, 7.034] [8.698, 8.789] (1.274) ❌ [6.411, 6.513] (0.943) ✅

✅ / ❌ = intervals disjoint from the donor's. The two runs agree to the third significant figure, and the same M128 w8 s3 entry costs 1.070–1.077 on head_dim 64 — which is the reason one fwd/default cannot serve both ranges.

LLM prefill shapes, causal=True, GQA, batch 1

Ratios against the donor, whose median in ms is given for scale.

All ten shapes: the entry this file routes to is the faster candidate in every row, significantly so in all ten — click for the table
shape head_dim donor ms M128 w4 s1 M128 w8 s3 routed by this file to
llama3-8B, sq 4096 128 3.885 1.331 ❌ 0.951 (overlap) default
llama3-8B, sq 16384 128 56.949 1.316 ❌ 0.904 ✅ default
mixtral-7B, sq 4096 128 3.950 1.331 ❌ 0.935 ✅ default
mixtral-7B, sq 16384 128 57.483 1.285 ❌ 0.895 ✅ default
kimi-k2.5 tp4, sq 4096 112 8.344 1.298 ❌ 0.892 ✅ default
kimi-k2.5 tp4, sq 16384 112 143.159 1.136 ❌ 0.775 ✅ default
deepseek-V3, sq 4096 56 8.746 0.828 ✅ 0.924 ✅ small_head
deepseek-V3, sq 16384 56 167.621 0.644 ✅ 0.714 ✅ small_head
glm-4.7-fp8 tp4, sq 4096 53 5.964 0.915 ✅ 1.019 (overlap) small_head
glm-4.7-fp8 tp4, sq 16384 53 93.844 0.852 ✅ 0.922 ✅ small_head

In every row the entry this file routes the shape to is the faster of the two candidates, and in nine of the ten it is significantly faster than the donor. Comparing the two candidates directly, their [min, max] intervals are disjoint in all ten rows — the routing decision is significant on every shape in the set, even where the winner's margin over the donor is not.

This set matters because #3560 tuned the donor on n=112 LLM shapes, so choosing fwd/default on diffusion shapes replaces an LLM-derived value with a diffusion-derived one. The measurement says the two families do not disagree on this card: both want num_warps 8 at head_dim 112–128, and the axis that separates them is head_dim, not workload.

head_dim 64, fwd/small_head

The LLM rows above are the evidence for this entry, and they are the only significant evidence for it.

On the diffusion shape it ranked first of 72 candidates in the sweep, but across 10 repeats it measures 0.933 and 0.936 with intervals that overlap the donor's in both runs, so no gain is claimed there.

The entry is nonetheless the right one for its range, and the comparison that establishes that is against the alternative rather than against the donor. Giving fwd/default to both ranges measures 1.070–1.077 on the diffusion head_dim 64 shape, and on the four head_dim 53–56 LLM shapes it is slower than M128 w4 s1 by 0.924 vs 0.828, 0.714 vs 0.644, 1.019 vs 0.915 and 0.922 vs 0.852 — disjoint intervals in all four.

What this does not establish

Batch > 1, the varlen / thd layout, sliding window, and the decode phase (seqlen_q = 1) were not measured. Forward path only; both backward sections are inherited. One card, one host and one software stack — no other RDNA or CDNA part is available to me, and nothing here is a claim about gfx1151, whose values this PR does not change.

A cross-attention shape (2, 4096, 77, 10, 64) was measured and is deliberately not reported: the donor's own timing on it varies by 31 % across repeats on this host, which is far above every other shape in the set, so it discriminates nothing in either direction.

A cross-architecture fallback for the same reachability defect is a separate change and is deliberately not bundled here. I am happy to file it separately if that is wanted.

@Ragua1
Ragua1 requested a review from a team July 31, 2026 21:42
@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:triton-300x Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
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: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
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 4493 --add-label <label>

@zufayu
zufayu requested a review from vgokhale August 5, 2026 02:48
@vgokhale
vgokhale requested a review from carlushuang August 5, 2026 15:54
carlushuang
carlushuang previously approved these changes Aug 6, 2026

@carlushuang carlushuang left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

LGTM

gfx1101 (RDNA3, e.g. RX 7800 XT) ships no MHA config, so `_get_config` in
`_triton_kernels/attention/mha.py` finds no
`configs/gfx1101/triton/attention/mha/DEFAULT.json` and every call to
`aiter.ops.triton.attention.mha.flash_attn_func` fails before a kernel runs.
Of the architectures in RDNA_ARCHS, only gfx1151 ships one today.

Nine of the eleven entries are taken verbatim from the gfx1151 donor
(RDNA3.5, added in ROCm#3423, tuned in ROCm#3560), which is the nearest tuned
architecture. Two forward entries are tuned on gfx1101 instead of inherited,
and they differ from each other only in `num_warps` and `num_stages`:

  fwd/default     BLOCK_M 128, num_warps 8, num_stages 3   (donor: 64 / 4 / 2)
  fwd/small_head  BLOCK_M 128, num_warps 4, num_stages 1

The split uses the `small_head` bucket added in ROCm#4414, which is opt-in per
architecture by the mere presence of the key, so this stays a data-only
change. It is needed because one `fwd/default` cannot serve both halves on
this card: measured against the donor, `M128 w4 s1` is 0.935x on head_dim 64
but 1.269x on head_dim 128, while `M128 w8 s3` is 0.914x on head_dim 128 but
1.062x on head_dim 64.

Measured on Windows native ROCm, triton 3.8.0, fp16, 10 independent repeats
of 20 iterations after 5 warmups, `torch.cuda.synchronize()` per iteration;
a result counts only when the [min, max] intervals across repeats are
disjoint. Ratios are against the gfx1151 donor entry, i.e. against what a
donor-inherited config would do.

  head_dim 128, `default`     Flux joint 0.914x / 0.891x / 0.892x on three
                              independent torch+ROCm stacks (2.11/7.15,
                              2.11/10.1, 2.15/10.1), all pinned to the same
                              triton; llama3-8B 0.882x, mixtral-7B 0.880x,
                              kimik25-tp4 0.849x at seqlen 16384
  head_dim <= 64, `small_head` SDXL self-attn 0.935x; deepseek-V3 0.658x and
                              glm47fp8-tp4 0.857x at seqlen 16384

The LLM shapes come from `op_tests/op_benchmarks/triton/utils/model_configs.json`
(prefill, causal, GQA, batch 1); the sequence lengths are not in that file and
are chosen here. Not covered: batch > 1, varlen/thd, sliding window, decode.

Note that the `small_head` comment in `_get_config` does not describe gfx1101.
It states that 16 < d <= 64 suffers a num_stages=1 pipelining pathology which
num_stages=3 cures. On this card the ordering is the opposite -- on the tuned
M128/N32/w4 tile, num_stages 1/2/3 measured 2.300 / 2.399 / 2.492 ms. The
bucket is still the right mechanism here, for a different reason: the two
head_dim ranges want a different `num_warps`, not a different `num_stages`.

Signed-off-by: Martin Domanský <ragua@email.cz>
@github-actions github-actions Bot changed the title [triton-mha] add gfx1101 tuning config [Triton/Gluon] [Config] add gfx1101 tuning config Sep 10, 2026
@Ragua1

Ragua1 commented Sep 10, 2026

Copy link
Copy Markdown
Contributor Author

The branch has been force-updated from c27cabd7 to ec3af272 and the description rewritten. GitHub dismissed the earlier approval on that push, and correctly so: it does not cover what is on the branch now, so this needs a fresh look rather than a merge on the strength of it. Apologies for the churn.

Two independent reasons, neither of which was visible when it was approved.

First, the path. The approved revision added configs/gfx1101-MHA-DEFAULT.json in the flat layout, which #5018 / #5019 removed without a fallback. _get_config now resolves configs/gfx1101/triton/attention/mha/DEFAULT.json, so merging the approved revision as-is would have added a file that nothing reads.

Second, and the reason this took longer: on re-measurement against today's main the approved values no longer hold either. The tuned fwd/default is 1.274x slower on head_dim 128, and on head_dim 64 it no longer establishes a gain at all — its intervals overlap the donor's. The approved values were tuned against an older kernel and an older Triton, and what moved since is the donor rather than the candidate: the inherited gfx1151 entry got faster on the current kernel while the tuned one stood still. Numerics did not move at all, so a correctness re-verification passes with digit-identical residuals; re-verifying rather than re-measuring would have merged the regression.

Rather than withdraw the tuning, this revision splits it across the small_head bucket added in #4414 (16 < head_dim_v <= 64) and fwd/default, which removes the regression instead of trading one shape against another. Two independent coarse sweeps, one per head-dim range, each selected the entry the file now ships for that range.

Both halves are then confirmed on ten LLM prefill shapes taken from op_tests/op_benchmarks/triton/utils/model_configs.json, which were not used for tuning. At head_dim 112–128 the default values give 0.775–0.935x against the donor on five of the six shapes, and 0.951x with overlapping intervals on the sixth. At head_dim 53–56 the small_head values give 0.644–0.915x on all four. Comparing the two candidates against each other rather than against the donor, their intervals are disjoint on all ten of those shapes, so the routing decision is supported on every shape in the set.

⚠️ Every number in the description comes from one current software stack — torch 2.15 on ROCm 10.1 with Triton 3.8.0 — and the whole series was run twice on it. Earlier stacks are deliberately not reported, so that no row describes a configuration that is no longer current. Protocol, the full tables and the shapes where nothing is claimed are in the description.

Still one new file, data only, no code touched, and no other architecture is affected — small_head is opt-in by the presence of the key.

One thing in the description I would rather flag here as well: the comment on that bucket in _get_config states that 16 < d <= 64 suffers a num_stages=1 pipelining pathology which num_stages=3 cures. On gfx1101 the measured ordering is the opposite, so this file uses the bucket for a different reason than the comment gives. If the comment is meant as an invariant rather than as an observation about the architectures tuned so far, please say so — that would change the answer.

@mgehre-amd — you added the gfx1151 MHA config in #3423 and tuned it in #3560, and nine of the eleven entries here are byte-identical to that file, so you are likely the right person to judge the two that are not. I have deliberately not asked @carlushuang to re-review: the values now on the branch are not the ones that were approved, and I would rather have that judged fresh than carried over.

@Ragua1 Ragua1 changed the title [Triton/Gluon] [Config] add gfx1101 tuning config [Triton/Gluon] [Config] Add a tuned gfx1101 MHA config, split small_head/default Sep 10, 2026
@Boss2002n
Boss2002n requested a lite review from Copilot September 11, 2026 01:05

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

Unresolved correctness, backward-coverage, and bf16-validation findings must be addressed before approval.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

Adds a gfx1101 Triton MHA configuration, making default dispatch reachable and splitting forward tuning by head dimension.

Changes:

  • Adds tuned fwd/default and fwd/small_head entries.
  • Adds inherited PE and backward configurations.
  • Adds architecture-specific MHA configuration data.

Review findings:

  • Critical (2 votes): Add d=8/16 correctness coverage or a safe route.
  • Moderate (2 votes): Validate or defer/gate gfx1101 backward support.
  • Moderate (1 vote): Measure bf16 or add dtype-specific selection.
File summaries
File Description
aiter/ops/triton/configs/gfx1101/triton/attention/mha/DEFAULT.json Adds gfx1101 MHA tuning and dispatch entries.
Review details

Suppressed comments (1)

aiter/ops/triton/configs/gfx1101/triton/attention/mha/DEFAULT.json:19

  • [verified] The new tiles are selected for both fp16 and bf16 because _get_config has no dtype-specific config branch, while every reported tuning run is fp16 and the benchmark defaults to bf16. A config that wins fp16 is not established to be a good bf16 choice, so this can regress the default model dtype on gfx1101. Author must measure the routed d64/d128 cases in bf16 as well, or add dtype-specific selection before shipping these values.
      "BLOCK_M": 128,
      "BLOCK_N": 32,
      "PRELOAD_V": false,
      "waves_per_eu": 1,
      "num_warps": 8,
      "num_ctas": 1,
      "num_stages": 3
  • Files reviewed: 1/1 changed files
  • Comments generated: 2
  • Review effort level: Lite

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@Ragua1

Ragua1 commented Sep 11, 2026

Copy link
Copy Markdown
Contributor Author

Both review findings are answered in their own threads, with measurements on gfx1101 rather than from the source alone. Neither finding leads to a change in the file. The values stand as submitted.

What remains unmeasured is listed here rather than left to be discovered:

  • fp8, on any section.
  • The three inherited forward sections fwd/dropout_or_fp32, fwd/pe and fwd/pe_dropout_or_fp32. They are byte-identical to the gfx1151 donor and were not measured on gfx1101.
  • Backward performance. The backward arm is a correctness smoke test only.

If any of these is required before approval, please say which and it will be measured.

One question is left open deliberately, because it is a decision about _get_config rather than about this file. It is raised in a separate comment, so it can be answered or declined independently of this PR.

@Ragua1

Ragua1 commented Sep 11, 2026

Copy link
Copy Markdown
Contributor Author

A question about _get_config's bucketing rather than about this PR's file, raised separately so it can be answered independently.

The small_head branch selects on a band, 16 < head_dim_v <= 64. The hazard its comment documents is attached to a value instead:

Using num_stages=3 recovers performance and is numerically verified for these dims, but regresses d128 and miscompiles d<=16, so only 16 < d <= 64 uses this path.

The band and the value are equivalent only as long as fwd/default does not itself carry num_stages 3.

Reading the four architectures that ship an MHA table on main at 12620102:

arch ships fwd/small_head fwd/default fwd/small_head
gfx942 no BLOCK_M 128, num_warps 4, num_stages 1 —
gfx1151 no BLOCK_M 64, num_warps 4, num_stages 2 —
gfx1250 no BLOCK_M 128, num_warps 4, num_stages 1 —
gfx950 yes BLOCK_M 128, num_warps 4, num_stages 1 BLOCK_M 128, num_warps 4, num_stages 3

On gfx950, the only architecture that has the bucket, fwd/default is at s1 and small_head at s3. The bucket is what raises num_stages for the middle band, and the d <= 16 clause is what stops that rise from reaching the small dimensions. The band and the hazard coincide there.

Tuning gfx1101 produced the opposite polarity. d64 selected num_warps 4 with num_stages 1, and d128 selected num_warps 8 with num_stages 3, so in this PR small_head holds s1 and default holds s3.

That the two architectures tune differently is not itself surprising. gfx950 is a CDNA part with Wave64 and matrix cores, gfx1101 is RDNA3 with Wave32, and different optima are expected. What does not carry across is the protection: with the polarity reversed, the same lower bound routes d <= 16 toward s3 rather than away from it.

So the observation is not that the values differ. It is that the bucket expresses a band, while the constraint it was introduced to enforce is a property of a value. Those two agree for only one of the two polarities.

This is not a defect report. On gfx1101, num_stages showed no numerical effect at any head dimension, in fp16 or bf16, including under causal masking and at dimensions that are not powers of two. The question is whether the dispatch is intended to express the constraint at all.

Three readings, and the code does not tell us which is intended:

  1. The band is the whole intent, and it is each architecture's own responsibility to keep num_stages 3 out of fwd/default. If so, that requirement is today discoverable only from a comment on the branch it does not apply to. A line in configs/CLAUDE.md would make it checkable by whoever adds the next table.
  2. The d <= 16 clause is meant to bind generally. A lower bucket would then express it in the dispatch, rather than relying on every architecture's default staying below s3.
  3. The clause is specific to gfx950 and the compiler it was measured on, and is not meant to constrain other architectures.

Which of the three is it? If the answer is the second, we are willing to prepare that change as a separate PR. It touches the elif chain and every architecture's table, so it should not ride along with a config-only change, and it needs your view on the intended shape first.

@cagrikymk

Copy link
Copy Markdown
Contributor

@Ragua1 I assume your measurements come from the latest aiter main vs your configs? Since there was some delay, I wanna make sure perf. measurements are not stale.

@Ragua1

Ragua1 commented Sep 26, 2026

Copy link
Copy Markdown
Contributor Author

@cagrikymk Yes. I re-ran the full campaign today against main @ 8253efc4, the base of your merge 518a4ed0. The software stack is the same as in the description (torch 2.15.0a0+rocm10.1.0a20260822, HIP 7.16.26332, triton-windows 3.8.0.post28), so aiter is the only thing that changed. The protocol is also unchanged: fp16, 5 warmups and 20 timed iterations per point, 10 independent repeats, and a gain counts only when the [min, max] intervals are disjoint. The results agree with the description and the file needs no change.

description (12620102) 8253efc4
no gfx1101 MHA config on main → FileNotFoundError yes yes
all six forward entries vs fp32 reference pass pass, identical max_abs
sweep winner, head_dim 64 / 128 small_head / default entries same two entries
previously approved fwd/default, head_dim 128 1.273 / 1.274 ❌ 1.265 ❌
fwd/default, head_dim 128 0.944 / 0.943 ✅ 0.927 ✅
small_head, diffusion head_dim 64 overlap, no gain claimed overlap, no gain claimed
small_head, deepseek-V3 sq 16384 0.644 ✅ 0.649 ✅
small_head, glm-4.7-fp8 tp4 sq 16384 0.852 ✅ 0.859 ✅
LLM shapes where the routed entry is significantly faster than the donor 9 of 10 9 of 10, same exception (llama3-8B sq 4096)
the two candidates' intervals, 10 LLM shapes disjoint in all 10 disjoint in all 10

Ratios are medians against the inherited gfx1151 values. ✅ / ❌ = intervals disjoint from the donor's.

@cagrikymk

Copy link
Copy Markdown
Contributor

@Ragua1 Thank you for the contribution!

@azaidy I think this PR is ready to be merged.

@azaidy azaidy 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.

LGTM!

@azaidy
azaidy merged commit f1f95ce into ROCm:main Sep 28, 2026
71 checks passed
@Ragua1
Ragua1 deleted the mha-config-gfx1101 branch September 28, 2026 16:34
vgokhale added a commit that referenced this pull request Oct 2, 2026
* [Config] Add gfx942 a8w8 blockscale GEMM tunings for Qwen3/Qwen3.5/GLM/DSV4 shapes (#5839)

MI325X (gfx942, 304 CUs): 286 (M, N, K) rows over 28 weight shapes that no
existing model_configs table covers, tuned on main with
gemm_a8w8_blockscale_tune.py --libtype ck --splitK. Keys already present in
model_configs and rows slower than the heuristic default are left out.

The rows go into each model's existing model_configs table; Qwen3-14B and
Qwen3.5-122B-A10B get new tables. A shape two models share is kept in one
table only, since the config merge rejects duplicate keys across files.

Co-authored-by: Cursor <cursoragent@cursor.com>
(cherry picked from commit 8260de6)

* [HIP] [OPUS] [FlyDSL] gfx950 MXFP8 e8m0 GEMM + DeepSeek-V4/V4.1 tuned configs (#5896)

* feat(flydsl): gfx950 MXFP8 e8m0 GEMM on FlyDSL + DeepSeek-V4/V4.1 configs

- New FlyDSL gfx950 MXFP8 batched GEMM (bmm_a8w8_mxscale_gfx950): scaled
  MFMA 16x16x128, async LDS DMA pipeline, 32x32 / 128x128 e8m0 blocks,
  split-K with same-XCD last-arrival reduction, B direct to registers, XCD
  tile order, non-temporal B, per-stage or preloaded scale panels, K % 64
  tail, column-major (blockscale) x_scale. Scale rows that are not whole
  dwords (e.g. 128-wide blocks at K = 384 / 768) are copied a byte per LDS
  slot with exact buffer bounds. check_bmm_config is the single legality
  source.
- Arch-neutral front door flydsl.batched_gemm_a8w8 dispatching by (arch,
  w_scale block, x_scale layout); gfx950 module with kernelName parsing,
  untuned-shape heuristic and a compiled-launcher cache (host 71 -> 21 us).
- gemm_a8w8_blockscale_bpreshuffle on gfx950 routes e8m0 x_scale + w_scale
  to it through the tuned table AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_MXSCALE_
  BPRESHUFFLE; FP32 scales keep CK / asm. gemm_a8w8_blockscale with
  isBpreshuffled forwards native group32 operands there. The V4 wo_a
  batched GEMM uses the same kernel.
- AOT precompiles kernelId=bmm rows (both x_scale layouts for 128x128).
- Tuned gfx950 configs: DeepSeek-V4-Pro 128x128 linears (980 rows, TP1/4/8,
  1.30x vs CK/asm), DeepSeek-V4.1 group32 (1224 rows, 1.22x vs Triton),
  V4 wo_a batched (680 rows).
- Also carries the other local changes in the tree: inverse_rope_group_quant,
  opus policy / bmm tune, gfx1250 bmm wrapper, tensor_shim helpers.

* refactor(flydsl): tidy gfx950 MXFP8 bmm kernel per FlyDSL cleanup guide

- Use fx.copy instead of fx.copy_atom_call for the single-atom copies.
- Build the async-LDS DMA destination from the LDS pointer with
  fx.add_offset + fx.to_llvm_ptr instead of ptrtoint/inttoptr with a
  hand-picked address space. The raw buffer_load_async_lds stays: the
  BufferLoadAsyncLDS atom has no 1-byte size and no cache-policy operand.
- Factor the compile-and-run + leaked ir.Context recovery out of
  tensor_shim._run_compiled into _compile_and_run, and use it for the
  bmm wrapper's per-config compile cache too.

Generated ISA is byte-identical on 9 reference configs.

(cherry picked from commit 40d524b)

* [Triton/Gluon] [ASM] [HIP] Mha v4: adds bf16 sparse, LSE support, KV varlen, fixes, etc (#5798)

Motivation: adds log-sum-exp output to MHA v4 gfx950 kernels (only dense variants now), so it can run under ring / context parallelism, which merges per-rank partials via LSE. Also extends block-sparse to the BF16 recipes, canonicalizes the MXFP4 rows, and improves K/V quantization accuracy.

# Kernels:
- New BF16 and BF16FP8 sorted-sparse kernels.
- Dense LSE epilogue on all ten hd128 recipes, gated at runtime on s_lse.
- f8f6, f6f4 and mxfp4 sparse rows moved to FP6-P V, matching their dense siblings.
- Dense mxfp4 claims the canonical FP6-P V order; the duplicate f4f4 row is disabled.

# Host:
- bugfix: per-channel V amax clamped so empty heads cannot quantize to NaN.
- Optional lse output on mha_v4 / mha_v4_packed, plumbed through asm_mha_v4_fwd.cu.
  + ABI unchanged: ptr_lse / s_lse / s_lse_Hs were already reserved in the kernarg.
- _LSE_CAPABLE_QV gates the supported format pairs; sorted-sparse still raises.
- K mean is subtracted before quantizing (K-smoothing), fused into the MX quantizer kernels.
- Dense MHA v4 accepts per-batch key lengths (ragged seqlen_k).

# Minor:
- Retired MXFP4 Q/K + FP8 V and the deprecated mha_v4_mxfp8 alias.
- bench_sage.py: improve input distributions, diffusion-calibrated default, BF16 sparse modes.
- Split the block-sparse cases into op_tests/test_mha_v4_sparse.py.

(cherry picked from commit 105615e)

* [tuner] Fail a task as soon as its worker process exits (#5841)

A GPU memory fault aborts the worker process, so its task never returns a
result. mp_tuner only noticed at the task timeout (1800s by default) or
never without one. Record which worker started each task and, when that
process is gone, fail the task and restart the pool right away, like the
existing accelerator-error path. Fixes #5840.

Co-authored-by: Cursor <cursoragent@cursor.com>
(cherry picked from commit bcb56d9)

* [CI] Update Aiter artifact downloads to v8.0.1 (#5908)

(cherry picked from commit c03689f)

* [MLA v4 nm] Test fix _run_one_point reading packed BF16 as FP32 partials (#5905)

* [MLA v4 nm] Fix _run_one_point reading packed BF16 as FP32 partials

Since #4311 the v4 nm dispatcher derives out_16_nosplit from num_kv_splits
and ignores the caller's value, so a single-split launch writes packed BF16
into the logits buffer. _run_one_point still read logits_buf[:, 0] as FP32
for num_kv_splits == 1, which compared every other BF16 element (plus the
never-written tail of the buffer) against the reference and printed spurious
`fp8_dequant_ref vs asm` checkAllclose failures in the script-mode sweep.

Read output_buf, which holds the final result for every split count.

Co-authored-by: Cursor <cursoragent@cursor.com>

* [MLA v4 nm] Gate accuracy checks so a failed! fails the test

checkAllclose only raises on a catastrophic delta; otherwise it logs
`failed!` and returns the mismatch fraction. Four of the six accuracy
checks in test_mla_v4_nm.py dropped that return value, which is how the
packed-BF16 readback bug fixed in the previous commit passed both pytest
and the script-mode CI run.

Route them through _gated_allclose, which asserts the mismatch fraction
against the same tol_err_ratio checkAllclose uses for `failed!`. The
script-mode sweep keeps going past a failing shape, prints a summary, and
exits non-zero so aiter_test.sh reports the file as failed.

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
(cherry picked from commit 049fae4)

* [AOT] Inline the FlyDSL FP8 FMHA head shapes, drop the config CSVs (#5901)

Follow-up to #5796. The AOT job list for the gfx950 FlyDSL FP8 flash
attention was driven by a header-only aiter/configs/fmha_fp8_aot.csv
merged with aiter/configs/model_configs/*_fmha_fp8_aot.csv. A single row
does not justify the CSV plumbing, so the head shapes now live in a
DEFAULT_SHAPES table in the module, the way mega_moe.py already does it,
and the comments introduced by #5796 are trimmed.

- fmha_fp8.py: DEFAULT_SHAPES replaces parse_csv()/DEFAULT_CSVS;
  default_jobs() replaces the CSV walk; --csv is gone (--shape still
  overrides). The cu_num column is dropped with the CSVs: the kernel is
  gfx950-only and AOT_ARCH is fixed, so every non-gfx950 row was warned
  about and skipped anyway.
- common.py: FMHA_FP8 returns default_jobs() next to MEGA_MOE instead of
  going through collect_aot_jobs().
- jit/core.py: drop AITER_CONFIG_FMHA_FP8_AOT and its config-file
  property, now unused.
- README.md: document the table instead of the CSVs.

_variant_space()/jobs_for_shape() are unchanged, so coverage is
unchanged: --list still emits the same 92 kernel names for Kimi-K3 TP8
(12:12:192:128 varlen_cross) as the CSV-driven version, and setup.py
source builds still compile them via run_aot().

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
(cherry picked from commit 71a31be)

* [FlyDSL] gfx942 fp8_mqa_logits: let _auto_variant choose rows_per_block (#4963)

* [FlyDSL] gfx942 fp8_mqa_logits: let _auto_variant choose rows_per_block

_auto_variant returned f"mfma_r2_w{wpb}", so rows_per_block was pinned at 2 and
seven of the nine registered variants -- including every member of the r4 family
-- could never be selected.

r4 amortizes each KV tile load over twice as many query rows and is faster from
seq_len 8 upward. The gap is widest exactly where it costs most: vLLM chunks
indexer prefill to fit VLLM_SPARSE_INDEXER_MAX_LOGITS_MB (512 MB), which caps
seq_len at 1024 when seq_len_kv is 131072, so the existing seq_len >= 2048 branch
cannot fire at long context and every such call took mfma_r2_w4 -- the median of
the nine by speed, with the best 1.65x faster.

Measured on MI325X (gfx942), seq_len_kv 131072, best variant vs the r2 pick:

  seq_len    1   r2 25.2 us    r4 79.1 us     r4 3.1x worse
  seq_len    4   r2 24.3 us    r4 31.5 us     r4 1.3x worse
  seq_len    8   r2 34.0 us    r4 32.9 us     r4 1.03x better
  seq_len   16   r2 56.5 us    r4 47.5 us     r4 1.19x better
  seq_len 1024   r2 2564.3 us  r4 1520.3 us   r4 1.69x better

Below seq_len 8 the host padding of seq_len up to a multiple of RPB dominates --
at seq_len 1 an r4 kernel computes 4 rows to obtain 1 -- so r2 is kept there and
behaviour for those shapes is unchanged.

Logits are bitwise identical across all variants at every shape tested, so this
is purely a blocking/occupancy change.

End to end on 8x MI325X, TP8, GLM-5.2-FP8, 131072 in / 1024 out, concurrency 8:
median TPOT improves 7.22% and output throughput 6.60%. This kernel is 16.8% of
GPU time at that point.

Signed-off-by: Jin Tao <jin.tao@amd.com>

* [FlyDSL] gfx942 fp8_mqa_logits: pick RPB on element count, and keep it a divisor

Refines the previous commit's rule after a 2-D sweep. That rule keyed RPB off
seq_len alone with a crossover measured only at seq_len_kv=131072; sweeping the
other contexts shows the crossover is not a seq_len threshold at all, and that a
second effect was being read as one.

RPB tracks the logits element count. Over seq_len 1..8192 x seq_len_kv
1024..262144 on MI325X, the boundaries land on the same element count at every
context: RPB=1 wins below 2**19 elements (27/27 shapes), RPB=2 at 2**19 (6/6),
RPB=4 from 2**21 up (38/38), with 2**20 a transition band split 3/3. Keying off
seq_len instead put the previous rule on the wrong side at low context: at
seq_len 16, seq_len_kv 1024 it chose RPB=4 and ran 1.26x slower than RPB=1.

RPB must also divide seq_len. When it does not, the launcher pads with four
torch.cat calls; that is a flat ~44 us of host-side overhead, independent of
seq_len_kv, and it is the whole of the "small seq_len" penalty the previous
commit attributed to wasted rows. At seq_len 1, seq_len_kv 131072: RPB=1 23.1 us,
RPB=2 67.8 us, of which the four cats are 44.1 us and pre-padding by hand
recovers all of it (21.9 us). So the penalty is not proportional to the padding
-- 1 wasted row of 2 costs the same as 3 of 4 -- and it applies to every odd
seq_len, which the old rule sent to RPB=2 unconditionally.

Stepping down to a divisor is only right while the kernel is cheap relative to
that fixed cost, so it is gated to the same 2**21 elements: at seq_len 1025,
seq_len_kv 131072 the dividing RPB=1 takes 3880 us against 2601 us for a padded
RPB=2.

Measured on MI325X, no FLYDSL_FP8_MQA_LOGITS_VARIANT set, old pick vs new:

  seq_len  seq_len_kv        old         new   speedup
     1024      131072   r2_w4 2598.7   r4_w4 1623.5    1.60x
     1025      131072   r2_w4 2609.3   r4_w4 1594.0    1.64x
      512      131072   r2_w4 1190.1   r4_w4  795.3    1.50x
      700       50000   r2_w4  661.6   r4_w4  440.3    1.50x
      333       12000   r2_w4  118.4   r4_w4   97.1    1.22x
       16      131072   r2_w4   58.3   r4_w4   47.4    1.23x
        3      131072   r2_w4   66.4   r1_w4   28.1    2.36x
        1      131072   r2_w4   69.4   r1_w4   23.2    2.99x
        1        1024   r2_w4   66.2   r1_w4   22.2    2.98x

No shape measured regresses; the smallest gain is 1.04x. Against the best of the
nine variants at each shape, pooled over held-out data (non-power-of-two shapes,
a fine seq_len sweep, and head counts 16 and 64), the geometric mean cost falls
from 1.45x to 1.03x and the worst case from 3.17x to 1.41x.

Logits are bitwise identical across all nine variants at all 180 shapes swept
(1620 timings), so this remains purely a blocking/occupancy change.

WPB is deliberately left alone. It is worth a few percent at most here, and
unlike RPB its optimum moves with the head count -- at 64 heads the current
WPB rule costs 1.61x worst case where a fixed WPB=4 costs 1.07x -- so it needs
its own sweep rather than a change fitted to one head count.

Signed-off-by: Jin Tao <jin.tao@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* [FlyDSL] gfx942 fp8_mqa_logits: trim _auto_variant comments per review

Move RPB element-count thresholds into _auto_variant and shorten the
docstring; logic unchanged.

Signed-off-by: Jin Tao <jin.tao@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* [FlyDSL] gfx942 fp8_mqa_logits: scope the _auto_variant step-down note

The docstring stated the divisor step-down as an unconditional rule, but it
only applies in the middle band: above the top threshold RPB stays 4 and the
padding is accepted. Say which band it applies to, and why the top band is
exempt. No logic change.

Co-authored-by: Cursor <cursoragent@cursor.com>

* [FlyDSL] gfx942 fp8_mqa_logits: unit-test the variant selector

The shape sweep in test_flydsl_fp8_mqa_logits.py never reaches r4 -- its
largest default shape is 1024 x 1560, below the 2**21 threshold -- and nothing
asserted _auto_variant directly, so a regression could re-pin RPB to 2 with
every correctness test still green.

Cover the RPB bands at both thresholds, that the edges track
seq_len * seq_len_kv rather than seq_len alone, the middle-band step-down on
odd seq_len (and that it stops above the top threshold), the production
1024 x 131072 shape, the unchanged WPB rule, and _resolve_variant precedence
so the auto path is confirmed to be the default. Pure shape arithmetic, so no
kernel launch.

Co-authored-by: Cursor <cursoragent@cursor.com>

* [FlyDSL] gfx942 fp8_mqa_logits: put selector tests on the CI path

CI only discovers op_tests/test_*.py, so the r1/r2/r4 selector coverage
never ran. Move it there, run pytest from __main__, and add two r4
shapes to the GPU sweep so auto-selected padding is launched.

Signed-off-by: Jin Tao <jin.tao@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* [FlyDSL] gfx942 fp8_mqa_logits: fold the selector pins into the op test

aiter op tests are plain scripts, not pytest, so drop
test_flydsl_fp8_mqa_logits_variant.py and pin the gfx942 auto-selected
variant in test_flydsl_fp8_mqa_logits.py instead: a host-only
verify_auto_variant table at both RPB band edges, either side of the
odd-seq_len step-down, the long-context prefill shape and the WPB switch.
It runs in the default verify scenario on gfx942, and its failures count
toward the same exit code as the kernel sweep. The selector import sits in
the file's existing ImportError guard.

Signed-off-by: Jin Tao <jin.tao@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Signed-off-by: Jin Tao <jin.tao@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: jin.tao@amd.com <jin.tao@amd.com@tus1-p15-g43.tus.tensorwave.lan>
Co-authored-by: Felix Li <felix.li@amd.com>
(cherry picked from commit bbaccdf)

* [Triton] [MHA] Add a tuned gfx1101 config, split small_head/default (#4493)

gfx1101 (RDNA3, e.g. RX 7800 XT) ships no MHA config, so `_get_config` in
`_triton_kernels/attention/mha.py` finds no
`configs/gfx1101/triton/attention/mha/DEFAULT.json` and every call to
`aiter.ops.triton.attention.mha.flash_attn_func` fails before a kernel runs.
Of the architectures in RDNA_ARCHS, only gfx1151 ships one today.

Nine of the eleven entries are taken verbatim from the gfx1151 donor
(RDNA3.5, added in #3423, tuned in #3560), which is the nearest tuned
architecture. Two forward entries are tuned on gfx1101 instead of inherited,
and they differ from each other only in `num_warps` and `num_stages`:

  fwd/default     BLOCK_M 128, num_warps 8, num_stages 3   (donor: 64 / 4 / 2)
  fwd/small_head  BLOCK_M 128, num_warps 4, num_stages 1

The split uses the `small_head` bucket added in #4414, which is opt-in per
architecture by the mere presence of the key, so this stays a data-only
change. It is needed because one `fwd/default` cannot serve both halves on
this card: measured against the donor, `M128 w4 s1` is 0.935x on head_dim 64
but 1.269x on head_dim 128, while `M128 w8 s3` is 0.914x on head_dim 128 but
1.062x on head_dim 64.

Measured on Windows native ROCm, triton 3.8.0, fp16, 10 independent repeats
of 20 iterations after 5 warmups, `torch.cuda.synchronize()` per iteration;
a result counts only when the [min, max] intervals across repeats are
disjoint. Ratios are against the gfx1151 donor entry, i.e. against what a
donor-inherited config would do.

  head_dim 128, `default`     Flux joint 0.914x / 0.891x / 0.892x on three
                              independent torch+ROCm stacks (2.11/7.15,
                              2.11/10.1, 2.15/10.1), all pinned to the same
                              triton; llama3-8B 0.882x, mixtral-7B 0.880x,
                              kimik25-tp4 0.849x at seqlen 16384
  head_dim <= 64, `small_head` SDXL self-attn 0.935x; deepseek-V3 0.658x and
                              glm47fp8-tp4 0.857x at seqlen 16384

The LLM shapes come from `op_tests/op_benchmarks/triton/utils/model_configs.json`
(prefill, causal, GQA, batch 1); the sequence lengths are not in that file and
are chosen here. Not covered: batch > 1, varlen/thd, sliding window, decode.

Note that the `small_head` comment in `_get_config` does not describe gfx1101.
It states that 16 < d <= 64 suffers a num_stages=1 pipelining pathology which
num_stages=3 cures. On this card the ordering is the opposite -- on the tuned
M128/N32/w4 tile, num_stages 1/2/3 measured 2.300 / 2.399 / 2.492 ms. The
bucket is still the right mechanism here, for a different reason: the two
head_dim ranges want a different `num_warps`, not a different `num_stages`.

Signed-off-by: Martin Domanský <ragua@email.cz>
Co-authored-by: Mehmet Cagri <mehmet.kaymak@amd.com>
(cherry picked from commit f1f95ce)

* [Triton/Gluon] Consolidate tuning harnesses (#5874)

* [Triton/Gluon] Consolidate tuning harnesses

* [Triton/Gluon] Simplify tuning harness

(cherry picked from commit af3514a)

* [Bugfix][Gluon][MLA] Follow-up: fix stale comments + test hardening (#5860)

Addresses review feedback on #5648 (kept separate to not disturb the approved PR):
- mla_gluon.py: the >2GB global_load calls now carry a bounds mask + other=0.0,
  so the old "No mask needed / in-bounds" comments above them were stale and
  gave the opposite (unsafe) guidance. Replace with an accurate one-liner.
- test_mla.py: allocate the output with device=q_nope.device instead of relying
  on the module default device; call torch.cuda.empty_cache() before the
  mem_get_info() free-memory gate to avoid allocator-state flakiness; and assert
  exact parity (atol=rtol=0) between the >2GB and <2GB paths, which read
  identical KV and must be bit-identical (a loose tolerance could hide a
  regression).

Signed-off-by: Rohan138 <rohanpotdar138@gmail.com>
(cherry picked from commit a966245)

* [Triton/Gluon] [Config] gemm_a16w16: gfx950 tuned per-shape configs (#5830)

* [Triton] gfx950: tuned defaults for gemm_a16w16, gmm, MHA fwd and PA decode

Tuned on MI355X (gfx950) and validated for correctness and no regressions on
broader shape sets than tuned. All changes are gated to gfx950 configs/arch.

- gemm_a16w16: per-shape configs for N,K = 2048/2048, 4096/4096, 8192/8192,
  10240/8192, 57344/8192, 8192/28672 (non-persistent and persistent). Only the
  tuned M bucket differs from DEFAULT.json. 1.07-2.5x.
- gmm: new "large_kn" config (8 warps), selected for K, N >= 4096 with
  >= 256 rows per group. 1.14-1.26x there; smaller problems keep "default".
- MHA fwd (Triton): new "mid_head" config (128x128, 2 stages) for bf16/fp16
  with 64 < head_dim <= 128. 1.03-1.16x; d<=64 and d>128 unchanged.
- PA decode (Triton): use v2 whenever there is more than one partition for
  bf16/fp16 KV. 1.9-5.3x for the batch/context sizes that previously hit v1.

Unit tests: test_gemm_a16w16 360 passed, test_gmm 48 passed,
test_pa_decode 816 passed, test_mha 2060 passed.

* [Triton] pa_decode: sort imports (ruff I001)

* Move gmm, MHA fwd and PA decode changes to their own PRs

Per review, each kernel gets its own PR; this PR keeps only the gemm_a16w16
tuned configs. The removed changes are on branches gfx950-tuned-gmm,
gfx950-tuned-mha-fwd and gfx950-pa-decode-v2.

* [Triton] gemm_a16w16_persistent: gfx950 tuned config for N=K=2048

(cherry picked from commit a651db0)

* replacing the fp32 mfma with two bf16 mfma. fn is split inside the kernel into hi = bf16(fn) and lo = bf16(fn -hi). a simple bf16 downcast for fn inside the kernel drops accuracy more and performs worse for large M. also change the concatenation of the four streams from column major to stream major (k = stream * TILE_K + column). retuned the gfx950 configs. (#5885)

(cherry picked from commit f10cd2a)

* perf(fused-moe): add tuned DSV4.1 TP4 configs (#5919)

(cherry picked from commit 977ae79)

* [Config] Add Qwen3.8-27B TP1 a8w8 blockscale GEMM tunings for gfx942 (#5585)

* [Config] Add Qwen3.8-27B TP1 a8w8 blockscale GEMM tunings for gfx942

#3324 covered this model family at TP=2/4/8 only, so the five widths
Qwen3.8-27B drives at TP1 have no tuned entries for gfx942/cu_num=304
and every call falls back to the default -- 275 "not found tuned
config" messages per profile run, now 4. This adds 641 rows in one
file, 249 decode (M <= 512) and 392 prefill (M 907-65536), from a shape
list extracted out of real vLLM server logs rather than generated as an
M ladder.

Decode, against AITER's own default, measured in a real vLLM serving
run -- Qwen3.8-27B-FP8 at 128K context, max-concurrency 1, on one
MI325X, byte-identical images, the config bind-mount the only variable,
three interleaved repeats per arm on an idle node:

  decode step   15.444 ms -> 12.449 ms   1.24x

Per-shape decode GEMM vs the default: down 2.33x, out_proj 1.51x,
in_proj 1.12x, qkv 1.00x, gate_up 1.00x. No decode shape regresses.
Decode-step spread across repeats was 0.037 ms and 0.012 ms.

The 392 prefill rows ship and are tuned, but their end-to-end effect is
currently unmeasurable on this workload: gemm_a8w8_blockscale silently
stores only M*N mod 2^31 output elements once M*N reaches 2^31, and at
a 65,536-token prefill chunk the 34,816-wide gate_up projection crosses
that limit. The truncation depends only on M*N and not on which kernel
the config selects -- AITER's default truncates identically -- so it is
neither introduced nor worsened by this change.

Tuned with gemm_a8w8_blockscale_tune.py --libtype ck --splitK, decode
rows then re-timed interleaved because the tuner's min() over short
samples is noisy on near-ties; only rows beating the default by >=3%
were kept, following #5421. test_csv_validation (15 tests) and
test_config_shape_collision (17) both pass, there are zero duplicate
(gfx, cu_num, M, N, K) keys against the 12 other sources in the runtime
merge set on main, and max relative Frobenius error is 4.11e-3 against
bf16's 2^-8 = 3.9e-3 rounding floor.

Scope: only the decode rows have in-situ evidence, and MI300X shares
the gfx942/304 key but is untested here. Two rows pin AITER's default
kernel (qkv and gate_up at M=1) after the trace measured the tuned
picks slower -- pinned, not deleted, because getPaddedM(1,N,K,0) == 16.
Six decode near-ties ship no row. gate_up ships 72 of 80 prefill rows;
the 8 missing have M*N > INT32_MAX and fault during tuning, but inherit
the tuned M=8192 row at runtime via the same padding collapse.

Signed-off-by: Pham Binh <phamhuuthanh.binh@amd.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* [Config] Thin Qwen3.8-27B TP1 GEMM rows to the padded M ladder

Lookup already rounds M through getPaddedM, so a row per traced M is redundant.
Keep the power-of-two M list used by the other gfx942 blockscale tables, and
omit shapes that main already ships.

---------

Signed-off-by: Pham Binh <phamhuuthanh.binh@amd.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: akii96 <aakif.nawaz@amd.com>
(cherry picked from commit 330f127)

* [HIP] [OPUS] [JIT] Unify gfx950 + gfx1250 MXFP4 paged MQA-logits into one module, schedule and API (#5761)

Fold the gfx950 (#5332) and gfx1250 (#5656) OPUS MXFP4 paged MQA-logits ops into one
implementation. The device math of both bodies is unchanged.

- One JIT module, module_pa_mqa_logits_mxfp4_opus. Sources move to
  csrc/kernels/opus_mqa_logits/pa_mqa_logits_mxfp4/:
  - _opus.h: ABI, kargs, traits;
  - _sched.cuh: shared builder;
  - _gfx950.cuh / _gfx1250.cuh: arch bodies, each an empty stub on the other arch;
  - _kernels.cu: launcher.
  Namespaces are opus_logits::gfx950 / ::gfx1250, with generic names (pa_mqa_logits_mxfp4_traits
  and _kernel). The per-arch sources and module_pa_mqa_logits_mxfp4_gfx1250_opus are deleted.
- One schedule: gfx1250's build_tiles + build_sched serve both arches. build_tiles is skipped at
  q_per_block == 1, where the cut is the identity, which saves one launch per forward on every
  gfx950 call and on gfx1250's MTP=1 decode.
- One runtime-arch dispatch: fwd_sched probes the arch once per process, then dispatches on
  (q_per_block, block_k). An unmatched config raises.
- Host-visible bounds on both arches:
  - num_rows is checked against q, weights and out, and against the lengths of local_ends,
    local_starts and row_to_batch;
  - block_tables width is checked in KV tiles;
  - both kernels drop records with row_id >= num_rows.
  The caller contract (four conditions) is documented in the module docstring and the header.
- One Python API (aiter.ops.opus.pa_mqa_logits_mxfp4: plan_buffers / plan / pa_mqa_logits_mxfp4),
  with per-arch MqaLogitsVariant instances. gfx1250: qlen4_kv64 / qlen1_kv64. gfx950: qlen1_kv64 /
  qlen1_kv256, cta_resident 1024.
  - Retired: #5332's gfx950 entry points (pa_mqa_logits_mxfp4_sched / _build_sched /
    _sched_slots / _sched_buffer_ints).
  - Kept: #5656's public gfx1250 API, except the gfx950-only plan(block_k=).
- -mllvm -enable-post-misched=1 is applied on gfx950 only. It is load-bearing there; on gfx1250
  it slowed chunked prefill.
- kargs drops six unread fields (144 -> 112 B).
- One op test for both arches, op_tests/test_pa_mqa_logits_mxfp4_opus.py. It adds ATOM-shaped
  padded decode, host-raise and raw row-guard cases, a real NaN control, and on gfx950 a FlyDSL
  cross-check.

(cherry picked from commit 9b3885f)

* [FlyDSL] feat(mega_moe/gfx1250): bind mori tokoff-ext allocator on the mori dispatch path (#5810)

* feat(mega_moe/gfx1250): bind mori tokoff-ext allocator on the mori dispatch path

mori's op layer builds the tokoff-ext slot allocator in
EpDispatchCombineOpHip.__init__; driving EpDispatchPlan directly bypasses
that and leaves EpArgs.tokOffPeers null, keeping dispatch on the
serializing cco-window atomic. Build it here from mori's TokOffExt, gated
exactly like mori (default on; MORI_EP_TOKOFF_EXT=0 opts out), pass its
peer-pointer array to plan.launch, and free it in close().

Needs mori with a public TokOffExt (dispatch_combine_v2.hip_backend).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* style(mega_moe/gfx1250): black-format the tok_off_peers ternary

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

---------

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: HaonanWang98 <hwang@amd.com>
(cherry picked from commit 475cf0f)

* [Triton] Do not write segment softmax state at NUM_SEGMENTS_PER_SEQ == 1 (#5935)

mla_decode_fwd splits each sequence's KV into NUM_SEGMENTS_PER_SEQ
segments and lets a reduce kernel merge the per-segment softmax max and
expsum. At one segment there is nothing to merge, so the host skips the
reduce kernel and hands both scratch pointers the output buffer itself:

    else:
        segm_output = out
        segm_max = out  # dummy ptr
        segm_expsum = out  # dummy ptr

The decode kernels store M and L unconditionally, so at one segment they
write the softmax state straight over the attention output. Nothing
faults and nothing warns; the result is simply wrong.

Guard both stores. NUM_SEGMENTS_PER_SEQ is a constexpr, so the branch
costs nothing when it is greater than one.

Only the gfx1250 Gluon kernel reaches this today: select_3d_config ends
its gfx12 branch at max(1, ...), while the other branch floors at
MIN_SEGMENTS >= 8. The segment count falls as batch x heads grows, so on
gfx1250 it reaches one at large batch -- a DeepSeek-R1 TP2 serving run at
--max-running-requests 256 sits in that range throughout, and its decode
is silently corrupted. The plain Triton kernel shares both the unguarded
store and the host-side aliasing, so guard it as well, before a future
tuning change makes it reachable.

test_mla_decode_fwd stays green either way, which is why this went
unnoticed. Its grid does reach one segment for the larger head count,
but the corruption lands only at out.flat[token * num_query_heads +
head] -- one element in kv_lora_rank, measured at 0.04% to 0.12% of the
output -- while the assertion allows a tol_err_ratio of 0.01. Catching
it needs a case with no error budget.

Checked on gfx1250. DeepSeek-R1 TP2 at page size 64, gsm8k over 2000
questions, scores 0.945 with the fix. A standalone sweep over batch 1 to
256, bf16 and fp8 e4m3 caches, sequence lengths that are not page
multiples and pages scattered through the pool gives a worst relative
error of 0.0033 for decode and prefill alike. Compared against a torch
reference with no error budget, a one-segment batch differs at a 0.08%
ratio before this change -- softmax maxima around 100 sitting where
attention outputs near 0.02 belong -- and matches to the last element
after it. The plain Triton kernel, forced on by disabling the gfx12
branch, matches the same reference at batch 1 to 256, confirming the
guard leaves the multi-segment reduce path intact. test_mla.py is
unchanged: its 128 pre-existing failures are all shuffled_kv_cache=True
on the pipelined kernel, identical before and after.

Signed-off-by: Lin, Soga <soga.lin@amd.com>
(cherry picked from commit 2b6ff3d)

* [Triton/Gluon] Enable unified-attention skip-mask for gfx950 hd256 FP8 prefill (#5598)

Long FP8 prefill at head size 256 matches `D_GEQ_256` (`BLOCK_M=16`, no `SPLIT_UNMASKED_LOOP`) because D outranks Q in the lookup, so it never reaches the hd128 skip key. Add the same `Q_GEQ_256` family under `D_GEQ_256`, with `SW` and `SHUF` siblings left skip-off, so the existing kernel fast path fires only for non-windowed, non-shuffled prefill.

(cherry picked from commit e289613)

* [Triton/Gluon] [Kimi-K3][ROCm] Add merged MoE front (#5321)

* feat(kimi-k3): add minimal large-M MoE front

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* feat(kimi-k3): enable merged front for M=7 decode

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* feat(kimi-k3): tune merged front for M=14 decode

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* perf(kimi-k3): tune merged front decode bucket

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* test(kimi-k3): validate the merged front across the MTP decode bucket

test_decode_full_front_matches_reference stopped at m=16, so the shapes an MTP
server actually decodes at were never checked against the torch.mm reference:
with num_speculative_tokens=3 a pure-decode step submits num_seqs * 4 tokens,
which is 32 and up for any server past four concurrent sequences.

Extend the parametrization to 32, 48, 64, 80, 96, 112, 128 and 192 -- the shapes
the companion vLLM change enables, and for which kimik3_bf16_tuned_gemm.csv
already ships tuned solutions at N=6016, K=7168, gfx950, cu_num=256. No kernel
or config change is needed; this only closes the correctness gap left by the
old parametrization.

gfx950: 9 passed -> 17 passed.

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* style(kimi-k3): satisfy the pinned ruff on this PR's own files

`ruff==0.16.0`, the version .github/workflows/pre-checks.yaml pins, reports two
errors in files this PR introduces: RUF022 on the `__all__` list in
kimi_k3_moe_front.py and I001 on the import block in its test. Both are
autofixes and neither changes behaviour -- `__all__` ordering only affects
`import *`, and every symbol here is imported by name.

The two EXE001 findings that remain in the repo are in unrelated flydsl files
and predate this PR.

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* perf(kimi-k3): retune the M=16 merged-front FP32 row

The decode shape is 4 * concurrency, and getPaddedM(gl=0) rounds every M <= 256
up to 16, so this row serves both the C1 (M=4) and C4 (M=16) decode buckets --
about 40% of iterations in the 8k/1k replay.

It was tuned against a single weight buffer. The merged front's weight is
6016x7168 BF16 = 86.2 MB and MI355X has 256 MB of LLC, so that regime serves
most of the GEMM from cache; the shipped 16.005 us is faster than anything the
shape can reach when the weight is actually streamed. In the real model 92
distinct MoE layers stream 7.9 GB per decode step and nothing is resident.

Re-searched all 681 hipblaslt solutions with the weight chained across 604 MB
of distinct buffers (past LLC) and timed inside a CUDA graph, which is how vLLM
runs decode:

  solidx 443935 (shipped)  22.364 us  3857 GB/s
  solidx 443486 (this)     19.279 us  4474 GB/s   -13.8%

443486 is not a new solution -- the table already carries it at M=7. The right
kernel was present and keyed to the wrong M.

Max relative error vs a torch.mm FP32 reference is 4.9e-07, so this is a
dispatch change only. The us/tflops/bw columns are the streamed measurement and
are therefore not comparable to the cache-hot numbers in neighbouring rows.

* perf(kimi-k3): size the front-GEMM config cache to the whitelist

_KIMI_K3_MERGED_FRONT_TOKEN_COUNTS admits 28 distinct token counts but the
config cache held 16, so the M values a mixed decode/prefill server cycles
through evicted each other.

Minor: the tuned table is cached upstream, so a miss costs a few get_padded_m()
extension calls rather than a CSV parse, and in the graphed decode path it is
paid at capture instead of replay. This just makes the cache cover the set it
was sized for.

* refactor(kimi-k3): separate Triton epilogue layers

Move the launchable epilogue kernel under _triton_kernels/moe, keep only the public wrapper in ops/triton/moe, and add a config-aware kernel repr and CUDA Graph benchmark. Remove Kimi-specific weight packing and hipBLASLt orchestration from the Triton module.

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* refactor(moe): generalize the SiTU epilogue

Rename the op to describe the SiTU epilogue it actually performs rather than implying that it projects the incoming activation. Make all branch widths caller-provided, mask partial tiles for arbitrary shapes, and reuse the shared Triton tanh helper.

Update the numerical and CUDA Graph coverage plus the benchmark for the generic API.

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* test(triton): use Triton arch helper for SiTU epilogue

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

* fix(triton): avoid branch-local mask type collisions

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>

---------

Signed-off-by: jiacao-amd <jiahui.cao@amd.com>
Co-authored-by: Jiahui Cao <jiacao@crs-m2m-cpu-spur-014.us-east2-a.compute.internal>
(cherry picked from commit aa84815)

* [Triton/Gluon] Add SonicMoE pure-Triton grouped GEMM MoE (#5725)

* [Triton] Add SonicMoE pure-Triton grouped GEMM MoE

Add a pure-Triton, expert-major grouped GEMM MoE with full autograd
support (SonicMoE), covering top-k and general routing, GLU/elementwise
activations, and blockwise FP8 scales for the grouped GEMM. Includes
gfx942/gfx950 tuned configs, unit tests, and a benchmark script.

Public entry points: moe_TC_softmax_topk_layer, moe_general_routing_inputs,
moe_pre_routed_inputs (aiter.ops.triton.sonicmoe).

This is the first of two PRs splitting the SonicMoE contribution; a
follow-up PR stacks the hipBLASLt/multistream grouped GEMM backend on
top of this pure-Triton implementation.

* [Triton] Drop unused SonicMoE stream_id dels and restating comments

Keep the positional stream argument for caller compatibility as
_stream_id, and remove comments that only restated the next GEMM call.

* [Triton] Move SonicMoE autograd API out of _triton_kernels

Keep @triton.jit kernels under _triton_kernels/moe/sonicmoe and put
the autograd wrappers plus public entry points in aiter.ops.triton.sonicmoe.

* [Triton] Use AMD copyright headers on SonicMoE kernels

Replace third-party author banners and drop external source-link
comments so new files match aiter's SPDX header.

* Address SonicMoE review feedback

* [Triton] Move SonicMoE routing kernels into moe_routing

* [Triton] Collapse SonicMoE host wrappers into one module

Keep the public API and tests in a single file instead of a set of sibling wrappers.

(cherry picked from commit 2e62094)

* [Triton/Gluon] Move the KDA_DECODE configs into the nested config layout (#5941)

(cherry picked from commit 92192fd)

* [Triton/Gluon] Move _triton_kernels/gated_delta_rule/ to _triton_kernels/gated_delta_net/ (#5943)

(cherry picked from commit cf89ffd)

* [Triton/Gluon] Move _gluon_kernels/gfx1250/norm/ to _gluon_kernels/gfx1250/normalization/ (#5944)

(cherry picked from commit 9ef0d08)

* [CI] Select impacted Triton and Gluon unit tests (#5878)

Select impacted Triton and Gluon unit tests

(cherry picked from commit 32b1cb0)

* [Triton/Gluon] [Config] Drop no-op kpack from RDNA GEMM configs (#5917)

* [Triton] [Config] Drop no-op kpack from RDNA GEMM configs

* [Triton] [Config] Drop redundant matrix_instr_nonkdim from selected RDNA GEMM configs

* [Triton] [Config] Extend RDNA matrix_instr_nonkdim cleanup

* [Triton] [Config] Retain matrix_instr_nonkdim for shared GEMM consumers

(cherry picked from commit 9ef3e41)

* mxfp8 gemm cga update, update (#5055)

Co-authored-by: Satya Nikhil Kodukula <nikhil.kodukula@gmail.com>
(cherry picked from commit 693701c)

* stop fused_bmm_rope_kv_cache from using batched_gemm_a8w8_a_per_token_group_prequant_w_per_batched_tensor_quant configs (#5877)

(cherry picked from commit a2650cd)

* [Triton/Gluon] Add Triton-based Conv3D kernels (#5952)

* conv3d implementation on Triton

* Remove unnecessary convolution compatibility aliases

* Update the README file.

* Add Conv2D weight-pack cache cleanup and updated the tests.

* Simplify the Conv2D Winograd test guard

* Move Conv3D device validation out of shape helper

* Add diagnostics to convolution test assertions

* Removing benchmark related tests from unit test

* Add context to Conv3D numerical assertions

* Deduplicate Conv2D and Conv3D blocked-layout kernels

* Simplified benchmark and fixed formatting issue on test file

* Split Conv2D and Conv3D cache-clear tests into their owning suites

* Deduplicate Conv2D and Conv3D Winograd transforms

* Deduplicate the Conv2D and Conv3D Winograd filter transform

* Consolidate Conv2D Winograd launch paths

* Deduplicate Conv2D and Conv3D prepack helpers and cache wrappers

* Black formatting

* Replace convolution test helper prints with logging

* Use safe configs for Conv3D on CDNA

(cherry picked from commit c8325e0)

---------

Signed-off-by: Jin Tao <jin.tao@amd.com>
Signed-off-by: Martin Domanský <ragua@email.cz>
Signed-off-by: Rohan138 <rohanpotdar138@gmail.com>
Signed-off-by: Pham Binh <phamhuuthanh.binh@amd.com>
Signed-off-by: Lin, Soga <soga.lin@amd.com>
Signed-off-by: jiacao-amd <jiahui.cao@amd.com>
Co-authored-by: siliangchen-amd <SiLiang.Chen@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Lingpeng Jin <103567126+valarLip@users.noreply.github.com>
Co-authored-by: Jesús Carabaño <jcaraban@users.noreply.github.com>
Co-authored-by: Leo <drleonid@amd.com>
Co-authored-by: liyjiang <liying.jiang@amd.com>
Co-authored-by: gbyu-amd <Guanbao.Yu@amd.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Co-authored-by: Jin Tao <jintao12@amd.com>
Co-authored-by: jin.tao@amd.com <jin.tao@amd.com@tus1-p15-g43.tus.tensorwave.lan>
Co-authored-by: Felix Li <felix.li@amd.com>
Co-authored-by: Martin Domanský <8312516+Ragua1@users.noreply.github.com>
Co-authored-by: Mehmet Cagri <mehmet.kaymak@amd.com>
Co-authored-by: Satya Nikhil Kodukula <nikhil.kodukula@gmail.com>
Co-authored-by: Rohan Potdar <rohanpotdar138@gmail.com>
Co-authored-by: Nimit Patel <61071220+NimitPtl@users.noreply.github.com>
Co-authored-by: Muhammad Ahmed <mm.ahmed2202@gmail.com>
Co-authored-by: yifehuan-amd <yifehuan@amd.com>
Co-authored-by: Pham Binh <phamhuuthanh.binh@amd.com>
Co-authored-by: akii96 <aakif.nawaz@amd.com>
Co-authored-by: li,xiangxiang <xiangxli@amd.com>
Co-authored-by: jhchouuu <jiahzhou@amd.com>
Co-authored-by: HaonanWang98 <hwang@amd.com>
Co-authored-by: sogalin_codegen <39478626+sogalin@users.noreply.github.com>
Co-authored-by: vorapolsiloai <115975949+vorapolsiloai@users.noreply.github.com>
Co-authored-by: jiacao-amd <jiahui.cao@amd.com>
Co-authored-by: Jiahui Cao <jiacao@crs-m2m-cpu-spur-014.us-east2-a.compute.internal>
Co-authored-by: WuLei-AMD <leiwu@amd.com>
Co-authored-by: Hyunjune Kim <132782704+hyjuunn@users.noreply.github.com>
Co-authored-by: Shao-Chun Lee <Shao-Chun.Lee@amd.com>
Co-authored-by: Saeid Rostami <123997133+saeid-rostami@users.noreply.github.com>
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