From e030db4dab90c6751ab1ff3b0d3ced21aa1b7f92 Mon Sep 17 00:00:00 2001 From: Tianlei Wu Date: Tue, 4 Aug 2026 18:02:36 +0000 Subject: [PATCH 1/5] update cuda plugin ep packaging --- cmake/onnxruntime_providers_cuda_plugin.cmake | 11 ++++++++ .../stages/plugin-linux-cuda-stage.yml | 13 +++++++++- .../stages/plugin-win-cuda-stage.yml | 25 +++++++++++++++++-- 3 files changed, 46 insertions(+), 3 deletions(-) diff --git a/cmake/onnxruntime_providers_cuda_plugin.cmake b/cmake/onnxruntime_providers_cuda_plugin.cmake index ffb783f1a7cc5..8d535e960612a 100644 --- a/cmake/onnxruntime_providers_cuda_plugin.cmake +++ b/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}) diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml index e665ac35dd8ed..f5613bbaf3dfb 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml @@ -125,8 +125,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 diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml index 2ec3e21391037..2a4db1c2d595f 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml @@ -269,8 +269,29 @@ stages: Write-Error "Expected plugin binary not found at '$pluginPath'. Failing build to avoid publishing an invalid package." exit 1 } - Write-Host "Verified plugin binary exists at: $pluginPath" - displayName: 'Verify plugin binary exists' + + $versionInfo = (Get-Item $pluginPath).VersionInfo + $requiredProperties = @( + 'FileDescription', + 'FileVersion', + 'OriginalFilename', + 'ProductName', + 'ProductVersion', + 'LegalCopyright', + 'Language' + ) + foreach ($property in $requiredProperties) { + if ([string]::IsNullOrWhiteSpace($versionInfo.$property)) { + throw "Plugin DLL version property '$property' is missing." + } + } + if ($versionInfo.OriginalFilename -ne 'onnxruntime_providers_cuda.dll') { + throw "Unexpected OriginalFilename '$($versionInfo.OriginalFilename)'." + } + + Write-Host "Verified plugin DLL version information:" + $versionInfo | Format-List $requiredProperties + displayName: 'Verify plugin binary and version information' - task: CopyFiles@2 displayName: 'Copy plugin binaries to staging directory' From 3ff27845bede0a920beea7c5351c784221942df6 Mon Sep 17 00:00:00 2001 From: Tianlei Wu Date: Tue, 4 Aug 2026 18:59:28 +0000 Subject: [PATCH 2/5] update arch list --- cmake/external/cuda_configuration.cmake | 2 +- .../azure-pipelines/plugin-cuda-pipeline.yml | 24 ++++++++------- .../stages/plugin-cuda-packaging-stage.yml | 30 ++++++++++++------- 3 files changed, 34 insertions(+), 22 deletions(-) diff --git a/cmake/external/cuda_configuration.cmake b/cmake/external/cuda_configuration.cmake index 9d7bfb0e5ad61..b9ae1d0dd5a13 100644 --- a/cmake/external/cuda_configuration.cmake +++ b/cmake/external/cuda_configuration.cmake @@ -203,7 +203,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$") diff --git a/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml b/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml index 4b713fbc1a3d7..08b72294dd7f5 100644 --- a/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml +++ b/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml @@ -72,9 +72,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 +106,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 +123,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: @@ -143,11 +143,13 @@ extends: ${{ 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' diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml index 5f9837fddb69e..700d8d9d99470 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml @@ -61,15 +61,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 x64' - default: '61-real;75-real;86-real;89-real;120-real;120-virtual' + displayName: 'CMAKE_CUDA_ARCHITECTURES for Windows x64' + default: '61-real;75-real;86-real;89-real;120-real' -- name: cmake_arm64_cuda_archs +- name: cmake_windows_arm64_cuda_archs type: string - displayName: 'CMAKE_CUDA_ARCHITECTURES for ARM64' - default: '110-real;120-real;121-real;120-virtual' + 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 Linux x64' + default: '75-real;80-real;86-real;89-real;90-real;120-real' + +- name: cmake_linux_aarch64_cuda_archs + type: string + 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,7 +157,7 @@ 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 }} @@ -160,7 +170,7 @@ 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 }} @@ -174,7 +184,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 }} @@ -192,7 +202,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 }} From a0b219da8d8a593754358e9a21302ff06fcff537 Mon Sep 17 00:00:00 2001 From: Tianlei Wu Date: Tue, 4 Aug 2026 19:12:56 +0000 Subject: [PATCH 3/5] Add compress size option --- cmake/CMakeLists.txt | 9 ++++++++- .../github/azure-pipelines/plugin-cuda-pipeline.yml | 9 +++++++++ .../stages/plugin-cuda-packaging-stage.yml | 9 +++++++++ .../stages/plugin-linux-cuda-stage.yml | 9 ++++++++- .../azure-pipelines/stages/plugin-win-cuda-stage.yml | 12 ++++++++++++ 5 files changed, 46 insertions(+), 2 deletions(-) diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index b9b63c0595c45..f875ad0fb64b6 100644 --- a/cmake/CMakeLists.txt +++ b/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") diff --git a/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml b/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml index 08b72294dd7f5..a8b5f3610d4a0 100644 --- a/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml +++ b/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml @@ -55,6 +55,14 @@ 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 + - name: cmake_build_type type: string default: 'Release' @@ -140,6 +148,7 @@ 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 }} ${{ 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' diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml index 700d8d9d99470..ebc369c13a5ae 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml @@ -51,6 +51,11 @@ parameters: - RelWithDebInfo - MinSizeRel +- name: enable_cuda_fatbin_size_compression + type: boolean + displayName: 'Compress CUDA fatbins with -compress-mode=size' + default: false + - name: docker_base_image type: string displayName: 'Linux x86_64 docker base image' @@ -162,6 +167,7 @@ stages: 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 }} # Windows ARM64 - ${{ if eq(parameters.build_windows_arm64, true) }}: @@ -175,6 +181,7 @@ stages: 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 }} # Linux x64 - ${{ if eq(parameters.build_linux_x64, true) }}: @@ -193,6 +200,7 @@ 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 }} # Linux aarch64 (CUDA 13.x only) - ${{ if eq(parameters.build_linux_aarch64, true) }}: @@ -211,6 +219,7 @@ 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 }} # NuGet packaging (runs after all platform builds) - template: plugin-cuda-nuget-packaging-stage.yml diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml index f5613bbaf3dfb..92793e3f8a109 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml @@ -57,6 +57,10 @@ parameters: type: string default: 'cuda_plugin_linux_x64' +- name: enable_cuda_fatbin_size_compression + type: boolean + default: false + stages: - stage: ${{ parameters.stage_name }} dependsOn: [] @@ -115,7 +119,10 @@ stages: workingDirectory: $(Build.SourcesDirectory) displayName: 'Build CUDA Plugin (Python ${{ parameters.python_version }}, ${{ parameters.arch }}, CUDA ${{ parameters.cuda_version }})' env: - EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine) + ${{ if eq(parameters.enable_cuda_fatbin_size_compression, true) }}: + EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine) onnxruntime_CUDA_FATBIN_COMPRESS_SIZE=ON + ${{ else }}: + EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine) - script: | set -e -x diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml index 2a4db1c2d595f..26a2618d74511 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml @@ -47,6 +47,10 @@ parameters: type: string default: '61-real;75-real;86-real;89-real;120-real;120-virtual' +- name: enable_cuda_fatbin_size_compression + type: boolean + default: false + stages: - stage: Win_plugin_cuda_${{ parameters.arch }}_Build dependsOn: [] @@ -85,6 +89,11 @@ stages: value: '-Dorg.gradle.daemon=false' - name: VSGenerator value: 'Visual Studio 17 2022' + - name: FatbinCompressOption + ${{ if eq(parameters.enable_cuda_fatbin_size_compression, true) }}: + value: '--cmake_extra_defines onnxruntime_CUDA_FATBIN_COMPRESS_SIZE=ON' + ${{ else }}: + value: '' steps: - checkout: self clean: true @@ -192,6 +201,7 @@ stages: --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES="${{ parameters.cmake_cuda_archs }}" --cmake_extra_defines onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON --cmake_extra_defines $(PluginEpVersionDefine) + $(FatbinCompressOption) $(TelemetryOption) workingDirectory: '$(Build.BinariesDirectory)' @@ -221,6 +231,7 @@ stages: --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES="${{ parameters.cmake_cuda_archs }}" --cmake_extra_defines onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON --cmake_extra_defines $(PluginEpVersionDefine) + $(FatbinCompressOption) $(TelemetryOption) workingDirectory: '$(Build.BinariesDirectory)' @@ -252,6 +263,7 @@ stages: --cmake_extra_defines onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON --cmake_extra_defines onnxruntime_BUILD_UNIT_TESTS=OFF --cmake_extra_defines $(PluginEpVersionDefine) + $(FatbinCompressOption) $(TelemetryOption) workingDirectory: '$(Build.BinariesDirectory)' From 1edd1f5c5a116d48fed9672d0eda8907327dfb99 Mon Sep 17 00:00:00 2001 From: Tianlei Wu Date: Tue, 4 Aug 2026 19:21:05 +0000 Subject: [PATCH 4/5] Add enable_fpa_intb_gemm parameter --- .../azure-pipelines/plugin-cuda-pipeline.yml | 9 +++++++++ .../stages/plugin-cuda-packaging-stage.yml | 9 +++++++++ .../stages/plugin-linux-cuda-stage.yml | 19 +++++++++++++++---- .../stages/plugin-win-cuda-stage.yml | 12 ++++++++++++ 4 files changed, 45 insertions(+), 4 deletions(-) diff --git a/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml b/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml index a8b5f3610d4a0..fade2cc416e49 100644 --- a/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml +++ b/tools/ci_build/github/azure-pipelines/plugin-cuda-pipeline.yml @@ -63,6 +63,14 @@ parameters: 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' @@ -149,6 +157,7 @@ extends: 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' diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml index ebc369c13a5ae..3522b92dd27be 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-cuda-packaging-stage.yml @@ -56,6 +56,11 @@ parameters: 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' @@ -168,6 +173,7 @@ stages: 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) }}: @@ -182,6 +188,7 @@ stages: 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) }}: @@ -201,6 +208,7 @@ stages: 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) }}: @@ -220,6 +228,7 @@ stages: 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 diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml index 92793e3f8a109..a1452b89ebb9a 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-linux-cuda-stage.yml @@ -61,6 +61,10 @@ parameters: type: boolean default: false +- name: enable_fpa_intb_gemm + type: boolean + default: true + stages: - stage: ${{ parameters.stage_name }} dependsOn: [] @@ -81,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 @@ -119,10 +133,7 @@ stages: workingDirectory: $(Build.SourcesDirectory) displayName: 'Build CUDA Plugin (Python ${{ parameters.python_version }}, ${{ parameters.arch }}, CUDA ${{ parameters.cuda_version }})' env: - ${{ if eq(parameters.enable_cuda_fatbin_size_compression, true) }}: - EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine) onnxruntime_CUDA_FATBIN_COMPRESS_SIZE=ON - ${{ else }}: - EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine) + EXTRA_CMAKE_DEFINES: $(PluginEpVersionDefine) $(FatbinCompressDefine) $(FpaIntBGemmDefine) - script: | set -e -x diff --git a/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml b/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml index 26a2618d74511..c16f40365f471 100644 --- a/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml +++ b/tools/ci_build/github/azure-pipelines/stages/plugin-win-cuda-stage.yml @@ -51,6 +51,10 @@ parameters: type: boolean default: false +- name: enable_fpa_intb_gemm + type: boolean + default: true + stages: - stage: Win_plugin_cuda_${{ parameters.arch }}_Build dependsOn: [] @@ -94,6 +98,11 @@ stages: value: '--cmake_extra_defines onnxruntime_CUDA_FATBIN_COMPRESS_SIZE=ON' ${{ else }}: value: '' + - name: FpaIntBGemmOption + ${{ if eq(parameters.enable_fpa_intb_gemm, true) }}: + value: '--cmake_extra_defines onnxruntime_USE_FPA_INTB_GEMM=ON' + ${{ else }}: + value: '--cmake_extra_defines onnxruntime_USE_FPA_INTB_GEMM=OFF' steps: - checkout: self clean: true @@ -202,6 +211,7 @@ stages: --cmake_extra_defines onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON --cmake_extra_defines $(PluginEpVersionDefine) $(FatbinCompressOption) + $(FpaIntBGemmOption) $(TelemetryOption) workingDirectory: '$(Build.BinariesDirectory)' @@ -232,6 +242,7 @@ stages: --cmake_extra_defines onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON --cmake_extra_defines $(PluginEpVersionDefine) $(FatbinCompressOption) + $(FpaIntBGemmOption) $(TelemetryOption) workingDirectory: '$(Build.BinariesDirectory)' @@ -264,6 +275,7 @@ stages: --cmake_extra_defines onnxruntime_BUILD_UNIT_TESTS=OFF --cmake_extra_defines $(PluginEpVersionDefine) $(FatbinCompressOption) + $(FpaIntBGemmOption) $(TelemetryOption) workingDirectory: '$(Build.BinariesDirectory)' From 528df3cda01e145a4a4633c535a1ca9b4c297c44 Mon Sep 17 00:00:00 2001 From: Tianlei Wu Date: Wed, 5 Aug 2026 00:30:59 +0000 Subject: [PATCH 5/5] fix windows arm64 cuda 13 build --- cmake/external/cuda_configuration.cmake | 7 +++++++ cmake/onnxruntime_providers_cuda_plugin.cmake | 6 ++++++ 2 files changed, 13 insertions(+) diff --git a/cmake/external/cuda_configuration.cmake b/cmake/external/cuda_configuration.cmake index b9ae1d0dd5a13..ff8b1a00aaea1 100644 --- a/cmake/external/cuda_configuration.cmake +++ b/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) diff --git a/cmake/onnxruntime_providers_cuda_plugin.cmake b/cmake/onnxruntime_providers_cuda_plugin.cmake index 8d535e960612a..061f776f4d879 100644 --- a/cmake/onnxruntime_providers_cuda_plugin.cmake +++ b/cmake/onnxruntime_providers_cuda_plugin.cmake @@ -357,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()