Repository navigation
Extend MoEPacked for Bolt ReLU-square routing and BF16 IQ2 precision - #33963
Rui Ren (rui-ren) wants to merge 7 commits into
Conversation
520f375 to
b9b776b
Compare
b9b776b to
c818493
Compare
736d86e to
a9bdbce
Compare
This comment was marked as resolved.
This comment was marked as resolved.
|
Standalone Bolt routing is now isolated in draft follow-up #33994 above #33969, extending native stack #33967 without rewriting this PR or the coworker's #32935. The existing integrated route remains compatible and uses the same shared CUDA implementation. The complete fresh production source build passed native tests, logits/cache/state and route comparisons, plus 240 exact same-mode generation/reset/counter checks. The repeated 12k/24k/32k/124k before/after timings and limitations are in #33994; those timings use the frozen extraction build, not a new benchmark of the production refactor. Preserved quantized control is not independent HF, and inherited ordinary-versus-k7 differences remain. |
|
The native kernel optimizations are now pushed in #34002 and registered as the Actual #33994-tip baseline versus the new native candidate, real 1028-1663-token The earlier ~92 tok/s private baseline already contained these optimizations: |
There was a problem hiding this comment.
🟡 Changes recommended
The BF16 IQ2 ReLU-square tests leave channel permutations and FC2 cancellation-sensitive accumulation order untested, requiring asymmetric reference cases and larger mixed-sign fixtures for both decode and prefill.
2 open issues
What changed in this PR
Extends the existing CUDA MoEPacked operator for Nemotron Bolt’s ReLU-square experts and routing requirements while retaining default SwiGLU behavior.
Changes:
- Adds preselected FP32 routes and projected sigmoid routing with selection-only correction bias.
- Introduces a BF16 IQ2 kernel preserving contiguous-K accumulation and rounding boundaries.
- Adds routing and precision regression tests and updates operator documentation.
| File | Description |
|---|---|
onnxruntime/test/contrib_ops/moe_packed_routing_test.cc |
Adds routing, malformed-route isolation, compatibility, and cancellation tests. |
onnxruntime/core/graph/contrib_ops/contrib_defs.cc |
Extends the schema with activation options and optional routing inputs. |
onnxruntime/contrib_ops/cuda/moe/moe_packed.h |
Adds activation and routing parameters. |
onnxruntime/contrib_ops/cuda/moe/moe_packed.cu |
Implements new routing paths, ReLU-square activation, and BF16 IQ2 arithmetic. |
onnxruntime/contrib_ops/cuda/moe/moe_packed.cc |
Validates new input combinations and configures CUDA dispatch. |
docs/OperatorKernels.md |
Updates the supported kernel inputs and types. |
docs/ContribOperators.md |
Documents the expanded operator schema. |
Review effort: Balanced
| const uint4 e_low = gguf_cuda::PackFloat8ToBf16x2(e0); | ||
| const uint4 e_high = gguf_cuda::PackFloat8ToBf16x2(e1); | ||
| const uint4 o_low = gguf_cuda::PackFloat8ToBf16x2(o0); | ||
| const uint4 o_high = gguf_cuda::PackFloat8ToBf16x2(o1); |
There was a problem hiding this comment.
Still open on e50e34972b, and a bit broader than channel permutations: the fixtures are also constant along K. Every block uses the same d, grid index 1 in all groups of all 8 sub-blocks, and each row's activation is a single value (factor / 32), so a mis-pairing of the A/B K fragments inside a thread's 16-element slice, or a wrong sub-block/half selection in the decode, would pass too.
For scoping: I ran the kernel text from this head outside the suite with random IQ2_XXS payloads (random bytes plus a random FP16 d per block, so all grid/sign/scale codes occur), random BF16 activations, and an independent reference (gguf_quant::DequantizeGgufRowCpu -> BF16 -> double accumulation -> the same BF16 rounding points), for hidden 256-2048, inter 256-1536 and k up to 32, in decode and prefill (4-subtile) tiles. FC1, FC2 and the finalize stage agree with the reference except for isolated BF16 rounding flips (about 0.02% of elements at most), and compute-sanitizer is clean, so I don't see a current permutation bug; this thread is about regression protection.
Suggested shape: any 66-byte pattern with a sane FP16 d is valid IQ2_XXS, so random packed bytes plus the host-side DequantizeGgufRowCpu (already in gguf_quant_tables.h) give an independent reference without extra infrastructure. Compare with a small ULP tolerance plus a cap on the fraction of mismatches: isolated elements differ by a few ULPs, mostly at BF16 rounding boundaries (and ReLU-square amplifies them in FC1).
| MoEPackedReluIq2ContiguousKKernel<false, 8><<<fc2_grid, 256, 0, stream>>>( | ||
| fc1_out, down_weights, sorted_routes, tiles, d_num_tiles, params, fc2_out); |
a9bdbce to
7b2b13a
Compare
Preserve default SwiGLU and internal routing. Add IQ2_XXS ReLU-square gate/down experts, unchanged preselected FP32 weights and expert IDs through k=32, and safe affected-token poisoning for malformed routes. Exercise analytic outputs, routing boundaries and existing defaults. Production-flag translation units compile. Runtime numerical, generation and throughput qualification remain pending. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: rui-ren <ruiren1225@gmail.com>
Add an explicit projected sigmoid routing mode with FP32 selection-only correction bias and unbiased normalized routing weights. Preserve the existing softmax/SwiGLU defaults and preselected routing. Retain the source reduction order, tie-breaking and route scaling, with FP16 and BF16 regression fixtures. Validation: all five focused MoE CUDA tests pass on H100; all eight combined projection/MoE tests pass compute-sanitizer memcheck with zero skips or errors. Source hashes match the tested v3 build. This is operator-level compatibility-overlay evidence; full-model strict bitwise parity currently fails and speculative/TPS/release acceptance remain outstanding. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: rui-ren <ruiren1225@gmail.com>
Use source-compatible FC1 and FC2 K partitions, MMA ordering, and BF16 store boundaries for the IQ2 ReLU-square path. Reuse existing expert routing and finalization, retain other dispatch paths, and exercise cancellation plus 32-route prefill subtiles. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: rui-ren <ruiren1225@gmail.com>
Keep unit contributions outside large-magnitude tensor-core partials so the exact expected result does not depend on Ampere versus Hopper internal MMA accumulation. Preserve the extreme cancellation input, kernel arithmetic and zero tolerance. The updated current packed test overlay executes all ten projection and MoE regressions on H100 without skips. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: rui-ren <ruiren1225@gmail.com>
Use the canonical CPU/CUDA/DirectML/WebGPU CI artifact for the MoE layer so schema and kernel documentation match the actual registrations. Preserve artifact provenance outside the repository. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: rui-ren <ruiren1225@gmail.com>
Keep upstream compiled-BF16 fallback, bounded launches, and checked workspace accounting while retaining contiguous ReLU-square reductions. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: f4d0e1a3-a389-4cc8-9c85-0fd0e3ffce5b Signed-off-by: rui-ren <ruiren1225@gmail.com>
Retain the upstream 16-expert softmax limit while allowing the existing preselected and projected-sigmoid paths to select up to 32 experts. Apply the same range to checked workspace sizing. Existing k=32 routing regressions expose the inherited guard mismatch. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: f4d0e1a3-a389-4cc8-9c85-0fd0e3ffce5b Signed-off-by: rui-ren <ruiren1225@gmail.com>
7b2b13a to
e50e349
Compare
Tianlei Wu (tianleiwu)
left a comment
There was a problem hiding this comment.
Reviewed this layer (stack #33967, layer 2) at e50e34972b.
Implementation. I traced the validation, workspace/chunk accounting and each new kernel, and ran the PR-head kernel code in a standalone H200 harness: random IQ2_XXS payloads against an independent dequantized reference (decode and prefill tiles), sigmoid top-k including ties and NaN/Inf up to 512 experts, and the preselected-route validator; compute-sanitizer (memcheck, initcheck, racecheck) was clean. I did not find an implementation defect. The remaining risk is in tests that cannot fail and in undocumented runtime-only constraints.
Should be fixed in this layer
MalformedRoutesPoisonOnlyTheirTokendoes not verify NaN poisoning in BF16:TensorCheck<BFloat16>checksisnan(expected)instead of the actual value (inline on the test).- The "FP32 weights" fixtures give identical BF16 expectations if the weights are rounded to BF16 (inline), so they cannot catch narrowing.
- The two coverage requests from the earlier review are still open (asymmetric/independent-reference fixtures, and an FC2-specific cancellation fixture with
inter_size >= 1024); the checklist item in the PR description is unchecked. I added details to the first thread. - Please clarify the
compute_75overlay result. The PR description and the author review say a compute_75-only PTX overlay passedReluIq2PreservesContiguousKCancellation. With compute_75 PTX the MMA body is compiled out (MmaBf16is guarded by__CUDA_ARCH__ >= 800),GetMoEPackedBf16KernelSupportreports false (ptxVersion < 80), and dispatch falls back to the SIMT kernels. Running the PR-head kernels on an H200, this fixture yields 1024 on the Tensor-Core path but 256 on the SIMT path for rows 1/8/17/129 (the SIMT FC1 lane tree adds the unit term at k=128 to2^25at k=0 before the cancellation, so FC1 is 1 instead of 2). So the expected value is specific to the Tensor-Core dispatch, the same situation asRoundedBf16ContiguousKCancellationin #33962, andHasCudaEnvironment(800)does not guarantee that path. Which dispatch did the overlay exercise?
Other requests (inline): negative tests for the ~10 new validation rules, naming (relu_squared vs the existing relu2), undocumented runtime-only constraints, the defaults-unchanged test not reaching prefill tiles, and unreachable relu_squared branches in the generic Tensor-Core SwiGLU kernels.
Minor / notes
- No test executes the >2048-row chunk loop (second-chunk offsets for preselected
indices/weightsand sigmoidrouter_input, and workspace reuse across chunks). I read the offset math as correct; a 2049-row case per routing mode (k=2 keeps it small) would pin it. - FP16
relu_squarednow rounds decoded weights to FP16 (p.relu_squared || is_bf16), whereas FP16 SwiGLU keeps FP32-decoded weights. That is a reasonable choice but is not mentioned in the PR text or the docs. - Error texts: with the default softmax routing and
router_input/router_weight/router_biassupplied, the check reports "sigmoid routing requires ...", which points at the wrong fix; and the constructor text "internally routed mode requires k <= 16" is only true for softmax (sigmoid routing is also internal but allows k <= 32). The// 1. Softmax + TopKcomment inmoe_packed.cuis stale now that it also covers sigmoid and preselected routing. docs/PackedWeightFormats.md(not part of this diff) still says MoEPacked always has three weight tensors with independently chosen formats. Withrelu_squaredthere is nouptensor and gate/down must both be IQ2_XXS, and the weight-rounding contract of the new BF16 IQ2 path (weights rounded to the activation dtype, FP32 accumulation with contiguous-K partitions, BF16 rounding of FC1/ReLU-square/FC2 outputs) is undocumented, unlike the MatMulPacked rounding table.
| if (malformed == 2) indices[1] = indices[0]; | ||
| if (malformed == 3) weights[0] = std::numeric_limits<float>::infinity(); | ||
| for (int col = 0; col < hidden_size; ++col) { | ||
| expected[col] = T(std::numeric_limits<float>::quiet_NaN()); |
There was a problem hiding this comment.
The BF16 NaN expectation here cannot fail, so this test does not verify the poisoning.
TensorCheck<BFloat16> (checkers.cc#L545-L549) does if (std::isnan(f_expected[i])) { EXPECT_TRUE(std::isnan(f_expected[i])) ... }, i.e. it tests the expected value rather than f_actual[i] (the MLFloat16 checker a few lines above uses f_actual). MalformedRoutesPoisonOnlyTheirToken only instantiates the default T = BFloat16, so any output for token 0 (finite garbage, stale workspace, zeros) satisfies the NaN expectation, and the test effectively only asserts that token 1 is unaffected. The checker is unchanged up to #34003, so this is not fixed later in the stack.
Suggested fix: a custom verifier replaces the default comparison, so it has to check both tokens (the Mamba2Verify test in #33969 works around the same checker behavior this way):
test.SetCustomOutputVerifier([&](const std::vector<OrtValue>& fetches, const std::string&) {
const auto out = fetches[0].Get<Tensor>().DataAsSpan<T>();
for (int col = 0; col < hidden_size; ++col) {
EXPECT_TRUE(std::isnan(out[col].ToFloat())) << "token 0 col " << col;
EXPECT_EQ(out[hidden_size + col].ToFloat(), expected[hidden_size + col].ToFloat()) << "token 1 col " << col;
}
});(for the malformed != 0 cases), and/or also run the malformed cases for MLFloat16, whose checker is correct. Fixing TensorCheck<BFloat16> itself to test f_actual[i] would cover every BF16 test, but it may surface other latent expectations, so it is probably better as a separate PR.
| const float weight = sigmoid ? (probability(expert) / (sigmoid_sum + 1e-20f) * 5.f) | ||
| : (rank == 0 ? .333333f : .5001f); |
There was a problem hiding this comment.
These fixtures cannot tell FP32 expert weights from BF16-rounded ones.
Replaying this helper's expected-value math with the real BFloat16/MLFloat16 conversions, the BF16 expectation is identical if weight is first rounded to T, for every case in ReluSquaredPreservesPreselectedFp32Routing (rows {1, 8, 9, 17} x k {2, 32}) and in the BF16 half of SigmoidCorrectionBiasUsesUnbiasedWeights: .333333f and .5001f are within a third of a BF16 ulp of their roundings, so the difference vanishes in the final BF16 rounding (the committed expectations are 8.1875 / 29.25 for k=2 and 119.5 / 474 for k=32 either way). A regression that narrows expert_weights (or the normalized sigmoid weights) to T would therefore pass. FP16 is only partly better: it differs for k=2 only.
Weights that do discriminate: {.3301f, .501f} give expected 8.1875 / 29.125 / 120 / 476 for (even row, k=2) / (odd row, k=2) / (even row, k=32) / (odd row, k=32) with FP32 weights, but 8.125 / 29 / 119.5 / 474 with BF16-rounded weights.
| ORT_RETURN_IF_NOT(sigmoid_routing | ||
| ? (router_input && router_weight && router_bias && !router_probs && !preselected && | ||
| normalize_routing_weights_ == 1 && relu_squared) | ||
| : (!router_input && !router_weight && !router_bias), | ||
| "MoEPacked sigmoid routing requires router_input, router_weight, router_bias, " | ||
| "relu_squared and normalize_routing_weights=1, with no other routing inputs"); | ||
| ORT_RETURN_IF_NOT(sigmoid_routing || | ||
| (preselected ? (indices && weights && !router_probs && normalize_routing_weights_ == 0) | ||
| : (router_probs && k_ <= 16)), | ||
| "MoEPacked preselected routing requires both indices and weights, no router_probs and " | ||
| "normalize_routing_weights=0; internal routing requires router_probs and k <= 16"); | ||
| ORT_RETURN_IF_NOT(relu_squared ? !up_w : up_w != nullptr, | ||
| "MoEPacked swiglu requires up weights; relu_squared requires omitted up weights"); |
There was a problem hiding this comment.
The new validation matrix has no negative tests.
This block adds about ten distinct rejection rules (the constructor adds more), but moe_packed_routing_test.cc has no OpTester::ExpectResult::kExpectFailure case, so relaxing any of them would go unnoticed. Several guard memory safety or silent truncation (router_width % 256, where the scores kernel would otherwise drop or over-read the row tail; the shape checks; and k <= 16 for the internal softmax kernels). Suggested cases, each a few lines on top of the existing helpers: (1) router_probs together with expert_indices; (2) preselected routing with normalize_routing_weights=1; (3) expert_indices without expert_weights; (4) sigmoid without relu_squared, with router_probs, with router_width % 256 != 0, and with a wrong router_weight/router_bias shape; (5) relu_squared with up_experts_weights present, or with non-IQ2_XXS gate/down; (6) internal softmax with k > 16; (7) swiglu without up_type; (8) router_scale != 1 with softmax.
Minor: the attribute-only rules (relu_squared => IQ2_XXS, sigmoid => relu_squared + normalize=1, router_scale) can be checked in the constructor so they fail at session creation rather than at the first Run.
| gate_ptrs[m], p.gate_type, sb32_idx, t, s_lut, g_vals); | ||
| gguf_cuda::DecodeGgufSubBlock8SmemLutTpl<kGateUpType>( | ||
| up_ptrs[m], p.up_type, sb32_idx, t, s_lut, u_vals); | ||
| if (!p.relu_squared) { |
There was a problem hiding this comment.
Nit: this relu_squared branch (and the matching ones at L581, L720 and L753 in the generic Tensor-Core SwiGLU kernels) is unreachable. The constructor limits relu_squared to IQ2_XXS gate and down weights, and DispatchMoEPackedCuda sends exactly that combination to MoEPackedReluIq2ContiguousKKernel on the BF16 Tensor-Core path (or to the SIMT kernels), so the generic Tensor-Core kernels never see p.relu_squared == true. The dead code has no test and costs a runtime branch plus an 8-register zero-init (float u_vals[8] = {}) in the hot SwiGLU decode/prefill kernels. Either drop these edits (the SIMT ones are live and should stay) or add the dispatch path and the test that need them.
| .Attr("normalize_routing_weights", "Whether to normalize top-k routing probabilities.", | ||
| AttributeProto::INT, static_cast<int64_t>(1)) | ||
| .Attr("activation_type", "Expert activation function ('swiglu').", | ||
| .Attr("activation_type", "Expert activation function ('swiglu' or 'relu_squared').", |
There was a problem hiding this comment.
Naming: MoE and QMoE already accept activation_type='relu2' for squared ReLU (contrib_defs.cc#L1711, parsed in moe_base.h#L52). Introducing relu_squared for MoEPacked spells the same activation two ways inside one operator family, and contrib schema strings are effectively frozen once released. Prefer relu2 (optionally also accepting relu_squared for readability), and update the kernel, tests and docs accordingly.
| .Input(7, "router_input", | ||
| "Projected sigmoid routing input [num_rows, router_width]; omit other routing inputs. " | ||
| "Requires relu_squared, routing_function='sigmoid' and normalize_routing_weights=1.", | ||
| "T", OpSchema::Optional) |
There was a problem hiding this comment.
Document the constraints that are only enforced at runtime. The schema and docs/ContribOperators.md do not say that: relu_squared currently requires IQ2_XXS gate and down weights (moe_packed.cc#L76-L78); router_width must be a positive multiple of 256 (moe_packed.cc#L138); the internal softmax path is limited to k <= 16 while sigmoid/preselected routing allow 32 (the k attribute text is unchanged); and router_scale/normalize_routing_weights are only meaningful with specific routing modes. Related wording: the Tweight constraint text ("Preselected routing weight type") is also used by router_bias, and the operator-level SetDoc still describes "gate, up, and down" blobs. Users otherwise only discover these through ORT_ENFORCE messages.
| constexpr int experts = 2; | ||
| const BFloat16 hidden(.125f * (.25f / (1.f + std::exp(-.25f)))); | ||
| const BFloat16 output(16.f * hidden.ToFloat()); | ||
| for (int rows : {4, 48}) { |
There was a problem hiding this comment.
This "defaults unchanged" test stays in the decode kernels, and SwiGLU + preselected routes is untested.
With k = 1, rows {4, 48} give 4 and 48 routes, both below the 128-route prefill threshold (use_prefill_tiles requires more than 128 routes), so only the decode Tensor-Core kernels run, while this PR also edits the prefill kernel MoEPackedFc1SwiGLUBf16TcKernel. Please add a case with >= 129 routes (e.g. rows = 129). Validation also admits SwiGLU with preselected routes and k up to 32 (nothing requires relu_squared for preselected mode), but no test exercises it; one small preselected SwiGLU case (k = 2 and k = 32) would pin that contract down.

Description
Extends
com.microsoft::MoEPacked(CUDA) for Nemotron Bolt's ReLU-square experts androuting contract without introducing a competing expert operator. Adds
activation_type="relu_squared", preselected expert IDs with FP32 mixture weights,and integrated projected-sigmoid routing. Correction bias affects selection only;
mixture weights come from unbiased sigmoid affinities and post-normalization scaling.
The existing softmax/SwiGLU defaults remain supported.
Preselected and projected-sigmoid routing support up to 32 selected experts; the
existing internal softmax path retains its 16-expert top-k limit. Validation keeps
routing modes mutually exclusive and isolates malformed routes to their token.
BF16 IQ2 FC1/FC2 preserve complete FP32 scale reconstruction, contiguous-K partitions,
BF16 storage boundaries, and FP32 routing/final reduction.
Dependency: #33962 in
native stack #33967;
the eventual landing target is
main. Packed weight layouts and GenAI source do not change.Summary of Changes
Operator Contract and Documentation
onnxruntime/core/graph/contrib_ops/contrib_defs.ccdocs/ContribOperators.md, docs/OperatorKernels.mdCUDA Implementation and Regression Tests
onnxruntime/contrib_ops/cuda/moe/moe_packed.{cc,h,cu}onnxruntime/test/contrib_ops/moe_packed_routing_test.ccTesting
Review snapshot (2026-11-09):
e50e34972b245acf51d3f259efe992a37bd6b24f.MoEPackedRoutingTest.*in a CUDA-enabled test build.compute_75-only PTX passed
ReluIq2PreservesContiguousKCancellation.bbac6c023794) passed all 19projection/routing regressions with zero skips. That is integration evidence for
the later stack, not a clean standalone build of this PR.
dequantized reference for decode and full 32-route prefill
(thread),
plus mixed-sign FC2 cancellation with inter_size >= 1024
(thread).
Existing fixtures do not establish those properties.
full-model/performance acceptance is claimed. Historical routing timings below
compare routing alternatives in one build, not this entire PR against its parent.
Detailed review.
Historical measurements and qualification evidence (not current-head acceptance)
These previously reported results are retained for provenance. They have not been rerun on the rebased head above; earlier CI/build status and benchmark results do not establish current-head CI, accuracy, or merge readiness.
Measured CUDA kernel latency - 2026-11-07
BF16, IQ2_XXS, ReLU-square experts, latent hidden dimension 1024, top-k 9. Internal softmax routing versus preselected FP32 indices/weights using identical routes and the same corrected expert arithmetic. Selected logits are +16 and others -16; selected weights are 1/9. Shared/distinct labels describe expert reuse across rows (the one-row variants are equivalent workload controls). Sixteen rows create 144 routes, crossing the 128-route prefill-tiles dispatch threshold.
The ratio measures routing-path overhead removed, not this entire PR versus its parent. All 24 complete BF16 outputs are bitwise equal in both runs. Benefits diminish for distinct-expert, bandwidth-heavy batches; all cases, including those smaller gains, are shown.
H100 80GB HBM3, CUDA 12.8.93, driver 580.159.04, GCC 11.4.0. CUDA-event median per dispatch; 64 dispatches per captured CUDA graph, 50 warmups, three graph warmups, and 15 paired timed batches alternating AB/BA. Four rotating allocation banks contain deterministic synthetic weights (seed 617); this is not a cold-cache guarantee or a checkpoint-weight benchmark. Raw p95 values describe batch-average latency, not individual-invocation tails. No profiler runs in the timing window. The assigned GPU had empty admission, one stable compute client during execution, and clean release afterward; clocks were not locked. Two independent 39-case runs passed, including the published replay runner.
Scope: same-build production CUDA dispatch alternatives, not historical parent-commit binaries. The measured runtime is an isolated CUDA provider relink with preserved core/support, not a clean full-source ORT build or end-to-end model qualification. Each measured operator's CUDA source hash matches its owning PR head. A clean source build is a separate replication variant and must be measured again. Kernel-only timings do not establish model TTFT/tok/s, MMLU/LCB, speculative integration, or full-model readiness.
Complete primary/repeat results, numerical metrics, source hashes and benchmark code. All 39 cases passed in both supervised runs; full-model work remains paused. No source push, stack change, merge, auto-merge or ready-for-review promotion is included.
Checklist