Sync with Microsoft ONNX Runtime - 28072026 - #1233
Merged
Merged
Conversation
…CompInt8 path (microsoft#29842) ### Description Adds `QuantizeARow_CompInt8_Fp16_avx512`, the fp16-input twin of `QuantizeARow_CompInt8_avx512`: it loads 16 fp16 elements and widens them with `_mm512_cvtph_ps`, then runs the identical max-abs, scale, round, and block-sum arithmetic. It uses no intrinsic beyond what the float quantizer already uses on this tier, so its output (block int8 data, per-block scale, and `scale * sum(a_i)`) is bit-for-bit what the float quantizer produces from the same values converted fp16 to fp32 first. It is wired into both `MlasSQNBitGemmDispatchAvx512` and `MlasSQNBitGemmDispatchAvx512vnni`; the VNNI table reuses the same function, the way it already reuses the float quantizer. `test_sqnbitgemm_fp16_quant_a.cpp` already covers this: it self-gates on `MlasQNBitGemmFp16DirectQuantASupported()`, so with this change it runs on AVX-512 and checks that the fp16 `A` path gives byte-identical GEMM output to the fp32-first path across 2, 4, and 8 bit, block lengths 16 through 256, M through 129 (the threaded and the single-threaded quantize paths), symmetric and asymmetric, with and without bias. ### Motivation and Context microsoft#29791 let the CompInt8 `MatMulNBits` path quantize a fp16 `A` in one pass instead of first materializing a full fp32 copy of `A`, but wired it only for the AVX2 and AVX-VNNI dispatch tables. On AVX-512 and AVX-512-VNNI hosts `QuantizeARowComputeBlkSum_CompInt8_Fp16` was left null, so `MlasQNBitGemmFp16DirectQuantASupported()` returned false and the fp16 path fell back to converting all of `A` to fp32 with `MlasConvertHalfToFloatBufferInParallel` into a separate `M*K` float buffer before the float quantizer ran. This wires the same single-pass path for AVX-512, so those hosts skip that extra full conversion of `A` and the fp32 `A` allocation.
…LASS) (microsoft#29811) ### Description Pilot for Phase A of the memory roadmap (microsoft#29775): give the CUDA EP a way to know a `MatMulNBits` node's transient CUTLASS scratch ("workspace") size *before* `Run()`, instead of relying on the accountant's blanket 1.5x heuristic that never sees the workspace term. Scoped to one kernel, in-tree only. The workspace formula is extracted into **one shared, stateless helper** that all three call sites funnel through, so the numbers are identical by construction: - **Shared formula** (`fpA_intB_gemm.h`): `ComputeFpAIntBGemmWorkspaceSize(m, n, k, sm, multi_processor_count)` — pure arithmetic, no device/instance state. Overflow-checked end-to-end with `SafeInt<size_t>` (both SM90 and non-SM90 branches, including ceil-div numerators); returns `nullopt` on overflow rather than throwing. The 5 tile constants are promoted `protected` → `public` so the free function can read them. `CutlassFpAIntBGemmRunner::getWorkspaceSize` now delegates to it — byte-for-byte identical on valid inputs. - **Eligibility single-source-of-truth** (`matmul_nbits.{h,cc}`): `CheckFpAIntBEligibility` + `EffectiveFpAIntBWorkspaceSm` centralize the multi-part fpA_intB path decision (dtype, option/env gate, nbits/block/alignment, SM, prepacked/SM90 constraints). The constructor is refactored to call it (behavior-preserving), so Level 1 can never disagree with the runtime path. - **Level 1 — partition-time** (`EstimateMatMulNBitsWorkspace`, wired into `CUDAExecutionProvider::GetCapability()`): a plain op-type-dispatched free function (no virtual, no registration) that estimates from node attributes alone, before any kernel instance exists. Exposed via a slim forward-declaring header; **log-only** for this pilot (does not alter the accountant's budget number). Both the call site and the `#include` are double-guarded with `#if !defined(DISABLE_CONTRIB_OPS) && defined(USE_FPA_INTB_GEMM)`. - **Level 2 — instance-level** (`MatMulNBits::DeclareWorkspaceRequirements`): new default no-op `virtual` on `OpKernel`. Because `matmul_nbits.cc` compiles into two hierarchies, the no-op is mirrored on both the in-tree and the plugin adapter `OpKernel` (following the `PrePack` precedent); the real CUTLASS-backed override is in-tree only (`#ifndef BUILD_CUDA_EP_AS_PLUGIN`). It feeds the **same effective arch** the runner resolves after `setArch()` (`FpAIntBPackingSmForKernel()`, not raw `sm_`), so it equals the runtime request when the queried shape matches. - **Boundary-safe struct**: `WorkspaceRequirement` (`size_bytes`, `slot_id`) lives in a new lightweight `core/framework/workspace_requirement.h` so both `OpKernel` hierarchies can include it without the adapter pulling in `core/framework/op_kernel.h`. This pilot uses one slot; the `InlinedVector`/`slot_id` generality is reserved for later multi-buffer kernels (e.g. Attention). ```cpp // getWorkspaceSize is now a thin delegate over the shared formula: auto ws = ComputeFpAIntBGemmWorkspaceSize(m, n, k, sm_, multi_processor_count_); ORT_ENFORCE(ws.has_value(), "fpA_intB workspace size overflow for m=", m, " n=", n); return *ws; ``` Unit tests cover the pure helpers: formula exact values (non-SM90 + SM90), overflow/negative-dim `nullopt`, `EffectiveFpAIntBWorkspaceSm` (drift guard A), and `CheckFpAIntBEligibility` (drift guard C). ### Motivation and Context The workspace size is currently computed live inside `ComputeInternal()` (`getWorkspaceSize(m, n, k)`), so the partition-time budget accountant has zero visibility into it and falls back to a 1.5x multiplier on `initializer + output` bytes — which both under- and over-estimates, since the CUTLASS scratch is fp32 and unrelated to output size. `MatMulNBits` was chosen as the pilot because its formula is real, non-trivial, kernel-instance-independent, and single-path, and the op dominates quantized-LLM projection cost. Notes for reviewers: - **ABI**: adding a `virtual` to in-tree `OpKernel` changes vtable layout — a known consideration for out-of-tree EPs compiled against older headers; the pilot does not solve it but flags it so it isn't discovered silently. - **Not addressed by design**: the plugin EP's separate budget loop still uses the 1.5x heuristic; substituting the estimate into the actual accept/reject decision (roadmap option (b)) and a real plugin-side override are follow-ups. - CUDA could not be built/run in this environment (no GPU); formula values were validated with a standalone host program and all changed files are clang-format clean. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: ade1ab12-1137-46bb-a92b-e730a6f28bf4
… build + track QNN materialize deps (microsoft#29577) Follow-up to the merged microsoft#29504 (ONNX node-test detachment). This PR consolidates **all** post-microsoft#29504 follow-up work into a single change set. Changes C/D/E address @tianleiwu's non-blocking post-merge review comments; A/B were the original scope of this PR. None change the behavior of the shipped feature — this is hardening + hygiene + doc detachment. ### (A) C# — loud-fail guard on the macOS node-test symlink `csharp/test/Microsoft.ML.OnnxRuntime.EndToEndTests/runtest.sh` symlinked the ONNX on-disk node-test corpus. Once the onnx pin advances past onnx/onnx#7959 (which deletes that corpus) the symlink target vanishes and the suite would silently pass with zero cases. Added a loud-fail guard: enforce a discovered-case floor (1000) and keep the immutable v1.13.1 pin, so corpus absence becomes a red instead of a silent-green. ### (B) JS — post-microsoft#7959 opset caveat in `prepare-onnx-node-tests.ts` Documented that once onnx ships without the on-disk node-test data (onnx/onnx#7959), post-microsoft#7959 opsets cannot use the `rel-*` archive download path — a caveat for the next opset bump. ### (C = N1) cmake — drop no-op ctest `DEPENDS` + document real ordering `cmake/onnxruntime_unittests.cmake`: the two `set_tests_properties(... PROPERTIES DEPENDS onnx_node_tests_materialized)` calls named an `add_custom_target(... ALL)`, not a registered test, so a CTest test-level `DEPENDS` on it is a silent no-op. Removed both and documented the real ordering guarantee: the `ALL` target materializes the corpus during the build (before ctest), and the `-m` / `--min-cases` tripwire fails loud on a missing/partial corpus. ### (D = N2) conf.py — hermetic in-memory doc model `docs/python/conf.py` urlretrieved `.../node/test_sigmoid/model.onnx` from GitHub at doc-build time — a latent 404 once onnx/onnx#7959 deletes that path. Replaced with an in-memory `onnx.helper` single-node Sigmoid graph (input `x` → output `y`, float32 `[3,4,5]`) written to the same destination. Keeps the doc build hermetic and drop-in compatible. ### (E = N3-5) requirements file for QNN/Android materialize deps + SKILL.md The four QNN/Android CI legs used inline `pip install onnx==... "numpy==...; ..."` pins (invisible to Dependabot) to feed the `onnxruntime_MATERIALIZE_ONNX_NODE_TESTS` gate. Added a dedicated minimal `requirements-materialize-onnx-node-tests.txt` (onnx + numpy only, with hybrid numpy env markers `2.2.6` for `python_version<'3.11'`, `2.4.2` for `>='3.11'`) and repointed all four legs to `pip install -r` (preserving `--user` on linux + android; the two Windows legs omit it). Placed at the repo root because `dependabot.yml` pins pip to `directory: "/"` (non-recursive). Pins stay locked in lockstep with the CMake gate, which HARD-FATALs on `installed != pinned`. Updated the `onnx-opset-bump-checklist` SKILL.md gotcha-p to reflect the new file (reversing its prior "inline by design" note) and added it to the numpy pin sync-site list. --- Consolidates the standalone microsoft#29746 (now closed) per author request. Triple-review passed on both change sets independently. Merge after microsoft#29504 (already merged), which this references. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
…nt and per-channel scales (microsoft#29900) ### Description Extends the XQA decode kernel to serve the quantized (INT8/FP8) KV-cache configurations that previously fell back to *dequantize-the-whole-cache-then-Flash-Attention*. Four changes, in `contrib_ops/cuda/bert/`: 1. **Attention sinks on quantized XQA.** `use_xqa_attention_sinks` no longer requires `!is_inputs_quantized`. The sink is folded into the softmax row sum by the same code the non-quantized path uses, and the dequant scale is already applied to the scores before the row max/sum, so sink and score live in the same domain — true for both the single-block and the multi-block (Flash Decoding) reduction, where the sink is added to the merged row sum. 2. **Independent `k_scale` / `v_scale`.** XQA previously required the two scale *tensors to be the same pointer*, which no builder-emitted model satisfies. The kernel now keeps them apart: `k_scale` → `qkScale` (applied to Q·Kᵀ), `v_scale` → `voScale` (applied to the P·V accumulator). 3. **`PER_CHANNEL` scales folded out of the kernel.** XQA only accepts a *scalar* dequant scale. But dequantization is linear and the channel index is the *contracted* dim for K and the *free* dim for V: ``` scores_t = sum_d q_d * (k_td * sk_d) = sum_d (q_d * sk_d) * k_td out_d = sum_t p_t * (v_td * sv_d) = (sum_t p_t * v_td) * sv_d ``` so the per-channel scales can be moved **out** of the kernel exactly: pre-scale Q by `sk`, post-scale the attention output by `sv`. `p_t` is unchanged because softmax only ever sees the (already correct) scores, and every step after the P·V accumulation is linear, so attention sinks, sliding window and the multi-block reduction all stay valid. XQA reads a null scale pointer as "scale = 1", which is how the caller signals it folded the scale itself. The Q side costs nothing extra: `UnpackRoPEAppend` already loads Q, applies RoPE and stores it to the very scratch buffer XQA reads, so it takes a `q_fold_scale` argument and applies the multiply to the registers it is about to store. Only the V side needs its own pass, since it consumes XQA's output. Both are `O(num_heads * head_size)` — **O(1) in context length**, versus the O(context) full-cache dequant they replace. 4. **Vectorized / length-aware dequant fallback** for what XQA still cannot serve (INT4, unsupported `head_size`/group size). The fallback dequant kernel now uses 32-byte vector loads, hoists per-channel scales out of the sequence loop, maps (batch, head) to `blockIdx.y` so the inner addressing is 32-bit, and walks only the rows inside the sequence instead of the padded cache capacity. `docs/contrib_ops/cuda/gqa.md` is updated to match: quantized-decode description, XQA eligibility conditions, backend-selection table, and the parity-test list. ### Motivation and Context The quantized KV-cache decode path was **~2.8x slower than an FP16 KV cache** at 16k context, and the gap grew with context length because the fallback rebuilt the entire dequantized K+V cache on *every* decode step. The binding gate was attention sinks: a model carrying `head_sink` sets `use_smooth_softmax`, which combined with a quantized cache disqualified XQA regardless of scale granularity, head size or group size. Scale-granularity workarounds do not help — all quantized configurations previously landed within 0.5% of each other, because none of them reached XQA. Decode latency, batch 1, 64 generated tokens, gpt-oss-20b geometry (24 layers, `head_size` 64, group 8), H200 / SM90, median of 5, idle GPU: | KV cache | 1k | 4k | **16k** | vs FP16 KV @16k | |---|--:|--:|--:|--:| | FP16 *(control)* | 2.38 -> 2.370 | 2.44 -> 2.434 | 2.57 -> 2.573 | — | | INT8 per-tensor | 3.22 -> 2.421 | 4.04 -> 2.490 | **7.42 -> 2.626** | +2.1% | | INT8 per-channel | 3.21 -> 2.476 | 4.07 -> 2.542 | **7.42 -> 2.683** | +4.3% | | FP8 per-tensor | 3.18 -> 2.390 | 4.08 -> 2.449 | **7.39 -> 2.561** | −0.5% | | FP8 per-channel | — -> 2.443 | — -> 2.492 | — -> **2.611** | +1.5% | (ms/token, `before -> after`. The FP16 control is unchanged, so the comparison is sound.) Quantized KV decode goes from **−66% to −4%** versus an FP16 cache at 16k, while keeping its **−26% TTFT** (749 ms vs 1017 ms @16k). Scale granularity now costs ~0.05 ms/token (2% at 16k), so per-channel — by far the most accurate INT8 option — no longer has to be traded away for decode speed. For configurations XQA still cannot serve, change 4 alone gives 7.44 -> 4.61 ms/token @16k (1.61x); its residual overhead over the FP16 non-XQA Flash path is now 0.44 ms/token (was 3.40), i.e. roughly 75% of theoretical HBM bandwidth for the traffic it must move. ### Testing - `onnxruntime/test/python/transformers/test_gqa.py`: **865 passed**, 384 of them new. Three new classes, each asserting *both* numerics and that the op actually reaches XQA (`SdpaKernel == "XQA"`), so a silent fallback fails the test rather than passing quietly: - `TestXQAQuantizedHeadSinkParity` (96) — quantized decode with an attention sink. - `TestXQASeparateKVScaleParity` (144) — distinct `k_scale` / `v_scale` tensors. - `TestXQAPerChannelScaleParity` (144) — per-channel K/V scales folded into Q and the output. Combined coverage: INT8 + FP8 x fp16 + bf16 x `head_size` 64/128/256 x group 4/8 x global / sliding-window / multi-block (past 512) x sink on/off x rotary on/off x prepacked vs runtime-converted sink. - `onnxruntime_provider_test --gtest_filter='*GroupQueryAttention*'`: 63 passed. - The `q_fold_scale` fusion in change 3 was verified to be a numerical no-op: op-level max-abs and RMS error against the fp32 reference are bit-identical to the unfused two-pass version across `head_size` 64/128 x past 4/512. - End to end, a gpt-oss-20b build with an INT8 per-channel KV cache produces a **token-identical** 128-token greedy continuation to the same model with an FP16 KV cache. - `lintrunner` clean.
…embler support (microsoft#28767) ### Description On toolchains where the assembler lacks AVX-VNNI support (e.g., RHEL 9.x with GCC 11 + binutils < 2.36), MLAS AVX2 builds fail with `inlining failed ... target specific option mismatch` and `no such instruction: vpdpbusd`. **CMake (`cmake/onnxruntime_mlas.cmake`)** - Replace the flawed compiler-version gate (`CMAKE_CXX_COMPILER_VERSION VERSION_GREATER "11"`, which incorrectly includes GCC 11.x since `"11.3.1" > "11.0.0"`) with a `check_cxx_source_compiles` probe that tests actual `-mavxvnni` support end-to-end (compiler + assembler). `-mavxvnni` is only added when the probe succeeds. **Header files** (`sqnbitgemm_kernel_avx2_int8_blklen{16,32,64}.h`, `sqnbitgemm_m1_sym_kernel_avx2_int8_blklen32.h`) - Replace `#if !defined(__GNUC__) || (__GNUC__ > 10)` with `#if !defined(__GNUC__) || defined(__AVXVNNI__)`. `__AVXVNNI__` is defined by GCC and Clang only when `-mavxvnni` is active, making this a reliable feature check rather than a version heuristic. Non-GCC compilers (MSVC) are unaffected since `__GNUC__` remains undefined for them. ### Motivation and Context RHEL 9.x ships GCC 11 with binutils that predate AVX-VNNI assembler support. The version-based heuristics in cmake and the C++ guards both incorrectly enabled the AVX-VNNI code path for GCC 11, causing build failures. Red Hat currently carries these as downstream patches; this upstreams a robust fix. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
## Description This PR optimizes the CUDA MatMulBlockScaledFp8 GEMV decode path by increasing memory-level parallelism and widening per-warp column work when N is large. The kernel now issues grouped global loads ahead of compute, uses vectorized FP8-to-half2 conversion for B tiles, and keeps FP32 accumulation semantics unchanged. The change improves decode performance for long-output shapes while preserving correctness through targeted coverage for wide-tile and ragged-tail dispatch. ## Summary of Changes ### CUDA Kernel Optimization | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/math/matmul_block_scaled_fp8.cu` | Reworked GEMV kernel loop into load-first/consume-second phases to increase in-flight memory requests and improve overlap. | | `onnxruntime/contrib_ops/cuda/math/matmul_block_scaled_fp8.cu` | Added vectorized macros for FP8 decode (`fp8x2 -> half2`) and fused 16-element dot-product accumulation in FP32. | | `onnxruntime/contrib_ops/cuda/math/matmul_block_scaled_fp8.cu` | Extended kernel shape with `ColsPerWarp` and `Unroll` template parameters plus per-column predication for ragged N tails. | | `onnxruntime/contrib_ops/cuda/math/matmul_block_scaled_fp8.cu` | Updated GEMV dispatch heuristic to select wider tiles for large N and to launch the corresponding specialized kernels. | ### Test Coverage | File | Change | |------|--------| | `onnxruntime/test/contrib_ops/matmul_block_scaled_fp8_test.cc` | Included `<cuda.h>` under CUDA guards to ensure `CUDA_VERSION` is available for test compilation gating. | | `onnxruntime/test/contrib_ops/matmul_block_scaled_fp8_test.cc` | Added `GemvDecodeWideTilesFp16` to validate wide-tile (`ColsPerWarp` 2/4) decode paths and ragged-tail column predication. | ### Documentation | File | Change | |------|--------| | `docs/contrib_ops/cuda/matmul_block_scaled_fp8_experiments.md` | Documented decode GEMV MLP optimization rationale, measurements, and tuning notes (including H200 section updates). | ## Testing - Targeted CUDA unit test coverage added in `MatMulBlockQuantizedFp8WeightOpTest.GemvDecodeWideTilesFp16`. - Existing MatMulBlockScaledFp8 CUDA tests remain in place for regression coverage. - Full local test execution was not run in this step. ## Motivation and Context The previous decode GEMV structure consumed loads immediately, which limited memory-level parallelism on large-N decode shapes. By restructuring to pre-issue loads and increasing columns processed per warp for large outputs, this change better utilizes memory bandwidth and reduces decode latency while keeping numerical behavior aligned with the prior FP32 accumulation flow. ## Checklist - [x] Tests added/updated - [x] No breaking changes - [x] Documentation updated (if applicable)
…osoft#29887) ## Description At decode (one token per sequence) the int4/int8 QMoE path launches six kernels per layer, two of which move data rather than compute: `expandInputRows` gathers each token into `top_k` permuted copies before FC1, and `finalizeMoeRouting` scatter-reduces those copies back afterwards. On a small GPU this is pure overhead — the MoE GEMVs themselves are already at ~94% of the measured DRAM roofline, so the remaining headroom is in the prologue/epilogue, not the math. This PR removes both kernels for the decode-shaped GEMV path: the finalize reduction is folded into the FC2 GEMV epilogue, and FC1 reads the unexpanded activations directly through the existing permutation map. Measured **+2.5% decode throughput** on gpt-oss-20b (int4) / RTX 5060 Ti. Both optimizations are gated behind support predicates and environment kill switches, and neither changes behavior for shapes outside the GEMV's profiled range. ## Summary of Changes ### 1. Fuse the MoE finalize into the FC2 GEMV epilogue `moe_gemv_fused_finalize_kernel` launches one block per `(token, n-tile)` and loops over the token's `experts_per_token` permuted rows, accumulating `row_scale * (dot + bias)` into a shared float accumulator before a single store. This replaces the separate `finalizeMoeRoutingKernelLauncher` pass, which re-read the FC2 output from global memory. - Gated by `is_moe_gemv_fused_finalize_supported()`, which additionally requires `num_rows * experts_per_token == expanded_num_rows` (i.e. no dropped tokens) on top of the existing GEMV shape checks. - Restricted to `OutputType == T`; the mixed-output-type and TMA-warp-specialized paths keep the old two-kernel flow. - Kill switch: `ORT_DISABLE_MOE_GEMV_FUSED_FINALIZE=1`. ### 2. Skip `expandInputRows` when the FC1 GEMV runs The FC1 interleaved-SwiGLU GEMV now takes an optional `permuted_row_to_source_row` map and resolves its activation row as `permuted_row_to_source_row[row] % num_rows`, so the expanded activation buffer never has to be materialized. Only the *activation iterator* is redirected — the output pointer still uses the permuted row index, so the FC1 output layout is unchanged. - `runMoe` decides this *before* `gemm1` runs, so the predicate must exactly match the one `gemm1` uses to select the GEMV. Both call the shared `moeGemvInterleavedSwiGLUWillRun()` / `MoeGemvFc1Disabled()` helpers rather than open-coding the condition, and an `ORT_ENFORCE` in `gemm1` catches any future divergence instead of silently computing on unexpanded data. - Only applies when the expansion would be a pure permutation (no FP4/FP8/AWQ/prequant-scale or per-expert activation scale variants). - Kill switch: `ORT_DISABLE_MOE_GEMV_SKIP_EXPAND=1`. ### 3. Supporting cleanup and test coverage - `is_moe_gemv_shape_supported(..., cta_n)` factors out the shared shape checks with the n-tiling requirement parameterized. The plain GEMVs pass `kCtaN`; the fused-finalize path passes the smaller `kFusedFinalizeCtaN` that its `grid.y` actually depends on, instead of inheriting the stricter constraint and then redundantly re-checking the looser one. - `quant_dequant_blockwise()` no longer emits an all-128 zero-point tensor for symmetric 8-bit. The offset is already baked into the weight converter, so the tensor was a numerical no-op — but it made `weight_zeros` non-null, and every MoE GEMV entry point rejects block-wise cases with non-null zeros. It was silently keeping all block-wise 8-bit configs off the GEMV. ### Files | File | Change | |------|--------| | `contrib_ops/cuda/llm/moe_gemm/moe_gemv.cu` | Fused-finalize kernel + launcher; source-row indirection in the split-K and interleaved-SwiGLU kernels; shape-predicate refactor | | `contrib_ops/cuda/llm/moe_gemm/moe_gemv.h` | Declarations for the fused-finalize launcher and support predicate | | `contrib_ops/cuda/llm/moe_gemm/moe_kernels.cu` | Dispatch for both optimizations, shared predicates, kill switches | | `contrib_ops/cuda/llm/moe_gemm/moe_kernels.h` | `gemm1` signature gains `permuted_row_to_source_row` | | `test/python/transformers/test_qmoe_cuda.py` | `test_swiglu_qmoe_gemv_parity`; symmetric zero-point fix | ## Benchmark results gpt-oss-20b int4, RTX 5060 Ti (sm_120, 16 GB) on WSL2, CUDA 13.0, CUDA graph enabled. `batch=1, prompt=512, gen=128, reps=5`, 5 interleaved passes. Arms are selected with the environment kill switches on a **single binary**, so the only variable is the code path taken. | Variant | TTFT (ms) | Token latency (ms) | Decode (tps) | vs. baseline | |---|---|---|---|---| | Neither optimization (baseline) | 122.17 | 6.398 | 156.31 | — | | Skip-expand only | 121.79 | 6.358 | 157.31 | +0.6% | | Both (this PR) | 122.37 | 6.242 | **160.23** | **+2.5%** | Per-pass decode tps (paired — same pass means same thermal/clock conditions): | Pass | none | skip-expand | both | |---|---|---|---| | 1 | 157.10 | 158.27 | 162.94 | | 2 | 156.10 | 158.35 | 160.11 | | 3 | 157.07 | 153.66 | 157.58 | | 4 | 156.07 | 157.99 | 159.43 | | 5 | 155.21 | 158.30 | 161.08 | Notes on interpretation: - "Both vs. neither" wins **5/5** paired passes (+2.5% mean), so the combined result is solid. - The **fused finalize is the larger contributor** (+1.9%, 5/5 paired wins over the skip-expand-only arm). **Skip-expand alone is +0.6%, which is at this harness's ~0.5% noise floor** (4/5 paired wins, and pass 3 is negative). It is worth keeping as a strict kernel-count reduction, but it should not be sold as a measurable win on its own. - TTFT is unchanged, as expected — prefill uses the grouped GEMM and neither path is touched. - Single machine, single model. GPUs where the MoE GEMV is not selected at all will see no change. ## Testing ```bash cd onnxruntime/test/python/transformers python -m pytest test_qmoe_cuda.py -q # 109 passed, 13 skipped ``` `test_swiglu_qmoe_gemv_parity` is new and is the only test that actually reaches the MoE GEMV: `is_moe_gemv_supported` requires every GEMM dim `>= 512`, while every pre-existing functional QMoE config uses `hidden_size=128, intermediate_size=256` and therefore exercises the CUTLASS grouped GEMM. The new case uses `hidden_size = intermediate_size = 512` with `num_experts_per_token = 4`, so FC1 sees `n=1024, k=512` and FC2 sees `n=512, k=512`; `batch * seq` of 1 and 2 gives 4 and 8 expanded rows, covering both sides of `kMaxProfiledExpandedRowsForSmallProblemDim`. The 2-token cases are what exercise the multi-token index math in both new code paths, which degenerates to a no-op when `num_rows == 1`. Coverage was verified by temporarily placing a conditional `ORT_ENFORCE` immediately before each GEMV launch and confirming which tests tripped it — armed one path at a time, since an FC1 throw preempts FC2. All 7 new cases reach both the FC1 GEMV and the FC2 fused finalize; no pre-existing test reached either. To verify the fallbacks are still correct, re-run with either kill switch set: ```bash ORT_DISABLE_MOE_GEMV_FUSED_FINALIZE=1 python -m pytest test_qmoe_cuda.py -q ORT_DISABLE_MOE_GEMV_SKIP_EXPAND=1 python -m pytest test_qmoe_cuda.py -q ``` ## Motivation and Context Profiling gpt-oss-20b decode on an RTX 5060 Ti attributes ~86% of decode time to QMoE. Modelling QMoE cost against bytes moved gives a slope of 2.748 µs/MiB (= 382 GB/s against a measured 427 GB/s sustained roofline) with a fixed ~120 µs intercept; after subtracting host-side `Run()` overhead that leaves ~41 µs per call of non-GEMV GPU work. That intercept, not the GEMV, is what this PR targets. Kernels per int4 QMoE decode call drop from six to four. One further prologue fusion was implemented and measured on this branch (folding softmax/top-k into the expert-map kernel, commits `2a1d49b4e8` / `67e67e1302`) and produced +0.24% — indistinguishable from noise at p≈0.23 over 7 passes — so it was reverted rather than kept. The launch-latency floor for the remaining prologue kernels appears to have been reached. The net diff of this branch contains only the two optimizations described above. ## Checklist - [x] Tests added/updated - [x] No breaking changes — both paths are opt-out gated and fall back to the existing kernels for any unsupported shape, dtype, or parallelism config - [ ] Documentation updated (not applicable — no public API or operator schema change)
…29818) ### Description Adds a new `com.microsoft` CUDA contrib operator **`MatMulBlockQuantizedFp4Weight `**: a weight-only NVFP4 (E2M1) matrix multiplication that computes `Y = A * dequant(B)^T (+ bias)`. The weight tensor `B` is stored as packed NVFP4 (two E2M1 values per byte, low nibble first). The dequantized weight is `e2m1(B) * weight_scale_2 * e4m3(weight_scale[n, k / block_size])`, where `weight_scale` holds one E4M3 scale per `block_size` (default 16) consecutive K values and `weight_scale_2` is a single global fp32 scale. The weight is dequantized to the activation type (FP16/BF16) and multiplied with the FP16/BF16 activation. #### Operator schema - **Inputs:** `A` (FP16/BF16 `[..., K]`), `B` (packed uint8 NVFP4 `[N, K/2]`), `weight_scale` (uint8 E4M3 `[N, ceil(K/block_size)]`), `weight_scale_2` (scalar fp32), optional `input_scale` (scalar fp32), optional `bias` (`[N]`). - **Attributes:** `K`, `N`, `block_size` (default 16). - **Output:** `Y` (`[..., N]`) in the activation type. #### Implementation - Architecture-independent dequant + GEMM/GEMV path that runs on Hopper (SM90) and Blackwell. - Native SM120 CUTLASS NVFP4 path (guarded by `ORT_ENABLE_BLOCKQUANT_SM120`). - ONNX schema, kernel registration, and CMake build wiring for the CUDA EP and the CUDA plugin EP. ### Motivation and Context Enables weight-only NVFP4 quantized matmul for LLM inference, reducing weight memory footprint while keeping FP16/BF16 activation precision. ### Testing - Added C++ unit tests (`matmul_block_scaled_fp4_test.cc`) covering FP16 and BF16 activations, with and without bias. - Added a Python profiling/verification script (`profile_matmul_block_scaled.py`). - Verified the CUDA EP test binary (`onnxruntime_provider_test`) builds cleanly with `ENABLE_FP4` for SM86 + SM120a.
### Description Implement the ONNX Max and Min operators natively in the WebGPU EP so models no longer fall back to the CPU EP for these ops. Implementation - Extract the two-input compute body of BinaryElementwise::ComputeInternal into a reusable RunBinaryProgram() helper. - Add a VariadicElementwise kernel that handles ONNX's variadic 1..N inputs with multidirectional (NumPy-style) broadcasting: 1 input copies to the output; N inputs are folded pairwise (acc = op(acc, input[i])) reusing the binary elementwise program. Intermediate results are held in a reserve()'d InlinedVector<Tensor> so pointers stay stable across RunProgram calls. - Register Max/Min for opsets 8-11, 12, and 13+ using WebGpuSupportedNumberTypes(), plus the kernel-create-info entries. - Normalize operands with vec4<input_*_element_t>(...) (as Equal does) so the WGSL max/min builtins work in vectorize-broadcast mode where one operand is scalar and the other is vec4. NaN propagation (ONNX opset 12+ requires it) - Float/float16 use a type-specific WGSL helper that detects NaN via an integer bitcast (exponent all ones, non-zero mantissa) instead of the x != x idiom, which the Tint/DXC fast-math path folds to false. - Integer types use the plain builtin (no NaN possible). Tests - Existing MathOpTest.Max*/Min* cases now also exercise the WebGPU EP for the supported types (float, float16, int32, uint32), covering broadcast and 3-input variadic shapes. - Add WebGPU-targeted NaN tests for Max/Min (elementwise, scalar-broadcast, variadic, and the float16 widen-then-bitcast path) to lock in NaN propagation on the WebGPU EP. ### Motivation and Context See above. --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
July 27, 2026 20:35
hdharpure9922
self-requested a review
July 28, 2026 04:57
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Automated daily backmerge from ORT main to ovep-develop. No conflicts detected. Do NOT squash or rebase - use merge commit only.