microsoft/onnxruntime plugin-ep-cuda/v0.4.0

ONNX Runtime CUDA Plugin EP 0.4.0

on GitHub

This release brings CPU/GPU hybrid expert offloading, native GGUF/block-packed weight execution, Tensor-Core
INT2/INT3/INT4 and offline K-packed NVFP4 QMoE kernels, long-context and speculative-decoding acceleration, and
reliability fixes to the ONNX Runtime CUDA Plugin EP. Highlights include session.moe.offload.* expert offloading for
Mixture-of-Experts models that exceed GPU memory, an opt-in persistent MatMulNBits tactic cache, new fused LLM
operators, ONNX opset 28 kernels, and expanded Linux aarch64 CUDA 13 GPU coverage.

Compatibility

The minimum ONNX Runtime core version remains 1.24.4, but meeting that minimum alone does not guarantee
compatibility for contributed operators. Core and the plugin must use the same contributed-operator schemas;
building both from the same ONNX Runtime revision is recommended. Schema mismatches can cause incorrect execution
or crashes and may not be detected when registering the plugin (#32855).

The matched ONNX Runtime version is 1.32.0 for this release.

• Features that require ONNX Runtime 1.32.0 core: CPU/GPU MoE offload (session.moe.offload.*), which relies on new core session-state support to keep expert initializers on the CPU; pipelined external-weight loading for plugin EPs; and the device memory capacity query (OrtApi::EpDevice_GetMemoryCapacity). Older runtimes ignore the plugin's new OrtEpFactory::GetDeviceMemoryCapacity callback, so it does not raise the 1.24.4 floor (#33451, #33839, #33447).
• New operators and opset 28: Models that use the contributed operators added in this release, or ONNX opset 28 (ONNX 1.23.2), need a core that registers those schemas, such as the matched 1.32.0 release (#33134).
• Not available in the plugin: The opt-in EP-owned runtime fusions added to the bundled CUDA EP in ONNX Runtime 1.32.0 (ep.cuda.runtime_fusion.unit_swish and ep.cuda.runtime_fusion.causal_conv_with_state) are not offered by CUDA plugin builds (#33757, #34021).
• Capacity-aware partitioning: With ONNX Runtime 1.32.0 core and session.resource_cuda_partitioning_settings, the core declines to assign GroupQueryAttention nodes to the plugin because route-aware workspace estimates cannot yet pass through the plugin accounting path. A warning is logged and those nodes are left for other EPs. Sessions without capacity-aware partitioning are unaffected (#33016).

Highlights

CPU/GPU Hybrid MoE Offload

• Added portable, session-scoped CPU/GPU expert offload for standard com.microsoft::QMoE nodes, configured entirely through session.moe.offload.* session options with no model changes. session.moe.offload.mode="auto" keeps every expert on the GPU when it fits and otherwise enables hybrid offload; "required" fails session creation when a node cannot be admitted. The planner budgets per-layer GPU expert slots (optionally guided by an expert-profile.bin popularity profile via session.moe.offload.popularity_file), runs cold experts on the CPU during decode through a doorbell path that works with CUDA Graph replay, streams cold experts to the GPU in waves during prefill, and promotes newly hot experts between steps with the default adaptive placement (#33451, #33452).
• Supported NVFP4 (weights_prepacked=1), SM80-prepacked INT4 (block_size=32) and INT2, and raw INT3 expert banks, including mixed INT4/INT3/INT2 models, on SM80 and newer GPUs, with Tensor-Core INT3/INT4 offload decode and grouped prefill (#33575, #33622, #33629, #33731, #33787).
• Extended hybrid offload to GGUF block-packed MoEPacked experts, with a three-stage pipelined prefill and multi-token cold-route execution (#33976, #33973, #34067, #34114).
• Added AVX-512 and AVX2+FMA+F16C cold-expert CPU kernels that read the same INT4/INT3/INT2/FP4 weight layout as the GPU without repacking, and memory-mapped expert storage with an optional bounded host cache (session.moe.offload.host_expert_bytes) for hosts with limited RAM (#33450, #33624, #33785, #33935).
• Improved offload on consumer GPUs and Windows with WDDM-aware VRAM budgeting that avoids spilling into shared GPU memory, Windows memory-mapping and pinned-memory accounting fixes, working-set retention for hot experts, ready-first scheduling of cached cold experts, and an opt-in zero-cold doorbell short-circuit (session.moe.offload.skip_zero_cold_doorbell=1) (#33916, #33935, #33699, #33624, #33629).
• Added opt-in approximate routing for offloaded decode: tail-expert pruning below a routing-probability threshold (session.moe.offload.min_expert_weight; by default only CPU-cold tail experts are pruned) and substitution of cold tail experts with GPU-resident experts (session.moe.offload.tail_substitution_threshold and session.moe.offload.tail_substitution_min_rank). Both are disabled by default and change routing results, so validate model accuracy before enabling them (#33916, #33944).
• Added session.moe.offload.gpu_device_mem_limit, a per-session cap on all device allocations made through the session's CUDA allocators (unlike gpu_mem_limit, which caps only the BFC arena), and session.moe.offload.budget_group for sharing one budget between target and draft sessions on the same GPU. The limit is not supported with external CUDA allocators or CUDA memory pools (#33451).
• Added ep.cuda.qmoe_in_warmup and ep.cuda.qmoe_prefill_new_turn run-config entries so applications can exclude warm-up runs from routing statistics and reset prefill state at the start of a conversation turn (#33452, #33622).

Quantized MoE Kernels

• Added support for offline K-packed NVFP4 QMoE weights (weights_prepacked=1), shared by the GEMV, grouped-GEMM, and dense fallback paths, with a register-decoded decode GEMV, single-CTA fused softmax/top-k routing for 256 and 512 experts, a single-pass finalize kernel, and programmatic dependent launch (PDL) overlap on SM90 and newer. Native E4M3 block-scale conversion is available as an opt-in with ORT_NVFP4_GEMV_NATIVE_SCALE=1 (#33094, #33354, #33366, #33545, #33323).
• Enabled fused block-scaled FP8 QMoE execution by default on SM80 and newer (ORT_ENABLE_FP8_FUSED=0 disables it), and vectorized the block-FP8 decode GEMV and prefill weight dequantization (#33094, #33817).
• Added fast INT3 and mixed-width (FC1, FC2) ∈ {2, 3, 4} × {2, 3, 4} QMoE kernels, offline SM80-prepacked INT2/INT4 weights (weights_prepacked=1), and BF16 Tensor-Core decode GEMVs for INT4, INT3, and INT2 experts. The Tensor-Core GEMVs are enabled by default; ep.cuda.disable_moe_gemv_tc=1 (INT4/INT2) and ep.cuda.disable_qmoe_int3_tc=1 (INT3) restore the SIMT kernels, and ep.cuda.moe_gemv_tc_min_rows and ep.cuda.qmoe_int3_tc_min_routes tune the dispatch thresholds (#33569, #33585, #33689, #33714, #33786).
• Added opt-in grouped BF16 INT3 prefill (ep.cuda.qmoe_int3_grouped_prefill=1), with optional Tensor-Core compatibility rounding (ep.cuda.qmoe_int3_tc_compatibility=1) (#34117).
• Extended fused routing to multi-token batches (S ≤ 8) with up to 512 experts, and added profile-free prefill tactic selection for deterministic MoE (ORT_FORCE_DETERMINISTIC_MOE=1), including INT4 with an Ampere tile fallback (#33596, #33367, #33595).
• Added tanh-GELU gating (activation_type="geglu") to MoE and integer QMoE, opt-in post-SiLU clipping (swiglu_limit_after_activation), and fixed gated GELU with separate FC3 weights (#33633, #33954, #33883).
• Shared successful MoE GEMM tactic-profiling results across identically shaped layers in the same process, reducing session-creation time for models with many identical MoE layers. Set ORT_DISABLE_MOE_SHARED_TACTIC_CACHE=1 to profile each node separately (#34169).
• Fixed MXFP4 QMoE initialization on SM80, allowed SM120 QMoE configurations on SM121 devices, and capped split-K SwiGLU GEMV partial buffers so prefill workspace no longer grows with prompt length (#32194, #33225, #34168).

Native GGUF and Block-Packed Weights

• Added GatherPacked, MatMulPacked, and MoEPacked, which consume block-packed quantized weights directly in their native layout without offline repacking or separate scale tensors. Supported formats include GGUF IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ4_NL, IQ4_XS, Q2_0, Q4_0, Q5_0, Q5_K, Q6_K, and Q8_0, plus F16, BF16, and F32 (#33953).
• Extended MatMulPacked with an NVFP4 global_scale, BF16 precision handling, and fused activation epilogues (activation_type: relu2/relu_squared or silu). Decoded weights are rounded to the activation type by default; ep.cuda.matmul_packed_allow_unrounded_gemv=1 lets eligible GEMV kernels skip that rounding (#33962, #34147).
• Extended MoEPacked with full-score top-k routing (separate selection scores and mixing weights), ReLU² experts, and fused routing normalization and scaling (routing_epsilon, routed_scaling_factor). Added MoEGroupRouter, a fused group-limited sigmoid router whose full-score outputs feed MoEPacked (#33963, #33817, #33994, #34147).
• Optimized packed projections and experts for multi-token verification (M = 2-8), large-chunk prefill of up to 4,096 rows, and the Q6_K, IQ3_XXS, IQ4_NL, IQ2_XXS, and NVFP4 formats (#34002, #34003, #34066, #34164, #34200).

Attention and KV Caching

• Added an optional query_seqlens input to GroupQueryAttention for right-padded, fixed-shape chunked prefill. Padded tokens and idle batch entries (query_seqlens[b] = 0) are skipped across the FlashAttention, memory-efficient, and unfused paths (#33503, #33832).
• Added q_norm_epsilon and k_norm_epsilon attributes so GQA pre-norm fusion preserves distinct Q and K normalization epsilons (#33363).
• Added FP16/BF16 head-size-512 GQA XQA decode kernels with local-window and attention-sink support (#33176).
• Added paged speculative XQA for BF16 queries with INT8 caches at head size 256 (GQA group size 6 or 16), and head-size-128, group-size-6 specializations for FP16/BF16 queries with INT8 caches and BF16 queries with per-channel INT4 caches, including INT4 single-token decode (#33538, #34117).
• Optimized INT8-cache PagedAttention prefill (window-aware gathering, vectorized dequantization, and architecture-qualified FlashAttention tiling) and decode, and added the ep.cuda.paged_attention_enable_xqa (default 1) and ep.cuda.paged_attention_enable_xqa_native_kv (default 0) session switches (#34117, #33817).
• Fused the small-batch PagedAttention sequence-length sanitizer into a single kernel, and added session.paged_attention.skip_input_sanitizer=1 (or ORT_PAGED_ATTENTION_SKIP_SANITIZER=1) for runtimes that already validate block_table, cumulative_sequence_length, and past_seqlens. Skipping sanitization is unsafe if those inputs can be invalid (#33597).

Sparse Attention and Long Context

• Added PackedSparseAttentionIndexerMerge for multi-token-prediction index sharing. Extended PackedSparseAttentionIndexer with index-sharing controls (indexshare_mode, max_output_entries, and indexshare_state_policy) and an optional state_status output that reports, per request, whether a state update was applied or rejected, so CUDA Graph-replayed loops can detect rejected steps (#33224, #33383).
• Added mixed-precision keys (BF16 activations with FP32 past_key_state), fused S ≤ 8 QSA refresh and selection, compact hierarchical TopK, and warp-specialized ragged prefill to PackedSparseAttentionIndexer, and optimized packed QSA scoring and multi-token verification (#33472, #33328).
• Added symmetric per-channel INT8 KV caches with per-batch committed_length validation to DynamicSparseAttention, plus FP16 KV caches and FP16 query/key inputs for BF16 models (#33472, #33997).
• Added Tensor-Core INT8-KV DynamicSparseAttention prefill and split-K decode kernels. Behavior change: these kernels are now enabled by default and accumulate in a different order, so results are not bit-identical to earlier releases. Set ep.cuda.dsa_prefill_int8_tc=0 or ep.cuda.dsa_decode_int8_tc=0 to restore the previous kernels. An FP32 fallback prevents earlier tiles from being lost when large value_scale values trigger rescaling (#33707, #33802, #33847).
• Optimized dense prefill for hybrid sparse-attention/linear-attention NVFP4 models, and tuned DynamicSparseAttention, indexer, hyper-connection, and QMoE kernels for Blackwell SM 12.x (100 KB shared memory per SM) and contexts of up to 256K tokens (#33367, #33802).

Linear Attention, State-Space Models, and Speculative Decoding

• Added Kimi Delta Attention to GatedDeltaNet (per-key-dimension decay with gate_activation="sigmoid", alias "kda"), with warp-specialized decode and split-chunked prefill. FP16 KDA prefill checks for state overflow by default; ep.cuda.gated_delta_net_fp16_state_guard=0 skips the check when the state is known to stay within FP16 range (#33817).
• Extended GatedDeltaNet and CausalConvWithState to speculative decoding with up to 8 tokens, packed projections, and compact state_update capture (state_update_capacity) (#33471).
• Added ep.cuda.gated_delta_net_bf16_precision for BF16 GatedDeltaNet chunked prefill: fp32 (default) preserves FP32 state precision, while opt-in fp16 (with range guards that fall back to the FP32-preserving path) and bf16 trade precision for speed. Chunked prefill also uses ldmatrix-based MMA loads and a faster GatedRMSNorm kernel (#33367, #33523, #33601).
• Extended VarlenCausalConvWithState with dynamic grouped two-tap convolutions and a faster tiled prefill (#34117, #33817).
• Added SpeculativePrefixCommit for single-launch, device-side speculative prefix verification and state commit across 3-8 verification rows, and Mamba2Verify for snapshot-aware Mamba2 verification (#33474, #33969).

New Fused LLM Operators

• Added SwiGLU for dense feed-forward and shared-expert paths; ONNX Runtime 1.32.0 core also fuses the decomposed SiLU-gated pattern into this operator (#32126).
• Added RMSNormRotary, which fuses RMS normalization with partial rotary embedding for sparse-attention query/key projections, with FP32 rotary caches and FP16 output (#33469, #33997).
• Added HyperConnectionLowRankMix, which fuses a hyper-connection block and its low-rank gate network into one node, with S ≤ 8 decode and fused prefill paths (#33384, #33548, #33738).
• Added RouterGateUpProj, which fuses the MoE router projection with the shared expert's gate/up projection and SwiGLU (#33650).
• Extended GatedAdd with an optional projected gate weight (gate_weight) (#33470).

MatMulNBits

• Added an opt-in, hardware-keyed persistent tactic cache for MatMulNBits fpA_intB GEMM. Set ep.cuda.gemm_tactic_cache_dir or ep.cuda.gemm_tactic_cache_prefix (or ORT_CUDA_GEMM_TACTIC_CACHE_DIR / ORT_CUDA_GEMM_TACTIC_CACHE_PREFIX), or set ep.cuda.gemm_tactic_cache_to_model=1 to store the cache next to the model. ep.cuda.gemm_tactic_cache_mode selects load_save (default), load, or save. Cached tactics are loaded and validated at session creation and saved at EP teardown; no file I/O occurs unless a location is configured (#29588, #33078).
• Behavior change: When deterministic compute is enabled, MatMulNBits selects a stable tactic (GEMV for small M, otherwise a fixed non-split-K GEMM) instead of the profiled choice, bypassing the tactic cache (#32898).
• Added a Tensor-Core symmetric INT8 MatMulNBits GEMV for M ≤ 8 on SM90 (ep.cuda.disable_fpa_intb_gemv_tc=1 disables it). For INT8 weights without zero points or bias, SM90 now by default dequantizes the weights once and runs a single cuBLAS GEMM when M ≥ 512; ep.cuda.fpa_intb_dequant_gemm_min_m adjusts the threshold, and 0 disables the path (#33689).
• Added profiler-selected INT4 group-32 M=1 decode geometries with checked profiler scratch sizing, extended the M=8 narrow-tile GEMV to BF16 activations and INT8 weights, and added an opt-in tail-safe INT4 path for N % 8 == 4 with M ≤ 4 (ORT_MATMULNBITS_UNALIGNED4_FAST_PATH=1) (#33177, #33600, #33299).

Standard ONNX Operators

• Added ONNX opset 28 support: native SpaceToDepth (DCR and CRD modes), floating-point floor-modulus Mod, signed-integer BitShift, and BFloat16 Compress, OneHot, and ReverseSequence, with opset 28 version splits for Cast, DepthToSpace, and Einsum (#32362, #32369, #32378, #33134, #33647).
• Added integer BitwiseAnd, BitwiseOr, BitwiseXor, and BitwiseNot kernels (opset 18 and later) and unsigned-integer BitShift for opsets 11-27 (#33645, #33646, #33647).
• Added BFloat16 support to ArgMax, ArgMin, and Pad (#33470, #34062).

Plugin Runtime and General Performance

• With ONNX Runtime 1.32.0 core, external initializers are loaded to the GPU through a multi-threaded, double-buffered pinned-staging pipeline. Tune it with session.external_data_loader.threads, session.external_data_loader.pinned_memory_bytes (default 32 MiB), session.external_data_loader.chunk_bytes (default 2 MiB), and session.external_data_loader.min_size_bytes (default 16 MiB); models below the minimum size use the pageable-memory path (#33839).
• Implemented OrtEpFactory::GetDeviceMemoryCapacity so applications can query total and currently available device memory through OrtApi::EpDevice_GetMemoryCapacity with ONNX Runtime 1.32.0 or later (#33447).
• Cached up to eight cuDNN 9 frontend Conv execution plans per kernel, and passed the cached SM version to cuDNN frontend graphs to avoid repeated device-property queries when shapes change (#33123, #33770).
• Zeroed the cuDNN SDPA graph-cache key before hashing to avoid duplicate cuDNN graph builds, and added a guarded reciprocal-square-root fast path for FLOAT Pow with exponent -0.5 (#34081, #33662).

Reliability and Correctness

• Validated PagedAttention sequence-length and block-table values, PackedAttention token offsets, and Attention projection and cache dimensions before use (#31651, #33330, #32946).
• Behavior change: Integer Mod now fails with Integer modulo by zero when a divisor is zero, for both fmod=0 and fmod=1, consistent with the CPU EP. Constant-initializer divisors are checked once when the kernel is created. Other integer divisors are checked on every call with a stream synchronization, which is not allowed during CUDA Graph capture; such models must use a constant-initializer divisor or disable CUDA Graphs (#32546).
• Fixed GroupQueryAttention fast decode ignoring explicit rotary positions, and XQA multi-block reduction races in shared-cache GQA decoding (#32951, #33703).
• Fixed variance cancellation in fused SkipLayerNormalization and EmbedLayerNormalization, ReLU² handling, RMSNormalization rounding, and numerical stability in Softplus, Sigmoid, QuickGelu, and SiLU paths (#33357, #33312).
• Fixed GatherBlockQuantized block indexing for non-block-aligned quantized axes and Split with num_outputs and an empty optional split input, and made MatMulNBits reject empty B weights with a clear error instead of reading an empty buffer during prepacking (#32665, #32972, #33559).

Build and Packaging

• Expanded the Linux aarch64 CUDA 13 packages to sm_87 (Jetson Orin), sm_100/sm_103 through the 100f family target, and sm_110 (Jetson Thor), and replaced the sm_120/sm_121 pair with the 120f family target. CUDA architecture gating now recognizes family-specific and other valid spellings, so SM-specific kernels are no longer compiled out (#33369, #33315, #33649, #33301, #34089).
• When SM80 is targeted, FlashAttention is no longer built separately for SM86/SM89 by default (onnxruntime_FLASH_ATTENTION_REUSE_SM80=ON); those GPUs run the SM80 binary, reducing package size (#34088).
• Added native sm_120 SASS to the plugin's LLM kernel target in Windows (MSVC) builds, fixing CUDA error 222 (unsupported PTX version) for fpA_intB kernels on RTX 50-series GPUs with older drivers and avoiding first-run JIT compilation (#33173).
• Bundled LICENSE and ThirdPartyNotices.txt in the NuGet package and set its project URL (#33405).
• Fixed CMAKE_CUDA_ARCHITECTURES cache handling, Windows CUDA builds with contrib operators disabled, a Windows NVCC error in the FP4 GEMV, and system cuDNN include-directory handling in plugin builds, and updated Eigen to 5.0.1 (#28019, #33702, #33554, #33736, #32475).

Contributors

Thanks to the 37 human contributors who contributed to this release:

@adrastogi,
@alexk201,
@apsonawane,
@arnej27959,
@baijumeswani,
@BBaditi,
@cboseak,
@cbourjau,
@chilo-ms,
@ConnorBaker,
@cursoragent,
@DevChuriwala,
@eserscor,
@GopalakrishnanN,
@gramalingam,
@gyagp,
@hanbitmyths,
@haoxiz-nvidia,
@javier-intel,
@jmcarcell,
@justinchuby,
@KafKafrnZ,
@keshavv27,
@kunal-vaishnavi,
@LinCanNerd,
@maxim-davgalev,
@mustjab,
@qjia7,
@rui-ren,
@sayanshaw24,
@selenayang888,
@silviot,
@tianleiwu,
@titaiwangms,
@Trisha-Sarvam,
@TrishaWadhwani,
@xadupre

This release covers changes from origin/plugin-ep-cuda/rel-0.3 to origin/plugin-ep-cuda/rel-0.4 affecting CUDA Plugin
EP code, shared CUDA kernels, build integration, and packaging. The highlights focus on user-facing changes rather
than listing every infrastructure-only update.

This summary was drafted with AI assistance from commit history and PR metadata.