Sync with Microsoft ONNX Runtime - 15112026 by ai-fw-intg · Pull Request #1395 · intel/onnxruntime

ai-fw-intg · 2026-11-14T20:40:19Z

Automated daily backmerge from ORT main to ovep-develop. No conflicts detected. Do NOT squash or rebase - use merge commit only.

### Description  Adds int8 data support to the WebGPU `Gather` kernel. - `int8` is added to the `T` type constraint of every WebGPU Gather registration (opsets 1-10, 11-12, 13+). - int8 uses the packed-byte path that bool and uint8 already use (four 1-byte elements per `u32`). Gather only moves data, so int8 is copied as raw bytes. `unpack4xU8` reads each byte, and the `Int8x4` `SetByOffset` keeps each lane's low byte. The two's-complement bit pattern, and so the sign, is preserved. - bool, uint8 and int64 behavior is unchanged. Tests in `onnxruntime/test/providers/webgpu/tensor/gather_op_test.cc` run only on the WebGPU EP with `kOrtSessionOptionsDisableCPUEPFallback` set, and skip if the EP is unavailable: - `Gather_int8_axis0_webgpu` / `Gather_int8_axis1_webgpu`: values include -128, -1, 0 and 127. They use int32 and int64 indices, negative indices, 2-D indices, and output sizes that do and don't fill the last packed word. - `Gather_int8_scalar_indices_webgpu`: 0-D indices, giving 0-D and 1-D outputs. - `Gather_int8_all_opsets_webgpu`: opsets 7, 11 and 13, one for each registration. - `Gather_int8_all_byte_values_webgpu`: all 256 int8 values, with 257 outputs that span two workgroups and end on a partial word. - `Gather_int8_then_cast_to_float_webgpu`: Gather(int8) -> Cast(float) graphs with runtime and initializer tables, 2-D and 0-D indices, and two Gathers on one table. This mirrors the WebNN conformance graphs and catches an int8 table that is read as unsigned. The existing all-EP CPU tests `Gather_axis1_indices2d_int8` and `Gather_axis1_neg_indices2d_int8` now also run on the WebGPU kernel. ### Motivation and Context  The new WPT page [`webnn/conformance_tests/gather-cast-int8-table.https.any.js`](https://wpt.live/webnn/conformance_tests/gather-cast-int8-table.https.any.js) fails 6 required subtests on the WebGPU EP. Each failing graph reads an int8 table with `gather` (constant or runtime table, followed by `cast` to float32 and sometimes `mul`/`reshape`, 0-D indices, or two gathers on one table). The WebGPU Gather kernel didn't register int8: ``` ep.cc:96 Ep::GetCapabilityImpl webgpu kernel not found in registries for Op type: Gather This session contains graph nodes that are assigned to the default CPU EP, but fallback to CPU EP has been explicitly disabled ``` WebNN disables CPU fallback, so these graphs failed to build. The cast-then-gather variants already passed. This is one of several WebGPU EP fixes for WebNN conformance. **Validation** - Built locally with the WebGPU EP. This commit was merged with main `121e34d92b` together with a separate Relu int32 change. - `onnxruntime_provider_test --gtest_filter="WebGpuActivationTest.*:ActivationOpTest.Relu*:GatherOpTest.*:CastOpTest.*:GatherElementsOpTest.*:GatherNDOpTest.*"`: **passed**, 268/268, 0 skipped. This includes all six new `Gather_int8_*_webgpu` tests and the existing `Gather_axis1_*_int8` tests, which now run on WebGPU. - Negative check: removing the `type_constraints.push_back(DataTypeImpl::GetTensorType<int8_t>());` line makes all six `Gather_int8_*` tests **fail** at session init with `fallback to CPU EP has been explicitly disabled`. This is the same error the WPT page hits. - `lintrunner -a` and `clang-format --dry-run -Werror` (20.1.8) on the changed files: clean. - WPT in the browser. Setup: plugin EP built from main `121e34d92b`, this PR and the Relu int32 PR (local merge `5c68bbb9f9`); Edge Canary 159.0.4372.0; RTX Spark N1X / D3D12; ORT runtime 1.28.3. - `webnn/conformance_tests/gather-cast-int8-table.https.any.js?gpu`: **PASS 9/9**. Before this change, 6 required subtests failed with "Failed to build graph". - `webnn/conformance_tests/gather.https.any.js?gpu`: still **PASS 39/39**. Co-authored-by: Aditya Rastogi <adityar@ntdev.microsoft.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description  Adds int32 support to the WebGPU EP `Relu` kernel for opset 14+. - `onnxruntime/core/providers/webgpu/math/unary_elementwise_ops.cc`: the `Relu` opset 14 kernel now uses the existing `SignedNumberTypes()` helper (float, float16, int32) instead of `WebGpuSupportedFloatTypes()`. Opsets 6-12 and 13 stay float-only because ONNX added the integer types in Relu-14. The shader is unchanged: `select(x_value_t(0), a, a > x_value_t(0))` is valid WGSL for `vec4<i32>`. - `onnxruntime/test/providers/webgpu/activation/activation_op_test.cc`: new `WebGpuActivationTest.Relu_Int32_Opset14_WebGpu`. It runs on the WebGPU EP only with `kOrtSessionOptionsDisableCPUEPFallback=1`, so a missing kernel fails the test instead of silently running on CPU (the CPU EP does register int32 Relu-14). It covers `INT32_MIN`, `-1`, `0`, `1` and `INT32_MAX` with shape `{7}` (partial trailing vec4) and `{1, 2, 2, 2}` (full vec4s). It skips before constructing `OpTester` when the WebGPU EP is unavailable. Out of scope: int8 and int64 Relu, which WGSL does not support natively. #### Validation Built with the WebGPU EP from a local, unpushed merge of this commit and a separate WebGPU Gather commit into `main` (121e34d). - `onnxruntime_provider_test --gtest_filter="WebGpuActivationTest.*:ActivationOpTest.Relu*:GatherOpTest.*:CastOpTest.*:GatherElementsOpTest.*:GatherNDOpTest.*"`: **passed**, 268/268 with 0 skipped. This includes `WebGpuActivationTest.Relu_Int32_Opset14_WebGpu`. - Negative check: with the Relu-14 registration reverted to `WebGpuSupportedFloatTypes()`, `WebGpuActivationTest.Relu_Int32_Opset14_WebGpu` **fails** at session init with "fallback to CPU EP has been explicitly disabled". This confirms the test catches a missing WebGPU kernel. - `lintrunner -a` and clang-format 20.1.8 on the changed files: no issues. - WPT `webnn/conformance_tests/relu.https.any.js?gpu`: **16/18 passed**, up from 15/18 before. "relu int32 4D tensor" now passes. The 2 remaining failures are the optional int8 and int64 subtests, which are out of scope. Tested with the WebGPU plugin EP built from `main` 121e34d plus this PR and the Gather int8 PR (local merge `5c68bbb9f9`), Edge Canary 159.0.4372.0, NVIDIA RTX Spark N1X on D3D12, and ORT runtime 1.28.3. ### Motivation and Context  The W3C WebNN WPT subtest `webnn/conformance_tests/relu.https.any.js` "relu int32 4D tensor" failed on the WebGPU EP with "Failed to build graph". WebNN disables CPU fallback, and WebGPU `Relu` had no int32 kernel, so the node was assigned to CPU and session creation failed. Co-authored-by: Aditya Rastogi <adityar@ntdev.microsoft.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description This change fixes shape mismatch handling for cases where the requested shape is effectively equivalent to the existing tensor shape at the element level, but differs in rank representation. In particular, it allows reuse when the tensor has a single element and the shapes are represented as `{1}` and `{}`. This covers the scalar/single-element case that was previously failing in WebNN applications and test cases. The implementation does **not** reuse buffers for arbitrary reshapes just because the total number of elements matches. For example, shapes like `{4,5}` and `{2,10}` are still treated as incompatible and are correctly rejected. ### Example - Allowed: `{1}` ↔ `{}` - Rejected: `{4,5}` ↔ `{2,10}` ### Impact This resolves failures in WebNN cases such as `softsign` with positive float32 0D tensors.

### Description This PR adds Python bindings for the Model Package APIs. Python applications can now create, inspect and edit model packages, and select a variant and create an inference session from one. It also exposes set_ep_context_binary_information on ModelCompiler. ## Changes - Add pybind bindings for model package authoring (ModelPackage) and for loading and selecting a variant (ModelPackageContext, ModelPackageOptions, ModelPackageComponentContext). - Add ModelCompiler.set_ep_context_binary_information. - Link model_package into the Python module and export the new classes from onnxruntime. - The bindings are not built in minimal builds. ### Motivation and Context The python API's are helpful in creating sample applications for public adoptions.

This pull request updates the GitHub Actions workflows to improve security when handling pull requests from forks. The main change is to ensure that jobs running on self-hosted runners with managed identities only execute trusted code from the main repository, preventing untrusted fork code from gaining elevated permissions. Security improvements for workflow jobs: * Updated the `build-linux-cuda-x64-release` and `test-linux-cuda-x64-release` jobs in `.github/workflows/linux_cuda_ci.yml` to only run on self-hosted runners if the pull request comes from the main repository, and moved job-level permissions under this condition. [[1]](diffhunk://#diff-04806013a5e7991ba5606145885d9b0fcd99a7df1f3bb96a2d38fc724ccd9b2aL16-R26) [[2]](diffhunk://#diff-04806013a5e7991ba5606145885d9b0fcd99a7df1f3bb96a2d38fc724ccd9b2aR61) * Added a similar conditional check to the `build_test_pipeline` job in `.github/workflows/reusable_linux_build.yml` to prevent fork-controlled code from running on self-hosted runners with managed identities.

Reverts microsoft#33376

…#33846) ### Description  Replace the `build-and-prep-ort-files` action invocation in `linux_minimal_build.yml` with explicit workflow steps for requirements installation, operator-registration validation, inference-only Debug wheel building and installation, followed by the local `tools/ci_build/prepare_minimal_build_test_data.py` script. The Python script does not install packages or modify the Python environment. The local script preserves source/package conversion passes, required-operator configs, standalone invoker ops, and custom-op conversion/cleanup. The preparation build omits `--enable_training_ops`. Install exactly one built wheel with `--force-reinstall --no-deps`, preserving requirements pins and rejecting missing/ambiguous wheels. Custom-op conversion failures are fatal and temporary models are cleaned up; a missing custom-op library still produces a warning. Remove the full-build preparation job's ccache cache, statistics commands, and `--use_cache`. The pinned shared `setup-build-tools` action still requires ccache version/hash inputs and downloads it, but `add-ccache-to-path: false` is set and this job does not use it. Other tool provisioning, artifact upload and downstream jobs remain unchanged. The test file `test_prepare_minimal_build_test_data.py` is excluded as requested. Validation after review changes: - `.\.venv\Scripts\python.exe -m ruff check tools\ci_build\prepare_minimal_build_test_data.py` — passed. - `.\.venv\Scripts\python.exe -m ruff format --check tools\ci_build\prepare_minimal_build_test_data.py` — passed. - `.\.venv\Scripts\python.exe -m py_compile tools\ci_build\prepare_minimal_build_test_data.py` — passed. - `git diff --check` — passed. - Inline Python checks of parsed workflow YAML, unchanged downstream jobs, preserved build flags, and mocked conversion/config generation, missing-library behavior, failure propagation and cleanup — passed. - Bash syntax checks for all preparation-job run steps and execution of wheel-selection logic with a stub pip — passed for one wheel; zero/multiple wheels correctly rejected before pip. - Previous commit `37b738543a` passed the full-build preparation job and all downstream Linux minimal-build jobs in CI. The updated commit still requires CI validation; no full Linux build or real model conversions were run locally. ### Motivation and Context  Keep minimal-build preparation logic in ONNX Runtime as ordinary Python instead of the external bundled JavaScript action, while removing training operators from the preparation build. Keep environment setup explicit in the workflow, as requested in review, and disable ccache for the preparation job. Existing training-retirement and companion actions PRs remain untouched. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…ng and plugin EP support (microsoft#33839) ### Description  This PR accelerates external initializer loading to GPU memory for both the built-in CUDA Execution Provider and the CUDA Plugin Execution Provider without introducing external storage dependencies (`libcufile` / GPUDirect Storage or DirectStorage): 1. **Lock-Free Multi-Worker Double-Buffered Pinned Staging & Async H2D Pipeline**: - Replaces the per-chunk global thread-pool barrier and per-tensor `cudaStreamSynchronize` barrier with a shared lock-free work queue (`ExecuteExternalDataLoadTasks` in `onnxruntime/core/framework/external_data_loader_thread_pool.h`). - Each reader worker owns a 4 KiB-aligned slice in both double-buffered pinned staging buffers and a pair of non-blocking CUDA streams, pulling chunks (`2 MiB` default) via atomic index and overlapping multi-threaded file reads (`pread` on POSIX / `OVERLAPPED` `ReadFile` on Windows) with `cudaMemcpyAsync` across tensor boundaries. 2. **CUDA Plugin EP (`PluginExecutionProvider`) Support**: - Implements `PluginExecutionProvider::GetExternalDataLoader()` using `host_accessible_memory_info` (`Reserve`/`Free` for pinned memory), lightweight copy-only `OrtSyncStreamImpl` streams (`copy_only=1`, skipping unnecessary cuBLAS/cuDNN handle creation), and `OrtDataTransferImpl` (`CopyTensorAsync`). - Ensures `CudaDataTransfer::CopyTensorsImpl` and `CudaSyncStream::FlushImpl` select and restore the target CUDA device when invoked from reader worker threads. 3. **Cross-Tensor Batch Lifecycle, Offset Sorting, and File-Handle Caching**: - Extends `IExternalDataLoader` and `ExternalDataLoaderManager` with `BeginLoad()`, `FinalizeLoad()`, and `EndLoad()` so external tensors in `SaveInitializedTensors` are pipelined across tensor boundaries and synchronized once in `FinalizeLoad()`. - Sorts initializers in `SaveInitializedTensors` by `(external_location, external_offset, ort_value_index)` for sequential file access instead of hash-bucket order, and caches validated `RandomAccessFile` handles per `SaveInitializedTensors` call. 4. **Conservative Defaults, Model-Size Gating, Session Options, and Minimal Build Exclusion**: - Computes the default reader thread count from physical CPU cores (`clamp(physical_cores / 2, 2, 8)`). - Uses a conservative **32 MiB** default total pinned staging allocation (two `16 MiB` buffers, 4x smaller than the previous 128 MiB default). - Defers pinned buffer, stream, and thread-pool allocation until `FinalizeLoad()` and gates it on a **16 MiB** minimum total external initializer size on the device (`session.external_data_loader.min_size_bytes`), so small models incur zero pinned-memory or stream-creation overhead. - Adds session configuration keys (`session.external_data_loader.threads`, `session.external_data_loader.pinned_memory_bytes`, `session.external_data_loader.chunk_bytes`, `session.external_data_loader.min_size_bytes`) alongside matching CUDA EP provider options (`external_data_loader_reading_threads`, `external_data_loader_pinned_memory_bytes`, `external_data_loader_chunk_bytes`). - Excludes the feature from minimal builds (`!defined(ORT_MINIMAL_BUILD)` and `!defined(USE_CUDA_MINIMAL)`). #### Summary of Changes | File | Change | |---|---| | `include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h` | Add `session.external_data_loader.{threads,pinned_memory_bytes,chunk_bytes,min_size_bytes}` keys | | `include/onnxruntime/core/providers/cuda/cuda_provider_options.h` | Add `external_data_loader_{pinned_memory_bytes,chunk_bytes}` and conservative defaults to `OrtCUDAProviderOptionsV2` | | `onnxruntime/core/framework/external_data_loader.h` | Add `ExternalDataLoaderOptions`, validation/parsing helpers, and `Configure`/`BeginLoad`/`FinalizeLoad`/`EndLoad` hooks | | `onnxruntime/core/framework/external_data_loader_thread_pool.h` | Shared portable `std::thread` worker pool and lock-free double-buffered `ExecuteExternalDataLoadTasks` pipeline | | `onnxruntime/core/framework/external_data_loader_manager.{h,cc}` | Add `BeginLoad`, `FinalizeLoad`, and `EndLoad` batch lifecycle management | | `onnxruntime/core/framework/session_state_utils.cc` | Sort external initializers by `(location, offset)`, cache validated file handles, and bracket loading with `BeginLoad`/`FinalizeLoad`/`EndLoad` | | `onnxruntime/core/framework/tensorprotoutils.{h,cc}` | Add `ExternalDataFileCache` to validate and open each external `.data` file once per `SaveInitializedTensors` pass | | `onnxruntime/core/providers/cuda/cuda_external_data_loader.{h,cc}` | Use shared lock-free pipeline, support configurable threads/pinned memory/chunk size/min model size, and guard minimal builds | | `onnxruntime/core/providers/cuda/cuda_execution_provider_info.{h,cc}`, `cuda_provider_factory.cc`, `cuda_execution_provider.cc`, `provider_bridge_ort.cc` | Parse, validate, and forward new external data loader options | | `onnxruntime/core/session/plugin_ep/ep_plugin_provider_interfaces.{h,cc}`, `onnxruntime/core/providers/cuda/plugin/*` | Implement `PluginExternalDataLoader`, copy-only sync streams, device-safe worker H2D copies, and option validation | | `docs/Model_Loading_Performance.md` | Document new session options, provider options, defaults, and model-size threshold | | `onnxruntime/test/framework/external_data_loader_test.cc`, `cuda_external_data_loader_test.cc`, `test_cuda_plugin_ep.py` | Add C++ and Python unit tests for session options, provider options, min-size gating, and invalid-option rejection | ### Motivation and Context  Loading large ONNX models with external weights (e.g., `qwen_flash_next_int4_int8_rc5_contrib_dflash2/model.onnx`, 68.46 GiB across 899 external tensors in 7 `.data` files) into GPU memory during `InferenceSession` creation had several bottlenecks: 1. `PluginExecutionProvider` (used when building the CUDA EP as a plugin) did not implement `GetExternalDataLoader()`, falling back to per-tensor `mmap` + single-threaded pageable `cudaMemcpy`. 2. In the existing pinned-buffer loader (microsoft#32437 / microsoft#32713 fallback when `nvidia-fs` is unavailable), a global thread-pool barrier per 64 MiB chunk and a `cudaStreamSynchronize` barrier at the end of every `LoadTensor()` call prevented CPU/GPU overlap across tensors `<= 64 MiB` (705 of 899 tensors), forced all 613 tensors `< 16 MiB` onto a single thread, iterated over initializers in arbitrary hash-bucket order, and reopened/revalidated `.data` files for every tensor. #### Benchmark Comparison (`qwen_flash_next_int4_int8_rc5_contrib_dflash2/model.onnx`, 68.46 GiB, 899 External Tensors, NVIDIA H200) | Configuration | Threads | Pinned Memory | End-to-End `InferenceSession` Time | Effective Throughput | Speedup vs Baseline | |---|---:|---:|---:|---:|---:| | **Baseline (`origin/main` CUDA Plugin EP)** | 1 (pageable) | 0 MiB | 14.481 s | 4.73 GiB/s | 1.00x | | **PR microsoft#32713 (`use_gds=1` -> 4-thread barrier fallback)** | 4 | 128 MiB | 5.110 s | 13.40 GiB/s | 2.83x | | **This PR (same 4 threads, 4x less pinned memory)** | **4** | **32 MiB** | **3.730 s** | **18.35 GiB/s** | **3.88x** | | **This PR (same 4 threads, same 128 MiB pinned memory)** | **4** | **128 MiB** | **3.657 s** | **18.72 GiB/s** | **3.96x** | | **This PR (Default: 8 threads from CPU cores, conservative 32 MiB pinned)** | **8** | **32 MiB** | **2.191 s** | **31.25 GiB/s** | **6.61x** | | **This PR (Opt-in: 16 threads, 64 MiB pinned)** | **16** | **64 MiB** | **1.967 s** | **34.81 GiB/s** | **7.36x** | Even at the **exact same 4-thread count and 4x smaller pinned memory (32 MiB vs 128 MiB)**, the lock-free cross-tensor pipeline + offset sorting + file-handle caching reduces session creation time from **5.110 s to 3.730 s (1.37x faster)**. With the CPU-core-derived default (`8` threads) and conservative `32 MiB` pinned memory, session creation completes in **2.191 s (31.25 GiB/s)**. #### Validation - `./onnxruntime_test_all --gtest_filter="*ExternalDataLoader*:*Initializer*"` — **127 tests ran, 126 passed, 1 skipped (32-bit platform check), 0 failed**. - `ORT_TEST_CUDA_PLUGIN_EP=1 python onnxruntime/test/python/transformers/test_cuda_plugin_ep.py` — **94 tests ran, 92 passed, 2 skipped, 0 failed** (including `test_external_data_loader_modes_and_validation`). - `lintrunner -a` — **passed (0 lint issues)**. --------- Co-authored-by: Xavier Dupré <xadupre@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…ment (microsoft#33481) ### Description Introduce generic `enableFeatures` and `disableFeatures` lists for named WebGPU EP features, initially supporting `subgroupMatrix`. The feature registry is independent of WebGPU device features. Each session’s execution provider filters kernel-facing feature queries according to its settings, including when using an externally supplied device. Disabling `subgroupMatrix` selects existing fallback kernels for that session. Sessions sharing a device configure features independently, regardless of creation order; omitted features use registry defaults. #### Usage: - JS (onnxruntime-web / onnxruntime-node): ``` executionProviders: [{ name: 'webgpu', disableFeatures: ['subgroupMatrix'] }] ``` - C++ provider options / session config entry: ``` session_options.AddConfigEntry( "ep.webgpuexecutionprovider.disableFeatures", "subgroupMatrix"); ``` ### Motivation and Context See above.

…t#32702) ### Description Add `FlashAttentionPrefillProgram`, a simplified variant of `FlashAttentionProgram` used for dense prefill on Intel GPUs. #### Kernel design (and key differences from `FlashAttentionProgram`): - One query row per thread. - The full-head K tile is staged before the QK^T + online-softmax stage. V is staged just before the PV matmul, in slices of up to 16 vec4 values, so global-memory latency overlaps with softmax compute and V needs less workgroup memory. The default kernel stages both K and V at full head size. - K and V tiles use a transposed layout (`[head_size_vec][kv_step]`), so adjacent lanes hit adjacent shared-memory addresses and avoid bank conflicts. The default kernel uses `[kv_step][head_size_vec]`. - Subgroup shuffle is a static code path chosen on the host from the adapter's guaranteed minimum subgroup size (`shuffle_chunk = min(kv_step, subgroup_min_size)`). Below 8, the kernel reads K/V directly from shared memory, which also covers adapters without subgroups. - The workgroup size and KV step are chosen per adapter and shapes. - No quantized KV cache, TurboQuant LUT, or compressed K/V decoding paths. #### Configuration (host-selected per adapter): - Default: 64 threads, 16 KV rows per step. - Xe2: 128 threads / 64 KV rows for query length <= 128, otherwise 256 threads / 32 KV rows. - Xe3: 128 threads / 64 KV rows. - If the K tile and V slice would exceed maxComputeWorkgroupStorageSize, fall back to the default configuration. #### Integration: - Adds CanApplyFlashAttentionPrefill and ApplyFlashAttentionPrefill. ApplyDensePrefill selects this kernel when the adapter is Intel and the KV cache is not quantized. Otherwise it uses the existing FlashAttentionProgram. - Quantized KV cache (int8 / TurboQuant) is not supported by this kernel yet. #### Tests: - Parameterize RunFp16HighMagnitudeAttention by head size and add head_size 80 and 96 cases, whose head_size / 4 is not a multiple of the V slice width. #### Performance (Phi-4-mini, Intel Panther Lake): | Metric | Baseline Prefill TPS | Optimized Prefill TPS | Improve | |---|---|---|---| | **Prefill-128** | 1961.46 | 1945.12 | -0.83% | | **Prefill-256** | 2481.54 | 2502.52 | 0.85% | | **Prefill-512** | 2577.39 | 2774.83 | 7.66% | | **Prefill-1024** | 2505.69 | 2751.47 | 9.81% | | **Prefill-2048** | 2346.94 | 2673.69 | 13.92% | | **Prefill-4096** | 1668.55 | 1985.91 | 19.02% | #### TODO (follow-up PRs): - Support quantized KV cache (int8 / TurboQuant). - Merge back into `FlashAttentionProgram`. ### Motivation and Context See above.

### Description  Add blocking GPU IO-binding and Release proxy runs, align native suites, and use descriptive browser profiles with automatic log discovery. ### Motivation and Context  Prepare for JSEP migration. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…icrosoft#33848) ### Description  Update the DLPack v1.3 dependency revision and hashes in CMake, vcpkg, and component governance. Add bidirectional mappings for all implemented ORT numeric/bool types, including all five float8 variants, packed float4, int4/uint4, and int2/uint2. Use native `kDLBool` metadata and remove `is_bool_tensor`, dtype inspection, and bool-as-uint8 workarounds throughout inference, ORTModule, Triton, and ATen scalar conversion. The low-level override API intentionally breaks; uint8 metadata now always imports as uint8. Export explicit contiguous strides, honor import byte offsets, allow nullable producer deleters, and add dtype/storage/lifetime/offset/invalid-metadata regression coverage. Retain legacy `DLManagedTensor` capsule exchange; versioned capsule negotiation is not added. ### Motivation and Context  Modern DLPack provides distinct bool and low-precision dtype codes, enabling correct zero-copy interchange without historical uint8 workarounds. float6 remains unsupported because ORT lacks its tensor storage/type implementation on this branch. Complex tensors are also unimplemented; string tensors have no compatible DLPack representation. --------- Co-authored-by: justinchuby <11205048+justinchuby@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…#33889) ### Description  - Stop requiring import libraries for dynamically loaded CUDA and TensorRT providers. - Detect the CUDA DLL in the CMake install tree's lib directory when staging headers. ### Motivation and Context  Fix packaging pipeline issues. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description  Move `test/providers/cuda/bitwise_ops_test.cc` to `test/providers/cpu/math/cuda_bitwise_ops_test.cc` under `USE_CUDA`, where both in-tree and plugin CUDA EP builds include it. Preserve all nine tests, integer types, opsets, broadcasting and edge cases unchanged, including explicit CUDA selection and disabled CPU fallback. The existing provider factory selects the plugin EP in plugin builds. Validation: evaluated the repository's CMake source-selection clauses for both build modes (suite included exactly once); verified the original test body is unchanged; MSVC `/WX /Zs` CPU-only compile check and `git diff --cached --check` passed. CUDA-enabled build/execution was not run: this local worktree has no build and the CUDA toolkit is not installed. ### Motivation and Context  Addresses [the review comment on microsoft#33646](microsoft#33646 (comment)): the CUDA-specific source glob excludes plugin builds, leaving these kernels untested by CUDA Plugin EP CI. This is a follow-up to merged microsoft#33646. Signed-off-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…t#33757) ### Description  Add one independent, default-off, EP-owned CUDA kernel replacement for both **FLOAT and BF16**: ```text ep.cuda.runtime_fusion.unit_swish = 0 | 1 ONNX Swish-24(x, alpha=1) -> com.microsoft.ScaledSiLU-1(x, alpha=1) ``` This PR is based directly on `main` at `dd543febefd37da323b11c4c472b4ecb6e81e159`, using the region-scoped `Run` API already merged in microsoft#33652. It **does not depend on microsoft#33705**. CausalConv remains a separate follow-on PR; there is no separate BF16 Swish PR. Only exact configuration strings `"0"` and `"1"` are accepted. Alpha may be absent/defaulted or a FLOAT exactly equal to 1. Non-unit alpha is not admitted: Swish is `x * sigmoid(alpha*x)`, while ScaledSiLU is `silu(alpha*x)`. The replacement uses the existing `rewrite::PatternRewriter` API, metadata-only discovery, a typed scratch CUDA-kernel probe, and owned metadata snapshots with graph-owned identity tokens revalidated before acceptance. It preserves graph outputs and fanout, skips ineligible anchors before matcher traversal, and checks the target's launch bound. The adapter separates the common admission gates in `Select` from `detail::SwishFusion::Discover` and the case-local `OfferSwish`. Offers carry deferred optimization callbacks rather than modifying the graph during discovery. Acceptance retains the existing region-scoped `Run`, the exactly-one-rewrite/one-target checks, and input/output identity revalidation. It does not import `RunOnce` or the unmerged CausalConv/multi-sink changes from microsoft#33705. Configuration strings, the internal `unit_swish_scaled_silu_v1` identifier, and diagnostics remain unchanged. Admission retains the main-graph/unfiltered-view, optimization-level, exact opset-import, full-CUDA-build, consumer, layout, resource, and ORT-loader guards. Symbolic, anonymous, scalar, and empty shapes are supported without guessing workload sizes. Learned initializer payloads and external files are not read during discovery. **No new CUDA kernel or arithmetic, operator schema, public API, input-model conversion, model-ID whitelist, CausalConv fusion, RMS, or Scan implementation is introduced.** Existing CUDA `ScaledSiLU` kernels are reused. **Validation** **Standalone adapter refactor:** built all objects and provider libraries afresh from this checkout, using only copied cached dependency sources. No previously built objects or libraries were reused. Release, H200/sm90, CUDA 13.3, QUICK_BUILD ON, native CUDA profiling OFF. | Exact command | Result | |---|---| | `cmake --build /mnt/data/titaiwang/runtime-fusion/swish-refactor-20261104/build/Release --target onnxruntime_optimizer onnxruntime_test_all onnxruntime_provider_test --parallel 24` | Passed; fresh build, including the new Swish translation unit | | `./onnxruntime_test_all --gtest_filter=UnitSwishHostTest.* --gtest_output=xml:/mnt/data/titaiwang/runtime-fusion/swish-refactor-20261104/host.xml` | **13 passed**, zero failures/errors/skips | | `./onnxruntime_provider_test --gtest_filter=CudaRuntimeFusionTest.UnitSwish* --gtest_output=xml:/mnt/data/titaiwang/runtime-fusion/swish-refactor-20261104/provider.xml` | **11 passed**, zero failures/errors/skips | | `clang-format --dry-run --Werror onnxruntime/core/optimizer/cuda_runtime_fusion.cc onnxruntime/core/optimizer/cuda_runtime_fusion_internal.h onnxruntime/core/optimizer/cuda_runtime_fusion_swish.cc` | Passed | | `git diff --check` | Passed | Both fresh test commands ran from `/mnt/data/titaiwang/runtime-fusion/swish-refactor-20261104/build/Release`, with `CUDA_VISIBLE_DEVICES=0` and `LD_LIBRARY_PATH=/mnt/data/titaiwang/runtime-fusion/swish-refactor-20261104/build/Release:/usr/local/cuda-13.3/lib64:/mnt/data/titaiwang/runtime-fusion/venv/lib/python3.12/site-packages/nvidia/cudnn/lib`. Provider library mappings, tested source hashes, generated CMake/Ninja source registration, XML counts, protected-build preservation, and owned cache cleanup were verified. No test inputs, assertions, or test source files changed. The bounded refactor was reviewed directly; no actionable findings were found. Existing predicates, diagnostics, callback lifetime/ownership, region `Run`, postconditions, and admission order were checked against the baseline. This fresh pass does not rerun the full provider suite, whole-model qualification, non-CUDA variants, or cross-architecture coverage. **Pre-refactor qualification:** source `74ca7dec28844f26504f92aac052591d4b77c456`. Proper offline Release build on H200/sm90 with CUDA 13.3 and QUICK_BUILD; existing cached dependencies only. The build was rerun after fixing standalone-test `-Werror=unused-function` errors caused by orphaned causal-only helpers. The archived results below describe that revision, not a rerun of the refactored adapter. | Exact command | Result | |---|---| | `cmake --build /mnt/data/titaiwang/runtime-fusion/build-bf16-unit-swish/Release --target onnxruntime onnxruntime_test_all onnxruntime_provider_test --parallel 8` | Passed | | `/mnt/data/titaiwang/runtime-fusion/build-bf16-unit-swish/Release/onnxruntime_test_all --gtest_filter=UnitSwishHostTest.*:CUDARuntimeFusionHostTest.*:PatternRewriterTests.*:PatternComponentTests.*:PatternCapabilityTests.*:ReplacementComponentTests.*:CheckAndBindComponentTests.*:MatchComponentTests.*:FunctionTest.* --gtest_output=xml:/mnt/data/titaiwang/runtime-fusion/swish-causal-cleanup-20261030/swish-final/host.xml` | **296 passed**, zero failures/errors/skips | | `/mnt/data/titaiwang/runtime-fusion/build-bf16-unit-swish/Release/onnxruntime_provider_test --gtest_filter=CudaRuntimeFusionTest.*:CudaRuntimeFusionProbe.ZeroBatch:CausalConvWithStateTest.*-*LocalAuthentic* --gtest_output=xml:/mnt/data/titaiwang/runtime-fusion/swish-causal-cleanup-20261030/swish-final/provider.xml` | **69 passed**, zero failures/errors/skips | Both test commands ran from the build's `Release` directory. The provider test loaded that build's CUDA/shared provider libraries; exact libraries were archived and SHA-256 recorded before subsequent work. The selected suites include unit-Swish host and CUDA tests. Coverage includes default-off/malformed/basic policy, typed-kernel rejection, metadata immutability and stale acceptance, external initializer headers, shape/launch bounds, graph outputs/fanout, NHWC/resource/consumer exclusions, ORT file/buffer compatibility, and both numeric dtypes. Numeric tests compare actual expanded CUDA source against recovered `ScaledSiLU` with CPU fallback disabled. Non-NaN values must match bitwise; NaN classification must agree. FLOAT raw-bit cases and all 65,536 BF16 raw patterns are covered; exhaustive BF16 enumeration is one rank-1 case, with seeded raw-bit coverage at additional ranks/shapes. `clang-format --dry-run --Werror` over changed C++ files and `git diff --check` passed. Full lintrunner was unavailable because `lintrunner_adapters` was missing; no dependencies were installed. **Additional unchanged-Qwen diagnostic** The original Qwen3.8-27B standard BF16 decoder and embedding model were loaded without conversion, re-export, or editing model files. Batch 1, 12-token prefill and one decode carrying each route's own states. Causal fusion was disabled; only unit-Swish differed. ```sh /mnt/data/titaiwang/runtime-fusion/venv/bin/python -B /mnt/data/titaiwang/runtime-fusion/swish-causal-parent-20261030/model.py /mnt/data/titaiwang/runtime-fusion/swish-causal-cleanup-20261030/swish-final/libraries 74ca7de swish-independent-model-final ``` Result: **passed**. All 129 outputs were finite and bitwise identical off/on at both prefill and decode. ORT profiles recorded 48 unique FLOAT and 112 unique BF16 recovered nodes. Raw Nsight/CUPTI records confirmed **96 FLOAT and 224 BF16 `ScaledSiLUKernel` launches** across the two steps; the disabled route had zero target events. Original model hashes and external-data file metadata were unchanged. Local harness/profiles are not committed; generated profiler SQLite files were removed and profiles compressed. This is an additional exercised-model result, not a required whole-model accuracy claim or proof for arbitrary models/generation lengths. **Review and exclusions** Pre-refactor five-role static review completed with actual runtime models verified. Critical/Major findings ledger: empty. Minority report: none. Nonblocking questions: Critical notes remaining whole-graph indexing per Swish candidate has not been scaling-benchmarked; Integration notes ORT serialization inlines source Swish, so the file/buffer round-trip test alone does not prove a preserved-Swish artifact takes the loader opt-out. Unsupported-consumer tests and direct loader wiring are covered separately. Not checked: other GPU architectures, broad performance/scaling, long-generation accuracy, and excluded/minimal/plugin build variants. The numeric results describe exercised inputs, not a universal cross-platform bitwise guarantee. No models, weights, profiler caches, or binaries are committed. ### Motivation and Context  Standard-op exports can preserve unit-alpha Swish even though CUDA already provides the equivalent ScaledSiLU kernel. Select that implementation through the EP's capability mechanism without changing the model or requiring exporter-specific graph rewrites, with one dtype-qualified replacement rather than separate FLOAT/BF16 PRs. --------- Co-authored-by: Ti-Tai Wang <18010845+titaiwangms@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…ort-circuit for QMoE offload (microsoft#33624) ## Motivation & Summary Builds on microsoft#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 microsoft#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) | microsoft#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. microsoft#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>

…icrosoft#33699) ### Mechanism In bounded host-cache offload mode (microsoft#33622), QMoEOffload::RunColdExperts fuses per-expert chunked demand loads (FileExpertStore::LoadExpertChunk and FinishLoadExpert) with FC1 and FC2 tile execution in a single parallel dispatch. Previously, the FC1 and FC2 stages indexed group slots in raw 0 .. group_size - 1 order, so a demand-missing slot (slot_ready == 0) at a lower slot index caused worker threads to spin on its readiness barrier before reaching already-cached ready slots (slot_ready == kReadyDoneVal). This change adds BuildReadyFirstSlotOrder and RunReadyFirstFusedColdExpertGroup in qmoe_shared_layout.h to stably partition group slots before parallel dispatch so initially ready cached slots execute before demand-loading slots in FC1 and FC2. Each slot retains its original activation input, lease binding, intermediate buffer, output buffer, and reduction order. When miss_load_count == 0 (all cache hits), reordering is bypassed. Pre-dispatch handling of in-flight prefetches (BeginLoadExpert returning -1) and lease cleanup behavior are unchanged. ### Base & Tested Configuration - **Base branch**: Applies cleanly on top of microsoft#33622 (commit 5855bd4); ready-first slot ordering is independent of CPU SIMD ISA. - **Tested configuration on Zen 2**: Both baseline and candidate benchmark builds merged microsoft#33622 with microsoft#33624 (commit 2ced7eaf86764c6cf25e38c2ba083c54251e9a45) so cold-expert INT4 GEMMs on the AVX2-only AMD EPYC 7V12 host executed Avx2Int4GemmRange rather than ScalarInt4GemmRange. AVX2 was part of both benchmark builds, not a functional requirement of ready-first scheduling. ### Scoped Benchmark Summary Autoregressive greedy single-token decode on Qwen3-30B-A3B INT4 (full_benchmark.onnx, 48 layers, 128 experts/layer, top_k=8, hidden=2048, intermediate=768, block_size=32), 1x NVIDIA A100-SXM4-80GB + AMD EPYC 7V12 (8 SpinPool CPU threads), tail_substitution_threshold=0.0, skip_zero_cold_doorbell=0. Evaluated across N = 8 balanced pairs (2 complete 4x4 cyclic Latin squares across 4 prompts x 32 generated tokens = 124 post-first tokens/run, counterbalanced across AB/BA orders = 32 paired requests/condition): | Condition (N = 8 Latin-Square pairs, 124 post-first tokens/run) | Baseline Decode Wall (ms, mean ± SD) | Candidate Decode Wall (ms, mean ± SD) | Paired Wall Diff (Cand - Base, mean ± 95% CI) | Baseline Sustained (tok/s, mean ± SD) | Candidate Sustained (tok/s, mean ± SD) | Paired Throughput Diff (mean ± 95% CI) | |---|---:|---:|---:|---:|---:|---:| | **Primary**: adaptive_promo_on, gpu_slots=32, host_cache_slots=16 (2.04 GB), prewarm=1 | 3592.00 ± 72.82 | 3533.61 ± 61.54 | **-58.38 ± 27.70 ms** (-1.62% ± 0.74%, 95% CI [-2.36%, -0.88%], p = 0.0016, 8/8 faster) | 34.534 ± 0.690 | 35.101 ± 0.613 | **+0.567 ± 0.256 tok/s** (**+1.65% ± 0.77%**, 95% CI [+0.88%, +2.42%], median +1.43%) | | **Cache-Hit Control**: static, gpu_slots=32, host_cache_slots=95 (12.11 GB), prewarm=1 | 3470.58 ± 60.32 | 3452.16 ± 39.59 | -18.42 ± 30.80 ms (-0.52% ± 0.88%, 95% CI [-1.40%, +0.37%]) | 35.738 ± 0.614 | 35.924 ± 0.410 | +0.186 ± 0.316 tok/s (+0.53% ± 0.89%, 95% CI [-0.36%, +1.42%]) | ### Output Parity & Validation - **Output parity**: Across all 64 paired requests (2,048 generated token steps), generated token sequences and final-step float32 logits SHA-256 hashes matched between baseline and candidate. - **Cache-hit control**: No statistically detectable regression in the tested cache-hit control (host_cache_slots=95, paired decode wall-time difference -0.52% ± 0.88%, 95% CI [-1.40%, +0.37%]). - **Unit & regression tests**: Added QMoESharedLayoutTest.ReadyFirstFusedColdExpertGroupCoversMixedHitsMissesDuplicatesAndFailures (13/13 non-AVX512 C++ tests in QMoESharedLayoutTest.* on the pure submission branch, 15/15 on the combined AVX2 branch, and all non-skipped Python tests in test_qmoe_offload_v4.py and test_qmoe_shared_layout.py pass). ### Limitations - Reorders compute slots within mixed ready + demand-loading groups only; synchronous LoadExpertChunk I/O (pread + TransposeInt4Scales) still executes on the same SpinPool CPU worker threads, yielding 1.65% higher throughput (95% CI [+0.88%, +2.42%]) in the tested Qwen INT4/A100/Zen 2 configuration. - Benchmark timings above were collected on the minimal ready-first binary (175145b1...); subsequent review-only hygiene edits (MSVC intrin.h guard, group_size <= MaxGroupSlots bounds check, test placement for clean apply on microsoft#33622, and reverting an unrelated 2-line GroupLeaseGuard edit for scope isolation) were validated with the C++ and Python regression suites without re-running the timed benchmark campaign. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

## Description Support batched right-padded GQA chunks and populated idle entries (`query_seqlens[b] = 0`) on CPU and CUDA, including single-token steps. Skip padded preprocessing and attention work while preserving physical tensor strides and the sliding-window history needed by the first valid query in each chunk. This follows microsoft#33503 and microsoft#33505, both already merged. Other execution providers continue to reject `query_seqlens`; the legacy multi-token/past-context batch restriction remains when the input is absent. ## Summary of Changes ### Length semantics and correctness - Accept `0 <= Q[b] <= min(S, T[b])` with `T[b] >= 1`. An idle entry keeps its absolute total and cache unchanged. Initially empty idle entries remain unsupported. - Honor explicit lengths at `S = 1`, rather than treating every entry as active. - Fix CPU flash dispatch for equal resident totals but different past lengths: the causal offset cannot be shared in that case. - Preserve staging/eviction behavior: eviction must not discard the preceding `W - 1` positions needed by the first token of a chunk. - Clarify that `total_sequence_length` bounds absolute totals, not windowed cache capacity. ### CPU - Bound unpack/transpose, RoPE, FP16 conversion, naive QK/softmax/SV and MLAS flash query work by the valid row count. - Keep padded physical strides for head, batch, output and bias addressing. - Skip idle attention while preserving required past-cache copies for separately allocated outputs. - Cover FP32/FP16 inputs and INT8/INT4 caches; retain default MLAS behavior when valid query counts are not supplied. ### CUDA - Guard padding before Q/K/V loads, QK-Norm, RoPE, quantization and cache append. - Skip idle aliased-cache staging seed/copy-back and unused ungroup/dequantization rows. - Bound FlashAttention/MEA query tiles and output stores by device-side valid lengths. - Skip padded unfused softmax/output rows and use masked V scratch to avoid `0 * NaN` contamination from unused cache rows. Include that scratch in runtime and bounded workspace calculations. - Bypass XQA, fast Flash decode and cuDNN SDPA with explicit query lengths because those paths do not represent idle query entries. No host synchronization is added to inspect query lengths. **Remaining work:** unfused CUDA cuBLAS QK/SV GEMMs still use dense physical query dimensions. Partially filled fused tiles retain tile-granular arithmetic. Scratch allocations and upstream Q/K/V projection nodes are not made ragged by this change. ### Regression coverage and documentation - Replay prompts `[513, 1025]` with `S = W = 512`, CPU cache capacity 1023 and CUDA capacity 512: | Chunk | Valid queries | Absolute totals | |---|---|---| | 1 | `[512, 512]` | `[512, 512]` | | 2 | `[1, 512]` | `[513, 1024]` | | 3 | `[0, 1]` | `[513, 1025]` | - Compare valid outputs and resident caches with independent unpadded single-entry full-cache calls. Include all-idle chunks, idle decode, resumed decode, mixed first/continuation entries, additional eviction, and equal-total/different-past cases. - Poison padded Q/K/V and unused floating-cache rows with NaNs. Prebind sentinel outputs and verify padded rows are not written. The sentinel check is an optimization regression, not a new public padding-output guarantee. - Exercise packed/unpacked inputs, RoPE, FP16/BF16 CUDA, FP32/FP16 CPU, and INT8/INT4 caches. - Assert native FlashAttention/MEA/math dispatch without CPU fallback, exercise MEA large-head accumulation, and replay CUDA graphs while changing device query lengths. - Update operator schema, generated operator documentation and CPU/CUDA guides. ## Testing Validated locally on H200 with CUDA 13 using the existing incremental build and feature-source overlays: - Rebuilt Python runtime, CPU/MLAS, CUDA provider and provider-test targets, including dependent FlashAttention/MEA translation units. Verified compiled feature-source and staged-provider parity. - GQA C++ suite: **123 passed, 14 skipped** (137 tests). - Focused MEA/unfused workspace GTests: **23 passed**. Both changed workspace test files compile; additional standalone layout/bounds checks passed. The main build has CUDA EP internal tests disabled, so the complete estimator/internal-test suite was not run. - Targeted Python query-length/windowed suites: **112 passed, 2 skipped**, plus **148 passing subtests**. - Final new optimization suite rerun: **36 passed**, plus **48 passing subtests**, including all three CUDA-graph backends. - Independent NumPy-oracle replay: **96 CPU calls and 156 CUDA calls passed**; CUDA dispatch verified for FlashAttention, MEA, math and automatic selection. - All configured lintrunner checks and `git diff --check` passed. The CPU output-sentinel test was also run before optimization and reproduced padded output writes for both flash settings and both cache-alias modes. ## Checklist - [x] CPU and CUDA changes included - [x] Regression tests added/updated - [x] Existing no-input behavior and other-EP rejection retained - [x] Documentation and backend limitations updated --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>

) Updates the requirements on [setuptools](https://github.com/pypa/setuptools) to permit the latest version. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/pypa/setuptools/blob/main/NEWS.rst">setuptools's changelog</a>.</em></p> <blockquote> <h1>v84.2.1</h1> <h2>Bugfixes</h2> <ul> <li>Changed wheel tags and stable-ABI validation to use sysconfig data. (<a href="https://redirect.github.com/pypa/setuptools/issues/5360">#5360</a>)</li> <li>Fixed unimportable extensions on free-threaded Python 3.13 and 3.14 when <code>Extension.py_limited_api</code> is enabled. These builds now use their version-specific extension suffix instead of using the unsupported <code>.abi3t</code> suffix. (<a href="https://redirect.github.com/pypa/setuptools/issues/5363">#5363</a>)</li> </ul> <h2>Improved Documentation</h2> <ul> <li>Replaced broken link to Python Eggs documentation with an archived version from the Wayback Machine. (<a href="https://redirect.github.com/pypa/setuptools/issues/1929">#1929</a>)</li> </ul> <h1>v84.2.0</h1> <h2>Features</h2> <ul> <li>Added support for building abi3t extensions and abi3.abi3t wheels on Python 3.15 and newer. See :pep:<code>803</code> and the <code>Python 3.15 "What's New" entry <https://docs.python.org/3.15/whatsnew/3.15.html#pep-803-abi3t-stable-abi-for-free-threaded-builds></code>_ for more details. (<a href="https://redirect.github.com/pypa/setuptools/issues/5193">#5193</a>)</li> <li>Updated vendored packaging version to 26.3, which adds support for the <code>abi3t</code> tag from :pep:<code>803</code>. (<a href="https://redirect.github.com/pypa/setuptools/issues/5361">#5361</a>)</li> </ul> <h2>Bugfixes</h2> <ul> <li>Use target sysconfig data for stable-ABI extension filenames. (<a href="https://redirect.github.com/pypa/setuptools/issues/5351">#5351</a>)</li> </ul> <h1>v84.1.0</h1> <h2>Features</h2> <ul> <li><code>setuptools.archive_util</code> now raises the new <code>UnsafeMember</code> exception when an archive member would be extracted outside of the extraction directory, where previously such a member was silently skipped. Aborting makes a malicious or malformed archive visible to the caller instead of yielding a quietly incomplete extraction, and matches the behavior of the standard</li> </ul>  </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/pypa/setuptools/commit/3895f9829fc0e51882981773a823a658eb1d1636"><code>3895f98</code></a> Bump version: 84.2.0 → 84.2.1</li> <li><a href="https://github.com/pypa/setuptools/commit/7c5cb7a14aaa61d5fec6144a4a08a3911257f0d2"><code>7c5cb7a</code></a> Merge pull request <a href="https://redirect.github.com/pypa/setuptools/issues/5362">#5362</a> from jameshilliard/fix/abi3t-target-config</li> <li><a href="https://github.com/pypa/setuptools/commit/00595bf726c0b813c0f5e0945bf9e5e8a0033744"><code>00595bf</code></a> Use sysconfig data for wheel tags and stable-ABI validation</li> <li><a href="https://github.com/pypa/setuptools/commit/873b29c064a6981cb9764c064f6c5fe0c56b0954"><code>873b29c</code></a> Merge pull request <a href="https://redirect.github.com/pypa/setuptools/issues/5363">#5363</a> from ngoldbaum/fix-3.14t</li> <li><a href="https://github.com/pypa/setuptools/commit/d8bc01b24240f8d6418f6b2855e87a2c523f6ef4"><code>d8bc01b</code></a> Merge pull request <a href="https://redirect.github.com/pypa/setuptools/issues/5330">#5330</a> from Topgearfanboy/main</li> <li><a href="https://github.com/pypa/setuptools/commit/e27271e1b2dc4342b057c3a0ea489380ab6d5231"><code>e27271e</code></a> Apply review comments from <a href="https://github.com/jameshillard"><code>@jameshillard</code></a></li> <li><a href="https://github.com/pypa/setuptools/commit/aa98f8596d2bb45323d4838ff48daed0578668e3"><code>aa98f85</code></a> add newsfragment</li> <li><a href="https://github.com/pypa/setuptools/commit/3301fa6500c3ea286581df83c038d80d8c5f9589"><code>3301fa6</code></a> Avoid .abi3t extension suffixes on Python 3.13t and 3.14t</li> <li><a href="https://github.com/pypa/setuptools/commit/301e27a484f17e175acf6b692fccc634f9a20c49"><code>301e27a</code></a> Bump version: 84.1.0 → 84.2.0</li> <li><a href="https://github.com/pypa/setuptools/commit/39a7c2ecd80bb1c393b3867d9673027d490322cb"><code>39a7c2e</code></a> Merge pull request <a href="https://redirect.github.com/pypa/setuptools/issues/5193">#5193</a> from ngoldbaum/abi3.abi3t</li> <li>Additional commits viewable in <a href="https://github.com/pypa/setuptools/compare/v61.0.0...v84.2.1">compare view</a></li> </ul> </details> <br /> Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

Bumps [actions/upload-artifact](https://github.com/actions/upload-artifact) from 7.0.1 to 7.0.2. <details> <summary>Release notes</summary> <p><em>Sourced from <a href="https://github.com/actions/upload-artifact/releases">actions/upload-artifact's releases</a>.</em></p> <blockquote> <h2>v7.0.2</h2> <h2>What's Changed</h2> <ul> <li>Improves artifact download retries when the service returns HTTP 429 (rate limiting), including honoring valid Retry-After headers.</li> <li>Updates <code>@actions/artifact</code> to v6.3.1.</li> </ul> <h2>New Contributors</h2> <ul> <li><a href="https://github.com/tunc-d"><code>@tunc-d</code></a> made their first contribution in <a href="https://redirect.github.com/actions/upload-artifact/pull/821">actions/upload-artifact#821</a></li> </ul> <p><strong>Full Changelog</strong>: <a href="https://github.com/actions/upload-artifact/compare/v7.0.1...v7.0.2">https://github.com/actions/upload-artifact/compare/v7.0.1...v7.0.2</a></p> </blockquote> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/actions/upload-artifact/commit/cf430e030ddbb5b0abf93d22962f4752f3646cd9"><code>cf430e0</code></a> Merge pull request <a href="https://redirect.github.com/actions/upload-artifact/issues/822">#822</a> from actions/tunc-d-prepare-upload-v7-0-2</li> <li><a href="https://github.com/actions/upload-artifact/commit/c2acdedd35a8824a1552c7137f8c18288c3ea469"><code>c2acded</code></a> chore: prepare v7.0.2 release</li> <li><a href="https://github.com/actions/upload-artifact/commit/4687c1bbaa0a58abb517d90b4bbade6754c80c1a"><code>4687c1b</code></a> Merge pull request <a href="https://redirect.github.com/actions/upload-artifact/issues/821">#821</a> from actions/tunc-d-upload-artifact-package-bump</li> <li><a href="https://github.com/actions/upload-artifact/commit/8cbe1bebdb858105d8fdd88f81404d4de4595719"><code>8cbe1be</code></a> Bump <code>@actions/artifact</code> to 6.3.1</li> <li>See full diff in <a href="https://github.com/actions/upload-artifact/compare/043fb46d1a93c77aae656e7c1c64a875d1fc6a0a...cf430e030ddbb5b0abf93d22962f4752f3646cd9">compare view</a></li> </ul> </details> <br /> [![Dependabot compatibility score](https://dependabot-badges.githubapp.com/badges/compatibility_score?dependency-name=actions/upload-artifact&package-manager=github_actions&previous-version=7.0.1&new-version=7.0.2)](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

### Description Raises the MLAS LinearAttention test tolerance from 1e-4 to 2e-4. The NEON kernel uses FP32 FMA reductions and reassociates the recurrence relative to the double-precision test oracle. On Windows ARM64, accumulated rounding differences can slightly exceed the previous tolerance. ### Motivation and Context Fixes microsoft#33862 Signed-off-by: Patryk Kaiser <patryk.kaiser@arm.com>

…33690) On Arm64 the NHWC transformer can rewrite a 1-D Conv into a float NhwcFusedConv, but that kernel accepted only 2-D kernels and failed at run time with conv.cc:269 Conv with channels_last layout currently supports 2D kernels. so session creation failed for every Conv1d model at ORT_ENABLE_ALL -- for example the two kernel_shape=[3] convolutions of the Whisper encoder. Windows ML creates its session at ORT_ENABLE_ALL with no optimization-level control, so callers had no in-band way around it. An NWC tensor [N, W, C] has exactly the memory layout of NHWC [N, 1, W, C]. The kernel's existing layout-conversion fallback therefore handles 1-D with a unit height: the channels-last spatial slice, output-shape reorder, fallback buffer sizing and NHWC<->NCHW conversions are generalized from a fixed four dimensions to N + kernel_rank + C. No new compute path is added. The NHWC fast path stays 2-D only, gated explicitly on kernel_rank == 2, and 2-D behavior is unchanged. Fixing the kernel rather than the transformer closes the failure regardless of which path produced the node. ValidateInputShape and ComputeKernelShape both run before the rejected check, so reaching it proves the failing node held a genuine rank-3 input and weight -- the exact state this change handles. Tests: four new FusedConvTest cases. Stride 1 and stride 2 match the shape classes of Whisper's /conv1/Conv and the failing /conv2/Conv; a fused-sum case covers the generalized manual-sum path; and Cpu_Conv1D_NcwCrossCheck runs the same data through the existing NCW FusedConv so the expected values do not rest on one reference. Before this change the three NHWC cases fail with the exact message above and the cross-check passes; after it all four pass. The existing 2-D NHWC tests also pass, through the generalized fallback path.

@edgchen1

### Description  Adds int4/uint4 support to the WebGPU EP for the W3C WebNN `quantizeLinear`/`dequantizeLinear` conformance (WPT) cases. Every gap was found by WPT running on the WebGPU EP with CPU fallback disabled (Edge Canary, RTX/D3D12). Commits: 1. **Support int4 and uint4 in Reshape.** - This is why the DQ cases failed. Chromium's WebNN ORT backend builds opset-21 models. It reshapes a broadcastable scale and zero point to the scalar or 1-D shape that (De)QuantizeLinear takes: `[1]` becomes a scalar, and `[1,1,2,1]` becomes `[2]` for per-axis. - The CPU Reshape kernels register only IRv9 types, which don't include int4, so constant folding can't fold a Reshape of an int4/uint4 zero point. WebGPU Reshape didn't claim the 4-bit types either, so the graph failed to build. - int8 zero points pass only because CPU folds their Reshape. The DequantizeLinear kernel already handled int4/uint4, including odd sizes, so it is unchanged. - Reshape is a pure copy/view, so int4/uint4 are added from opset 21, the first opset whose schema allows them. 2. **Add int4 and uint4 QuantizeLinear output.** - int4/uint4 `y`/`y_zero_point` for opsets 21-22, 23, 24, 25-27 and 28. - Each invocation packs 8 nibbles into one u32 word, with the first element in the low nibble. Odd element counts work, and packed output bytes are never shared between invocations. - Results saturate to [-8, 7] / [0, 15]. The 4-bit zero point is read from packed u32 storage and sign-extended for int4. - Unchanged from microsoft#33683: the exact-quotient rounding, the `precision` filter in both GetCapability paths, and the int8/uint8 WGSL. The plugin EP looks kernels up in the same registry, so it claims the new types without changes to `ep/ep.cc`. Semantics and packing (two 4-bit values per byte, sign extension, blocked quantization, saturation ranges) follow @edgchen1's unmerged branch [`edgchen1/webgpu_quantizelinear`](https://github.com/microsoft/onnxruntime/tree/edgchen1/webgpu_quantizelinear). That branch refactors QL/DQ into a WGSL template plus `quantization_utils`, but it predates microsoft#33683. This PR applies the same semantics to main's current kernel instead of porting the refactor. Tests are pinned to WebGPU with CPU EP fallback disabled: - `onnxruntime/test/providers/webgpu/tensor/quantize_linear_test.cc`: - The WPT int4/uint4 QL and DQ vectors (f32 and f16), both directly and with the Reshape nodes Chromium inserts. - Rounding and saturation, multi-word per-axis and blocked zero points with odd sizes, all opsets, `output_dtype` without a zero point, and a Q -> DQ round trip. - `onnxruntime/test/providers/webgpu/tensor/reshape_op_test.cc`: int4/uint4 Reshape, including to a scalar, from both an input and an initializer. The existing CPU `QuantizeLinearOpTest.Int4/UInt4/OddLarge_*` tests now also run on WebGPU. Validation (Windows ARM64, NVIDIA RTX Spark N1X / D3D12, at 50323bf): - `onnxruntime_provider_test --gtest_filter="WebGpuQuantizeLinearTest.*:WebGpuDequantizeLinearTest.*:TensorOpTest.Reshape*:QuantizeLinearOpTest.*:DequantizeLinearOpTest.*:QuantizeLinearOp21BlockedTest.*:QuantizeLinearOpMLFloat16Test.*:QuantizeLinearPrecisionTest.*"`: 188 tests, 184 passed, 0 failed, 4 skipped. The skips are DML-only. - Sensitivity checks, both applied and rebuilt: - With the Reshape 4-bit types disabled (`since_version >= 99`), `TensorOpTest.Reshape_int4_webgpu` and `WebGpuDequantizeLinearTest.Int4WebNNConformance_Float/Float16` fail at session initialization, because the node lands on CPU. `Int4PerAxisMultiWordZeroPoint`, which has no Reshape, still passes. - With the nibble order swapped in the QuantizeLinear shader, every `WebGpuQuantizeLinearTest.Int4*` test fails on values. - WPT, using the plugin EP (`--use_webgpu shared_lib`) built from 50323bf with Edge Canary 159.0.4372.0 and ORT runtime 1.28.3: - `quantizeLinear`: 29/29 pass (was 15/29) - `dequantizeLinear`: 33/33 pass (was 27/33) - `reshape`: 67/67 pass - `qdq_subgraph`: 43/43 pass ### Motivation and Context  This fixes the WebNN conformance failures on the WebGPU EP: - `quantizeLinear.https.any.js`, 14 subtests (f32 and f16 each): - int4 zeroPoint of odd size and of even size - 2D with int4 zeroPoint of even size - 2D with int4 zeroPoint and block_size = [3, 2] - uint4 zeroPoint of odd size and of even size - 1D with uint4 zeroPoint and block_size = 3 - `dequantizeLinear.https.any.js`, 6 subtests: - uint4 and int4 1D tensors of odd size, with f32 and f16 1D scale - uint4 4D tensor with broadcasting f32/f16 4D scale and uint4 4D zeroPoint WebNN disables CPU fallback, so an unclaimed node makes the graph fail to build. --------- Co-authored-by: Aditya Rastogi <adityar@ntdev.microsoft.com> Co-authored-by: edgchen1 <18449977+edgchen1@users.noreply.github.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description Disable x86-specific tests and fix LabelEncoder.ExternalDataInKeysTensorOpset4 test on s390x. ### Motivation and Context These changes clean up some testsuite failures on s390x. --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

…icrosoft#29713) (microsoft#29792) ### Description Fixes microsoft#29713. On the CUDA EP, a device copy scheduled for a weight retained as an in-memory-external `OrtValue` initializer can have a source and destination that are the **same physical host buffer**. The host->host branch of `GPUDataTransfer::CopyTensor`/`CopyTensorAsync` did not handle this: - **Non-cuda-graph:** `ORT_ENFORCE(dst_data != src_data)` fired -> hard failure. - **CUDA-graph capture:** the `HOST_ACCESSIBLE` `cudaStreamSynchronize` is illegal during capture -> `CUDA 900` -> `cudaStreamEndCapture` `CUDA 901`, invalidating the capture. This became reachable once microsoft#29158 (initializer retention) landed. ### Fix Early-return the no-op self-copy at the top of the host->host branch in both methods, **before** the `HOST_ACCESSIBLE` stream sync (the key to keeping capture valid). This matches `CPUDataTransfer::CopyTensor` and the existing GPU->GPU guard. ### Testing Adds a CUDA EP internal regression test (`gpu_data_transfer_test.cc`) covering sync, async, and cuda-graph-capture self-copies. Verified failing before the fix (ORT_ENFORCE / CUDA 900->901) and passing after, on an RTX 3050 with a CUDA 12.8 `sm_86` build. --------- Co-authored-by: Trisha-sarvam <trishawadhwani@sarvam.ai> Co-authored-by: Cursor <cursoragent@cursor.com>

### Description  Remove ACL EP code. ### Motivation and Context  The ACL EP was already deprecated. Following up with the removal. ### Release notes draft ``` Removed the deprecated ACL (Arm Compute Library) execution provider, including its build options and standalone C registration API. Python's `get_all_providers()` no longer lists `ACLExecutionProvider`. Applications should remove this name from explicit provider lists. Requests containing it now follow unknown-provider handling: the default `InferenceSession` constructor can retry with fallback providers (usually CPU), while `set_providers()` or construction with fallback disabled raises an error. ``` --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…NT3 cold experts (microsoft#33916) ## Summary This PR fixes Windows WDDM and bounded-host-cache issues in QMoE CPU/GPU offload, adds tail-expert probability threshold pruning (`session.moe.offload.min_expert_weight`), and optimizes prefill and decode throughput for mixed-precision (`INT4` / `INT3` / `INT2`) QMoE models on consumer GPUs (e.g., RTX 4090 `24 GB` VRAM + `30 GiB` RAM): 1. **Fix Windows `FileExpertStore::InitializeMapped` & `QMoEOffload::InitializeBoundedHostCache` (`qmoe_offload.h`, `qmoe_offload.cc`, `qmoe_shared_layout.cc`)**: - Propagate `external_file_path` and `external_file_offset` from `MoeExpertSourceDescriptor` into `QMoEOffload::InitializeBoundedHostCache()`. - On Windows (`#elif defined(_WIN32)` in `FileExpertStore::InitializeMapped`), accept committed `MEM_MAPPED` (or `read_opt_enabled`) mappings via `VirtualQuery` when `source.file_path == nullptr`. 2. **Fix pinned memory budget accounting & persistent decode pinned starvation (`qmoe_offload.cc`)**: - Add `SessionPromotionBudget::HeldBytes()` and compute `reserved = bounce_held + std::max(decode_held, reserve_target)` in `QMoEOffload::ComputeShared` so promotion bounce allocations (`16 MiB`) are not conflated with persistent decode headroom (`24 MiB`). - In `QMoEOffload::EnsurePersistentDecodeState`, release idle promotion bounce and landing-ring pinned staging when `try_reserve_decode_pinned()` needs headroom. 3. **Automatic machine-aware session planning (`qmoe_session_controller.h`)**: - On Windows hosts where `ullTotalPhys < host_expert_bytes + 24 GiB`, automatically prefer direct OS file-backed mapping (`use_bounded_file_cache = false` unless `ORT_QMOE_FORCE_BOUNDED_HOST_CACHE=1`), avoiding multi-GiB private heap commits and enabling 2-bank parallel pinned wave staging. - For full-model sessions (`num_layers >= 24`, `total_vram >= 16 GiB`, `total_full_gpu_bytes >= 8 GiB`), automatically raise implicit `device_working_bytes` to `560 MiB` and implicit `pinned_bytes` to `160 MiB`, and adapt `admitted_gpu_expert_budget` against `platform_free_vram` (DXGI local budget) and `SessionBudgetAccount` without spilling into WDDM non-local shared memory. - Enable `EvictUploadedHostPages` (`VirtualUnlock`) on Windows `MEM_MAPPED` views. 4. **8K single-tile streaming prefill, 2-bank double-buffered wave staging, and default Windows `PrefetchVirtualMemory` (`qmoe_offload.cc`)**: - Default `stream_rows` to `8192` across `INT2`, `INT3`, and full-budget `INT4` layers. - Automatically halve `wave_experts` (`32 -> 16`) in `ComputeShared` when needed so `INT3` layers fit `stage_weight_banks = 2` (full double-buffered CPU staging + PCIe H2D DMA + GPU Tensor-Core GEMM overlap). - Extend `ORT_QMOE_WIN_PREFETCH_MMAP` (default `1` on Windows) to all file-mapped integer-bank layers. 5. **Tail-Expert Probability Threshold Pruning (`session.moe.offload.min_expert_weight`, `qmoe_kernels.cu`, `moe_kernels.cu`, `qmoe_offload.cc`)**: - Add `kOrtSessionOptionsMoeOffloadMinExpertWeight` (`session.moe.offload.min_expert_weight`, env fallback `ORT_QMOE_MIN_EXPERT_WEIGHT`) to skip non-top-1 routed experts (`rank >= 1`) whose probability after `SoftmaxTopK` is strictly below `min_expert_weight` (`topk_scales[route] = 0.0f`, `topk_indices[route] = -2`), always keeping at least the top-1 expert (`rank == 0`). - Support `ORT_QMOE_SKIP_TAIL_COLD_ONLY=1` to retain GPU-resident (`slot >= 0`) tail experts during decode at zero wall-time cost while pruning only CPU-cold (`slot < 0`) tail experts inside `QMoEOffloadClassifyKernel`. - Skip zero-weight routes in `QMoEOffloadMergeFinalizeCombineKernel` and `finalizeMoeRoutingKernel` without reading uninitialized contribution rows, and skip `input_dev -> host_input_out` PCIe copies in `QMoEOffloadClassifyKernel` when `has_doorbell && cold_count == 0`. 6. **CPU AVX2 `INT2` / `INT3` cold-expert & `SM89` dense `INT8` Tensor-Core GEMV optimizations (`qmoe_shared_layout_int2_avx2.cc`, `qmoe_shared_layout_int3_avx2.cc`, `dispatcher.h`, `paged_attention_impl.cu`)**: - Add 8-column `Int2Group8PrepackedAvx2` so all 8 columns of each 128-byte SM80 interleaved tile share one activation load while scaling weights before products. - Hoist activation vector loads outside `step_col` in `Int3Group4Avx2` and implement a true 8-column `Int3Group8Avx2`. - Enable `CanUseInt8TensorCore` on `params.device_sm >= 89` (`SM89`) and fix MSVC `/WX` `C4459` variable shadowing in `paged_attention_impl.cu`. ## Testing - `test_qmoe_offload_v4.py` targeted unit tests: **PASS**. - End-to-end Qwen3.8-Flash-Next RC5 (`qwen_flash_next_int4_int8_rc5_contrib_dflash2_qmoe_2b375_x02_y14_z32_dp23_b64`) benchmarks on RTX 4090 (`24 GB` VRAM) + Intel i9-13900 (`30 GiB` RAM): - **`2K` (`2,048` tokens)**: TTFT `8.92 s` (`7.55 s` at `min_expert_weight=0.08`), decode `12.15–23.43 tok/s` (`24.63–29.63 tok/s` with `min_expert_weight=0.065–0.08`). - **`8K` (`8,192` tokens)**: TTFT `14.42–15.32 s` (`10.84–12.30 s` with `min_expert_weight=0.065–0.08`), decode `23.04 tok/s` median / `29.49 tok/s` peak. - **`32K` (`32,768` tokens)**: TTFT `71.71–73.16 s`, decode `8.22–9.48 tok/s`. - **`128K` (`131,072` tokens)**: TTFT `536.21–579.64 s`, decode `7.42–15.27 tok/s`, peak VRAM `21,345 MiB`. - **100-Sample MMLU-Pro (`shard 0 of 8`, `max-new 8192`)**: - Default (`min_expert_weight=0`): **86/100 (`86.0%`)**, `0` truncated, **`27.39 tok/s` aggregate decode** (`3,196 s` wall time). - Tail-skip (`min_expert_weight=0.065`, `skip_cold_only=1`): **84/100 (`84.0%`)**, `0` truncated, **`34.70 tok/s` aggregate decode** (`43.21 tok/s` max, **`2,099 s` wall time**; `43.10 tok/s` avg / `53.42 tok/s` max at `0.080` on 20-item subset). --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

## Summary Allow native WebGPU hosts to persist Dawn shader/pipeline blobs without supplying their own device or relying on a global cache bridge. - Add a versioned public `OrtDawnCacheInterface` and the `dawnCacheInterface` provider option. - Attach host load/store callbacks and an optional isolation key when ORT creates the Dawn device, with bundled or external Dawn. - Validate trusted in-process pointer configuration, callback availability, interface version, external-device incompatibility, and shared-context interface identity. - Defer shared allocator context acquisition until first allocation, retaining the configured context for the allocator lifetime. - Serialize default-context acquisition and configured creation under one mutex, eliminating the lookup/fallback race without weakening explicit cache-identity checks. - Document ownership, thread safety, configuration order, and the companion GenAI requirement; add isolated configuration regression tests that handle exceptions and no-exceptions builds. - Exercise real cache callbacks through bundled Dawn, external Dawn, and the dynamic WebGPU plugin, including cold stores, warm hits, and isolation-key separation. The interface caches compiled shaders/pipelines, not model weights. Existing accelerated weight loading continues to use the ORT-owned device. Companion GenAI draft: microsoft/onnxruntime-genai#2893. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…soft#34002) ## Description Optimizes the existing CUDA MatMulPacked and IQ2_XXS MoEPacked paths without changing operator schemas, packed-weight layouts, or exported graphs. NVFP4 uses vectorized group decoding, paired BF16 operands, and half-chunk staging for eligible large-N Hopper projections. Q8 uses aligned packed-record windows with stride/alignment fallbacks; dense IQ2 reconstructs complete BF16 operands using the original K partitions. IQ2 experts use exact magnitude-codebook reconstruction and launch-selected aligned FC1 activation loads. FP32 scale reconstruction, sign-before-BF16 rounding, FC1/FC2 warp partitions, ReLU-square boundaries, and final reduction order are preserved. The shared LoadPacked8 fallback from the standalone-router extraction is reused. **Dependency:** microsoft#33994 in [native stack #33967](https://github.com/microsoft/onnxruntime/stack/33967). Mamba and standalone/integrated routing are unchanged. No requantization, GenAI inference change, graph conversion, or acceptance-threshold adjustment. ## Summary of Changes ### Packed Projection and Expert Kernels | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/quantization/gguf_quant_cuda.cuh` | Add vectorized NVFP4/Q8 loads and dense IQ2 BF16 operand reconstruction. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cu` | Select eligible optimized paths while retaining architecture/alignment fallbacks. | | `onnxruntime/contrib_ops/cuda/moe/moe_packed_iq2.cuh` | Add IQ2 expert codebook/operand helpers with exact FP32 scale reconstruction. | | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cu` | Use the expert helpers and aligned FC1 activation-load specialization. | ### Regression Tests | File | Change | |------|--------| | `onnxruntime/test/contrib_ops/matmul_packed_precision_test.cc` | Cover production projection widths, NVFP4 operand extremes, Q8 record windows/stride fallback, and IQ2 complete operands. | | `onnxruntime/test/contrib_ops/moe_packed_routing_test.cc` | Add production-width and IQ2 codebook/scale reconstruction cases. | | `onnxruntime/test/contrib_ops/cuda_kernels/matmul_packed_kernel_test.cc` | Pin the compiled TC contiguous-128-K reduction for small/staged NVFP4 paths. | | `onnxruntime/test/python/transformers/test_gguf_packed_ops.py` | Exercise default-rounded byte-offset IOBinding fallbacks without a model precision attribute. | ## Testing **Session-policy compatibility update:** head `6294c1fd47b493fc5dbbbed6cfad3cc406ca476e`, rebased onto the default-rounded MatMulPacked contract from microsoft#33962. The removed node attribute is no longer emitted by test model factories or assigned through internal launch parameters. Default numerical references round weights for both activation dtypes and every dispatch. **Validation at optimized stack tip `a3f2076b65e0d1a5c7e9b82165ffab014e233876`:** 15 native CUDA + 15 compute75-only PTX projection tests, 5 Python format/policy/alignment tests, and 1 actual internal CUDA contiguous-128-K cancellation test; zero failures/skips. CLANGFORMAT, Ruff/format and whitespace checks pass. These are focused ABI-matched overlays with fresh changed objects, **not a clean full-source build, full CI, performance remeasurement, or new full-model qualification**. The previous broader review qualification and historical performance results are preserved below. <details> <summary>Historical measurements and qualification evidence (not current-head acceptance)</summary> ### Earlier review qualification (historical) **Review-fix snapshot:** `74847ad6021d5b5431fa5984efb7fa95f6a310cd`. Added offset-buffer IOBinding fallbacks (NVFP4/Q8 and IQ2 ReLU-square MoE), nonuniform single-product and bitwise 8-vs-9-row MoE oracles, all-256-position IQ2 one-hot coverage, NVFP4 high-nibble coverage and an actual internal contiguous-128-K cancellation regression. Removed duplicate IQ2 table/decoder logic and dead branches; lowered the operand-extremes fixture gate to SM80. Seven wrong-mapping mutations produce nonzero output differences in the new oracle. **Owning-layer validation:** 20 C++ + 2 Python IOBinding + 1 actual internal CUDA cancellation test; zero failures/skips. **Integrated tip regression:** 51 C++ provider/schema tests, 5 actual internal CUDA tests and 2 Python IOBinding tests passed, zero failures/skips. All owning-layer provider cases are retained. The final overlay recompiles the changed CUDA/host/schema/test sources over ABI-matched preserved ORT support; **not a clean full-source build, full-CI approval, independent HF parity or new end-to-end qualification**. Targeted configured clang-format/Ruff and whitespace checks passed. These previously reported results are retained for provenance. They have not been rerun on the rebased head above; earlier CI/build status and benchmark results do not establish current-head CI, accuracy, or merge readiness. Original comparison pins: parent microsoft#33994 at `a8ff5642f5c85519561964928006413a4fde7859`; candidate `1b5063c4af651e12bf955bc978c9312e2734a9f9`. **This is the actual stack-tip comparison**, not the earlier private-branch versus microsoft#33940 comparison. That distinction matters: the remote stack tip did not yet contain the optimized native kernels. | Input tokens | Stack ordinary tok/s | Candidate ordinary tok/s | Stack k=7 tok/s | Candidate k=7 tok/s | |---|---:|---:|---:|---:| | 178-254 | 54.36 | **92.42** | 223.84 | **355.42** | | 427-932 | 54.20 | **92.08** | 185.24 | **293.62** | | 1028-1663 | 54.25 | **92.04** | 201.48 | **318.91** | Approximately **70% ordinary** and **58-59% k=7** improvement in this controlled comparison. This does not claim that all remaining latency is resolved. The separate microsoft#33940 equivalent-package result of approximately 109 ordinary / 370 k=7 tok/s on the longest bin is **not** a measurement of this candidate; transferring that remaining gain is still outstanding. Protocol: original frozen 30 real prompts, ten per bin, unchanged warmups and per-prompt mode order, natural EOS/512-new-token cap, two reversed-order passes, same physical H100 per baseline/candidate shard. Eight logical shards ran in waves over GPUs 1,3,4,5,6,7. All admitted runs passed the foreign-client and resource guards. No profiler ran in the timing window; clocks were not locked. Rates pool decoded tokens over elapsed decode time, excluding the first token, not arithmetic means of per-request TPS. ## Correctness and build validation - Exact same native model package, packed bytes and GenAI libraries in both arms. Graph hashes match the preserved source-value audit; full package-file hashes were recorded again. No exported checkpoint/package was modified. - **22 native CUDA regressions passed, zero skips**, covering projections, complete operand rounding/cancellation, IQ2 production widths/codebook scales, routing and the standalone router. Changed C++ files passed CLANGFORMAT; git whitespace checks passed. - **1,782 complete live tensor/state hashes per arm**, bitwise equal and finite in ordinary and k=7 modes, also equal to the preserved quantized control. - Independent full LCB145 ordinary/k=7 probes passed all **512 generated tokens** and complete counters. - The accepted timing run passed **240 complete measured generation sequences** against frozen ordinary-control tokens. Within-mode complete speculative counter dictionaries matched between arms and repeats. - Isolated incremental rebuild on the archived clean microsoft#33994 build; every affected compilation unit/header dependent was rebuilt and the provider/test binary relinked with the production flags. Untouched objects were reused: **this is not a new clean full-source build**. - CUDA 12.8.93, driver 580.159.04, H100 80GB. Common corrected GenAI source: `c0fd01914265d5157d6da1376c38b4df521a76c5` ([microsoft/onnxruntime-genai#2916](microsoft/onnxruntime-genai#2916)). Both arms use that same source; its correctness fix is not credited with the kernel speedup. Accepted statistics SHA256: `0b9b7e2385b2dd9d7e5f7d459c65ef74338b5f707d1536454e7e27fb9a407c0f`. </details> ## Checklist - [x] Review feedback addressed with focused regressions - [x] Documentation and supported contracts updated - [x] Native stack order and immediate-parent dependencies preserved - [ ] Clean full-source/full-CI qualification (not claimed by focused overlays) ### Python IQ2 alignment regression correction Commit `f6e5af95c1` migrates `test_moe_packed_relu_iq2_aligned_input_fallback` from the removed seven-input compact-routing interface to six-input full-score routing: `[X, R, G, "", D, W]`, with FP32 selection scores and mixing values. Explicit `normalize_routing_weights=0`, buffer offsets and the exact bitwise oracle remain unchanged. No production code changes. The original test was reproduced failing graph resolution (7 inputs; schema maximum 6). The corrected actual Python unittest passed on H100 with zero skips, including all four combinations of rows 8/9 and activation offsets 0/2. Profiles confirm four `CUDAExecutionProvider/MoEPacked` executions. Ruff lint/format and whitespace checks passed. This run used the retained six-input full-score ORT 1.32.0 runtime built from `3bf5237cd65131701f2fb05435774672afaf1f96`; it is not a fresh runtime build of this PR, a full Python suite, or full CI approval. The corrected test file is byte-identical in the synchronized microsoft#34003 head. --------- Signed-off-by: rui-ren <ruiren1225@gmail.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: f4d0e1a3-a389-4cc8-9c85-0fd0e3ffce5b

…osoft#34003) ## Description Reduces CUDA packed decode tile/loader overhead above microsoft#34002. Narrows eligible eight-row IQ2 FC1 tiles, adds aligned record-window reconstruction for eligible IQ2 FC2 rows, and uses full-tile NVFP4 paths only when their shape/alignment guards hold. Unsupported geometries retain the existing fallback. **Dependency:** microsoft#34002 in [native stack #33967](https://github.com/microsoft/onnxruntime/stack/33967). No schema, packed layout, cache/mask precision workaround, or GenAI source change. The R8-prefill/R1-ordinary hybrid session/package work is in the companion microsoft/onnxruntime-genai#2916, **not implemented by this native-kernel PR**. ## Summary of Changes ### Packed Decode Specializations | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cu` | Select narrower FC1 tiles and eligible aligned FC2 loader variants. | | `onnxruntime/contrib_ops/cuda/quantization/gguf_quant_cuda.cuh`, `onnxruntime/contrib_ops/cuda/moe/moe_packed_iq2.cuh` | Reuse the shared IQ2 decoder for aligned word windows, including two-byte-offset records. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cu` | Specialize eligible full NVFP4 output tiles while preserving guarded fallbacks. | | `onnxruntime/test/contrib_ops/moe_packed_routing_test.cc` | Exercise eight-row width/routing boundaries and larger intermediate widths. | ## Testing **Review-fix snapshot:** `ed6e6e1991fd944f65f1d97472d3f29046d97361`. Bound aligned-window FC2 to four resident CTAs, kept the R8-specific narrow FC1 gate decode-only (`!use_prefill_tiles`), simplified the FC2 guard to its real row/base alignment invariant, and reused the lower shared decoder. Nonuniform regression protection is introduced in microsoft#33963/microsoft#34002. Mamba replay changes were moved into their owning PR microsoft#33969; this layer no longer changes Mamba numerics or tests. **Owning-layer validation:** Integrated stack regression run and exact-production-kernel H100 A/B below. **Integrated tip regression:** 51 C++ provider/schema tests, 5 actual internal CUDA tests and 2 Python IOBinding tests passed, zero failures/skips. All owning-layer provider cases are retained. The final overlay recompiles the changed CUDA/host/schema/test sources over ABI-matched preserved ORT support; **not a clean full-source build, full-CI approval, independent HF parity or new end-to-end qualification**. Targeted configured clang-format/Ruff and whitespace checks passed. ### Review-requested kernel A/B (H100, CUDA 12.8.93) Exact production kernel/helper text from the integrated head; random nonuniform IQ2, 512 experts, 16 rotating route sets, R=8. All variant outputs are bit-identical across all route sets. CUDA-event median: 20 warmups, 9 alternating batches x32 launches; no full-model/cold-cache guarantee. `k=9` is decode, `k=17` is four-subtile prefill; the latter remains excluded from the narrow dispatch. | Hidden | Inter | k | FC1 wide→narrow μs | FC2 plain / window / hinted μs | |---:|---:|---:|---:|---:| | 1024 | 1024 | 9 | 15.035 → 10.888 | 12.874 / 11.966 / 12.062 | | 1024 | 1024 | 17 | 16.264 → 14.610 | 18.149 / 15.968 / 16.188 | | 1024 | 2048 | 9 | 18.041 → 15.920 | 21.786 / 20.282 / 20.458 | | 1024 | 2048 | 17 | 26.279 → 23.879 | 32.182 / 28.613 / 28.801 | | 2048 | 1024 | 9 | 27.187 → 18.500 | 20.663 / 19.135 / 19.234 | | 2048 | 1024 | 17 | 28.614 → 25.400 | 29.040 / 25.962 / 26.176 | | 2048 | 2048 | 9 | 31.466 → 27.515 | 34.251 / 32.346 / 32.446 | | 2048 | 2048 | 17 | 48.462 → 42.877 | 51.977 / 45.898 / 46.270 | | 4096 | 1024 | 9 | 49.654 → 32.181 | 30.966 / 29.015 / 29.088 | | 4096 | 1024 | 17 | 53.387 → 46.004 | 50.201 / 45.097 / 45.371 | | 4096 | 2048 | 9 | 57.550 → 50.024 | 53.390 / 50.104 / 50.318 | | 4096 | 2048 | 17 | 91.847 → 79.396 | 90.444 / 80.277 / 80.838 | The local 12.8 compiler **did not reproduce** the reported 79-register/3-CTA window allocation: unhinted and hinted window kernels both use 64 registers/4 CTAs here. The selective hint is retained as a portability cap for the reviewer's newer-toolchain result, with its local timing tradeoff shown rather than hidden. These kernel-only measurements do not attribute the older hybrid end-to-end gains below to the later narrow-tile commit. <details> <summary>Historical measurements and qualification evidence (not current-head acceptance)</summary> These previously reported results are retained for provenance. They have not been rerun on the rebased head above; earlier CI/build status and benchmark results do not establish current-head CI, accuracy, or merge readiness. The following results include the companion GenAI runtime and hybrid package, not just this PR's native kernel changes. They use the original pinned binaries and package hashes below, not the current rebased ORT head. **Unprofiled frozen real30, six H100s, two reversed same-GPU passes, 240 complete measured generations; original warmups, mode order, EOS/512 cap.** | Actual input tokens | Before ordinary tok/s | Hybrid ordinary tok/s | Gain | Original R8/k7 tok/s | |---|---:|---:|---:|---:| | 178-254 | 94.40 | 119.95 | +27.1% | 361.46 → 361.57 | | 427-932 | 94.01 | 119.59 | +27.2% | 298.68 → 298.94 | | 1028-1663 | 93.89 | 118.90 | +26.6% | 324.25 → 324.31 | Ordinary TTFT medians: short 357.35→359.00ms, medium 826.97→831.29ms, long 2153.15→2161.92ms. The earlier static-R1 version's prefill regression is **not** the package being qualified here. Separate same-GPU synthetic-2048 CUDA profiles, 64 captured target passes: IQ2 expert work **3.259→1.496ms/pass**, Mamba **0.824→0.423ms/pass**. These are device-work attribution, not an additional throughput claim. **Correctness:** all240 exact same-mode sequences and complete counters match; 388 complete live committed-state boundaries (66.43GB compared), capture off/on, two resets, prefix through130; original R8 ordinary/k7 retains1782 complete finite tensor hashes. Three native kernel tests, seven Mamba operator tests, fifteen GenAI tests execute without skips. The new replay test also passes CUDA memcheck with zero errors (100 repetitions). **Scope:** hybrid is ordinary-only. Both speculative arms use the original R8 package. The benchmark explicitly keeps two per-mode models resident in each arm; it is not a unified hybrid-speculative package or a k7 speedup. No precision/cache/mask workaround, no weight repack, no independent HF rerun, no clean-build/full-CI claim. Reproduce the exact tested graph/config with: ```bash python examples/python/nemotron-bolt-hybrid.py --input <packed-r8> --output <fresh-same-filesystem-hybrid> ``` Requires onnx-ir/numpy/ml-dtypes and both updated runtimes. External-data validation stays enabled; payloads are hardlinked inside the output package, not copied. Pins: - ORT: `0357a7dfe47d3c383bf36be62015162b6a612d5e` - Compiled/common GenAI: `d73a6e627d6d20a5a1688b9fbd46237069a7be89` - Published GenAI: `db6ffd29fc96e319f645ff461efa05c643235f29` (converter/docs-only child) - CUDA provider: `fac33c75d0336236736aa96456268057b628d93d0b9901d49598e275bf335275` - GenAI shared library: `f977660a408bc0a01a31151c7280282ba443fa256785680ae3e841654eec9be7` - Tested hybrid graph: `aaf5d6ae06b6740ab5b3af523359581930d2b6c42c38a18529a22737c2cfe530` - Config: `77f2d920d2b04e17884c3d2ee691d9053225b957c047454ebd0a4ea3649dac44` - Unchanged packed payload: `06d3d508e8d11235d0b770484b6b59764c2f297c3383b6d0de75648d26c11867` Full commands, per-case exact tokens/counters, raw kernel traces, provider maps, UUID/admission receipts, and rejection evidence remain on the validation host under `/opt/bolt-runtime-fixes-20261108-v1/evidence/ordinary-r1-v1`. Sealed manifest SHA256: `412b3898f09565557b01e82906358f017f9df3227db2bf330850245a52d32bbc`. Root-disk binaries were reloaded and verified with full ordinary/k7 generation; the large model payload still resides on the existing tmpfs, not reboot-persistent. Original package/deployment untouched; no merge. </details> ## Checklist - [x] Review feedback addressed with focused regressions - [x] Documentation and supported contracts updated - [x] Native stack order and immediate-parent dependencies preserved - [ ] Clean full-source/full-CI qualification (not claimed by focused overlays) --------- Signed-off-by: rui-ren <ruiren1225@gmail.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: f4d0e1a3-a389-4cc8-9c85-0fd0e3ffce5b

…ron Bolt (microsoft#34147) ## Description This PR stacks on top of microsoft#34003 to close the remaining prefill and decode latency gap between native `com.microsoft` contrib ops (`MoEGroupRouter`, `MoEPacked`, `MatMulPacked`) and the `v11_gguf` custom-op baseline on the 88-layer `NVIDIA-Nemotron-3.5-Bolt-74B-A9B` model (`rc7`). It extends `MoEPacked` and `MatMulPacked` with optional fusion attributes (`routed_scaling_factor`, `routing_epsilon`, `activation_type`) so per-layer Top-K routing normalization/scaling and shared-expert ReLU² epilogues fuse directly into the operator kernels, and optimizes the CUDA `MoEGroupRouter`, `MoEPacked`, and `MatMulPacked` kernels on Hopper (`sm_90`). ## Summary of Changes ### Operator Schema & Documentation | File | Change | |------|--------| | `onnxruntime/core/graph/contrib_ops/contrib_defs.cc` | Add `routed_scaling_factor` (`FLOAT`, default `1.0`) and `routing_epsilon` (`FLOAT`, default `0.0`) to `com.microsoft::MoEPacked(1)`, and `activation_type` (`STRING`, default `"none"`, supporting `"none"`, `"relu2"`/`"relu_squared"`, `"silu"`) to `com.microsoft::MatMulPacked(1)`. | | `docs/ContribOperators.md` | Update generated contrib operator schema documentation for `MatMulPacked` and `MoEPacked`. | ### CUDA Kernel Optimizations | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_group_router_impl.cu` | Fuse `AddSelectionBiasKernel` into `SigmoidScoresKernel`, `SigmoidScoresR8Kernel`, and `SigmoidLogitsKernel` (eliminating 40 separate kernel launches per step); add `SigmoidScoresR8Kernel<CudaT, kAligned>` (`<<<experts, 128>>>`, 4 warps splitting `K`) so `rows == 8` speculative verify reads `projection_weight[e, :]` once for all 8 rows (`11.89 us -> 5.76 us/layer`). | | `onnxruntime/contrib_ops/cuda/moe/moe_packed.h`<br>`onnxruntime/contrib_ops/cuda/moe/moe_packed.cc`<br>`onnxruntime/contrib_ops/cuda/moe/moe_packed.cu` | Fuse `(selected_weight / (sum + routing_epsilon)) * routed_scaling_factor` into `FullScoreRoutingKernel` and `FullScoreRoutingGenericKernel`; launch 1 warp/CTA (`<<<rows, 32>>>`) when `rows <= 64` (`18.18 us -> 9.81 us/layer`); hoist weight row pointers and use 2-byte aligned `LoadIq2SubBlock16Stage` + `DecodeIq2StageToBf16` with `__launch_bounds__(KW * 32, ...)` in `MoEPackedReluIq2ContiguousKKernel` across all `K` (including `K=1280` 5-block FC2). | | `onnxruntime/contrib_ops/cuda/quantization/gguf_quant_cuda.cuh` | Add `LoadIq2SubBlock16Stage<kAlignedU16>` and `DecodeIq2StageToBf16` (using `bf16x2` sign-mask bit injection before FP32 scale multiplication and single `bfloat162` rounding) and `LoadQ8SubBlock16AlignedStage` + `DecodeQ8StageToBf16` to eliminate intermediate float arrays and reduce register pressure. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_packed.h`<br>`onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cc`<br>`onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cu` | Add `activation_type` epilogue fusion (`ApplyMatMulPackedActivation`) across Tensor Core, GEMV, and prefill paths, and use `LoadIq2SubBlock16Stage`/`DecodeIq2StageToBf16` and `LoadQ8SubBlock16AlignedStage`/`DecodeQ8StageToBf16` (`Q8_0` LM head `310.8 us -> 266.3 us`, `IQ2_XXS` `sh_dn` `19.30 us -> 14.77 us`). | ### Unit Tests | File | Change | |------|--------| | `onnxruntime/test/python/transformers/test_gguf_packed_ops.py` | Add `test_matmul_packed_activation_types`, `test_moe_packed_routed_scaling_factor_and_routing_epsilon` (including `inter_size=1280` 5-block IQ2_XXS), and `test_moe_group_router_r8_and_fused_bias`. | | `onnxruntime/test/contrib_ops/matmul_packed_precision_test.cc` | Add `ActivationTypeRelu2AndSiluAcrossDispatchPaths`. | | `onnxruntime/test/contrib_ops/moe_packed_routing_test.cc` | Add validation tests for nonfinite `routed_scaling_factor` and negative/nonfinite `routing_epsilon`. | ## Motivation and Context On NVIDIA H200 (`sm_90`, CUDA 13.0) with `NVIDIA-Nemotron-3.5-Bolt-74B-A9B` (`rc7`, 88 layers, `k=7` MTP verify `rows=8`, CUDA graph enabled): | Metric | `rc7` Baseline (PR microsoft#34003) | `rc7` + This PR | `v11_gguf` Custom Ops | Delta vs PR microsoft#34003 | |---|---:|---:|---:|---:| | **ONNX Graph Nodes** | 2,912 | **1,272** | 1,155 | **-1,640 nodes (-56.3%)** | | **CUDA Graph Kernel + Memcpy Launches / Round** | 2,009 | **729** | 690 | **-1,280 launches (-63.7%)** | | **`nsys` CUDA Graph Replay GPU Time / Round** | 16.562 ms | **12.408 ms** | 11.596 ms | **-4.154 ms (-25.1%)** | | **MMLU-Pro 200 Accuracy (`k=7`)** | 119/200 (59.50%) | **121/200 (60.50%)** | 119/200 (59.50%) | **+1.00%** | | **Prefill Throughput (`500` prompts, `k=7`)** | 536.5 tok/s | **729.3 tok/s** | 685.6 tok/s | **+35.9% (+6.4% vs `v11_gguf`)** | | **Decode Throughput (`k=7`, normal EOS)** | 268.88 tok/s | **349.91 tok/s** | 378.89 tok/s | **+30.1%** | | **Decode Throughput (`k=7`, `--disable-eos`)** | 304.67 tok/s | **394.65 tok/s** | 428.88 tok/s | **+29.5%** | | **Decode Throughput (`k=0`, no MTP)** | 77.60 tok/s | **106.10 tok/s** | 109.76 tok/s | **+36.7% (96.7% of `v11_gguf`)** | ## Testing - Ran `onnxruntime/test/python/transformers/test_gguf_packed_ops.py` on H200 (`sm_90`), including all existing and newly added `MatMulPacked`, `MoEPacked`, and `MoEGroupRouter` tests. - Validated MMLU-Pro 200 accuracy (`121/200 = 60.50%`), prefill (`500` prompts), and decode (`100` prompts, `k=7` and `k=0`) on `NVIDIA-Nemotron-3.5-Bolt-74B-A9B` (`rc7`). ## Checklist - [x] Tests added/updated - [x] No breaking changes (all new attributes have backward-compatible defaults) - [x] Documentation updated (`docs/ContribOperators.md`) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…2565) ### Description Make the min-blocks-per-SM hint architecture-conditional via `__CUDA_ARCH__`, which nvcc redefines on each per-architecture device-code compilation pass. Below SM89 the hint drops to `1` (a no-op occupancy hint on dead code); at SM89+ it stays `3`, preserving the residency-pinning behavior from microsoft#32433 unchanged. No behavior change on SM89+. No currently-selected shape or dispatch path is affected outside of fixing compilation for pre-Ada targets. ### Motivation and Context `MatMulBlockScaledFp8MmaGemvKernelPinned` (added in microsoft#32433) carries `__launch_bounds__(32 * KSplit, 3)`. For the only instantiation that's ever launched (`KSplit = 16`), that's `__launch_bounds__(512, 3)`, which requires 512 * 3 = 1536 threads resident per SM. `Fp8MmaGemvPinsResidency` in `matmul_block_scaled_fp8_tiling.h` gates this entry point to SM89+ (Ada), where the max resident threads/SM is exactly 1536 — the intended fit. But nvcc still codegens the template instantiation for every architecture passed to `--generate-code`, and on Turing (`sm_75`, max 1024 threads/SM) 1536 exceeds the hardware limit. `ptxas` rejects it: ``` ptxas error : Value of threads per SM for entry ...MatMulBlockScaledFp8MmaGemvKernelPinned... is out of range. .minnctapersm will be ignored ptxas fatal : Ptx assembly aborted due to errors ``` This breaks any build that includes `sm_75` in its target architectures, even though the kernel is unreachable there at runtime. Note that Nvidia T4 (entry-level GPU available in AWS and GCP) is sm_75. ### Tests Rebuilt `matmul_block_scaled_fp8.cu.o` for `sm_75`, `sm_80`, `sm_89` — compiles clean (previously failed with the `ptxas` error above on `sm_75`). Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com>

@zadesai

…crosoft#34159) ### Description Forward the initializer-data callback into `OrtGraphToProto` subgraph recursion, keeping the same mutable callback state throughout a traversal. Preserve the empty-callback path and both public function signatures. The recursive call currently omits the callback, so nested initializers silently stay inline and nested callback errors are never reported. Passing the `std::function` by value instead would copy mutable captures into each branch, potentially reusing external-data offsets. A conditional `std::ref` wrapper shares the active callback during synchronous recursion without wrapping an empty function. Adds generated, CPU-only regressions in the existing EP graph suite: - root, sibling and depth-two initializers, including same-named initializers with different values; - a mutable value-captured offset, checking every external-data location/offset/length and the corresponding bytes; - external-file reload and actual inference for every `If` branch; - both GraphProto/ModelProto overloads, omitted/empty callbacks, and exact nested error propagation with no subsequent callbacks. #### Validation Built and tested on macOS arm64, Apple Clang 17, CPU-only Release shared-library build, based on `a8246c0d094f1ab0d6bf236228e5e67683321cd7`. No minimal build or operator exclusions were used. ```sh source .venv/bin/activate python tools/ci_build/build.py \ --build_dir build/MacOS --config Release --update --build --parallel 4 \ --target onnxruntime_ep_graph_test --build_shared_lib \ --cmake_generator Ninja --skip_submodule_sync --skip_tests \ --cmake_extra_defines FETCHCONTENT_TRY_FIND_PACKAGE_MODE=NEVER # Target-only builds do not run the general test target's fixture-copy step. cmake -E copy_directory onnxruntime/test/testdata build/MacOS/Release/testdata cd build/MacOS/Release ./onnxruntime_ep_graph_test --gtest_filter='EpGraphTest.*' ``` - **Before:** both added regressions fail against the unmodified header. The nested fixture serializes eight initializers but invokes the callback once; nested rejection incorrectly returns success. - **After:** all **20 EP graph tests pass**, including the two added regressions; no skips. - **Mutation checks:** forwarding the callback by value fails the offset/data/output assertions; unconditionally wrapping the empty callback fails the inline serialization control. The intended fix was restored, rebuilt and the full suite rerun afterward. - `lintrunner` and `git diff --check` pass. <details> <summary>Selected output from the actual local runs</summary> Baseline: ```text Expected equality of these values: callback_calls Which is: 1 inline_initializers.size() Which is: 8 Value of: status.IsOK() Actual: true Expected: false [ PASSED ] 0 tests. [ FAILED ] 2 tests ``` Restored final patch: ```text [==========] 20 tests from 1 test suite ran. (48 ms total) [ PASSED ] 20 tests. ``` </details> Windows execution has not been tested locally. Existing macOS telemetry duplicate-class and model-fixture warnings were not suppressed. ### Motivation and Context Fixes microsoft#34137. Thanks to @zadesai for the reproducer and the reference-forwarding approach. This fixes the existing header-only EP utility. microsoft#33921 concerns FLOAT16 tensor attributes, while microsoft#34014 adds a separate serialization C API; neither fixes this callback omission. AI assistance: OpenAI Codex assisted with implementation, independent source review, and execution of the reported local tests. The results above are from those actual automated runs. Signed-off-by: Silu Panda <31051721+SiluPanda@users.noreply.github.com>

… libraries (microsoft#34126) Windows directory junction checks in model_package currently rely on the MSVC-specific filesystem file_type junction enumerator, which fails to compile on Windows builds using libc++ or other C++ standard libraries that only expose standard file_type values and classify NTFS junctions as directories. This change queries Win32 file attributes and the reparse tag directly on Windows when symlink_status does not already report a symbolic link, rejecting name-surrogate reparse points such as NTFS junctions across all C++ standard libraries while preserving fast early return on ordinary files and directories. Validated on Windows by building and running the model_package test suites, including root, nested, and shared_assets junction rejection tests, and by verifying clean warning-free compilation under clang-cl with libc++. Co-authored-by: Mustapha Jaber <mustjab+microsoft@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…icrosoft#34099) ### Description Since `huggingface_hub` 1.32, the Hugging Face Hub cache stores Xet files once per cache in a [blob store shared across repos](https://huggingface.co/docs/huggingface_hub/guides/manage-cache#shared-blobs-across-repos) (`<cache>/blobs/<hash prefix>/<hash>`), and the repo's `blobs/<etag>` becomes a relative symlink to it. A model and its external data in one snapshot can then resolve into different folders (for example `blobs/28/…` and `blobs/d4/…`, or one in the repo's `blobs/` and one in the shared store). The symlink fallback from microsoft#27374 only accepts external data next to the resolved model, so models with external data loaded from a default `hf_hub_download` path fail with `External data path escapes model directory`. This PR teaches both external data validators about the shared store without trusting anything outside the cache's blob folders. ### Summary of Changes #### Validation (`onnxruntime/core/framework/tensorprotoutils.cc`) New Pass 3 in `ValidateExternalDataPath`, reached only when the model is a symlink and Passes 1–2 fail. If the canonical model directory is in a snapshot of a cache repo (`<cache>/<type>s--<repo id>/snapshots/<revision>/…`, also from a subfolder such as `onnx/`), external data is accepted when: 1. its path does not leave the model directory before symlinks are resolved (no `..` escape), and 2. both the model and the external data resolve to files under the blob stores of that cache: the repo's `blobs/` folder or the cache-wide `blobs/` folder. `ValidateOpenedExternalDataPath` applies the same blob-store check to the canonical path of the opened file, so the opened-file validation from microsoft#33063 also covers the new case. The cost is one extra canonicalization, and only after the existing passes fail. | File | Change | |------|--------| | `onnxruntime/core/framework/tensorprotoutils.cc` | Add `IsInHuggingFaceCacheBlobStores` and use it in both validators. | | `onnxruntime/test/framework/tensorutils_test.cc` | `HuggingFaceCachePathValidationTest`: 11 tests on a cache built like `huggingface_hub` does, with two-hop relative symlinks (6 accepted layouts, 5 rejected). | | `onnxruntime/test/python/onnxruntime_test_python_symlink_data.py` | End-to-end `InferenceSession` test for the shared store and for a mixed repo-local/shared layout. | #### Security notes - Paths are compared in canonical form. A canonical path has no symlinks, so a file whose canonical path is under `<cache>/blobs` is stored in that real folder. If a `blobs/` folder is itself a symlink or junction, nothing matches (`huggingface_hub` also never uses a symlinked `blobs/` folder as the shared store). - The model itself must resolve into the blob stores. This rejects symlinked models that point outside the cache, and `symlink/..` tricks in the model path. - The no-`..` requirement keeps a malicious model from reaching other cached repos' files. In a real cache, the entries of a snapshot folder are symlinks that `huggingface_hub` creates to files of the same repo. - Someone who can write into the cache root gains nothing new. With the existing Pass 2 (for example ORT 1.31.1), they can already make a model read any blob by putting the model file in the same shard folder (verified). `huggingface_hub` also gives the store and shard folders the same permissions as the cache root. - Unchanged: empty, absolute, and in-memory-tag locations are still rejected before any fallback, and Passes 1–2 behave as before. #### Comparison with microsoft#33974 microsoft#33974 targets the same issue. It also checks the `.huggingface-shared-blobs` marker (version `1`) and blob names (hex, 40/64 characters, 2-character shard), and requires a `models--` repo folder. These checks are not security boundaries, since anyone who can create the layout can also create the marker and names. They do tie ORT to `huggingface_hub` internals, so a future layout change would break loading again. I built both and loaded the same synthetic caches: | Scenario | microsoft#33974 | This PR | |---|---|---| | Shared store, model and data in different shards | loads | loads | | Mixed repo-local/shared blobs (both directions) | loads | loads | | `datasets--…` repo | fails | loads | | Model folder reached through a symlinked directory | fails | loads | | Shared store marker version other than `1` | fails | loads | | Location `../x`; data symlink outside the cache; symlinked `blobs/` | rejected | rejected | ### Motivation and Context Fixes microsoft#33950. Downstream projects currently work around it by pinning `huggingface_hub<=1.31`, setting `HF_HUB_DISABLE_SHARED_BLOBS=1`, or copying files out of the cache. ### Testing - Real cache: reproduced the issue with `onnxruntime==1.31.1` and `huggingface_hub==2.4.0` (`onnx-community/sam2.1-hiera-tiny-ONNX`, fresh `HF_HOME`). With this PR, the same cache loads. The two mixed layouts created by downloading one of the files with `HF_HUB_DISABLE_SHARED_BLOBS=1` also fail with 1.31.1 and load with this PR. - `onnxruntime_test_all --gtest_filter='HuggingFaceCachePathValidationTest.*:PathValidationTest.*:GetExtDataFromTensorProtoTest.*'`: all pass. Without the `tensorprotoutils.cc` change, the 6 tests for accepted layouts fail. Removing any one of the conditions (no `..` escape, model in a blob store, `blobs/` folder bound, repo folder name, symlinked `blobs/`, opened-file check) makes the targeted test fail. - `python -m pytest onnxruntime/test/python/onnxruntime_test_python_symlink_data.py`: 4 passed (plus 2 subtests). - Full `onnxruntime_test_all` on Linux x64 (CPU, Release): 2897 passed. `lintrunner`: no issues. - Not run locally: Windows and WASM. The new tests skip when symlinks cannot be created, and the load test is excluded on WASM like the other external-file tests in that file. ### Checklist - [x] Tests added/updated - [x] No breaking changes - [ ] Documentation updated (not applicable) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

microsoft#34164) ## Description This PR builds on microsoft#34147 (`main`) to enable and accelerate large-chunk (`R` up to `4,096`) prefill across the native `com.microsoft` CUDA contrib operators (`Mamba2Verify`, `MoEPacked`, `MoEGroupRouter`, `MatMulPacked`, and `GroupQueryAttention`) for `NVIDIA-Nemotron-3.5-Bolt-74B-A9B`. Combined with the companion `onnxruntime-genai` dynamic-prefill runtime in [microsoft/onnxruntime-genai#2961](microsoft/onnxruntime-genai#2961), prefill throughput on NVIDIA H200 (`sm_90`) reaches **`10,185 tok/s` (`k=7`) / `10,429 tok/s` (`k=0`) at `L = 2048`** and **`11,424 tok/s` (`k=7`) / `11,587 tok/s` (`k=0`) at `L = 3968`** — **`2.45×–2.67×` faster than `llama.cpp` peak (`ub=4096`)** and **`13.98×–15.73×` faster than fixed `R=8` prefill**. ## Summary of Changes ### `Mamba2Verify` Large-`R` Prefill & Right-Aligned Snapshots | File | Change | |------|--------| | `onnxruntime/core/graph/contrib_ops/fused_state_space_defs.h`<br>`onnxruntime/contrib_ops/cuda/sequence/fused_state_space.cc`<br>`docs/ContribOperators.md` | Lift the `R <= 16` and `slots >= R` constraints (`R >= 1`, `slots >= 1`). When `R <= slots`, behavior is unchanged (`0..R-1` written, `R..slots-1` preserved). When `R > slots`, only the trailing `slots` tokens (`R - slots .. R - 1`) write snapshots `0 .. slots - 1` so slot `slots - 1` holds the final recurrent state without writing intermediate HBM snapshots. | | `onnxruntime/contrib_ops/cuda/sequence/fused_state_space_impl.cu` | Add 2D sequence-and-channel tiled `MambaConvPrefillKernel`, race-free `MambaConvWriteSnapshotsKernel` (safe under in-place `conv_state` aliasing), and `MambaSsmPrefillKernel<4, 16>` (register-resident state scan with 36 KiB tiled shared-memory staging and exact 4-partition readout parity). Aliasable state pointers are not `__restrict__`, and the grouped RMS normalization launch is sliced into 65,535-row grid-y chunks so any `R <= INT_MAX` launches. | ### `MoEPacked`, `MoEGroupRouter`, `MatMulPacked` & `GroupQueryAttention` Prefill Optimizations | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_packed.h`<br>`onnxruntime/contrib_ops/cuda/moe/moe_packed.cu` | Raise `kMoEPackedMaxChunkRows` to `4096` so `L <= 4096` executes in a single chunk, and add `BuildMoEPackedTilesKernel<64>` + `MoEPackedReluIq2PrefillTcKernel<kFc1, 8>` so each IQ2_XXS weight sub-block is decoded once per up-to-64 routes in shared memory instead of once per 8 routes. Each MMA accumulator chain covers one 64-wide K step and partials are combined with FP32 adds, matching the rounding of the warp-split decode kernels (~1.4% kernel time on A100). | | `onnxruntime/contrib_ops/cuda/moe/moe_group_router_impl.h`<br>`onnxruntime/contrib_ops/cuda/moe/moe_group_router_impl.cu`<br>`onnxruntime/contrib_ops/cuda/bert/linear_attention_gates.cc` | Add a zero-workspace `cublasGemmEx` Tensor Core GEMM + in-place `SigmoidLogitsKernel` path for `rows > 16`, while keeping `SigmoidScoresR8Kernel` for `rows == 8`. The schema now states that the FP32 projection reduction order depends on the row count. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cu` | Add templated `MatMulPackedBf16DequantFastKernel` (`IQ2_XXS`, `NVFP4`, `Q8_0`), widen BF16 prefill tile width up to `18560` so projections execute in a single `cublasGemmEx` call, and apply `NVFP4` `global_scale` via `CUBLAS_POINTER_MODE_DEVICE` to eliminate intermediate FP32 buffers and `M=128` sub-tiling. The cuBLAS epilogue can round `global_scale` products one ulp differently from an FP32 multiply followed by the cast (documented in the schema). The activation epilogue is sliced into 65,535-row grid-y chunks. | | `onnxruntime/contrib_ops/cuda/bert/group_query_attention.cc`<br>`onnxruntime/contrib_ops/cuda/bert/group_query_attention_impl.cu`<br>`onnxruntime/contrib_ops/cuda/bert/attention_data.h`<br>`onnxruntime/contrib_ops/cuda/bert/cudnn_fmha/cudnn_flash_attention.h`<br>`onnxruntime/contrib_ops/cuda/bert/cudnn_fmha/cudnn_flash_attention.cc` | Prefer `FlashAttention` for small `R <= 16` shared-KV verify steps in auto mode (so `R=8` verify uses `FlashAttention` without `ORT_ENABLE_CUDNN_FLASH_ATTENTION=0` while large-`R` prefill uses Hopper `cuDNN SDPA`), keep `cudnn_sdpa` during CUDA graph capture when its plan is already cached (`cudnn_sdpa::has_cached_graph`) and fall back only when a non-capturable plan build would be needed, and route cuDNN SDPA workspace allocations through arena-pooled `GetScratchBuffer` (eliminating synchronous `cudaMalloc`/`cudaFree` stalls). | ### Follow-ups from microsoft#34147 | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cu` | Use zero coefficients when `sum + routing_epsilon == 0` (previously Inf/NaN rows), matching the documented normalization contract. | | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cc`<br>`onnxruntime/core/graph/contrib_ops/contrib_defs.cc` | Reject non-default `routed_scaling_factor`/`routing_epsilon` without `router_weights`; legacy softmax routing silently ignored them. | | `onnxruntime/core/graph/contrib_ops/bert_defs.cc`<br>`docs/ContribOperators.md` | Document the row-count-dependent `MoEGroupRouter` projection reduction order (8-row split-K kernel, cuBLAS for `rows > 16`). | ### Unit Tests | File | Change | |------|--------| | `onnxruntime/test/contrib_ops/fused_state_space_op_test.cc` | Update schema validation test for `slots >= 1`. | | `onnxruntime/test/contrib_ops/cuda_kernels/fused_state_space_test.cc` | Add `MambaNormRowsBeyondGridYLimit` (65,537 rows) and `MambaPrefillInPlaceStatesMatchSeparateOutputs` (aliased vs separate state buffers for `R = 20/37`, including `state_row = 0`). | | `onnxruntime/test/contrib_ops/moe_packed_routing_test.cc` | Add legacy-routing rejection cases for `routed_scaling_factor` and `routing_epsilon`. | | `onnxruntime/test/contrib_ops/matmul_packed_precision_test.cc` | Allow one ulp for BF16 NVFP4 `global_scale` prefill rows (decode and FP32 prefill stay bit-exact). | | `onnxruntime/test/python/transformers/test_gqa.py` | Add `test_cuda_graph_capture_reuses_warm_cudnn_sdpa_plan`. | | `onnxruntime/test/python/transformers/test_gguf_packed_ops.py` | Add `test_mamba2_verify_prefill_right_aligned_snapshots` (`R = 37, slots = 8` vs sequential `R = 8` chunks) and `m = 65` Tensor Core GEMM coverage in `test_moe_group_router_r8_and_fused_bias`; add `test_mamba2_verify_prefill_rows_beyond_grid_y_limit`, `test_moe_packed_routing_epsilon_zero_denominator` and `test_matmul_packed_activation_rows_beyond_grid_y_limit`. | ## Motivation and Context ### Prefill Throughput on NVIDIA H200 (`sm_90`, CUDA 13.0, `NVIDIA-Nemotron-3.5-Bolt-74B-A9B`) | Prompt `L` | `llama.cpp` Default (`ub=512`) | `llama.cpp` Peak (`ub=4096`) | Prior ORT (`R=8`, `k=7`) | **This PR (`k=0`)** | **This PR (`k=7`)** | **Speedup (`k=7`) vs `llama.cpp` Peak / vs Prior ORT** | |---:|---:|---:|---:|---:|---:|---:| | **128** | 85.96 ms (`1,489.0 tok/s`) | 86.20 ms (`1,484.9 tok/s`) | 175.28 ms (`730.3 tok/s`) | **52.80 ms (`2,424.3 tok/s`)** | **55.66 ms (`2,299.6 tok/s`)** | **`1.55×`** / **`3.15×`** | | **512** | 166.26 ms (`3,079.5 tok/s`) | 166.86 ms (`3,068.4 tok/s`) | 700.32 ms (`731.1 tok/s`) | **83.82 ms (`6,108.6 tok/s`)** | **87.66 ms (`5,840.7 tok/s`)** | **`1.90×`** / **`7.99×`** | | **1024** | 332.31 ms (`3,081.4 tok/s`) | 271.43 ms (`3,772.6 tok/s`) | 1,391.56 ms (`735.9 tok/s`) | **120.56 ms (`8,493.6 tok/s`)** | **125.06 ms (`8,187.8 tok/s`)** | **`2.17×`** / **`11.13×`** | | **2048** | 671.54 ms (`3,049.7 tok/s`) | 493.12 ms (`4,153.2 tok/s`) | 2,810.28 ms (`728.8 tok/s`) | **196.39 ms (`10,428.5 tok/s`)** | **201.08 ms (`10,184.8 tok/s`)** | **`2.45×`** / **`13.98×`** | | **3968** | 1,320.22 ms (`3,005.6 tok/s`) | 915.41 ms (`4,334.7 tok/s`) | 5,462.94 ms (`726.3 tok/s`) | **342.46 ms (`11,586.9 tok/s`)** | **347.34 ms (`11,423.8 tok/s`)** | **`2.64×`** / **`15.73×`** | ### `nsys` Target Kernel-Family Breakdown at `L = 2048` (`192.76 ms` vs `483.81 ms` in `llama.cpp` = **`2.51×` faster**) - **IQ2_XXS MoE Routed GEMMs** (`40` layers): `98.51 ms` vs `137.61 ms` (**`1.40×` faster**) - **Mamba2 Conv / SSM Scan / Norm** (`40` layers): `46.42 ms` vs `47.82 ms` (**`1.03×` faster**; conv `2.43 ms` vs `3.42 ms`) - **NVFP4 Mamba2 Projections** (`40` layers): `25.69 ms` vs `123.21 ms` (**`4.80×` faster**) - **IQ2_XXS MoE Shared GEMMs** (`40` layers): `11.63 ms` vs `35.56 ms` (**`3.06×` faster**) - **Q8_0 Projections + `M=1` LM Head**: `4.64 ms` vs `13.22 ms` (**`2.85×` faster**) - **MoE Routing / Tile Build / Finalize**: `2.39 ms` vs `77.68 ms` (**`32.45×` faster**) - **GQA Attention + RoPE / KV**: `1.08 ms` vs `2.76 ms` (**`2.55×` faster**, `0` synchronous `cudaMalloc`/`cudaFree` calls) ## Testing - Ran `onnxruntime/test/python/transformers/test_gguf_packed_ops.py` on H200 (`sm_90`), including `test_mamba2_verify_prefill_right_aligned_snapshots`, `test_moe_group_router_r8_and_fused_bias`, `test_moe_packed_routed_scaling_factor_and_routing_epsilon`, and `test_matmul_packed_activation_types`. - Verified end-to-end MMLU-Pro 200 accuracy (`119 / 200 = 59.50%`, `97.0%` exact choice agreement with microsoft#34147) and speculative decode throughput (`k=7`) on `NVIDIA-Nemotron-3.5-Bolt-74B-A9B`. - Verified formatting and linting with `lintrunner -a` (`No lint issues`). - On A100 (`sm_80`, CUDA 13.0): `test_gguf_packed_ops.py` (50 tests), `onnxruntime_provider_test --gtest_filter='MoEPacked*:MatMulPacked*:MoEGroupRouter*:FusedStateSpace*:Mamba*:GatedDeltaNet*:GroupQueryAttention*:*MultiHeadAttention*:AttentionTest*:PagedAttention*'` (478 tests), `FusedStateSpaceKernelTest.*` via `CUDA_EP_Unittest.All`, and the `test_gqa.py` cuDNN/CUDA graph subset. The new boundary, zero-denominator and CUDA graph tests fail on dc36d14. ## Checklist - [x] Tests added/updated - [x] No breaking changes (backward-compatible with `1 <= R <= slots`) - [x] Documentation updated (`docs/ContribOperators.md`) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…crosoft#28457) ### Description This adds `onnxruntime.set_default_logger_callback(callback, severity=2)`. Python applications can use it to send ONNX Runtime log messages to any Python callable, such as the standard `logging` module, instead of stderr. ```python def set_default_logger_callback( callback: Callable[[int, str, str, str, str], None] | None, severity: int = 2, # 0=Verbose, 1=Info, 2=Warning, 3=Error, 4=Fatal ) -> None: ... ``` **Example:** ```python import logging import onnxruntime as ort def ort_log_to_python(severity: int, category: str, logid: str, code_location: str, message: str) -> None: level = [logging.DEBUG, logging.INFO, logging.WARNING, logging.ERROR, logging.CRITICAL][severity] logging.log(level, "[ORT:%s] %s", category, message) ort.set_default_logger_callback(ort_log_to_python, severity=0) # capture all sess = ort.InferenceSession("model.onnx") ort.set_default_logger_callback(None) # restore the platform default (stderr/OutputDebugString) ``` **Behavior:** - The callback receives messages from the environment's logging manager, including messages from existing sessions. - Installing a callback sets the default logger severity. Session and run log severities still apply, and existing sessions keep their configured severity. Passing `None` restores the platform sink and leaves the severity unchanged. - The callback runs synchronously, with the GIL held, on the thread that emits the message. - Messages from threads that Python does not know about, such as ORT thread-pool workers and EP-internal threads, go to the platform sink. Waiting for the GIL on those threads can deadlock session teardown or abort the process at interpreter exit. - If the callback logs, those nested messages go to the platform sink, as do messages whose callback raises. Python exceptions are reported through `sys.unraisablehook`. - If log text contains bytes that are not valid UTF-8, those bytes are backslash-escaped. The message still reaches the callback. - The API requires an ORT environment created by the Python module. Otherwise it raises an error. #### Summary of Changes | File | Change | |------|--------| | `onnxruntime/core/session/ort_env.h`, `ort_env.cc` | Add an optional `sink_factory` function pointer to `OrtEnv::LoggingManagerConstructionInfo`. When set, it creates the environment's default sink and takes precedence over `logging_function`. It is a plain function pointer rather than `std::function`, so the struct stays trivially destructible and minimal builds do not grow. | | `onnxruntime/python/onnxruntime_pybind_state.cc` | Add `PythonCallbackSink`, a sink whose callback can be replaced and that falls back to the platform sink. Add the `set_default_logger_callback` binding. | | `onnxruntime/python/onnxruntime_pybind_module.cc`, `onnxruntime_pybind_module_functions.h` | Install the sink through `sink_factory` when the Python module creates the `OrtEnv`. | | `orttraining/orttraining/python/orttraining_python_module.cc` | Make the same change in the training module. | | `onnxruntime/__init__.py`, `docs/python/api_summary.rst` | Export and document the new API. | | `onnxruntime/test/python/onnxruntime_test_python.py` | Add tests for the new callback (listed under Testing). | #### Testing Run the tests from a build with `--enable_pybind`: ``` python -m pytest onnxruntime/test/python/onnxruntime_test_python.py -k "logger_callback or logging_callback" ``` The tests cover: - delivery of records and their types - reset with `None` - invalid arguments - a raising callback, observed through `sys.unraisablehook` - a callback that replaces itself, and re-entrancy from a finalizer - replacing the callback while sessions log on two Python threads - log text that is not valid UTF-8 - on Linux, a subprocess check that ORT pool threads logging during session teardown and interpreter exit neither deadlock nor abort ### Motivation and Context Fixes microsoft#28437. The Python binding creates the ORT environment without a logging function. As a result, Python users could only redirect ORT logs by manipulating file descriptors. The C API already supports a custom logger through `CreateEnvWithCustomLogger`, and this change gives the Python API the same capability. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…#34117) ### Description - Extend CUDA `VarlenCausalConvWithState` with dynamic grouped two-tap convolutions, including per-token coefficient deltas, boundary masking, and schema validation. - Add opt-in grouped BF16 INT3 MoE prefill, with a compatible SIMT path and guarded tensor-core acceleration using canonical recomputation for rounding-sensitive outputs. - Extend H128/GQA6 XQA support with BF16 INT8/INT4 speculative verification and BF16 INT4 single-token decoding. - Optimize INT8 KV prefill through sliding-window-aware gathering, vectorized dequantization, and revised FlashAttention query tiling. - Add BF16 shadow caching for global INT8 KV prefill. Retain a dequantized prefix across prefill chunks and convert only newly appended tokens, avoiding repeated full-prefix gathering. The optional `prefill_cache_control` input supplies an epoch, capacity, and final-chunk flag to explicitly govern reuse. - Add shadow-cache lifecycle safeguards: epoch invalidation, CUDA-stream synchronization, failure handling, and final-chunk buffer release. Missing controls, CUDA graph capture, multi-request batches, and unsupported configurations retain stateless behavior. - Expand correctness, dispatch, fallback, and cache-lifecycle coverage; update operator documentation and CUDA build dependencies. ### Motivation and Context There are existing bottlenecks in INT3 expert computation and repeated quantized-KV gathering, alongside gaps in BF16 quantized-cache speculative attention. These changes provide the CUDA primitives needed to reduce prefill overhead and support lower-memory speculative decoding. Optimizations are restricted to eligible configurations, with grouped INT3 prefill remaining opt-in. Prefix reuse requires explicit caller-managed lifetimes and trades temporary BF16 workspace for reduced dequantization work. The companion GenAI changes supply the necessary graph lowering and lifecycle controls. Model weights are not requantized; quantized-KV accuracy still requires application-specific qualification. ### Benchmarks Notes: 1. H200, batch 1, 512 output tokens, one warmup and five measured requests. The below benchmarks compare BF16 KV cache and INT8 KV cache versions of the RCs. TTFT and decode are medians; peak GPU memory includes weights and reserved pools, in decimal GB. INT8 KV cache versions use the optimized tile64 + shadow path; draft KV remains BF16 throughout. 2. Memory usage is not a fair comparison between ORT GenAI and llama.cpp because ORT GenAI tries to capture as much GPU memory as it can for future conversations. Memory usage is simply noted here for recording purposes. **INT3 RC** | Context | BF16 TTFT (s) | INT8 TTFT (s) | BF16 Decode (tok/s) | INT8 Decode (tok/s) | BF16 Peak (GB) | INT8 Peak (GB) | |---|---:|---:|---:|---:|---:|---:| | 8K* | 1.313 | 1.325 | 446.396 | **480.024** | 105.429 | 89.860 | | 16K* | 2.684 | 2.706 | 227.500 | **469.024** | 105.429 | 89.860 | | 32K* | 5.600 | 5.594 | 195.555 | **453.775** | 105.429 | 90.933 | | 64K | 12.157 | 12.041 | 302.338 | **428.886** | 105.429 | 92.141 | | 128K | 28.123 | 27.547 | 211.742 | **379.787** | 105.429 | 95.363 | **INT2 RC** | Context | BF16 TTFT (s) | INT8 TTFT (s) | BF16 Decode (tok/s) | INT8 Decode (tok/s) | BF16 Peak (GB) | INT8 Peak (GB) | |---|---:|---:|---:|---:|---:|---:| | 8K | 1.236 | 1.225 | 280.496 | **420.968** | 84.122 | 68.622 | | 16K | 2.490 | 2.502 | 405.926 | **441.012** | 84.122 | 68.620 | | 32K | 5.184 | 5.174 | 357.272 | **429.573** | 84.122 | 69.763 | | 64K | 11.297 | 11.127 | 275.144 | **391.317** | 84.122 | 70.834 | | 128K | 26.356 | 25.621 | 197.655 | **346.274** | 84.122 | 74.056 | **llama.cpp (2-bit)** | Prompt | TTFT (s) | Decode (tok/s) | Peak GPU (GB) | |---|---:|---:|---:| | 8K | 1.387 | 319.586 | 60.888 | | 16K | 2.839 | 330.673 | 61.250 | | 32K | 5.919 | 324.807 | 61.924 | | 64K | 13.035 | 317.357 | 63.268 | | 128K | 31.620 | 290.775 | 65.952 |

…ft#33957) ## What changes Return a clear error when an operation has no assigned execution backend, or names a backend that is not registered, instead of failing while preparing data copies. This prevents cases where ops could not be assigned an ep from crashing the browser GPU process when ORT WebGPU is used as the backend for webnn. The error identifies the affected node. Unsupported operations return `NOT_IMPLEMENTED`; invalid backend names return `INVALID_ARGUMENT`. Validation happens before copy nodes are inserted. Valid CPU-assigned nodes keep their existing behavior. `ModifyGraph` now returns `Status` and reports graph modification through a `bool&` output parameter. Assignment checks run in its existing initial node scan, before collecting each node's copy requirements; there is no separate validation traversal in `ApplyImpl`. `ApplyImpl` propagates failures with `ORT_RETURN_IF_ERROR`. These are internal helper changes, with no public API/ABI change. The existing non-CPU-provider guard and subgraph recursion remain unchanged. ## Native tests **3 tests passed, with no skips, after the refactor.** They cover an unassigned node, an unavailable named backend, an unchanged valid CPU-assigned node, and successful CPU-to-GPU-to-CPU copy insertion with `modified == true`. Invalid assignments are tested after an earlier valid GPU-assigned node has been processed; node assignments and connections remain unchanged on failure. Regression-sensitivity checks used disposable faulty source copies with the unchanged tests. Dropping the modification flag, omitting a GPU-to-CPU copy, and changing an earlier valid node's assignment each failed the intended assertion. The production source was then rebuilt and all three tests rerun successfully. From an existing Release build configured with WebGPU: ```powershell cmake --build build\latest-webgpu\Release --target onnxruntime_test_all --parallel 8 Push-Location build\latest-webgpu\Release .\onnxruntime_test_all.exe --gtest_filter=TransformerTest.MemcpyTransformer* Pop-Location ``` The tests use a small test backend and do not depend on GPU hardware. --------- Co-authored-by: Sushanth Rajasankar <Sushanth.Rajasankar@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Adds signed int8 and int64 support for Abs, Neg, Relu, and Sign to improve WebGPU EP coverage on WebNN WPT tests. The measured WebNN WPT pass rate increased from **96.6% (1,691/1,751)** to **97.1% (1,701/1,751)** in combined validation. Eight newly passing integer unary cases correspond to this change; the other two gains are not attributed to this PR. The selection excludes float16 input/output cases. Int64 unary kernels respect the `enableInt64` opt-in; int8 support is unaffected. Validation: all 12 native unary tests pass with no skips, covering enabled/disabled int64 configurations with CPU fallback disabled. WPT measurements predate the subsequent kernel refactor; WPT was not rerun afterward. --------- Co-authored-by: Sushanth Rajasankar <Sushanth.Rajasankar@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…16C/Batch4/40-core cold decode (microsoft#34114) ## Description Stacks on microsoft#33976 (and integrates microsoft#34066 and microsoft#34067). Optimizes `com.microsoft::MoEPacked` CPU/GPU hybrid expert offload and `com.microsoft::MatMulPacked` multi-token (`M=2..8`) MTP verification for block-packed GGUF MoE models (`Qwen3.8-Flash-Next` `UD-IQ3_XXS`) based on Nsight Systems (`nsys --cuda-graph-trace=node`) timeline profiling on NVIDIA H200: 1. **Prefill**: Eliminates cross-stream H2D Copy-Engine head-of-line blocking by double-buffering per-wave route/tile metadata (`d_wave_sorted_routes[2]`, `d_wave_tiles[2]`, `d_wave_num_tiles[2]`) on `wave_staging.copy_stream` and decoupling host `std::memcpy` (`h2d_done[b]`) from device compute (`compute_done[b]`), increasing true 2-stream compute/H2D overlap **3.77x** (`6.6% -> 24.9%`, `154.4 ms -> 478.8 ms` at `8K`) and cutting `2K–128K` TTFT by **`-14.7% to -25.9%`** (`+17.3% to +35.0%` prefill tok/s, saving `8.67s–9.72s` at `128K`). 2. **Decode**: Replaces 129 software `MLFloat16::ToFloat()` scale conversions in `gguf_vec_dot_cpu.h` AVX2 loops with hardware F16C `_mm_cvtph_ps`, adds 4-route (`Batch-4`) fused AVX2 INT8 (`vpmaddubsw`) dot-product kernels (`AccumIq2S32Q8Batch4Avx2`, `AccumIq3Xxs32Q8Batch4Avx2`, `AccumIq4Nl32Q8Batch4Avx2`, `VecDotGgufRow4xQ8Batch4Avx2`, `VecDotGgufGateUpRow2xQ8Batch4Cpu`), scales `MoEPackedParallelPool` to `40` physical cores, and dispatches `MatMulPackedBf16TcWarpKSplit<4, 1>` for `N <= 2560`, cutting `MoEPackedDoorbellWaitKernel` by **`-49.6%`** (`53.40 -> 26.91 ms/cycle` exact; **`-81.9%`** to `9.66 ms/cycle` with `tail_sub=0.02`) and accelerating decode throughput by **`+60.1% to +79.4%`** (exact routing) and **`+102.8% to +152.5%` (`2.03x–2.53x`)** with tail expert substitution. ## Summary of Changes ### 1. `MoEPacked` 3-Stage Prefill Pipeline & Cold-Expert Decode (`onnxruntime/contrib_ops/cuda/moe/moe_packed.cc`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cc` | Double-buffer `d_wave_sorted_routes[2]`, `d_wave_tiles[2]`, and `d_wave_num_tiles[2]` in `MoEPackedSharedPrefillWaveStaging` and copy them on `wave_staging.copy_stream` prior to recording `h2d_done[b]`; synchronize host `std::memcpy` only on `h2d_done[b]` while `copy_stream` waits on `compute_done[b]` on the device timeline; parallelize host expert staging across `w_cnt * 3` tasks; raise `MoEPackedParallelPool::ResolveThreads()` default cap from `24` to `40` cores; integrate `Batch-4` + `Batch-2` Q8 cold-expert execution in `RunColdDecodeStep()` (opt-in via `ep.cuda.moe_packed_cpu_q8_dot=1` or `ORT_MOE_PACKED_CPU_Q8_DOT=1` so existing FP16/BF16 models retain full FP32 cold-expert dot precision by default). | ### 2. Hardware F16C & 4-Route (`Batch-4`) AVX2 INT8 GGUF Kernels (`onnxruntime/contrib_ops/cpu/quantization/gguf_vec_dot_cpu.h`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cpu/quantization/gguf_vec_dot_cpu.h` | Add `Fp16ToFloatF16c` (`_mm_cvtph_ps`) under `ORT_GGUF_CPU_HAS_AVX2_TARGET` and replace 129 software `Fp16BitsToFloatCpu` calls inside AVX2 loops; add 4-route fused AVX2 INT8 dot-product kernels (`AccumIq2S32Q8Batch4Avx2`, `AccumIq3Xxs32Q8Batch4Avx2`, `AccumIq4Nl32Q8Batch4Avx2`, `VecDotGgufRow4xQ8Batch4Avx2`, `VecDotGgufRow4xQ8Batch4Cpu`, `VecDotGgufGateUpRow2xQ8Batch4Cpu`) plus 2-route `Batch-2` kernels from microsoft#34067. | ### 3. `MoEPacked` Decode FC2 & `MatMulPacked` Multi-Token GEMV (`moe_packed.cu`, `matmul_packed.cu`, `gguf_quant_cuda.cuh`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cu` | Integrate microsoft#34067 4-warp (`64` output channels/CTA) `MoEPackedFc2Bf16TcKernel<1>` decode geometry with shared-memory `s_lut` GGUF lookup tables and direct `__ldg(fc1_row)` activation loads. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cu` | Integrate microsoft#34066 `(NR=2, KW=2)` multi-token GEMV and dispatch `MatMulPackedBf16TcWarpKSplit<4, 1>` when `N <= 2560` (`M in 2..8`) to double CTA grid occupancy (`1,280 -> 2,560` CTAs across 132 H200 SMs, `-45.4%` per-call latency on 72 GDN projection calls/cycle). | | `onnxruntime/contrib_ops/cuda/quantization/gguf_quant_cuda.cuh` | Integrate microsoft#34066 uniform `uint16x2` `__ldg` loads in `ReadU32Dev` to eliminate 50% warp divergence on 2-byte-aligned GGUF blocks (`Q8_0`, `Q6_K`, `IQ4_NL`, `IQ4_XS`). | --- ## Performance & Accuracy Evaluation (`Qwen3.8-Flash-Next` `UD-IQ3_XXS` + `dflash2` `S=8` MTP on H200) ### 1. Setting 1: `30 GiB VRAM + 50 GiB RAM` (`device_expert_bytes = 12.50 GiB`, `6,463` GPU Resident Slots, Peak VRAM = `29.82 GiB`, Peak `VmRSS` = `35.44–36.09 GiB`) | Configuration | `2K` TTFT / Decode | `8K` TTFT / Decode | `32K` TTFT / Decode | `128K` TTFT / Decode | MMLU-Pro Smoke (`20`) Acc | MMLU Tok-Wtd / Med TPS | CPU `ms/step` (`us/exp`) | |---|---|---|---|---|---:|---|---:| | **1. Baseline (`origin/main` `6cc3c09506`)** | `1.061s` (`1,931 t/s`) / `39.18 t/s` | `2.433s` (`3,367 t/s`) / `38.87 t/s` | `9.156s` (`3,579 t/s`) / `39.28 t/s` | `38.790s` (`3,379 t/s`) / `38.31 t/s` | `100.0%` (`20/20`) | `41.36` / `43.84 t/s` | `61.58 ms` (`43.72 us`) | | **2. Merged PRs (microsoft#33975+microsoft#33976+microsoft#34066+microsoft#34067, FP32 CPU)** | `1.052s` (`1,947 t/s`) / `43.48 t/s` | `2.432s` (`3,369 t/s`) / `44.26 t/s` | `8.961s` (`3,657 t/s`) / `43.09 t/s` | `38.357s` (`3,417 t/s`) / `38.55 t/s` | `100.0%` (`20/20`) | `42.57` / `48.49 t/s` | `61.65 ms` (`43.52 us`) | | **3. Merged PRs (microsoft#33975+microsoft#33976+microsoft#34066+microsoft#34067, Q8 Batch-2 CPU)** | `1.035s` (`1,978 t/s`) / `50.17 t/s` | `2.395s` (`3,420 t/s`) / `50.37 t/s` | `8.942s` (`3,665 t/s`) / `49.13 t/s` | `38.318s` (`3,421 t/s`) / `47.88 t/s` | `100.0%` (`20/20`) | `49.50` / `57.69 t/s` | `47.85 ms` (`33.35 us`) | | **4. This PR (Exact Routing)** | **`0.786s` (`2,607 t/s`, `+35.0%`) / `62.73 t/s` (`+60.1%`)** | **`1.807s` (`4,534 t/s`, `+34.6%`) / `69.30 t/s` (`+78.3%`)** | **`6.813s` (`4,809 t/s`, `+34.4%`) / `69.45 t/s` (`+76.8%`)** | **`30.124s` (`4,351 t/s`, `+28.8%`) / `61.87 t/s` (`+61.5%`)** | `95.0%` (`19/20`) | **`67.12` / `79.32 t/s` (`+80.9%`)** | **`32.80 ms` (`22.89 us`)** | | **5. This PR + TailSub (`thr=0.02, min_rank=2`)** | **`0.853s` (`2,400 t/s`, `+24.3%`) / `97.11 t/s` (`+147.9%`)** | **`1.969s` (`4,160 t/s`, `+23.5%`) / `88.70 t/s` (`+128.2%`)** | **`7.495s` (`4,372 t/s`, `+22.2%`) / `81.71 t/s` (`+108.0%`)** | **`32.830s` (`3,993 t/s`, `+18.2%`) / `77.68 t/s` (`+102.8%`)** | **`100.0%` (`20/20`)** | **`97.76` / `108.06 t/s` (`+146.5%`)** | **`14.13 ms` (`25.95 us`)** | ### 2. Setting 2: `22 GiB VRAM + 24 GiB RAM` (`device_expert_bytes = 5.60 GiB`, `2,843` GPU Resident Slots, Peak VRAM = `21.71 GiB`, Peak `RssAnon` = `0.91–0.95 GiB`) | Configuration | `2K` TTFT / Decode | `8K` TTFT / Decode | `32K` TTFT / Decode | `128K` TTFT / Decode | MMLU-Pro Smoke (`20`) Acc | MMLU Tok-Wtd / Med TPS | CPU `ms/step` (`us/exp`) | |---|---|---|---|---|---:|---|---:| | **1. Baseline (`origin/main` `6cc3c09506`)** | `1.216s` (`1,684 t/s`) / `29.39 t/s` | `4.125s` (`1,986 t/s`) / `31.86 t/s` | `15.739s` (`2,082 t/s`) / `31.76 t/s` | `66.065s` (`1,984 t/s`) / `30.32 t/s` | `95.0%` (`19/20`) | `34.00` / `39.65 t/s` | `74.14 ms` (`42.04 us`) | | **2. Merged PRs (microsoft#33975+microsoft#33976+microsoft#34066+microsoft#34067, FP32 CPU)** | `1.197s` (`1,712 t/s`) / `34.54 t/s` | `4.042s` (`2,027 t/s`) / `35.58 t/s` | `15.762s` (`2,079 t/s`) / `33.70 t/s` | `66.006s` (`1,986 t/s`) / `34.26 t/s` | `95.0%` (`19/20`) | `37.03` / `42.90 t/s` | `73.66 ms` (`42.24 us`) | | **3. Merged PRs (microsoft#33975+microsoft#33976+microsoft#34066+microsoft#34067, Q8 Batch-2 CPU)** | `1.230s` (`1,665 t/s`) / `37.75 t/s` | `4.028s` (`2,034 t/s`) / `41.31 t/s` | `15.600s` (`2,101 t/s`) / `36.93 t/s` | `66.186s` (`1,980 t/s`) / `41.27 t/s` | `95.0%` (`19/20`) | `43.87` / `49.14 t/s` | `57.36 ms` (`32.47 us`) | | **4. This PR (Exact Routing)** | **`0.980s` (`2,090 t/s`, `+24.1%`) / `52.61 t/s` (`+79.0%`)** | **`3.345s` (`2,449 t/s`, `+23.3%`) / `52.73 t/s` (`+65.5%`)** | **`13.205s` (`2,482 t/s`, `+19.2%`) / `56.98 t/s` (`+79.4%`)** | **`56.343s` (`2,326 t/s`, `+17.3%`) / `50.16 t/s` (`+65.4%`)** | `95.0%` (`19/20`) | **`62.68` / `70.61 t/s` (`+78.1%`)** | **`38.21 ms` (`21.71 us`)** | | **5. This PR + TailSub (`thr=0.02, min_rank=2`)** | **`0.962s` (`2,130 t/s`, `+26.5%`) / `74.21 t/s` (`+152.5%`)** | **`3.321s` (`2,467 t/s`, `+24.2%`) / `74.47 t/s` (`+133.7%`)** | **`13.191s` (`2,484 t/s`, `+19.3%`) / `68.41 t/s` (`+115.4%`)** | **`56.890s` (`2,304 t/s`, `+16.1%`) / `72.87 t/s` (`+140.3%`)** | `95.0%` (`19/20`) | **`89.00` / `100.57 t/s` (`+153.6%`)** | **`17.55 ms` (`23.32 us`)** | ### 3. Nsight Systems (`nsys --cuda-graph-trace=node`) Before-vs-After Verification (`30g_50g`, `8192` Context) | Nsight Systems Metric (`8192` Context, `17` `S=8` MTP Cycles) | Pre-Optimization | Post-Opt (Exact) | Post-Opt (`tail-sub=0.02`) | Delta (Post Exact / Tail-Sub vs. Pre) | |---|---:|---:|---:|---| | **8K Prefill GPU Window (`Chunk 0 + Chunk 1`)** | `2,354.70 ms` | **`1,922.92 ms`** | **`1,919.95 ms`** | **`-431.78 ms` (`-18.3%`) / `-434.75 ms` (`-18.5%`)** | | **True 2-Stream Compute/H2D Overlap (`stream` & `copy_stream` both active)** | `154.43 ms` (`6.6%`) | **`478.83 ms` (`24.9%`)** | **`480.18 ms` (`25.0%`)** | **`+324.40 ms` (`3.10x` time / `3.77x` share)** | | **Both GPU Streams Idle in Prefill** | `293.64 ms` (`12.5%`) | **`84.03 ms` (`4.4%`)** | **`83.47 ms` (`4.3%`)** | **`-209.61 ms` (`-71.4%` idle bubble)** | | **48-Layer Target Verify CUDA Graph Duration (`ms/pass`)** | `84.11 ms/pass` | **`53.36 ms/pass`** | **`38.91 ms/pass`** | **`-30.75 ms` (`-36.6%`) / `-45.20 ms` (`-53.7%`)** | | **`MoEPackedDoorbellWaitKernel` (`48` calls/cycle)** | `53.40 ms/cyc` (`60.6%`) | **`26.91 ms/cyc` (`47.2%`)** | **`9.66 ms/cyc` (`22.7%`)** | **`-49.6%` (`1.98x` shorter) / `-81.9%` (`5.53x` shorter)** | | **`MatMulPackedBf16TcWarpKSplit<4, 1>` (`72` GDN calls/cycle, `N <= 2560`)** | `9.44 ms/cyc` (`131.1 us/call` on `<4,2>`) | **`5.15 ms/cyc` (`71.6 us/call`)** | **`5.17 ms/cyc` (`71.8 us/call`)** | **`-4.29 ms/cyc` (`-45.4%` per-call speedup)** | ## Testing - `pytest onnxruntime/test/python/transformers/test_gguf_packed_ops.py`: passes CPU/CUDA/hybrid offload numerical parity and tail-substitution tests on H200 (`sm_90`). - End-to-end `2K / 8K / 32K / 128K` context benchmarks and 20-item MMLU-Pro smoke evaluation under both `30 GiB VRAM + 50 GiB RAM` and `22 GiB VRAM + 24 GiB RAM` settings. ## Checklist - [x] Tests verified (`test_gguf_packed_ops.py` and MMLU-Pro smoke evaluation) - [x] No breaking changes - [x] Documentation updated --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description  Replace the Android setup action's `coreutils`/`ninja-build` installation and the separate React Native Ninja installation with offline Ninja version and path checks. ### Motivation and Context  The Linux runner image already provides coreutils and provisions Ninja independently of the Debian package. Installing `ninja-build` can overwrite the image-provided executable and unnecessarily depends on Ubuntu mirrors; [DNS failures blocked Android jobs](https://github.com/microsoft/onnxruntime/actions/runs/43042030824/job/130596064083) at this step. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description  Move DirectML compilation in `.github/workflows/windows_dml.yml` to `ort-github-win-16-amd-scus-ss`, with a dependent test job on the existing `ort-github-win-6-a10-scus-ss` GPU pool. Preserve the existing `Windows GPU DML CI Pipeline` check name on the GPU test job. - CPU: explicit `--update --build --skip_tests`, retaining the original Visual Studio generator, WCOS/BinSkim settings, vcpkg cache, Python 3.12, C# configuration, documentation/delegate validation, NuGet setup, and DirectML wheel packaging. - GPU: install the produced wheel, then run explicit `build.py --test` with all original feature/test flags. No configure, compilation, wheel rebuild, or NuGet restore target runs here. Fix `setup_dml_build` so test-only invocations do not implicitly call `cmake --build --target RESTORE_PACKAGES`. - Hand off the runtime output tree, nested CTest registrations, and any ONNX backend fixtures present in the original build. Relocate producer build/source/Python/CMake paths, retain runtime DLLs/Python packages/test fixtures (including hidden files), and exclude native objects/libraries/debug symbols and dependency build trees. Fail on missing required test binaries, Python/DML runtime files, registrations, or a missing/ambiguous wheel. - Run build-stage regressions on the CPU runner. Coverage includes build-only wheel packaging/C# configuration and test-only execution without build commands. The separate `test_dml_test_artifacts.py` module was removed as requested; the production `dml_test_artifacts.py` helper remains. C# scope: the existing workflow uses `--build_csharp` but not `--build_nuget`. In the current build driver/CMake code this generates C# build properties; it does not invoke NuGet packaging or `dotnet test`. That behavior and the delegate-attribute validation are preserved, not replaced with a different C# test surface. **Observed performance:** the [successful split run](https://github.com/microsoft/onnxruntime/actions/runs/42848703495) reduced GPU-runner occupancy by approximately **45%** and end-to-end execution time by approximately **24%**, compared with the average of three successful main runs ([42838643765](https://github.com/microsoft/onnxruntime/actions/runs/42838643765), [42861308845](https://github.com/microsoft/onnxruntime/actions/runs/42861308845), [42863264161](https://github.com/microsoft/onnxruntime/actions/runs/42863264161)). | Metric | Before: combined GPU job (three-run average) | After: split jobs | Difference | | --- | --- | --- | --- | | GPU-runner time | 1h 48m 53s | 1h 00m 21s | 48m 32s less (45%) | | End-to-end execution | 1h 48m 53s | 1h 22m 29s | 26m 24s faster (24%) | | CPU-runner time | None | 21m 22s | Build moved off GPU | The CPU build step took **16m 18s**; the GPU test step took **56m 31s**. Artifact staging, upload, download, and relocation took **1m 19s** combined. End-to-end is measured from the first job start through the final job completion, including the split's 46-second inter-job gap but excluding initial queue time. This is an observed comparison, not a controlled benchmark; there is one successful split sample so far. **Workflow scan:** examined all 47 workflow YAML files and all 13 local composite actions. The only active combined DirectML build/test workflow is `windows_dml.yml`. The DirectML-capable runtime build in `windows_gpu_doc_gen.yml` already runs on the CPU pool with a dependent GPU documentation-generation job. Other DirectML strings in CPU workflows are package cleanup/reference strings, not DML-enabled build variants. Other candidates found, deliberately left unchanged: | Path | Job / remaining GPU-side work | | --- | --- | | `.github/workflows/windows_webgpu.yml` | `webgpu_static_plugin_build_x64_RelWithDebInfo`: combined native build/test on GPU (separate work already planned; this PR does not depend on that branch). | | `.github/workflows/windows_webgpu.yml` | `webgpu_external_dawn_build_x64_RelWithDebInfo`: configure/build, external-Dawn host tests, and package export on GPU. | | `.github/workflows/windows_webgpu.yml` | `webgpu_dawn_package_build_x64_RelWithDebInfo`: package-backed ORT/plugin builds, negative configure checks, and GPU host tests. | | `.github/workflows/windows-web-ci-workflow.yml` | `build_onnxruntime_web`: WASM is already downloaded from a separate build, but NPM packaging and browser tests still share the GPU runner. This is a JS packaging candidate, not another unsplit native WASM build. | Windows CUDA/no-cuDNN/CUDA-plugin/TensorRT and the corresponding Linux workflows already use CPU-build/GPU-test handoffs. Windows WebGPU built-in/shared-plugin variants also already have split jobs. No unrelated workflow/composite action, compiler-cache adoption, or generator migration is included. **Validation** (isolated local Python environment; commands run from the repository root): - `.\.venv\Scripts\python.exe -m unittest tools.ci_build.test_build_args -q` — passed, 24 tests after removing the artifact test module. - `.\.venv\Scripts\python.exe -m ruff check tools\ci_build\build.py tools\ci_build\test_build_args.py tools\ci_build\github\windows\dml_test_artifacts.py` — passed. - `.\.venv\Scripts\python.exe -m ruff format --check tools\ci_build\build.py tools\ci_build\test_build_args.py tools\ci_build\github\windows\dml_test_artifacts.py` — passed, three files already formatted. - `git diff --check` and `git diff --cached --check` — passed before commit. - Parsed the revised DirectML workflow with PyYAML and asserted the production staging/prepare helper calls remain, the deleted test module is absent, and the remaining build-stage regression invocation is intact — passed. - Initial split validation also parsed all 60 workflow/composite YAML files and all six DirectML PowerShell steps. Assertions comparing original build options (only stage/skip flags differ), triggers/concurrency/environment/check name, runner/dependency wiring, artifact identity/guards, wheel ordering, and unchanged unrelated workflows/actions passed. - Locally reproduced the telemetry launcher failure with a missing producer CMake path, then confirmed both nested CTest launchers executed after relocation, including forward-slash and escaped Windows paths. Missing producer/consumer CMake checks failed explicitly as intended. - [DirectML CI run 42848703495](https://github.com/microsoft/onnxruntime/actions/runs/42848703495) at commit `0b1507622bc29e9cd1e757e17a65a8653b08bc96` — passed both the CPU build and GPU test jobs, verifying the artifact handoff and fixing the four telemetry shutdown tests that previously retained the CPU runner's absolute CMake path. Full native DirectML compilation, production wheel installation, GPU test execution, and the actual GitHub artifact transfer were not run locally; they were verified by the successful CI run above. No workflow was manually dispatched and no live cache was modified. ### Motivation and Context  Use the established 16-core Windows CPU pool for compilation instead of occupying the six-core GPU pool. Reserve the GPU runner for backend execution while retaining the existing DirectML CI coverage and validation settings. The observed successful run saves approximately 49 GPU-runner minutes per invocation and completes approximately 26 minutes sooner than the recent main-run average. This is an independent PR based on main, separate from the WebGPU/static-plugin and macOS caching work. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…microsoft#34199) ### Description  Move the existing per-type `HasCudaEnvironment(is_float16 ? 530 : 0)` check to the start of `RunModelWithRandomInput`, before constructing the random generator, allocating token offsets, or generating inputs, weights, and biases. Unsupported environments now report `GTEST_SKIP()` before any `OpTester` is constructed. The shared guard covers `TestWithRandomData`, `TestWithRandomDataWithRBP`, and `TestWithRandomDataLargeSeq`. Supported GPU FP32/FP16/RBP/large-sequence dimensions, seeds, inputs, tolerances, reference comparisons, and minimum CUDA architectures are unchanged. The guard remains per invocation: lack of FP16 support does not prevent the mixed-type cases from exercising supported FP32 inputs. No CI, build optimization, cache, or image changes. **Validation — head `b8eda7e79f5c9787ae23670802a2179fc0a2418f`:** - **Linux ARM64 Debug CI: passed.** The [draft PR job](https://github.com/microsoft/onnxruntime/actions/runs/43470505744/job/132040829102) built and tested successfully on `ort-github-linux-16-arm64-eus2-ss`, the same CPU runner pool as the reference. Its provider-test command was `/onnxruntime_src/build/Debug/Debug/onnxruntime_provider_test "--gtest_output=xml:/onnxruntime_src/build/Debug/Debug/onnxruntime_provider_test.Debug.results.xml"` — passed; **7,024 tests selected**, including all three random-data cases, each explicitly reported skipped. This is not a zero-match pass. - `git diff --check` — passed. - `.\.venv\Scripts\clang-format.exe --dry-run --Werror onnxruntime\test\contrib_ops\packed_attention_op_test.cc` — passed with clang-format 20.1.8. - **Windows x64 CPU-only Debug:** in an isolated `.venv`, after `call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvars64.bat"`, ran `.venv\Scripts\python.exe tools\ci_build\build.py --build_dir build\n --config Debug --update --build --parallel 12 --skip_submodule_sync --cmake_generator Ninja --target onnxruntime_provider_test --disable_memleak_checker` — passed. Build log confirms `packed_attention_op_test.cc.obj` recompiled and the executable relinked. Cache confirms Debug and CUDA OFF. - From `build\n\Debug`: `onnxruntime_provider_test.exe --gtest_filter=PackedAttentionTest.TestWithRandomData* --gtest_list_tests` — exit 0; listed exactly the three existing random-data tests. - From the same directory: `onnxruntime_provider_test.exe --gtest_filter=PackedAttentionTest.TestWithRandomData* --gtest_output=xml:C:\Users\erscor\.copilot\session-state\893b4970-454c-4884-9b9c-2c936909fed7\files\packed-attention-cpu-debug.xml` — exit 0; **3 selected, 3 skipped, 0 failures**. Local x64 XML reports 10 ms / 4 ms / 0 ms for random data / RBP / large sequence; gtest reports 16 ms suite time, 19 ms including global setup/teardown. These are separate x64 results, not the ARM64 performance comparison below. Local build notes: Visual Studio generation first hit Windows' 260-character MSBuild path limit, so validation used Ninja in this worktree. The initial default Debug executable compiled successfully and selected/skipped all three tests, but its experimental Windows memory-leak checker exited -1 at shutdown; even `--gtest_list_tests` alone reported shutdown leaks. The clean local validation above uses the repository-supported `--disable_memleak_checker` option locally only; Debug optimization settings are unchanged. No repository build settings were edited. CUDA hardware/runtime is unavailable locally. GPU execution was not performed locally; the original architecture expression and supported GPU path were inspected for unchanged behavior. CPU-only runs verify early gating, not GPU numerical correctness. ### Motivation and Context  The CUDA-only random-data tests previously generated large CPU fixtures before checking whether CUDA was available. **Measured Linux ARM64 Debug comparison**, using the completed [reference job](https://github.com/microsoft/onnxruntime/actions/runs/43458960142/job/132000464739) and successful [after-change draft PR job](https://github.com/microsoft/onnxruntime/actions/runs/43470505744/job/132040829102): | Test (`PackedAttentionTest.` prefix) | Before | After | Saved | | --- | ---: | ---: | ---: | | `TestWithRandomData` | 266,632 ms | 4 ms (skipped) | 266,628 ms | | `TestWithRandomDataWithRBP` | 133,346 ms | 3 ms (skipped) | 133,343 ms | | **Two reported slow tests** | **399,978 ms** | **7 ms** | **399,971 ms (~6m40s)** | | `TestWithRandomDataLargeSeq` | 10,943 ms | 0 ms reported (skipped) | 10,943 ms | | **All three helper callers** | **410,921 ms** | **7 ms** | **410,914 ms (~6m51s)** | The reported slow pair's runtime fell by **99.998%**, measured on the same ARM64 Debug CPU runner pool. Millisecond timings are rounded by gtest; 0 ms does not mean literally zero elapsed time. The whole provider-test CTest entry went from **2,112.21 s to 1,706.48 s**, an observed reduction of **405.73 s (19.21%)**. Treat that as descriptive, not a controlled whole-suite attribution: the reference selected **7,019** tests and the PR merge selected **7,024**, and these are different CI runs/commits. The pair alone accounted for 18.94% of the reference provider-test time; all three accounted for 19.45%. The supplied cold-run reference pair is consistent (266,686 + 133,286 ms); the supplied ARM64 Release reference pair was 2,333 + 1,170 ms. The additional ARM64 Release measurements are reported below; no GPU performance timings are claimed. Independent change based on `main`; does not modify or stack on microsoft#34163 or the ccache experiment. #### Cross-platform impact Combined gtest time for all three existing random-data tests (including the large-sequence case): | Configuration | Before | After | Saved | | --- | ---: | ---: | ---: | | Linux ARM64 Debug | 410,921 ms | 7 ms | 410,914 ms | | Linux x64 Debug + ASan | 56,309 ms | 6 ms | 56,303 ms | | macOS ARM64 Debug | 14,912 ms | 0 ms reported | 14,912 ms | | Windows x64 Release (two executions) | 6,625 ms | 2 ms | 6,623 ms | | Linux ARM64 Release | 3,598 ms | 1 ms | 3,597 ms | | Linux x64 Release | 2,442 ms | 0 ms reported | 2,442 ms | | Linux x64 Release / Python 3.14 | 2,451 ms | 0 ms reported | 2,451 ms | **Total measured aggregate test execution removed: 497,242 ms (8m17.242s)** across these jobs, not an equivalent decrease in elapsed PR/CI time. Every after-change occurrence explicitly reports skipped; all compared jobs succeeded. Windows runs the provider tests twice (regular and example-plugin executions), and both executions are included. Gtest rounds timings to milliseconds. Sources: Linux [before](https://github.com/microsoft/onnxruntime/actions/runs/43458960142) / [after](https://github.com/microsoft/onnxruntime/actions/runs/43470505744); macOS [before](https://github.com/microsoft/onnxruntime/actions/runs/43467069151/job/132028944327) / [after](https://github.com/microsoft/onnxruntime/actions/runs/43470505763/job/132040830226); Windows [before](https://github.com/microsoft/onnxruntime/actions/runs/43467069133/job/132028943543) / [after](https://github.com/microsoft/onnxruntime/actions/runs/43470505843/job/132040829952). Linux compares the earlier experiment reference with this PR; macOS/Windows use successful main runs at this branch's base commit, 6080d7c. These are separate runs, not controlled repeated benchmarks. #### Total workflow elapsed time: estimate versus observation **No net end-to-end workflow reduction was observed in these particular run pairs.** Using the first job's `started_at` through the final job's `completed_at` (excluding initial runner queue delay, but including later job scheduling gaps), GitHub reports: | Workflow | Before execution span | After execution span | Observed change | | --- | ---: | ---: | ---: | | Linux CI | 46m31s | 51m06s | 4m35s longer | | MacOS CI Pipeline | 86m42s | 94m44s | 8m02s longer | | Windows x64 Release | 33m54s | 34m04s | 10s longer | Including initial queue time, creation-to-final-job-completion was Linux 49m48s -> 54m33s, macOS 86m46s -> 94m49s, and Windows 43m14s -> 66m28s. In particular, the Windows after-run waited about 32m24s before its job started, versus 9m20s before; this is not test runtime. The fixture savings are directly visible in individual test timings, but build/configuration/cache state, scheduling, other tests, and different commits confound whole-workflow comparisons. The ARM64 Debug provider-test entry itself decreased by 405.73s (19.21%), while the entire ARM64 Debug job increased from 43m38s to 50m22s; other work more than offset the test savings in this pair. **All-else-equal critical-path estimate:** ARM64 Debug was the last-finishing Linux job in both runs. Removing its measured 410.914s of fixture work would shorten the reference Linux workflow execution span from 46m31s to about 39m40s: **approximately 6m51s, or 14.7%**, provided other timings remain fixed and no other job becomes the bottleneck. This is a modeled potential wall-clock saving, not an observed workflow speedup. Savings in parallel Linux jobs do not add to that critical-path reduction. The measured macOS CPU Debug job was not the final-finishing job (the ARM64 iPhone simulator was), so its 14.912s saving alone does not shorten that workflow's critical path. For the single-job Windows workflow, the isolated all-else-equal saving is about 6.623s (0.33% of baseline execution time). Builds that do not execute these tests receive no benefit. Supported CUDA paths still generate and exercise the original fixtures; no GPU computation speedup is claimed. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…oft#33921) ### Description - Add the missing `FLOAT16` case to tensor attribute serialization in `OrtGraphToProto`. - Add a `ConstantOfShape` regression test covering `raw_data` and `int32_data`, verifying the serialized payload and inference with optimizations disabled. - All 19 EP graph tests pass. ### Motivation and Context Valid FP16 tensor attributes were rejected with `Unexpected ONNXTensorElementDataType with value 10`. Fixes microsoft#33919.

…soft#34163) ### Description  Repair portable, rolling Linux ccache snapshots without adopting sccache or changing compilers, PCH, build/security flags, or test coverage. Retain shared snapshots where cold compilation takes at least five minutes; leave shorter builds on their existing local ccache without GitHub snapshot restore/save overhead. - Archive workspace-relative `.cache/ccache` instead of the reusable workflow's `~/.cache/cache` typo or home-relative paths that depend on runner workspace depth. - Share a setup action with a fresh `ccache-linux-v2` namespace. Scope compatibility to build variant/dependency identity, architecture, and actual local Docker image ID or host compiler/ccache versions. Publish unique run-ID/attempt snapshots with only a compatibility-scoped restore prefix; no broad fallback. - Publish snapshots only after successful main builds. PRs can restore but never publish; compatible restores no longer suppress publication of updated entries. Historical archives are not deleted; normal GitHub eviction applies. - Explicitly route the pinned Docker build action to `/onnxruntime_src/.cache/ccache` through its existing environment prefix while preserving caller Python setup. Limit cross-UID permissions and umask changes to compiler-cache contents. - Apply the measured cost cutoff in `655926d690efbc406ba4ce5c14a4fe381315cbb6`: remove both snapshot setup/restore and save from six minimal variants whose **cold build phase was under five minutes**: custom ops, type reduction, globally allowed types, minimal without optional features, extended minimal without optional features, and Android NNAPI. Keep snapshots for exceptions-disabled minimal, extended minimal, regular without optional features, and the unchanged reusable Linux callers. Existing local ccache tools/flags, VCPKG policy, tests, and other settings remain unchanged. The cutoff is a static policy based on cold measurements, not warm timings or a runtime threshold. - Add namespace, publication-policy and eligibility regression tests, portable archive/cross-UID smoke tests, and cache-contract documentation. Declare the tests' PyYAML dependency in `requirements-dev.txt`. **Real GitHub cache-service cold/warm validation — November 13** Ran both existing workflows twice through normal branch-specific push triggers, preserving all six Linux CPU variants, all nine previously snapshot-enabled minimal variants, the full-build/test-data prerequisite, original compiler/PCH/security options and tests. The experiment was isolated on `eserscor-2026/11/13-branch-ccache-experiment`, not added to this PR. | Workflow | Cold run | Warm run | Created-to-last-job completion, cold → warm | | --- | --- | --- | ---: | | Linux CI | [passed](https://github.com/microsoft/onnxruntime/actions/runs/43448889588) | [passed](https://github.com/microsoft/onnxruntime/actions/runs/43458960142) | 59:02 → 49:48 | | Linux CPU Minimal Build E2E | [passed](https://github.com/microsoft/onnxruntime/actions/runs/43448889207) | [passed](https://github.com/microsoft/onnxruntime/actions/runs/43458959669) | 27:17 → 28:47 | **All 32 jobs passed.** All 15 cold compiler-cache restores missed. Every cold snapshot save was confirmed in job logs and the GitHub caches API before the second push. All 15 warm restores matched their corresponding cold publication exactly, with no toolchain/Docker-identity mismatch; all 15 warm saves were also confirmed. This exercised real pinned `actions/cache` restore/save, **not artifact-based compiler-cache transport**. The ordinary `test_data` artifact remained unchanged. Cold commit: `602870c80f0baac6e9c6bd1178dfd91a9f163dc6`. Warm commit: `bf9737d205e44946522a7e7b3d085f4cd98b3bcb`, an intentionally empty commit. Both have identical source tree `571f761d828096a281d5be672f1834ea4c90dd42`. An experiment-only cache prefix kept these snapshots separate from production main keys. No cache deletion, custom maintenance, shared-image modification, `actions:write` grant, manual rerun or GPU workflow expansion was used. Times below are **cold → warm**, in minutes:seconds. Build means the standalone Actions compilation-step wall time; starred native composite/Android rows use `build.py` CMake-build command to test-start/build-complete timestamps within a combined step. Whole job includes setup, cache operations, tests and post steps, but excludes queue/dependency waiting. Hit rates are among cacheable compiler calls, not all tool invocations. Every measured variant restored cold=miss/warm=hit, saved both times, and passed both times. The measurements precede the final six-variant snapshot-policy reduction. | Variant | Build time | Whole-job time | Warm hits / misses | Hit rate | Final shared-snapshot policy | | --- | ---: | ---: | ---: | ---: | --- | | x64 Debug ASan | 17:05 → 2:37 | 41:46 → 27:22 | 1,364 / 1 | 99.93% | Keep | | x64 Release | 16:11 → 0:55 | 33:07 → 16:44 | 1,411 / 2 | 99.86% | Keep | | x64 Release Python 3.14 | 16:03 → 1:03 | 32:18 → 20:05 | 1,411 / 2 | 99.86% | Keep | | x64 MinSizeRel, exceptions disabled | 7:20 → 1:06 | 9:30 → 4:05 | 638 / 1 | 99.84% | Keep | | ARM64 Debug | 17:54 → 3:54 | 57:14 → 43:38 | 1,570 / 1 | 99.94% | Keep | | ARM64 Release | 17:47 → 0:53 | 28:22 → 11:37 | 1,615 / 2 | 99.88% | Keep | | Minimal, exceptions disabled | 5:24 → 1:17 | 8:31 → 5:03 | 622 / 1 | 99.84% | Keep | | Minimal, custom ops* | 4:17.2 → 0:44.6 | 6:46 → 3:27 | 694 / 1 | 99.86% | Omit: cold build <5m | | Minimal, type reduction* | 4:17.0 → 0:34.4 | 7:16 → 3:15 | 686 / 1 | 99.85% | Omit: cold build <5m | | Minimal, globally allowed types* | 4:27.3 → 0:55.0 | 7:14 → 4:18 | 686 / 1 | 99.85% | Omit: cold build <5m | | Extended minimal | 5:58 → 1:36 | 9:17 → 5:01 | 734 / 1 | 99.86% | Keep | | Regular, no optional features | 8:46 → 2:40 | 11:55 → 6:41 | 705 / 2 | 99.72% | Keep | | Minimal, no optional features | 4:33 → 0:40 | 7:12 → 3:53 | 449 / 1 | 99.78% | Omit: cold build <5m | | Extended minimal, no optional features | 4:37 → 0:42 | 6:49 → 4:02 | 470 / 1 | 99.79% | Omit: cold build <5m | | Android NNAPI* | 3:57.3 → 0:13.2 | 6:36 → 3:09 | 1,100 / 1 | 99.91% | Omit: cold build <5m | The unchanged full-ORT prerequisite has no compiler snapshot: its build step took 14:15 → 15:02 and whole job 16:11 → 17:27 ([cold](https://github.com/microsoft/onnxruntime/actions/runs/43448889207/job/131965536019), [warm](https://github.com/microsoft/onnxruntime/actions/runs/43458959669/job/132000462302)). The minimal workflow was **90 seconds slower overall** despite faster snapshot-enabled compilation: that prerequisite was 76 seconds slower, and warm dependent jobs waited 3:53–4:43 for runners versus 8–12 seconds cold. Compiler acceleration is not end-to-end workflow acceleration. This is one sequential pair, not a statistical benchmark; runner, Docker/setup, existing VCPKG-cache and test variability remain. **Why ARM64 Debug remains the long tail** [Warm ARM64 Debug job](https://github.com/microsoft/onnxruntime/actions/runs/43458960142/job/132000464739): compilation fell from **17:54 to 3:54**, but tests remained **35:18 cold / 35:17 warm**, about **81% of the warm job**. Ccache does not accelerate test execution; cache restore/save took roughly ten seconds each. CTest already uses `--parallel 16`, but the **7,019 provider cases execute inside one unsharded `onnxruntime_provider_test` process**. That entry took 2,112.85 seconds cold / 2,112.21 seconds warm; `onnxruntime_test_all` finished in 340.30 seconds warm. Eight provider cases account for 1,417 seconds, approximately 67% of its wall time. - `MatMulNBits` cases total about **14m13s**; `Float16_4b_Accuracy4` alone takes **4m48s** and includes a large 369×6144×2048 regression case. - Two `PackedAttentionTest` random-data cases take **6m40s combined**, despite this being CPU-only. [`RunModelWithRandomInput`](https://github.com/microsoft/onnxruntime/blob/0acfa78a809f0719280f913861a53692252bb8ad/onnxruntime/test/contrib_ops/packed_attention_op_test.cc#L798-L925) generates inputs, weights and biases **before** its `HasCudaEnvironment` check. The two loops make 120 calls and generate approximately 212 million Gaussian weight values even when CUDA is unavailable. The same two tests take about 3.5 seconds combined in the measured ARM64 Release job. - `ReduceSum_batch_by_seq_by_128` takes **3m24s** across 90 large-shape cases; `Gather_overflow_check` takes **1m49s** and deliberately creates an output just over 2 GiB plus reference/output copies. This is reproducible Debug test/fixture cost, not a cache miss or a demonstrated runner stall. Moving CUDA availability checks before unused fixture generation and sharding the provider-test executable are separate potential follow-ups; this cache PR does not skip tests or change Debug optimization/coverage. **Validation** Final eligibility policy at `655926d690efbc406ba4ce5c14a4fe381315cbb6`: - `wsl -d Ubuntu-24.04 --exec bash -c 'source /tmp/ort-ccache-ci-305f4061-venv/bin/activate && python3 -m unittest tools.ci_build.test_linux_ccache -v'` — **7 tests run: 6 passed, 1 optional Docker skip**. - Normalized YAML comparison against the original main baseline — **passed**: both workflows retain identical non-snapshot build/test settings, compiler/PCH/security flags, tool setup, VCPKG policy and coverage. - `actionlint -shellcheck= -pyflakes= .github/workflows/reusable_linux_build.yml .github/workflows/linux_minimal_build.yml` with actionlint 1.7.12, compared with the pre-cutoff baseline — **no new diagnostics**; the same 10 existing custom-runner-label diagnostics remain. Shellcheck and pyflakes disabled for this comparison. - `git diff --check` — **passed**. Final policy committed and pushed; normal PR CI is triggered, not claimed complete here. The four successful experiment runs above validate the repair before the final eligibility-only reduction, not this new head's CI. Earlier repair validation retained as supporting evidence: - `python -m pip install -r requirements-dev.txt` — **passed in clean Linux and Windows virtual environments**, including PyYAML. Both tooling CI jobs and Python format passed after the dependency fix: https://github.com/microsoft/onnxruntime/actions/runs/43354291711 . - `python -m pytest -c .github/workflows/tooling_tests.pytest.ini --rootdir=. -v` — **198 passed, 3 skipped, 206 subtests passed** on the native Linux source copy before the eligibility reduction. - `ORT_CCACHE_DOCKER_IMAGE=ort-ccache-smoke:305f4061 python3 -m unittest tools.ci_build.test_linux_ccache -v` — original **6 tests passed, no skips**, covering differing workspace depths, external-action/direct-Docker environments, UIDs 1001/1002 and host access to container-created entries. - `python -m ruff check tools/ci_build/test_linux_ccache.py` and `python -m ruff format --check tools/ci_build/test_linux_ccache.py` — **passed** on the earlier repair, including repository-pinned ruff 0.12.12 validation. - A prior [artifact-transfer Debug extended-minimal experiment](https://github.com/microsoft/onnxruntime/actions/runs/43443808349) achieved 605.728 seconds cold / 76.928 seconds warm, 1,402 direct hits / zero misses; its temporary workflow/helper were removed. The full-workflow cache-service measurements above supersede that narrower experiment as the principal evidence. Raw logs, exact per-job cache keys, compiler counters, configure/build/test/cache timing JSON and CSV were retained with the experiment report. The main-only publication policy is regression-tested; the experiment validates branch-scoped cache-service publication/consumption, not main-to-PR visibility after merge. No GPU performance claim is made, and unmeasured reusable callers retain their existing policy. ### Motivation and Context  The reusable workflow saves a path that does not contain its compiler cache, while minimal builds can restore immutable historical archives containing relative traversals to an obsolete runner home/workspace layout. Static exact-match keys also prevent populated snapshots from rolling forward. Fix these root causes while preserving build behavior, and avoid spending shared snapshot resources on short builds that do not determine the workflow's long tail. Independent fix from main, not stacked on draft microsoft#34109; that PR and its branch are unchanged. The app branch control only supported a dated kebab-case name with its automatic prefix (`eserscor-2026-11-12-linux-ccache-repair`), not the requested slash-separated branch format. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…ft#34127) ### Description  Adds a full (unreduced) native WebGPU EP Asyncify WebAssembly build, `ort-wasm-simd-threaded-full.asyncify.{wasm,mjs}`. ### Motivation and Context  JSEP migration work. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…32 GDN kernels (microsoft#34200) ## Summary This PR optimizes CUDA prefill and speculative-decode (`M = 8` verify) kernels across `MatMulPacked`, `MoEPacked`, `DynamicSparseAttention` (`fp16_kv`), `HyperConnectionLowRankMix`, `GatedDeltaNet` (`fp32` state), and `QMoE` helper launches for Qwen3.8-Flash-Next GGUF (`UD-IQ3_XXS`) and `QMoE` (`2b375_fp32`) models: 1. **`MatMulPacked` `Q6_K` Contiguous Half-Superblock (`kChunk = 128`) & Aligned SWAR Unpack** (`gguf_quant_cuda.cuh`, `matmul_packed.cu`): - Assigns contiguous 128-element `Q6_K` half-superblocks per warp (`LoadQ6KHalf128Stage` + `DecodeQ6KHalf128GroupBf16x2` SWAR unpack) in `MatMulPackedBf16TcWarpKSplitKernel` (`KW=10, MT=1` for `N<=1024`; `KW=8` for `K=6144, N<=2560`; `KW=4, MT=2` for `N>2560`) and enables `Q6_K` in `MatMulPackedBf16DequantFastKernel`. - Cuts `MatMulPacked` decode time from `19.60 -> 13.48 ms/cycle` (`-6.12 ms/cyc`; `N=512` `Q6_K` latency drops `5.5x` from `33.8 us -> ~6.1 us`). 2. **`MoEPacked` Contiguous `kChunk = 128` (`IQ3_XXS`, `IQ2_S`, `IQ4_XS`) FC1 Decode & Aligned 36B Pair (`IQ4_NL`) FC2 Decode** (`gguf_quant_cuda.cuh`, `moe_packed.cu`): - Hoists superblock scale/sign loads across 4 sub-blocks per warp in `MoEPackedFc1DecodeWarpKSplitKernel` with 4-way bank-replicated `iq3xxs_grid` (`s_lut[1024]`), and decodes 4-byte-aligned 36B (`2 x block_iq4_nl`) pairs (`DecodeIq4NlPair64AlignedBf16x2`) in `MoEPackedFc2Bf16TcKernel<1, kGgmlTypeIQ4_NL>` (`21.47 -> 19.76 ms/cyc`, reaching `295 GB/s` vs `301 GB/s` peak on RTX Spark). 3. **FP16-KV Tensor-Core `DynamicSparseAttentionFp16TcKernel<kSplitK, kCandTile>`** (`dynamic_sparse_attention_prefill.cuh`, `dynamic_sparse_attention_impl.cu`): - Replaces the scalar `GroupedDecodeDynamicSparseAttentionKernel` fallback on `fp16_kv` with a split-K `cp.async` Tensor-Core `m16n8k16` FP16-KV kernel (`kCandTile = 32` on 100 KB/SM GPUs, `P_hi + P_residual` 22-bit probability split) for decode and two-buffer `<false, 64>` for prefill. - Speeds up FP16-KV attention decode by **`5.4x` (`2.71 -> 0.50 ms/cyc`)** and 8K prefill by **`1.49x` (`486.0 -> 325.3 ms`)**. 4. **Single-Wave `m16n8k16` Tensor-Core `HcDownTc8Kernel` & `HcUpPreTc8Kernel` (`rows = 8`) + `NormVec8H2560Kernel`** (`hyper_connection_low_rank_mix_impl.cu`): - Replaces 2-wave scalar `LaunchFusedRows<T, 4, 2>` (`3.81 GB/cyc` weight reads) on `rows=8, branches=4, hidden=2560, rank=320` (BF16) with single-wave BF16 Tensor-Core kernels (`grid=(6,8,1)` and `grid=(80,1,1)`) with PDL weight prefetch (`1.27 GB/cyc` = **`3.0x` less weight DRAM traffic**, saving **`~3.5–4.5 ms/cyc`**), plus 320-thread `NormVec8H2560Kernel` for prefill. 5. **Fused Dual-A 3-Slice FP32-State `GatedDeltaNetChunkedKernel`** (`gated_delta_net_mma.cuh`, `gated_delta_net_impl.cu`): - Hoists `LoadFragA` across slices and fuses `(k @ S0, Qg @ S0, R epilogue)` (`SmemGemmFp32StateDualS0Bf16`) and `(P @ U, Kd^T @ U)` (`SmemGemmFp32StateDualUBf16`), reducing 3-slice FP32->BF16 passes per chunk from `5 -> 3` (`35 -> 12` barriers) while preserving exact 24-bit IEEE-754 precision (`427.0 -> 343.3 ms` at 8K prefill, `-19.6%`). 6. **Column-Matched CTA Widths in `QMoE` `expandInputRowsKernel` & `doActivationKernel`** (`moe_kernels.cu`, `moe_gemm_activation_kernels.cuh`): - Matches CTA block widths to column vector counts (`320` for `hidden_size=2560`; `64`/`96` for `inter_size=512/640`), eliminating 69%–75% idle tail threads (`-31 ms` at 8K prefill). --- ## End-to-End Performance & Accuracy (`2K–255K + 1K` on NVIDIA RTX Spark `SM 12.1`) ### 1. ORT `gguf_ud_iq3_xxs_q8_engram` (`DFlash2 K=7`) vs. Baseline & `llama.cpp + MTP` (`UD-IQ3_XXS + Q8_0`) | Context (`Prompt + Gen`) | ORT Baseline TTFT / Decode (`ms/cyc`) | **ORT Optimized TTFT / Decode (`ms/cyc`)** | `llama.cpp + MTP` TTFT / Decode | **ORT Opt vs. `llama.cpp + MTP`** | |---|---:|---:|---:|---:| | **`2K + 1K`** | `1.509 s` / `52.71 tok/s` (`67.62 ms`) | **`1.468 s` / `66.70 tok/s` (`52.53 ms`)** | `2.582 s` / `57.81 tok/s` | **`1.76x` prefill / `+15.4%` decode** | | **`8K + 1K`** | `5.468 s` / `46.42 tok/s` (`67.81 ms`) | **`5.231 s` / `66.49 tok/s` (`53.80 ms`)** | `9.616 s` / `54.22 tok/s` | **`1.84x` prefill / `+22.6%` decode** | | **`16K + 1K`** | `10.822 s` / `51.75 tok/s` (`67.24 ms`) | **`10.425 s` / `61.24 tok/s` (`54.06 ms`)** | `19.331 s` / `58.48 tok/s` | **`1.85x` prefill / `+4.7%` decode** | | **`32K + 1K`** | `21.663 s` / `48.44 tok/s` (`68.34 ms`) | **`20.929 s` / `63.54 tok/s` (`54.76 ms`)** | `40.822 s` / `54.69 tok/s` | **`1.95x` prefill / `+16.2%` decode** | | **`64K + 1K`** | `44.858 s` / `53.26 tok/s` (`70.09 ms`) | **`42.684 s` / `58.04 tok/s` (`57.23 ms`)** | `81.592 s` / `50.67 tok/s` | **`1.91x` prefill / `+14.5%` decode** | | **`128K + 1K`** | `94.444 s` / `42.02 tok/s` (`72.02 ms`) | **`91.288 s` / `59.12 tok/s` (`58.06 ms`)** | `172.053 s` / `45.89 tok/s` | **`1.88x` prefill / `+28.8%` decode** | | **`255K + 1K`** | `206.564 s` / `42.15 tok/s` (`81.71 ms`) | **`201.523 s` / `45.41 tok/s` (`67.65 ms`)** | `378.801 s` / `33.42 tok/s` | **`1.88x` prefill / `+35.9%` decode** | ### 2. ORT `rc6_2b375_fp32` (`DFlash2 K=7`, `fp16_kv + FP32 GDN`) Before vs. After | Context (`Prompt + Gen`) | Baseline TTFT (`Prefill tok/s`) | **Optimized TTFT (`Prefill tok/s`)** | Baseline Decode (`ms/cyc`) | **Optimized Decode (`ms/cyc`)** | |---|---:|---:|---:|---:| | **`2K + 1K`** | `1.232 s` (`1,661.9`) | **`1.204 s` (`1,700.5 tok/s`)** | `62.78 tok/s` (`54.68 ms`) | **`81.20 tok/s` (`47.91 ms`, `-12.4%`)** | | **`8K + 1K`** | `4.607 s` (`1,778.1`) | **`4.352 s` (`1,882.3 tok/s`)** | `63.91 tok/s` (`54.44 ms`) | **`76.47 tok/s` (`48.65 ms`, `-10.6%`)** | | **`32K + 1K`** | `18.642 s` (`1,757.8`) | **`17.589 s` (`1,862.9 tok/s`)** | `58.34 tok/s` (`56.02 ms`) | **`65.65 tok/s` (`50.11 ms`, `-10.5%`)** | | **`128K + 1K`** | `81.111 s` (`1,616.0`) | **`77.106 s` (`1,699.9 tok/s`)** | `48.82 tok/s` (`58.37 ms`) | **`51.38 tok/s` (`53.52 ms`, `-8.3%`)** | | **`255K + 1K`** | `179.976 s` (`1,450.9`) | **`172.099 s` (`1,517.3 tok/s`)** | `51.17 tok/s` (`68.93 ms`) | **`50.52 tok/s` (`63.08 ms`, `-8.5%`)** | ## Validation - Multi-arch CUDA compilation verified cleanly across `sm_75`, `sm_80`, `sm_86`, `sm_90`, and `sm_120f` with `-Werror all-warnings --Werror default-stream-launch`. - Speculative (`M = 8`) vs. `--force-reject` (`M = 1`) greedy equivalence verified `128/128` (`100%` bit-identical) on both `rc6_2b375_fp32` and `gguf_ud_iq3_xxs_q8_engram`. - MMLU-Pro shard `0/80` (`10` items): `9/10` (`90%`, matching baseline item-for-item) on `rc6_2b375_fp32` and `10/10` (`100%`) on `gguf_ud_iq3_xxs_q8_engram` with `0` truncations. - Added unit tests in `matmul_packed_precision_test.cc`, `gguf_quant_packed_test.cc`, and `dynamic_sparse_attention_op_test.cc`. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…icrosoft#34169) ### Description Each QMoE/MoE kernel instance owns a MoeGemmProfiler with a private (GemmId, M bucket) cache. A model whose MoE layers all have the same shape therefore profiles the identical problem once per layer for every new M bucket. Qwen3-30B-A3B has 48 such QMoE nodes. This change shares successful profiling results process-wide. The key holds everything a profiling run reads: - the runner type; - runner state that changes the candidate tactics (SM80 FP4 family) or the profiled GEMM2 epilogue (deterministic Hopper reduce); - the device; - every GemmProfilerBackend::init parameter; - the GemmId and the M bucket. This is the in-process part of the QMoE cross-node dedup planned in docs/contrib_ops/cuda/gemm_profiler_cache.md (section 8, phase 2). It follows the process-wide profile map that MatMulNBits shares through GemmPluginProfilerManager (matmul_nbits.cc), with two differences: - Coordination is per key. A short map lock hands out the key's entry, and only that entry's mutex is held while it is profiled. Concurrent first uses of one key profile it once, while other problems, M buckets and devices profile in parallel. - Only successful results are shared, as the design doc requires for the MatMulNBits cache. A run that returns no config (for example after a transient allocation failure) is kept by that profiler instance as before, and the next profiler retries instead of reusing the failure. The cache is intentionally leaked, like the MatMulNBits tactic cache registry, because a CUDA EP owned by another static can be destroyed after this library's statics. ORT_DISABLE_MOE_SHARED_TACTIC_CACHE=1 restores per-instance profiling and is documented in gemm_profiler_cache.md. MoeGemmProfiler::ProfilingRunCountForTesting() is a read-only counter that only the tests use. Tests (moe_gemm_profiler_shared_cache_test.cc): - IdenticalProblemsAreProfiledOncePerProcess and RunnerStateThatChangesProfilingIsPartOfTheKey: sharing works, and runner state that changes tactics keeps results apart. - FailedProfilingIsNotShared: a failed run is retried by the next profiler. - ConcurrentFirstUsesOfOneKeyProfileOnce: four threads that reach one key profile it once. - OtherKeysProfileWhileOneIsInFlight: a second problem profiles to completion while the first is held inside its workspace allocation. FailedProfilingIsNotShared and OtherKeysProfileWhileOneIsInFlight fail against an earlier, unpublished draft of this change that held one global lock while profiling and cached failed results. With sharing forced off, the four sharing tests fail. Measured on an A100 80GB, main (6596e0c) vs this change: - Qwen3-30B-A3B-Instruct-2507 through ORT GenAI, 32,640-token prompt, builds run in A-B-B-A order: - Cold time to first token is 82,465 and 83,704 ms on main vs 5,524 and 5,630 ms with this change. - Warm time to first token changes by -0.21% (95% CI -0.26% to -0.17%, 10 requests). - Generated tokens are identical, and peak memory is the same. - Decode measured -0.65% (95% CI -1.12% to -0.18%). Decode at 8 expanded rows runs the GEMV path, which uses no profiled tactic, so I read this as run-to-run drift rather than an effect of the change. - Multi-turn chat in one session, with prompts of 2,048, 8,192, 16,384 and 32,640 tokens. Time to first token for turns 1 and 2 is 798 and 2,347 ms instead of 22,922 and 79,956 ms. Turn 4 completes; on main it runs out of memory, because each node's profiling request leaves a 1 GiB arena region (fixed separately in microsoft#34167). - 48 same-shape QMoE nodes built from layer 0 of that model, first run with profiling at 2,048 rows: 22,916 ms to 811 ms (10 interleaved pairs), with identical outputs. ### Motivation and Context On Qwen3-30B-A3B-Instruct-2507 at a 32,640-token prompt on an A100, profiling dominated the first request: 82.5 s cold time to first token versus about 3.4 s warm. Most of that was the same two GEMMs being profiled 48 times. Limits: - Results are shared only inside one process, with no disk persistence. That is still future work in the design doc. - A key that never yields a config is profiled again by every node, as each node did before. - Tested on SM80 Linux. CI covers SM86, Windows and the plugin EP build. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Discovered when trying to use RandomNormal (opset 22) in a unit test: RandomNormal, RandomNormalLike, RandomUniform, RandomUniformLike and Multinomial were revised in ONNX opset 22, but their CPU kernels stayed registered at their original start version with no end version, which matches that version exactly. As a result these ops had no CPU kernel at opset 22+. Cap the old registrations at 21 and add new ones at 22, the same split EyeLike and the CUDA Random kernels already use. Type constraints are unchanged (float and double). ### Description Splits the CPU kernel registrations for `RandomNormal`, `RandomNormalLike`, `RandomUniform`, `RandomUniformLike` and `Multinomial` at opset 22: the existing registrations are capped with an end version of 21, and a new registration is added at 22 for each op. `docs/OperatorKernels.md` is regenerated to match, and each op gets an opset 22 test. Type constraints are unchanged — the opset 22 registrations declare the same `float`/`double` constraints as the earlier ones. ### Motivation and Context All five ops were revised in ONNX opset 22, but their CPU kernels stayed registered at their original start version (1, or 7 for `Multinomial`) with no end version. A kernel registered that way matches its start version *exactly*: `VerifyVersion` in `kernel_registry.cc` only takes the range branch when `kernel_end_version != INT_MAX`, on the reasoning that without an end version ORT cannot know whether a later op revision added something the kernel mishandles. Since the version that gets matched is the op's schema `since_version` — 1 for opsets 1-21 and 22 for opsets 22+ — these ops had no CPU kernel at all at opset 22 and later. Version matching happens before type matching, so this is not type-specific: an opset 22 `RandomNormal` with `dtype=float`, squarely inside the registered type constraints, still failed with `Could not find an implementation for RandomNormal(22)`. `EyeLike` was revised in the same opset and was updated at the time; it shows the intended pattern of capping the old registration and adding a new one. CUDA already has that split for the four Random ops, so this was CPU-side drift. This PR applies the same treatment to all five ops. Opset 22 also widened these schemas to the full set of float types by adding `bfloat16`. That is deliberately not implemented here — supporting a new dtype is a separate feature, whereas this change restores a working kernel lookup for the types already supported. #### On the tests `OpTester`'s default `Run()` pre-assigns every node to an EP and, when no EP has a matching kernel, logs a warning and returns without failing. The obvious version of these tests therefore *passed* against the unfixed code. The new tests run with an explicitly specified CPU EP, which leaves node assignment to the graph partitioner, so a missing registration surfaces as a session initialization failure. All five new tests were verified failing before this change and passing after. Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>

…ll rows (microsoft#34168) ### Description CutlassMoeFCRunner reserves an FP32 split-K partials buffer for the fused SwiGLU MoE GEMV (FP16 activations, INT4 weights, swiglu_fusion == 1, see MoeGemvSplitK2SwiGLUSupported) and sized it for every expanded row of the call. Its only reader is the split-K two-pass GEMV. That path runs only when is_moe_gemv_supported() accepts the problem, which needs expanded_num_rows <= moe_gemv::kMaxProfiledExpandedRows (80), and only with FP32 GEMV accumulation (ORT_MOE_GEMV_FP32_ACCUM=1). Above 80 expanded rows the grouped GEMM runs and the buffer is never read, so at prefill sizes the workspace carried a buffer that grows with the prompt and is never used. This change caps the partials at 80 expanded rows. Calls with up to 80 expanded rows get exactly the same size as before; larger calls keep a fixed 2 * 80 * 2 * inter_size * 4 bytes (0.94 MiB for inter_size 768). Tests: - QMoEWorkspaceTest.SplitKSwiGluPartialsDoNotScaleWithPrefillRows fails on main (per-token growth at GEMV sizes 176,640 B, needs at least 225,857 B) and passes with this change. - QMoEWorkspaceTest.SplitKSwiGluPartialsCoverGemvRows checks that GEMV-sized calls still get the full buffer. Validation on A100 80GB (SM80), using layer 0 of Qwen3-30B-A3B-Instruct-2507 (INT4, block 128, top-8, 128 experts) as a standalone QMoE model, main vs this branch: - Outputs are bitwise identical at 8, 80, 88 and 1024 expanded rows, both with ORT_MOE_GEMV_FP32_ACCUM=1 and with the default. - With ORT_MOE_GEMV_FP32_ACCUM=1, Nsight Systems shows moe_gemv_splitk_partials_kernel and moe_gemv_splitk_reduce_swiglu_kernel launching at 8 and 80 expanded rows, and no GEMV kernels at 88 and 1024. - compute-sanitizer memcheck reports 0 errors on those runs. memcheck alone would not catch an undersized buffer, because the partials share one workspace allocation with the other buffers. As a control, halving the cap changes the output at 80 expanded rows while memcheck still reports 0 errors, so the parity check is the meaningful one. - End to end with the full model through ORT GenAI at a 32,640-token prompt and ORT_MOE_GEMV_FP32_ACCUM=1 (one process per build): process peak 47,248 MiB on main vs 43,152 MiB with this change, identical generated tokens, decode 126.2-126.6 vs 126.3-126.6 tokens/s. The buffer is reserved the same way with the default accumulation. An earlier run of the same change on main at e4234e5 measured a peak of 47,254 vs 43,158 MiB in that mode (two processes per build, identical in each). ### Motivation and Context For Qwen3-30B-A3B at a 32,640-token prompt the partials buffer was 2 * 261,120 * 1,536 * 4 B = 3,059 MiB of workspace that was never read. Limits: - Only the FP16 + INT4 + fused-SwiGLU runner allocates this buffer, so other dtypes and activations are unaffected. - The small capped buffer is still reserved when FP32 accumulation is off. Dropping it entirely would tie the workspace layout to that environment variable, so this PR leaves it. - Tested on SM80 Linux only. CI covers SM86 and Windows. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description Replace the CUDA device-side `printf` used for INT3 QMoE compatibility counters with a host-side device-to-host counter readback and flushed stdout report. Update the test to require zero missed corrections in both FC1 and FC2 reports. ### Motivation and Context The Windows CUDA CI test selects the Tensor Core path but does not capture the device-side `missed=0` report, causing `test_cuda_int3_tensor_core_compatibility_shadow` to fail. Host-side reporting makes the verification output capturable by the test while preserving its checks. Validation: Python syntax check and `git diff --check` passed. The CUDA regression test was not run locally; the available T1000 does not support the BF16 Tensor Core path required by this test. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…#34167) ### Description MoeGemmProfiler::runProfiling took its one-off profiling workspace from the temp-space allocator with a stream-less Alloc(). With the default CUDA EP that allocator is the BFC arena, which keeps every region it grows. A stream-less request also cannot reuse a chunk that the compute stream has used since (BFCArena::FindChunkPtr). So each MoE node that profiles a new M bucket can leave another workspace-sized region in the arena for the rest of the session. This change reserves the scratch instead, the same way the fpA_intB GEMM profiler already does (IAllocator::MakeUniquePtr with use_reserve = true in contrib_ops/cuda/llm/gemm_profiler.h). BFCArena::Reserve takes the memory straight from the device allocator and Free returns it. Allocators without a Reserve override fall back to Alloc(), so they behave as before. One example is the plain CUDAAllocator the CUDA plugin EP uses for temp space. Because reserved memory goes back to the device on Free, the workspace guard now synchronizes the profiling stream before freeing. It is declared after the stream guard, so a private stream outlives it on every exit path, including exceptions from backend_.prepare() or a failing tactic. Test: MoeGemmProfilerTest.ProfilingScratchDoesNotGrowArena profiles three nodes at two M buckets through a real BFC arena. On main the arena grows by two extensions (3,145,728 bytes) with no reserves. With this change it grows by 0 bytes and uses Reserve. Measured on an A100 80GB, main (6596e0c) vs this change: - Multi-turn chat with Qwen3-30B-A3B-Instruct-2507 through ORT GenAI, with prompts of 2,048, 8,192, 16,384 and 32,640 tokens in one session. After turn 2 the model arena is 17,642 MiB instead of 66,282 MiB. Turn 4 completes instead of failing with out-of-memory, and the arena after turn 4 is 23,786 MiB. - Single 32,640-token prompt, builds run in A-B-B-A order, 10 warm requests per build. Time to first token changes by +0.04% (95% CI -0.01% to +0.09%) and decode by -0.04%. Generated tokens are identical, and process peak is 47,266 MiB on main vs 46,946 MiB with this change. - Profiling cost, on 48 same-shape QMoE nodes built from layer 0 of that model: the first run, which includes all 96 profiling runs, takes 96 ms longer (22,916 to 23,012 ms; 95% CI of the difference 64 to 129 ms; 10 interleaved pairs). Arena growth during that run drops from 1,610,612,736 to 536,870,912 bytes in every pair, and outputs are identical. ### Motivation and Context In a multi-turn chat with Qwen3-30B-A3B (48 QMoE layers) on an A100 80GB, each new prompt-length bucket made all 48 nodes profile. Each node's roughly 770 MiB profiling request at the largest bucket left a 1 GiB arena region behind, so the arena grew by about 48 GiB in one turn and the fourth turn ran out of memory. Limits: - Reserve() bypasses BFCArena binning and its gpu_mem_limit check. The memory is still counted in the arena's total while profiling and released right after, and the CUDA device memory budget, when configured, still applies through CUDAAllocator. - Each profiling run now does one cudaMalloc/cudaFree pair. Profiling happens once per node and M bucket, and never during CUDA graph capture. - Tested on SM80 Linux. CI covers SM86, Windows and the plugin EP build. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

adrastogi and others added 30 commits November 4, 2026 13:02

Revert "Fix/gate cuda ci registry auth" (microsoft#33894)

4177f55

Reverts microsoft#33376

rui-ren and others added 26 commits November 12, 2026 21:33

Merge remote-tracking branch 'origin/master' into sync_msft_15112026

700837f

ai-fw-intg requested review from Jaswanth51, ankitm3k, jatinwadhwa921 and vthaniel November 14, 2026 20:40

Navigation Menu

Repository navigation

Sync with Microsoft ONNX Runtime - 15112026 - #1395

Sync with Microsoft ONNX Runtime - 15112026#1395
ai-fw-intg wants to merge 140 commits into
ovep-developfrom
sync_msft_15112026

ai-fw-intg commented Nov 14, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

20 participants

Conversation

ai-fw-intg commented Nov 14, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

20 participants