Sync with Microsoft ONNX Runtime - 07082026

.agents/skills/ort-release-notes/SKILL.md

@@ -0,0 +1,151 @@
+---
+name: ort-release-notes
+description: Draft ONNX Runtime release notes using preset configurations for full ORT or scoped component releases. Use when generating highlights with PR links, compiling human contributor acknowledgments from compile_contributors.py output, and applying preset path filtering.
+argument-hint: "preset base_ref target_ref [version] [output_dir]"
+---
+
+# ONNX Runtime Release Notes
+
+Use this skill to produce a consistent release-note draft from commit history and contributor metadata.
+
+## When To Use
+
+Use this skill when you need to:
+
+- Draft release notes for a full ONNX Runtime release
+- Draft release notes for a scoped component (e.g., in-tree plugin EP) release
+- Select a release profile by preset name instead of manually supplying path/version files
+- Add PR links to highlight bullets
+- Build a human-only contributor acknowledgment list from contributor metadata
+
+## Required Inputs
+
+Collect these inputs from the user or infer from context:
+
+1. `preset`: release profile name (for example, `ort`, `webgpu-plugin-ep`, `cuda-plugin-ep`)
+2. `base_ref`: previous release tag
+3. `target_ref`: release commit/tag/branch tip
+
+Optional inputs:
+
+- `version` override
+- `output_dir` override
+
+## Presets
+
+Read preset definitions from [presets.json](./presets.json).
+
+The config defines shared output defaults:
+
+1. `outputDirPattern`
+2. `draftFileName`
+
+Each preset defines:
+
+1. `displayName`: reader-facing product or component name
+2. `versionFile`
+3. `pathsFile` (nullable): file of git pathspecs to filter to, one per line. Use `:(top)` to anchor an entry at repo root.
+
+Example presets:
+
+1. `ort` (full ONNX Runtime)
+2. `webgpu-plugin-ep` (scoped WebGPU Plugin EP)
+3. `cuda-plugin-ep` (scoped CUDA Plugin EP)
+
+### CUDA Plugin EP Scope
+
+The `cuda-plugin-ep` preset uses the pathspecs in `plugin-ep-cuda/paths.txt` to scope release-note changes.
+
+## Workflow
+
+1. Determine release mode.
+   - Select preset and load configuration from [presets.json](./presets.json).
+   - Use the preset's `displayName` whenever the release-note content names the product or component. The preset
+     key is internal and must not appear in published content.
+   - If preset has `pathsFile`, run in scoped mode. Otherwise run full mode.
+2. Resolve version, in this order:
+   1. explicit `version` input
+   2. value from preset `versionFile`
+3. Resolve output directory, referred to as `resolved_output_dir` after this step.
+   The output directory contains contributor artifacts and the release notes draft.
+   Resolve it in this order:
+   1. explicit `output_dir` input
+   2. shared `outputDirPattern` rendered with preset name and resolved version
+4. Gather metadata.
+   - If the output directory is missing or lacks contributor artifacts, generate them with
+     `tools/python/compile_contributors.py`.
+   - Generation can take a while because it scans commit history and fetches PR metadata.
+   - Use `--paths-file` only when preset has a `pathsFile`.
+   - If existing contributor artifacts are reused, verify `resolved_output_dir/logs.txt` matches base/target before trusting them.
+5. Read `resolved_output_dir/detail.csv` as the primary source for PR numbers, titles, authors, target commits, and cherry-pick mapping.
+   - Use `resolved_output_dir/logs.txt` for contributor summary context and base/target verification.
+   - Use `git log` only as a fallback sanity check when artifacts are present but incomplete or suspect.
+    - Check PRs with unexpectedly large author lists for rebased history that imported unrelated commits. Replace those
+       authors with the actual PR author or authors before building contributor acknowledgments; do not credit authors
+       solely because they authored an unrelated imported commit. For example, PR #28299, the rebased history contains unrelated commits and co-author metadata.
+6. Build highlight categories.
+   - Full ORT example categories: performance, model/operator support, execution providers, API/languages,
+     reliability/security, build/packaging/tooling, docs/dev workflow.
+   - Scoped mode: narrow categories to the component domain.
+7. Draft markdown.
+   - Write the release-note draft to `resolved_output_dir/<draftFileName>`.
+   - Contents:
+     - Intro sentence
+     - `## Highlights`
+     - Inline PR links on every highlight bullet
+     - `## Contributors`
+     - Optional scope note for scoped-component releases
+       - Use the preset's reader-facing `displayName`, not the internal preset key.
+       - Describe the scope in reader-facing terms, such as "commits affecting WebGPU Plugin EP code and
+         packaging."
+     - AI disclaimer if AI drafted
+   - Do not mention presets, `pathsFile`, configuration files, or other release-note-generation implementation
+     details in the release-note content.
+   - Do not refer to the release notes as a "draft" in their content. "Draft" is only an internal workflow
+     and file-naming concept.
+8. Build contributors section.
+   - Start from `detail.csv` output
+   - Include humans only
+   - Exclude bots/agents (for example: `github-actions[bot]`, `app/copilot-swe-agent`, `claude`)
+   - Sort alphabetically
+9. Validate draft quality.
+   - Every highlight bullet has at least one PR link
+   - PRs are traceable to metadata or git history
+   - Contributor list is human-only and alphabetical
+   - Scope is correct for full vs component release
+
+## PowerShell Command Patterns
+
+### compile_contributors.py
+
+Full ORT metadata:
+
+```powershell
+python .\tools\python\compile_contributors.py \
+  --base <previous_tag> \
+  --target <target_ref> \
+  --dir <resolved_output_dir>
+```
+
+Scoped metadata:
+
+```powershell
+python .\tools\python\compile_contributors.py \
+   --base <previous_tag> \
+   --target <target_ref> \
+   --dir <resolved_output_dir> \
+   --paths-file <paths_file>
+```
+
+## Style and Policy
+
+Default policy unless release owners override:
+
+1. Treat the range as changes since the previous release.
+2. Keep PR links inline with highlight claims.
+3. Keep contributor acknowledgments human-only and best effort.
+4. Include an AI disclaimer when highlights are AI drafted.
+5. Use GitHub Releases pages as preferred style references.
+   E.g., [ORT 1.28 release page](https://github.com/microsoft/onnxruntime/releases/tag/v1.28.0).
+6. Prefer preset-driven configuration over ad-hoc path/version arguments.
+7. Use a single `output_dir` for contributor artifacts, logs, and the release-note draft.

.agents/skills/ort-release-notes/presets.json

@@ -0,0 +1,22 @@
+{
+  "defaultPreset": "ort",
+  "outputDirPattern": "release-notes-output.{preset}.{version}",
+  "draftFileName": "release-notes-draft.md",
+  "presets": {
+    "ort": {
+      "displayName": "ONNX Runtime",
+      "versionFile": "VERSION_NUMBER",
+      "pathsFile": null
+    },
+    "webgpu-plugin-ep": {
+      "displayName": "ONNX Runtime WebGPU Plugin EP",
+      "versionFile": "plugin-ep-webgpu/VERSION_NUMBER",
+      "pathsFile": "plugin-ep-webgpu/paths.txt"
+    },
+    "cuda-plugin-ep": {
+      "displayName": "ONNX Runtime CUDA Plugin EP",
+      "versionFile": "plugin-ep-cuda/VERSION_NUMBER",
+      "pathsFile": "plugin-ep-cuda/paths.txt"
+    }
+  }
+}

.github/workflows/windows_cuda_no_cudnn.yml

@@ -121,6 +121,7 @@ jobs:
             --use_vcpkg `
             --use_vcpkg_ms_internal_asset_cache `
             --enable_cuda_profiling `
+            --cmake_extra_defines onnxruntime_USE_TRT_FUSED_ATTENTION=OFF `
             --cmake_extra_defines onnxruntime_QUICK_BUILD=ON `
             --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 `
             --cmake_extra_defines onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON

cmake/CMakeLists.txt

@@ -123,12 +123,18 @@ option(onnxruntime_USE_VSINPU "Build with VSINPU support" OFF)
 cmake_dependent_option(onnxruntime_USE_FLASH_ATTENTION "Build flash attention kernel for scaled dot product attention" ON "onnxruntime_USE_CUDA" OFF)
 option(onnxruntime_USE_LEAN_ATTENTION "Build lean attention kernel for scaled dot product attention" OFF)
 cmake_dependent_option(onnxruntime_USE_MEMORY_EFFICIENT_ATTENTION "Build memory efficient attention kernel for scaled dot product attention" ON "onnxruntime_USE_CUDA" OFF)
+# The TensorRT fused MHA kernels are prebuilt cubins for sm70-sm89 that mainly benefit BERT-style
+# encoder models. Turning this OFF drops ~14MB of embedded cubin data; attention falls back to
+# flash / memory efficient / cuDNN / unfused kernels.
+cmake_dependent_option(onnxruntime_USE_TRT_FUSED_ATTENTION "Build TensorRT fused multi-head attention cubin kernels" ON "onnxruntime_USE_CUDA" OFF)
 option(onnxruntime_USE_FP4_QMOE "Build CUDA QMoE FP4 kernels" OFF)
 option(onnxruntime_USE_FP8_QMOE "Build CUDA QMoE FP8 kernels" OFF)
 cmake_dependent_option(onnxruntime_USE_FPA_INTB_GEMM "Build FpA IntB gemm cuda kernels" ON "onnxruntime_USE_CUDA" OFF)
 option(onnxruntime_USE_INT4_KV_CACHE "Build cuda kernels for int4 kv cache" OFF)
 option(onnxruntime_USE_FP8_KV_CACHE "Build cuda kernels for fp8 kv cache" ON)
 option(onnxruntime_QUICK_BUILD "Speed up build by skipping some kernels for faster development" OFF)
+# Raises the minimum driver to the CUDA 12.4 level (Linux >= 550.54.14, Windows >= 551.61); always on for CUDA >= 13.0.
+cmake_dependent_option(onnxruntime_CUDA_FATBIN_COMPRESS_SIZE "Compress CUDA fatbins with -compress-mode=size" OFF "onnxruntime_USE_CUDA" OFF)
 
 option(onnxruntime_BUILD_FOR_NATIVE_MACHINE "Enable this option for turning on optimization specific to this machine" OFF)
 option(onnxruntime_USE_AVX "Use AVX instructions" OFF)
@@ -788,6 +794,7 @@ else()
   set(onnxruntime_USE_FLASH_ATTENTION OFF)
   set(onnxruntime_USE_LEAN_ATTENTION OFF)
   set(onnxruntime_USE_MEMORY_EFFICIENT_ATTENTION OFF)
+  set(onnxruntime_USE_TRT_FUSED_ATTENTION OFF)
 endif()
 
 if (onnxruntime_USE_CUDA)
@@ -809,6 +816,11 @@ if (onnxruntime_USE_CUDA)
       list(APPEND ORT_PROVIDER_FLAGS -DUSE_MEMORY_EFFICIENT_ATTENTION=1)
     endif()
 
+    if (onnxruntime_USE_TRT_FUSED_ATTENTION)
+      message( STATUS "Enable TensorRT fused multi-head attention for CUDA EP")
+      list(APPEND ORT_PROVIDER_FLAGS -DUSE_TRT_FUSED_ATTENTION=1)
+    endif()
+
     if (onnxruntime_USE_FPA_INTB_GEMM)
       message( STATUS "Enable FpA IntB Gemm for CUDA EP")
       list(APPEND ORT_PROVIDER_FLAGS -DUSE_FPA_INTB_GEMM=1)
@@ -1526,7 +1538,12 @@ if (onnxruntime_USE_CUDA)
     message(FATAL_ERROR "onnxruntime_USE_FP4_QMOE requires CUDA Toolkit version 12.8 or newer")
   endif()
 
-  if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "13.0")
+  if(onnxruntime_CUDA_FATBIN_COMPRESS_SIZE AND CMAKE_CUDA_COMPILER_VERSION VERSION_LESS "12.8")
+    message(FATAL_ERROR "onnxruntime_CUDA_FATBIN_COMPRESS_SIZE requires CUDA Toolkit version 12.8 or newer")
+  endif()
+
+  if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "13.0" OR onnxruntime_CUDA_FATBIN_COMPRESS_SIZE)
+    message(STATUS "Compressing CUDA fatbins with -compress-mode=size (requires a CUDA 12.4 or newer driver)")
     set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xfatbin=-compress-all -compress-mode=size")
   else()
     set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xfatbin=-compress-all")

cmake/external/cuda_configuration.cmake

@@ -164,6 +164,7 @@ macro(setup_cuda_architectures)
   unset(ORT_HAS_SM80_OR_LATER)
   unset(ORT_HAS_SM90_OR_LATER)
   unset(ORT_HAS_SM100_OR_LATER)
+  unset(ORT_HAS_SM120_OR_LATER)
   foreach(CUDA_ARCH IN LISTS CMAKE_CUDA_ARCHITECTURES_ORIG)
     if(CUDA_ARCH MATCHES "^([0-9]+)")
       if(CMAKE_MATCH_1 GREATER_EQUAL 80)
@@ -175,6 +176,9 @@ macro(setup_cuda_architectures)
       if(CMAKE_MATCH_1 GREATER_EQUAL 100)
         set(ORT_HAS_SM100_OR_LATER ON)
       endif()
+      if(CMAKE_MATCH_1 GREATER_EQUAL 120)
+        set(ORT_HAS_SM120_OR_LATER ON)
+      endif()
     endif()
   endforeach()
   if(ORT_HAS_SM80_OR_LATER)
@@ -186,6 +190,9 @@ macro(setup_cuda_architectures)
   if(ORT_HAS_SM100_OR_LATER)
     add_definitions("-DHAS_SM100_OR_LATER")
   endif()
+  if(ORT_HAS_SM120_OR_LATER)
+    add_definitions("-DHAS_SM120_OR_LATER")
+  endif()
 
   set(ARCHITECTURES_WITH_KERNELS "80" "86" "89" "90" "100" "110" "120")
   foreach(CUDA_ARCH IN LISTS ARCHITECTURES_WITH_KERNELS)
@@ -203,7 +210,7 @@ macro(setup_cuda_architectures)
   endforeach()
 
   # Enable accelerated features (like WGMMA, TMA and setmaxnreg) for SM >= 90.
-  set(ARCHITECTURES_WITH_ACCEL "90" "100" "101" "110" "120")
+  set(ARCHITECTURES_WITH_ACCEL "90" "100" "103" "110" "120" "121")
   unset(CMAKE_CUDA_ARCHITECTURES_NORMALIZED)
   foreach(CUDA_ARCH IN LISTS CMAKE_CUDA_ARCHITECTURES)
     if(CUDA_ARCH MATCHES "^([0-9]+)f$")

cmake/onnxruntime_providers_cuda.cmake

@@ -67,6 +67,14 @@
 
   include(onnxruntime_cuda_source_filters.cmake)
   onnxruntime_filter_cuda_cu_sources(onnxruntime_cuda_contrib_ops_cu_srcs)
+
+  if (NOT onnxruntime_USE_TRT_FUSED_ATTENTION)
+    # Drop the prebuilt TensorRT fused MHA cubin blobs. cudaDriverWrapper is kept because
+    # sparse attention depends on it.
+    list(FILTER onnxruntime_cuda_contrib_ops_cc_srcs EXCLUDE REGEX
+      ".*/bert/tensorrt_fused_multihead_attention/.*(\\.cubin\\.cc|_kernel\\.sm[0-9]+\\.cc)$")
+  endif()
+
   onnxruntime_extract_sm_specific_cuda_sources(onnxruntime_cuda_contrib_ops_cu_srcs
     SM90_SOURCES onnxruntime_cuda_sm90_tma_srcs
     SM120_SOURCES onnxruntime_cuda_sm120_tma_srcs

cmake/onnxruntime_providers_cuda_plugin.cmake

@@ -44,6 +44,13 @@ list(FILTER CUDA_PLUGIN_EP_CU_SRCS EXCLUDE REGEX "onnxruntime/contrib_ops/cuda/c
 list(FILTER CUDA_PLUGIN_EP_CC_SRCS EXCLUDE REGEX "onnxruntime/contrib_ops/cuda/aten_ops/.*")
 list(FILTER CUDA_PLUGIN_EP_CC_SRCS EXCLUDE REGEX "onnxruntime/contrib_ops/cuda/collective/.*")
 
+if (NOT onnxruntime_USE_TRT_FUSED_ATTENTION)
+  # Drop the prebuilt TensorRT fused MHA cubin blobs. cudaDriverWrapper.cc is kept because
+  # sparse attention depends on it.
+  list(FILTER CUDA_PLUGIN_EP_CC_SRCS EXCLUDE REGEX
+    ".*/bert/tensorrt_fused_multihead_attention/.*(\\.cubin\\.cc|_kernel\\.sm[0-9]+\\.cc)$")
+endif()
+
 # Exclude files that include cuda_execution_provider.h (directly or transitively),
 # which conflicts with the adapter shim CUDAExecutionProvider class.
 list(FILTER CUDA_PLUGIN_EP_CC_SRCS EXCLUDE REGEX ".*/cuda_execution_provider\\.cc$")
@@ -123,6 +130,17 @@ onnxruntime_add_shared_library_module(onnxruntime_providers_cuda_plugin
     ${CUDA_PLUGIN_EP_CU_SRCS}
 )
 
+if(WIN32)
+  # Add version information to the packaged plugin DLL.
+  target_sources(onnxruntime_providers_cuda_plugin PRIVATE
+      "${ONNXRUNTIME_ROOT}/core/providers/cuda/onnxruntime_providers_cuda.rc")
+  target_compile_definitions(onnxruntime_providers_cuda_plugin PRIVATE
+      FILE_NAME=\"onnxruntime_providers_cuda.dll\")
+elseif(UNIX AND NOT APPLE)
+  # The build output is packaged directly, so do not embed the build machine's CUDA path.
+  set_target_properties(onnxruntime_providers_cuda_plugin PROPERTIES SKIP_BUILD_RPATH TRUE)
+endif()
+
 # Mirror directory structure in the Visual Studio solution tree under "onnxruntime".
 source_group(TREE ${ONNXRUNTIME_ROOT} PREFIX "onnxruntime" FILES ${CUDA_EP_CC_SRCS} ${CUDA_EP_CU_SRCS})
 source_group(TREE ${ONNXRUNTIME_ROOT} PREFIX "onnxruntime" FILES ${CUDA_CONTRIB_OPS_CC_SRCS} ${CUDA_CONTRIB_OPS_CU_SRCS})
@@ -346,6 +364,12 @@ if(NOT onnxruntime_DISABLE_CONTRIB_OPS)
   if(_cuda_plugin_llm_srcs)
     if(MSVC AND NOT onnxruntime_USE_FP4_QMOE)
       onnxruntime_filter_cuda_archs(_plugin_llm_cuda_architectures MIN_SM 75 EXCLUDE_SM120_REAL)
+      # A native-only Windows ARM64 build has no lower architecture left after the
+      # MSVC SM120 exclusion. Emit PTX privately for this object library so its host
+      # launchers and device kernels are still linked into the plugin.
+      if(NOT _plugin_llm_cuda_architectures AND ORT_HAS_SM120_OR_LATER)
+        set(_plugin_llm_cuda_architectures "120-virtual")
+      endif()
     else()
       onnxruntime_filter_cuda_archs(_plugin_llm_cuda_architectures MIN_SM 75)
     endif()