[CUDA] Add win-arm64 packaging and size options to plugin pipeline

cmake/CMakeLists.txt

@@ -129,6 +129,8 @@ cmake_dependent_option(onnxruntime_USE_FPA_INTB_GEMM "Build FpA IntB gemm cuda k
 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)
@@ -1526,7 +1528,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_plugin.cmake

@@ -123,6 +123,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 +357,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()

tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml

@@ -55,6 +55,22 @@ parameters:
   - dev
   default: dev
 
+# CUDA 13.x always compresses fatbins with -compress-mode=size. Enabling this makes the
+# CUDA 12.8 build match it, which shrinks the binary a lot but raises the minimum driver
+# to the CUDA 12.4 level (Linux >= 550.54.14, Windows >= 551.61).
+- name: enable_cuda_fatbin_size_compression
+  displayName: 'Compress CUDA fatbins for size (CUDA 12.8 only; needs r550+ driver)'
+  type: boolean
+  default: false
+
+# The fpA_intB GEMM/GEMV kernels are the largest single component of the binary after
+# flash attention. They are opt-in at run time (ORT_FPA_INTB_GEMM / ep.cuda.fpa_intb_gemm),
+# but MatMulNBits still uses them automatically when the weights are prepacked.
+- name: enable_fpa_intb_gemm
+  displayName: 'Build fpA_intB GEMM CUDA kernels'
+  type: boolean
+  default: true
+
 - name: cmake_build_type
   type: string
   default: 'Release'
@@ -72,9 +88,9 @@ variables:
   # Non-dev package versions (release, RC) must use Release build type
   - name: invalidBuildTypeConfig
     value: ${{ and(ne(parameters.package_type, 'dev'), ne(parameters.cmake_build_type, 'Release')) }}
-  # aarch64 is only available for CUDA 13.x
-  - name: invalidAArch64Config
-    value: ${{ and(eq(parameters.build_linux_aarch64, true), ne(parameters.cuda_version, '13.x')) }}
+  # ARM64 is only available for CUDA 13.x
+  - name: invalidArm64Config
+    value: ${{ and(or(eq(parameters.build_windows_arm64, true), eq(parameters.build_linux_aarch64, true)), ne(parameters.cuda_version, '13.x')) }}
 
 extends:
   template: v1/1ES.Official.PipelineTemplate.yml@1esPipelines
@@ -106,7 +122,7 @@ extends:
 
     stages:
       # Validate parameter combinations
-      - ${{ if or(eq(variables['invalidBuildTypeConfig'], 'True'), eq(variables['invalidAArch64Config'], 'True')) }}:
+      - ${{ if or(eq(variables['invalidBuildTypeConfig'], 'True'), eq(variables['invalidArm64Config'], 'True')) }}:
         - stage: Validate_Parameters
           displayName: 'Validate Parameters'
           dependsOn: []
@@ -123,11 +139,11 @@ extends:
                   echo "##vso[task.logissue type=error]Non-dev package version requires Release build type."
                   exit 1
                 displayName: 'ERROR: Non-dev package version requires Release build type'
-            - ${{ if eq(variables['invalidAArch64Config'], 'True') }}:
+            - ${{ if eq(variables['invalidArm64Config'], 'True') }}:
               - script: |
-                  echo "##vso[task.logissue type=error]Linux aarch64 build is only available for CUDA 13.x."
+                  echo "##vso[task.logissue type=error]Windows ARM64 and Linux aarch64 builds are only available for CUDA 13.x."
                   exit 1
-                displayName: 'ERROR: aarch64 requires CUDA 13.x'
+                displayName: 'ERROR: ARM64 requires CUDA 13.x'
       - ${{ else }}:
         - template: stages/plugin-cuda-packaging-stage.yml
           parameters:
@@ -140,14 +156,18 @@ extends:
             package_type: ${{ parameters.package_type }}
             version_file: ${{ variables.epVersionFile }}
             cmake_build_type: ${{ parameters.cmake_build_type }}
+            enable_cuda_fatbin_size_compression: ${{ parameters.enable_cuda_fatbin_size_compression }}
+            enable_fpa_intb_gemm: ${{ parameters.enable_fpa_intb_gemm }}
             ${{ if eq(parameters.cuda_version, '12.8') }}:
               python_package_name: 'onnxruntime-ep-cuda12'
               docker_base_image: 'onnxruntimebuildcache.azurecr.io/internal/azureml/onnxruntime/build/cuda12_x64_almalinux8_gcc14:20251017.1'
-              cmake_x64_cuda_archs: '61-real;75-real;86-real;89-real;120-real;120-virtual'
-              cmake_arm64_cuda_archs: '61-real;75-real;86-real;89-real;120-real;120-virtual'
+              cmake_windows_x64_cuda_archs: '61-real;75-real;86-real;89-real;120-real'
+              cmake_linux_x64_cuda_archs: '75-real;80-real;86-real;89-real;90-real;120-real'
             ${{ if eq(parameters.cuda_version, '13.x') }}:
               python_package_name: 'onnxruntime-ep-cuda13'
               docker_base_image: 'onnxruntimebuildcache.azurecr.io/internal/azureml/onnxruntime/build/cuda13_x64_almalinux8_gcc14:20251107.1'
               docker_base_image_aarch64: 'onnxruntimebuildcache.azurecr.io/public/azureml/onnxruntime_build_cuda13_aarch64_almalinux9_gcc14:20260323.1'
-              cmake_x64_cuda_archs: '75-real;80-real;86-real;89-real;90-real;120-real;120-virtual'
-              cmake_arm64_cuda_archs: '110-real;120-real;121-real;120-virtual'
+              cmake_windows_x64_cuda_archs: '75-real;80-real;86-real;89-real;120-real'
+              cmake_windows_arm64_cuda_archs: '120-real;121-real'
+              cmake_linux_x64_cuda_archs: '75-real;80-real;86-real;89-real;90-real;120-real'
+              cmake_linux_aarch64_cuda_archs: '89-real;90-real;100-real;103-real;120-real;121-real'

tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml

@@ -51,6 +51,16 @@ parameters:
    - RelWithDebInfo
    - MinSizeRel
 
+- name: enable_cuda_fatbin_size_compression
+  type: boolean
+  displayName: 'Compress CUDA fatbins with -compress-mode=size'
+  default: false
+
+- name: enable_fpa_intb_gemm
+  type: boolean
+  displayName: 'Build fpA_intB GEMM CUDA kernels'
+  default: true
+
 - name: docker_base_image
   type: string
   displayName: 'Linux x86_64 docker base image'
@@ -61,15 +71,25 @@ parameters:
   displayName: 'Linux aarch64 docker base image'
   default: ''
 
-- name: cmake_x64_cuda_archs
+- name: cmake_windows_x64_cuda_archs
+  type: string
+  displayName: 'CMAKE_CUDA_ARCHITECTURES for Windows x64'
+  default: '61-real;75-real;86-real;89-real;120-real'
+
+- name: cmake_windows_arm64_cuda_archs
+  type: string
+  displayName: 'CMAKE_CUDA_ARCHITECTURES for Windows ARM64'
+  default: '120-real;121-real'
+
+- name: cmake_linux_x64_cuda_archs
   type: string
-  displayName: 'CMAKE_CUDA_ARCHITECTURES for x64'
-  default: '61-real;75-real;86-real;89-real;120-real;120-virtual'
+  displayName: 'CMAKE_CUDA_ARCHITECTURES for Linux x64'
+  default: '75-real;80-real;86-real;89-real;90-real;120-real'
 
-- name: cmake_arm64_cuda_archs
+- name: cmake_linux_aarch64_cuda_archs
   type: string
-  displayName: 'CMAKE_CUDA_ARCHITECTURES for ARM64'
-  default: '110-real;120-real;121-real;120-virtual'
+  displayName: 'CMAKE_CUDA_ARCHITECTURES for Linux aarch64'
+  default: '89-real;90-real;100-real;103-real;120-real;121-real'
 
 - name: python_version
   type: string
@@ -147,11 +167,13 @@ stages:
         parameters:
           arch: 'x64'
           cuda_version: ${{ parameters.cuda_version }}
-          cmake_cuda_archs: ${{ parameters.cmake_x64_cuda_archs }}
+          cmake_cuda_archs: ${{ parameters.cmake_windows_x64_cuda_archs }}
           package_version: ${{ parameters.package_type }}
           version_file: ${{ parameters.version_file }}
           python_package_name: ${{ parameters.python_package_name }}
           cmake_build_type: ${{ parameters.cmake_build_type }}
+          enable_cuda_fatbin_size_compression: ${{ parameters.enable_cuda_fatbin_size_compression }}
+          enable_fpa_intb_gemm: ${{ parameters.enable_fpa_intb_gemm }}
 
     # Windows ARM64
     - ${{ if eq(parameters.build_windows_arm64, true) }}:
@@ -160,11 +182,13 @@ stages:
           arch: 'arm64'
           cuda_version: ${{ parameters.cuda_version }}
           arm64_cuda_version: ${{ parameters.arm64_cuda_version }}
-          cmake_cuda_archs: ${{ parameters.cmake_arm64_cuda_archs }}
+          cmake_cuda_archs: ${{ parameters.cmake_windows_arm64_cuda_archs }}
           package_version: ${{ parameters.package_type }}
           version_file: ${{ parameters.version_file }}
           python_package_name: ${{ parameters.python_package_name }}
           cmake_build_type: ${{ parameters.cmake_build_type }}
+          enable_cuda_fatbin_size_compression: ${{ parameters.enable_cuda_fatbin_size_compression }}
+          enable_fpa_intb_gemm: ${{ parameters.enable_fpa_intb_gemm }}
 
     # Linux x64
     - ${{ if eq(parameters.build_linux_x64, true) }}:
@@ -174,7 +198,7 @@ stages:
           arch: 'x64'
           machine_pool: 'onnxruntime-Ubuntu2404-AMD-CPU'
           cuda_version: ${{ parameters.cuda_version }}
-          cmake_cuda_archs: ${{ parameters.cmake_x64_cuda_archs }}
+          cmake_cuda_archs: ${{ parameters.cmake_linux_x64_cuda_archs }}
           package_version: ${{ parameters.package_type }}
           version_file: ${{ parameters.version_file }}
           python_package_name: ${{ parameters.python_package_name }}
@@ -183,6 +207,8 @@ stages:
           python_version: ${{ parameters.python_version }}
           docker_python_exe_path: ${{ parameters.docker_python_exe_path }}
           artifact_name: cuda_plugin_linux_x64
+          enable_cuda_fatbin_size_compression: ${{ parameters.enable_cuda_fatbin_size_compression }}
+          enable_fpa_intb_gemm: ${{ parameters.enable_fpa_intb_gemm }}
 
     # Linux aarch64 (CUDA 13.x only)
     - ${{ if eq(parameters.build_linux_aarch64, true) }}:
@@ -192,7 +218,7 @@ stages:
           arch: 'aarch64'
           machine_pool: 'onnxruntime-linux-ARM64-CPU-2019'
           cuda_version: ${{ parameters.cuda_version }}
-          cmake_cuda_archs: ${{ parameters.cmake_arm64_cuda_archs }}
+          cmake_cuda_archs: ${{ parameters.cmake_linux_aarch64_cuda_archs }}
           package_version: ${{ parameters.package_type }}
           version_file: ${{ parameters.version_file }}
           python_package_name: ${{ parameters.python_package_name }}
@@ -201,6 +227,8 @@ stages:
           python_version: ${{ parameters.python_version }}
           docker_python_exe_path: ${{ parameters.docker_python_exe_path }}
           artifact_name: cuda_plugin_linux_aarch64
+          enable_cuda_fatbin_size_compression: ${{ parameters.enable_cuda_fatbin_size_compression }}
+          enable_fpa_intb_gemm: ${{ parameters.enable_fpa_intb_gemm }}
 
     # NuGet packaging (runs after all platform builds)
     - template: plugin-cuda-nuget-packaging-stage.yml

tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml

@@ -57,6 +57,14 @@ parameters:
   type: string
   default: 'cuda_plugin_linux_x64'
 
+- name: enable_cuda_fatbin_size_compression
+  type: boolean
+  default: false
+
+- name: enable_fpa_intb_gemm
+  type: boolean
+  default: true
+
 stages:
 - stage: ${{ parameters.stage_name }}
   dependsOn: []
@@ -77,6 +85,16 @@ stages:
         artifactName: ${{ parameters.artifact_name }}
     variables:
       - template: ../templates/common-variables.yml
+      - name: FatbinCompressDefine
+        ${{ if eq(parameters.enable_cuda_fatbin_size_compression, true) }}:
+          value: 'onnxruntime_CUDA_FATBIN_COMPRESS_SIZE=ON'
+        ${{ else }}:
+          value: ''
+      - name: FpaIntBGemmDefine
+        ${{ if eq(parameters.enable_fpa_intb_gemm, true) }}:
+          value: 'onnxruntime_USE_FPA_INTB_GEMM=ON'
+        ${{ else }}:
+          value: 'onnxruntime_USE_FPA_INTB_GEMM=OFF'
     steps:
       - checkout: self
         clean: true
@@ -115,7 +133,7 @@ stages:
         workingDirectory: $(Build.SourcesDirectory)
         displayName: 'Build CUDA Plugin (Python ${{ parameters.python_version }}, ${{ parameters.arch }}, CUDA ${{ parameters.cuda_version }})'
         env:
-          EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine)
+          EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine) $(FatbinCompressDefine) $(FpaIntBGemmDefine)
 
       - script: |
           set -e -x
@@ -125,8 +143,19 @@ stages:
             echo "Error: Expected plugin binary not found at '$plugin_path'. Failing build to avoid publishing an invalid package."
             exit 1
           fi
+
+          dynamic_search_path="$(readelf -d "$plugin_path" | sed -nE 's/.*\((RPATH|RUNPATH)\).*\[([^]]*)\].*/\2/p')"
+          if [[ -n "$dynamic_search_path" && ":$dynamic_search_path:" == *::* ]]; then
+            echo "Error: Plugin binary contains an empty RPATH/RUNPATH component: '$dynamic_search_path'"
+            exit 1
+          fi
+          if [[ "$dynamic_search_path" == *cuda* ]]; then
+            echo "Error: Plugin binary contains a hard-coded CUDA RPATH/RUNPATH: '$dynamic_search_path'"
+            exit 1
+          fi
+
           cp "$plugin_path" "$(Build.ArtifactStagingDirectory)/bin/"
-        displayName: 'Copy plugin binaries'
+        displayName: 'Verify and copy plugin binaries'
 
       - script: |
           set -e -x