[CUDA] Add MoERouter and SwiGLU contrib operators - #32126
Open
Tianlei Wu (tianleiwu) wants to merge 3 commits into
Open
[CUDA] Add MoERouter and SwiGLU contrib operators#32126Tianlei Wu (tianleiwu) wants to merge 3 commits into
Tianlei Wu (tianleiwu) wants to merge 3 commits into
Conversation
MoERouter turns a gate GEMM's output into the two tensors an expert-parallel QMoE call needs. Every rank sees the same scores and computes the same global top-k independently, so routing needs no collective; each rank then emits only its local expert columns in the log domain, and `weight_scale` multiplies QMoE's local softmax normalisation back out. A token with no local expert gets a zero scale. SwiGLU exposes the clamped gated activation MoE and QMoE already apply internally, with the same limit/alpha/beta contract, for the dense feed-forward and shared-expert paths that do not run through a grouped expert GEMM.
Tianlei Wu (tianleiwu)
requested review from
Akshay Sonawane (apsonawane),
Justin Chu (justinchuby) and
kunal-vaishnavi
August 18, 2026 16:59
Contributor
|
SwiGLU is also declared at https://onnx.ai/onnx/operators/onnx__SwiGLU.html. We could register for the standard domain too? |
| </dl> | ||
|
|
||
|
|
||
| ### <a name="com.microsoft.MoERouter"></a><a name="com.microsoft.moerouter">**com.microsoft.MoERouter**</a> |
Contributor
There was a problem hiding this comment.
This will be a very useful op for expert parallelism. Thanks for adding it! Should we consider expanding its usage for other types of routing besides top-k and hash routing?
| </dl> | ||
|
|
||
|
|
||
| ### <a name="com.microsoft.SwiGLU"></a><a name="com.microsoft.swiglu">**com.microsoft.SwiGLU**</a> |
Contributor
There was a problem hiding this comment.
We already have a partial C++ graph fusion for SwiGLU that runs on existing models. Can we update it to work with the SwiGLU op?
| constexpr int kThreads = 256; | ||
|
|
||
| // ORT's Sigmoid keeps the exponent non-positive on both branches. | ||
| __device__ __forceinline__ float Sigmoid(float a) { |
Contributor
There was a problem hiding this comment.
Can we move this into a shared utility for common CUDA kernel calculations?
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
Adds two CUDA contrib operators for the MoE block.
MoERouterturns a gate GEMM's output into the two tensors an expert-parallelQMoEcall needs, andSwiGLUexposes the clamped SwiGLU activation thatMoE/QMoEalready apply internally, for the dense and shared-expert paths that do not run through a grouped expert GEMM.MoERouteris the piece that makes expert parallelism expressible in ONNX today: the existingQMoEoperator takesrouter_probsand applies its own softmax, so a rank holding only a slice of the experts needs a router row constructed such that that softmax still produces the globally correct weights. Doing that with primitives requires a top-k, a scatter and a masked log per layer.Summary of Changes
MoERouteronnxruntime/contrib_ops/cuda/math/moe_router.{h,cc}OpKernel— attributes, validation, output shapesonnxruntime/contrib_ops/cuda/math/moe_router_impl.{h,cu}Scoring and selection are separate, orthogonal attributes so the operator covers the common router designs rather than one model's:
Ties go to the lower expert index. Supplying the optional
expert_idsinput overridesselectionand fixes the choice per token (hash routing);biasis then ignored. Either way the weights are the affinities of the chosen experts, normalized to sum to one.Expert parallelism is the reason for the two outputs:
router_probscarrieslog(w_e)for a chosen local expert and a large negative value elsewhere (-1e30, or-1e4for float16 which-1e30does not survive), soQMoE's own softmax over the local block returnsw_e / W_local.weight_scaleisroute_scale * W_local, which multiplies that factor back out of the expert output before the all-reduce. A token with no local expert gets a zero scale, which annihilates the degenerate uniform softmax an all-negative row would otherwise produce.local_expert_start/local_expert_countdescribe the rank's slice; a single-rank model just sets them to the full expert range.SwiGLUonnxruntime/contrib_ops/cuda/math/swiglu.{h,cc}OpKernelonnxruntime/contrib_ops/cuda/math/swiglu_impl.{h,cu}T. Alimitof zero or less disables both clamps.MoE/QMoEapply internally via theirswiglu_limit/activation_alpha/activation_betaattributes, so a model can use the same three values for its routed experts and its dense path.upis optional. When omitted,gatecarries both halves of one[.., 2 * inter]projection (gate first, then up) and the split is done internally, so a fused sibling GEMM needs noSplitnode.Why a standalone operator rather than folding the path into
MoE/QMoE: the motivating case is a shared expert that is tensor-parallel sharded while the routed experts are expert-parallel sharded. It therefore has a different (here 8x larger) local intermediate size than a routed expert, is ungated, and runs on every rank — so it cannot be represented as one more always-on expert slot inside the grouped GEMM without putting dense per-token work on a single rank. Accepting TP-sharded shared-expert weights as separateQMoEinputs is a reasonable follow-up; this PR is the piece needed to express the path at all.Shared
onnxruntime/core/graph/contrib_ops/contrib_defs.cconnxruntime/core/graph/contrib_ops/ms_opset.honnxruntime/contrib_ops/cuda/cuda_contrib_kernels.ccdocs/ContribOperators.md,docs/OperatorKernels.mdTesting
onnxruntime/test/python/transformers/test_moe_router_swiglu.py— PyTorch references for both operators over float32/float16/bfloat16: scoring/selection combinations, hash routing viaexpert_ids, the no-local-expert token, the float16-1e4sentinel,SwiGLUwith and withoutup, with and without the clamp, and non-defaultactivation_alpha/activation_beta.Additive only — no existing operator, schema or kernel is modified;
MoE/QMoEare untouched.Checklist