[CUDA] Add INT4 and INT3 QMoE CPU/GPU offload support by tianleiwu · Pull Request #33575 · microsoft/onnxruntime

Tianlei Wu (tianleiwu) · 2026-10-24T16:55:27Z

Description

Based on main after #33452 and #33450 merged (with already-merged #33569 commits dropped):

SM80-Prepacked INT4 & Raw INT3 Offload Admission (moe_quantization.cc): Admits SM80-prepacked INT4 (quant_type="int", expert_weight_bits=4, weights_prepacked=1, block_size=32) and raw INT3 (quant_type="int", expert_weight_bits=3, weights_prepacked=0, block_size in {32, 64, 128}) com.microsoft::QMoE nodes into the v4 MoE offload pipeline.
GPU Int3MoeParams Cold-Slot Skipping (qmoe_kernels.{h,cu}): Extends FusedFc1FastKernel, Fc2RouteFastKernel, and LaunchInt3Typed to skip negative cold-miss slots (expert < 0), support source_route_ids (hit_ids) indirection, and emit wave_output partials.
int4_bank_mode_ & int3_bank_mode_ in QMoEOffload (qmoe_offload.{h,cc}): Adds INT4 and INT3 offload support across doorbell hybrid decode, adaptive promotion, byte-preserving FileExpertStore host caching (scales_are_block_major = false on CPU), and wave-streaming prefill, along with test_v4_int4_and_int3_offload_modes_and_numerical_parity in test_qmoe_offload_v4.py.

PR stack

Merge in this order:

[CUDA] Add INT4 and INT3 QMoE CPU/GPU offload support #33575 - INT4/INT3 QMoE CPU/GPU offload (base: main)
[CUDA] Optimize INT4 QMoE offload prefill/decode and S=8 doorbell execution #33622 - INT4 offload prefill/decode and S=8 doorbell optimizations (base: [CUDA] Add INT4 and INT3 QMoE CPU/GPU offload support #33575)
[CUDA] Add 8-route grouped INT3 offload prefill, zero-cold doorbell short-circuit, and S=8 streaming optimizations #33629 - Grouped INT3 offload prefill and S=8 streaming optimizations (base: [CUDA] Optimize INT4 QMoE offload prefill/decode and S=8 doorbell execution #33622)
[CUDA] Enable Tensor-Core INT3/INT4 QMoE offload decode/prefill and 8K single-tile streaming for Qwen3.8 RC5 #33731 - Tensor-core INT3/INT4 offload decode/prefill (base: [CUDA] Add 8-route grouped INT3 offload prefill, zero-cold doorbell short-circuit, and S=8 streaming optimizations #33629)

Prerequisites #33452, #33450, #33689, and #33650 have merged into main.
Retarget each dependent PR to main after its predecessors merge.

github-actions

You can commit the suggested changes from lintrunner.

Copilot

🟡 Changes recommended

Unresolved offload issues can cause null-pointer dereferences, out-of-bounds scale reads, silent FP16/BF16 misinterpretation, and streaming that exceeds the configured expert-wave limit.

4 open issues

What changed in this PR

Extends the existing QMoE CPU/GPU offload backend to support prepacked INT4, raw INT3, and models combining both formats through session.moe.offload.* options.

Changes:

Adds integer-format admission and per-layer expert memory accounting.
Integrates CPU cold-expert execution, GPU resident/streaming dispatch, and INT4 scale transposition.
Adds numerical parity tests and documents supported integer layouts.

File	Description
`onnxruntime/test/python/test_qmoe_offload_v4.py`	Adds integer model builders, offload parity tests, and Windows provider setup.
`onnxruntime/contrib_ops/cuda/moe/qmoe_offload.h`	Declares integer bank modes and the scale-transposition helper.
`onnxruntime/contrib_ops/cuda/moe/qmoe_offload.cc`	Integrates integer expert loading, CPU execution, GPU streaming, and promotion handling.
`onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.h`	Adds INT3 route-indirection and gathered-input parameters.
`onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.cu`	Extends fast INT3 kernels for cold-slot skipping and partial outputs.
`onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc`	Enables integer offload registration and format-aware memory accounting.
`docs/contrib_ops/cuda/moe_qmoe.md`	Documents INT4/INT3 offload eligibility and tensor layouts.

_{Review effort: Balanced}

github-actions

You can commit the suggested changes from lintrunner.

Ti-Tai Wang (titaiwangms) · 2026-10-30T17:28:21Z

Static review: offload correctness and Olive/Mobius interoperability

Reviewed head: e622696689efd7bf4fba22929420854fa37d003b (verified as the current head before posting).

Recommendation: address the two major findings below before qualifying this path. This is a source-level review, not a CUDA reproduction. No combined build of #33575, #33622, and #33689 or end-to-end inference was run.

Major: INT4 admission and backend disagree about `weights_prepacked=-1`

Locations: moe_quantization.cc:423–444, 918–936, and qmoe_offload.cc:2938–2945.

The CUDA-side normalization treats -1 and 1 as the SM80 integer layout, and the new INT4 admission uses that normalized boolean. Backend initialization instead requires the original attribute to equal 1. A graph declaring -1 can therefore pass admission and subsequently fail initialization.

Please normalize the declaration consistently, or reject it explicitly at admission rather than admitting it and failing later. Cover omitted/default, explicit -1, and explicit 1 declarations with offload enabled. This finding does not assume that Mobius emits -1.

Major: zero experts reaches division by zero in per-expert accounting

Location: moe_quantization.cc:942–1000.

The new accounting reads num_e from fc1_w.Shape()[0] and divides by it without first validating that it is positive. Rank/CPU-placement admission checks do not exclude an INT4 initializer with E=0; the pinned non-INT3 shape-inference path propagates output type/shape without validating this dimension. The neighboring record_slice guard does not protect the separate raw_slot_bytes division.

Please validate positive, representable expert/feature dimensions before descriptor construction and per-expert arithmetic. Add an E=0 rejection test that returns a normal model-validation error rather than executing undefined arithmetic.

Smaller follow-ups

test/python/test_qmoe_offload_v4.py:1627–1628 still expects quant_type='int' offload requires, while the constructor now emits QMoE offload mode='required' rejected node. Update the assertion and cover the auto-mode behavior described by the test comment. This mismatch was inspected, not executed.
Please document a separate offload eligibility matrix for projection widths, block size, prepack declaration, fusion, and zero-point inputs, together with the changed async/stream defaults. Ordinary QMoE support is not the same as offload eligibility.

Producer/consumer contract

Olive RTN and KQuant save the same native packed checkpoint format; KQuant changes quantization-parameter selection, not the storage ABI. The inspected Mobius revision, onnxruntime/mobius@7f61c0d7820628f288b98c285badca22eaa66d3d, consumes that format and interleaves INT3 FC1 gate/up rows and companion scales before export. I am not reporting a missing-interleave defect in Mobius.

There is an important distinction for integration:

Mobius exports raw QMoE weights with weights_prepacked=0. That matches the raw INT3 contract, subject to the other eligibility checks.
This INT4 offload path instead requires actual SM80-prepacked weights, block size 32, and the corresponding declaration. Changing the attribute alone is not a valid conversion.
Asymmetric/mixed configurations outside the advertised offload contract are not themselves regressions.

Not checked: CUDA execution, numerical parity, CUDA graphs, performance, concurrent sessions, or the combined PR heads. ModelBuilder/GenAI integration is outside this review's scope.

Copilot

🟡 Changes recommended

Unresolved critical INT4/INT3 offload validation and format-handling defects can cause out-of-bounds reads from bias, scale, and weight buffers.

11 open issues

4 resolved since last review

_{Review effort: Max}

- Admit SM80-prepacked INT4 (quant_type="int", expert_weight_bits=4, weights_prepacked=1, block_size=32) and raw INT3 (quant_type="int", expert_weight_bits=3, weights_prepacked=0, block_size in {32, 64, 128}) QMoE nodes into the v4 MoE offload pipeline in SessionState and QMoE::BuildOffloadDescriptors. - Validate positive expert counts and bank/scale dimensions via ResolveNvfp4BankDims and ResolveInt3BankDims before per-expert byte division, and require explicit weights_prepacked=1 for INT4 offload across admission and backend initialization. - Extend GPU Int3MoeParams (FusedFc1ScalarKernel, Fc2RouteScalarKernel, FusedFc1FastKernel, Fc2RouteFastKernel, LaunchInt3Typed) to skip negative cold-miss slots (expert < 0), support source_route_ids (hit_ids) indirection, and emit wave_output partials. - Add int4_bank_mode_ and int3_bank_mode_ in QMoEOffload across doorbell hybrid decode, adaptive promotion, bounded FileExpertStore host caching, and wave-streaming prefill while honoring prefill_wave_experts. - Strengthen INT3/INT4 offload numerical parity tests (nonzero reference signal, per-token relative L2 and max-abs checks, controlled cold-route check, FP16 + block_size=128, auto-mode fallback for raw/omitted/-1 INT4, and E=0 rejection) and update offload eligibility docs. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

- Gate offload admission on SM80+ device capability before CPU initializer relocation in SessionState and QMoE::BuildOffloadDescriptors while retaining clear rejection for mode='required'. - Require H and I to be divisible by 64 for prepacked INT4 across SessionState pre-placement validation, ResolveNvfp4BankDims, QMoE::BuildOffloadDescriptors, and QMoEOffload::InitializeNvfp4Offload. - Validate optional FC1/FC2 biases as [E, 2*I] and [E, H] with matching activation dtype before admission and upload. - Exclude integer banks from NVFP4 prefill lookahead registration and tail-slab lending. - Ignore unused global-scale inputs (slots 15/16) on integer banks and synthesize E unit scales before host-cache construction and uploads. - Scope group_calls=64 and fc1_cols=128/fc2_cols=256 defaults to integer banks so NVFP4 layers retain their previous CpuWork footprint. - Preserve explicit stream_min_tokens / ORT_QMOE_STREAM_MIN_TOKENS overrides on integer banks. - Stage transposed INT4 scale banks in event-retained buffers instead of synchronizing the stream after every scale bank upload. - Track async decode launch sequence across both doorbell and callback paths so repeated uncaptured 8-token runs after a 2-token doorbell run continue advancing adaptive promotions. - Add unit and parity test coverage for all new validation and runtime cases and accept the INT3 schema diagnostic in rejection tests. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Copilot

🟡 Changes recommended

Unresolved critical issues include incorrect FP16 results, promotion data races and unsafe buffer reuse, and allocations that exceed configured memory budgets.

10 open issues

11 resolved since last review

_{Review effort: Max}

… feedback - Replace std::getenv with ORT Env/ParseEnvironmentVariableWithDefault in CUDAExecutionProvider::SupportsMoeOffload and PluginExecutionProvider::SupportsMoeOffload to resolve Windows MSVC C4996 deprecation errors. - Size integer streaming waves and tiles against configured GPU working and pinned budgets (both CPU-enabled and CPU-disabled streaming). - Bound and recycle the scale-transposition staging pool during batch restores and uploads (kMaxStagedScaleBuffers = 8) and trim excess buffers once restores complete. - Use session-wide QMoESessionController::CurrentRunId() for promotion epochs and retain each SessionPromotionBudget bounce buffer via CUDA events until its H2D copies complete. - Skip intermediate gate/up FP16 projection rounding before SwiGLU in Int4Fc1SwiGluRange while retaining final activation rounding. - Use DoorbellService::WaitForNodeIdle to cancel pending cross-layer promotions under the service mutex. - Apply original-expert-indexed FC2 bias before FP16/BF16 rounding in INT3 offload (both GPU LaunchInt3Moe and CPU Int3Fc2Range) and suppress duplicate finalizer bias addition. - Restore NVFP4 streaming promotion fallback to resident_promotions_ when max_promotions == 0. - Add test coverage for FP16 INT4 parity, INT4/INT3 CUDA Graph capture and replay, FC2 bias cancellation, and SwiGLU FP16 intermediate projection overflow. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Baiju Meswani (baijumeswani)

This is a useful extension, but I do not think it is ready for approval yet.

The PR adds INT4 and INT3 expert weights to CPU/GPU QMoE offload, so more quantized MoE models can use the hybrid path without keeping every expert on the GPU. The stronger shape checks, budget-aware streaming, staging-buffer reuse, and FP16 tests are good improvements. The latest changes also address the Windows environment-variable build issue and the missing promotion-counter checks.

I reviewed ea4b66e with two independent reviewers. Before approval, please fix the async INT3 bias mapping, keep promotion-copy source buffers alive on every error path, and reject overridable expert banks before moving their defaults to CPU. The three inline comments explain these concerns.

Fresh validation included CPU tests, GPU bias reproductions and corrected controls, 24 admission checks under UBSan, 1,200 layout comparisons, and 960 budget-policy cases. GPU memcheck reported zero errors. These checks do not replace a full current-provider/model run or real CUDA Graph qualification; those remain unverified in this review.

CI failures: - Fix the arm64 -Werror=unused-variable in Int4Fc1SwiGluRange: the 'use_f16c = false' fallback is dead when AVX2 is not compiled in. - Exclude MoE offload from minimal builds. Guard the new virtual IExecutionProvider::SupportsMoeOffload() and its overrides with !ORT_MINIMAL_BUILD so no vtable slots or relocations are added (Android minimal-build binary size check). Review feedback: - INT3 async path: pass a null fc2_bias_route_ids. Async rows are already indexed by original route; hit_ids holds token indices, so with top_k > 1 the second expert picked up the first expert's FC2 bias. - Pinned bounce buffer: do not mark it idle when a wait fails, hold it through every exit path of PromoteResident, and quarantine it (keeping the budget charge) when copy completion cannot be confirmed. - Placement: verify each expert bank is a constant initializer with GraphViewer::GetConstantInitializer before relocating it to CPU, so initializers that are also graph inputs stay on GPU. 'required' mode now fails early with a clear message. - CodeQL: drop the unused local variable assignments in the test. Add regression tests for the async INT3 FC2-bias routing and for overridable expert initializers. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

ORT_QUICK_BUILD does not compile the native BF16 INT4/INT8 QMoE runner, so creating a BF16 INT4 session fails with: failed to construct MoE runner for quant_type='int'. The quick-build CUDA CI jobs (Windows CUDA, Windows CUDA Plugin EP) hit this in test_qmoe_offload_v4.py. - Probe once in setUpClass whether a BF16 INT4 session can be created and skip the checks that need one when it cannot. - Split the INT4/INT3 monolith at the BF16 INT4 prepacked sections (prepacked parity, mixed INT4+INT3, unused global scales, INT4 CUDA graph replay) so the INT3 and FP16 INT4 coverage still runs in quick builds. - In the mxfp4 rejection test, skip only the BF16 INT4 success sub-case; the required-mode rejections still run. - Pick a seed for the INT3 test whose T=264 hybrid-stream check keeps about 30% headroom below its L2 bound. That check is a heavy-tailed max over per-token errors, so arbitrary seeds land between 0.7x and 1.4x of the bound. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Create the reference and CUDA graph sessions inside the replay helper so they are released when it returns. This drops the trailing del statements that CodeQL flagged as unnecessary, while still releasing model file handles before the temporary directory is removed. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The wheel registers its bundled CUDA plugin EP at import. setUpClass then unregistered it and registered the byte-identical copy from ORT_CUDA_PLUGIN_PATH. The first copy stays mapped, so the CUDA runtime saw every kernel twice ("Duplicate entry kernels") and later launches failed with "invalid argument" or hung until the job timeout in the Linux CUDA Plugin EP CI. Reuse the registered library when it is the same file or an identical copy, and only replace it when it differs. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The wheel registers its bundled CUDA plugin EP at import. In the plugin CI the bundled copy is repaired after the build, so it is not byte-identical to ORT_CUDA_PLUGIN_PATH and the content comparison never matched. The test then unregistered the bundled copy and loaded a second one. The first copy stays mapped, so the shared CUDA runtime sees every kernel twice and kernel launches fail with "invalid argument" or hang. Reuse any registered CUDA plugin EP and only register ORT_CUDA_PLUGIN_PATH when none is registered. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The Windows GPU Doc Gen build enables the EP API adapters, so CMake defines ORT_USE_EP_API_ADAPTERS=1 for every target, including the in-tree CUDA provider. The QMoE offload code used that macro to select the adapter-only OpKernelInfo::GetKernelInfo() and OpKernelContext::ParallelFor(), which do not exist on the in-tree types, so the in-tree CUDA provider failed to compile (C2039) in that configuration. Only the CUDA plugin target defines BUILD_CUDA_EP_AS_PLUGIN, and that target also defines ORT_USE_EP_API_ADAPTERS=1. Keying the four sites on BUILD_CUDA_EP_AS_PLUGIN leaves the plugin build and the standard in-tree build unchanged and fixes the combined configuration. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…cution (#33622) Stacked on #33575. INT4 offload prefill can use an SM80 CUTLASS runner for resident and streamed routes, with up to 4,096 routes per chunk. Workspace is sized for eligible invocations and bounded by the configured working-memory limit; smaller chunks or GEMV handle constrained budgets. Streaming respects the configured maximum wave size and preserves the base branch's byte-preserving host cache, scale transposition, and DMA lease lifetimes. Hybrid decode supports eight-token doorbell execution and fuses parallel chunked FileExpertStore miss loading with CPU expert work. A single scope guard owns acquisition-failure cleanup. Validation includes dispatch-verified FP16/BF16 parity across the 4,096-route chunk boundary, bounded-cache streaming, scratch-budget reduction and fallback, configured wave limits, and eight-token hybrid CUDA graph replay with resident hits and cold misses. Verified: CUDA 13.0 Release build on H200; all 26 QMoE offload tests across focused runs; `lintrunner -a` and `git diff --check`. Inherited admission and warmup assertions were updated for the current base policy. ### PR stack Merge in this order: 1. #33575 - INT4/INT3 QMoE CPU/GPU offload (base: `main`) 2. #33622 - INT4 offload prefill/decode and S=8 doorbell optimizations (base: #33575) 3. #33629 - Grouped INT3 offload prefill and S=8 streaming optimizations (base: #33622) 4. #33731 - Tensor-core INT3/INT4 offload decode/prefill (base: #33629) Prerequisites #33452, #33450, #33689, and #33650 have merged into `main`. Retarget each dependent PR to `main` after its predecessors merge. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…hort-circuit, and S=8 streaming optimizations (#33629) ### Description Stacked on #33622 (with CPU AVX2 INT3 kernels moved to #33450): 1. **CUDA INT3 8-Route Grouped Offload Prefill (`qmoe_kernels.{h,cu}`, `qmoe_offload.cc`)**: Adds `BuildInt3ExpertTilesKernel`, `FusedFc1GroupedFastKernel`, and `Fc2GroupedFastKernel` in `LaunchInt3Typed` and wires `int3_grouped_ws` in `QMoEOffload::ComputeShared` for `total_routes >= 32`. 2. **Zero-Cold Doorbell GPU Short-Circuit (`qmoe_kernels.{h,cu}`, `qmoe_offload.cc`)**: Adds `skip_zero_cold_doorbell` GPU short-circuit and device-memory `QMoEOffloadTelemetry` (co-authored with `@BBaditi`, #33624). 3. **Multi-Threaded Pinned Wave Staging (`qmoe_offload.cc`)**: Adds parallel `SpinPool` host-to-pinned wave staging (`wave_w1_pinned`, `wave_w2_pinned`, `wave_s1_pinned`, `wave_s2_pinned`) and overlaps `run_cpu_tail_experts()` with GPU streaming waves. ### PR stack Merge in this order: 1. #33575 - INT4/INT3 QMoE CPU/GPU offload (base: `main`) 2. #33622 - INT4 offload prefill/decode and S=8 doorbell optimizations (base: #33575) 3. #33629 - Grouped INT3 offload prefill and S=8 streaming optimizations (base: #33622) 4. #33731 - Tensor-core INT3/INT4 offload decode/prefill (base: #33629) Prerequisites #33452, #33450, #33689, and #33650 have merged into `main`. Retarget each dependent PR to `main` after its predecessors merge. --------- Co-authored-by: Bhuvaneswar Baditi <bbaditi@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…K single-tile streaming for Qwen3.8 RC5 (#33731) > **Stacked on #33629** (`tlwu/20261026/qmoe-offload-avx2-int3-s8-opt`, head `8fa39307bd`). Retarget to `main` after the offload stack merges. > > Rebased on 2026-11-01 by replaying only this PR's last three commits. The former prerequisite PRs #33689 and #33650 are now in `main`; their old integration merges are not carried forward. The **Files changed** tab now shows only this PR's own changes: [base to head](8fa3930...2aca198) (6 files, +345/-85). `git range-diff` confirms all three patches are unchanged. > > | Previous commit | Rebased commit | What | > |:--|:--|:--| > | `c24eb4e60b` | `e84cc7c960` | Wire offload-style INT3 launches into the tensor-core and vectorized kernels | > | `3c505d4032` | `54fd82b8ef` | Enable Tensor-Core INT3/INT4 QMoE offload decode/prefill and 8K single-tile streaming for Qwen3.8 RC5 | > | `ec62a2799b` | `2aca198e7f` | Force the max-shared carveout only on SM 12.x | ### Description  This PR optimizes `QMoEOffload` prefill and decode execution for the Qwen3.8-Flash-Next RC5 (`com.microsoft::RouterGateUpProj` + `DFlash2` $S=8$ speculative decoding) INT4 (`candA`) and mixed INT4/INT3 (`top08_44_33_b64`) packages on CUDA: 1. **Tensor-Core INT3 Offload Decode & Grouped Prefill (`qmoe_int3_tc.cuh`, `qmoe_kernels.cu`, `qmoe_kernels.h`)**: - Extends `int3_tc::Int3TcKernel` (`mma.sync.m16n8k16`, rolled weight pipeline, subnormal `fma.rn.bf16x2` dequantization) to support `QMoEOffload` decode (`source_route_ids`, `input_is_gathered`, `total_routes_override`, `params.output == nullptr`) and grouped prefill waves (`TryLaunchBf16Grouped` consuming `BuildInt3ExpertTilesKernel<8>`'s `sorted_routes` and `Int3ExpertTileDesc* tile_descs` for arbitrary `total_routes >= 32`). - Enables `FusedFc1FastVecKernel` and `Fc2RouteFastVecKernel` (`DequantizeInt3x8Bf16`) for offload calls when `routes < 33`, with upfront `expert < 0` cold-route skip. 2. **Tensor-Core INT4 Offload Decode (`moe_gemv_tc.cuh`)**: - Extends `MoeGemvTcInt4Kernel` and `TryLaunchInt4Bf16` to support unsorted route-to-expert maps (`expert_first_token_offset == nullptr` for `expanded_num_rows <= 256`) using in-kernel warp-ballot (`__ballot_sync`) grouping of up to 8 hit routes per resident expert into a single `mma.sync.m16n8k16` CTA. 3. **8K Single-Tile GPU Prefill & Zero-CPU-Tail Pinned Staging (`qmoe_offload.cc`)**: - _These `qmoe_offload.cc` changes are inherited from the #33629 offload stack, not part of this PR's own three-commit diff._ - Raises default `stream_rows_` for `int4_bank_mode_` / `int3_bank_mode_` from `2048` to `8192` (`max_wave_routes = 2048` for `use_int4_runner` so 8K GPU scratch fits in `458 MiB < 512 MiB`), executing 8K prefill chunks (`chunk = 8192` in RC5) in 1 GPU streaming pass per layer instead of 4 passes. - Sets `eff_stream_min_rows = 1` for `tile_rows >= 2048` and omits `host_input` / `host_result` pinned allocations and `D2H` activation copies, preserving `102.4 MiB` of pinned budget for 2-bank INT3 (`60.9 MiB`) and 1-bank INT4 (`49.7 MiB`) parallel pinned wave staging within the `128 MiB` `pinned_bytes` cap while keeping `eff_stream_min_rows = 8` for short/medium prompts (`L < 2048`). - Lowers the `int4_bank_mode_` CUTLASS `runner_->runMoe` prefill threshold from `num_r >= 256 && num_r >= num_exp * 8` to `num_r >= 64 && num_r >= num_exp * 4` and overlaps 1-bank INT4 `run_cpu_tail_experts()` at `wave == 0` with PCIe H2D DMA. 4. **Vectorized `SplitActivationVec8Kernel` (`router_gate_up_proj_impl.cu`) & MSVC Plugin Build Fix (`onnxruntime_providers_cuda_plugin.cmake`)**: - Adds 128-bit `uint4` vectorized router/gate/up split + SwiGLU epilogue for `RouterGateUpProj` generic prefill path (`rows > 8`) when `router` and `inter` are multiples of 8 and 16-byte aligned. - Adds `/bigobj` for `CXX` compilation in `onnxruntime_providers_cuda_plugin.cmake` to prevent MSVC `C1128` on `cudnn_flash_attention.cc`. ### Motivation and Context  In `QMoEOffload`, resident GPU hits during $S=8$ decode and streaming prefill waves previously bypassed the RC5 Tensor-Core INT3 (`Int3TcKernel`) and INT4 (`MoeGemvTcInt4Kernel`) kernels, and 8K prefill chunks (`chunk = 8192`) were split into 4 tiles of `2048` rows that streamed cold experts 4 times per layer over PCIe. #### Historical restack notes (previous #33689-based branch) The notes and H200 measurements below describe the previous integration base (`5593c1ff75`) and original PR head (`ec62a2799b`), not a new benchmark of the #33629-based branch. The commit mapping above identifies the equivalent rebased patches. - Item 3 (`qmoe_offload.cc`: 8K single-tile prefill, zero-CPU-tail pinned staging, lower CUTLASS threshold) is no longer in this PR's own diff. It is already in the offload chain that is merged into this branch (`default_stream_rows` 8192 for the INT bank modes, `eff_stream_min_rows`, `num_r >= 64 && num_r >= num_exp * 4`), and `qmoe_offload.cc` is unchanged by the three commits above. - `qmoe_kernels.{cu,h}` (item 1: `Int3ExpertTileDesc` moved to the header, the INT3 decode selector no longer returns the CUDA-core `kFast` kernels for offload calls, `LaunchInt3Typed` calls `int3_tc::TryLaunchBf16Grouped`, `expert < 0` early-out in the `FastVec` kernels) is now a separate commit on top of #33689 + the chain. - Conflict resolution in `moe_gemv_tc.cuh`: #33689's compact 80-route INT4 path (`kCompact80`, only used with a sorted `expert_first_token_offset`) is kept. This PR's warp-ballot grouping is used when the route-to-expert map is unsorted (`expert_first_token_offset == nullptr`). - Not part of this stack: the SM80 direct-INT4 path (`weights_are_sm80`, `Sm80Packed`, `shared_int4_`, `direct_shared_layout_`) and the removal of the `PromoteFileMappingToSharedReadOnly` call in `QMoE::QMoE` (`moe_quantization.cc`). They were in the RC5 commit of the old integration branch (`tlwu/qmoe_offload_rc5_base`) but are not part of this PR's description. #### Operator micro-benchmark on H200 (added when this PR was restacked onto #33689) **Summary** (H200, BF16, shapes of `qwen_flash_next_int4_int8_rc5_contrib_dflash2_qmoe_top16_44_33_b64`; *base* = `5593c1ff75`, i.e. #33689 + offload chain + #33650, *this PR* = `ec62a2799b`): - **INT3 decode (`S = 8`)**: expert GEMM kernels 203.8 → 43.6 µs per layer (**4.7×**); ORT-level latency with all experts on the GPU −71%. - **INT4 decode (`S = 8`)**: 93.0 → 34.1 µs per layer (**2.7×**); ORT-level −52%. - **INT3 prefill (512 / 2048 / 8192 rows)**: expert GEMM kernels **7.0×** faster; change of the ORT-level layer latency: - all experts on the GPU (`off_all`): -73% / -76% / -77% - 19% of the experts in the GPU bank, the rest on the CPU (`off_cpu`): -14% / -38% / -68% - 60% of the experts in the GPU bank, the rest on the CPU (`off_cpu`): -14% / -55% / -71% - cold experts streamed over PCIe (`off_gpu`, PCIe-bound): -1% / -1% / -1% - **INT4 prefill**: unchanged (±1%): the CUTLASS runner path is not touched. - **Cases bound by the CPU cold path or PCIe** (`off_cpu` decode, `off_gpu` prefill) show no measurable change; the realized gain scales with the GPU hit rate. - **`RouterGateUpProj`**: the vectorized split kernel is 2.3× faster at 8192 rows (node latency −12%) but about 1.5 µs/node slower at 9 – 64 rows. - **Carveout**: forcing the maximum shared-memory carveout for the INT3 tensor-core kernels costs +10% (decode) / +15% (prefill) kernel time on H200, so the last commit applies it only on SM 12.x. - **Numerics**: `resident` is bit-identical; at `S = 8` the offload output is now bit-identical to the resident output; everything else differs from base by bf16 rounding noise. **Compared.** *base* = this branch without the three commits above, *this PR* = `ec62a2799b`. Both CUDA plugins were built with the same toolchain and flags (Release, CUDA 13.0, `sm_90a` cubins) and were run on an H200; the RTX 5090 of the validation below was not available for this run. The *this PR* plugin used for the runs was built from an earlier iteration of this stack; `cuobjdump -sass` of its kernels is identical to a plugin built from `ec62a2799b` (687 functions in the cubins that hold `Int3TcKernel`, `MoeGemvTcInt4Kernel`, `SplitActivationVec8Kernel`, `FusedFc1*`, `Fc2*Fast*`, `BuildInt3ExpertTiles` and `moe_gemv`: 0 differences). **Workload.** Shapes and attributes of the `QMoE` and `RouterGateUpProj` nodes of the model: hidden 2560, intermediate 640, 512 experts, top-k 10, SwiGLU, BF16 activations, INT4 (block 32) and INT3 (block 64) weights. Each benchmark model has two `QMoE` layers with random weights and uniformly random routing (kernel time does not depend on the values) and is run through ORT with the RC5 `QMoEOffload` options. `S = 8` decode is `T = 8` rows. Decode runs under CUDA graph (offload: `T ≤ 8`, resident: `T ≤ 16`; offload with more rows cannot be captured), prefill runs eagerly. Cells are **ms per layer** (median over 200 iterations for `T ≤ 16`, 60 / 30 / 20 iterations for 512 / 2048 / 8192 rows, divided by the 2 layers), shown as *base → this PR (change)*. **Modes.** | Mode | Meaning | |:--|:--| | `resident` | ordinary `QMoE`, no offload (control: no change expected; the output is bit-identical) | | `off_all` | `QMoEOffload` with every expert in the GPU bank, so only GPU kernels run (isolates this PR) | | `off_cpu` | offload with ~19% or ~60% of the experts in the GPU bank, the rest executed on the CPU. The measured GPU hit rate of the decode routes is 18–20% and 59–60% (`ORT_QMOE_SHARED_CPU_TIMING`); prefill streams to the GPU and does not count | | `off_gpu` | ~19% bank, cold experts streamed over PCIe (`cpu_execution=disabled`) | ##### GPU kernel time per layer (Nsight Systems, `off_all`) | Case (`off_all`) | expert GEMM/GEMV kernels per layer, base (µs) | this PR (µs) | speedup | this PR with the max-shared carveout forced on every GPU, as in the original commit (µs) | |:--|--:|--:|--:|--:| | INT3 decode, S=8 (80 routes) | 203.8 | 43.6 | 4.67× | 48.1 | | INT4 decode, S=8 (80 routes) | 93.0 | 34.1 | 2.73× | - | | INT3 prefill, 512 rows | 6,526 | 932.1 | 7.00× | 1,078 | | INT3 prefill, 2048 rows | 26,172 | 3,748 | 6.98× | 4,322 | | INT3 prefill, 8192 rows | 104,627 | 14,943 | 7.00× | 17,319 | | INT4 prefill, 512 rows | 2,685 | 2,676 | 1.00× | - | | INT4 prefill, 2048 rows | 10,729 | 10,638 | 1.01× | - | | INT4 prefill, 8192 rows | 42,351 | 42,341 | 1.00× | - | The INT3 grouped prefill kernels went from `FusedFc1GroupedFastKernel` 2,060 µs and `Fc2GroupedFastKernel` 1,207 µs per call to `Int3TcKernel<true,4,2,3>` 276 µs and `Int3TcKernel<false,4,1,5>` 192 µs. The INT4 prefill path (CUTLASS `MoeFCGemm` plus `moe_gemv`) is not touched by this PR and is identical within 1%. ##### ORT latency per layer Median per-layer latency of the whole `QMoE` node in ms, *base → this PR (change)*. `†`: the per-round medians of one variant differ by more than 25% across the three rounds of the multi-token matrix. `§`: replaced by a single-case measurement (one process per run, variant order alternating, 3–4 repetitions; min–max in the last table), see "Noise" below. `off_gpu` is only measured for prefill. <details> <summary>INT3 (32 of the 48 layers of the model)</summary> | T (rows) | resident (control) | `off_all` (all hit) | `off_cpu`, ~19% hit | `off_cpu`, ~60% hit | `off_gpu` (PCIe streaming) | |---:|:--|:--|:--|:--|:--| | 1 | 0.032 → 0.032 (+1%) | 0.062 → 0.043 (-30%) | 0.164 → 0.165 (+1%) | 0.121 → 0.118 (-3%) | - | | 8 | 0.054 → 0.054 (±0%) | 0.224 → 0.065 (-71%) § | 0.674 → 0.694 (+3%) § | 0.363 → 0.358 (-1%) § | - | | 16 | 0.348 → 0.344 (-1%) | 0.410 → 0.088 (-79%) | 5.838 → 5.862 (±0%) § | 5.304 → 5.318 (±0%) § | - | | 512 | 1.648 → 1.647 (±0%) | 7.714 → 2.106 (-73%) § | 17.82 → 15.31 (-14%) § | 9.331 → 8.061 (-14%) § | 28.63 → 28.30 (-1%) | | 2048 | 2.607 → 2.609 (±0%) | 29.67 → 7.201 (-76%) | 32.93 → 20.29 (-38%) § | 28.92 → 13.14 (-55%) § | 113.2 → 112.3 (-1%) | | 8192 | 6.466 → 6.470 (±0%) | 117.7 → 27.54 (-77%) | 93.25 → 29.64 (-68%) § | 91.88 → 26.63 (-71%) § | 448.0 → 445.7 (-1%) | </details> <details> <summary>INT4 (16 of the 48 layers of the model)</summary> | T (rows) | resident (control) | `off_all` (all hit) | `off_cpu`, ~19% hit | `off_cpu`, ~60% hit | `off_gpu` (PCIe streaming) | |---:|:--|:--|:--|:--|:--| | 1 | 0.038 → 0.038 (±0%) | 0.042 → 0.042 (±0%) | 0.219 → 0.223 (+2%) § | 0.110 → 0.109 (-1%) | - | | 8 | 0.047 → 0.047 (±0%) | 0.116 → 0.056 (-52%) | 0.741 → 0.747 (+1%) § | 0.387 → 0.382 (-1%) § | - | | 16 | 0.161 → 0.161 (±0%) | 0.205 → 0.085 (-58%) | 5.907 → 5.868 (-1%) § | 5.305 → 5.272 (-1%) § | - | | 512 | 1.008 → 1.008 (±0%) | 4.038 → 4.016 (-1%) § | 19.11 → 18.96 (-1%) | 9.925 → 9.948 (±0%) § | 40.92 → 40.93 (±0%) | | 2048 | 1.822 → 1.826 (±0%) | 14.66 → 14.63 (±0%) | 29.00 → 28.95 (±0%) | 17.35 → 17.19 (-1%) | 158.7 → 158.2 (±0%) | | 8192 | 4.551 → 4.545 (±0%) | 57.30 → 56.87 (-1%) | 38.89 → 38.47 (-1%) § | 29.19 → 28.53 (-2%) | 624.2 → 630.9 (+1%) | </details> <details> <summary>Single-case re-measurements (one process per run)</summary> | Case | rows | base, ms/layer: median (min–max) | this PR, ms/layer: median (min–max) | Δ | runs per variant | |:--|---:|:--|:--|--:|--:| | INT3 `off_all`, all hit | 8 | 0.224 (0.224–0.225) | 0.065 (0.065–0.066) | -70.9% | 3 | | INT3 `off_all`, all hit | 512 | 7.714 (7.706–8.059) | 2.106 (2.098–2.237) | -72.7% | 4 | | INT3 `off_cpu`, ~19% hit | 8 | 0.674 (0.671–0.686) | 0.694 (0.674–0.709) | +2.9% | 3 | | INT3 `off_cpu`, ~19% hit | 16 | 5.838 (5.765–6.320) | 5.862 (5.810–5.871) | +0.4% | 3 | | INT3 `off_cpu`, ~19% hit | 512 | 17.82 (17.69–17.83) | 15.31 (15.29–15.41) | -14.1% | 3 | | INT3 `off_cpu`, ~19% hit | 2048 | 32.93 (32.66–33.34) | 20.29 (20.26–20.38) | -38.4% | 3 | | INT3 `off_cpu`, ~19% hit | 8192 | 93.25 (92.84–93.85) | 29.64 (29.46–29.69) | -68.2% | 3 | | INT3 `off_cpu`, ~60% hit | 8 | 0.363 (0.363–0.365) | 0.358 (0.357–0.362) | -1.5% | 3 | | INT3 `off_cpu`, ~60% hit | 16 | 5.304 (5.267–5.783) | 5.318 (5.288–6.213) | +0.3% | 3 | | INT3 `off_cpu`, ~60% hit | 512 | 9.331 (9.263–9.360) | 8.061 (8.024–8.078) | -13.6% | 4 | | INT3 `off_cpu`, ~60% hit | 2048 | 28.92 (28.67–28.97) | 13.14 (13.08–13.82) | -54.6% | 3 | | INT3 `off_cpu`, ~60% hit | 8192 | 91.88 (91.51–92.34) | 26.63 (26.52–26.68) | -71.0% | 3 | | INT4 `off_all`, all hit | 512 | 4.038 (4.037–4.265) | 4.016 (4.016–4.287) | -0.5% | 3 | | INT4 `off_cpu`, ~19% hit | 1 | 0.219 (0.219–0.223) | 0.223 (0.220–0.225) | +2.0% | 3 | | INT4 `off_cpu`, ~19% hit | 8 | 0.741 (0.736–1.556) | 0.747 (0.744–0.777) | +0.9% | 3 | | INT4 `off_cpu`, ~19% hit | 16 | 5.907 (5.871–6.361) | 5.868 (5.860–6.403) | -0.7% | 3 | | INT4 `off_cpu`, ~19% hit | 8192 | 38.89 (38.85–38.95) | 38.47 (38.34–38.54) | -1.1% | 3 | | INT4 `off_cpu`, ~60% hit | 8 | 0.387 (0.384–0.393) | 0.382 (0.382–0.390) | -1.2% | 3 | | INT4 `off_cpu`, ~60% hit | 16 | 5.305 (5.299–5.316) | 5.272 (5.260–5.324) | -0.6% | 3 | | INT4 `off_cpu`, ~60% hit | 512 | 9.925 (9.896–9.939) | 9.948 (9.892–9.985) | +0.2% | 4 | </details> **Reading the numbers.** - **Decode with all experts on the GPU (`off_all`, `S = 8`)**: before this PR an offloaded layer could not use the tensor-core kernels for decode (INT3: always the CUDA-core `kFast` kernels; INT4: no support for the unsorted route-to-expert map). It now runs the same kernels as the resident path: INT3 203.8 → 43.6 µs per layer (4.7×), INT4 93.0 → 34.1 µs (2.7×), ORT-level −71% / −52%, and the output is bit-identical to the resident output. - **Prefill, INT3**: the grouped tensor-core kernels are 7.0× faster at every size. How much of that shows in the layer latency depends on how much of it is GPU kernel time: −73% … −77% when every expert is on the GPU, −14% (512 rows) to −71% (8192 rows) when part of the experts is on the CPU, and −1% when the layer is bound by PCIe streaming (`off_gpu`: 28 / 113 / 448 ms per layer at 512 / 2048 / 8192 rows). - **Prefill, INT4**: unchanged (±1%), as expected: the CUTLASS runner path is not touched. - **Decode with experts on the CPU (`off_cpu`)**: the layer latency is set by the CPU cold path (about 0.4 ms per layer at 60% hit and 0.7 ms at 19% hit at `S = 8`, against 0.06 ms for `off_all` with this PR), so the faster GPU hits are hidden: all `off_cpu` decode cells are within ±3%. Both variants print the pre-existing "GPU memory was allocated during CUDA graph capture" warning in this mode. - **Upper bound for a decode cycle** of the 48-layer model (32 INT3 + 16 INT4 layers) when every route hits the GPU bank: 32 × (203.8 − 43.6) + 16 × (93.0 − 34.1) µs ≈ 6.1 ms. The realized gain scales with the GPU hit rate (decode hit rates are 18 – 20% and 59 – 60% for the 19% and 60% banks). - **Noise**: in the multi-token matrix (all token counts of a model kind in one process) the 512-row and some CPU-bound decode cells were bimodal from process to process (about +3 – 5 ms per layer at 512 rows, ×2 on single decode cells), independent of the variant. Those cells are marked † and were re-measured in single-case processes (§); the median is reported and the min–max is in the last table. The absolute latency of `off_cpu` cases that run the CPU cold path eagerly also depends on the process: at 16 rows a fresh single-case process takes 5.3 – 5.9 ms per layer, whereas the same case after the 1- and 8-row cases in one process took 1.2 – 1.8 ms (both variants alike), and the table shows the fresh-process value for every 16-row `off_cpu` cell. Base and this PR are always compared in the same kind of process, so the relative change is not affected. The absolute saving of this PR on INT3 `off_all` at 512 rows is the same in every run (5.6 ms per layer). ##### `RouterGateUpProj` (`SplitActivationVec8Kernel`, 48 nodes in the real model) | rows | latency per call incl. 2 Cast nodes, base → this PR (µs; CUDA graph for ≤ 16 rows) | split + SwiGLU kernel, base → this PR (µs) | all GPU kernels per call (GEMM + split + casts), base → this PR (µs) | |---:|:--|:--|:--| | 1 | 18.9 → 19.5 (+3%) | none (fused `RouterGateUpProjKernel`, unchanged) | 8.6 → 8.6 (±0%) | | 8 | 21.7 → 21.7 (±0%) | none (fused `RouterGateUpProjKernel`, unchanged) | 9.4 → 9.3 (-1%) | | 9 | 22.7 → 24.7 (+9%) | 1.44 → 3.11 | 10.9 → 12.4 (+14%) | | 16 | 23.8 → 25.3 (+6%) | 1.47 → 3.18 | 10.9 → 12.5 (+15%) | | 64 | 42.7 → 45.0 (+5%) | 1.66 → 3.13 | 11.4 → 12.9 (+13%) | | 512 | 68.8 → 67.8 (-1%) | 4.56 → 3.87 | 19.6 → 19.1 (-3%) | | 2048 | 112.4 → 107.6 (-4%) | 14.31 → 7.79 | 58.2 → 52.0 (-11%) | | 8192 | 265.6 → 233.2 (-12%) | 59.12 → 25.50 | 218.1 → 184.6 (-15%) | Outputs are bit-identical to the previous kernel at every row count. The GEMM (`nvjet_*`) time is unchanged; the vectorized split + SwiGLU kernel is latency-bound at small row counts (about 1.5 µs slower per node for 9 – 64 rows, break-even between 64 and 512 rows) and 2.3× faster at 8192 rows (59.1 → 25.5 µs per node). Decode (`rows ≤ 8`) never takes this path: it runs the fused `RouterGateUpProjKernel`, which is unchanged (so the +3% at 1 row is measurement noise). A row-count guard between 64 and 512 rows (for example `rows >= 256`; the break-even was not measured more finely) would keep the old kernel for the small chunks; this PR does not change that. ##### Carveout (`ec62a2799b`) `EnsureInt3TcCarveout()` asks for the maximum shared-memory carveout for the raw INT3 tensor-core kernels on every GPU. On H200 the default carveout already gives the same CTAs/SM (4 for FC1, 5 for FC2), and forcing the maximum shrinks the L1 that the weight streams rely on: decode FC1 + FC2 48.1 µs vs 43.6 µs per layer (+10%), grouped prefill +15–16% (last column of the kernel table), INT3 `off_all` at 512 rows 4.5 → 4.2 ms for two layers (−7%). `ec62a2799b` therefore applies the carveout only on SM 12.x (`ComputeCapabilityMajor == 12`), where it keeps the behavior of the original commit. SM 12.x was not measured here; if the carveout is not needed there either, the commit can simply be dropped in favor of removing the call. ##### Numerics Output of this PR against base, relative L2 `‖new − base‖ / ‖base‖` of the `[T, 2560]` BF16 output for identical inputs (random weights, so the output magnitudes are large and only the relative error is meaningful): | Case | INT3 | INT4 | |:--|:--|:--| | `resident`, any `T` | bit-identical | bit-identical | | `off_all`, `T` = 1 | bit-identical | bit-identical | | `off_all`, `T` = 8 / 16 | 6.5e-4 / 4.4e-4 | 1.9e-3 / 1.9e-3 | | `off_all`, `T` = 512 / 2048 / 8192 | 4.6e-4 / 3.7e-4 / 3.9e-4 | bit-identical | | `off_cpu`, `T` = 1 / 8 / 16 | bit-identical / 9.8e-8 / 3.8e-5 | bit-identical / 1.4e-3 / 1.7e-3 | | `off_cpu`, `T` = 512 / 2048 / 8192 | 4.7e-4 / 3.7e-4 / 3.9e-4 | bit-identical / 1.3e-4 / bit-identical | | `off_gpu`, `T` = 512 / 2048 / 8192 | 4.6e-4 / 3.7e-4 / 3.9e-4 | 7.9e-4 / 8.5e-4 / 8.5e-4 | Against the `resident` output, base and this PR are equally far away (equal within 0.1e-3) in every other offload case: INT3 prefill 1.6e-3 – 1.7e-3, INT4 prefill 1.6e-3 – 1.9e-3 (3.4e-4 and 4.2e-5 for `off_cpu` at 2048 / 8192 rows, where most experts run on the CPU). These are the bf16 accumulation-order differences between the offload engine and the resident path that existed before, so the differences in the table are of the same size as that pre-existing noise. The exception is `off_all` at `S = 8`, where the offload output of this PR is now **bit-identical to the resident output** (base: 6.5e-4 for INT3 and 1.9e-3 for INT4 away), because decode uses the same tensor-core kernels as the resident path. ##### Validation on the H200 plugin build - `onnxruntime/test/python/transformers/test_qmoe_cuda.py -k "not benchmark and not latency"`: 119 passed, 10 skipped, 2 failed. The two failures are identical on the base and are not related to this PR: `test_int4_fp16_qmoe_cuda_graph` ("CUDA Provider not available": plugin-only build) and `test_s8_e512_k10_fused_routing_and_fallback_parity_1_bf16` (1 of 4096 elements off by 0.00195). - `onnxruntime/test/python/test_qmoe_offload_v4.py`: 15 passed, 1 failed (`test_v4_zero_cold_doorbell_cuda_graph_replay`, "CUDA Provider not available": plugin-only build). - `clang-format` 20.1.8 reports no changes for the six files; `router_gate_up_proj_op_test` was not run (no gtest binary in the plugin-only build, covered by the bit-exact A/B above). ##### Not covered / follow-ups - Only an H200 (SM 9.0) was available. The RTX 5090 numbers of the validation below are the author's; the SM 12.x carveout path is unchanged from the original commit and untested here. - The offload parity tests use `hidden = 256` and therefore never reach the INT3 tensor-core kernels (`hidden % 512 == 0`, BF16, block 64). The only coverage of that path is the bit-exact offload-vs-resident comparison of the benchmark at `hidden = 2560`; a unit test would be a good follow-up. - Random weights and uniform routing: real routing is skewed (hot experts), which changes the hit rate and the tile grouping but not the per-kernel cost. #### Validation & Benchmark Summary (Windows 11 RTX 5090 32 GB + Ryzen 9 9950X 16C/32T AVX-512, 64 GB RAM) - **`smoke.py` (`top08_44_33_b64` & `candA`)**: - Command: `python tools/qwen38_graph/smoke.py <model_dir> --provider <model_dir>/provider/onnxruntime_providers_cuda.dll` - Result: **PASSED** on both models (`1.16 s` warm TTFT / `36.9–48.6 tok/s` on `top08_44_33_b64`; `2.18 s` warm TTFT / `23.6 tok/s` on `candA`, `80.3%` draft acceptance). - **`engine_behaviours.py` (27-Check Qualification Suite)**: - Command: `python tools/qwen38_graph/engine_behaviours.py <model_dir> --provider <model_dir>/provider/onnxruntime_providers_cuda.dll --out <out.json>` - Result: **PASSED (27 / 27 checks)** on both `top08_44_33_b64` (`94.12%` draft acceptance) and `candA` (`93.63%` draft acceptance). - **`bench_ttft_engine.py` (`L = 128, 256, 512, 2048, 8192`, `15.0 GB` Adaptive)**: - Command: `python tools/qwen38_graph/bench_ttft_engine.py <model_dir> --provider <provider.dll> --lengths 128,256,512,2048,8192 --text-pool mmlu_pro_samples.jsonl --reps 3 --out <out.json>` - Result (**PASSED**): - `top08_44_33_b64`: `L=128`: **`1.445 s`** (`1.357 s` warm); `L=256`: **`1.490 s`** (`1.459 s` warm); `L=512`: **`1.808 s`** (`1.725 s` warm); `L=2048`: **`4.701 s`** (`4.639 s` warm); `L=8192`: **`17.162 s` (`13.853 s` warm, `1.66x` faster than RC4 Opt `22.95 s` and `1.41x` faster than RC5 Base `19.58 s`)**. - `candA`: `L=128`: **`1.756 s`** (`1.653 s` warm); `L=256`: **`1.971 s`** (`1.896 s` warm); `L=512`: **`2.471 s`** (`2.192 s` warm, `1.24x` faster than RC5 Base `3.072 s`); `L=2048`: **`6.751 s`** (`6.153 s` warm); `L=8192`: **`27.436 s` (`21.878 s` warm, `1.63x` faster than RC4 Opt `35.68 s` and `1.41x` faster than RC5 Base `36.40 s`)**. - **`bench_engine_mtp.py` & Context Scaling (`2K` / `8K` / `32K` / `64K` / `128K`)**: - Command: `python bench_rc5_adaptive_context.py --model <model_dir> --tiers "8k:2048,8192;32k:2048,8192,32768;64k:65536;128k:131072"` - Result (**PASSED**): - `top08_44_33_b64`: **`99.33 tok/s`** (`58.04–60.89 ms/cyc`) at `ctx=2048`; **`92.91 tok/s`** (`52.57 ms/cyc`, `8.49–19.20 s` TTFT) at `ctx=8192`; **`53.01 s` TTFT (`1.36x` faster than RC4 Opt `71.85 s`) & `62.84 tok/s`** at `32K`; **`118.83 s` TTFT (`1.42x` faster than RC4 Opt `168.52 s`) & `69.88 tok/s`** at `64K`; **`192.12 s` TTFT (`1.57x` faster than RC4 Opt `300.77 s`, `-108.7 s`)** at `128K`. - `candA`: **`59.78–65.35 tok/s`** at `ctx=2048`; **`57.89 tok/s`** (`23.67–33.66 s` TTFT) at `ctx=8192`; **`103.64 s` TTFT (`1.37x` faster than RC4 Opt `141.62 s`) & `40.71 tok/s`** at `32K`; **`239.87 s` TTFT (`1.35x` faster than RC4 Opt `323.58 s`)** at `64K`; **`443.64 s` TTFT (`1.45x` faster than RC4 Opt `643.07 s`, `-199.4 s`)** at `128K`. - **`mmlu_engine_eval.py` (`--shards 80 --shard 0`)**: - Command: `python tools/qwen38_graph/mmlu_engine_eval.py <model_dir> --provider <provider.dll> --prompts mmlu_pro_800_prompts.jsonl --checkpoint <model_dir> --shards 80 --shard 0 --out <out.jsonl>` - Result: **PASSED** (`8/10` on `top08_44_33_b64` in **`46.30 s` wall**, up to **`133.21 tok/s`** decode; `9/10`–`10/10` on `candA` in **`81.16 s` wall**, `-29.3%` faster than RC4 Opt `114.81 s`). ### Restack validation (2026-11-01) - All three original patches preserved exactly (`git range-diff`); six changed files, +345/-85. - CUDA 13.0 Release rebuild of the CUDA provider, operator tests, and Python bindings succeeded (`sm_80;sm_90`). - NVIDIA A100: all 33 focused tensor-core QMoE and RouterGateUpProj C++ tests passed. - NVIDIA A100: all six INT3/INT4 offload Python tests passed, covering numerical parity, bounded streaming, scratch/pinned budgets, wave limits, and eight-token hybrid CUDA graph replay. - `lintrunner` on all six changed files and `git diff --check` passed. This validation used the ordinary CUDA EP on SM80, not the plugin EP, SM12.x, or MSVC. The H200 benchmarks above remain historical measurements of the previous base. ### PR stack Merge in this order: 1. #33575 - INT4/INT3 QMoE CPU/GPU offload (base: `main`) 2. #33622 - INT4 offload prefill/decode and S=8 doorbell optimizations (base: #33575) 3. #33629 - Grouped INT3 offload prefill and S=8 streaming optimizations (base: #33622) 4. #33731 - Tensor-core INT3/INT4 offload decode/prefill (base: #33629) Prerequisites #33452, #33450, #33689, and #33650 have merged into `main`. Retarget each dependent PR to `main` after its predecessors merge. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…ort-circuit for QMoE offload (#33624) ## Motivation & Summary Builds on #33575 (`[CUDA/CPU] Add INT4 and INT3 QMoE CPU/GPU offload support`) with two decode-path optimizations (numerically equivalent to within `~1e-5` for the AVX2 cold-expert kernel, and bitwise identical for the zero-cold doorbell short-circuit) and one C++ test fixture fix: 1. **AVX2 + FMA + F16C cold-expert INT4 kernel (`qmoe_shared_layout.cc`):** On x86-64 hosts with `AVX2+FMA+F16C` that do not support `AVX-512` (such as AMD EPYC Zen 2 on Azure NDv4 A100 VMs and Intel Haswell–Alder Lake), `CpuHasAvx512()` returns `false` and cold-expert INT4 dequant-GEMM and FP16/BF16 row conversions fall back to scalar C++ loops (`46.36 ms/tok` of CPU compute at `25%` GPU residency on 48-layer `Qwen3-30B-A3B`). This PR adds runtime `CPUID` + `XGETBV` detection (`CpuHasAvx2FmaF16c()`) and an `AVX2+FMA+F16C` implementation for FP16/BF16 row conversions and `Avx2Int4GemmRange` (supporting both prepacked SM80 INT4 layout and raw INT4 layout, FP16 and BF16 scales/biases), plus an $O(1)$ lookup table in `ReorderedPosition`. 2. **Zero-PCIe doorbell short-circuit (`session.moe.offload.skip_zero_cold_doorbell`):** During partial-residency hybrid decode ($M \le 4$), when all $M \times k$ selected experts for a layer are already GPU-resident (`cold_count == 0`), the default doorbell path still copies the activation tile and expert indices to pinned host memory and wakes the CPU doorbell worker thread. When `session.moe.offload.skip_zero_cold_doorbell="1"` (or `ORT_QMOE_SKIP_ZERO_COLD_DOORBELL=1`) is enabled under static placement without cross-layer promotion, `QMoEOffloadClassifyKernel` counts cold selections via `__syncthreads_count(is_cold)` before issuing any host writes and, when `cold_count == 0`, advances `doorbell_done_host` and `doorbell_req_host` directly on the GPU (`__threadfence_system()`) so `QMoEWaitAndAccumulateKernel` proceeds without PCIe host staging or CPU worker wakeups. 3. **INT3 CUDA test fixture update (`moe_test.cc`):** Updates `MoETest.QMoECudaInt3TiledLogHasOneCounterUpdate` to match the SwiGLU `QMoE` contract introduced in #33575 (`activation_type = "swiglu"`, `swiglu_fusion = 1`, `fc1` rows `= 2 * inter_size`, and unsigned INT3 zero-point code `4`). ## Key Changes - `onnxruntime/contrib_ops/cpu/moe/qmoe_shared_layout.{h,cc}`: - Add `CpuHasAvx2FmaF16c()`, `CompilerHasAvx2()`, and C ABI `qmoe_shared_cpu_has_avx2()` / `qmoe_shared_set_disable_avx512_for_testing()`. - Add `Avx2HalfBitsToFloatRow`, `Avx2FloatToHalfBitsRow`, `Avx2BFloat16BitsToFloatRow`, `Avx2FloatToBFloat16BitsRow`, `Avx2RoundFp16InPlace`, and `Avx2Int4GemmRange`. - Replace the linear scan in `ReorderedPosition` with a 32-entry constant lookup table. - `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.{h,cu}`, `onnxruntime/contrib_ops/cuda/moe/qmoe_offload.{h,cc}`, `onnxruntime/core/framework/moe_offload_plan.h`, `include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h`: - Add `kOrtSessionOptionsMoeOffloadSkipZeroColdDoorbell` (`session.moe.offload.skip_zero_cold_doorbell`, default `"0"`). - Short-circuit PCIe host staging and CPU doorbell wakeups when `cold_count == 0` and cross-layer promotion is inactive; retire pre-completed sequence numbers cleanly in `DoorbellService`. - Allocate optional `QMoEOffloadTelemetry` (`ORT_QMOE_OFFLOAD_TELEMETRY=1`) in device scratch memory and copy back once in `~QMoEOffload()`. - `onnxruntime/test/contrib_ops/moe_test.cc`, `onnxruntime/test/contrib_ops/qmoe_shared_layout_test.cc`, `onnxruntime/test/python/transformers/test_qmoe_shared_layout.py`, `onnxruntime/test/python/test_qmoe_offload_v4.py`: - Fix `MoETest.QMoECudaInt3TiledLogHasOneCounterUpdate`. - Add C++ and Python unit tests covering `AVX2+FMA+F16C` vs. scalar numerical parity (`max_abs_diff < 1e-5`, including all 65,536 FP16 half patterns and partial-range vector interiors) and zero-cold doorbell bitwise parity, sequence counter transitions, and cross-layer promotion fallback. - `docs/contrib_ops/cuda/moe_qmoe.md`: - Document `AVX-512 -> AVX2+FMA+F16C -> scalar` cold-expert dispatch, `session.moe.offload.skip_zero_cold_doorbell`, and environment variables (`ORT_QMOE_SKIP_ZERO_COLD_DOORBELL`, `ORT_QMOE_OFFLOAD_TELEMETRY`, `ORT_QMOE_CPU_DISABLE_AVX512`). ## Verification & Benchmarks ### Unit & Provider Tests - `onnxruntime_test_all --gtest_filter="*MoETest*:*QMoE*"`: **6 / 6 passed** - `onnxruntime_provider_test --gtest_filter="QMoESharedLayoutTest.*:MoETest.*"`: **91 / 91 passed** - `pytest onnxruntime/test/python/test_qmoe_offload_v4.py onnxruntime/test/python/transformers/test_qmoe_shared_layout.py onnxruntime/test/python/transformers/test_qmoe_cuda.py`: **153 / 153 passed** ### End-to-End 48-Layer `Qwen3-30B-A3B` (`QMoE` INT4, `block_size=32, weights_prepacked=1`) Decode Latency (`ORT_QMOE_OFFLOAD_TELEMETRY=0`, 5 Trials `mean ± stdev`) Measured on 1x NVIDIA A100-SXM4-80GB + AMD EPYC 7V12 (Zen 2, `AVX2+FMA+F16C`, no `AVX-512`), 5 independent trials of 10 warmup + 30 decode steps ($M = 1$, $k = 8$, 128 experts/layer) with `ORT_QMOE_OFFLOAD_TELEMETRY=0` (and a separate `ORT_QMOE_OFFLOAD_TELEMETRY=1` pass to record wakeup/short-circuit counts): | GPU Residency (Experts / Layer) | #33575 Baseline (Scalar Cold CPU, ms/tok) | AVX2 Cold Kernel (`skip_zero_cold=0`, ms/tok) | AVX2 + Zero-Cold Doorbell (`skip_zero_cold=1`, ms/tok) | Speedup vs. #33575 (`skip=0` / `skip=1`) | Zero-Cold Short-Circuits / Token (`skip=1`) | CPU Doorbell Wakes / Token (`skip=0` $\to$ `skip=1`) | Min Logit Cosine vs. Scalar | Top-1 Token Parity | | :--- | :---: | :---: | :---: | :---: | :---: | :---: | :---: | :---: | | **32 / 128 (25%)** | 54.46 ms (18.4 tok/s) | **21.60 ± 0.37 ms** (46.3 tok/s) | **21.89 ± 0.55 ms** (45.7 tok/s) | **2.52x / 2.49x** | 7.78 / 48 | 48.00 $\to$ **40.23** (-16.2%) | 0.999561 | **30 / 30 (100%)** | | **64 / 128 (50%)** | 15.52 ms (64.4 tok/s) | **10.80 ± 0.10 ms** (92.6 tok/s) | **10.66 ± 0.23 ms** (93.8 tok/s) | **1.44x / 1.46x** | 36.65 / 48 | 48.00 $\to$ **11.35** (-76.4%) | 0.999818 | **30 / 30 (100%)** | | **96 / 128 (75%)** | 9.04 ms (110.6 tok/s) | **8.30 ± 0.19 ms** (120.4 tok/s) | **8.56 ± 0.41 ms** (116.8 tok/s) | **1.09x / 1.06x** | 46.98 / 48 | 48.00 $\to$ **1.02** (-97.9%) | 0.999883 | **30 / 30 (100%)** | | **128 / 128 (100%)** | 8.31 ms (120.4 tok/s) | **7.55 ± 0.06 ms** (132.5 tok/s) | **7.61 ± 0.26 ms** (131.4 tok/s) | **1.10x / 1.09x** | — | 0.00 $\to$ **0.00** | 1.000000 | **30 / 30 (100%)** | *(Note: `skip_zero_cold_doorbell="1"` is 100% bitwise identical (`max_abs_diff = 0.0`, `cosine = 1.000000`) to `skip_zero_cold_doorbell="0"` across all residencies. On an uncontended host, single-stream decode latency for `skip=0` vs. `skip=1` is within noise because GPU resident GEMVs overlap the zero-cold doorbell wakeup, while `skip=1` eliminates the 4 KB/layer PCIe host staging writes and reduces CPU doorbell wakeups by 76.4% at 50% residency and 97.9% at 75% residency.)* --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Tianlei Wu (tianleiwu) added this pull request to stack #33495 October 24, 2026 22:17

github-actions Bot reviewed Oct 24, 2026

View reviewed changes

Comment thread onnxruntime/contrib_ops/cpu/moe/qmoe_shared_layout.h Resolved

Comment thread onnxruntime/test/python/test_qmoe_offload_v4.py Resolved

Comment thread onnxruntime/test/python/test_qmoe_offload_v4.py Resolved

Tianlei Wu (tianleiwu) mentioned this pull request Oct 26, 2026

[CUDA] Optimize INT4 QMoE offload prefill/decode and S=8 doorbell execution #33622

Merged

This was referenced Oct 26, 2026

[CUDA/CPU] Add AVX2 INT4 cold-expert kernel and zero-cold doorbell short-circuit for QMoE offload #33624

Merged

[CUDA] Add opt-in tail expert substitution for hybrid QMoE decode #33625

Closed

Tianlei Wu (tianleiwu) added the qmoe_offload label Oct 29, 2026

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch from 05605d2 to a5af696 Compare October 29, 2026 20:55

Copilot AI balanced review requested due to automatic review settings October 29, 2026 20:55

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) October 29, 2026 20:55

View session

Tianlei Wu (tianleiwu) changed the title ~~[CUDA/CPU] Add INT4 and INT3 QMoE CPU/GPU offload support~~ [CUDA] Add INT4 and INT3 QMoE CPU/GPU offload support Oct 29, 2026

Copilot AI reviewed Oct 29, 2026

View reviewed changes

Comment thread onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.cu Resolved

Comment thread onnxruntime/contrib_ops/cuda/moe/qmoe_offload.cc Resolved

Comment thread onnxruntime/contrib_ops/cuda/moe/qmoe_offload.cc Resolved

Comment thread onnxruntime/contrib_ops/cuda/moe/qmoe_offload.cc Resolved

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch from a5af696 to 237bed6 Compare October 29, 2026 23:04

github-actions Bot reviewed Oct 29, 2026

View reviewed changes

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch from 237bed6 to 430bf4a Compare October 29, 2026 23:36

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch 2 times, most recently from 6f480c4 to 60d8510 Compare October 30, 2026 01:06

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch from 60d8510 to f8e4146 Compare October 30, 2026 01:09

This was referenced Oct 30, 2026

[Qwen38] Size the QMoE offload budgets from the context tier and run DFlash2 warmup in offload warmup mode microsoft/onnxruntime-genai#2784

Merged

Add the Qwen3.8-Flash QMoE offload package builder microsoft/onnxruntime-genai#2821

Closed

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch from f8e4146 to e622696 Compare October 30, 2026 06:36

Ti-Tai Wang (titaiwangms) mentioned this pull request Oct 30, 2026

[CUDA] Tensor-core QMoE INT4/INT3 decode, INT8 MatMulNBits GEMV and large-M dequantize + cuBLAS GEMM #33689

Merged

3 tasks done

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch 2 times, most recently from 79e51fb to 3db1426 Compare October 31, 2026 04:27

This was referenced Nov 1, 2026

[CUDA] Add 8-route grouped INT3 offload prefill, zero-cold doorbell short-circuit, and S=8 streaming optimizations #33629

Merged

[CUDA] Enable Tensor-Core INT3/INT4 QMoE offload decode/prefill and 8K single-tile streaming for Qwen3.8 RC5 #33731

Merged

Copilot AI reviewed Nov 1, 2026

View reviewed changes

github-advanced-security Bot found potential problems Nov 2, 2026

View reviewed changes

Comment thread onnxruntime/test/python/test_qmoe_offload_v4.py Fixed

Tianlei Wu (tianleiwu) requested a max review from Copilot November 2, 2026 04:57

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) November 2, 2026 04:58

View session

Tianlei Wu (tianleiwu) and others added 2 commits November 1, 2026 21:11

Tianlei Wu (tianleiwu) force-pushed the tlwu/20261024/qmoe-offload-int4-int3 branch from 797c07e to c3e255a Compare November 2, 2026 05:11

Copilot AI reviewed Nov 2, 2026

View reviewed changes

github-advanced-security Bot found potential problems Nov 2, 2026

View reviewed changes

Comment thread onnxruntime/test/python/test_qmoe_offload_v4.py Fixed

Comment thread onnxruntime/test/python/test_qmoe_offload_v4.py Fixed

Baiju Meswani (baijumeswani) reviewed Nov 2, 2026

View reviewed changes

Comment thread onnxruntime/contrib_ops/cuda/moe/qmoe_offload.cc Resolved

Comment thread onnxruntime/contrib_ops/cuda/moe/qmoe_offload.cc Resolved

Comment thread onnxruntime/core/framework/session_state.cc Resolved

Baiju Meswani (baijumeswani) mentioned this pull request Nov 2, 2026

Add the Qwen3.8-Flash CUDA data plane library and its host loader microsoft/onnxruntime-genai#2726

Merged

3 tasks done

Tianlei Wu (tianleiwu) requested a review from Baiju Meswani (baijumeswani) November 2, 2026 08:12

github-advanced-security Bot found potential problems Nov 2, 2026

View reviewed changes

Comment thread onnxruntime/test/python/test_qmoe_offload_v4.py Fixed

Tianlei Wu (tianleiwu) and others added 3 commits November 2, 2026 10:02

Tianlei Wu (tianleiwu) enabled stacked auto-merge (squash) November 2, 2026 16:15

Baiju Meswani (baijumeswani) approved these changes Nov 2, 2026

View reviewed changes

Tianlei Wu (tianleiwu) merged commit b69bd01 into main Nov 2, 2026

Tianlei Wu (tianleiwu) deleted the tlwu/20261024/qmoe-offload-int4-int3 branch November 2, 2026 22:50

Conversation

Tianlei Wu (tianleiwu) commented Oct 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

PR stack

Uh oh!

github-actions Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

github-actions Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Ti-Tai Wang (titaiwangms) commented Oct 30, 2026

Static review: offload correctness and Olive/Mobius interoperability

Major: INT4 admission and backend disagree about weights_prepacked=-1

Major: zero experts reaches division by zero in per-expert accounting

Smaller follow-ups

Producer/consumer contract

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Baiju Meswani (baijumeswani) left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Reviewers

Uh oh!

Tianlei Wu (tianleiwu) commented Oct 24, 2026 •

edited

Loading

Major: INT4 admission and backend disagree about `weights_prepacked=-1`