Sync with Microsoft ONNX Runtime - 17072026#1202
Merged
Merged
Conversation
Builds are blocked by onnxruntime-github-vs2022-latest. Try unblock our limited PRs for release. Co-authored-by: GitHub Copilot <copilot@example.com>
… (sm_120) (microsoft#29706) ### Summary GroupQueryAttention's XQA decode kernel failed on consumer Blackwell GPUs (RTX 50-series, sm_120) with `cudaErrorInvalidValue`, while working fine on A100 (sm_80) and H200 (sm_90). This adds a runtime shared-memory capability check so XQA is only selected when the device can actually satisfy the kernel's dynamic shared-memory request, and otherwise falls back to cuDNN SDPA / Flash. A100/H200 behavior is unchanged. ### Root cause XQA bakes its shared-memory layout at compile time from `__CUDA_ARCH__`: - sm_80 / sm_87 / sm_90 use the large K/V-tile layout (`preferedKHeadPartBytes=128`, `cacheVTileSeqLen=64`) → up to ~140 KB of dynamic shared memory for head_size 128/256. - sm_86 / sm_89 / sm_120 use the small layout (64 / 32) → ~78–96 KB. Release/packaging binaries are built with a maximum arch of `90-virtual` (compute_90 PTX only, no native sm_120 SASS). On sm_120 the driver JIT-compiles that sm_90 PTX, so the kernel's `smemSize` carries the Hopper value (~140 KB). `launchMHA` then calls `cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, size)`, which exceeds sm_120's ~99 KB per-block opt-in limit (`sharedMemPerBlockOptin`) and returns `cudaErrorInvalidValue`. A100 (163 KB) and H200 (227 KB) have enough room, so they were unaffected. ### Key changes | File | Change | |---|---| | `xqa/xqa_impl_gen.cuh` | Add `GetSmemSize()` host helper that reads the per-kernel `smemSize` device symbol (accurate even for a PTX kernel JIT-compiled for the running SM). | | `xqa/xqa_loader_fp16_impl.cuh` | Add `GetXQAKernelSmemBytes(group_size)` head-dim dispatcher. The non-quantized fp16 footprint is an upper bound for the int8/fp8/bf16 variants (smaller cache element), so one query covers all XQA paths. | | `xqa/xqa_loader_fp16.cu`, `xqa/xqa_loader.h` | Expose `GetXQARequiredSharedMemoryBytes(device_prop, head_size, num_heads, kv_num_heads)`; a single non-templated entry point used by both the fp16 and bf16 GQA kernels. | | `group_query_attention.cc`, `group_query_attention.h` | Gate XQA selection on `required_smem <= device_prop.sharedMemPerBlockOptin`; fall back to cuDNN SDPA / Flash when it does not fit. Result is cached per node (`xqa_shared_memory_ok_`) since head_size/group are constant. | | `xqa/mha_impl.cuh` | Defensive backstop in `launchMHA`: if the requested shared memory still exceeds the device limit, throw an actionable message (which SM to build for / how to disable XQA) instead of the opaque `cudaErrorInvalidValue`. | ### CUDA graph safety `GetXQARequiredSharedMemoryBytes` uses `cudaMemcpyFromSymbol`, which synchronizes and is illegal during CUDA graph capture. The query is: - **cached** per node, so it runs at most once; - **guarded** with `onnxruntime::llm::common::isCapturing(Stream(context))` so the synchronizing call is only issued when the compute stream is not capturing; - resolved during ORT's non-captured warm-up run(s) before capture begins. If the value is somehow still unresolved while capturing, XQA is conservatively skipped for that run (safe fallback) without caching, so a later non-capturing run can resolve it. Warm-up and capture therefore make the same XQA/fallback decision, keeping the captured graph consistent with replay. ### Testing notes - Built the affected TUs (GQA dispatcher + fp16/bf16/int8/fp8 XQA loaders) with `CMAKE_CUDA_ARCHITECTURES="80;90"` (the configuration that reproduces the failure); all compile cleanly. - To validate the fix end-to-end, run a fp16/bf16 GQA decode workload on an sm_120 GPU (e.g. RTX 5090): it should now run (via fallback) instead of returning `cudaErrorInvalidValue`. Set `ORT_ENABLE_ATTENTION_KERNEL_DEBUG_INFO=1` to confirm the selected backend. - To actually run XQA (the fast path) on Blackwell, build with native arch `120` in `CMAKE_CUDA_ARCHITECTURES` (and `100` for datacenter Blackwell). With a native sm_120 cubin the layout is ~80 KB and fits, so XQA is selected.
To avoid hard dependency on nvrtc dll even when it is not used for some models.
Drop 52-real; 90-virtual Add 120-real; 120-virtual Ensure 86-real is included Q: Why not add 100-real to cuda 12.8 build? A: We assume that those machines will have cuda 13.x for best performance. Q: Why drops 52-real A: Many applications require float16 support, while 52-real cannot support it.
…lugin EP (microsoft#29620) ### Summary Phase 2 of the CUDA plugin execution provider "no-cuDNN" work. It lets single last-axis `ArgMax`/`ArgMin` run through a small custom CUDA kernel instead of cuDNN, fixes `LogSoftmax` classification in the plugin adapter, and adds a non-throwing cuDNN handle accessor so reduction kernels can fall back gracefully when cuDNN is disabled. ### Key Changes | Area | Change | |---|---| | `reduction_functions.cu` / `.h` | New `arg_min_max_last_axis<TIn, IsArgMax>` kernel (instantiated for `half`, `float`, `double`) that computes ArgMax/ArgMin indices over the last dimension of a row-major matrix without cuDNN. | | `reduction_ops.cc` | In `ReduceComputeCore`, route a single last-axis ArgMax/ArgMin (`CUDNN_REDUCE_TENSOR_FLATTENED_INDICES`) to the custom kernel when shapes fit `int`; otherwise fall through to the existing cuDNN path. `ReduceKernel::ComputeImpl` now uses `TryGetCudnnHandle`. | | `cuda_kernel.h` (native) / `cuda_kernel_adapter.h` (plugin) | Add `TryGetCudnnHandle`, which returns the cuDNN handle when available and `nullptr` otherwise (instead of throwing at handle acquisition). | | `softmax.h` | Detect `LogSoftmax` from `node.OpType()` instead of `info.GetKernelDef().OpName()`, so the plugin EP adapter classifies it correctly. | | `test_cuda_plugin_ep.py` | Add `LogSoftmax` and `ArgMin` tests; drop the `@requires_cudnn` gate from `ArgMax`, `ReduceMean`, `ReduceSum`; reduce over the last axis to exercise the cuDNN-free paths. | | `docs/cuda_plugin_ep/QUICK_START.md` | Drop `ArgMax` and reductions from the list of ops that still require cuDNN. | ### Correctness Notes - `select_last_index == 1` is already rejected on the CUDA EP, so the kernel keeping the first matching index (strict `>` / `<`) is spec-correct for the supported case. - The custom path guards `n > 0`, returns early for `m == 0`, computes the row offset in `int64_t`, and only engages when `m` and `n` fit in `int` (`gsl::narrow_cast`); larger tensors fall back to cuDNN. ### Testing - `python -m pytest onnxruntime/test/python/transformers/test_cuda_plugin_ep.py -k "log_softmax or argmax or argmin or reduce_mean or reduce_sum"` - Plugin no-cuDNN validation: `bash .env/cuda_130_plugin_no_cudnn.sh --build --test_plugin` - `onnxruntime_provider_test --gtest_filter='*Reduce*:*ArgM*'`
…9) + retiring equivalence proof + corpus-collapse tripwire (microsoft#29504) ### Motivation ONNX onnx/onnx#7959 removes the on-disk `onnx/backend/test/data/node/` corpus (targeted for ONNX 1.23), replacing it with on-the-fly in-memory generation. ORT's C++ `onnx_test_runner` reads that corpus from disk — after the deletion it would load **0 node tests and silently exit 0** (green CI on a corpus that no longer exists). This PR detaches ORT from the deleted artifacts *ahead* of the bump. ### What this does (3 pieces) 1. **Detach** — a build-time materializer (`tools/python/materialize_onnx_node_tests.py`) regenerates the node corpus to disk from ONNX's surviving generator (`collect_testcases`), so the C++ runner keeps reading from disk unchanged. EP-agnostic (CPU/CUDA/QNN share one materialized tree). 2. **Equivalence proof** — a *retiring* test (`onnxruntime/test/python/onnx_node_test_equivalence_test.py` + `tools/python/compare_node_test_corpora.py`) proves the materialized corpus is byte-identical to the original (modulo a documented ULP band); it auto-skips once the ONNX on-disk oracle disappears post-microsoft#7959. 3. **Cause-agnostic tripwire** — build-time `--min-cases` gate (FATALs the build on every MATERIALIZE=ON leg if generation < floor) + a runtime `-m` floor on the CPU node ctest, turning silent-green-on-empty into a loud red. ### Key decisions - **numpy**: cmake configure = HARD FATAL on off-pin numpy (build-env reproducibility across CI legs); Python materializer = SOFT WARN (standalone advisory). Different layers, different questions. The onnx pin stays HARD on both sides. - **QNN legs**: repointed to the materialized tree + pinned onnx/numpy installed pre-build. Node-dir runs carry a **low `-m 1` collapse sentinel** (NOT the 1500 build floor): `-e qnn` legitimately reduces the collected set to ~1529, so a 1500 runtime floor would be a false-red timebomb on the next opset bump — the strict count is enforced at build time via `--min-cases`, while `-m 1` (zero false-red risk) catches a per-leg `MATERIALIZE=OFF` that still runs the node dir. The android leg's single-case `cp -r` source is repointed to the materialized tree (the old `data/node` source vanishes post-microsoft#7959). ### Testing - Runtime tripwire empirically verified: empty/truncated corpus → runner FATAL (nonzero); full → exit 0; default (no `-m`) unchanged. - Equivalence: 1799-case byte-probe (dir-set match; byte-identical modulo the documented ULP band). - cmake configure verified to generate cleanly on latest main (ep_context + node-test regions coexist, no collisions). - **[needs-run]** on a pinned-numpy CI box: the `MATERIALIZE=ON` inner path (actual materialization + the two add_test bodies) — this dev box is off-pin so the numpy gate FATALs by design. Relates to: onnx/onnx#7959 --- ## Cross-consumer viability note (re: onnx/onnx#7959) While detaching, we assessed all of ORT's node-test consumers. The in-memory generator approach behind microsoft#7959 works cleanly for **single-version** consumers (C++, C#, docs), but has a gap for consumers needing a **historical multi-opset matrix**: - ORT's JS/web tests pull node data for opset 7–21 from 15 immutable `rel-*` release archives. Since onnx's generator is single-version, old opsets can't be regenerated from a new onnx. - Those consumers stay green post-microsoft#7959 only because released branches are immutable — there is no generator-based path to add a *new* post-microsoft#7959 opset for them. Net: workable for single-version consumers; the multi-version case would need a small per-version serialization utility or a documented migration note upstream. --- --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
July 16, 2026 20:35
hdharpure9922
self-requested a review
July 17, 2026 10:14
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
5 participants
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.