From da53553e3386c88c2e09fbae7ae23a1acc351434 Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 15:28:38 +0000 Subject: [PATCH 01/19] Fix Windows CUDA internal-test target dependency cycle Keep the module-to-host link dependency on Windows and remove the reverse build-order edge from the provider test executable. Document the module target as the targeted Windows internal-test build entry point. Fixes #32804 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_unittests.cmake | 6 ++++++ docs/contrib_ops/cuda/matmul_nbits.md | 12 ++++++++++++ 2 files changed, 18 insertions(+) diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index dfcbfdc7e186b..b7460187188c6 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1450,6 +1450,12 @@ block() ) set(onnxruntime_provider_test_deps ${onnxruntime_test_providers_dependencies}) + if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) + # The module links against this executable's import library on Windows, so it must + # build after the executable. It remains part of the default build; for a targeted + # internal-test build, build onnxruntime_providers_cuda_ut to get both artifacts. + list(REMOVE_ITEM onnxruntime_provider_test_deps onnxruntime_providers_cuda_ut) + endif() AddTest( TARGET onnxruntime_provider_test diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index 5af0bf21b8a63..0592bc7c1b06f 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -479,6 +479,18 @@ present. `ComputeInternal` then: ./onnxruntime_provider_test --gtest_filter=CUDA_EP_Unittest.* ``` + For the non-plugin CUDA EP, configure with `onnxruntime_USE_CUDA=ON`, + `onnxruntime_BUILD_UNIT_TESTS=ON`, and `onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON`. + On Windows, a targeted build must use the module target: + + ```powershell + cmake --build --config Release --target onnxruntime_providers_cuda_ut + ``` + + This builds `onnxruntime_provider_test.exe` first, then links the internal-test + DLL against its import library. Building only the executable does not build + the DLL on Windows. The default build includes both artifacts. + This wrapper executes the internal CUDA-UT shared library and covers the fpA_intB / MatMulNBits groupwise GEMM tests under [onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc](../../../onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc) From ae177b55a1051d7334a30cd79afc9e1286d54acf Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 15:30:45 +0000 Subject: [PATCH 02/19] Avoid adding the Windows CUDA test back dependency Guard the shared dependency-list append at its source instead of filtering the provider-test copy later. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_unittests.cmake | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index b7460187188c6..6a473f4397193 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1056,7 +1056,12 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_ "$<$>:/wd4100>") endif() - list(APPEND onnxruntime_test_providers_dependencies onnxruntime_providers_cuda_ut) + # On Windows, the module links against onnxruntime_provider_test's import library + # and must build after the executable. Adding the reverse dependency would create + # a cycle. The module remains part of the default build on all platforms. + if (NOT WIN32) + list(APPEND onnxruntime_test_providers_dependencies onnxruntime_providers_cuda_ut) + endif() endif() if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_PLUGIN AND @@ -1450,12 +1455,6 @@ block() ) set(onnxruntime_provider_test_deps ${onnxruntime_test_providers_dependencies}) - if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) - # The module links against this executable's import library on Windows, so it must - # build after the executable. It remains part of the default build; for a targeted - # internal-test build, build onnxruntime_providers_cuda_ut to get both artifacts. - list(REMOVE_ITEM onnxruntime_provider_test_deps onnxruntime_providers_cuda_ut) - endif() AddTest( TARGET onnxruntime_provider_test From c7970cff9307e6cb4b9e433d88832665abe03af3 Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 15:35:40 +0000 Subject: [PATCH 03/19] Remove the CUDA test module build prerequisite on all platforms Build the dynamically loaded module explicitly for targeted builds instead of making it a prerequisite of the test executable. Retain the Windows module-to-executable import-library link. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_unittests.cmake | 6 ------ docs/contrib_ops/cuda/matmul_nbits.md | 13 +++++++------ 2 files changed, 7 insertions(+), 12 deletions(-) diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index 6a473f4397193..3d5efb6c5a3b6 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1056,12 +1056,6 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_ "$<$>:/wd4100>") endif() - # On Windows, the module links against onnxruntime_provider_test's import library - # and must build after the executable. Adding the reverse dependency would create - # a cycle. The module remains part of the default build on all platforms. - if (NOT WIN32) - list(APPEND onnxruntime_test_providers_dependencies onnxruntime_providers_cuda_ut) - endif() endif() if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_PLUGIN AND diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index 0592bc7c1b06f..30aec74391e85 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -481,15 +481,16 @@ present. `ComputeInternal` then: For the non-plugin CUDA EP, configure with `onnxruntime_USE_CUDA=ON`, `onnxruntime_BUILD_UNIT_TESTS=ON`, and `onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON`. - On Windows, a targeted build must use the module target: + For a targeted build, explicitly build both the executable and the module: - ```powershell - cmake --build --config Release --target onnxruntime_providers_cuda_ut + ```bash + cmake --build --config Release --target onnxruntime_provider_test onnxruntime_providers_cuda_ut ``` - This builds `onnxruntime_provider_test.exe` first, then links the internal-test - DLL against its import library. Building only the executable does not build - the DLL on Windows. The default build includes both artifacts. + Building only the executable does not build the dynamically loaded module on + any platform. The default build includes both artifacts. On Windows, the module + links against the executable's import library, so CMake builds the executable + first even when only the module target is requested. This wrapper executes the internal CUDA-UT shared library and covers the fpA_intB / MatMulNBits groupwise GEMM tests under From 64dccfb0755469ac95c47683faae37690a3b96b3 Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 15:44:42 +0000 Subject: [PATCH 04/19] Preserve CUDA internal-test builds with a Windows aggregate target Separate runtime module prerequisites from binary build dependencies. In the Windows non-plugin internal-test configuration, keep onnxruntime_provider_test as the public aggregate target and link the module against a separate executable target with the original output filename. Preserve CTest names, environment, timeout, reporting and executable PCH settings. Other configurations retain their executable target and module build dependency. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_test_pch.cmake | 2 +- cmake/onnxruntime_unittests.cmake | 60 ++++++++++++++++++--------- docs/contrib_ops/cuda/matmul_nbits.md | 12 +++--- 3 files changed, 48 insertions(+), 26 deletions(-) diff --git a/cmake/onnxruntime_test_pch.cmake b/cmake/onnxruntime_test_pch.cmake index 4a8735a9c346c..88598b65c1ff3 100644 --- a/cmake/onnxruntime_test_pch.cmake +++ b/cmake/onnxruntime_test_pch.cmake @@ -6,7 +6,7 @@ if(CMAKE_CXX_COMPILER_ID MATCHES "MSVC") "${CMAKE_CURRENT_SOURCE_DIR}/test_pch.h" ) if (TARGET onnxruntime_provider_test) - target_precompile_headers(onnxruntime_provider_test PRIVATE + target_precompile_headers(${onnxruntime_provider_test_target} PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/test_pch.h" ) endif() diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index 3d5efb6c5a3b6..37a7791592662 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -58,7 +58,10 @@ function(onnxruntime_disable_gtest_character_conversion_as_error target_name) endfunction() function(AddTest) - cmake_parse_arguments(_UT "DYN" "TARGET" "LIBS;SOURCES;DEPENDS;TEST_ARGS" ${ARGN}) + cmake_parse_arguments(_UT "DYN" "TARGET;TEST_NAME" "LIBS;SOURCES;DEPENDS;TEST_ARGS" ${ARGN}) + if (NOT _UT_TEST_NAME) + set(_UT_TEST_NAME ${_UT_TARGET}) + endif() list(REMOVE_DUPLICATES _UT_SOURCES) filter_test_srcs(_UT_SOURCES) @@ -272,7 +275,7 @@ function(AddTest) if (onnxruntime_ENABLE_WEBASSEMBLY_THREADS) list(APPEND TEST_NPM_FLAGS "--wasm-threads") endif() - add_test(NAME ${_UT_TARGET} + add_test(NAME ${_UT_TEST_NAME} COMMAND ${NPM_CLI} test -- ${TEST_NPM_FLAGS} --entry=${_UT_TARGET} ${TEST_ARGS} WORKING_DIRECTORY $ ) @@ -290,20 +293,20 @@ function(AddTest) set(NODE_EXECUTABLE node) endif() - add_test(NAME ${_UT_TARGET} + add_test(NAME ${_UT_TEST_NAME} COMMAND ${NODE_EXECUTABLE} ${TEST_NODE_FLAGS} ${_UT_TARGET}.js ${TEST_ARGS} WORKING_DIRECTORY $ ) endif() # Set test timeout to 3 hours. - set_tests_properties(${_UT_TARGET} PROPERTIES TIMEOUT 10800) + set_tests_properties(${_UT_TEST_NAME} PROPERTIES TIMEOUT 10800) else() - add_test(NAME ${_UT_TARGET} + add_test(NAME ${_UT_TEST_NAME} COMMAND ${_UT_TARGET} ${TEST_ARGS} WORKING_DIRECTORY $ ) # Set test timeout to 3 hours. - set_tests_properties(${_UT_TARGET} PROPERTIES TIMEOUT 10800) + set_tests_properties(${_UT_TEST_NAME} PROPERTIES TIMEOUT 10800) endif() endif() endfunction(AddTest) @@ -998,6 +1001,7 @@ set(all_tests ${onnxruntime_test_lora_src} ) +set(onnxruntime_test_providers_runtime_dependencies) if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_AS_PLUGIN) if (NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_REDUCED_OPS_BUILD) set(onnxruntime_test_cuda_kernels_src_patterns "${TEST_SRC_DIR}/contrib_ops/cuda_kernels/*.cc") @@ -1055,7 +1059,7 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_ target_compile_options(onnxruntime_providers_cuda_ut PRIVATE "$<$:SHELL:--compiler-options /wd4100>" "$<$>:/wd4100>") endif() - + list(APPEND onnxruntime_test_providers_runtime_dependencies onnxruntime_providers_cuda_ut) endif() if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_PLUGIN AND @@ -1126,7 +1130,7 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P endif() endif() -set(all_dependencies ${onnxruntime_test_providers_dependencies} ) +set(all_dependencies ${onnxruntime_test_providers_dependencies} ${onnxruntime_test_providers_runtime_dependencies}) if (onnxruntime_ENABLE_TRAINING) list(APPEND all_tests ${onnxruntime_test_training_src}) @@ -1420,6 +1424,10 @@ endif() # Execution provider-related tests. # These also have some support for dynamically specified plugin EPs. if (NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_REDUCED_OPS_BUILD) +set(onnxruntime_provider_test_target onnxruntime_provider_test) +if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) + set(onnxruntime_provider_test_target onnxruntime_provider_test_executable) +endif() block() set(supporting_test_srcs ${TEST_SRC_DIR}/common/cuda_op_test_utils.cc @@ -1449,15 +1457,29 @@ block() ) set(onnxruntime_provider_test_deps ${onnxruntime_test_providers_dependencies}) + if (onnxruntime_provider_test_target STREQUAL "onnxruntime_provider_test") + list(APPEND onnxruntime_provider_test_deps ${onnxruntime_test_providers_runtime_dependencies}) + endif() AddTest( - TARGET onnxruntime_provider_test + TARGET ${onnxruntime_provider_test_target} + TEST_NAME onnxruntime_provider_test SOURCES ${onnxruntime_provider_test_srcs} LIBS ${onnxruntime_provider_test_libs} DEPENDS ${onnxruntime_provider_test_deps} ) - onnxruntime_apply_test_target_workarounds(onnxruntime_provider_test) + if (NOT onnxruntime_provider_test_target STREQUAL "onnxruntime_provider_test") + # Keep the public build target responsible for both runtime artifacts without + # making the executable depend on the module that imports its symbols. + set_target_properties(${onnxruntime_provider_test_target} PROPERTIES OUTPUT_NAME onnxruntime_provider_test) + add_custom_target(onnxruntime_provider_test ALL) + add_dependencies(onnxruntime_provider_test + ${onnxruntime_provider_test_target} ${onnxruntime_test_providers_runtime_dependencies}) + set_target_properties(onnxruntime_provider_test PROPERTIES FOLDER "ONNXRuntimeTest") + endif() + + onnxruntime_apply_test_target_workarounds(${onnxruntime_provider_test_target}) onnxruntime_set_plugin_ep_test_environment(onnxruntime_provider_test) # The CUDA EP internal unit tests (onnxruntime_providers_cuda_ut) are built as a shared-library @@ -1466,7 +1488,7 @@ block() # InferenceSession, whose symbols are statically linked into this executable. Mirror # onnxruntime_test_all and export them so the dlopen'd module can resolve them at load time; # without this the module fails to load with an undefined-symbol error. - set_target_properties(onnxruntime_provider_test PROPERTIES ENABLE_EXPORTS 1) + set_target_properties(${onnxruntime_provider_test_target} PROPERTIES ENABLE_EXPORTS 1) # On Windows, ENABLE_EXPORTS makes CMake emit an import library (onnxruntime_provider_test.lib) # for the exported symbols, but a MODULE library (onnxruntime_providers_cuda_ut, built via @@ -1477,17 +1499,17 @@ block() # On Linux the runtime -rdynamic export path (above) is sufficient, so this is Windows-only. # Note: onnxruntime_providers_cuda_ut only exists in the non-plugin CUDA-EP-internal-tests path. if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) - target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE onnxruntime_provider_test) + target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE ${onnxruntime_provider_test_target}) endif() if (onnxruntime_USE_CUDA AND onnxruntime_BUILD_CUDA_EP_AS_PLUGIN) - target_compile_definitions(onnxruntime_provider_test PRIVATE + target_compile_definitions(${onnxruntime_provider_test_target} PRIVATE ORT_UNIT_TEST_CUDA_PLUGIN_EP_LIBRARY_PATH="$" ORT_UNIT_TEST_HAS_CUDA_PLUGIN_EP=1) endif() if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_PLUGIN) - target_link_libraries(onnxruntime_provider_test PRIVATE + target_link_libraries(${onnxruntime_provider_test_target} PRIVATE CUDA::cudart CUDA::cublas CUDA::cublasLt @@ -1503,19 +1525,19 @@ block() target_include_directories(qnn_sdk_headers_include INTERFACE ${onnxruntime_QNN_HOME}/include ${onnxruntime_QNN_HOME}/include/QNN) - target_link_libraries(onnxruntime_provider_test PRIVATE qnn_sdk_headers_include) + target_link_libraries(${onnxruntime_provider_test_target} PRIVATE qnn_sdk_headers_include) endif() # enable dynamic plugin EP usage - target_compile_definitions(onnxruntime_provider_test PRIVATE ORT_UNIT_TEST_ENABLE_DYNAMIC_PLUGIN_EP_USAGE) - onnxruntime_apply_emscripten_test_link_settings(onnxruntime_provider_test) + target_compile_definitions(${onnxruntime_provider_test_target} PRIVATE ORT_UNIT_TEST_ENABLE_DYNAMIC_PLUGIN_EP_USAGE) + onnxruntime_apply_emscripten_test_link_settings(${onnxruntime_provider_test_target}) if (IOS) add_custom_command( - TARGET onnxruntime_provider_test POST_BUILD + TARGET ${onnxruntime_provider_test_target} POST_BUILD COMMAND ${CMAKE_COMMAND} -E copy_directory ${TEST_DATA_SRC} - $/testdata) + $/testdata) endif() endblock() endif() diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index 30aec74391e85..f1b49da146f04 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -481,16 +481,16 @@ present. `ComputeInternal` then: For the non-plugin CUDA EP, configure with `onnxruntime_USE_CUDA=ON`, `onnxruntime_BUILD_UNIT_TESTS=ON`, and `onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON`. - For a targeted build, explicitly build both the executable and the module: + Building the provider-test target also builds the internal-test module: ```bash - cmake --build --config Release --target onnxruntime_provider_test onnxruntime_providers_cuda_ut + cmake --build --config Release --target onnxruntime_provider_test ``` - Building only the executable does not build the dynamically loaded module on - any platform. The default build includes both artifacts. On Windows, the module - links against the executable's import library, so CMake builds the executable - first even when only the module target is requested. + On Windows, this is an aggregate build target: it builds the executable first, + then the module that links against the executable's import library. The executable + remains named `onnxruntime_provider_test.exe`, and the CTest name remains + `onnxruntime_provider_test`. The default build also includes both artifacts. This wrapper executes the internal CUDA-UT shared library and covers the fpA_intB / MatMulNBits groupwise GEMM tests under From 82428e11acd3359338c612f43c0d66e1df3a6ba5 Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 15:57:58 +0000 Subject: [PATCH 05/19] Require Windows CI to build and execute CUDA internal tests Evaluate the internal-test option after its prerequisites. Reject a missing test module in the Windows build job and explicitly run the wrapper on the GPU runner, requiring a fresh completed-test XML result so absent or skipped tests cannot pass silently. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .github/workflows/windows_cuda.yml | 23 +++++++++++++++++++++++ cmake/CMakeLists.txt | 11 ++++++----- docs/contrib_ops/cuda/matmul_nbits.md | 4 ++++ 3 files changed, 33 insertions(+), 5 deletions(-) diff --git a/.github/workflows/windows_cuda.yml b/.github/workflows/windows_cuda.yml index 7eaf7e06b2df5..be0287eb1f2eb 100644 --- a/.github/workflows/windows_cuda.yml +++ b/.github/workflows/windows_cuda.yml @@ -121,6 +121,11 @@ jobs: exit $lastExitCode } + $cudaTestModule = Join-Path $env:OnnxRuntimeBuildDirectory "RelWithDebInfo\RelWithDebInfo\onnxruntime_providers_cuda_ut.dll" + if (-not (Test-Path -LiteralPath $cudaTestModule -PathType Leaf)) { + throw "CUDA EP internal tests were requested, but the build did not produce $cudaTestModule" + } + # Clean up the output directory before uploading artifacts $outputDir = "${{ runner.temp }}\build\RelWithDebInfo" Write-Host "Cleaning up files from $outputDir..." @@ -230,6 +235,24 @@ jobs: shell: pwsh run: nvidia-smi + - name: Verify CUDA EP internal tests execute + working-directory: ${{ runner.temp }}\build\RelWithDebInfo\RelWithDebInfo + shell: pwsh + run: | + if (Test-Path -LiteralPath cuda_ep_internal_ci.xml) { + Remove-Item -LiteralPath cuda_ep_internal_ci.xml + } + .\onnxruntime_provider_test.exe --gtest_filter=CUDA_EP_Unittest.All --gtest_output=xml:cuda_ep_internal_ci.xml + if ($LASTEXITCODE -ne 0) { + exit $LASTEXITCODE + } + + [xml]$results = Get-Content -LiteralPath cuda_ep_internal_ci.xml -Raw + $executedTest = $results.SelectNodes("/testsuites/testsuite[@name='CUDA_EP_Unittest']/testcase[@name='All' and @status='run' and @result='completed']") + if ($executedTest.Count -ne 1) { + throw "CUDA_EP_Unittest.All did not execute: the test must be present and must not be skipped." + } + - name: Run Tests working-directory: ${{ runner.temp }} run: | diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index 21f1ee7c3c4eb..e9eccba2d5bf2 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -70,11 +70,6 @@ option(onnxruntime_ENABLE_PYTHON "Enable python bindings" OFF) option(onnxruntime_ENABLE_MEMLEAK_CHECKER "Experimental: Enable memory leak checker in Windows debug build" OFF) option(onnxruntime_ENABLE_CONVSYMKERNELAVX2_SAT_CHECKER "Experimental: Enable ConvSymKernelAvx2 assembly saturation checker in build" OFF) option(onnxruntime_USE_CUDA "Build with CUDA support" OFF) -# Enable ONNX Runtime CUDA EP's internal unit tests that directly access the EP's internal functions instead of through -# OpKernels. When the option is ON, we will have two copies of GTest library in the same process. It is not a typical -# use. If you hit any problem with that, please do not report it to GTest. Turn OFF the following build option instead. -cmake_dependent_option(onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS "Build with CUDA unit tests" OFF "onnxruntime_USE_CUDA;onnxruntime_BUILD_UNIT_TESTS" OFF) - cmake_dependent_option(onnxruntime_USE_CUDA_NHWC_OPS "Build CUDA with NHWC op support" ON "onnxruntime_USE_CUDA" OFF) cmake_dependent_option(onnxruntime_BUILD_CUDA_EP_AS_PLUGIN "Build CUDA EP as a separate plugin shared library instead of the legacy in-tree provider" OFF "onnxruntime_USE_CUDA" OFF) option(onnxruntime_BUILD_CUDA_QUANT_PREPROCESS "Build CUDA weight-packing module onnxruntime_cuda_quant_preprocess.so" OFF) @@ -97,6 +92,12 @@ option(onnxruntime_USE_ARM_NEON_NCHWC "Build with ARM Neon NCHWc kernels in MLAS option(onnxruntime_USE_KLEIDIAI "Build with KleidiAI integration in MLAS" OFF) option(onnxruntime_USE_QMX_KLEIDIAI_COEXIST "Build with QMX and Arm KLEIDIAI libraries" OFF) option(onnxruntime_BUILD_UNIT_TESTS "Build ONNXRuntime unit tests" ON) +# Declare both prerequisites before evaluating this dependent option on a fresh configure. +# Enable ONNX Runtime CUDA EP's internal unit tests that directly access the EP's internal functions instead of through +# OpKernels. When the option is ON, we will have two copies of GTest library in the same process. It is not a typical +# use. If you hit any problem with that, please do not report it to GTest. Turn OFF the following build option instead. +cmake_dependent_option(onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS "Build with CUDA unit tests" OFF "onnxruntime_USE_CUDA;onnxruntime_BUILD_UNIT_TESTS" OFF) + # Materialize the ONNX node-test corpus from ONNX's Python generators into the build tree at # configure/build time, instead of depending on the on-disk corpus shipped in the ONNX source # archive. This detaches ORT from ONNX PR #7959 (which deletes onnx/backend/test/data/node). diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index f1b49da146f04..a6f11120f168b 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -492,6 +492,10 @@ present. `ComputeInternal` then: remains named `onnxruntime_provider_test.exe`, and the CTest name remains `onnxruntime_provider_test`. The default build also includes both artifacts. + Windows CUDA CI requires the internal-test DLL to be produced and runs + `CUDA_EP_Unittest.All` explicitly. Its XML report must contain the completed + wrapper test; an absent or skipped test fails the check. + This wrapper executes the internal CUDA-UT shared library and covers the fpA_intB / MatMulNBits groupwise GEMM tests under [onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc](../../../onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc) From bae016ab2dea7ab27c328975d2249c511b5ed1bd Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 16:12:33 +0000 Subject: [PATCH 06/19] Verify targeted provider-test builds before the full Windows build Separate configure from the full build and first build only the public provider-test target. Require fresh executable and CUDA test module outputs so the full build cannot mask a missing aggregate dependency. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .github/workflows/windows_cuda.yml | 40 +++++++++++++++++++++++---- docs/contrib_ops/cuda/matmul_nbits.md | 8 ++++-- 2 files changed, 39 insertions(+), 9 deletions(-) diff --git a/.github/workflows/windows_cuda.yml b/.github/workflows/windows_cuda.yml index be0287eb1f2eb..5ac6c569238a8 100644 --- a/.github/workflows/windows_cuda.yml +++ b/.github/workflows/windows_cuda.yml @@ -108,22 +108,50 @@ jobs: $buildDir = Join-Path ${{ runner.temp }} "build" echo "OnnxRuntimeBuildDirectory=$buildDir" >> $env:GITHUB_ENV - - name: Build and Clean Binaries + - name: Configure build working-directory: ${{ runner.temp }} run: | npm install -g typescript if ($lastExitCode -ne 0) { exit $lastExitCode } - # Execute the build process - python.exe ${{ github.workspace }}\tools\ci_build\build.py --update --build --config RelWithDebInfo --build_dir build --skip_submodule_sync --build_csharp --parallel --nvcc_threads 4 --flash_nvcc_threads 4 --use_binskim_compliant_compile_flags --cmake_generator "Visual Studio 17 2022" --build_shared_lib --build_wheel --build_java --use_cuda --cuda_home="$env:RUNNER_TEMP\v12.8" --enable_cuda_profiling --use_vcpkg --use_vcpkg_ms_internal_asset_cache --enable_transformers_tool_test --cmake_extra_defines onnxruntime_QUICK_BUILD=ON --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 --cmake_extra_defines onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON + python.exe ${{ github.workspace }}\tools\ci_build\build.py --update --config RelWithDebInfo --build_dir build --skip_submodule_sync --build_csharp --parallel --nvcc_threads 4 --flash_nvcc_threads 4 --use_binskim_compliant_compile_flags --cmake_generator "Visual Studio 17 2022" --build_shared_lib --build_wheel --build_java --use_cuda --cuda_home="$env:RUNNER_TEMP\v12.8" --enable_cuda_profiling --use_vcpkg --use_vcpkg_ms_internal_asset_cache --enable_transformers_tool_test --cmake_extra_defines onnxruntime_QUICK_BUILD=ON --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 --cmake_extra_defines onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON if ($lastExitCode -ne 0) { exit $lastExitCode } + shell: pwsh - $cudaTestModule = Join-Path $env:OnnxRuntimeBuildDirectory "RelWithDebInfo\RelWithDebInfo\onnxruntime_providers_cuda_ut.dll" - if (-not (Test-Path -LiteralPath $cudaTestModule -PathType Leaf)) { - throw "CUDA EP internal tests were requested, but the build did not produce $cudaTestModule" + - name: Verify targeted provider-test build + shell: pwsh + run: | + $buildDir = Join-Path $env:OnnxRuntimeBuildDirectory "RelWithDebInfo" + $binaryDir = Join-Path $buildDir "RelWithDebInfo" + $testBinaries = @("onnxruntime_provider_test.exe", "onnxruntime_providers_cuda_ut.dll") + foreach ($binary in $testBinaries) { + $path = Join-Path $binaryDir $binary + if (Test-Path -LiteralPath $path) { + throw "The targeted-build check requires fresh outputs, but $path already exists." + } + } + + cmake --build "$buildDir" --config RelWithDebInfo --target onnxruntime_provider_test --parallel + if ($LASTEXITCODE -ne 0) { + exit $LASTEXITCODE + } + + foreach ($binary in $testBinaries) { + $path = Join-Path $binaryDir $binary + if (-not (Test-Path -LiteralPath $path -PathType Leaf)) { + throw "Building only onnxruntime_provider_test did not produce $path" + } + } + + - name: Build and Clean Binaries + working-directory: ${{ runner.temp }} + run: | + python.exe ${{ github.workspace }}\tools\ci_build\build.py --build --config RelWithDebInfo --build_dir build --skip_submodule_sync --build_csharp --parallel --nvcc_threads 4 --flash_nvcc_threads 4 --use_binskim_compliant_compile_flags --cmake_generator "Visual Studio 17 2022" --build_shared_lib --build_wheel --build_java --use_cuda --cuda_home="$env:RUNNER_TEMP\v12.8" --enable_cuda_profiling --use_vcpkg --use_vcpkg_ms_internal_asset_cache --enable_transformers_tool_test --cmake_extra_defines onnxruntime_QUICK_BUILD=ON --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 --cmake_extra_defines onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON + if ($lastExitCode -ne 0) { + exit $lastExitCode } # Clean up the output directory before uploading artifacts diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index a6f11120f168b..2cc6e7dc3b0f0 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -492,9 +492,11 @@ present. `ComputeInternal` then: remains named `onnxruntime_provider_test.exe`, and the CTest name remains `onnxruntime_provider_test`. The default build also includes both artifacts. - Windows CUDA CI requires the internal-test DLL to be produced and runs - `CUDA_EP_Unittest.All` explicitly. Its XML report must contain the completed - wrapper test; an absent or skipped test fails the check. + Before the full build, Windows CUDA CI builds only `onnxruntime_provider_test` + and requires both the executable and internal-test DLL to be produced from + fresh outputs. It then runs `CUDA_EP_Unittest.All` explicitly on the GPU runner. + Its XML report must contain the completed wrapper test; an absent or skipped + test fails the check. This wrapper executes the internal CUDA-UT shared library and covers the fpA_intB / MatMulNBits groupwise GEMM tests under From e4d41c85e6e1dcacd100f25741b607e72846dc9b Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 17:09:45 +0000 Subject: [PATCH 07/19] Fix MSVC warnings exposed by CUDA internal-test builds Preserve size_t workspace sizes, use a float comparison literal, and widen the tactic-pruning threshold multiplication so the FLT_MAX no-best sentinel cannot overflow during constant folding. Keep warnings-as-errors enabled. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h | 3 ++- .../contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc | 4 ++-- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h b/onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h index b2d9c93b76188..eb21a43f0990a 100644 --- a/onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h +++ b/onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h @@ -62,7 +62,8 @@ inline int GetProfileTimedRuns(float probe_ms, float best_ms) { if (probe_ms * kMaxRuns <= kTimedBudgetMs) { return kMaxRuns; } - if (probe_ms > kPruneRatio * best_ms) { + // best_ms can be FLT_MAX when no tactic has been profiled yet. + if (probe_ms > kPruneRatio * static_cast(best_ms)) { return 0; } return std::max(1, static_cast(kTimedBudgetMs / probe_ms)); diff --git a/onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc b/onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc index f646bc3f39276..6d29da2f1e957 100644 --- a/onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc +++ b/onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc @@ -102,7 +102,7 @@ float compare(void* a, void* b, size_t size, float scale) { float total_diff = 0.f; float max_val = 0.f; int diff_count = 0; - float threshold = 1e-7; + float threshold = 1e-7f; for (size_t n = 0; n < size; ++n) { float va = static_cast(pa[n]); float vb = static_cast(pb[n]); @@ -462,7 +462,7 @@ class KernelTestFixture : public ::testing::Test { } #endif auto& gemm_runner = *runner; - int ws_bytes = gemm_runner.getWorkspaceSize(m_, n_, k_); + const size_t ws_bytes = gemm_runner.getWorkspaceSize(m_, n_, k_); CudaBuffer ws_buffer(ws_bytes); char* ws_ptr = reinterpret_cast(ws_buffer.data()); From 5b6b3eaefbfb4a2dfa53bfbdafc36fdc72d44591 Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 28 Sep 2026 10:44:09 +0000 Subject: [PATCH 08/19] WIP: export Windows provider-test symbols and check internal test count Checkpoint the current incomplete fix. Automatic exports do not cover linked static-library symbols, and the test-count check runs before Google Test selects tests. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_unittests.cmake | 12 ++++-------- .../providers/cuda/test_cases/cuda_test_provider.cc | 2 ++ 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index 37a7791592662..c0b1c0be69c28 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1490,15 +1490,11 @@ block() # without this the module fails to load with an undefined-symbol error. set_target_properties(${onnxruntime_provider_test_target} PROPERTIES ENABLE_EXPORTS 1) - # On Windows, ENABLE_EXPORTS makes CMake emit an import library (onnxruntime_provider_test.lib) - # for the exported symbols, but a MODULE library (onnxruntime_providers_cuda_ut, built via - # onnxruntime_add_shared_library_module) cannot have unresolved externals at *link* time the way - # a dlopen'd .so can on Linux. Since tests compiled into onnxruntime_providers_cuda_ut (e.g. the - # MatMulNBits end-to-end workspace test) call into InferenceSession symbols owned by this - # executable, link the module against that import library so those symbols resolve at link time. - # On Linux the runtime -rdynamic export path (above) is sufficient, so this is Windows-only. - # Note: onnxruntime_providers_cuda_ut only exists in the non-plugin CUDA-EP-internal-tests path. + # A Windows executable needs explicitly exported symbols before CMake can produce its import + # library. Export the symbols from its objects and linked static libraries, then link the CUDA + # test module against that import library. Linux uses the runtime -rdynamic export path above. if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) + set_target_properties(${onnxruntime_provider_test_target} PROPERTIES WINDOWS_EXPORT_ALL_SYMBOLS 1) target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE ${onnxruntime_provider_test_target}) endif() diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc index 01c7573b9de14..48d29809c9725 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc @@ -126,6 +126,8 @@ struct ProviderInfo_CUDA_TestImpl : ProviderInfo_CUDA { char* argv[] = {mock_exe_name, nullptr}; // char* argv[] = {mock_exe_name, "--gtest_filter=ReductionFunctionsTest.*", nullptr}; ::testing::InitGoogleTest(&argc, argv); + ORT_ENFORCE(::testing::UnitTest::GetInstance()->test_to_run_count() > 0, + "CUDA EP internal-test module contains no runnable tests."); ORT_ENFORCE(RUN_ALL_TESTS() == 0); } }; From 236645d6eab2d04e3c1a57804bc5e40ec1d8f7ae Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 28 Sep 2026 12:18:25 +0000 Subject: [PATCH 09/19] Limit Windows CUDA test host exports to required symbols Compile internal-test objects before the host link and derive targeted exports from module references and static host libraries. Keep tests in the DLL, add export-generation regressions and CI artifact checks, and verify successful internal tests after execution. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .github/workflows/windows_cuda.yml | 7 +- cmake/onnxruntime_test_exports.cmake | 57 ++++++ cmake/onnxruntime_unittests.cmake | 50 ++++-- docs/contrib_ops/cuda/matmul_nbits.md | 13 +- .../cuda/test_cases/cuda_test_provider.cc | 4 +- tools/ci_build/gen_test_exports.py | 104 +++++++++++ tools/ci_build/test_gen_test_exports.py | 170 ++++++++++++++++++ 7 files changed, 379 insertions(+), 26 deletions(-) create mode 100644 cmake/onnxruntime_test_exports.cmake create mode 100644 tools/ci_build/gen_test_exports.py create mode 100644 tools/ci_build/test_gen_test_exports.py diff --git a/.github/workflows/windows_cuda.yml b/.github/workflows/windows_cuda.yml index 5ac6c569238a8..cf252e2ae7135 100644 --- a/.github/workflows/windows_cuda.yml +++ b/.github/workflows/windows_cuda.yml @@ -124,9 +124,14 @@ jobs: - name: Verify targeted provider-test build shell: pwsh run: | + python.exe ${{ github.workspace }}\tools\ci_build\test_gen_test_exports.py + if ($LASTEXITCODE -ne 0) { + exit $LASTEXITCODE + } + $buildDir = Join-Path $env:OnnxRuntimeBuildDirectory "RelWithDebInfo" $binaryDir = Join-Path $buildDir "RelWithDebInfo" - $testBinaries = @("onnxruntime_provider_test.exe", "onnxruntime_providers_cuda_ut.dll") + $testBinaries = @("onnxruntime_provider_test.exe", "onnxruntime_provider_test.lib", "onnxruntime_providers_cuda_ut.dll") foreach ($binary in $testBinaries) { $path = Join-Path $binaryDir $binary if (Test-Path -LiteralPath $path) { diff --git a/cmake/onnxruntime_test_exports.cmake b/cmake/onnxruntime_test_exports.cmake new file mode 100644 index 0000000000000..c244b9525c124 --- /dev/null +++ b/cmake/onnxruntime_test_exports.cmake @@ -0,0 +1,57 @@ +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. + +function(onnxruntime_export_test_symbols host) + cmake_parse_arguments(EXPORTS "" "OBJECT_TARGET" "HOST_LIBS;MODULE_LIBS" ${ARGN}) + find_package(Python COMPONENTS Interpreter REQUIRED) + + set(module_files $) + set(host_files) + set(library_dependencies) + foreach(side HOST MODULE) + foreach(library IN LISTS EXPORTS_${side}_LIBS) + if(TARGET ${library}) + get_target_property(library_type ${library} TYPE) + if(library_type STREQUAL "STATIC_LIBRARY") + if(side STREQUAL "HOST") + list(APPEND host_files $) + else() + list(APPEND module_files $) + endif() + list(APPEND library_dependencies ${library}) + elseif(side STREQUAL "MODULE") + if(library_type STREQUAL "OBJECT_LIBRARY") + list(APPEND module_files $) + list(APPEND library_dependencies ${library}) + elseif(library_type STREQUAL "UNKNOWN_LIBRARY") + list(APPEND module_files $) + list(APPEND library_dependencies ${library}) + endif() + endif() + endif() + endforeach() + endforeach() + if(NOT host_files) + message(FATAL_ERROR "No static host libraries found for ${host}'s test exports") + endif() + list(REMOVE_DUPLICATES host_files) + list(REMOVE_DUPLICATES module_files) + list(REMOVE_DUPLICATES library_dependencies) + + set(export_dir "${CMAKE_CURRENT_BINARY_DIR}/${host}_exports/$") + file(GENERATE OUTPUT "${export_dir}/host.rsp" CONTENT "\"$\"\n") + file(GENERATE OUTPUT "${export_dir}/module.rsp" CONTENT "\"$\"\n") + set(export_script "${REPO_ROOT}/tools/ci_build/gen_test_exports.py") + set(export_file "${export_dir}/exports.def") + add_custom_command(OUTPUT "${export_file}" + COMMAND ${Python_EXECUTABLE} "${export_script}" + --linker "${CMAKE_LINKER}" + --host "${export_dir}/host.rsp" + --module "${export_dir}/module.rsp" + --output "${export_file}" + DEPENDS "${export_script}" "${export_dir}/host.rsp" "${export_dir}/module.rsp" + ${EXPORTS_OBJECT_TARGET} ${library_dependencies} ${host_files} ${module_files} + VERBATIM) + set_target_properties(${host} PROPERTIES ENABLE_EXPORTS ON WINDOWS_EXPORT_ALL_SYMBOLS OFF) + target_sources(${host} PRIVATE "${export_file}") +endfunction() diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index c0b1c0be69c28..907b881aa3b31 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1026,14 +1026,30 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_ "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_plugin_test_shims.cc") # onnxruntime_providers_cuda_ut is only for unittests. - onnxruntime_add_shared_library_module(onnxruntime_providers_cuda_ut ${onnxruntime_test_providers_cuda_ut_src} $) - config_cuda_provider_shared_module(onnxruntime_providers_cuda_ut) - target_compile_options(onnxruntime_providers_cuda_ut PRIVATE "$<$:SHELL:--threads \"${onnxruntime_NVCC_THREADS}\">") - onnxruntime_add_include_to_target(onnxruntime_providers_cuda_ut GTest::gtest GTest::gmock) - add_dependencies(onnxruntime_providers_cuda_ut onnxruntime_test_utils) - target_include_directories(onnxruntime_providers_cuda_ut PRIVATE ${ONNXRUNTIME_ROOT}/core/mickey) - target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE GTest::gtest GTest::gmock ${ONNXRUNTIME_MLAS_LIBS} - onnxruntime_test_utils ${PROTOBUF_LIB}) + set(onnxruntime_cuda_ut_compile_targets onnxruntime_providers_cuda_ut) + if (WIN32) + # Inspect the compiled tests before linking the host, without depending on the module's link. + onnxruntime_add_object_library(onnxruntime_providers_cuda_ut_objects ${onnxruntime_test_providers_cuda_ut_src}) + set(onnxruntime_cuda_ut_sources $) + list(APPEND onnxruntime_cuda_ut_compile_targets onnxruntime_providers_cuda_ut_objects) + else() + set(onnxruntime_cuda_ut_sources ${onnxruntime_test_providers_cuda_ut_src}) + endif() + onnxruntime_add_shared_library_module(onnxruntime_providers_cuda_ut ${onnxruntime_cuda_ut_sources} $) + foreach(cuda_ut_target IN LISTS onnxruntime_cuda_ut_compile_targets) + config_cuda_provider_shared_module(${cuda_ut_target}) + target_compile_options(${cuda_ut_target} PRIVATE "$<$:SHELL:--threads \"${onnxruntime_NVCC_THREADS}\">") + onnxruntime_add_include_to_target(${cuda_ut_target} GTest::gtest GTest::gmock) + add_dependencies(${cuda_ut_target} onnxruntime_test_utils) + target_include_directories(${cuda_ut_target} PRIVATE ${ONNXRUNTIME_ROOT}/core/mickey) + target_link_libraries(${cuda_ut_target} PRIVATE GTest::gtest GTest::gmock ${ONNXRUNTIME_MLAS_LIBS} + onnxruntime_test_utils ${PROTOBUF_LIB}) + if (MSVC) + # Cutlass code has an issue with warning C4100: 'magic': unreferenced formal parameter. + target_compile_options(${cuda_ut_target} PRIVATE "$<$:SHELL:--compiler-options /wd4100>" + "$<$>:/wd4100>") + endif() + endforeach() # Link architecture-specific OBJECT libraries (same as onnxruntime_providers_cuda). if(TARGET onnxruntime_providers_cuda_sm90_tma) target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE onnxruntime_providers_cuda_sm90_tma) @@ -1053,12 +1069,6 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_ if(TARGET onnxruntime_providers_cuda_llm_fp4) target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE onnxruntime_providers_cuda_llm_fp4) endif() - if (MSVC) - # Cutlass code has an issue with the following: - # warning C4100: 'magic': unreferenced formal parameter - target_compile_options(onnxruntime_providers_cuda_ut PRIVATE "$<$:SHELL:--compiler-options /wd4100>" - "$<$>:/wd4100>") - endif() list(APPEND onnxruntime_test_providers_runtime_dependencies onnxruntime_providers_cuda_ut) endif() @@ -1490,11 +1500,15 @@ block() # without this the module fails to load with an undefined-symbol error. set_target_properties(${onnxruntime_provider_test_target} PROPERTIES ENABLE_EXPORTS 1) - # A Windows executable needs explicitly exported symbols before CMake can produce its import - # library. Export the symbols from its objects and linked static libraries, then link the CUDA - # test module against that import library. Linux uses the runtime -rdynamic export path above. + # Export only host symbols referenced by the module, including definitions in static libraries. + # Exporting every provider-test symbol exceeds the Windows import-library limit. if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) - set_target_properties(${onnxruntime_provider_test_target} PROPERTIES WINDOWS_EXPORT_ALL_SYMBOLS 1) + include(onnxruntime_test_exports) + get_target_property(cuda_ut_link_libraries onnxruntime_providers_cuda_ut LINK_LIBRARIES) + onnxruntime_export_test_symbols(${onnxruntime_provider_test_target} + OBJECT_TARGET onnxruntime_providers_cuda_ut_objects + HOST_LIBS ${onnxruntime_provider_test_libs} + MODULE_LIBS onnxruntime_providers_cuda_obj ${cuda_ut_link_libraries}) target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE ${onnxruntime_provider_test_target}) endif() diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index 2cc6e7dc3b0f0..8da5e70ff0062 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -487,16 +487,19 @@ present. `ComputeInternal` then: cmake --build --config Release --target onnxruntime_provider_test ``` - On Windows, this is an aggregate build target: it builds the executable first, - then the module that links against the executable's import library. The executable + On Windows, this is an aggregate build target: it compiles the internal-test + objects, exports only the host symbols they require from the executable's static + libraries, then links the executable. The module links against its import library. + The tests remain in the module; there is no additional runtime interface. The executable remains named `onnxruntime_provider_test.exe`, and the CTest name remains `onnxruntime_provider_test`. The default build also includes both artifacts. Before the full build, Windows CUDA CI builds only `onnxruntime_provider_test` - and requires both the executable and internal-test DLL to be produced from + and requires the executable, its import library, and the internal-test DLL to be produced from fresh outputs. It then runs `CUDA_EP_Unittest.All` explicitly on the GPU runner. - Its XML report must contain the completed wrapper test; an absent or skipped - test fails the check. + Its XML report must contain the completed wrapper test. The wrapper also requires + at least one successful internal test after execution; an empty or entirely skipped + internal run fails the check. This wrapper executes the internal CUDA-UT shared library and covers the fpA_intB / MatMulNBits groupwise GEMM tests under diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc index 48d29809c9725..b1fc2c1bebff2 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc @@ -126,9 +126,9 @@ struct ProviderInfo_CUDA_TestImpl : ProviderInfo_CUDA { char* argv[] = {mock_exe_name, nullptr}; // char* argv[] = {mock_exe_name, "--gtest_filter=ReductionFunctionsTest.*", nullptr}; ::testing::InitGoogleTest(&argc, argv); - ORT_ENFORCE(::testing::UnitTest::GetInstance()->test_to_run_count() > 0, - "CUDA EP internal-test module contains no runnable tests."); ORT_ENFORCE(RUN_ALL_TESTS() == 0); + ORT_ENFORCE(::testing::UnitTest::GetInstance()->successful_test_count() > 0, + "CUDA EP internal-test module must execute at least one non-skipped test."); } }; ProviderInfo_CUDA_TestImpl g_test_info; diff --git a/tools/ci_build/gen_test_exports.py b/tools/ci_build/gen_test_exports.py new file mode 100644 index 0000000000000..f4165ffb79502 --- /dev/null +++ b/tools/ci_build/gen_test_exports.py @@ -0,0 +1,104 @@ +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. + +"""Export only symbols consumed by a test module from its Windows host executable.""" + +import argparse +import re +import shutil +import subprocess +import sys +import tempfile +from collections.abc import Iterable +from dataclasses import dataclass, field +from pathlib import Path + + +@dataclass +class Symbols: + undefined: set[str] = field(default_factory=set) + defined: set[str] = field(default_factory=set) + functions: set[str] = field(default_factory=set) + + +SYMBOL = re.compile( + r"^\s*[0-9A-F]+\s+([0-9A-F]+)\s+(UNDEF|SECT[0-9A-F]+|ABS)\s+" + r"(.+?)\s+(External|WeakExternal)\s+\|\s+(\S+)", + re.IGNORECASE, +) + + +def read_symbols(output: Iterable[str]) -> Symbols: + symbols = Symbols() + for line in output: + match = SYMBOL.match(line) + if not match: + continue + value, section, symbol_type, storage, name = match.groups() + if storage == "WeakExternal" or section == "ABS": + continue + if section == "UNDEF" and int(value, 16) == 0: + symbols.undefined.add(name) + else: + symbols.defined.add(name) + if "()" in symbol_type: + symbols.functions.add(name) + if not symbols.defined and not symbols.undefined: + raise ValueError( + "No COFF external symbols found; /GL objects or an unexpected linker output are not supported." + ) + return symbols + + +def select_exports(host: Symbols, module: Symbols) -> dict[str, bool]: + exports = {} + for reference in sorted(module.undefined - module.defined): + name = reference.removeprefix("__imp_") + if name not in host.defined or name in module.defined: + continue + is_data = name not in host.functions + if is_data and reference == name: + raise ValueError(f"Test module references host data {name} without __declspec(dllimport).") + exports[name] = is_data + if not exports: + raise ValueError("No host symbols are required by the test module; refusing to create an empty import library.") + if len(exports) > 65535: + raise ValueError(f"{len(exports)} required exports exceed the Windows limit of 65535.") + return exports + + +def write_exports(path: Path, exports: dict[str, bool]) -> None: + lines = ["EXPORTS"] + lines.extend(f' "{name}"' + (" DATA" if exports[name] else "") for name in sorted(exports)) + path.write_text("\n".join(lines) + "\n", encoding="utf-8") + + +def dump_symbols(linker: str, response_file: Path) -> Symbols: + # CUDA object symbol dumps can be large; do not retain the dump and split lines in memory. + with tempfile.TemporaryFile(mode="w+t") as output: + result = subprocess.run( + [linker, "/dump", "/nologo", "/symbols", f"@{response_file}"], + check=False, + stdout=output, + ) + output.seek(0) + if result.returncode: + shutil.copyfileobj(output, sys.stderr) + result.check_returncode() + return read_symbols(output) + + +def main() -> None: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--linker", required=True) + parser.add_argument("--host", required=True, type=Path) + parser.add_argument("--module", required=True, type=Path) + parser.add_argument("--output", required=True, type=Path) + args = parser.parse_args() + exports = select_exports(dump_symbols(args.linker, args.host), dump_symbols(args.linker, args.module)) + write_exports(args.output, exports) + print(f"Exporting {len(exports)} host symbols required by the CUDA internal-test module.") + + +if __name__ == "__main__": + main() diff --git a/tools/ci_build/test_gen_test_exports.py b/tools/ci_build/test_gen_test_exports.py new file mode 100644 index 0000000000000..2c43f36d33c61 --- /dev/null +++ b/tools/ci_build/test_gen_test_exports.py @@ -0,0 +1,170 @@ +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. + +import subprocess +import sys +import tempfile +import unittest +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parent)) + +from gen_test_exports import Symbols, read_symbols, select_exports, write_exports + + +class TestTestExports(unittest.TestCase): + def test_parse_coff_symbols(self): + symbols = read_symbols( + """ +Dump of file tests.obj +00A 00000000 UNDEF notype () External | ?Run@Session@@QEAAHXZ (public: int __cdecl Session::Run(void)) +00B 00000000 SECT2 notype () External | ?Local@@YAHXZ +00C 00000000 SECT10 notype External | ?data@@3HA +00D 00000004 UNDEF notype External | common_data +00E 00000000 UNDEF notype External | __imp_?data@@3HA +00F 00000000 SECT3 notype () Static | local_function +010 00000000 ABS notype External | @feat.00 +011 00000000 UNDEF notype WeakExternal | weak_fallback +""".splitlines() + ) + self.assertEqual(symbols.undefined, {"?Run@Session@@QEAAHXZ", "__imp_?data@@3HA"}) + self.assertEqual(symbols.defined, {"?Local@@YAHXZ", "?data@@3HA", "common_data"}) + self.assertEqual(symbols.functions, {"?Local@@YAHXZ"}) + + def test_static_library_definitions_and_duplicate_references(self): + module = read_symbols( + """ +001 00000000 UNDEF notype () External | required +002 00000000 UNDEF notype () External | required +003 00000000 UNDEF notype () External | module_local +004 00000000 UNDEF notype () External | system_function +Archive member name at 123: helpers.obj +005 00000000 SECT1 notype () External | module_local +""".splitlines() + ) + host = read_symbols( + """ +Archive member name at 12: session.obj +001 00000000 SECT1 notype () External | required +Archive member name at 34: helpers.obj +002 00000000 SECT2 notype () External | module_local +003 00000000 SECT3 notype () External | unrelated +""".splitlines() + ) + self.assertEqual(select_exports(host, module), {"required": False}) + + def test_dllimport_functions_and_data(self): + host = Symbols(defined={"function", "data"}, functions={"function"}) + module = Symbols(undefined={"__imp_function", "__imp_data"}) + self.assertEqual(select_exports(host, module), {"data": True, "function": False}) + + def test_direct_data_reference_is_rejected(self): + with self.assertRaisesRegex(ValueError, "without __declspec"): + select_exports(Symbols(defined={"data"}), Symbols(undefined={"data"})) + + def test_unrelated_symbols_do_not_exhaust_import_library(self): + functions = {f"unrelated_{i}" for i in range(70000)} | {"required"} + self.assertEqual( + select_exports(Symbols(defined=functions, functions=functions), Symbols(undefined={"required"})), + {"required": False}, + ) + + def test_excess_required_exports_are_rejected(self): + functions = {f"required_{i}" for i in range(65536)} + with self.assertRaisesRegex(ValueError, "exceed the Windows limit"): + select_exports(Symbols(defined=functions, functions=functions), Symbols(undefined=functions)) + + def test_empty_or_unrecognized_inputs_are_rejected(self): + with self.assertRaisesRegex(ValueError, "No COFF external symbols"): + read_symbols(["Microsoft linker: unknown object format"]) + with self.assertRaisesRegex(ValueError, "No host symbols"): + select_exports(Symbols(defined={"unrelated"}), Symbols(undefined={"required"})) + + def test_definition_file_is_sorted_and_preserves_decorated_names(self): + with tempfile.TemporaryDirectory() as directory: + output = Path(directory) / "exports.def" + write_exports(output, {"?Run@Session@@QEAAHXZ": False, "?Data@@3HA": True}) + self.assertEqual( + output.read_text(encoding="utf-8"), + 'EXPORTS\n "?Data@@3HA" DATA\n "?Run@Session@@QEAAHXZ"\n', + ) + + @unittest.skipUnless(sys.platform == "win32", "Requires the Windows MSVC toolchain") + def test_windows_targeted_build_and_incremental_exports(self): + repo_root = Path(__file__).resolve().parents[2] + with tempfile.TemporaryDirectory(prefix="ort test exports ") as directory: + source = Path(directory) + build = source / "build" + (source / "host.cc").write_text( + "namespace test_runtime {\nint required(int value) { return value + 1; }\n" + + "\n".join(f"int unused_{i}(int value) {{ return value + {i}; }}" for i in range(70000)) + + "\n}\n", + encoding="utf-8", + ) + (source / "main.cc").write_text("int main() { return 0; }\n", encoding="utf-8") + module = source / "module.cc" + module.write_text( + "namespace test_runtime { int required(int); }\n" + 'extern "C" __declspec(dllexport) int run_tests() { return test_runtime::required(41); }\n', + encoding="utf-8", + ) + (source / "CMakeLists.txt").write_text( + f""" +cmake_minimum_required(VERSION 3.28) +project(TestExports LANGUAGES CXX) +set(REPO_ROOT "{repo_root.as_posix()}") +include("${{REPO_ROOT}}/cmake/onnxruntime_test_exports.cmake") +add_library(runtime STATIC host.cc) +target_compile_options(runtime PRIVATE /bigobj) +add_library(test_objects OBJECT module.cc) +add_executable(provider_test_executable main.cc) +set_target_properties(provider_test_executable PROPERTIES OUTPUT_NAME provider_test) +target_link_libraries(provider_test_executable PRIVATE runtime) +onnxruntime_export_test_symbols(provider_test_executable OBJECT_TARGET test_objects HOST_LIBS runtime) +add_library(test_module MODULE $) +target_link_libraries(test_module PRIVATE provider_test_executable) +add_custom_target(provider_test ALL DEPENDS provider_test_executable test_module) +file(GENERATE OUTPUT "${{CMAKE_BINARY_DIR}}/linker.txt" CONTENT "${{CMAKE_LINKER}}") +""", + encoding="utf-8", + ) + subprocess.run( + ["cmake", "-S", str(source), "-B", str(build), "-G", "Visual Studio 17 2022", "-A", "x64"], + check=True, + timeout=120, + ) + command = ["cmake", "--build", str(build), "--config", "Release", "--target", "provider_test", "--parallel"] + subprocess.run(command, check=True, timeout=240) + output_dir = build / "Release" + for artifact in ("provider_test.exe", "provider_test.lib", "test_module.dll"): + self.assertTrue((output_dir / artifact).is_file(), artifact) + definition = build / "provider_test_executable_exports/Release/exports.def" + self.assertEqual( + definition.read_text(encoding="utf-8").splitlines(), + [ + "EXPORTS", + ' "?required@test_runtime@@YAHH@Z"', + ], + ) + linker = (build / "linker.txt").read_text(encoding="utf-8") + imports = subprocess.check_output( + [linker, "/dump", "/nologo", "/imports", str(output_dir / "test_module.dll")], + text=True, + timeout=30, + ) + self.assertIn("provider_test.exe", imports) + self.assertIn("?required@test_runtime@@YAHH@Z", imports) + + module.write_text( + "namespace test_runtime { int required(int); int unused_0(int); }\n" + 'extern "C" __declspec(dllexport) int run_tests() {\n' + " return test_runtime::required(41) + test_runtime::unused_0(1);\n}\n", + encoding="utf-8", + ) + subprocess.run(command, check=True, timeout=120) + self.assertEqual(len(definition.read_text(encoding="utf-8").splitlines()), 3) + self.assertIn("?unused_0@test_runtime@@YAHH@Z", definition.read_text(encoding="utf-8")) + + +if __name__ == "__main__": + unittest.main() From 664e0f3c2f08d3135b91f64d05762e59cf82db8f Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 28 Sep 2026 12:56:09 +0000 Subject: [PATCH 10/19] Fix CUDA test export helper include path Resolve the helper relative to the including CMake file instead of searching CMAKE_MODULE_PATH. Add a regression exercising the production include from a separate build directory with no module search path. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_unittests.cmake | 2 +- tools/ci_build/test_gen_test_exports.py | 22 ++++++++++++++++++++++ 2 files changed, 23 insertions(+), 1 deletion(-) diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index 907b881aa3b31..ec5a38ecad32d 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1503,7 +1503,7 @@ block() # Export only host symbols referenced by the module, including definitions in static libraries. # Exporting every provider-test symbol exceeds the Windows import-library limit. if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) - include(onnxruntime_test_exports) + include("${CMAKE_CURRENT_LIST_DIR}/onnxruntime_test_exports.cmake") get_target_property(cuda_ut_link_libraries onnxruntime_providers_cuda_ut LINK_LIBRARIES) onnxruntime_export_test_symbols(${onnxruntime_provider_test_target} OBJECT_TARGET onnxruntime_providers_cuda_ut_objects diff --git a/tools/ci_build/test_gen_test_exports.py b/tools/ci_build/test_gen_test_exports.py index 2c43f36d33c61..f38038bd2b3eb 100644 --- a/tools/ci_build/test_gen_test_exports.py +++ b/tools/ci_build/test_gen_test_exports.py @@ -1,6 +1,7 @@ # Copyright (c) Microsoft Corporation. All rights reserved. # Licensed under the MIT License. +import shutil import subprocess import sys import tempfile @@ -13,6 +14,27 @@ class TestTestExports(unittest.TestCase): + def test_production_include_without_module_search_path(self): + cmake_dir = Path(__file__).resolve().parents[2] / "cmake" + includes = [ + line.strip() + for line in (cmake_dir / "onnxruntime_unittests.cmake").read_text(encoding="utf-8").splitlines() + if line.strip().startswith("include(") and "onnxruntime_test_exports" in line + ] + self.assertEqual(len(includes), 1) + with tempfile.TemporaryDirectory(prefix="ort export include ") as directory: + source = Path(directory) + build = source / "build" + build.mkdir() + shutil.copyfile(cmake_dir / "onnxruntime_test_exports.cmake", source / "onnxruntime_test_exports.cmake") + script = source / "include.cmake" + script.write_text( + 'set(CMAKE_MODULE_PATH "")\n' + includes[0] + "\nif(NOT COMMAND onnxruntime_export_test_symbols)\n" + ' message(FATAL_ERROR "Test export helper was not loaded")\nendif()\n', + encoding="utf-8", + ) + subprocess.run(["cmake", "-P", str(script)], cwd=build, check=True, timeout=30) + def test_parse_coff_symbols(self): symbols = read_symbols( """ From af229d1c7597ca315ae417d56a2a7e4ef6585782 Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 28 Sep 2026 13:33:30 +0000 Subject: [PATCH 11/19] Force inclusion of selected Windows test host exports Emit a generated source containing /INCLUDE linker directives for the selected exports. This ensures MSVC extracts the required static-library members even when the host itself does not reference them, without exporting or forcing inclusion of unrelated symbols. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_test_exports.cmake | 7 +++++-- docs/contrib_ops/cuda/matmul_nbits.md | 4 +++- tools/ci_build/gen_test_exports.py | 8 ++++++++ tools/ci_build/test_gen_test_exports.py | 22 +++++++++++++++++++++- 4 files changed, 37 insertions(+), 4 deletions(-) diff --git a/cmake/onnxruntime_test_exports.cmake b/cmake/onnxruntime_test_exports.cmake index c244b9525c124..62a89dc86b23a 100644 --- a/cmake/onnxruntime_test_exports.cmake +++ b/cmake/onnxruntime_test_exports.cmake @@ -43,15 +43,18 @@ function(onnxruntime_export_test_symbols host) file(GENERATE OUTPUT "${export_dir}/module.rsp" CONTENT "\"$\"\n") set(export_script "${REPO_ROOT}/tools/ci_build/gen_test_exports.py") set(export_file "${export_dir}/exports.def") - add_custom_command(OUTPUT "${export_file}" + set(force_include_file "${export_dir}/force_include.cc") + add_custom_command(OUTPUT "${export_file}" "${force_include_file}" COMMAND ${Python_EXECUTABLE} "${export_script}" --linker "${CMAKE_LINKER}" --host "${export_dir}/host.rsp" --module "${export_dir}/module.rsp" --output "${export_file}" + --force-include "${force_include_file}" DEPENDS "${export_script}" "${export_dir}/host.rsp" "${export_dir}/module.rsp" ${EXPORTS_OBJECT_TARGET} ${library_dependencies} ${host_files} ${module_files} VERBATIM) set_target_properties(${host} PROPERTIES ENABLE_EXPORTS ON WINDOWS_EXPORT_ALL_SYMBOLS OFF) - target_sources(${host} PRIVATE "${export_file}") + set_source_files_properties("${force_include_file}" PROPERTIES SKIP_PRECOMPILE_HEADERS ON) + target_sources(${host} PRIVATE "${export_file}" "${force_include_file}") endfunction() diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index 8da5e70ff0062..d5955561e2651 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -489,7 +489,9 @@ present. `ComputeInternal` then: On Windows, this is an aggregate build target: it compiles the internal-test objects, exports only the host symbols they require from the executable's static - libraries, then links the executable. The module links against its import library. + libraries, then links the executable. Generated `/INCLUDE` directives ensure MSVC + extracts those symbols even when only the module references them. + The module links against its import library. The tests remain in the module; there is no additional runtime interface. The executable remains named `onnxruntime_provider_test.exe`, and the CTest name remains `onnxruntime_provider_test`. The default build also includes both artifacts. diff --git a/tools/ci_build/gen_test_exports.py b/tools/ci_build/gen_test_exports.py index f4165ffb79502..21f795b372c8e 100644 --- a/tools/ci_build/gen_test_exports.py +++ b/tools/ci_build/gen_test_exports.py @@ -73,6 +73,12 @@ def write_exports(path: Path, exports: dict[str, bool]) -> None: path.write_text("\n".join(lines) + "\n", encoding="utf-8") +def write_force_includes(path: Path, exports: dict[str, bool]) -> None: + # MSVC must extract exported symbols from static archives even when the host never references them. + lines = [f'#pragma comment(linker, "/include:{name}")' for name in sorted(exports)] + path.write_text("\n".join(lines) + "\n", encoding="utf-8") + + def dump_symbols(linker: str, response_file: Path) -> Symbols: # CUDA object symbol dumps can be large; do not retain the dump and split lines in memory. with tempfile.TemporaryFile(mode="w+t") as output: @@ -94,9 +100,11 @@ def main() -> None: parser.add_argument("--host", required=True, type=Path) parser.add_argument("--module", required=True, type=Path) parser.add_argument("--output", required=True, type=Path) + parser.add_argument("--force-include", required=True, type=Path) args = parser.parse_args() exports = select_exports(dump_symbols(args.linker, args.host), dump_symbols(args.linker, args.module)) write_exports(args.output, exports) + write_force_includes(args.force_include, exports) print(f"Exporting {len(exports)} host symbols required by the CUDA internal-test module.") diff --git a/tools/ci_build/test_gen_test_exports.py b/tools/ci_build/test_gen_test_exports.py index f38038bd2b3eb..c0cbcbb7a83ee 100644 --- a/tools/ci_build/test_gen_test_exports.py +++ b/tools/ci_build/test_gen_test_exports.py @@ -10,7 +10,7 @@ sys.path.insert(0, str(Path(__file__).resolve().parent)) -from gen_test_exports import Symbols, read_symbols, select_exports, write_exports +from gen_test_exports import Symbols, read_symbols, select_exports, write_exports, write_force_includes class TestTestExports(unittest.TestCase): @@ -111,6 +111,16 @@ def test_definition_file_is_sorted_and_preserves_decorated_names(self): 'EXPORTS\n "?Data@@3HA" DATA\n "?Run@Session@@QEAAHXZ"\n', ) + def test_force_includes_preserve_function_and_data_names(self): + with tempfile.TemporaryDirectory() as directory: + output = Path(directory) / "force_include.cc" + write_force_includes(output, {"?Run@Session@@QEAAHXZ": False, "?Data@@3HA": True}) + self.assertEqual( + output.read_text(encoding="utf-8"), + '#pragma comment(linker, "/include:?Data@@3HA")\n' + '#pragma comment(linker, "/include:?Run@Session@@QEAAHXZ")\n', + ) + @unittest.skipUnless(sys.platform == "win32", "Requires the Windows MSVC toolchain") def test_windows_targeted_build_and_incremental_exports(self): repo_root = Path(__file__).resolve().parents[2] @@ -161,6 +171,11 @@ def test_windows_targeted_build_and_incremental_exports(self): for artifact in ("provider_test.exe", "provider_test.lib", "test_module.dll"): self.assertTrue((output_dir / artifact).is_file(), artifact) definition = build / "provider_test_executable_exports/Release/exports.def" + force_includes = definition.with_name("force_include.cc") + self.assertEqual( + force_includes.read_text(encoding="utf-8"), + '#pragma comment(linker, "/include:?required@test_runtime@@YAHH@Z")\n', + ) self.assertEqual( definition.read_text(encoding="utf-8").splitlines(), [ @@ -186,6 +201,11 @@ def test_windows_targeted_build_and_incremental_exports(self): subprocess.run(command, check=True, timeout=120) self.assertEqual(len(definition.read_text(encoding="utf-8").splitlines()), 3) self.assertIn("?unused_0@test_runtime@@YAHH@Z", definition.read_text(encoding="utf-8")) + self.assertEqual(len(force_includes.read_text(encoding="utf-8").splitlines()), 2) + self.assertIn( + '#pragma comment(linker, "/include:?unused_0@test_runtime@@YAHH@Z")', + force_includes.read_text(encoding="utf-8"), + ) if __name__ == "__main__": From 6f6c9d57430677ec9ecc1805a80165b09194a30e Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 28 Sep 2026 14:01:15 +0000 Subject: [PATCH 12/19] Emit unquoted COFF names in Windows test exports Match CMake's DEF writer by emitting decorated symbol names directly. Add native linker library diagnostics and dump the generated DEF and EXP symbol table on fixture failure so export-name mismatches can be distinguished from archive resolution failures. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- tools/ci_build/gen_test_exports.py | 2 +- tools/ci_build/test_gen_test_exports.py | 18 +++++++++++++----- 2 files changed, 14 insertions(+), 6 deletions(-) diff --git a/tools/ci_build/gen_test_exports.py b/tools/ci_build/gen_test_exports.py index 21f795b372c8e..64ffc45fdecb8 100644 --- a/tools/ci_build/gen_test_exports.py +++ b/tools/ci_build/gen_test_exports.py @@ -69,7 +69,7 @@ def select_exports(host: Symbols, module: Symbols) -> dict[str, bool]: def write_exports(path: Path, exports: dict[str, bool]) -> None: lines = ["EXPORTS"] - lines.extend(f' "{name}"' + (" DATA" if exports[name] else "") for name in sorted(exports)) + lines.extend(f" {name}" + (" DATA" if exports[name] else "") for name in sorted(exports)) path.write_text("\n".join(lines) + "\n", encoding="utf-8") diff --git a/tools/ci_build/test_gen_test_exports.py b/tools/ci_build/test_gen_test_exports.py index c0cbcbb7a83ee..e6a38eba74094 100644 --- a/tools/ci_build/test_gen_test_exports.py +++ b/tools/ci_build/test_gen_test_exports.py @@ -108,7 +108,7 @@ def test_definition_file_is_sorted_and_preserves_decorated_names(self): write_exports(output, {"?Run@Session@@QEAAHXZ": False, "?Data@@3HA": True}) self.assertEqual( output.read_text(encoding="utf-8"), - 'EXPORTS\n "?Data@@3HA" DATA\n "?Run@Session@@QEAAHXZ"\n', + "EXPORTS\n ?Data@@3HA DATA\n ?Run@Session@@QEAAHXZ\n", ) def test_force_includes_preserve_function_and_data_names(self): @@ -152,6 +152,7 @@ def test_windows_targeted_build_and_incremental_exports(self): add_executable(provider_test_executable main.cc) set_target_properties(provider_test_executable PROPERTIES OUTPUT_NAME provider_test) target_link_libraries(provider_test_executable PRIVATE runtime) +target_link_options(provider_test_executable PRIVATE /VERBOSE:LIB) onnxruntime_export_test_symbols(provider_test_executable OBJECT_TARGET test_objects HOST_LIBS runtime) add_library(test_module MODULE $) target_link_libraries(test_module PRIVATE provider_test_executable) @@ -166,11 +167,19 @@ def test_windows_targeted_build_and_incremental_exports(self): timeout=120, ) command = ["cmake", "--build", str(build), "--config", "Release", "--target", "provider_test", "--parallel"] - subprocess.run(command, check=True, timeout=240) output_dir = build / "Release" + definition = build / "provider_test_executable_exports/Release/exports.def" + linker = (build / "linker.txt").read_text(encoding="utf-8") + result = subprocess.run(command, check=False, timeout=240) + if result.returncode: + if definition.is_file(): + print(definition.read_text(encoding="utf-8"), flush=True) + export_object = output_dir / "provider_test.exp" + if export_object.is_file(): + subprocess.run([linker, "/dump", "/nologo", "/symbols", str(export_object)], check=True, timeout=30) + result.check_returncode() for artifact in ("provider_test.exe", "provider_test.lib", "test_module.dll"): self.assertTrue((output_dir / artifact).is_file(), artifact) - definition = build / "provider_test_executable_exports/Release/exports.def" force_includes = definition.with_name("force_include.cc") self.assertEqual( force_includes.read_text(encoding="utf-8"), @@ -180,10 +189,9 @@ def test_windows_targeted_build_and_incremental_exports(self): definition.read_text(encoding="utf-8").splitlines(), [ "EXPORTS", - ' "?required@test_runtime@@YAHH@Z"', + " ?required@test_runtime@@YAHH@Z", ], ) - linker = (build / "linker.txt").read_text(encoding="utf-8") imports = subprocess.check_output( [linker, "/dump", "/nologo", "/imports", str(output_dir / "test_module.dll")], text=True, From 263ce0e573b0a9ad0f9829e9c1806a4ac78616b3 Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 28 Sep 2026 15:18:38 +0000 Subject: [PATCH 13/19] Fix CUDA internal-test host/provider boundaries Keep tests in the module and call provider Node/Tensor implementations through statically linked opaque-pointer test adapters. Initialize the module's C++ API consistently, import host test utilities instead of duplicating ort_env-dependent objects, and link Windows ONNX protobuf definitions directly. Collect exports from actual host link inputs, including imported and object libraries. Extend the MSVC fixture to exercise class/struct adapters, separate API initialization modes, DLL loading, and incremental exports. Validation: rebuilt the actual Linux CUDA test module and passed all 126 selected tests with no skips. Generator regressions and Windows COFF linking passed locally; full native Windows validation remains for CI. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_test_exports.cmake | 27 +++--- cmake/onnxruntime_unittests.cmake | 11 ++- docs/contrib_ops/cuda/matmul_nbits.md | 7 ++ .../cuda_external_data_loader_test.cc | 5 +- .../cuda/test_cases/cuda_test_bridge.cc | 60 ++++++++++++++ .../cuda/test_cases/cuda_test_bridge.h | 49 +++++++++++ .../cuda/test_cases/cuda_test_provider.cc | 1 + ...query_attention_workspace_estimate_test.cc | 9 +- .../matmul_nbits_e2e_workspace_test.cc | 37 +++++---- ...acked_attention_workspace_estimate_test.cc | 5 +- tools/ci_build/test_gen_test_exports.py | 83 +++++++++++++++++-- 11 files changed, 242 insertions(+), 52 deletions(-) create mode 100644 onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.cc create mode 100644 onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.h diff --git a/cmake/onnxruntime_test_exports.cmake b/cmake/onnxruntime_test_exports.cmake index 62a89dc86b23a..3214b901ed2b9 100644 --- a/cmake/onnxruntime_test_exports.cmake +++ b/cmake/onnxruntime_test_exports.cmake @@ -12,22 +12,19 @@ function(onnxruntime_export_test_symbols host) foreach(library IN LISTS EXPORTS_${side}_LIBS) if(TARGET ${library}) get_target_property(library_type ${library} TYPE) - if(library_type STREQUAL "STATIC_LIBRARY") - if(side STREQUAL "HOST") - list(APPEND host_files $) - else() - list(APPEND module_files $) - endif() - list(APPEND library_dependencies ${library}) - elseif(side STREQUAL "MODULE") - if(library_type STREQUAL "OBJECT_LIBRARY") - list(APPEND module_files $) - list(APPEND library_dependencies ${library}) - elseif(library_type STREQUAL "UNKNOWN_LIBRARY") - list(APPEND module_files $) - list(APPEND library_dependencies ${library}) - endif() + if(library_type STREQUAL "STATIC_LIBRARY" OR library_type STREQUAL "UNKNOWN_LIBRARY") + set(library_files $) + elseif(library_type STREQUAL "OBJECT_LIBRARY") + set(library_files $) + else() + continue() endif() + if(side STREQUAL "HOST") + list(APPEND host_files ${library_files}) + else() + list(APPEND module_files ${library_files}) + endif() + list(APPEND library_dependencies ${library}) endif() endforeach() endforeach() diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index ec5a38ecad32d..0a51adf0be172 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1039,11 +1039,15 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_ foreach(cuda_ut_target IN LISTS onnxruntime_cuda_ut_compile_targets) config_cuda_provider_shared_module(${cuda_ut_target}) target_compile_options(${cuda_ut_target} PRIVATE "$<$:SHELL:--threads \"${onnxruntime_NVCC_THREADS}\">") - onnxruntime_add_include_to_target(${cuda_ut_target} GTest::gtest GTest::gmock) + onnxruntime_add_include_to_target(${cuda_ut_target} GTest::gtest GTest::gmock onnxruntime_test_utils) add_dependencies(${cuda_ut_target} onnxruntime_test_utils) target_include_directories(${cuda_ut_target} PRIVATE ${ONNXRUNTIME_ROOT}/core/mickey) + target_compile_definitions(${cuda_ut_target} PRIVATE ORT_API_MANUAL_INIT) target_link_libraries(${cuda_ut_target} PRIVATE GTest::gtest GTest::gmock ${ONNXRUNTIME_MLAS_LIBS} - onnxruntime_test_utils ${PROTOBUF_LIB}) + ${PROTOBUF_LIB}) + if(WIN32) + target_link_libraries(${cuda_ut_target} PRIVATE onnx_proto) + endif() if (MSVC) # Cutlass code has an issue with warning C4100: 'magic': unreferenced formal parameter. target_compile_options(${cuda_ut_target} PRIVATE "$<$:SHELL:--compiler-options /wd4100>" @@ -1504,10 +1508,11 @@ block() # Exporting every provider-test symbol exceeds the Windows import-library limit. if (WIN32 AND TARGET onnxruntime_providers_cuda_ut) include("${CMAKE_CURRENT_LIST_DIR}/onnxruntime_test_exports.cmake") + get_target_property(provider_test_link_libraries ${onnxruntime_provider_test_target} LINK_LIBRARIES) get_target_property(cuda_ut_link_libraries onnxruntime_providers_cuda_ut LINK_LIBRARIES) onnxruntime_export_test_symbols(${onnxruntime_provider_test_target} OBJECT_TARGET onnxruntime_providers_cuda_ut_objects - HOST_LIBS ${onnxruntime_provider_test_libs} + HOST_LIBS ${provider_test_link_libraries} MODULE_LIBS onnxruntime_providers_cuda_obj ${cuda_ut_link_libraries}) target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE ${onnxruntime_provider_test_target}) endif() diff --git a/docs/contrib_ops/cuda/matmul_nbits.md b/docs/contrib_ops/cuda/matmul_nbits.md index d5955561e2651..c3339378d4eb0 100644 --- a/docs/contrib_ops/cuda/matmul_nbits.md +++ b/docs/contrib_ops/cuda/matmul_nbits.md @@ -496,6 +496,13 @@ present. `ComputeInternal` then: remains named `onnxruntime_provider_test.exe`, and the CTest name remains `onnxruntime_provider_test`. The default build also includes both artifacts. + Tests using core `Node` and `Tensor` objects call provider implementations through + statically linked test adapters with opaque, borrowed pointers. The adapters use + the existing provider-host accessors, avoiding the distinct core/provider C++ types + in cross-translation-unit signatures. The module uses manual C++ API initialization + consistently and imports host test utilities instead of linking a second copy that + depends on the executable's `ort_env` global. On Windows it links `onnx_proto` directly. + Before the full build, Windows CUDA CI builds only `onnxruntime_provider_test` and requires the executable, its import library, and the internal-test DLL to be produced from fresh outputs. It then runs `CUDA_EP_Unittest.All` explicitly on the GPU runner. diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc index 9c688e2788561..cd496c6f90959 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc @@ -25,6 +25,7 @@ #include "core/session/onnxruntime_session_options_config_keys.h" #include "cuda_runtime.h" #include "gtest/gtest.h" +#include "test/providers/cuda/test_cases/cuda_test_bridge.h" #include "test/test_environment.h" #include "test/unittest_util/framework_test_utils.h" #include "test/util/include/default_providers.h" @@ -256,8 +257,8 @@ TEST(CudaExternalDataLoaderTest, NormalizesBoolWithPinnedAndPageableFallback) { } Tensor tensor(DataTypeImpl::GetType(), TensorShape({static_cast(input.size())}), *allocator); - ASSERT_STATUS_OK(loader->LoadTensor( - Env::Default(), path, kFilePrefixSize, input.size(), tensor)); + ASSERT_STATUS_OK(LoadCudaExternalDataForTest( + loader.get(), Env::Default(), path, kFilePrefixSize, input.size(), &tensor)); std::array output{}; ASSERT_EQ(cudaSuccess, cudaMemcpy( output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.cc new file mode 100644 index 0000000000000..c2ada1734d96d --- /dev/null +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.cc @@ -0,0 +1,60 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "core/providers/shared_library/provider_api.h" +#include "core/providers/cuda/cuda_external_data_loader.h" +#include "test/providers/cuda/test_cases/cuda_test_bridge.h" + +namespace onnxruntime::test { + +#if !defined(USE_CUDA_MINIMAL) && !defined(DISABLE_CONTRIB_OPS) && !defined(BUILD_CUDA_EP_AS_PLUGIN) +std::optional EstimateGroupQueryAttentionWorkspaceForTest( + const void* node, gsl::span input_shapes, + const cudaDeviceProp& device_prop, const AttentionKernelOptions& kernel_options, + bool head_sink_is_constant_initializer) { + return contrib::cuda::EstimateGroupQueryAttentionWorkspace( + *static_cast(node), input_shapes, device_prop, kernel_options, + head_sink_is_constant_initializer); +} + +std::optional EstimatePackedAttentionWorkspaceForTest( + const void* node, gsl::span input_shapes, + const cudaDeviceProp& device_prop, const AttentionKernelOptions& kernel_options) { + return contrib::cuda::EstimatePackedAttentionWorkspace( + *static_cast(node), input_shapes, device_prop, kernel_options); +} +#endif + +#if !defined(DISABLE_CONTRIB_OPS) && defined(USE_FPA_INTB_GEMM) && USE_FPA_INTB_GEMM +std::optional EstimateMatMulNBitsMemoryForTest( + const void* node, const cudaDeviceProp& device_prop, + contrib::cuda::MatMulNBitsMemoryEstimateOptions options) { + return contrib::cuda::EstimateMatMulNBitsMemory(*static_cast(node), device_prop, options); +} + +std::optional EstimateMatMulNBitsMemoryForTest( + const void* node, gsl::span input_shape, const cudaDeviceProp& device_prop, + contrib::cuda::MatMulNBitsMemoryEstimateOptions options) { + return contrib::cuda::EstimateMatMulNBitsMemory( + *static_cast(node), input_shape, device_prop, options); +} + +std::optional EstimateMatMulNBitsWorkspaceForTest( + const void* node, const cudaDeviceProp& device_prop) { + return contrib::cuda::EstimateMatMulNBitsWorkspace(*static_cast(node), device_prop); +} + +std::optional EstimateMatMulNBitsWorkspaceForTest( + const void* node, gsl::span input_shape, const cudaDeviceProp& device_prop) { + return contrib::cuda::EstimateMatMulNBitsWorkspace(*static_cast(node), input_shape, device_prop); +} +#endif + +common::Status LoadCudaExternalDataForTest( + const void* loader, const Env& env, const std::filesystem::path& path, + FileOffsetType offset, SafeInt length, void* tensor) { + return static_cast(loader)->LoadTensor( + env, path, offset, length, *static_cast(tensor)); +} + +} // namespace onnxruntime::test diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.h b/onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.h new file mode 100644 index 0000000000000..7d96e4121f896 --- /dev/null +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_test_bridge.h @@ -0,0 +1,49 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include "core/framework/external_data_loader.h" + +// Include after the caller's core or provider headers have declared Node. +#if !defined(USE_CUDA_MINIMAL) && !defined(DISABLE_CONTRIB_OPS) && !defined(BUILD_CUDA_EP_AS_PLUGIN) +#include "contrib_ops/cuda/bert/group_query_attention_workspace_estimate.h" +#include "contrib_ops/cuda/bert/packed_attention_workspace_estimate.h" +#endif +#if !defined(DISABLE_CONTRIB_OPS) && defined(USE_FPA_INTB_GEMM) && USE_FPA_INTB_GEMM +#include "contrib_ops/cuda/quantization/matmul_nbits_workspace_estimate.h" +#endif + +namespace onnxruntime::test { + +// Node and Tensor are borrowed core objects. The provider-side accessors forward +// their addresses back to ProviderHost; neither type crosses this link boundary. +#if !defined(USE_CUDA_MINIMAL) && !defined(DISABLE_CONTRIB_OPS) && !defined(BUILD_CUDA_EP_AS_PLUGIN) +std::optional EstimateGroupQueryAttentionWorkspaceForTest( + const void* node, gsl::span input_shapes, + const cudaDeviceProp& device_prop, const AttentionKernelOptions& kernel_options, + bool head_sink_is_constant_initializer = false); + +std::optional EstimatePackedAttentionWorkspaceForTest( + const void* node, gsl::span input_shapes, + const cudaDeviceProp& device_prop, const AttentionKernelOptions& kernel_options); +#endif + +#if !defined(DISABLE_CONTRIB_OPS) && defined(USE_FPA_INTB_GEMM) && USE_FPA_INTB_GEMM +std::optional EstimateMatMulNBitsMemoryForTest( + const void* node, const cudaDeviceProp& device_prop, + contrib::cuda::MatMulNBitsMemoryEstimateOptions options = {}); +std::optional EstimateMatMulNBitsMemoryForTest( + const void* node, gsl::span input_shape, const cudaDeviceProp& device_prop, + contrib::cuda::MatMulNBitsMemoryEstimateOptions options = {}); +std::optional EstimateMatMulNBitsWorkspaceForTest( + const void* node, const cudaDeviceProp& device_prop); +std::optional EstimateMatMulNBitsWorkspaceForTest( + const void* node, gsl::span input_shape, const cudaDeviceProp& device_prop); +#endif + +common::Status LoadCudaExternalDataForTest( + const void* loader, const Env& env, const std::filesystem::path& path, + FileOffsetType offset, SafeInt length, void* tensor); + +} // namespace onnxruntime::test diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc index b1fc2c1bebff2..12e78834987c5 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc @@ -137,6 +137,7 @@ struct CUDA_Test_Provider : Provider { void* GetInfo() override { return &g_test_info; } void Initialize() override { + InitProviderOrtApi(); InitializeRegistry(); } diff --git a/onnxruntime/test/providers/cuda/test_cases/group_query_attention_workspace_estimate_test.cc b/onnxruntime/test/providers/cuda/test_cases/group_query_attention_workspace_estimate_test.cc index 95d3ef0304477..451a638d81f29 100644 --- a/onnxruntime/test/providers/cuda/test_cases/group_query_attention_workspace_estimate_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/group_query_attention_workspace_estimate_test.cc @@ -22,6 +22,7 @@ #include "core/session/onnxruntime_session_options_config_keys.h" #include "contrib_ops/cpu/bert/attention_common.h" #include "contrib_ops/cuda/bert/group_query_attention_workspace_estimate.h" +#include "test/providers/cuda/test_cases/cuda_test_bridge.h" #include "test/test_environment.h" #include "test/util/include/asserts.h" #include "test/util/include/inference_session_wrapper.h" @@ -147,8 +148,8 @@ std::optional EstimateFromNode( const std::vector inputs{&query, &key, &value, &past_key}; const std::vector outputs; Node node{"gqa", "GroupQueryAttention", "", inputs, outputs, &attributes, kMSDomain}; - return EstimateGroupQueryAttentionWorkspace( - node, input_shapes, Device(), options); + return EstimateGroupQueryAttentionWorkspaceForTest( + &node, input_shapes, Device(), options); } void SetValueInfo(ONNX_NAMESPACE::ValueInfoProto& value_info, @@ -297,8 +298,8 @@ TEST(GroupQueryAttentionWorkspaceEstimateTest, GetCapabilityBudgetUsesLevel1Esti ASSERT_EQ(node->GetExecutionProviderType(), kCudaExecutionProvider); auto shapes = SeparateShapes(); shapes[11] = Known({8}); - estimate = EstimateGroupQueryAttentionWorkspace( - *node, gsl::make_span(shapes), cuda_ep->GetDeviceProp(), + estimate = EstimateGroupQueryAttentionWorkspaceForTest( + node, gsl::make_span(shapes), cuda_ep->GetDeviceProp(), *cuda_ep->GetAttentionKernelOptions(), /*head_sink_is_constant_initializer=*/true); ASSERT_TRUE(estimate.has_value()); diff --git a/onnxruntime/test/providers/cuda/test_cases/matmul_nbits_e2e_workspace_test.cc b/onnxruntime/test/providers/cuda/test_cases/matmul_nbits_e2e_workspace_test.cc index b96023e5443e8..82360ceee93dd 100644 --- a/onnxruntime/test/providers/cuda/test_cases/matmul_nbits_e2e_workspace_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/matmul_nbits_e2e_workspace_test.cc @@ -23,7 +23,7 @@ // This translation unit runs a real InferenceSession, so it includes the core framework headers. // Those cannot coexist with the CUDA-provider (shared-provider bridge) headers in one TU, so the two // provider-world pieces it needs (the Level-1 estimate and the runtime probe) are reached through -// slim, bridge-free declarations. It lives in the CUDA-only unit-test module because that is the only +// statically linked test adapters with opaque Node pointers. It lives in the CUDA-only unit-test module because that is the only // place these provider-internal symbols are linkable. Requires a real CUDA device; skips otherwise. #include "gtest/gtest.h" @@ -56,6 +56,7 @@ #include "contrib_ops/cuda/quantization/matmul_nbits_workspace_estimate.h" +#include "test/providers/cuda/test_cases/cuda_test_bridge.h" #include "test/providers/cuda/test_cases/matmul_nbits_workspace_test_probe.h" #include "test/test_environment.h" #include "test/unittest_util/framework_test_utils.h" @@ -309,7 +310,7 @@ std::optional EstimateWorkspaceFromGraphProtoShape( device_prop.major = 8; device_prop.minor = 0; device_prop.multiProcessorCount = 100; - return onnxruntime::contrib::cuda::EstimateMatMulNBitsWorkspace(node, device_prop); + return EstimateMatMulNBitsWorkspaceForTest(&node, device_prop); } } // namespace @@ -352,8 +353,8 @@ TEST(MatMulNBitsWorkspace, GetCapabilityBudgetChargesLazyProfileScratch) { ASSERT_NE(mm_node, nullptr); EXPECT_EQ(mm_node->GetExecutionProviderType(), kCudaExecutionProvider); - const auto estimate = onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory( - *mm_node, cuda_ep->GetDeviceProp(), + const auto estimate = EstimateMatMulNBitsMemoryForTest( + mm_node, cuda_ep->GetDeviceProp(), {/*fpa_intb_gemm=*/std::string_view{"1"}, /*profile_m=*/std::string_view{"1"}}); ASSERT_TRUE(estimate.has_value()); ASSERT_TRUE(estimate->runtime_workspace_bytes.has_value()); @@ -426,15 +427,15 @@ TEST(MatMulNBitsWorkspace, MaxShapeBudgetChargesMissingSmallerProfileBucket) { EXPECT_EQ(mm_node->GetExecutionProviderType(), kCudaExecutionProvider); const TensorShape planning_shape({256, kE2eK}); - const auto estimate = onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory( - *mm_node, planning_shape.GetDims(), cuda_ep->GetDeviceProp(), + const auto estimate = EstimateMatMulNBitsMemoryForTest( + mm_node, planning_shape.GetDims(), cuda_ep->GetDeviceProp(), {/*fpa_intb_gemm=*/std::string_view{"1"}, /*profile_m=*/std::string_view{"256"}, /*input_shape_is_upper_bound=*/true}); ASSERT_TRUE(estimate.has_value()); const TensorShape missing_bucket_shape({128, kE2eK}); - const auto missing_bucket_estimate = onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory( - *mm_node, missing_bucket_shape.GetDims(), cuda_ep->GetDeviceProp(), + const auto missing_bucket_estimate = EstimateMatMulNBitsMemoryForTest( + mm_node, missing_bucket_shape.GetDims(), cuda_ep->GetDeviceProp(), {/*fpa_intb_gemm=*/std::string_view{"1"}, /*profile_m=*/std::string_view{"256"}}); ASSERT_TRUE(missing_bucket_estimate.has_value()); @@ -603,12 +604,12 @@ TEST(MatMulNBitsWorkspace, EndToEndWorkspaceAgreement) { // ---- Level 1: the estimator function GetCapability() uses, invoked directly on the node + device // properties and the concrete estimation shape for this run. ---- const std::optional level1 = - onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory( - *mm_node, positive_a_shape.GetDims(), cuda_ep->GetDeviceProp()); + EstimateMatMulNBitsMemoryForTest( + mm_node, positive_a_shape.GetDims(), cuda_ep->GetDeviceProp()); ASSERT_TRUE(level1.has_value()) << "Level-1 estimate returned nullopt for an eligible node."; const auto static_shape_with_bound_option = - onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory( - *mm_node, positive_a_shape.GetDims(), cuda_ep->GetDeviceProp(), + EstimateMatMulNBitsMemoryForTest( + mm_node, positive_a_shape.GetDims(), cuda_ep->GetDeviceProp(), {/*fpa_intb_gemm=*/std::string_view{"1"}, /*profile_m=*/std::string_view{"256"}, /*input_shape_is_upper_bound=*/true}); @@ -734,8 +735,8 @@ TEST(MatMulNBitsWorkspace, EndToEndWorkspaceAgreement) { // ---- Empty-output parity on the same dynamic-shape session and kernel. ---- // Level 1 knows that m == 0 and must return a known zero, including for native SM90. const std::optional zero_m_level1 = - onnxruntime::contrib::cuda::EstimateMatMulNBitsWorkspace( - *mm_node, zero_m_a_shape.GetDims(), cuda_ep->GetDeviceProp()); + EstimateMatMulNBitsWorkspaceForTest( + mm_node, zero_m_a_shape.GetDims(), cuda_ep->GetDeviceProp()); ASSERT_TRUE(zero_m_level1.has_value()); EXPECT_EQ(*zero_m_level1, 0u); @@ -835,7 +836,7 @@ TEST(MatMulNBitsWorkspace, FixedShapeViaFreeDimensionOverride) { // ---- Level 1: estimator reads the (now-overridden) node shape directly. ---- const std::optional level1 = - onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory(*mm_node, cuda_ep->GetDeviceProp()); + EstimateMatMulNBitsMemoryForTest(mm_node, cuda_ep->GetDeviceProp()); ASSERT_TRUE(level1.has_value()) << "Level-1 estimate returned nullopt after the fixed override made the shape static."; ASSERT_TRUE(level1->runtime_workspace_bytes.has_value()); @@ -936,7 +937,7 @@ TEST(MatMulNBitsWorkspace, DynamicShapeNoOverrideFallsBack) { // ---- Level 1: dynamic leading dim leaves runtime workspace unknown, while shape-independent // prepack memory remains estimable. ---- const std::optional level1 = - onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory(*mm_node, cuda_ep->GetDeviceProp()); + EstimateMatMulNBitsMemoryForTest(mm_node, cuda_ep->GetDeviceProp()); ASSERT_TRUE(level1.has_value()); EXPECT_FALSE(level1->runtime_workspace_bytes.has_value()) << "Level-1 runtime workspace must be unknown for a dynamic (symbolic) leading dim."; @@ -948,8 +949,8 @@ TEST(MatMulNBitsWorkspace, DynamicShapeNoOverrideFallsBack) { // modifying its canonical shape metadata. const TensorShape max_input_shape({256, kE2eK}); const std::optional bounded_level1 = - onnxruntime::contrib::cuda::EstimateMatMulNBitsMemory( - *mm_node, max_input_shape.GetDims(), cuda_ep->GetDeviceProp()); + EstimateMatMulNBitsMemoryForTest( + mm_node, max_input_shape.GetDims(), cuda_ep->GetDeviceProp()); ASSERT_TRUE(bounded_level1.has_value()); EXPECT_TRUE(bounded_level1->runtime_workspace_bytes.has_value()); diff --git a/onnxruntime/test/providers/cuda/test_cases/packed_attention_workspace_estimate_test.cc b/onnxruntime/test/providers/cuda/test_cases/packed_attention_workspace_estimate_test.cc index 9430dc7b4793f..786db96d14bff 100644 --- a/onnxruntime/test/providers/cuda/test_cases/packed_attention_workspace_estimate_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/packed_attention_workspace_estimate_test.cc @@ -21,6 +21,7 @@ #include "core/providers/cuda/cuda_execution_provider_info.h" #include "contrib_ops/cpu/bert/attention_common.h" #include "contrib_ops/cuda/bert/packed_attention_workspace_estimate.h" +#include "test/providers/cuda/test_cases/cuda_test_bridge.h" #include "test/test_environment.h" #include "test/util/include/asserts.h" #include "test/util/include/inference_session_wrapper.h" @@ -326,8 +327,8 @@ std::optional EstimateFromNode( const std::vector inputs{&input}; const std::vector outputs; Node node{"packed_attention", op_type, "", inputs, outputs, &attributes, kMSDomain}; - return EstimatePackedAttentionWorkspace( - node, input_shapes, Sm80Device(), options); + return EstimatePackedAttentionWorkspaceForTest( + &node, input_shapes, Sm80Device(), options); } void SetValueInfo(ONNX_NAMESPACE::ValueInfoProto& value_info, diff --git a/tools/ci_build/test_gen_test_exports.py b/tools/ci_build/test_gen_test_exports.py index e6a38eba74094..bdf7201980671 100644 --- a/tools/ci_build/test_gen_test_exports.py +++ b/tools/ci_build/test_gen_test_exports.py @@ -121,6 +121,40 @@ def test_force_includes_preserve_function_and_data_names(self): '#pragma comment(linker, "/include:?Run@Session@@QEAAHXZ")\n', ) + def test_host_link_inputs_include_imported_and_object_libraries(self): + repo_root = Path(__file__).resolve().parents[2] + with tempfile.TemporaryDirectory(prefix="ort host inputs ") as directory: + source = Path(directory) + build = source / "build" + (source / "main.cc").write_text("int main() { return 0; }\n", encoding="utf-8") + (source / "library.cc").write_text("int required() { return 42; }\n", encoding="utf-8") + (source / "CMakeLists.txt").write_text( + f""" +cmake_minimum_required(VERSION 3.28) +project(TestHostInputs LANGUAGES CXX) +set(REPO_ROOT "{repo_root.as_posix()}") +include("${{REPO_ROOT}}/cmake/onnxruntime_test_exports.cmake") +add_library(host_runtime STATIC library.cc) +add_library(host_objects OBJECT library.cc) +add_library(external UNKNOWN IMPORTED) +set_target_properties(external PROPERTIES IMPORTED_LOCATION "${{CMAKE_CURRENT_SOURCE_DIR}}/external.lib") +add_library(test_objects OBJECT library.cc) +add_executable(host main.cc) +target_link_libraries(host PRIVATE host_runtime host_objects external) +get_target_property(host_libraries host LINK_LIBRARIES) +onnxruntime_export_test_symbols(host OBJECT_TARGET test_objects HOST_LIBS ${{host_libraries}}) +""", + encoding="utf-8", + ) + subprocess.run(["cmake", "-S", str(source), "-B", str(build)], check=True, timeout=120) + responses = list((build / "host_exports").rglob("host.rsp")) + self.assertTrue(responses) + for response in responses: + inputs = response.read_text(encoding="utf-8").splitlines() + self.assertEqual(len(inputs), 3) + for name in ("host_runtime", "host_objects.dir", "external.lib"): + self.assertTrue(any(name in entry for entry in inputs), (name, inputs)) + @unittest.skipUnless(sys.platform == "win32", "Requires the Windows MSVC toolchain") def test_windows_targeted_build_and_incremental_exports(self): repo_root = Path(__file__).resolve().parents[2] @@ -128,16 +162,46 @@ def test_windows_targeted_build_and_incremental_exports(self): source = Path(directory) build = source / "build" (source / "host.cc").write_text( + '#pragma detect_mismatch("ORT_API_MANUAL_INIT", "disabled")\n' "namespace test_runtime {\nint required(int value) { return value + 1; }\n" + "\n".join(f"int unused_{i}(int value) {{ return value + {i}; }}" for i in range(70000)) + "\n}\n", encoding="utf-8", ) - (source / "main.cc").write_text("int main() { return 0; }\n", encoding="utf-8") + (source / "main.cc").write_text( + "#include \n#include \n" + "int main(int argc, char** argv) {\n" + " if (argc != 2) return 2;\n" + ' HMODULE module = LoadLibraryW(L"test_module.dll");\n' + " if (!module) return 3;\n" + ' auto run = reinterpret_cast(GetProcAddress(module, "run_tests"));\n' + " int result = run && run() == std::atoi(argv[1]) ? 0 : 4;\n" + " FreeLibrary(module);\n" + " return result;\n}\n", + encoding="utf-8", + ) + (source / "bridge.cc").write_text( + '#pragma detect_mismatch("ORT_API_MANUAL_INIT", "enabled")\n' + 'extern "C" int CoreNodeValue(const void*);\n' + "namespace test_runtime {\n" + "struct Node { int Value() const { return CoreNodeValue(this); } };\n" + "int Estimate(const Node& node) { return node.Value(); }\n}\n" + "int EstimateForTest(const void* node) {\n" + " return test_runtime::Estimate(*static_cast(node));\n}\n", + encoding="utf-8", + ) + module_preamble = ( + '#pragma detect_mismatch("ORT_API_MANUAL_INIT", "enabled")\n' + "namespace test_runtime { class Node { public: int value; }; int required(int); int unused_0(int); }\n" + 'extern "C" int CoreNodeValue(const void* node) {\n' + " return static_cast(node)->value;\n}\n" + "int EstimateForTest(const void*);\n" + ) module = source / "module.cc" module.write_text( - "namespace test_runtime { int required(int); }\n" - 'extern "C" __declspec(dllexport) int run_tests() { return test_runtime::required(41); }\n', + module_preamble + 'extern "C" __declspec(dllexport) int run_tests() {\n' + " test_runtime::Node node{41};\n" + " return test_runtime::required(EstimateForTest(&node));\n}\n", encoding="utf-8", ) (source / "CMakeLists.txt").write_text( @@ -148,12 +212,13 @@ def test_windows_targeted_build_and_incremental_exports(self): include("${{REPO_ROOT}}/cmake/onnxruntime_test_exports.cmake") add_library(runtime STATIC host.cc) target_compile_options(runtime PRIVATE /bigobj) -add_library(test_objects OBJECT module.cc) +add_library(test_objects OBJECT module.cc bridge.cc) add_executable(provider_test_executable main.cc) set_target_properties(provider_test_executable PROPERTIES OUTPUT_NAME provider_test) target_link_libraries(provider_test_executable PRIVATE runtime) target_link_options(provider_test_executable PRIVATE /VERBOSE:LIB) -onnxruntime_export_test_symbols(provider_test_executable OBJECT_TARGET test_objects HOST_LIBS runtime) +get_target_property(host_libraries provider_test_executable LINK_LIBRARIES) +onnxruntime_export_test_symbols(provider_test_executable OBJECT_TARGET test_objects HOST_LIBS ${{host_libraries}}) add_library(test_module MODULE $) target_link_libraries(test_module PRIVATE provider_test_executable) add_custom_target(provider_test ALL DEPENDS provider_test_executable test_module) @@ -199,11 +264,12 @@ def test_windows_targeted_build_and_incremental_exports(self): ) self.assertIn("provider_test.exe", imports) self.assertIn("?required@test_runtime@@YAHH@Z", imports) + subprocess.run([str(output_dir / "provider_test.exe"), "42"], cwd=output_dir, check=True, timeout=30) module.write_text( - "namespace test_runtime { int required(int); int unused_0(int); }\n" - 'extern "C" __declspec(dllexport) int run_tests() {\n' - " return test_runtime::required(41) + test_runtime::unused_0(1);\n}\n", + module_preamble + 'extern "C" __declspec(dllexport) int run_tests() {\n' + " test_runtime::Node node{41};\n" + " return test_runtime::required(EstimateForTest(&node)) + test_runtime::unused_0(1);\n}\n", encoding="utf-8", ) subprocess.run(command, check=True, timeout=120) @@ -214,6 +280,7 @@ def test_windows_targeted_build_and_incremental_exports(self): '#pragma comment(linker, "/include:?unused_0@test_runtime@@YAHH@Z")', force_includes.read_text(encoding="utf-8"), ) + subprocess.run([str(output_dir / "provider_test.exe"), "43"], cwd=output_dir, check=True, timeout=30) if __name__ == "__main__": From d57c3d4c0fafaf835320053566628000cb81f274 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 29 Sep 2026 09:38:27 +0000 Subject: [PATCH 14/19] Guard CUDA test temporary files Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../providers/cuda/test_cases/cuda_external_data_loader_test.cc | 2 ++ 1 file changed, 2 insertions(+) diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc index cd496c6f90959..730b2dd15aba3 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc @@ -47,6 +47,7 @@ void CreateExternalDataFile(size_t length, PathString& path, FILE* file = nullptr; path = ORT_TSTR("cuda_external_data_loader_XXXXXX"); CreateTestFile(file, path); + ASSERT_NE(file, nullptr); std::vector chunk(1024 * 1024); ASSERT_EQ(kFilePrefixSize, fwrite(chunk.data(), 1, kFilePrefixSize, file)); @@ -324,6 +325,7 @@ TEST(CudaExternalDataLoaderTest, ExternalInitializerSessionMatchesWithLoaderEnab PathString model_path = ORT_TSTR("cuda_external_initializer_model_XXXXXX"); FILE* model_file = nullptr; ASSERT_NO_FATAL_FAILURE(CreateTestFile(model_file, model_path)); + ASSERT_NE(model_file, nullptr); ScopedFileDeleter model_deleter{model_path}; std::unique_ptr model_file_owner(model_file, fclose); const auto model_bytes = model.SerializeAsString(); From 7a37727ee73c5a52921414bc0d9698eee4dd4ff4 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 29 Sep 2026 09:57:24 +0000 Subject: [PATCH 15/19] Raise Android minimal binary size threshold Account for the sparse initializer validation added on main. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .github/workflows/android.yml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.github/workflows/android.yml b/.github/workflows/android.yml index b3d6b895a97fb..5142f466bfbde 100644 --- a/.github/workflows/android.yml +++ b/.github/workflows/android.yml @@ -79,8 +79,8 @@ jobs: run: | set -e -x BINARY_SIZE_THRESHOLD_ARGS="" - echo "Binary size threshold in bytes: 1596416" - BINARY_SIZE_THRESHOLD_ARGS="--threshold_size_in_bytes 1596416" + echo "Binary size threshold in bytes: 1599488" + BINARY_SIZE_THRESHOLD_ARGS="--threshold_size_in_bytes 1599488" # Ensure ANDROID_NDK_HOME is available and get its real path if [ -z "$ANDROID_NDK_HOME" ]; then From 3dfc0ce5936653157b7472f5108eca9e1c73d8bc Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Tue, 29 Sep 2026 10:09:22 +0000 Subject: [PATCH 16/19] Limit Windows ARM64 workflow token permissions Co-authored-by: xadupre <22452781+xadupre@users.noreply.github.com> --- .github/workflows/windows_arm64_python.yml | 3 +++ 1 file changed, 3 insertions(+) diff --git a/.github/workflows/windows_arm64_python.yml b/.github/workflows/windows_arm64_python.yml index f7132a7cc8ab6..a0be7cfc66327 100644 --- a/.github/workflows/windows_arm64_python.yml +++ b/.github/workflows/windows_arm64_python.yml @@ -11,6 +11,9 @@ on: - rel-* workflow_dispatch: +permissions: + contents: read + concurrency: group: ${{ github.workflow }}-${{ github.event_name == 'pull_request' && github.ref || github.sha }} cancel-in-progress: true From a27bfddd304181296b1565fe0f41ec30f61b1125 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 29 Sep 2026 11:54:04 +0000 Subject: [PATCH 17/19] Fix QMoE routing snapshot build Replace the stale routing record reference in the packed INT GEMV path and preserve expert selection collection. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc b/onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc index 23f8c69fed824..c36167c6faacc 100644 --- a/onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc +++ b/onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc @@ -1772,6 +1772,12 @@ Status QMoE::ComputeInternal(OpKernelContext* context) const { const auto fused_routing = route_tile(0, num_rows); ORT_ENFORCE(fused_routing.router_logits == nullptr, "QMoE packed INT GEMV requires materialized routing outputs."); +#if !defined(BUILD_CUDA_EP_AS_PLUGIN) && !defined(ORT_MINIMAL_BUILD) + if (routing_snapshot_) { + ORT_RETURN_IF_ERROR(routing_snapshot_->Capture( + expert_indices, SafeInt(num_rows) * SafeInt(k_), stream)); + } +#endif const bool expert_maps_built = ck::fusedBuildExpertMapsSortFirstToken( expert_indices, p_r2u, unpermuted_row_to_permuted_row, p_exp, p_efto, num_rows, num_experts, static_cast(k_), 0, num_experts, stream); @@ -1836,9 +1842,8 @@ Status QMoE::ComputeInternal(OpKernelContext* context) const { workspace_size, total_scratch_bytes); } #if !defined(BUILD_CUDA_EP_AS_PLUGIN) && !defined(ORT_MINIMAL_BUILD) - if (routing_record != nullptr) { - ORT_RETURN_IF_ERROR(routing_record->CaptureTile( - expert_indices, expert_scales, 0, SafeInt(num_rows) * SafeInt(k_), true, stream)); + if (routing_snapshot_) { + ORT_RETURN_IF_ERROR(routing_snapshot_->Consume()); } #endif return Status::OK(); From d978b23426aaa97f2d35eefaab85ee9ae469b179 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 29 Sep 2026 12:16:30 +0000 Subject: [PATCH 18/19] Address provider target review feedback Use the executable target for WebGPU definitions and cover packed INT GEMV expert counting. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_unittests.cmake | 2 +- .../framework/moe_expert_counting_test.cc | 55 +++++++++++++++---- 2 files changed, 45 insertions(+), 12 deletions(-) diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index 91921c4e192c9..b3527dd02b569 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1536,7 +1536,7 @@ block() endif() if (onnxruntime_USE_WEBGPU AND onnxruntime_USE_EP_API_ADAPTERS) - target_compile_definitions(onnxruntime_provider_test PRIVATE + target_compile_definitions(${onnxruntime_provider_test_target} PRIVATE ORT_UNIT_TEST_HAS_WEBGPU_PLUGIN_EP=1) endif() diff --git a/onnxruntime/test/framework/moe_expert_counting_test.cc b/onnxruntime/test/framework/moe_expert_counting_test.cc index 8563dc1223332..3fb9a61c6c17c 100644 --- a/onnxruntime/test/framework/moe_expert_counting_test.cc +++ b/onnxruntime/test/framework/moe_expert_counting_test.cc @@ -12,6 +12,7 @@ #include "core/graph/onnx_protobuf.h" #include "core/session/onnxruntime_session_options_config_keys.h" #include "gtest/gtest.h" +#include "test/common/cuda_op_test_utils.h" #include "test/test_environment.h" #include "test/unittest_util/framework_test_utils.h" #include "test/util/include/asserts.h" @@ -49,7 +50,8 @@ void AddZeroInitializer(GraphProto& graph, const char* name, int type, tensor->set_raw_data(std::string(size, '\0')); } -void PopulateMoeGraph(GraphProto& graph, bool quantized, bool cuda, bool subgraph, int64_t rows) { +void PopulateMoeGraph(GraphProto& graph, bool quantized, bool cuda, bool subgraph, int64_t rows, + bool packed_int_gemv = false) { graph.set_name("expert_counting"); SetValue(subgraph ? *graph.add_value_info() : *graph.add_input(), "input", TensorProto_DataType_FLOAT16, {rows, kWidth}); @@ -57,14 +59,21 @@ void PopulateMoeGraph(GraphProto& graph, bool quantized, bool cuda, bool subgrap TensorProto_DataType_FLOAT16, {rows, kExperts}); SetValue(*graph.add_output(), "output", TensorProto_DataType_FLOAT16, {rows, kWidth}); const int weight_type = quantized ? TensorProto_DataType_UINT8 : TensorProto_DataType_FLOAT16; - AddZeroInitializer(graph, "w1", weight_type, {kExperts, kWidth, quantized ? kWidth / 2 : kWidth}, + const int64_t fc1_rows = packed_int_gemv ? 2 * kWidth : kWidth; + const int64_t packed_width = packed_int_gemv ? kWidth / 4 : kWidth / 2; + AddZeroInitializer(graph, "w1", weight_type, {kExperts, fc1_rows, quantized ? packed_width : kWidth}, quantized ? 1 : 2); - AddZeroInitializer(graph, "w2", weight_type, {kExperts, kWidth, quantized ? kWidth / 2 : kWidth}, + AddZeroInitializer(graph, "w2", weight_type, {kExperts, kWidth, quantized ? packed_width : kWidth}, quantized ? 1 : 2); if (quantized) { const int scale_type = cuda ? TensorProto_DataType_FLOAT16 : TensorProto_DataType_FLOAT; - AddZeroInitializer(graph, "s1", scale_type, {kExperts, kWidth}, cuda ? 2 : 4); - AddZeroInitializer(graph, "s2", scale_type, {kExperts, kWidth}, cuda ? 2 : 4); + if (packed_int_gemv) { + AddZeroInitializer(graph, "s1", scale_type, {kExperts, fc1_rows, kWidth / 64}, cuda ? 2 : 4); + AddZeroInitializer(graph, "s2", scale_type, {kExperts, kWidth, kWidth / 64}, cuda ? 2 : 4); + } else { + AddZeroInitializer(graph, "s1", scale_type, {kExperts, kWidth}, cuda ? 2 : 4); + AddZeroInitializer(graph, "s2", scale_type, {kExperts, kWidth}, cuda ? 2 : 4); + } } for (int i = 0; i < 2; ++i) { auto* node = graph.add_node(); @@ -86,11 +95,24 @@ void PopulateMoeGraph(GraphProto& graph, bool quantized, bool cuda, bool subgrap auto* activation = node->add_attribute(); activation->set_name("activation_type"); activation->set_type(AttributeProto_AttributeType_STRING); - activation->set_s("relu"); + activation->set_s(packed_int_gemv ? "swiglu" : "relu"); + if (packed_int_gemv) { + auto add_int_attribute = [node](const char* name, int64_t value) { + auto* attribute = node->add_attribute(); + attribute->set_name(name); + attribute->set_type(AttributeProto_AttributeType_INT); + attribute->set_i(value); + }; + add_int_attribute("swiglu_fusion", 1); + add_int_attribute("expert_weight_bits", 2); + add_int_attribute("weights_prepacked", 0); + add_int_attribute("block_size", 64); + } } } -std::string MakeCountingModel(bool quantized = false, bool cuda = false, bool subgraphs = false, int64_t rows = 3) { +std::string MakeCountingModel(bool quantized = false, bool cuda = false, bool subgraphs = false, int64_t rows = 3, + bool packed_int_gemv = false) { ModelProto model; model.set_ir_version(ONNX_NAMESPACE::Version::IR_VERSION); auto* opset = model.add_opset_import(); @@ -101,7 +123,7 @@ std::string MakeCountingModel(bool quantized = false, bool cuda = false, bool su opset->set_version(1); auto& graph = *model.mutable_graph(); if (!subgraphs) { - PopulateMoeGraph(graph, quantized, cuda, false, rows); + PopulateMoeGraph(graph, quantized, cuda, false, rows, packed_int_gemv); } else { graph.set_name("conditional_counting"); SetValue(*graph.add_input(), "input", TensorProto_DataType_FLOAT16, {rows, kWidth}); @@ -116,7 +138,7 @@ std::string MakeCountingModel(bool quantized = false, bool cuda = false, bool su auto* attr = node->add_attribute(); attr->set_name(branch); attr->set_type(AttributeProto_AttributeType_GRAPH); - PopulateMoeGraph(*attr->mutable_g(), quantized, cuda, true, rows); + PopulateMoeGraph(*attr->mutable_g(), quantized, cuda, true, rows, packed_int_gemv); } } return model.SerializeAsString(); @@ -271,7 +293,8 @@ void TestDisabledRecording(bool quantized) { } } -void TestCounting(bool quantized, bool cuda, bool tiled = false, int64_t rows = 3) { +void TestCounting(bool quantized, bool cuda, bool tiled = false, int64_t rows = 3, + bool packed_int_gemv = false) { auto options = CountingOptions(); if (tiled) { ASSERT_STATUS_OK(options.config_options.AddConfigEntry("ep.cuda.qmoe_row_tile_size", "1")); @@ -279,6 +302,10 @@ void TestCounting(bool quantized, bool cuda, bool tiled = false, int64_t rows = if (cuda) { ASSERT_STATUS_OK(options.config_options.AddConfigEntry(kOrtSessionOptionsDisableCPUEPFallback, "1")); } + if (packed_int_gemv) { + ASSERT_STATUS_OK( + options.config_options.AddConfigEntry("ep.cuda.qmoe_int_dequant_max_scratch_bytes", "1")); + } InferenceSessionWrapper session(options, GetEnvironment()); if (cuda) { auto provider = DefaultCudaExecutionProvider(); @@ -290,7 +317,7 @@ void TestCounting(bool quantized, bool cuda, bool tiled = false, int64_t rows = } ASSERT_STATUS_OK(session.RegisterExecutionProvider(std::move(provider))); } - const auto model = MakeCountingModel(quantized, cuda, false, rows); + const auto model = MakeCountingModel(quantized, cuda, false, rows, packed_int_gemv); ASSERT_STATUS_OK(session.Load(model.data(), static_cast(model.size()))); ASSERT_STATUS_OK(session.Initialize()); const auto* state = session.GetSessionState().GetMoeExpertState(); @@ -406,6 +433,12 @@ TEST(MoeExpertCountingTest, CudaMoE) { TestCounting(false, true); } TEST(MoeExpertCountingTest, CudaQMoE) { TestCounting(true, true); } TEST(MoeExpertCountingTest, CudaQMoETiled) { TestCounting(true, true, true); } TEST(MoeExpertCountingTest, CudaQMoESingleToken) { TestCounting(true, true, false, 1); } +TEST(MoeExpertCountingTest, CudaQMoEPackedIntGemv) { + if (!HasCudaEnvironment(800)) { + GTEST_SKIP() << "CUDA device with compute capability 8.0 or newer is required."; + } + TestCounting(true, true, false, 3, true); +} #endif TEST(MoeExpertCountingTest, DisabledAndIndependentSessions) { From 347b0884619d043a813629beeb3db2991654a9d4 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 29 Sep 2026 13:03:01 +0000 Subject: [PATCH 19/19] Fix packed INT GEMV counting test shape Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../framework/moe_expert_counting_test.cc | 48 +++++++++++-------- 1 file changed, 28 insertions(+), 20 deletions(-) diff --git a/onnxruntime/test/framework/moe_expert_counting_test.cc b/onnxruntime/test/framework/moe_expert_counting_test.cc index 3fb9a61c6c17c..ccc49233446fd 100644 --- a/onnxruntime/test/framework/moe_expert_counting_test.cc +++ b/onnxruntime/test/framework/moe_expert_counting_test.cc @@ -24,6 +24,7 @@ namespace onnxruntime::test { namespace { using namespace ONNX_NAMESPACE; constexpr int64_t kWidth = 128; +constexpr int64_t kPackedIntGemvWidth = 512; constexpr int64_t kExperts = 4; void SetValue(ValueInfoProto& value, const std::string& name, int type, @@ -52,27 +53,28 @@ void AddZeroInitializer(GraphProto& graph, const char* name, int type, void PopulateMoeGraph(GraphProto& graph, bool quantized, bool cuda, bool subgraph, int64_t rows, bool packed_int_gemv = false) { + const int64_t width = packed_int_gemv ? kPackedIntGemvWidth : kWidth; graph.set_name("expert_counting"); SetValue(subgraph ? *graph.add_value_info() : *graph.add_input(), "input", - TensorProto_DataType_FLOAT16, {rows, kWidth}); + TensorProto_DataType_FLOAT16, {rows, width}); SetValue(subgraph ? *graph.add_value_info() : *graph.add_input(), "router", TensorProto_DataType_FLOAT16, {rows, kExperts}); - SetValue(*graph.add_output(), "output", TensorProto_DataType_FLOAT16, {rows, kWidth}); + SetValue(*graph.add_output(), "output", TensorProto_DataType_FLOAT16, {rows, width}); const int weight_type = quantized ? TensorProto_DataType_UINT8 : TensorProto_DataType_FLOAT16; - const int64_t fc1_rows = packed_int_gemv ? 2 * kWidth : kWidth; - const int64_t packed_width = packed_int_gemv ? kWidth / 4 : kWidth / 2; - AddZeroInitializer(graph, "w1", weight_type, {kExperts, fc1_rows, quantized ? packed_width : kWidth}, + const int64_t fc1_rows = packed_int_gemv ? 2 * width : width; + const int64_t packed_width = packed_int_gemv ? width / 4 : width / 2; + AddZeroInitializer(graph, "w1", weight_type, {kExperts, fc1_rows, quantized ? packed_width : width}, quantized ? 1 : 2); - AddZeroInitializer(graph, "w2", weight_type, {kExperts, kWidth, quantized ? packed_width : kWidth}, + AddZeroInitializer(graph, "w2", weight_type, {kExperts, width, quantized ? packed_width : width}, quantized ? 1 : 2); if (quantized) { const int scale_type = cuda ? TensorProto_DataType_FLOAT16 : TensorProto_DataType_FLOAT; if (packed_int_gemv) { - AddZeroInitializer(graph, "s1", scale_type, {kExperts, fc1_rows, kWidth / 64}, cuda ? 2 : 4); - AddZeroInitializer(graph, "s2", scale_type, {kExperts, kWidth, kWidth / 64}, cuda ? 2 : 4); + AddZeroInitializer(graph, "s1", scale_type, {kExperts, fc1_rows, width / 64}, cuda ? 2 : 4); + AddZeroInitializer(graph, "s2", scale_type, {kExperts, width, width / 64}, cuda ? 2 : 4); } else { - AddZeroInitializer(graph, "s1", scale_type, {kExperts, kWidth}, cuda ? 2 : 4); - AddZeroInitializer(graph, "s2", scale_type, {kExperts, kWidth}, cuda ? 2 : 4); + AddZeroInitializer(graph, "s1", scale_type, {kExperts, width}, cuda ? 2 : 4); + AddZeroInitializer(graph, "s2", scale_type, {kExperts, width}, cuda ? 2 : 4); } } for (int i = 0; i < 2; ++i) { @@ -125,11 +127,12 @@ std::string MakeCountingModel(bool quantized = false, bool cuda = false, bool su if (!subgraphs) { PopulateMoeGraph(graph, quantized, cuda, false, rows, packed_int_gemv); } else { + const int64_t width = packed_int_gemv ? kPackedIntGemvWidth : kWidth; graph.set_name("conditional_counting"); - SetValue(*graph.add_input(), "input", TensorProto_DataType_FLOAT16, {rows, kWidth}); + SetValue(*graph.add_input(), "input", TensorProto_DataType_FLOAT16, {rows, width}); SetValue(*graph.add_input(), "router", TensorProto_DataType_FLOAT16, {rows, kExperts}); SetValue(*graph.add_input(), "condition", TensorProto_DataType_BOOL, {}); - SetValue(*graph.add_output(), "output", TensorProto_DataType_FLOAT16, {rows, kWidth}); + SetValue(*graph.add_output(), "output", TensorProto_DataType_FLOAT16, {rows, width}); auto* node = graph.add_node(); node->set_op_type("If"); node->add_input("condition"); @@ -144,12 +147,14 @@ std::string MakeCountingModel(bool quantized = false, bool cuda = false, bool su return model.SerializeAsString(); } -NameMLValMap CountingFeeds(bool subgraphs = false, bool condition = true, int64_t rows = 3) { +NameMLValMap CountingFeeds(bool subgraphs = false, bool condition = true, int64_t rows = 3, + bool packed_int_gemv = false) { ORT_ENFORCE(rows == 1 || rows == 3); + const int64_t width = packed_int_gemv ? kPackedIntGemvWidth : kWidth; auto allocator = TestCPUExecutionProvider()->CreatePreferredAllocators()[0]; OrtValue input, router, cond; - const std::vector values(rows * kWidth, MLFloat16(1.0f)); - CreateMLValue(allocator, {rows, kWidth}, values, &input); + const std::vector values(rows * width, MLFloat16(1.0f)); + CreateMLValue(allocator, {rows, width}, values, &input); const std::vector routing{ MLFloat16(9.f), MLFloat16(1.f), MLFloat16(0.f), MLFloat16(0.f), MLFloat16(8.f), MLFloat16(1.f), MLFloat16(0.f), MLFloat16(0.f), @@ -165,14 +170,17 @@ NameMLValMap CountingFeeds(bool subgraphs = false, bool condition = true, int64_ } Status ExecuteCountingModel(InferenceSession& session, std::vector& outputs, - bool subgraphs = false, bool condition = true, int64_t rows = 3) { + bool subgraphs = false, bool condition = true, int64_t rows = 3, + bool packed_int_gemv = false) { const std::array output_names{"output"}; - return session.Run(RunOptions{}, CountingFeeds(subgraphs, condition, rows), output_names, &outputs); + return session.Run(RunOptions{}, CountingFeeds(subgraphs, condition, rows, packed_int_gemv), + output_names, &outputs); } -void RunCountingModel(InferenceSession& session, bool subgraphs = false, bool condition = true, int64_t rows = 3) { +void RunCountingModel(InferenceSession& session, bool subgraphs = false, bool condition = true, int64_t rows = 3, + bool packed_int_gemv = false) { std::vector outputs; - ASSERT_STATUS_OK(ExecuteCountingModel(session, outputs, subgraphs, condition, rows)); + ASSERT_STATUS_OK(ExecuteCountingModel(session, outputs, subgraphs, condition, rows, packed_int_gemv)); ASSERT_EQ(outputs.size(), 1U); for (auto value : outputs[0].Get().DataAsSpan()) { EXPECT_EQ(value.ToFloat(), 0.f); @@ -335,7 +343,7 @@ void TestCounting(bool quantized, bool cuda, bool tiled = false, int64_t rows = } EXPECT_EQ(expert_ids.size(), 2U * kExperts); for (int run = 1; run <= 2; ++run) { - RunCountingModel(session, false, true, rows); + RunCountingModel(session, false, true, rows, packed_int_gemv); for (size_t node_index : node_indices) { const double expected = run == 1 ? 0.1 : 0.19; InlinedVector counters;