Skip to content

[HIP] [Perf] Fold shared-expert gate GEMV into topk_softmax kernel - #5652

Merged
yzhou103 merged 10 commits into
ROCm:mainfrom
Emmanuel0612:opt/folding_fused_shared_expert_softmax_topk
Sep 30, 2026
Merged

yzhou103 merged 10 commits into
ROCm:mainfrom
Emmanuel0612:opt/folding_fused_shared_expert_softmax_topk

Conversation

@Emmanuel0612

@Emmanuel0612 Emmanuel0612 commented Sep 18, 2026 •

Copy link
Copy Markdown
Contributor

Add fuse-gate mode to topk_softmax so the shared-expert sigmoid gate is computed inside the routed softmax top-k kernel, exposing the new hidden_states/gate_weight/shared_expert_scale/shared_expert_base args through the pybind registration. Adds an op_test with a perf+correctness sweep.

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Add fuse-gate mode to topk_softmax so the shared-expert sigmoid gate is
computed inside the routed softmax top-k kernel, exposing the new
hidden_states/gate_weight/shared_expert_scale/shared_expert_base args
through the pybind registration. Adds an op_test with a perf+correctness
sweep.

Co-Authored-By: Claude <noreply@anthropic.com>
@Emmanuel0612
Emmanuel0612 requested a review from a team September 18, 2026 02:08
@github-actions

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

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

Extended tests (opt-in via labels):

Label Tests
ci:gfx1250-ffm-triton Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
multigpu Aiter multi-GPU tests on the 8-GPU runner
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: 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 5652 --add-label <label>

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

@github-actions github-actions Bot changed the title [Perf] Fold shared-expert gate GEMV into topk_softmax kernel [HIP] [Perf] Fold shared-expert gate GEMV into topk_softmax kernel Sep 18, 2026
@github-actions github-actions Bot added the HIP label Sep 18, 2026
Emmanuel0612 and others added 2 commits September 18, 2026 11:08
Satisfies the pinned ruff 0.16.0 (I001, UP045) and black CI style jobs;
no functional change.

Co-Authored-By: Claude <noreply@anthropic.com>
…perts)

The gate-staging __shared__ buffer was sized NUM_SHARED_EXPERTS * 4096 *
sizeof(DTYPE); for fp32 with 8 shared experts that is 128KB, exceeding the
64KB LDS limit and failing the PREBUILD_KERNELS build (only surfaces there,
not in JIT op_tests). Gate the LDS staging on a compile-time fit check: when
the buffer would exceed 64KB, size it to 1 and read the gate from global
instead. Result is identical; only the caching optimization is skipped for
that combo. Qwen (bf16, 1 shared expert) still stages into LDS unchanged.

Co-Authored-By: Claude <noreply@anthropic.com>
@zufayu
zufayu requested review from yzhou103 and zufayu September 20, 2026 01:11
@yzhou103

yzhou103 commented Sep 21, 2026 •

Copy link
Copy Markdown
Contributor

Hi I'd like to understand the baseline this is being compared against, because downstream the shared-expert gate already appears to be folded into the router GEMM.
In Qwen3-Next — a gated shared expert with sigmoid scoring, i.e. exactly this PR's target
shape — the router projection is a merged linear:

self.gate = MergedReplicatedLinear(
    config.hidden_size,
    [config.num_experts, config.n_shared_experts],   # routed | shared, one weight
    bias=False, quant_config=None, ...)
...
logits = self.gate(hidden_states)   # single GEMM -> [M, E + S]
router_logits = logits              # trailing columns are the shared gate logit
  1. Which model/config can't use the merged-gate approach, and why? The gate is
    unquantized bf16 (quant_config=None), so concatenating the [S, H] rows at load time
    seems straightforward, and the infrastructure is already generic.
  2. Could you share measurements against the merged-gate baseline — rather than against
    a standalone GEMV — over a realistic range of M? The test added here only reports the
    fused timing, with no baseline column.

If there's no case the merged gate can't cover, I'd rather not add a second way to do this:
it's a new API mode, a new weight-plumbing path, and it moves work off the matrix cores.

Happy to go through the kernel-level findings (static LDS sizing, vector-load alignment vs.
the tail loop, missing validation) once the motivation is settled — but they're moot if the
merged gate already covers thi

Address review on ROCm#5652: instead of extending topk_softmax (which added an
unconditional s_gate LDS buffer that penalized existing/legacy shared-expert
callers), add a standalone op topk_softmax_fused_shared_gate with its own
kernel + dispatch. topk_softmax and its kernel are reverted to base verbatim,
so non-fuse-gate callers carry no LDS penalty.

- New: topkGatingSoftmaxFusedGate kernel + launcher + LAUNCH_SOFTMAX_FUSED*
  macros + applySharedExpertSigmoidFusedGate + dispatch + host op
  topk_softmax_fused_shared_gate (+ moe_op.h decl, rocm_ops.hpp pybind,
  moe_op.py wrapper).
- op_test: separate fused / unfused unit benchmarks for fused-vs-not-fused perf.
- Guards: require contiguous hidden_states/gate_weight; meaningful index-width
  check; fix the scale-wording comment (sigmoid, then scale).

Co-Authored-By: Claude <noreply@anthropic.com>
@Emmanuel0612

Copy link
Copy Markdown
Contributor Author

Hi I'd like to understand the baseline this is being compared against, because downstream the shared-expert gate already appears to be folded into the router GEMM. In Qwen3-Next — a gated shared expert with sigmoid scoring, i.e. exactly this PR's target shape — the router projection is a merged linear:

self.gate = MergedReplicatedLinear(
    config.hidden_size,
    [config.num_experts, config.n_shared_experts],   # routed | shared, one weight
    bias=False, quant_config=None, ...)
...
logits = self.gate(hidden_states)   # single GEMM -> [M, E + S]
router_logits = logits              # trailing columns are the shared gate logit
  1. Which model/config can't use the merged-gate approach, and why? The gate is
    unquantized bf16 (quant_config=None), so concatenating the [S, H] rows at load time
    seems straightforward, and the infrastructure is already generic.
  2. Could you share measurements against the merged-gate baseline — rather than against
    a standalone GEMV — over a realistic range of M? The test added here only reports the
    fused timing, with no baseline column.

If there's no case the merged gate can't cover, I'd rather not add a second way to do this: it's a new API mode, a new weight-plumbing path, and it moves work off the matrix cores.

Happy to go through the kernel-level findings (static LDS sizing, vector-load alignment vs. the tail loop, missing validation) once the motivation is settled — but they're moot if the merged gate already covers thi

Hi, thank you for the careful read.
Let me explain why this PR exist.

  1. Our target is Qwen3.5-397B-A17B-MXFP4.
    The router gate is MXFP4 (fp4) but shared_expert_gate is in the quant exclude list (bf16, all 60 layers), so they can't share one GEMM without re-quantizing the bf16 gate to fp4 and changing calibrated numerics.

  2. Since the merged gate isn't available for this recipe, the fair comparison is the separate bf16 GEMV + append. Please have a check of following tables.
    without folding:

image with folding: image
  1. We also restructured per your first note: this is now a SEPARATE op
    (topk_softmax_fused_shared_gate) with its own kernel. The existing
    topk_softmax and its kernel are reverted to base verbatim, so non-fuse-gate
    callers carry NO shared-gate LDS allocation (the earlier static s_gate penalty
    is gone from the shared kernel).

@yzhou103

Copy link
Copy Markdown
Contributor

Thanks! please look at these:
2. Alignment is still unchecked (UB). is_contiguous() actually makes this concrete:
with a contiguous hidden_states, hs_row = base + row * H, so any H not a multiple of
ELTS_PER_LDG makes every row after the first misaligned for the 64B AccessType load
(same for gw_row at s >= 1). Please either AITER_CHECK(H * sizeof(DTYPE) % BYTES_PER_LDG == 0)
and drop the tail loop, or make the tail path actually handle misalignment.

  1. Empty shared_expert_scoring_func is still a silent no-op. The check is
    if(!empty()) AITER_CHECK(== "sigmoid"), so "" passes, SCORING_FUNC resolves to NONE,
    and no shared weights or ids are written at all — uninitialized output, no error. For this
    op the scoring function is mandatory; please require it (and num_shared_experts > 0).

  2. The duplication has already diverged. The header says "Bodies are verbatim copies
    with renamed identifiers", and in the copy the fuse branch got float score = 0.0f; while
    the legacy branch still has the uninitialized float score;. That legacy branch is also
    dead code here — the host always passes a non-null gate_weight — so it's an unreachable
    path carrying an unfixed bug. Suggest deleting the legacy branch (plus the now-redundant
    shared_gate_weight != nullptr and shared_expert_base >= 0 runtime checks) from the new
    kernel, and factoring the shared top-k body (argmax + multithread_reduce + clear-to--INFINITY

  • renorm) into a device helper in a header so the two copies can't drift further.
  1. token_expert_indices width is now documented incorrectly. The header says
    [num_tokens, topk + num_shared_experts], but the kernel writes
    source_rows[k * thread_row + k_idx] with a hardcoded stride of topk. With a
    topk + S wide buffer the writes land in the wrong columns. Either thread a stride
    through, or document it as [num_tokens, topk].

  2. Still open from the previous round:

  • indices[...] = NUM_EXPERTS (the out-of-range marker) equals shared_expert_base in the
    intended usage, so a filtered routed slot reads back as shared expert 0.
  • applySharedExpertSigmoidFusedGate still does a serial for(h = 0; h < H; ++h) per
    thread on the non-power-of-2 path. Given this op requires a gate, I'd just
    AITER_CHECK that path unsupported rather than ship a GEMV that's slower than the
    unfused version.
  • Test coverage is still E=512 / H=4096 / S in {1,2} — the risky configs (S=8,
    H > 4096 where the LDS staging silently falls back, H not divisible by
    ELTS_PER_LDG, fp32) are all untested.
  • The file is named test_.py but contains only bench_ functions, so pytest collects
    it and runs nothing. Either rename to bench_*.py or add a real parametrized test.

- Alignment (UB): host op AITER_CHECKs 64B-aligned row stride; drop dead tail loop
- Mandatory scoring: require num_shared_experts > 0 and sigmoid (was silent no-op)
- Fuse-only kernel: remove dead legacy branch + redundant null/base guards
- token_expert_indices: size [num_tokens, topk] to match the kernel's stride
- Marker: use NUM_EXPERTS + NUM_SHARED_EXPERTS so a filtered slot can't alias shared expert 0
- Non-power-of-2 experts: reject via AITER_CHECK; delete slow serial GEMV helper
- Tests: real pytest test_correctness (S {1,2,8}, fp32, H=8192) + alignment/non-pow2 negatives
- Bad input now raises RuntimeError instead of abort()

Co-Authored-By: Claude <noreply@anthropic.com>
@Emmanuel0612

Copy link
Copy Markdown
Contributor Author

Thanks! please look at these: 2. Alignment is still unchecked (UB). is_contiguous() actually makes this concrete: with a contiguous hidden_states, hs_row = base + row * H, so any H not a multiple of ELTS_PER_LDG makes every row after the first misaligned for the 64B AccessType load (same for gw_row at s >= 1). Please either AITER_CHECK(H * sizeof(DTYPE) % BYTES_PER_LDG == 0) and drop the tail loop, or make the tail path actually handle misalignment.

  1. Empty shared_expert_scoring_func is still a silent no-op. The check is
    if(!empty()) AITER_CHECK(== "sigmoid"), so "" passes, SCORING_FUNC resolves to NONE,
    and no shared weights or ids are written at all — uninitialized output, no error. For this
    op the scoring function is mandatory; please require it (and num_shared_experts > 0).
  2. The duplication has already diverged. The header says "Bodies are verbatim copies
    with renamed identifiers", and in the copy the fuse branch got float score = 0.0f; while
    the legacy branch still has the uninitialized float score;. That legacy branch is also
    dead code here — the host always passes a non-null gate_weight — so it's an unreachable
    path carrying an unfixed bug. Suggest deleting the legacy branch (plus the now-redundant
    shared_gate_weight != nullptr and shared_expert_base >= 0 runtime checks) from the new
    kernel, and factoring the shared top-k body (argmax + multithread_reduce + clear-to--INFINITY
  • renorm) into a device helper in a header so the two copies can't drift further.
  1. token_expert_indices width is now documented incorrectly. The header says
    [num_tokens, topk + num_shared_experts], but the kernel writes
    source_rows[k * thread_row + k_idx] with a hardcoded stride of topk. With a
    topk + S wide buffer the writes land in the wrong columns. Either thread a stride
    through, or document it as [num_tokens, topk].
  2. Still open from the previous round:
  • indices[...] = NUM_EXPERTS (the out-of-range marker) equals shared_expert_base in the
    intended usage, so a filtered routed slot reads back as shared expert 0.
  • applySharedExpertSigmoidFusedGate still does a serial for(h = 0; h < H; ++h) per
    thread on the non-power-of-2 path. Given this op requires a gate, I'd just
    AITER_CHECK that path unsupported rather than ship a GEMV that's slower than the
    unfused version.
  • Test coverage is still E=512 / H=4096 / S in {1,2} — the risky configs (S=8,
    H > 4096 where the LDS staging silently falls back, H not divisible by
    ELTS_PER_LDG, fp32) are all untested.
  • The file is named test__.py but contains only bench__ functions, so pytest collects
    it and runs nothing. Either rename to bench_*.py or add a real parametrized test.

Thanks for the second review. Every point is fixed in the new op and kernel. Here is what changed,

  1. The host op now checks that each row's size in bytes (hidden size times the data-type size) is a multiple of 64. If it is not, the op raises an error. Because of this check, the hidden size is always an exact multiple of the vector width, so we also removed the old scalar "tail" loop — with the check in place it could never run.

  2. The op now always requires two things — at least one shared expert (num_shared_experts > 0) and scoring_func == "sigmoid". If either is missing, it raises an error instead of returning garbage.

  3. We removed the dead "legacy" branch (the one with the uninitialized score) and the two runtime guards that could never be false (the null-gate check and the shared_expert_base >= 0 check). The new kernel now has only the fused path. About your suggestion to move the shared top-k code into one helper that both kernels include: we chose to keep the two kernels separate instead. The reason is the design goal from round 1 — the original topk_softmax kernel stays exactly the same as upstream, so every existing caller that does not use the shared gate keeps working with zero risk.

  4. This buffer is scratch space: the op writes it, but nothing reads it back. The kernel fills it with a fixed stride of topk (routed slots only). So we now size it [num_tokens, topk] instead of [num_tokens, topk + num_shared], and fixed the header comment to match.

  • The routed code marks a "skip this slot" entry with a special number. That number used to be num_experts, which is exactly the id of the first shared expert — so a skipped slot could be read back as shared expert 0. We changed the marker to num_experts + num_shared_experts, which is outside the range of all real ids, so it can no longer be confused with any expert.
  • That path used a slow, one-element-at-a-time GEMV that was actually slower than not fusing at all. Instead of shipping it, the op now rejects non-power-of-2 expert counts with a clear error, and we deleted that slow code.
  • The test file used to contain only benchmark functions, so pytest collected the file but ran no real test. We added a real test (test_correctness) that compares the fused op and the unfused baseline against a reference for several shapes: 1, 2, and 8 shared experts; float32 inputs; and a larger hidden size that triggers the "read the gate from global memory" fallback. We also added two tests that confirm the new checks actually fire — one for a misaligned hidden size, one for a non-power-of-2 expert count.

@yzhou103

yzhou103 commented Sep 24, 2026 •

Copy link
Copy Markdown
Contributor
  1. bench_unfused is not a representative baseline, so the reported speedup is inflated.**

The timed region is ~9 kernels, and most of them are artifacts of the torch plumbing
rather than of "not fusing":

topk_softmax(r_w, r_i, r_tei, gating, renorm)                            # 1  necessary
shared_w = torch.sigmoid(hs.float() @ gate_weight.float().t()) * scale   # 2 upcasts + mm + sigmoid + mul
w[:, :topk] = r_w;   w[:, topk:] = shared_w                              # 2 copies
ids[:, :topk] = r_i; ids[:, topk:] = shared_ids                          # 2 copies
  • The two .float() casts materialize an [M, H] fp32 temporary (128MB at M=8192,
    H=4096). A bf16 F.linear already accumulates in fp32 on MFMA, which is what the fused
    kernel does — so the upcast also makes the baseline more accurate than the op under
    test, which is the wrong way round.
  • The two routed copies are avoidable. The kernel indexes its outputs through
    topk_weights.stride(0) / topk_indices.stride(0), so w[:, :topk] and ids[:, :topk]
    can be passed directly as views and the routed columns get written in place. That
    strided-output path is already exercised in-repo by op_tests/test_moeTopkSoftmax.py,
    which calls the base op with torch.empty_strided buffers.
  • shared_ids depends only on (num_routing_experts, num_shared_experts), not on the
    input, so it belongs outside the timed region.

Suggested rewrite (~4 kernels, and the sigmoid/scale/store could be folded further):

w = torch.empty(tokens, total, dtype=dtypes.fp32)
ids = torch.empty(tokens, total, dtype=dtypes.i32)
tei = torch.empty(tokens, topk, dtype=dtypes.i32)
w_routed, ids_routed = w[:, :topk], ids[:, :topk] # views, written in place
ids[:, topk:] = base + torch.arange(num_shared, dtype=dtypes.i32) # input-independent
logit = torch.empty(tokens, num_shared, dtype=dtype) # reused, not reallocated

def fn():
topk_softmax(w_routed, ids_routed, tei, gating, renorm)
torch.nn.functional.linear(hs, gate_weight, out=logit) # bf16, no upcast
torch.sigmoid(logit, out=logit)
w[:, topk:] = logit * scale

  1. The M sweep is decode-only. --tokens defaults to [1, 4, 17, 64]. The in-kernel
    GEMV runs on VALU with THREADS_PER_ROW = 16 lanes per row and adds a full [M, H] HBM
    pass, so the interesting question is at the other end. Please extend to
    [1, 4, 16, 64, 256, 1024, 4096, 16384] so we can see whether there's a crossover, and
    where it is.

  2. Missing axes in test_correctness. It currently pins tokens=64, num_experts=512, topk=10, scale=1.0, renorm=True, so several branches are untested:

  • need_renorm=False — a separate epilogue, currently never executed under pytest.
  • shared_expert_scale != 1.0 — applied in the kernel; would pass unnoticed if dropped.
  • num_experts=256 — at E < 512 the launcher picks BYTES_PER_LDG = 32, i.e. a
    different ELTS_PER_LDG and a different vector width in the GEMV. Only the 64B path is
    covered today.
  • Token-count edge cases: M = 1, and an M that isn't a multiple of ROWS_PER_CTA
    (8 at E=512) such as M = 65, to cover the tail block. M = 0 currently yields
    num_blocks == 0 and a failed launch — worth deciding whether that should be a guarded
    no-op.

Comment thread aiter/ops/moe_op.py Outdated
- Make the unfused benchmark a fair baseline: write routed columns in place
  via strided views, run the gate GEMV in bf16 (no fp32 upcast), and precompute
  the input-independent shared ids outside the timed region.
- Extend the token sweep to [1, 4, 16, 64, 256, 1024, 4096, 16384] to expose
  the decode->prefill crossover.
- Add the missing test_correctness axes (need_renorm=False, scale!=1.0,
  num_experts=256, M=1/M=65 tail), plus a guarded M=0 no-op and its test.
- Drop the dead softmax_workspace parameter from the fused op (C++ entry,
  moe_op.h, pybind, wrapper); the base topk_softmax keeps it.

Co-Authored-By: Claude <noreply@anthropic.com>
@yzhou103

Copy link
Copy Markdown
Contributor

please update your forked aiter

@Emmanuel0612

Copy link
Copy Markdown
Contributor Author
  1. bench_unfused is not a representative baseline, so the reported speedup is inflated.**

The timed region is ~9 kernels, and most of them are artifacts of the torch plumbing rather than of "not fusing":

topk_softmax(r_w, r_i, r_tei, gating, renorm)                            # 1  necessary
shared_w = torch.sigmoid(hs.float() @ gate_weight.float().t()) * scale   # 2 upcasts + mm + sigmoid + mul
w[:, :topk] = r_w;   w[:, topk:] = shared_w                              # 2 copies
ids[:, :topk] = r_i; ids[:, topk:] = shared_ids                          # 2 copies
  • The two .float() casts materialize an [M, H] fp32 temporary (128MB at M=8192,
    H=4096). A bf16 F.linear already accumulates in fp32 on MFMA, which is what the fused
    kernel does — so the upcast also makes the baseline more accurate than the op under
    test, which is the wrong way round.
  • The two routed copies are avoidable. The kernel indexes its outputs through
    topk_weights.stride(0) / topk_indices.stride(0), so w[:, :topk] and ids[:, :topk]
    can be passed directly as views and the routed columns get written in place. That
    strided-output path is already exercised in-repo by op_tests/test_moeTopkSoftmax.py,
    which calls the base op with torch.empty_strided buffers.
  • shared_ids depends only on (num_routing_experts, num_shared_experts), not on the
    input, so it belongs outside the timed region.

Suggested rewrite (~4 kernels, and the sigmoid/scale/store could be folded further):

w = torch.empty(tokens, total, dtype=dtypes.fp32) ids = torch.empty(tokens, total, dtype=dtypes.i32) tei = torch.empty(tokens, topk, dtype=dtypes.i32) w_routed, ids_routed = w[:, :topk], ids[:, :topk] # views, written in place ids[:, topk:] = base + torch.arange(num_shared, dtype=dtypes.i32) # input-independent logit = torch.empty(tokens, num_shared, dtype=dtype) # reused, not reallocated

def fn(): topk_softmax(w_routed, ids_routed, tei, gating, renorm) torch.nn.functional.linear(hs, gate_weight, out=logit) # bf16, no upcast torch.sigmoid(logit, out=logit) w[:, topk:] = logit * scale

  1. The M sweep is decode-only. --tokens defaults to [1, 4, 17, 64]. The in-kernel
    GEMV runs on VALU with THREADS_PER_ROW = 16 lanes per row and adds a full [M, H] HBM
    pass, so the interesting question is at the other end. Please extend to
    [1, 4, 16, 64, 256, 1024, 4096, 16384] so we can see whether there's a crossover, and
    where it is.
  2. Missing axes in test_correctness. It currently pins tokens=64, num_experts=512, topk=10, scale=1.0, renorm=True, so several branches are untested:
  • need_renorm=False — a separate epilogue, currently never executed under pytest.
  • shared_expert_scale != 1.0 — applied in the kernel; would pass unnoticed if dropped.
  • num_experts=256 — at E < 512 the launcher picks BYTES_PER_LDG = 32, i.e. a
    different ELTS_PER_LDG and a different vector width in the GEMV. Only the 64B path is
    covered today.
  • Token-count edge cases: M = 1, and an M that isn't a multiple of ROWS_PER_CTA
    (8 at E=512) such as M = 65, to cover the tail block. M = 0 currently yields
    num_blocks == 0 and a failed launch — worth deciding whether that should be a guarded
    no-op.

Thanks for the detailed review. All 3 points are fixed and re-tested on
gfx950 (build is clean, every benchmark row shows err == 0, and the pytest suite
is 13 passed). Details below, point by point.

1) The unfused baseline was not a fair comparison

It was true: the old bench_unfused was doing extra work that had nothing to
do with "not fusing," so it made the speedup look bigger than it really is. We
rewrote it so it does the same work the fused kernel does, and nothing more:

  • The routed top-k now writes straight into the wide output through views
    (w[:, :topk] and ids[:, :topk]). Before, it wrote into separate buffers
    and then copied them over — those two copies are gone.
  • The shared-gate GEMV now runs in bf16 with no up-cast. Before, the code did
    hs.float() @ gate_weight.float().t(), which built a big [M, hidden] fp32
    temporary AND (as you noted) made the baseline more accurate than the fused
    kernel, which is backwards. A bf16 matmul already accumulates in fp32 on the
    matrix cores, so this now matches the fused op.
  • The shared-expert ids depend only on (base, num_shared), not on the input,
    so they are computed once, outside the timed region.

Small note: your snippet used torch.nn.functional.linear(..., out=logit), but
F.linear has no out= argument, so we used torch.mm(hs, gate_weight.t(), out=logit) instead. Same math (hs @ gate_weight.T), same bf16, no up-cast.

Effect: the unfused baseline got faster (roughly 45 us -> 28 us at M = 64), so
the reported speedup is smaller now — but it is honest. err is still 0.

2) The token (M) sweep was decode-only

Agreed. The default sweep now runs
tokens = [1, 4, 16, 64, 256, 1024, 4096, 16384], so we can see decode all the
way up to a large prefill batch.

What the data shows (Qwen shape: experts=512, hidden=4096, top_k=10,
num_shared=1, bf16):

  • The fused op is faster across the WHOLE range — there is no crossover up to
    M = 16384.
  • The gap does shrink as M grows, exactly for the reason you gave (the
    in-kernel GEMV adds a full [M, hidden] pass over memory). In the stable
    large-M rows it goes from about 1.9x at M = 4096 to about 1.5x at M = 16384.
  • The fused op stays ahead even at large M because it also skips a second
    kernel launch and avoids writing the shared logits out to memory and reading
    them back (the unfused path does both).

*** PROVISIONAL — single run, REPLACE with medians before posting ***
(num_shared=1, renorm=True, bf16; us = microseconds, lower is better)

M unfused us fused us fused speedup
1 46.1 14.6 3.2x
4 22.0 13.9 1.6x
16 35.5 15.2 2.3x
64 28.7 14.8 1.9x
256 52.9 15.8 3.4x
1024 39.3 16.6 2.4x
4096 44.5 23.5 1.9x
16384 65.6 44.2 1.5x

The rows marked "unfused noisy" bounced between runs; treat the exact ratios as
provisional. The steady story is the large-M end (M = 4096 -> 16384: 1.9x ->
1.5x), which shows the gap narrowing but never crossing.

One honesty note for our own numbers: the UNFUSED column is noisy run-to-run at
small and mid M (for example M = 1 measured 46 us in one run and 22 us in
another). The fused column is stable. Before we quote exact ratios we will re-run
the sweep a few times and post medians, so a single noisy row does not
mis-state the result. The trend above (fused wins, gap narrows with M) is stable
across runs.

3) Missing axes in test_correctness

The test was pinned at tokens=64, experts=512, scale=1.0, renorm=True, which left
several code paths untested. We added a case for each one you listed (each case
changes just one thing from the baseline, so a failure points right at the path
it exercises):

  • need_renorm = False -> exercises the separate normalization epilogue.
  • shared_expert_scale != 1.0 (0.5) -> would catch it if the scale were dropped.
  • num_experts = 256 -> at E < 512 the launcher picks a smaller vector width
    (BYTES_PER_LDG = 32) for the GEMV, which was never covered before.
  • M = 1 and M = 65 -> M = 65 is not a multiple of ROWS_PER_CTA (8 at E=512),
    so it exercises the leftover "tail" block.

On M = 0: you asked whether it should be a guarded no-op. We decided yes. An
empty batch would otherwise launch a 0-block grid and fail. The host op now
returns early when num_tokens == 0 (after the input checks still run), and there
is a test (test_zero_tokens_noop) that confirms it. If you would rather this
raise an error instead, that is a one-line change — just let us know.

All of these now run under pytest and pass.
image

Comment thread op_tests/test_topk_softmax_shared_gate.py Outdated
Emmanuel0612 and others added 3 commits September 29, 2026 15:40
Follow the aiter op_test standard (op_tests/test_quant.py):
- Merge bench_unfused/bench_fused into one @benchmark fn with an
  {unfused, fused} candidates dict, so both paths are timed and compared
  in a single markdown table.
- Correctness is the per-candidate err (checkAllclose vs a torch reference),
  now asserted (err == 0) so it runs AND gates under `python3 <file>` --
  the way CI invokes op_tests. The previous pytest test_* functions were
  never executed in CI; they are removed.
- Fold the extra correctness axes (num_experts=256, hidden=8192, fp32,
  num_shared=8, M=65) and the negative/guard behaviors (misaligned raise,
  non-power-of-2 raise, M=0 no-op) into main() so they run in CI too.

Co-Authored-By: Claude <noreply@anthropic.com>
Put first-party aiter imports in their own group after third-party, so
ruff's import-order check passes.

Co-Authored-By: Claude <noreply@anthropic.com>
Other aiter op_tests (e.g. op_tests/test_quant.py) record the per-candidate
err in the table and let checkAllclose log mismatches; they do not assert on
numerical err. Follow that convention: correctness still runs under
`python3 <file>` (CI) and the err is reported in the table and logs, but a
numerical mismatch no longer hard-fails via exit code. The guard asserts
(misaligned / non-power-of-2 raise, M=0 no-op) stay, since they check
raise/no-op behavior.

Co-Authored-By: Claude <noreply@anthropic.com>

@yzhou103 yzhou103 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

@yzhou103
yzhou103 merged commit 6f6d66a into ROCm:main Sep 30, 2026
56 checks passed
@Emmanuel0612
Emmanuel0612 deleted the opt/folding_fused_shared_expert_softmax_topk branch September 30, 2026 02:29
valarLip added a commit that referenced this pull request Oct 4, 2026
… experts (#6125)

Add tuned rows for DeepSeek-V4.1-Flash routed experts sharded across TP
(model_dim 5120, per-rank inter 640 at TP4 / 1152 at TP2; target 384x6 and
DSpark draft 128x3; tokens 1..32768) with MXFP8 activations on interleaved
gate/up rows (A8W4), the precision the checkpoint's reference runs.

The default tuner sweep (TUNE_ONLY=flydslv2) only generates the fused-fp8
stage1 + moe2_layout family. At 32-64 tokens that family loses to the
heuristic's bf16-out `_gui` stage1 + non-layout stage2 by up to 1.25x, and a
v2-only table raised V4.1 TP4 decode TPOT from 8.2 ms to 12.8 ms. The rows here
take, per shape, the faster of a v2 sweep and a TUNE_ONLY=flydsl sweep (v1
family). Each row was then checked against the heuristic on the production
fused_moe op. The 7 shapes where the tuned pick measured >3% slower there carry
the heuristic's own kernel pair instead; their us is the measured op time and
us1/us2/err are left empty, since the tuner never timed that pair.

End to end (V4.1 TP4, DSpark-5, ISL/OSL 1024/512, conc 8): 2597 vs 2552 tok/s
for the heuristic, within run-to-run spread.

Also fix a build error from #5652 in topk_softmax_kernels.cu:
`__shared__ alignas(64) T x[]` is rejected by ROCm 7.1.1's clang
("'alignas' attribute cannot be applied to types"), which breaks
module_moe_asm and with it the MoE tuner. Use
`__attribute__((aligned(64)))`, as elsewhere in csrc.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants