[CUDA] Fuse MoE softmax/top-k into the expert-map prologue at decode - #29919
Merged
Conversation
This reverts commit 67e67e1.
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Contributor
There was a problem hiding this comment.
Pull request overview
This PR reduces decode-time overhead in the CUDA MoE/QMoE path by fusing the standalone softmax/top-k routing step into the expert-map prologue when the configuration is a single token with ≤32 experts (and no expert parallelism), with an environment-variable kill switch to revert to the prior two-kernel behavior.
Changes:
- Adds a fused single-warp “routing + expert-map build” prologue for the single-token case and wires it through
runMoeviaFusedRoutingParams. - Moves softmax/top-k helper math into
topk_warp_sort.cuhso the standalone and fused routing paths share bit-identical implementations. - Updates MoE/QMoE call sites to optionally skip
LaunchSoftmaxTopKand instead pass raw router logits into the runner for fused routing.
Reviewed changes
Copilot reviewed 6 out of 6 changed files in this pull request and generated no comments.
Show a summary per file
| File | Description |
|---|---|
| onnxruntime/core/providers/cuda/cu_inc/topk_warp_sort.cuh | Centralizes shared softmax/top-k helper math in a CUDA header for bit-identical reuse. |
| onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.cu | Removes duplicated helper implementations and reuses the shared helpers from topk_warp_sort.cuh. |
| onnxruntime/contrib_ops/cuda/moe/moe.cc | Passes default (disabled) fused-routing parameters into runMoe for the non-quantized MoE path. |
| onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc | Adds gating logic to skip standalone softmax/top-k at supported decode shapes and passes FusedRoutingParams into runMoe. |
| onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_kernels.h | Extends the runner interface with FusedRoutingParams and declares isFusedMoeRoutingSupported(). |
| onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_kernels.cu | Implements the fused routing kernel, kill switch, support predicate, and integrates the fused path into runMoe. |
kunal-vaishnavi
approved these changes
Jul 29, 2026
tianleiwu
added a commit
that referenced
this pull request
Jul 31, 2026
…ode (#31156) ### Description Cuts the cost of the MoE routing prologue and epilogue at small batch, where those kernels stop being negligible relative to the expert GEMMs. Touches `moe_gemm/moe_kernels.cu`, `moe/qmoe_kernels.cu` and `core/providers/cuda/cu_inc/topk_warp_sort.cuh`. ### Motivation and Context At decode batch sizes the routing kernels run on a handful of blocks (grid 1 and grid 4 in the table below), so they are latency-bound, not throughput-bound, and their fixed cost is a measurable share of the step. Once the expert GEMVs themselves were optimized, the routing family became the next item on the profile. `BlockRadixRank` keeps a private set of packed digit counters *per thread*, which is what makes the sort-based path expensive at these sizes; the replacement is a warp-level top-k in `core/providers/cuda/cu_inc/topk_warp_sort.cuh`. **Tie-breaking is unchanged** — the same packed stable key is used, so routing decisions are identical. ### Measured impact 1x H200 SXM (SM90, 132 SM, ~4.8 TB/s HBM), CUDA 13.0, Qwen3.6-35B-A3B-NVFP4 + MTP `N=3` (verify batch `M=4`). Per kernel (graph OFF): | kernel | before (µs) | after (µs) | before (ms/step) | after (ms/step) | |---|---:|---:|---:|---:| | `finalizeMoeRoutingKernel` (grid 4) | 4.64 | **4.18** | 0.185 | 0.169 | | `SoftmaxTopKMergeKernel` -> `WarpTopKKernel` | 4.52 | **3.68** | 0.181 | 0.150 | | `fusedBuildExpertMapsSortFirstTokenKernel` (grid 1) | 4.16 | **2.22** | 0.167 | 0.089 | | `expandInputRowsKernel` (grid 1056) | 2.49 | — | 0.100 | 0.100 (unchanged) | | **routing family total** | | | **0.656** | **0.529** | All kernels (both families): 6.033 -> 5.893 ms/step (−0.140 ms/step). End-to-end (graph ON, 5 interleaved reps per arm): * before: 8.511 / 8.535 / 8.501 / 8.541 / 8.566 * after: 8.467 / 8.382 / 8.370 / 8.374 / 8.445 **8.531 -> 8.408 ms/step (−1.4%).** Left on the table (~0.05 ms/step, judged low ROI): 32-bit packed keys in `WarpBitonicTopN`, and splitting `finalizeMoeRouting`'s columns across more blocks to use more than 4 SMs' L1 bandwidth. ### Tests Existing MoE/QMoE unit tests; routing output is unchanged by construction (same stable key), so these are parity tests rather than new coverage. ### Methodology note End-to-end deltas are quoted as **ms/step** from a fixed-step measurement, never tok/s: any numerics change alters the generated sequence and therefore the MTP acceptance rate, which swamps the speed delta. Per-kernel durations are taken with CUDA graphs **off** — `nsys --cuda-graph-trace=node` inflates durations ~35% globally and up to 3.8x for large-grid kernels. > [!NOTE] > Rebased onto `main` after #29919 ("Fuse MoE softmax/top-k into the expert-map prologue at > decode") merged. The two are **complementary, not competing** — no textual or semantic > conflict: > > * #29919's `fusedRoutingBuildExpertMapsSingleTokenKernel` is gated by > `isFusedMoeRoutingSupported()`, which requires `num_tokens == 1`, `ep_size == 1` and > `num_experts <= kWarpBitonicMaxSize` (32). Qwen3.6-35B-A3B has **256 experts**, and an MTP > verify step has **4 tokens**, so that path never engages here. > * `fusedBuildExpertMapsSortFirstTokenKernel` (the general prologue) still uses > `cub::BlockRadixRank` on `main` — this PR's `BlockRadixRankMatch` swap applies to it. > * The `num_experts <= 256` dispatch still lands on `SoftmaxTopKMergeKernel<T, 128, 2>` — this > PR's `SoftmaxTopKWarpTopKKernel` replaces it for `num_experts <= 256 && k <= 8`. > * `finalizeMoeRoutingKernel` was untouched by #29919. > > #29919 also moved `WarpReduceMax` / `WarpReduceSum` / `SafeInvSum` / `TopKNormalizeDenom` into > `topk_warp_sort.cuh` and re-exported them into `qmoe_kernels.cu` with `using` declarations, so > the unqualified calls added here still resolve. --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
At decode (one token per sequence) the MoE/QMoE path launches a standalone softmax/top-k kernel and
then a second single-block kernel that builds the expert permutation maps. Both are one-block
kernels whose cost is launch and drain latency rather than work, so the second launch is nearly pure
overhead. This PR merges them.
fusedRoutingBuildExpertMapsSingleTokenKernelhandles the single-token case with one warp: laneeowns expert
e's logit, warp reductions give the softmax, andWarpBitonicSortDescendinggives thetop-k. For a single token the permutation is just "the k selected experts ordered by expert id",
which every lane derives with shuffles, so no CUB block rank and no shared memory are needed. The
kernel writes
token_selected_experts,token_final_scalesand all three permutation maps.Summary of Changes
launchFusedRoutingBuildExpertMapsreplaces theLaunchSoftmaxTopK+fusedBuildExpertMapsSortFirstTokenpair on the supported decode shape.SoftmaxTopKWarpBitonicKernel, soSoftmaxScale,SafeInvSum,TopKNormalizeDenom,WarpReduceMaxandWarpReduceSummove fromqmoe_kernels.cuintotopk_warp_sort.cuh, andboth translation units now share one definition.
isFusedMoeRoutingSupported()is host-only anddeterministic. The QMoE op uses it to decide whether to skip its own softmax/top-k launch, and
runMoere-checks it, so the two can never disagree about who computed the routing. It requiresone token, no expert parallelism, and
num_experts <= 32. The FP4 family is excluded in the opbecause its decode fast path consumes
expert_indices/expert_scalesitself instead of callingrunMoe.ORT_DISABLE_MOE_FUSED_ROUTING=1restores the two-kernel prologue.Performance
The win scales with how large a launch is relative to the rest of the call, so it is much bigger on
a smaller GPU. Isolated QMoE node at the gpt-oss-20b decode shape (hidden 2880, intermediate 2880,
32 experts, top-4, batch 1, fp16, int4), arms interleaved within every pass via the kill switch:
RTX 3060 (sm_86), per-channel int4, 12 passes x 1000 reps
ORT_DISABLE_MOE_FUSED_ROUTING=1Fused wins 12 of 12 paired passes (median +7.0%).
RTX 3060 (sm_86), block-wise int4 (
block_size=64), 8 passes x 1000 repsORT_DISABLE_MOE_FUSED_ROUTING=1Fused wins 8 of 8 paired passes (median +5.0%).
On an RTX 5060 Ti (sm_120) the same change measured +0.24% end-to-end decode throughput on
gpt-oss-20b — inside that harness's noise floor. A tripwire build confirmed the fused path really
does engage at decode there, so the sm_120 number is a genuine null result rather than a dead path;
the faster GPU simply amortizes the extra launch better.
Testing
On the RTX 3060 (sm_86), with and without
ORT_DISABLE_MOE_FUSED_ROUTING=1:onnxruntime/test/python/transformers/test_qmoe_cuda.py: 109 passed, 13 skippedonnxruntime/test/python/transformers/test_moe_cuda.py: 48 passed, 15 skipped