[Perf] [AMD] Qwen3-Next: route GDN decode at AITER's fused gfx950 kernel by rbrugaro-amd · Pull Request #39554 · sgl-project/sglang

rbrugaro-amd · 2026-09-15T06:28:57Z

[AMD] Qwen3-Next: route GDN decode at AITER's fused gfx950 kernel

Requires AITER ≥ dd6b8907 — ROCm/aiter#5520
([Triton/Gluon] Fused Qwen3-Next GDN decode op for gfx950) merged
2026-10-06 and added fused_gdn_decode_qkvz /
fused_gdn_decode_qkvz_supported to aiter.ops.triton.gated_delta_net.
This PR only consumes them.

With the env var unset (the default) this PR does not import AITER at all.

Motivation

Qwen3-Next's Gated DeltaNet decode step runs four kernels per layer, per step,
across 36 GDN layers:

fused_qkvzba_split_reshape_cat  ->  causal_conv1d_update
  ->  fused_recurrent_gated_delta_rule_packed_decode  ->  gated RMSNorm

AITER ships a single Gluon kernel that does all four and consumes the packed
qkvz/ba projections directly. This wires Qwen3-Next to it.

+4.7% to +8.6% output throughput on 4x MI355X / TP4, no accuracy
regression. Off by default; opt in with
SGLANG_QWEN3_NEXT_GDN_FUSED_BACKEND=aiter.

Note the fused proj+Conv1D path that already exists in the GDN backend
(fused_qkvzba_causal_conv1d_update_contiguous) is Qwen3.5-only —
_gdn_decode_fused_proj_conv appears in qwen3_5.py and nowhere else — so
Qwen3-Next still runs the full unfused chain today.

Modifications

The fusion spans SGLang's model/backend boundary, which is why this is not a
one-line dispatch change. Qwen3GatedDeltaNet.forward owns the qkvz split
(stage 1) and the gated RMSNorm (stage 5); GDNAttnBackend.forward_decode owns
the conv and recurrence (stages 2-4) and the state pools. Neither side holds all
the inputs: the backend receives post-split mixed_qkv and never sees z or
the norm weight, while the model never touches the mamba pools.

This follows the attempt-and-verify stash that Kimi-K3 already uses for its
fused KDA decode (kimi_k3.py / kda_backend.py): the model offers the
output-norm weight on the attention layer, the backend consumes it only if a
covered kernel engages, and the model applies its own norm when the stash comes
back unconsumed. Purely additive — no signature changes, and it fails safe.

New:

kernels/ops/attention/gdn_fused_decode_aiter.py — fail-closed adapter:
enabled() (env), available() (platform, AITER import, Triton Gluon),
covered() (per-call contract) and run(). covered() delegates the tensor
contract to AITER's own fused_gdn_decode_qkvz_supported() so the two cannot
drift, and adds the one thing AITER cannot see — the model's output-gate
activation, which the kernel hard-codes to SiLU.
test/srt/test_gdn_fused_decode_aiter.py — the opt-in, the capability probe
and every decline path. The decline cases need neither a GPU nor AITER, which
is the point: they are exactly what CI can check on any runner.

Modified:

srt/layers/attention/linear/gdn_backend.py — _try_aiter_fused_gdn_decode()
inside the existing isinstance(mixed_qkv, tuple) branch. Returns post-norm
output and sets the consumed flag, or None to fall through to today's chain.
Logs once on engage and once on reject-with-reason.
srt/models/qwen3_next.py — the decode-only stash and fast path,
_gdn_fused_decode_ready (static gate: platform, opt-in, head ratio, head
dims, conv width), and _prepare_gdn_fused_decode() called from
load_weights to publish the bf16 norm weight before graph capture. A zero
conv bias is pre-allocated at init for models without one, so nothing is
allocated inside a captured decode graph.

The backend's tuple input (projected_qkvz, projected_ba) already exists for
Qwen3.5's fused path; Qwen3-Next simply had not opted in. No new mechanism.

Scope

The CUDA path is unaffected. The fused path requires is_hip() and gfx950,
and with the env var unset the probe returns False before AITER is imported,
leaving the existing chain byte-identical. Every gate degrades to the current
behaviour:

gate	on failure
env var unset (default)	stock chain
not HIP / not gfx950	stock chain
AITER or Triton Gluon unavailable	stock chain, logged once
Triton < 3.8 (see below)	stock chain, logged once with the reason
shape/dtype not covered	stock chain, logged once with the reason
output-gate activation is not SiLU	stock chain

Following _can_fuse_accept_state's framing in kda_backend.py: this is a
capability, not a mode — falling short of the contract falls back rather than
raising.

SGLang authors no Gluon; this only consumes AITER's, and Gluon availability is
probed rather than pinned, as aiter_mla_gluon.py does.

A note on Triton 3.7 images

Some ROCm builds ship Triton 3.7 carrying a backported Gluon that imports
cleanly but miscompiles the layout ops this kernel uses. AITER therefore gates
on triton >= 3.8 rather than on an import probe. On such an image this feature
declines and falls back — correct, but inert. Concretely, of the 2026-10-12
nightlies, rocm724 ships Triton 3.7.0 and rocm10 ships 3.8.0.

Environment

Hardware: 4x AMD Instinct MI355X (gfx950), TP4 / EP1.
Image: lmsysorg/sglang-rocm:v0.5.21-rocm10-mi35x-20261012.
Model: Qwen/Qwen3-Next-80B-A3B-Instruct-FP8.

component	version	note
sglang	2026-10-12	matches main at time of testing
Triton	3.8.0	native — no manual install needed on this image
AITER	2026-09-25 in-image	predates #5520; install AITER ≥ `dd6b8907`

Both arms differ by exactly this PR's commits — the baseline is this
branch's own merge base (f92199b569), not an older image default, so no
unrelated upstream drift is folded into the comparison. Random dataset,
ISL/OSL 1024/1024, --random-range-ratio 1,
num-prompts == max-concurrency. All numbers un-profiled.

Accuracy Tests

python3 -m sglang.test.few_shot_gsm8k --num-questions 1319 --parallel 256
(the old benchmark/gsm8k/bench_sglang.py has been removed from main):

	decode cap 64	decode cap 256
baseline	0.943	0.940
this PR	0.946	0.942
Invalid (both arms)	0.000	0.000

Standard error at n=1319, p~0.94 is 0.0065, so this is "no detected
regression" rather than a measured equality. One run per arm per
configuration.

Kernel-level correctness lives in the AITER PR: the op is verified
bit-identical to its reference implementation on all five outputs across
batch 1-256, and its test suite covers the pad-slot contract, every dispatch
branch and determinism.

Speed Tests

conc	baseline tok/s	this PR tok/s	delta	baseline TPOT	this PR TPOT
1	368.20	391.42	+6.3%	5.36 ms	5.04 ms
8	2712.32	2850.68	+5.1%	5.75 ms	5.47 ms
32	8301.50	8987.28	+8.3%	7.40 ms	6.82 ms
64	13240.56	14377.92	+8.6%	9.20 ms	8.43 ms

Engagement is confirmed structurally from the server log — one
AITER fused GDN decode engaged line and zero not covered lines — so the
throughput above is attributable to the fused path actually running, not to a
silent fallback.

Raised graph cap

AITER's op carries a second decomposition above batch 128, which only fires when
--cuda-graph-max-bs-decode exceeds its default of 64. Same two arms at a cap
of 256:

conc	baseline tok/s	this PR tok/s	delta	baseline TPOT	this PR TPOT
128	21701.42	22863.24	+5.4%	11.05 ms	10.44 ms
256	31343.49	32821.80	+4.7%	14.98 ms	14.25 ms

The gain narrows at the top of the range, as expected — the decode kernel is a
smaller share of step time once the batch is large enough to be bandwidth-bound.

At this cap SGLang captures 36 decode graph sizes and the op covers all of
them; selection inside AITER is a range check, not a shape table.

Kernel level

Per decode step, 36 GDN layers, TP-0, from torch profiler traces. These traces
were taken on the earlier ROCm 7.0 image and have not been re-collected on the
image above — the ratio is the claim, not the absolute microseconds:

kernel	baseline	this PR
`fused_recurrent_gated_delta_rule_packed_decode_kernel`	36x, 361.6 us	0
`fused_qkvzba_split_reshape_cat_kernel`	36x, 163.5 us	0
`_causal_conv1d_update_kernel`	36x, 155.7 us	0
`layer_norm_gated_fwd_kernel`	36x, 149.9 us	0
AITER fused GDN decode	0	36x, 279.2 us
total	830.8 us	279.2 us (2.98x)

Methodology notes

Both arms were run with an idle-GPU assertion (no server processes, <3% VRAM)
before each arm, not just between them. An earlier ladder produced a spurious
-80% baseline result from contention with a stale server; it reproduced within
that session, which is precisely why the guard now runs per arm.

Absolute throughput here is not comparable to the figures previously quoted
in this PR: those came from a ROCm 7.0 image with Triton 3.8 installed manually,
these from ROCm 10 with native 3.8 — a different compiler and a different
absolute scale. The arm-vs-arm deltas are the like-for-like comparison, and they
are unchanged across that move and a ~1700-commit jump in main.

Known gaps

The FP8 epilogue is not wired up. The AITER op can emit per-head group-128
FP8 activations that a block-FP8 out_proj consumes directly, removing a
further dynamic_per_group_scaled_quant launch per GDN layer. This PR calls
it with quant_dtype=None. Measured against an out-of-tree reference
implementation that does use it, that leaves roughly 2 percentage points
of throughput on the table. It is deferred because no linear-attention backend
currently returns an (fp8, scales) tuple into a Linear, and it deserves its
own gate and its own review.

Reproduction

docker run --rm -it --device /dev/kfd --device /dev/dri --network host \
  --group-add video --ipc host --shm-size 32g \
  -v $HF_HOME:/hf lmsysorg/sglang-rocm:v0.5.21-rocm10-mi35x-20261012

# The image's AITER predates #5520; install one that contains it.
pip install -U "git+https://github.com/ROCm/aiter.git@main"   # >= dd6b8907

export SGLANG_USE_AITER=1
export SGLANG_QWEN3_NEXT_GDN_FUSED_BACKEND=aiter   # omit for the baseline arm

python3 -m sglang.launch_server \
  --model-path Qwen/Qwen3-Next-80B-A3B-Instruct-FP8 \
  --tensor-parallel-size 4 --ep-size 1 \
  --dtype bfloat16 --quantization fp8 --kv-cache-dtype fp8_e4m3 \
  --attention-backend aiter --moe-runner-backend aiter \
  --enable-aiter-allreduce-fusion \
  --disable-radix-cache --disable-shared-experts-fusion \
  --mem-fraction-static 0.80 \
  --context-length 8192 --chunked-prefill-size 16384 \
  --cuda-graph-max-bs-decode 64 \
  --port 8888 --trust-remote-code

Checklist

Off by default; unconditional fallback to the existing chain
pre-commit clean on all changed files
DCO sign-off on every commit
CUDA path explicitly unaffected (see Scope)
Registered test (test/srt/test_gdn_fused_decode_aiter.py, 9 tests)
Companion AITER PR merged (dd6b8907) and the minimum version stated
3x GSM8K per arm

CI States

Latest PR Test (Base): ❌ Run #39374448593
Latest PR Test (Extra): ❌ Run #39374448586
Latest PR Test (AMD ROCm 10): ❌ Run #39374448597

Replaces the four-kernel Qwen3-Next GDN decode chain (fused_qkvzba_split_reshape_cat -> causal_conv1d_update -> fused_recurrent_gated_delta_rule_packed_decode -> gated RMSNorm) with a single AITER Gluon launch that consumes the packed projections directly and folds the output gated RMSNorm into the recurrence. Off by default. Opt in with SGLANG_QWEN3_NEXT_GDN_FUSED_BACKEND=aiter. Follows the attempt-and-verify handoff Kimi-K3 already uses for its fused KDA decode (kimi_k3.py / kda_backend.py): the model offers the output-norm weight on the attention layer, the backend consumes it only if a covered kernel engages, and the model applies its own norm when the stash comes back unconsumed. The handoff is needed because the fusion spans the model/backend boundary -- the backend owns the conv and recurrent state pools but receives post-split mixed_qkv and never sees z or the norm weight, while the model has the reverse. Reuses the GDN backend's existing tuple input (projected_qkvz, projected_ba), which already exists for Qwen3.5's fused proj+Conv1D path; Qwen3-Next simply had not opted in. Everything is probed and nothing assumed: platform, env var, AITER import, Triton Gluon availability, and a per-call covered() contract delegated to AITER's own predicate so the two cannot drift, plus the one thing AITER cannot see -- the model's output-gate activation, which the kernel hard-codes to SiLU. A miss logs once and falls back rather than raising: a capability, not a mode. The CUDA path is unaffected: the fused path requires is_hip() and gfx950, and with the env var unset the probe returns False before any AITER import, leaving the existing chain byte-identical. Kernel side: ROCm/aiter fused_gdn_decode_qkvz, verified bit-identical to its upstream source on every output at M = 1, 32 and 256 on MI355X. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

The reject path already logged its reason once; the success path logged nothing, so an arm that silently fell back was indistinguishable from one that fused. Both directions are now observable from the server log, which is the structural evidence a perf claim needs. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

Covers SGLang's side of the contract -- the opt-in, the capability probe and the decline paths -- rather than the kernel's numerics, which AITER's own op tests own. The decline cases need neither a GPU nor AITER, which is the point: the fused path must degrade to the existing chain everywhere it is not supported, and that is exactly the behaviour CI can check on any runner. Included is the one condition AITER's predicate cannot see, the model's output-gate activation, since the kernel folds a SiLU gate and anything else would be silently wrong numerics. Also asserts every decline carries a non-empty reason, because the caller logs it once on the fallback path and an empty string there is useless. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

available() imported triton.experimental.gluon and stopped there. A Triton that ships the Gluon module but not gl.amd.cdna3 -- the one arch-specific primitive this kernel uses -- would import cleanly and then fail at first kernel compile with an opaque error, which is precisely the failure the gate exists to prevent. Now checks for gl.amd.cdna3.buffer_load as well. SGLang does not pin Triton, so this feature-probes rather than version-compares, following aiter_mla_gluon.py's inspect.signature idiom for the same reason. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

Wire the Qwen3-Next GDN (gated delta-net) prefill to AITER's fused single-kernel `fused_gdn_prefill_qkvz` behind the shared opt-in flag `SGLANG_QWEN3_NEXT_GDN_FUSED_BACKEND=aiter`. This is the prefill sibling of the decode consumer (sgl-project#39554): it replaces the four-kernel prefill chain (causal_conv1d_split_qkv -> fused gating -> chunked delta scan -> gated RMSNorm) with the fused kernel's tight Gluon launches, and folds the per-head group-128 FP8 quant into the epilogue so a block-FP8 out_proj consumes the activations directly. - adapter `kernels/ops/attention/gdn_fused_prefill_aiter.py`: fail-closed available()/covered()/run() -- platform, opt-in env, Triton>=3.8 Gluon probe, and a per-call shape/dtype contract delegated to AITER's predicate. Any miss runs the ordinary unfused chain. - GDNAttnBackend.try_fused_gdn_prefill: build the ragged-batch args from forward_metadata + the mamba conv/state pools, covered()->run(), stash the post-norm output and the group-128 FP8 activations; decline MIS / target-verify / tracked-conv batches. - RadixLinearAttention.try_fused_gdn_prefill: unwrap HybridLinearAttnBackend to the GDN backend; guard extend / not-verify / not-piecewise. - qwen3_next.py: static readiness gate (available + num_v_heads==2*num_k_heads + head_k==head_v==128 + conv==4), publish the bf16 norm weight after load, attempt the fused path first in forward; when out_proj is block-FP8 feed it the kernel's group-128 FP8 activations directly, else the bf16 normalized fallback. Requires the AITER fused prefill kernel (ROCm/aiter#5606). E2E (Qwen3-Next-80B FP8, MI355X, TP4): prefill-focused (ISL 6144/OSL 8) Median TTFT -27%, input throughput +35%; balanced 1k/1k TTFT -7%, decode-neutral (TPOT flat); GSM8K parity. Off by default; unsupported platform / uncovered call is unchanged. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

ROCm/aiter#5520 review widened fused_gdn_decode_qkvz_supported() to take A_log, dt_bias and norm_weight, which it previously never saw -- it could return covered for a launch whose parameter tensors were on the CPU or wrongly shaped, after which the kernel dereferenced them at fixed offsets. Threads those through covered() and the backend call site so the delegation stays complete, which is the point of delegating: the two contracts cannot drift if SGLang forwards everything AITER checks. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

yichiche · 2026-09-24T00:23:23Z

+        fused_gdn = (
+            self._gdn_fused_decode_ready and forward_batch.forward_mode.is_decode()
+        )
+        if fused_gdn:


Please add a guard with if_hip or if_rocm to make sure your change doesn't affect the NVIDIA common path. Use if_hip if this kernel can be used by all AMD GPUs; use use_aiter if this kernel imports from the AITER library.

@yichiche thanks for comment. I added _use_aiter = get_bool_env_var("SGLANG_USE_AITER") and is_hip() and gated both the forward-path branch and the constructor's readiness check on it. Used _use_aiter rather than is_hip since the kernel imports from AITER, per your guidance. The NVIDIA path is unaffected — with the flag false nothing is probed, allocated, or dispatched.

The kernel imports from AITER, so gate on the AITER opt-in as well as the platform. The NVIDIA path is unaffected. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

yichiche · 2026-09-30T01:26:16Z

@rbrugaro-amd Please resolve conflict issue.

Resolves the one conflict, in gdn_backend.py: upstream added target-verify qkv routing helpers (_init_target_verify_qkv_routing, _replayssm_fold_uses_cutedsl, _target_verify_supports_strided_qkv) at the same point in GDNAttnBackend where this branch adds _try_aiter_fused_gdn_decode. Both are kept; the additions are independent. AITER sgl-project#5520 landed upstream on 2026-10-06 (dd6b8907), so the fused decode op this branch consumes is now in aiter main unchanged. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

rbrugaro-amd · 2026-10-12T20:14:14Z

Rebased onto current main and refreshed every number. Description updated with the full tables; summary of what changed:

Conflict resolved. Merged main — no force-push, commit history preserved. The one conflict was in gdn_backend.py: upstream added target-verify qkv routing helpers at the same point in GDNAttnBackend where this PR adds _try_aiter_fused_gdn_decode. Both kept, since the additions are independent.

The AITER dependency is no longer pending. ROCm/aiter#5520 merged on 2026-10-06 (dd6b8907), so the description now pins a concrete minimum version instead of referencing an open PR.

@yichiche — on your earlier note about guarding the NVIDIA path: that is addressed in 40a88fd19e. _use_aiter = get_bool_env_var("SGLANG_USE_AITER") and is_hip() now gates both the forward-path branch and the constructor's readiness check. I used _use_aiter rather than is_hip per your guidance, since the kernel imports from AITER. With the flag unset nothing is probed, allocated or dispatched.

Re-validated end to end on a current nightly (lmsysorg/sglang-rocm:v0.5.21-rocm10-mi35x-20261012, sglang 2026-10-12). Both arms on the same image in the same session, differing by exactly this PR's commits — the baseline is this branch's own merge base rather than an older image default, so no unrelated upstream drift is folded in.

concurrency	1	8	32	64	128¹	256¹
throughput delta	+6.3%	+5.1%	+8.3%	+8.6%	+5.4%	+4.7%

¹ with --cuda-graph-max-bs-decode 256, which is what enables AITER's large-batch decomposition.

GSM8K (1319 questions): 0.943 → 0.946 at the default cap and 0.940 → 0.942 at the raised cap, Invalid: 0.000 in all four runs. The fused path logs its engage line exactly once per server with zero not covered fallbacks, so the measured arm really is running the kernel rather than quietly falling back.

These deltas are within a point of what this PR originally claimed, now carried across a ~1700-commit jump in main, a merged AITER dependency, and a move to ROCm 10 with native Triton 3.8.

Two things worth flagging for reviewers:

Absolute throughput is not comparable to the figures previously in this PR. Those came from a ROCm 7.0 image with Triton 3.8 installed manually; these come from ROCm 10 with native 3.8 — different compiler, different absolute scale. The arm-vs-arm deltas are the like-for-like comparison.
On Triton 3.7 images this feature is inert. Some ROCm builds ship a backported Gluon that imports cleanly but miscompiles the layout ops this kernel uses, so AITER gates on triton >= 3.8 rather than on an import probe. On those images the path declines and falls back cleanly rather than failing. Of the 2026-10-12 nightlies, rocm724 ships 3.7.0 and rocm10 ships 3.8.0.

The reproduction block in the description also gained three flags that were used for every measurement but had never been published (--disable-radix-cache, --disable-shared-experts-fusion, --mem-fraction-static 0.80), and the GSM8K command changed, since benchmark/gsm8k/bench_sglang.py no longer exists in main — it is python3 -m sglang.test.few_shot_gsm8k now.

Ready for another look.

The test was added at test/srt/test_gdn_fused_decode_aiter.py, which no longer reaches CI: run_suite.py discovers tests by globbing test/registered/**/*.py, and test/srt is referenced by no workflow -- only by the stale-PR notice that documents the move. The directory is down to four pre-migration stragglers against 2260 files under test/registered, so the file was collected by nothing and the suite never ran it. Move it to test/registered/attention/, alongside the other GDN and KDA attention tests, and add the registration marker that discovery requires; the glob is strict, so a file placed there without one fails collection rather than being skipped. Registered for stage-b-test-1-gpu-small-amd-mi35x: the op is gfx950-only, so mi35x is the suite where the device-dependent class runs instead of skipping. No CUDA registration -- the fused path is unreachable off ROCm. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

rbrugaro-amd · 2026-10-13T03:48:17Z

@yichiche The registered test was not actually reaching CI, and now is. (addressed in commit 5d220d9) It had been added under test/srt/, which run_suite.py no longer discovers — collection globs test/registered/**/*.py, and test/srt is referenced by no workflow. The directory is down to four pre-migration files against ~2260 under test/registered/, so nothing collected ours and the suite silently never ran it. Moved to test/registered/attention/ alongside the other GDN and KDA attention tests, with the registration marker discovery requires (the glob is strict, so a file placed there without one fails collection rather than being skipped). Registered for stage-b-test-1-gpu-small-amd-mi35x: the op is gfx950-only, so that is the suite where the device-dependent class runs instead of skipping, and there is no CUDA registration because the fused path is unreachable off ROCm. Verified via collect_tests(..., sanity_check=True) that it is now picked up, and the 9 tests pass from the new path.

Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

Resolves one conflict in gdn_backend.py's imports. Upstream's kernels regroup (sgl-project#45112) moved fused_gdn_gating from ops/attention/fla/ to ops/attention/linear/; this branch's adjacent line adds the gdn_fused_decode_aiter import. Took upstream's relocated path and kept the adapter import. The refactor did not move ops/attention/gdn_fused_decode_aiter.py, and the adapter imports nothing from the regrouped tree. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

rbrugaro-amd · 2026-10-20T05:24:56Z

@yichiche can you please take another look? the aiter dependent PR ROCm/aiter#5520 merged 2 weeks ago

rbrugaro-amd added 3 commits September 13, 2026 05:28

github-actions Bot added the jit-kernel label Sep 15, 2026

karverma-amd mentioned this pull request Sep 17, 2026

[Perf] [AMD] qwen3-next: Consume the AITER fused GDN prefill kernel #39873

Open

rbrugaro-amd marked this pull request as ready for review September 17, 2026 23:09

rbrugaro-amd requested review from BBuf, DarkSharpness, Fridge003, HaiShaw, HydraQYH, Qiaolin-Yu, celve, hebiao064, ispobock, merrymercy and yuan-luo as code owners September 17, 2026 23:09

kkHuang-amd added amd run-ci CI: run the baseline test suite on this PR labels Sep 18, 2026

rbrugaro-amd requested review from hanming-lu and yizhang2077 as code owners September 22, 2026 17:57

rbrugaro-amd force-pushed the dev/gdn-fused-decode branch from 625f185 to d969914 Compare September 22, 2026 19:19

yichiche requested changes Sep 24, 2026

View reviewed changes

[AMD] Guard the fused GDN decode path on _use_aiter

40a88fd

The kernel imports from AITER, so gate on the AITER opt-in as well as the platform. The NVIDIA path is unaffected. Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

weitliao mentioned this pull request Sep 28, 2026

[AMD] Qwen3-Next: fused QK Gemma RMSNorm + partial RoPE on gfx950 #41519

Closed

5 of 8 tasks

Merge remote-tracking branch 'origin/main' into dev/gdn-fused-decode

c271409

This comment has been minimized.

Sign in to view

rbrugaro-amd changed the title ~~[AMD] Qwen3-Next: route GDN decode at AITER's fused gfx950 kernel~~ [Perf] [AMD] Qwen3-Next: route GDN decode at AITER's fused gfx950 kernel Oct 19, 2026

rbrugaro-amd added 2 commits October 19, 2026 23:02

Merge remote-tracking branch 'origin/main' into dev/gdn-fused-decode

1a55642

Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

Navigation Menu

Repository navigation

[Perf] [AMD] Qwen3-Next: route GDN decode at AITER's fused gfx950 kernel - #39554

[Perf] [AMD] Qwen3-Next: route GDN decode at AITER's fused gfx950 kernel#39554
rbrugaro-amd wants to merge 11 commits into
sgl-project:mainfrom
rbrugaro-amd:dev/gdn-fused-decode

rbrugaro-amd commented Sep 15, 2026 •

edited by github-actions Bot

Loading

Uh oh!

yichiche Sep 24, 2026

Uh oh!

rbrugaro-amd Sep 24, 2026 •

edited

Loading

Uh oh!

yichiche commented Sep 30, 2026

Uh oh!

rbrugaro-amd commented Oct 12, 2026

Uh oh!

rbrugaro-amd commented Oct 13, 2026

Uh oh!

This comment has been minimized.

rbrugaro-amd commented Oct 20, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

rbrugaro-amd commented Sep 15, 2026 • edited by github-actions Bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

[AMD] Qwen3-Next: route GDN decode at AITER's fused gfx950 kernel

Motivation

Modifications

Scope

A note on Triton 3.7 images

Environment

Accuracy Tests

Speed Tests

Raised graph cap

Kernel level

Methodology notes

Known gaps

Reproduction

Checklist

CI States

Uh oh!

yichiche Sep 24, 2026

Choose a reason for hiding this comment

Uh oh!

rbrugaro-amd Sep 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

yichiche commented Sep 30, 2026

Uh oh!

rbrugaro-amd commented Oct 12, 2026

Uh oh!

rbrugaro-amd commented Oct 13, 2026

Uh oh!

This comment has been minimized.

rbrugaro-amd commented Oct 20, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

rbrugaro-amd commented Sep 15, 2026 •

edited by github-actions Bot

Loading

rbrugaro-amd Sep 24, 2026 •

edited

Loading